CUDA学习笔记2——全局内存的合理使用（以矩阵转置为例，讲述全局内存的合并访问以及__ldg()）

6. 全局内存的合理使用

6.1 全局内存的合并与非合并访问

对全局内存的访问会触发内存事务（memory transaction），也就是数据传输（data transfer）。在启用了L1缓存的情况下，对全局内存的读取首先尝试经过L1缓存；如果未命中，则尝试通过L2缓存；如果还未命中，则直接从DRAM读取。一次数据传输的数据量在默认情况下是32字节。

关于全局内存的访问模式，有合并（coalesced）与非合并（uncoalesced）两种。合并访问指的是一个warp对全局内存的一次访问请求导致最少数量的数据传输。定量的说，可以定义一个合并度（degree of coalescing），它等于warp请求的字节数除以该请求导致的所有数据传输处理的字节数。如果所有数据传输处理的数据都是该warp所需要的，那么合并度为100%，即合并访问。

举个栗子

考虑一个warp访问float类型的全局内存变量的情况，一个float占4字节，因此该warp请求128字节的数据。在理想情况下（合并度为100%时），这将仅触发128/32=4次数据传输。那么，什么情况下会触发多于4次的数据传输呢？（PS：在一次数据传输中，转移的一片内存的首地址一定是32的整数倍，例如一次数据传输只能从全局内存读取地址为0到31字节、32到63字节等片段的数据；使用CUDA运行时API分配的内存的首地址至少是256的整数倍。）

• 不对齐的非合并访问

  __global__ void add_offset(float *x, float *y, float *z) {
      int tid = blockIdx.x * blockDim.x + threadIdx.x + 1;
      z[tid] = x[tid] + y[tid];
  }
  
  add_offset<<<128, 32>>>(x, y, z);
  

  第一个线程块中的warp访问x的1到32号元素（标号从0开始），假设x的首地址是256字节，那么该warp访问的数据对应的地址是260到387字节。这将触发5次数据传输，对应的内存地址分别为：256到287字节、288到319字节、320到351字节、352到383字节、384到415字节。合并度为4/5=80%。

• 广播式的非合并访问

  __global__ void add_broadcast(float *x, float *y, float *z) {
      int tid = blockIdx.x * blockDim.x + threadIdx.x;
      z[tid] = x[0] + y[tid];
  }
  
  add_offset<<<128, 32>>>(x, y, z);
  

  每次计算只需要x的0号元素（4字节），但是一次数据传输的数据量为32字节，因此合并度为4/32=12.5%。这样的访问（如果是读数据的话），非常适合使用常量内存。

6.2 矩阵转置

考虑对一个N*N的方阵A进行转置，转置后的矩阵记为B。

将矩阵A分割成一个个32*32的小方阵进行处理，示例代码如下：

__global__ void transpose(const double *A, double *B, const int N) {
    const int nx = blcokIdx.x * blockDim.x + threadIdx.x;
    const int ny = blockIdx.y * blockDim.y + threadIdx.y;
    
    if (nx < N && ny < N) {
        B[ny * N + nx] = A[nx * N + ny];
    }
}

const int grid_size_x = (N % 32 == 0) ? (N / 32) : (N / 32 + 1);
const int grid_size_y = grid_size_x;
const dim3 block_size(32, 32);
const dim3 grid_size(grid_size_x, grid_size_y);

transpose<<<grid_size, block_size>>>(A, B, N);

在这里，对A的访问是非合并的，对B的访问是合并的（相邻线程访问相邻矩阵元素）。

也有另一种转置写法：

B[nx * N + ny] = A[ny * N + nx];

在这里，对A的访问是合并的，对B的访问是非合并的。

从直觉上来看，这两种方式的执行速度应该差不多，因为都是一个合并访问、一个非合并访问。但事实上，第一种的执行速度约是第二种的两倍。 这是因为：第一种方式中，读取数据（访问A）是非合并访问、写入数据（访问B）是合并访问，从帕斯卡架构开始，如果编译器能够判断一个全局内存变量在整个核函数范围内都只可读（比如A），则会自动使用函数***__ldg()***读取全局内存（__ldg()是只读数据缓存的加载函数），从而对数据的读取进行 缓存，缓解非合并访问带来的影响。但是，对于全局内存的写入则没有类似的技术，因此第一种方式比第二种方式快。