Sign in to follow this  
smallGame

OPenCL global memory read write understanding

Recommended Posts

smallGame    211
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
lerlinghagen    257
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

Create an account or sign in to comment

You need to be a member in order to leave a comment

Create an account

Sign up for a new account in our community. It's easy!

Register a new account

Sign in

Already have an account? Sign in here.

Sign In Now

Sign in to follow this