OpenCL 内存带宽/合并

Posted 2023-02-19

【中文标题】OpenCL 内存带宽/合并

【英文标题】：OpenCL Memory Bandwidth/Coalescing

【发布时间】：2021-01-13 07:14:37

【问题描述】：

总结：

我正在尝试编写一个内存绑定 OpenCL 程序，该程序接近我 GPU 上宣传的内存带宽。实际上，我偏离了约 50 倍。

设置：

我只有一张比较旧的 Polaris Card (RX580)，所以我不能使用 CUDA，现在只能选择 OpenCL。我知道这是次优的，我无法让任何调试/性能计数器工作，但这是我所拥有的。

我是 GPU 计算的新手，想体验一下我可以期待的一些性能 从 GPU 与 CPU。对我来说首先要做的是内存带宽。

我编写了一个非常小的 OpenCL 内核，它从跨步内存位置读取，我希望波前中的所有工作人员一起在一个大内存段上执行连续内存访问，合并访问。然后内核对加载的数据所做的所有事情就是将这些值相加，并在最后将和写回另一个内存位置。代码（大部分都是我从各种来源无耻地复制在一起的）非常简单

__kernel void ThroughputTestKernel(
                     __global float* vInMemory,
                     __global float* vOutMemory,
                     const int iNrOfIterations,
                     const int iNrOfWorkers
                   )

    const int gtid = get_global_id(0);
    
    __private float fAccumulator = 0.0;
    
    for (int k = 0; k < iNrOfIterations; k++) 
        fAccumulator += vInMemory[gtid + k * iNrOfWorkers];
    
    
    vOutMemory[gtid] = fAccumulator;

我生成iNrOfWorkers 这些内核并测量它们完成处理所需的时间。对于我的测试，我设置了iNrOfWorkers = 1024 和iNrOfIterations = 64*1024。根据处理时间和iMemorySize = iNrOfWorkers * iNrOfIterations * sizeof(float)，我计算出大约 5GByte/s 的内存带宽。

期望：

我的问题是内存访问似乎比我被认为可用的 256GByte/s 慢一到两个数量级。

GCN ISA 手册 [1] 让我假设我有 36 个 CU，每个 CU 包含 4 个 SIMD 单元，每个单元处理 16 个元素的向量。因此，我应该有 36416 = 2304 个可用的处理元素。

我生成的数量少于该数量，即 1024 个全局工作单元（“线程”）。线程按顺序访问内存位置，相隔 1024 个位置，因此在循环的每次迭代中，整个波前访问 1024 个连续元素。因此，我认为 GPU 应该能够产生连续的内存地址访问，而不会中断。

我的猜测是，它只产生很少的线程，而不是 1024，可能每个 CU 一个？这样一来，它就必须一遍又一遍地重新读取数据。不过，我不知道如何验证这一点。

[1]http://developer.amd.com/wordpress/media/2013/12/AMD_GCN3_Instruction_Set_Architecture_rev1.1.pdf

【参考方案1】：

您的方法存在一些问题：

您没有使 GPU 饱和。要获得最佳性能，您需要启动比 GPU 执行单元更多的线程。更多意味着 >10000000。 您的循环包含索引整数计算（用于结构数组合并访问）。在这里，这可能不足以让您进入计算限制，但通常最好使用 #pragma unroll 展开小循环；然后编译器已经完成了所有的索引计算。您还可以通过 C++ 字符串连接或硬编码将常量 iNrOfIterations 和 iNrOfWorkers 直接烘焙到具有 #define iNrOfIterations 16 / #define iNrOfWorkers 15728640 的 OpenCL 代码中。

根据您的访问模式，有 4 种不同的内存带宽：合并/未对齐的读取/写入。 Coalesced 比未对齐快得多，并且未对齐读取的性能损失小于未对齐写入。只有合并的内存访问才能让您接近广告带宽。您测量 iNrOfIterations 合并读取和 1 个合并写入。要分别测量所有四种类型，您可以使用：

#define def_N 15728640
#define def_M 16
kernel void benchmark_1(global float* data) 
    const uint n = get_global_id(0);
    #pragma unroll
    for(uint i=0; i<def_M; i++) data[i*def_N+n] = 0.0f; // M coalesced writes

kernel void benchmark_2(global float* data) 
    const uint n = get_global_id(0);
    float x = 0.0f;
    #pragma unroll
    for(uint i=0; i<def_M; i++) x += data[i*def_N+n]; // M coalesced reads
    data[n] = x; // 1 coalesced write (to prevent compiler optimization)

kernel void benchmark_3(global float* data) 
    const uint n = get_global_id(0);
    #pragma unroll
    for(uint i=0; i<def_M; i++) data[n*def_M+i] = 0.0f; // M misaligned writes

kernel void benchmark_4(global float* data) 
    const uint n = get_global_id(0);
    float x = 0.0f;
    #pragma unroll
    for(uint i=0; i<def_M; i++) x += data[n*def_M+i]; // M misaligned reads
    data[n] = x; // 1 coalesced write (to prevent compiler optimization)

这里data 数组的大小为N*M，每个内核都在N 范围内执行。对于带宽计算，每个内核执行几百次（更好的平均值）并获得平均执行时间time1、time2、time3 和time4。然后像这样计算带宽：

合并读取带宽 (GB/s) = 4.0E-9f*M*N/(time2-time1/M) 合并写入带宽 (GB/s) = 4.0E-9f*M*N/( time1 ) 未对齐的读取带宽 (GB/s) = 4.0E-9f*M*N/(time4-time1/M) 未对齐的写入带宽 (GB/s) = 4.0E-9f*M*N/(time3 )

作为参考，here 是使用此基准测量的一些带宽值。

编辑：如何测量内核执行时间：

时钟

#include <thread>
class Clock 
private:
    typedef chrono::high_resolution_clock clock;
    chrono::time_point<clock> t;
public:
    Clock()  start(); 
    void start()  t = clock::now(); 
    double stop() const  return chrono::duration_cast<chrono::duration<double>>(clock::now()-t).count(); 
;

K 内核执行的时间测量

const int K = 128; // execute kernel 128 times and average execution time
NDRange range_local  = NDRange(256); // thread block size
NDRange range_global = NDRange(N); // N must be divisible by thread block size
Clock clock;
clock.start();
for(int k=0; k<K; k++) 
    queue.enqueueNDRangeKernel(kernel_1, NullRange, range_global, range_local);
    queue.finish();

const double time1 = clock.stop()/(double)K;

【讨论】：

非常感谢您提供这段代码！我运行了所有这些，它们的速度都在 4.5 到 5.5 GByte/s 之间！合并的在上端，未对齐的在下端，但它们与我的预期相差甚远！不过，我确实对您的带宽计算有一些疑问。为什么要相互减去各个时间，又为什么要除以 M？

注意：我在合并写入 (CW) 带宽方程中有一个错字（这里没有 /M）。因为benchmark_2 和benchmark_4 各包含 1 个 CW，所以存在读取带宽的减法。它们的执行时间是 M 合并/未对齐读取 + 1 CW 的时间。必须减去 1 个 CW 的时间才能获得孤立的合并/未对齐读取的时间。内核benchmark_1 包含M CW，因此为了获得1 CW 的时间，我将time1 除以M。 4 用于每 32 位浮点 4 字节，而 1E-9 用于转换 Byte->GigaByte。

如何测量内核执行时间？这也可能是错误的来源。我将在我的帖子中添加一个示例，说明我是如何做到的。

我刚才真的想通了！问题似乎是一个简单的驱动程序问题。我无法在 Linux 上超过 5 GByte/s，但使用 Windows 和普通 Radeon 驱动程序（以及您的一些建议）我得到大约 200 GByte/s。我会做更多的测试并更新最初的帖子。非常感谢@ProjectPhysX，您的回答极大地帮助了我理解如何从中获得更多性能！