•      Sign In
• Create Account

### #ActualsmallGame

Posted 13 December 2012 - 10:28 AM

Hi,

I wrote a simple kernel, I read and write in global memory, but I have some troubles.
From my understanding any work item A could write in the global memory, and any work item B could read what was written by A even if they are not in the same work group. (of course if synchronization is correct)

__kernel void Test(volatile __global int* volatileGlobalMemory) // CL_MEM_READ_WRITE memory flag, size 10
{
size_t globalId = get_global_id(0);
// Init array to zero
if (globalId < 10)
{
volatileGlobalMemory[globalId] = 0;
}
barrier(CLK_GLOBAL_MEM_FENCE);
// atomic histogram
int indexBucket = globalId % 10;
atomic_inc(&volatileGlobalMemory[indexBucket]);
// Here everything is fine we have the expected hitogram inside all work items and work group
barrier(CLK_GLOBAL_MEM_FENCE);
// Copy just once; this seems to work only on the first work group, I  don't understand why ???!!!
if (globalId == 0)
{
for (int i = 1; i < 10; i++)
{
volatileGlobalMemory[i] += volatileGlobalMemory[i - 1];
}
}
barrier(CLK_GLOBAL_MEM_FENCE);
// This value is correct on the first work group, wrong on the other ones, why ?
int a0 = volatileGlobalMemory[0];
}


Thanks for your help,

### #1smallGame

Posted 13 December 2012 - 10:27 AM

Hi,

I wrote a simple kernel, I read and write in global memory, but I have some troubles.
From my understanding any work item A could write in the global memory, and any work item B could read what was writtenby A even if they are not in the same work group. (of course if synchronization is correct)

__kernel void Test(volatile __global int* volatileGlobalMemory) // CL_MEM_READ_WRITE memory flag, size 10
{
size_t globalId = get_global_id(0);
// Init array to zero
if (globalId < 10)
{
volatileGlobalMemory[globalId] = 0;
}
barrier(CLK_GLOBAL_MEM_FENCE);
// atomic histogram
int indexBucket = globalId % 10;
atomic_inc(&volatileGlobalMemory[indexBucket]);
// Here everything is fine we have the expected hitogram inside all work items and work group
barrier(CLK_GLOBAL_MEM_FENCE);
// Copy just once; this seems to work only on the first work group, I  don't understand why ???!!!
if (globalId == 0)
{
for (int i = 1; i < 10; i++)
{
volatileGlobalMemory[i] += volatileGlobalMemory[i - 1];
}
}
barrier(CLK_GLOBAL_MEM_FENCE);
// This value is correct on the first work group, wrong on the other ones, why ?
int a0 = volatileGlobalMemory[0];
}


Thanks for your help,

PARTNERS