__local atom in opencl
About atomic access to __local variables:
I know that it is slower to do global operations compared to local ones. In this sense, I would like to do atomic access to some variables.
I know I can do atomic operations in OpenCL:
// Program A:
#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable
__kernel void test(global int * num)
{
atom_inc(&num[0]);
}
How can I share atomic data between work-itens within a given workgroup?
for ex: I would like to do something like this:
// Program B: (it doesn't work, just to show how I'd like it to be)
#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable
__kernel void test(global int * num, const int numOperations)
{
__local int num;
if (get_global_id(0) < numOperations) {
atom_inc(&num);
}
}
The num value should appear at the end: numOperations - 1;
Isn't that impossible? If not, how can I do this?
+3
source to share
1 answer
Typically, you have a single thread that initializes a shared (local) atom followed by some kind of barrier. That is, your kernel starts up like this:
__local int sharedNum;
if (get_local_id (0) == 0) {
sharedNum = 0;
}
barrier (CLK_LOCAL_MEM_FENCE);
// Now, you can use sharedNum
while (is_work_left ()) {
atomic_inc (&sharedNum);
}
There's not much magic there - all items in the workgroup can see the same local variables, so you can just access it as usual.
+3
source to share