• Advertisement
Sign in to follow this  

OPenCL global memory read write understanding

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

If you intended to correct an error in the post then please contact us.

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)


[CODE]
__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];
}
[/CODE]

Thanks for your help, Edited by smallGame

Share this post


Link to post
Share on other sites
Advertisement
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 this post


Link to post
Share on other sites
Thanks lerlinghagen [img]http://public.gamedev.net//public/style_emoticons/default/smile.png[/img] !!

Share this post


Link to post
Share on other sites
Sign in to follow this  

  • Advertisement