c - 进行最终还原的策略

我正在尝试实现一个OpenCL版本来减少浮点数数组。
为了实现这一目标，我采用了以下在网络上找到的代码片段:

__kernel void sumGPU ( __global const double *input,
                       __global double *partialSums,
               __local double *localSums)
 {
  uint local_id = get_local_id(0);
  uint group_size = get_local_size(0);

  // Copy from global memory to local memory
  localSums[local_id] = input[get_global_id(0)];

  // Loop for computing localSums
  for (uint stride = group_size/2; stride>0; stride /=2)
     {
      // Waiting for each 2x2 addition into given workgroup
      barrier(CLK_LOCAL_MEM_FENCE);

      // Divide WorkGroup into 2 parts and add elements 2 by 2
      // between local_id and local_id + stride
      if (local_id < stride)
        localSums[local_id] += localSums[local_id + stride];
     }

  // Write result into partialSums[nWorkGroups]
  if (local_id == 0)
    partialSums[get_group_id(0)] = localSums[0];
 }

此内核代码运行良好，但我想通过将每个工作组的所有部分和相加来计算最终和。
当前，我通过一个简单的循环和迭代nWorkGroups由CPU执行最后求和的这一步骤。
我还看到了另一个带有原子函数的解决方案，但它似乎是为int而不是浮点数实现的。我认为只有CUDA可以为float提供原子功能。
我还看到，我可以使用另一个内核代码来执行sum的此运算，但是为了保持简单易读的源，我想避免使用此解决方案。也许我不能没有这个解决方案...
我必须告诉你，我在Radeon HD 7970 Tahiti 3GB上使用OpenCL 1.2(由clinfo返回)(我的卡不支持OpenCL 2.0)。
更一般而言，我想获得有关使用我的显卡模型和OpenCL 1.2执行最后最后求和的最简单方法的建议。

最佳答案

如果该浮点的数量级小于exa比例，则:

代替

if (local_id == 0)
  partialSums[get_group_id(0)] = localSums[0];

你可以用

if (local_id == 0)
{
    if(strategy==ATOMIC)
    {
        long integer_part=getIntegerPart(localSums[0]);
        atom_add (&totalSumIntegerPart[0] ,integer_part);
        long float_part=1000000*getFloatPart(localSums[0]);
         // 1000000 for saving meaningful 7 digits as integer
        atom_add (&totalSumFloatPart[0] ,float_part);
    }
}

这将溢出float部分，因此当您在另一个内核中将其除以1000000时，它可能具有大于1000000的值，因此您将获得其整数部分并将其添加到实整数部分:

   float value=0;
   if(strategy==ATOMIC)
   {
       float float_part=getFloatPart_(totalSumFloatPart[0]);
       float integer_part=getIntegerPart_(totalSumFloatPart[0])
       + totalSumIntegerPart[0];
       value=integer_part+float_part;
   }

仅有一些原子操作对整个内核时间无效。

其中一些get___part可以使用floor和类似的函数轻松编写。有些需要除以1M。

关于c - 进行最终还原的策略，我们在Stack Overflow上找到一个类似的问题：https://stackoverflow.com/questions/36879187/