代码之家 › 专栏 › 技术社区 › 3Dave

矢量类型的OpenCL原子加法?

atomic opencl vector

3Dave · 技术社区 · 6 年前

我正在从两个通道更新缓冲区中的一个元素,需要一个用于float4类型的原子。(更具体地说,我启动的线程是缓冲区元素的两倍,并且每个连续的线程对更新同一个元素。)

int idx = get_global_id(0);
int mapIdx = floor (idx / 2.0);

float4 toAdd;
// ...
if (idx % 2)
{
    toAdd = (float4)(0,1,0,1);
}
else
{
    toAdd = float3(1,0,1,0);
}

// avoid race condition here?
// I'd like to atomic_add(map[mapIdx],toAdd);
map[mapIdx] += toAdd;

在这个例子中, map[0] 应该递增 (1,1,1,1) (0,1,0,1) 从线程0,和 (1,0,1,0) 从线程1。

建议?我在CL文档中没有找到任何关于矢量原子的参考。我想我可以分别对每个向量分量执行此操作:

atomic_add(map[mapIdx].x, toAdd.x);
atomic_add(map[mapIdx].y, toAdd.y);
atomic_add(map[mapIdx].z, toAdd.z);
atomic_add(map[mapIdx].w, toAdd.w);

... 但那感觉是个坏主意。(由于没有浮动原子,因此需要cmpxchg破解。

建议?

2 回复 | 直到 6 年前

doqtor 6 年前

或者,您可以尝试使用这样的本地内存:

__local float4 local_map[LOCAL_SIZE/2];

if(idx < LOCAL_SIZE/2) // More optimal would be to use work items together than every second (idx%2) as they work together in a warp/wavefront anyway, otherwise that may affect the performance
    local_map[mapIdx] = toAdd;
barrier(CLK_LOCAL_MEM_FENCE);
if(idx >= LOCAL_SIZE/2)
    local_map[mapIdx - LOCAL_SIZE/2] += toAdd;
barrier(CLK_LOCAL_MEM_FENCE);

更快的是什么-原子或本地内存-或可能的(本地内存的大小可能太大)取决于实际的内核,所以您需要基准测试并选择正确的解决方案。

更新:

if(idx < LOCAL_SIZE/2)
    map[mapIdx] = local_map[mapIdx];

或者可以尝试不引入本地缓冲区,直接写入全局缓冲区:

if(idx < LOCAL_SIZE/2)
    map[mapIdx] = toAdd;
barrier(CLK_GLOBAL_MEM_FENCE); // <- notice that now we use barrier related to global memory
if(idx >= LOCAL_SIZE/2)
    map[mapIdx - LOCAL_SIZE/2] += toAdd;
barrier(CLK_GLOBAL_MEM_FENCE);

除此之外,我现在可以看到索引的问题。要使用我的答案中的代码,前面的代码应该如下所示:

if(idx < LOCAL_SIZE/2)
{
    toAdd = (float4)(0,1,0,1);
}
else
{
    toAdd = (float4)(1,0,1,0);
}

如果你需要使用 id%2 尽管如此,所有的代码都必须遵循这一点,否则您将不得不执行一些索引运算,以便将值放入 map

Boris Ivanov 6 年前

如果我正确地理解了这个问题,我会做下一步。通过使用偏移创建数组来消除ifs

float4[2] = {(1,0,1,0), (0,1,0,1)}

并使用idx%2作为偏移量

map 进入本地内存并使用 mem_fence(CLK_LOCAL_MEM_FENCE)