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[