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 += 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,
!!
