时间戳

首页 > 解决方案 > 为什么 OpenCL 工作组大小对 GPU 的性能影响巨大？

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

问题描述

我在 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

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

标签： c++cachingopenclgpgpumemory-access

解决方案

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

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

各种设备的合并/未对齐内存带宽

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

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

#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];
}

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

推荐阅读