最终减少的策略

Question

我正在尝试实现一个 OpenCL 版本来减少浮点数组。

为了实现它，我拿了以下在网上找到的sn-p代码：

__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 执行最后最终求和的最简单方法的建议。

Answer 1

如果该浮点数的数量级小于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);
    }
}

这会溢出浮点部分，所以当你在另一个内核中将它除以 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。

Answer 2

抱歉之前的代码。 它也有问题。

CLK_GLOBAL_MEM_FENCE 仅影响当前工作组。 我很困惑。 =[

如果你想通过 GPU 减少总和，你应该在 clFinish(commandQueue) 之后通过 NDRangeKernel 函数将减少内核入队。

请记住概念。

__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];

    barrier(CLK_GLOBAL_MEM_FENCE);

      if(get_group_id(0)==0){
          if(local_id < get_num_groups(0)){  // 16384
            for(int n=0 ; n<get_num_groups(0) ; n+= group_size )
               localSums[local_id] += partialSums[local_id+n];
            barrier(CLK_LOCAL_MEM_FENCE);

            for(int s=group_size/2;s>0;s/=2){
               if(local_id < s)
                  localSums[local_id] += localSums[local_id+s];
               barrier(CLK_LOCAL_MEM_FENCE);
            }
            if(local_id == 0)
               partialSums[0] = localSums[0];
          }
       }
 }