c++ - OpenCL 内存带宽/合并

Question

概括:
我正在尝试编写一个内存绑定(bind) 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 个元素的 vector 。因此我应该有 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

Answer 1

您的方法存在一些问题:

您不会使 GPU 饱和。要获得最佳性能，您需要启动比 GPU 执行单元更多的线程。更多意味着> 10000000。

您的循环包含索引整数计算(用于结构数组合并访问)。在这里，这可能不足以让您进入计算限制，但通常最好使用 #pragma unroll 展开小循环。 ;然后编译器已经完成了所有的索引计算。您还可以烘焙常量 iNrOfIterations和 iNrOfWorkers使用 #define iNrOfIterations 16 直接进入 OpenCL 代码/#define iNrOfWorkers 15728640通过 C++ 字符串连接或硬编码。

根据您的访问模式，有 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;