双重减少opencl教程

Question

我是一个OpenCl的新手。

我需要在一维双精度数组上运算一个约化（和运算符）。

我一直在网上游荡，但我发现的例子很混乱。 任何人都可以发布一个容易阅读（也可能是高效的）教程实施？

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

Answer 1

您提到您的问题涉及60k双打，这将不适合您的设备的本地内存。我把一个内核放在一起，将你的矢量减少到10-30个左右的值，你可以将它与你的主机程序相加。我在我的机器上遇到双打问题，但是如果您启用双打并将“浮动”更改为“双倍”，则此内核应该可以正常工作。我将调试我遇到的双重问题，并发布更新。

PARAMS：

• 全球浮动* inVector - 花车源总结
• 全球浮动* outVector - 花车的列表，每个工作组
• const int的inVectorSize - 总inVector持有的花车数量
• 本地浮动* 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卡上。