OpenCL内存带宽/合并

问题描述

摘要：

我正在尝试编写一个内存绑定的OpenCL程序，该程序接近GPU上公布的内存带宽。实际上，我差了50倍。

设置：

我只有一个相对较旧的Polaris卡（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单元，每个SIMD单元的处理向量为16个元素。因此，我应该有36个 4 16 = 2304个处理元素可用。

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

我的猜测是，它生成的线程数很少，而不是1024，每个CU可能只有一个？这样，它将不得不一次又一次地重新读取数据。不过，我不知道该如何验证。

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

解决方法

您的方法存在一些问题：

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

根据您的访问模式，有4种不同的内存带宽：合并/未对齐的读/写。合并比未对齐的速度快得多，并且未对齐的读取的性能损失小于未对齐的写入。只有合并的内存访问才能使您到达广告带宽附近的任何位置。您测量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是使用此基准测试测得的一些带宽值。

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

1. 时钟

#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(); }
};

1. 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;