c++ - 为什么 OpenCL 工作组大小对 GPU 性能影响巨大？

Question

我正在 Qualcomm Adreno 630 GPU 上对一个简单的矩阵转置内核进行基准测试，我试图了解不同工作组大小的影响，但令人惊讶的是，我得到了一些我无法解释的有趣结果。这是我的内核代码:

__kernel void transpose(__global float *input, __global float *output, const int width, const int height)
    int i = get_global_id(0);
    int j = get_global_id(1);
    output[i*height + j] = input[j*width + i];
}

宽度和高度均为6400，实验结果为(执行时间为END事件和START事件的差值):

work group size      execution time
x     y
4    64              24ms
64   4               169ms
256  1               654ms
1    256             34ms
8    32              27ms
1    1024            375ms
1024 1               657ms
32   32              26ms

在此之后，我做了另一个实验，将宽度和高度从 6400 更改为 6401(以及 NDRangeKernel 调用中的全局工作大小)，结果更加有趣:

work group size      execution time
x     y
4    64              28ms
64   4               105ms
256  1               359ms
1    256             31ms
8    32              32ms
1    1024            99ms
1024 1               358ms
32   32              32ms

大多数场景的执行时间显着下降。我知道内存合并或缓存可以在这里发挥作用，但我无法完全解释这一点。

Answer 1

当连续线程访问 128 字节对齐段内连续全局内存地址处的数据时，就会发生内存合并。然后内存访问被合并为一个，从而显着减少整体延迟。

在 2D 范围内，合并仅沿 get_global_id(1) 或您的情况下的 j 方向发生。在 output[i*height + j] = input[j*width + i]; 行中，input[j*width + i]; 是未对齐的(非-coalesced) 读取，output[i*height + j] 是合并写入。合并内存访问通常比未对齐访问快得多，但合并/未对齐读取的性能损失可能与合并/未对齐写入有很大不同。在大多数桌面 GPU 架构上，未对齐读取和合并写入的组合比相反的方式更快，请参见下图。因此您的实现应该已经是更快的变体了。

由于合并访问只能沿着 j 索引进行，因此如果您的范围为 (x=256,y=1) (i code> 沿 x 方向，j 沿 y 方向)，您不会得到任何合并。对于 (x=8,y=32)，j 每个线程 block 以 32 为一组合并 8 次，因此内存带宽相当饱和，性能良好。

如果您想要最大可能的性能，我建议您使用一维索引。这样您就可以完全控制合并，并且合并发生在整个线程 block 上。您的矩阵转置内核将如下所示:

#define width 6400
__kernel void transpose(__global float *input, __global float *output) {
    const int n = get_global_id(0);
    int i = n/width;
    int j = n%width;
    output[i*height + j] = input[j*width + i];
}

您可以在 C++ 运行时以及 OpenCL 编译时之前通过字符串连接将 width 烘焙到 OpenCL Ccode 中。