OpenCL - 双重原子操作 - 工作到极限

Question

按照这个link，我尝试实现一个计算双精度数组总和的原子函数，因此我实现了自己的atom_add 函数（用于双精度）。

这是使用的内核代码：

#pragma OPENCL EXTENSION cl_khr_fp64: enable
#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable

void atom_add_double(__global double *val, double delta)
{
  union {
  double f;
  ulong  i;
  } old, new;

  do
  {
   old.f = *val;
   new.f = old.f + delta;
  } 
  while (atom_cmpxchg((volatile __global ulong *)val, old.i, new.i) != old.i);

}  

__kernel void sumGPU ( __global const double *input, 
               __global double *finalSum
                 )
{
  // Index of current workItem
  uint gid = get_global_id(0);

  // Init sum
  *finalSum = 0.0;

  // Compute final sum 
  atom_add_double(finalSum, input[gid]);

}                   

我的问题是内核代码会产生良好的结果，直到我达到 input 数组大小的大约 100000 个元素。

超过这个限制，计算不再有效（我可以很容易地检查结果，因为在我的测试用例中，我用循环填充输入数组for(i=0;i<sizeArray;i++) input[i]=i+1;，所以总和等于sizeArray*(sizeArray+1)/2） .

我可以定义像atom_add_double这样的函数并将其放入内核代码中吗？

Answer 1

@huseyin 的回答是正确的，可以解决问题。

但是，我忍不住要说“不要使用原子来减少。”

甚至更糟糕的原子会锁定在一个while循环中并直接访问全局数据。我们可能正在谈论至少 10 倍的性能损失。

如果可以，请使用proper automatic reduction (CL 2.0+)。

__kernel void sumGPU(__global const double *input, __global double *finalSum)
{
  // Index of current workItem
  uint gid = get_global_id(0);

  // Sum locally without atomics
  double sum = work_group_scan_inclusive_add(input[gid]);

  // Compute final sum using atomics
  // but it is even better if just store them in an array and do final sum in CPU
   // Only add the last one, since it contains the total sum
  if (get_local_id(0) == get_local_size(0) - 1) { 
    atom_add_double(finalSum, sum);
  }
} 

Answer 2

*finalSum = 0.0;

是所有进行中线程的竞争条件。它使我的计算机的结果为零。删除它，从主机端初始化它。如果您的 gpu 非常好，则运行中线程的数量可能高达 50000 甚至更多，并且每个线程在任何开始原子函数之前都达到 finalSum = 0.0 但是当您超过该限制时，第 50001 个（只是一个微不足道的数字）线程将其重新初始化为零。

那么，所有元素的总和不等于 size*(size+1)/2 因为它是从零开始的（第零个元素为零）所以它实际上是

(size-1)*(size)/2

当我从内核中删除 finalSum =0.0 时，我的计算机给出了正确的结果。