高效的全局内存访问需要保证两点对齐访问和相邻的线程访问相邻的数据。This is oversimple, but the correct way to do it is just have consecutive threads access consecutive memory addresses.GPU在访问全局内存的时候通常一次性传输32 bytes或者128 bytes如果L1和L2 cache都被使用了那么一次性传输128 bytes如果只有L2 cache被使用那么一次性传输32 bytesL1和L2 cache是否被使用和GPU的架构等都有关系L1 cache行和设备内存中一个128字节对齐的内存区域对应。对齐访问(Aligned memory accesss)要求访问的设备内存的首地址是缓存大小的偶数倍Aligned memory accesses occur when the first address of a device memory transaction is an even multiple of the cache granularity being used to service the transaction (either 32 bytes for L2 cache or 128 bytes for L1 cache). Performing a misaligned load will cause wasted bandwidth.。注此处有问题详情可见https://stackoverflow.com/questions/75457768/does-an-aligned-memory-accesses-in-cuda-need-to-address-an-even-multiple-of?rq1正确的应该是这样的Global memory resides in device memory and device memory is accessed via 32-, 64-, or 128-byte memory transactions. These memory transactions must be naturally aligned: Only the 32-, 64-, or 128-byte segments of device memory that are aligned to their size (i.e., whose first address is a multiple of their size) can be read or written by memory transactions.合并内存访问则要求warp内的线程访问内存的连续块Coalesced memory accesses occur when all 32 threads in a warp access a contiguous chunk of memory.。如下图就是一个既满足对齐访问又满足合并访问的示例下图则是不的对齐不合并访问的内存访问示例In this case, there may be as many as three 128-byte memory transactions to read the data from device memory: one starting at offset 0 to include the data being read below the contiguous region, one at offset 256 to read the data being read above the contiguous region, and one at offset 128 that fetches the bulk of the data. Note that most of the bytes fetched by the lower and upper memory transactions will not be used, leading to wasted bandwidth.讲解一为什么“每个线程连续访问”反而可能不是合并访问假设现在为了简化说明只看 4 个线程thread0 thread1 thread2 thread3有一个二维数组0 1 2 3 4 5 6 7 8 9 a b在 C/CUDA 中二维数组默认采用row-major行优先存储因此它在物理内存中的排列实际上是一维连续的0 1 2 3 4 5 6 7 8 9 a b也就是说可以把它理解成floatdata[12];其中data[0] 0 data[1] 1 data[2] 2 data[3] 3 data[4] 4 ... data[11] b现在考虑两种不同的线程访问方式。第一种访问方式每个线程负责一段连续数据访问模式如下thread0: 0, 1, 2 thread1: 3, 4, 5 thread2: 6, 7, 8 thread3: 9, a, b对应的 CUDA 代码可以写成__global__voidkernel1(constfloat*data){inttidthreadIdx.x;for(inti0;i3;i){floatxdata[tid*3i];// 使用 x}}乍一看这种写法似乎非常合理。因为从单个线程的角度来看thread0: data[0], data[1], data[2] thread1: data[3], data[4], data[5] thread2: data[6], data[7], data[8] thread3: data[9], data[10], data[11]每个线程访问的地址都是连续的。但是 CUDA 的 global memory coalescing 并不是这样判断的。合并访问关注的是同一个 warp 中的线程在执行同一条 global memory load/store 指令时它们访问的地址是否集中。因此我们应该按照每一次循环来看。第一次循环i 0所有线程执行floatxdata[tid*30];于是thread0 - data[0] thread1 - data[3] thread2 - data[6] thread3 - data[9]也就是0 3 6 9 ↑ ↑ ↑ ↑ T0 T1 T2 T3线程之间访问的地址不是连续的而是以 stride 3 的方式跳跃。第二次循环floatxdata[tid*31];访问thread0 - data[1] thread1 - data[4] thread2 - data[7] thread3 - data[a]即1 4 7 a第三次循环则是2 5 8 b因此从 warp 的视角来看真实访问模式实际上是同一条 load 指令thread0thread1thread2thread3i 00369i 1147ai 2258b所以虽然每个线程自己访问的是连续地址但是同一时刻多个线程访问的地址却是分散的。这通常会导致 GPU 需要访问更多的 memory segment产生更多 global memory transaction最终降低有效显存带宽。第二种访问方式每个线程负责一列现在换一种映射thread0: 0, 4, 8 thread1: 1, 5, 9 thread2: 2, 6, a thread3: 3, 7, b对应 CUDA 代码__global__voidkernel2(constfloat*data){inttidthreadIdx.x;for(inti0;i3;i){floatxdata[i*4tid];// 使用 x}}从单个线程来看这一次反而是不连续访问thread0: 0 - 4 - 8 thread1: 1 - 5 - 9 thread2: 2 - 6 - a thread3: 3 - 7 - b但是这并不重要。因为我们仍然应该从warp 执行同一条 load 指令的角度来看。第一次循环i 0执行floatxdata[0*4tid];于是thread0 - data[0] thread1 - data[1] thread2 - data[2] thread3 - data[3]即0 1 2 3 ↑ ↑ ↑ ↑ T0 T1 T2 T34 个线程访问连续地址。第二次循环thread0 - data[4] thread1 - data[5] thread2 - data[6] thread3 - data[7]即4 5 6 7第三次循环8 9 a b整理以后同一条 load 指令thread0thread1thread2thread3i 00123i 14567i 289ab可以看到每次执行 global memory load 时相邻线程访问的都是相邻地址。因此这种模式更容易被 GPU 合并为较少的 memory transactions显存利用率也更高。为什么 CUDA 要这样判断CUDA GPU 并不是简单地先让 thread0 读数据 然后让 thread1 读数据 然后 thread2 ...而是以warp为基本执行单位。真正的 CUDA warp 中有 32 个线程。当这 32 个线程一起执行floatxdata[index];这条 global memory load 时memory subsystem 会同时看到整个 warp 的地址请求然后判断这些线程请求的数据落在多少个 global memory segment 中如果访问类似data[0]data[1]data[2]...data[31]假设元素是float32 threads × 4 bytes 128 bytes这些数据集中在一段连续的 128B 内存区域中因此只需要覆盖这一小片连续地址。这就是典型的合并访问。而如果代码类似data[tid*Ni]且N很大那么固定某一次i后thread0 - data[i] thread1 - data[N i] thread2 - data[2N i] thread3 - data[3N i] ...32 个线程请求的地址可能散落在完全不同的 memory segment 中。于是一次 warp load 就可能需要大量 memory transactions。虽然每个线程真正需要的数据可能只有32 × 4B 128B但是显存系统实际搬运的数据量可能远大于 128B因此有效带宽会明显下降。一个非常重要、也很反直觉的结论判断 CUDA global memory 是否合并访问时不要沿着单个线程的执行过程看地址。例如下面这种thread0: 0 - 1 - 2看起来很连续但这并不能说明它是合并访问。真正应该固定某一条 memory instruction然后横向观察 warp 中所有线程thread0 thread1 thread2 thread3 第一次 load 0 3 6 9 第二次 load 1 4 7 a 第三次 load 2 5 8 b这种访问实际上是 stride access。相反thread0: 0 - 4 - 8单线程看非常不连续但是thread0 thread1 thread2 thread3 第一次 load 0 1 2 3 第二次 load 4 5 6 7 第三次 load 8 9 a b每一条 load 指令对应的 warp 地址却非常集中因此反而是更好的 global memory access pattern。总结可以把这个例子的核心记成一句话CUDA 的合并内存访问看的是“同一个 warp 中的线程在同一条 global memory 指令上访问哪些地址”而不是“某一个线程前后访问的地址是否连续”。因此对于下面两种代码// stride accessdata[tid*Ni];和// contiguous access across threadsdata[i*Ntid];在典型情况下第二种访问模式更加符合 CUDA global memory coalescing 的要求。当然严格来说“相邻线程访问相邻元素”只是最典型、最容易获得高效率的访问模式。现代 NVIDIA GPU 最终关心的是一个 warp 的地址请求需要覆盖多少个 memory segment。地址越集中需要的 memory transactions 通常越少地址越分散需要的 transaction 越多有效显存带宽也就越低。来源https://stackoverflow.com/questions/5041328/in-cuda-what-is-memory-coalescing-and-how-is-it-achieved讲解二①一个warp里面的所有线程在任何时刻执行相同的指令注意一个warp里面的线程的线程id是连续递增的。②如果相邻的数据被同时访问那么就会引发DRAM bursts此时相邻的数据会被以极高的速度传输所谓的合并内存访问就是引发DRAM bursts。③由于一个warp里面的所有线程执行相同的指令,因此任何时刻他们都在相同的循环中也就是说如果线程内有for、whild循环等那么他们在任意时刻的循环次数都是一样的也就是都是第X次循环)。上图为两种内存访问方式矩阵M也就是图(A)的访问方式这种访问方式中对于每个线程第i次循环会访问某一行的第i个元素。矩阵N也就是图(B)的访问方式在这种访问方式中对于每个线程第i次循环会访问某一列的第i个元素。N的详解图如下由于相同的循环也就是相同的时间访问了相邻的数据因此其速度会很快符合合并内存访问模式。而对于矩阵M的访问模式如下图矩阵M的访问模式则如上图可见其在循环0的时候T0、T1、T2和T3分别访问M[0]、M[4]、M[8]和M[12]不符合相邻线程访问相邻内存的模式。来源https://nichijou.co/cuda5-coalesce/https://zhuanlan.zhihu.com/p/300785893讲解三二维的合并内存访问模式如《CUDA C权威编程指南》所述从逻辑角度来看线程块是线程的集合它们可以被组织为一维、二维或三维布局。从硬件角度来看线程块是一维线程束的集合。在线程块中线程被组织成一维布局 每32个连续线程组成一个线程束。在一个块中每个线程都有一个唯一的ID。用x维度作为最内层的维度y维度作为第二个维度z作为最外层的维度则二维或三维线程块的逻辑布局可以转化为一维物理布局。例如对于一个给定的二维线程块在 一个块中每个线程的独特标识符都可以用内置变量threadIdx和blockDim来计算对于一个三维线程块计算如下而一个warp由全局ID连续的线程组成因此二维情况下其访问模式为tid_in_block threadIdx.x threadIdx.y * blockDim.x; //线程在block中的线程全局ID bid_in_grid blockIdx.x blockIdx.y * gridDim.x; //现在所在的block的block全局ID threads_per_block blockDim.x * blockDim.y; //一个block中的线程数 tid_in_grid tid_in_block thread_per_block * bid_in_grid; //二维情形下的线程全局ID global_memory[tid_in_grid] ...;来源https://stackoverflow.com/questions/15459230/how-to-properly-coalesce-writes-from-global-memory-into-global-memory