opencl 内核线程中的竞争条件

Question

如果多个线程同时写入一个内存位置，就会出现竞争条件，对吧？？ 在我的情况下，同样的事情正在发生..

考虑来自 'reduce.cl' 的模块

int i = get_global_id(0);
int n,j;

n = keyMobj[i];                       // this n is the key..It can be either 0 or 1.
for(j=0; j<2; j++)
      sumMobj[n*2+j] += dataMobj[i].dattr[j];        //summing operation.

这里，内存位置
sumMobj===> [...0..., ....1...] 被 4 个线程同时访问 & sumMobj===> [....3..., ....4...] 被6个线程同时访问..

有什么方法可以让它保持并行，比如使用锁定或信号量？因为这个求和是我算法中非常重要的一部分......

Answer 1

我可以给你一些提示，因为我也面临类似的问题。

我可以想到三种不同的方法来实现类似的目标：

考虑一个简单的内核，假设您启动了 4 (0-3) 个线程

_kernel void addition (int *p)
{
int i = get_local_id(0);
     p[4]+= p[i];
}

您希望将值 p[0]、p[1]、p[2]、p[3]、p[4] 相加，并将最终总和存储在 p[4] 中。正确的？即：

p[4]= p[0] + p[1] + p[2] + p[3] + p[4] 

方法-1（无并行）

将此作业仅分配给 1 个线程（无并行性）：

int i = get_local_id(0);
if (i==0)

{

p[4]+= p[i];

} 

方法 2（具有并行性）

表达你的问题如下：

p[4]= p[0] + p[1] + p[2] + p[3] + p[4] + 0  

这是一个归约问题

因此启动 3 个线程：i=0 到 i=2。在第一次迭代中

 i=0 finds p[0] + p[1]
 i=1 finds p[2] + p[3]  
 i=2 finds p[4] + 0

现在你有了三个数字，你应用与上面相同的逻辑并将这些数字相加（适当的填充 0 使其成为 2 的幂）

方法-3原子操作

如果还需要原子实现，可以使用atomic_add()：

  int fsfunc atomic_add (   volatile __global int *p ,int val)

说明

读取存储在指向位置的 32 位值（称为旧值） 由 p。计算 (old + val) 并将结果存储在 p 指向的位置。 该函数返回旧的。

这是假设数据是 int 类型。否则，您可以看到上面建议的link。