3

我是 OpenCl 的新手。

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

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

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

4

1 回答 1

6

您提到您的问题涉及 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。有关传入本地 (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 卡上。

于 2012-04-16T14:01:18.997 回答