提高非常简单的 OpenCL 内核的速度答案

【问题标题】：Improve speed of very simple OpenCL kernel提高非常简单的 OpenCL 内核的速度
【发布时间】：2016-05-22 08:35:50
【问题描述】：

明天我有一个关于 OpenCL 的测试。我们有一些示例测试，但没有解决方案。给出的是这段代码：

void scalar_add(int n, double *a, double *b, double *result) {
    for (int i=0; i<n; i++)
        result[i] = a[i] + b[i];
}

第一个任务是编写一个 OpenCL 内核。所以我的解决方案：

__kernel void scalar_add(
                          __global double *a,
                          __global double *b,
                          __global double *result
) {
       size_t i = get_global_id(0;)
       result[i] = a[i] + b[i];
    }

对于每个元素，我从 A 读取一次，从 B 读取一次，向 C 写入一次。我不知道如何使用私有或本地内存来提高速度。下一个问题是关于如何通过一个简单的更改来提高速度（“Welche kleine Änderung könnte auf einer Standard-Grafikkarte zu einer deutlichen Leistungssteigerung führen?”）。 有没有办法提高速度？

内核只从 A 和 B 读取，所以也许可以使用它。我尝试对参数 A 和 B 使用“__local”，但这不会编译或运行。

【问题讨论】：

我唯一能想到的就是将计算批处理在一起，因为内核只是一个计算。因此，在一个内核中计算多个元素以提高一致性并减少调度内核时的开销。我不知道这是否算作“一个简单的改变”..
在许多 GPU 中，使用向量类型无济于事，但实际上会生成不能很好地合并的内存访问模式。同时这里的算术强度是如此之低，以至于内存带宽无论如何都会成为限制。这里最好的性能提示实际上是：不要使用 GPU，使用 CPU 上的 SIMD 指令进行标量加法。也许答案是用浮点数替换双精度？
@JanLucas wat
@Hubert：最大化传输中的数据量通常是个好主意，但在这里它不太可能有帮助。该内核将使用如此少的寄存器，无论如何您都将获得完美的占用，并且每个线程已经在运行中产生了 16（读取）+8（写入）字节。我们可以使用little's law来计算我们需要多少个飞行字节（例如K20）：208 GB/s * 300 Cycles / (706 MCycle/s) = 88 kB in flight。 K20 上的完全占用是 13*2048 个线程，因此每个线程需要 3.3 个字节的运行时间 -> 我们使用所有内存带宽，因为每个线程有更多的 8 个字节在运行
使用浮点数会减少两倍所需的带宽，同时，浮点数的算术吞吐量远高于几乎所有 GPU 的两倍（至少两倍，但通常是 8x-32x）跨度>

标签： performance opencl

【解决方案1】：

下面的代码应该会为您加快速度。您需要使用单个工作组调用它。尝试 32 的倍数的组大小 - 64 通常是好的。反复试验是为您的硬件找到组大小最佳位置的一种很好的方法。我假设您为此内核使用 GPU。

我还添加了 maxSize 作为循环计数器的上限，它应该等于向量的长度。您可以将其作为参数传递，或者根据需要将其编码为常量。

__kernel void scalar_add(
                          __global double *a,
                          __global double *b,
                          __global double *result
) 
{
    size_t gid = get_global_id(0);
    size_t groupSize = get_global_size(0);
    for(int i = gid; i< maxSize; i+= groupSize){
        result[i] = a[i] + b[i];
    }
}

不同的是内存合并。当每个工作项读取其下一个元素时，这些读取由硬件在低级别组合。每次读取应该能够获得 32 个字节或更多字节，因此 1 的价格有 4 个 double 值。

【讨论】：

您的意思是 get_local_size(0) 吗？因为使用 get_global_size 我们可以访问超出范围的索引。
全局正确。这仅适用于一个工作组调用，因此两者将具有相同的值。

【解决方案2】：

这里一个值得注意的延迟是内核开销。仅针对“+”的一个内核执行是过大的（但仍然更快），除非它已经在 gpu 中（没有从 cpu 馈送到 gpu 并且有很大的优化空间）。

如果是gpu，那么你可以简单地将这个加法操作移到另一个内核，并获得大约几百微秒到几毫秒的时间。当然，当其他内核项完全独立于其邻居的结果时，这是适用的，因为整个项数组没有同步（只完成计算单元范围的同步，而不是所有项）

__kernel void scalar_add(
                          __global double *a,
                          __global double *b,
                          __global double *otherParameters

) 
{
    size_t i = get_global_id(0);
    double result[i] = a[i] + b[i];
    // other computations that are independent from neighbours' result 
    // or a,b are not dependant from other things of neighbour items.
    // or
    otherParameters[i]=sin(result[i]);  // ok for example
    otherParameters[i]=cos(result[i+1]) // not ok 
}

如果连续的内核似乎是依赖的，尝试将其中一个分解为两个更轻的组件，其中一个必须独立于这个添加，另一部分必须独立于第三个内核，合并两个组件应该给出两个而不是三个内核执行。这种分解可能如下所示：

  // 3 unbalanced kernels
  c=a+b                        // old 1st kernel

  d[i]=sin(c[i+1])+cos(c[i-1]) // old 2nd kernel but dependent to neighbours

  e[i]=raytrace(d[i]);         // old 3rd kernel that is time consuming

  // 2 more balanced kernels     
  c=a+b;                         // new kernel-1
  d0[i]=sin(c[i]);               // new kernel-1 (decomposed 2nd old-kernel part 1)


  d1[i]=cos(c[i]);               // new kernel-2 (decomposed 2nd old-kernel part 2)
  e[i+1]=raytrace(d1[i]+d2[i+2]);// new kernel-2

如果你能做到这一点，你就可以摆脱对 *result 数组的不必要的全局内存访问。

【讨论】：