# OPenCL global memory read write understanding

This topic is 1865 days old which is more than the 365 day threshold we allow for new replies. Please post a new topic.

## Recommended Posts

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 += 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, Edited by smallGame

##### Share on other sites
There is no intra-workgroup synchronization within a kernel. The barrier
instruction only performs synchronization within a given workgroup. The OpenCL specification does not guarantee consistency between different workgroups - see section 3.3.1 of the specification.

##### Share on other sites
Thanks lerlinghagen !!