opencl - __local 原子在 opencl

Question

关于 __local 变量的原子访问:

我知道与本地操作相比，进行全局操作的速度很慢。从这个意义上说，我想对某些变量进行原子访问。

我知道我可以在 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]);
}

如何在给定工作组内的工作项之间共享原子数据？

例如:我想做这样的事情:

// 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);
    }
}

最后 num 值应该返回: numOperations - 1;

这不可能吗？如果不是，我怎么能做到呢？

Answer 1

通常，您有一个线程初始化共享(本地)原子，然后是一些屏障。 IE。你的内核是这样启动的:

__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);
}

它没有什么神奇之处——工作组中的所有项目都可以看到相同的局部变量，因此您可以像往常一样访问它。