双倍和减法OpenCL教程 [英] double sum reduction opencl tutorial

问题描述

我完全是OpenCL的新手。

我需要对一维双精度数组进行减法运算(求和运算符)。

我一直在网上转来转去，但我找到的例子很让人困惑。 任何人都可以发布易于阅读(并且可能高效)的教程实现吗？

其他信息： -我可以访问一台GPU设备； -我使用C编写内核代码

推荐答案

您提到您的问题涉及60k Double，无法放入设备的本地内存。我组装了一个内核，它可以将您的向量降低到10-30左右的值，您可以将其与主机程序相加。我在我的机器上使用Double遇到了问题，但是如果您启用了Double并在找到的地方将"Float"更改为"Double"，那么这个内核应该工作得很好。我将调试我遇到的双重问题，然后发布更新。

参数：

• 全局浮点*inVector-要求和的浮点源
• 全局浮点*outVector-浮点列表，每个工作组一个浮点
• const int inVectorSize-inVector持有的浮点数总数
• LOCAL FLOAT*ResultScratch-每个工作组要使用的本地内存。您需要为组中的每个工作项分配一个浮点。预期大小=sizeof(Cl_Floate)*get_local_size(0)。例如，如果每个组使用64个工作项，这将是64个浮点数=256个字节。切换到双精度将使其为512字节。OpenCL规范定义的最小LDS大小为16KB。See this question有关传入本地(空)参数的详细信息。

用法：

1. 为输入和输出缓冲区分配内存。
2. 为设备上的每个计算单元创建工作组。
3. 确定最佳工作组大小，并使用此大小计算"result tScratch"的大小。
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)，显示此内核的双精度版本实际上是5870和6970卡上绑定的ALU。

这篇关于双倍和减法OpenCL教程的文章就介绍到这了，希望我们推荐的答案对大家有所帮助，也希望大家多多支持IT屋！