双倍减法opencl教程

Question

我是 OpenCl 的新手。

我需要对一维双精度数组进行归约（求和运算符）。

我一直在网上闲逛，但我发现的例子很混乱。 任何人都可以发布一个易于阅读（并且可能高效）的教程实现吗？

附加信息： - 我可以使用一台 GPU 设备； - 我使用 C 作为内核代码

Answer 1

您提到您的问题涉及 60k 双打，这不适合您设备的本地内存。我组装了一个内核，它将您的向量减少到 10-30 左右的值，您可以将其与您的主机程序相加。我在我的机器上遇到了双打问题，但是如果你启用双打并将'float'更改为'double'，这个内核应该可以正常工作。我将调试我遇到的双重问题，并发布更新。

参数：

• global float* inVector - 求和的浮点数来源
• global float* outVector - 一个浮动列表，每个工作组一个
• const int inVectorSize - inVector 持有的浮点数总数
• local float* resultScratch - 每个工作组使用的本地内存。您需要为组中的每个工作项分配一个浮点数。预期大小 = sizeof(cl_float)*get_local_size(0)。例如，如果每个组使用 64 个工作项，则这将是 64 个浮点数 = 256 个字节。切换到双精度将使其变为 512 字节。 openCL 规范定义的最小 LDS 大小为 16kb。 See this question 了解有关传入本地 (NULL) 参数的更多信息。

用法：

1. 为输入和输出缓冲区分配内存。
2. 为设备上的每个计算单元创建一个工作组。
3. 确定最佳工作组大小，并使用它来计算“resultScratch”的大小。
4. 调用内核，读出向量回主机
5. 循环遍历您的 outVector 副本并相加得到最终总和。

潜在的优化：

1. 像往常一样，您想调用带有大量数据的内核。太少的数据不值得传输和设置时间。

2. 使 inVectorSize（和向量）成为（工作组大小）*（工作组数）的最高倍数。仅使用此数据量调用内核。内核均匀地分割数据。在等待回调时计算主机上任何剩余数据的总和（或者，为 cpu 设备构建相同的内核并仅将剩余数据传递给它）。在上面的步骤 #5 中添加 outVector 时，从这个总和开始。这种优化应该使工作组在整个计算过程中保持均匀饱和。

   __kernel void floatSum(__global float* inVector, __global float* outVector, const int inVectorSize, __local float* resultScratch){
       int gid = get_global_id(0);
       int wid = get_local_id(0);
       int wsize = get_local_size(0);
       int grid = get_group_id(0);
       int grcount = get_num_groups(0);
   
       int i;
       int workAmount = inVectorSize/grcount;
       int startOffest = workAmount * grid + wid;
       int maxOffest = workAmount * (grid + 1);
       if(maxOffset > inVectorSize){
           maxOffset = inVectorSize;
       }
       resultScratch[wid] = 0.0;
       for(i=startOffest;i<maxOffest;i+=wsize){
               resultScratch[wid] += inVector[i];
       }
       barrier(CLK_LOCAL_MEM_FENCE);
   
       if(gid == 0){
               for(i=1;i<wsize;i++){
                       resultScratch[0] += resultScratch[i];
               }
               outVector[grid] = resultScratch[0];
       }
   

   }

另外，启用双打：

#ifdef cl_khr_fp64
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
#else
#ifdef cl_amd_fp64
#pragma OPENCL EXTENSION cl_amd_fp64 : enable
#endif
#endif

更新：AMD APP KernelAnalyzer 获得了更新 (v12)，显示此内核的双精度版本实际上是 ALU 绑定在 5870 和 6970 卡上。