OpenCL - 使用原子减少双重

Question

我知道不推荐使用 OpenCL-1.x 的原子函数，但我只是想了解一个原子示例。

以下内核代码运行不正常，它会生成随机最终值，用于计算所有数组值的总和（求和）：

#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable

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

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

}  

__kernel void sumGPU ( __global const double *input, 
               __local double *localInput,
               __global double *finalSum
                 )
{

  uint lid = get_local_id(0); 
  uint gid = get_global_id(0);
  uint localSize = get_local_size(0);
  uint groupid = get_group_id(0);
  local double partialSum;
  local double finalSumTemp;

 // Initialize sums
  if (lid==0)
  {
   partialSum = 0.0;
   finalSumTemp = 0.0;
  }
  barrier(CLK_LOCAL_MEM_FENCE);

  // Set in local memory
  int idx = groupid * localSize + lid;
  localInput[lid] = input[idx];

  // Compute atom_add into each workGroup
  barrier(CLK_LOCAL_MEM_FENCE);
  atom_add_double(&partialSum, localInput[lid]);
  // See and Check if barrier below is necessary
  barrier(CLK_LOCAL_MEM_FENCE);

  // Final sum of partialSums
  if (lid==0)
  {
   atom_add_double(&finalSumTemp, partialSum);
   *finalSum = finalSumTemp;
  }

}                   

带有global id 策略的版本效果很好，但上面的版本通过使用local memory（共享内存），并没有给出预期的结果（*finalSum 的值对于每次执行都是随机的)。

这里是我在主机代码中放入的缓冲区和内核参数：

 // Write to buffers
  ret = clEnqueueWriteBuffer(command_queue, inputBuffer, CL_TRUE, 0,
        nWorkItems * sizeof(double), xInput, 0, NULL, NULL);
  ret = clEnqueueWriteBuffer(command_queue, finalSumBuffer, CL_TRUE, 0,
                      sizeof(double), finalSumGPU, 0, NULL, NULL);

 // Set the arguments of the kernel
  clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&inputBuffer);
  clSetKernelArg(kernel, 1, local_item_size*sizeof(double), NULL);
  clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&finalSumBuffer);

最后，我阅读了finalSumBuffer 以获取总和值。

我认为我的问题来自内核代码，但我找不到错误在哪里。

如果有人能看出问题所在，很高兴告诉我。

谢谢

更新 1：

我几乎设法完成了这种减少。按照 huseyin tugrul buyukisik 的建议，我修改了内核代码如下：

#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable

void atom_add_double(volatile __local double *val, double delta)
{
  union {
  double d;
  ulong  i;
  } old, new;

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

}  

__kernel void sumGPU ( __global const double *input, 
               __local double *localInput,
               __local double *partialSum,
               __global double *finalSum
                 )
{

  uint lid = get_local_id(0); 
  uint gid = get_global_id(0);
  uint localSize = get_local_size(0);
  uint groupid = get_group_id(0);

  // Initialize partial sums
  if (lid==0)
    partialSum[groupid] = 0.0; 


  barrier(CLK_LOCAL_MEM_FENCE);
  // Set in local memory
  int idx = groupid * localSize + lid;
  localInput[lid] = input[idx];

  // Compute atom_add into each workGroup
  barrier(CLK_LOCAL_MEM_FENCE);
  atom_add_double(&partialSum[groupid], localInput[lid]);
  // See and Check if barrier below is necessary
  barrier(CLK_LOCAL_MEM_FENCE);

  // Compute final sum
  if (lid==0)
    *finalSum += partialSum[groupid]; 

}                   

正如 huseyin 所说，我不需要对所有部分和的最终总和使用原子函数。

所以我最后做了：

// Compute final sum
  if (lid==0)
    *finalSum += partialSum[groupid]; 

但不幸的是，最终总和并没有给出预期的值，并且该值是随机的（例如，使用nwork-items = 1024 和size-WorkGroup = 16，我得到的随机值是[1e+3 - 1e+4] 而不是5.248e+05预计。

以下是主机代码中的参数设置：

 // Set the arguments of the kernel
  clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&inputBuffer);
  clSetKernelArg(kernel, 1, local_item_size*sizeof(double), NULL);
  clSetKernelArg(kernel, 2, nWorkGroups*sizeof(double), NULL);
  clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *)&finalSumBuffer);

你能看到我在内核代码中的错误在哪里吗？

谢谢

Answer 1

不是错误而是逻辑问题：

atom_add_double(&finalSumTemp, partialSum);

每个组只工作一次（由零本地索引线程）。

所以你只是在做

finalSumTemp = partialSum

所以这里不需要原子。

---

有竞争条件

*finalSum = finalSumTemp;

在每个零索引本地线程写入相同地址的工作组之间。 所以这应该是原子加法（用于学习目的），或者可以写在不同的单元格上以添加到主机端，例如 sum_group1+sum_group2+... = 总和。

---

int idx = groupid * localSize + lid;
localInput[lid] = input[idx];

这里使用 groupid 对多设备求和是可疑的。因为每个设备都有自己的全局范围和工作组 id 索引，所以两个设备可以为两个不同的组具有相同的组 id 值。当使用多个设备时，应该使用一些设备相关的偏移量。如：

idx= get_global_id(0) + deviceOffset[deviceId];

---

此外，如果原子操作是不可避免的，并且如果恰好操作了 N 次，则可以将其移动到单个线程（例如 0 索引线程）并在第二个内核中循环 N 次（可能更快），除非该原子操作操作延迟无法通过其他方式隐藏。