双减少opencl教程

Question

我是OpenCl的新手.

我需要在一维双精度数组上运算减少(求和运算符).

我一直在网上徘徊,但我发现的例子很混乱.任何人都可以发布易于阅读(并且可能有效)的教程实现吗？

附加信息: - 我可以访问一个GPU设备; - 我使用C作为内核代码

Answer 1

你提到你的问题涉及60k双打,它不适合你设备的本地内存.我整理了一个内核,它会将你的向量减少到10-30左右,你可以将它与你的宿主程序相加.我在我的机器上遇到双打问题,但如果启用双打并且在找到它时将'float'更改为'double',则此内核应该可以正常工作.我将调试我遇到的双重问题,并发布更新.

PARAMS:

• global float*inVector - 求和的浮点数的来源
• global float*outVector - 一个浮点列表,每个工作组一个浮点数
• const int inVectorSize - inVector持有的浮点数总数
• local float*resultScratch - 要使用的每个工作组的本地内存.您需要为组中的每个工作项分配一个浮点数.expected size = sizeof(cl_float)*get_local_size(0).例如,如果每组使用64个工作项,则这将是64个浮点数= 256个字节.切换到双精度将使其成为512字节.openCL规范定义的最小LDS大小为16kb.有关传入本地(NULL)参数的详细信息,请参阅此问题.

用法:

1. 为输入和输出缓冲区分配内存.
2. 为设备上的每个计算单元创建一个工作组.
3. 确定最佳工作组大小,并使用它来计算'resultScratch'的大小.
4. 调用内核,将outVector读回主机
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卡上.