Memory Abstraction & Hierarchy

Performance Limits & Roofline Model 性能限制以及带宽限制模型

对于GPU内存而言,由于不同级别间GPU数据访问速度不一致,以及最大带宽有限,性能限制很大程度上来自于内存分配与内存带宽的利用.在Programming Massively Parallel Processors中也提到了,性能与最大带宽以及计算密度之间有着极大关联,如下图所示的Roofline Model

尤其是在GPU最广泛的运用——GEMM计算中,往往存在着多次内存访问,因而内存访问速度以及带宽限制对性能影响是不可忽略的,这也是GEMM的重大优化方向.

Maximum Memory Bandwidth 最大化内存带宽

对于kernal函数,为了优化对于内存的访问,我们要尽量做到以下几点来尽量跑满内存带宽:

做到完美的coalesced(内存合并,即内存连续访问)
没有divergence,即减少warp内线程的分支差异
100% occupancy,即SM的线程数跑满
没有数据依赖,确保所有计算可并行

考虑到Latency,在延迟的时候也要尽量载入数据来”隐藏”延迟,增加传输,即有如下

为了极致压榨GPU内存带宽,尽可能”隐藏”延迟,需要采用循环展开(Loop Unrolling)

Loop Unrolling 循环展开

由于GPU计算耗时较访问内存极短,所以在一次访问中,我们可以尽可能多的读取/写入数据,然后再来做多个计算,减少内存带宽”空等”(stall)的时间,将这部分时间用来加载另一个元素或者进行计算.

以下是一个展开的例子:

__global__
void kernal(const float* __restrict__ a, const float* __restrict__ b, float* __restrict__ c){
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;

    #pragma unroll 2
    for (int i = 0; i < 2; i++){
        const int idx = tid + i * stride;
        c[idx] += a[idx] * b[idx];
    }
}

以上示例实质上是在告诉我们,不要只让一个线程老老实实算一个数据.让一个线程一次拿两三个数据,通过展开循环把内存加载和计算“重叠”起来,才能不让 GPU 闲着.

Occupancy 占用率

对于占用率,即SM实际使用线程比上最大支持线程,其限制因素主要在于SM的寄存器大小、共享内存大小的限制,需要加以注意.

Claescing 合并访问

在读取内存的时候,实际上会连续读取若干元素(通常是128bytes)到缓存中,如果这些元素被利用,就不会再从内存中读取而是从更快的缓存中.可悲的是,缓存极易在后续指令中被覆盖,因此,我们需要尽量连续合并访问内存来充分利用连续读取这一缓存机制,加速数据的读取.

Vectorized Memory Access 向量化访存

我们已经知道,连续访问内存会更加快速,所以我们可以使用向量化访存,比如类型float4,如下:

__global__ void saxpy_vectorized(const float4* __restrict__ x,
                                 float4* __restrict__ y, float a, int n) {
    const int i = blockIdx.x * blockDim.x + threadIdx.x;
    const int n4 = n / 4;
    
    if (i < n4) {
        const float4 xv = x[i];
        const float4 yv = y[i];
        
        y[i] = make_float4(a * xv.x + yv.x, a * xv.y + yv.y,
                           a * xv.z + yv.z, a * xv.w + yv.w);
    }
}