进行最终还原的策略

Strategy for doing final reduction

我正在尝试实现一个 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];
 }                  

此内核代码运行良好,但我想通过将每个工作组的所有部分和相加来计算最终总和。 目前,我通过 CPU 使用简单的循环和迭代 nWorkGroups.

来完成最终总和的这一步

我还看到了另一个具有原子函数的解决方案,但它似乎是为 int 实现的,而不是为浮点数实现的。我认为只有CUDA提供了float的原子函数。

我还看到我可以使用另一个内核代码来执行此求和运算,但我想避免使用此解决方案以保持简单可读的源代码。也许我离不开这个解决方案...

我必须告诉你,我在 Radeon HD 7970 Tahiti 3GB 上使用 OpenCL 1.2(由 clinfo 返回)(我认为我的显卡不支持 OpenCL 2.0)。

更一般地说,我想获得有关使用我的显卡型号和 OpenCL 1.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];
          }
       }
 }

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