How_to_Access_Global_Memory_Efficiently

How to Access Global Memory Efficiently in CUDA C/C++ Kernels

2020 Jan 28th CookieLau

Source: https://devblogs.nvidia.com/how-access-global-memory-efficiently-cuda-c-kernels/

[toc]

testCode

分别用一下代码测试 offset 访问和 stride 访问的间隔对带宽的影响：

__global__ void offset(T* a, int s)
{
  int i = blockDim.x * blockIdx.x + threadIdx.x + s;
  a[i] = a[i] + 1;
}

__global__ void stride(T* a, int s)
{
  int i = (blockDim.x * blockIdx.x + threadIdx.x) * s;
  a[i] = a[i] + 1;
}

Misaligned Data Access

C870: Compute Capability 1.0 C1060: Compute Capability 1.3 C2050: Compute Capability 2.0

在 device 中分配的数组都被 cuda Driver 按照 256字节 对齐，当访存 global Memory 的时候可以通过 32字节、64字节 或 128字节等分块进行数据交换。

对于C870这种 compute capability 在 1.0 及以下的 GPU，其 warp size 只有16，而且当发生 misaligned access 的时候，会对每个 misaligned 的 data 单独进行存取，所以会从原来的 16 thread 变成 16次的 32-bytes 访存。对于存取 float 数据而言，每次取 32-bytes 的数据中只有 4-bytes 是有效的，所以带宽减少到原来(offset=0) 的 1/8.

对于C1060这种 compute capability 较好的，对带宽的影响没有 C870 那么严重，只要 misaligned 落在访存的 segment，如 32,64,128-bytes 只会降低部分的性能。

对于C2050这种 compute capability 大于等于 2.0 的设备而言，其每个 multiprocessor 都配有一个 128-bytes 的 L1 cache，所以offset的改变几乎不会对带宽产生影响。

Stride Memory Access

对于C870完全不能处理misaligned，只能处理linear+aligned的架构来说，除了stride=1的情况，都出现了 7/8 的落差，变为正常情况下的 1/8。

对于CC(Compute Capability>1.0) 的来说，都能处理部分 misaligned，所以带宽曲线是 smoothly 下降，但对于相隔很远的访问也无能为力。

但是我们又时常需要进行 stride 访问，比如 grid-stride，则可以通过 shared memory 来解决这一问题。 Shared memory 是 on-chip 的、被一个 thread block 中的所有 threads 共享的一部分内存。

举一个例子：将 2D 的数组的数据加载到 shared memory 里面进行访存降低对带宽的伤害。在共享内存中是没有 stride access penalty 的。