文章目录摘要1. 为什么要学习 Memory Coalescing2. 什么是 Memory Coalescing2.1 合并访问Coalesced Access2.2 非合并访问Uncoalesced Access3. 实验代码设计3.1 Coalesced kernel3.2 Uncoalesced kernel4. 实验设置5. 实验结果5.1 原始结果6. 实验结果解读6.1 Coalesced 访问为什么快6.2 Uncoalesced 访问为什么慢7. 为什么 stride 越大slowdown 越明显8. 为什么 coalesced 时间也会变化9. 有效带宽估算9.1 Coalesced 有效带宽9.2 Uncoalesced 有效带宽10. 和矩阵、图像、深度学习有什么关系10.1 矩阵访问10.2 图像处理10.3 深度学习 Tensor Layout11. 实践优化建议11.1 让 threadIdx.x 对应连续内存11.2 二维数组优先保证行方向连续11.3 对非连续访问进行重排本课结论摘要第 7 课围绕 CUDA 中非常关键的Memory Coalescing即内存合并访问展开。实验通过对比连续访问data[idx]和跨步访问data[idx * stride]验证了 Global 内存访问模式对性能的巨大影响。结果显示在 Tesla T4 上当stride32时非合并访问比合并访问慢约25.5 倍当stride16时慢约14.6 倍当stride8时慢约8.0 倍。这说明 GPU 性能不仅取决于“算多少”更取决于 warp 内线程是否访问连续地址。1. 为什么要学习 Memory Coalescing在前几课中我们已经学习了Pinned Memory优化 CPU ↔ GPU 数据传输 Shared Memory优化 GPU kernel 内部数据复用 CUDA Stream让传输与计算重叠第 7 课关注的是另一个核心问题同样访问 Global 内存为什么有的访问方式很快有的访问方式会慢几十倍GPU 的 Global 内存访问不是以单个线程为单位孤立执行而是以warp为基本调度单位。一个 warp 通常有32 个线程如果这 32 个线程访问的是连续地址GPU 可以把这些访问合并成较少的内存事务。如果这 32 个线程访问的是分散地址GPU 就需要发起更多内存事务带宽利用率会急剧下降。2. 什么是 Memory CoalescingMemory Coalescing 可以理解为warp 内线程访问连续内存地址时GPU 将多个线程的访存请求合并成较少的内存事务从而提高 Global 内存带宽利用率。2.1 合并访问Coalesced Access例如thread 0 - data[0] thread 1 - data[1] thread 2 - data[2] ... thread 31 - data[31]这些地址是连续的。对于float来说每个元素 4 字节32 个线程访问的数据正好是连续的一段32 × 4 bytes 128 bytes这类访问非常适合 GPU 内存系统合并处理。2.2 非合并访问Uncoalesced Access如果访问模式变成thread 0 - data[0] thread 1 - data[32] thread 2 - data[64] thread 3 - data[96] ... thread 31 - data[992]这时线程之间访问地址相隔很远。对于stride32相邻线程地址间隔 32 × 4 bytes 128 bytes也就是说一个 warp 内的每个线程几乎都落在不同的内存区域GPU 很难把它们合并成少数内存事务。结果就是访问请求变多 带宽利用率下降 kernel 时间显著增加3. 实验代码设计本次实验设计两个 kernel。3.1 Coalesced kernel__global__voidcoalesced_access(float*data,size_t ops){size_t idx(size_t)blockIdx.x*blockDim.xthreadIdx.x;if(idxops){data[idx]1.0f;}}这个 kernel 的访问模式是data[0], data[1], data[2], ...也就是连续访问。3.2 Uncoalesced kernel__global__voiduncoalesced_access(float*data,size_t ops,intstride){size_t idx(size_t)blockIdx.x*blockDim.xthreadIdx.x;if(idxops){size_t access_idxidx*(size_t)stride;data[access_idx]1.0f;}}这个 kernel 的访问模式是data[0], data[stride], data[2 * stride], ...也就是跨步访问。4. 实验设置实验设置如下Array size: 4096 MB block size: 256 threads 测试 stride: 8、16、32 计时方式: cudaEvent 统计范围: kernel 执行时间代码中constsize_t n1ULL30;constsize_t bytesn*sizeof(float);constsize_t opsn/stride;这里的设计很重要ops n / stride这样可以保证 uncoalesced kernel 不越界。同时在同一组 stride 下coalesced kernel 和 uncoalesced kernel 执行相同数量的线程操作区别只在于访问模式不同。5. 实验结果5.1 原始结果StrideEffective accessesCoalesced timeUncoalesced timeSlowdown3233,554,4321.09542 ms27.9479 ms25.5133x1667,108,8642.15622 ms31.4834 ms14.6012x8134,217,7284.47146 ms35.6195 ms7.96598x这个结果符合课程预期stride 越大 warp 内线程访问越分散 memory coalescing 越差 性能下降越明显6. 实验结果解读6.1 Coalesced 访问为什么快以stride32那组为例coalesced kernel 访问的是data[0], data[1], data[2], ...warp 内线程访问连续地址GPU 可以高效合并访问。所以它只用了1.09542 ms这说明 GPU 对连续 Global 内存访问的带宽利用率很高。6.2 Uncoalesced 访问为什么慢同样是stride32uncoalesced kernel 访问的是data[0], data[32], data[64], data[96], ...一个 warp 内相邻线程间隔32 × 4 bytes 128 bytes这会导致原本可以合并的访问被拆散内存事务数量显著增加cache line / memory sector 利用率下降GPU 为了满足分散的线程访问不得不搬来一整段内存数据但每段里只有很少几个字节真正被线程使用结果是内存事务变多、有效带宽下降、kernel 变慢实际带宽被浪费所以它耗时27.9479 ms相比 coalesced 慢了25.5133x这是非常显著的性能差距。7. 为什么 stride 越大slowdown 越明显slowdown 是StrideSlowdown87.96598x1614.6012x3225.5133x这个趋势很清晰stride8 线程间地址间隔 32 bytes stride16 线程间地址间隔 64 bytes stride32 线程间地址间隔 128 bytesstride 越大warp 内线程访问越分散。因此合并访问能力越差 内存事务越多 带宽浪费越严重 性能越差所以 slowdown 从约 8 倍逐渐扩大到约 25 倍。8. 为什么 coalesced 时间也会变化coalesced 的时间分别是StrideopsCoalesced time3233,554,4321.09542 ms1667,108,8642.15622 ms8134,217,7284.47146 ms你会发现ops 翻倍coalesced time 也大致翻倍原因是opsn/stride;所以 stride 越小实际访问次数越多。这说明 coalesced kernel 的行为比较稳定主要受有效访问次数控制。9. 有效带宽估算因为data[idx]1.0f;大致可以看作一次读 一次写所以可粗略估算有效数据量useful_bytes ops × sizeof(float) × 2有效带宽bandwidth useful_bytes / time9.1 Coalesced 有效带宽StrideCoalesced time估算有效带宽321.09542 ms约 245 GB/s162.15622 ms约 249 GB/s84.47146 ms约 240 GB/sCoalesced 访问的有效带宽比较稳定说明连续访问能充分利用 Global 内存带宽。9.2 Uncoalesced 有效带宽StrideUncoalesced time估算有效带宽3227.9479 ms约 9.6 GB/s1631.4834 ms约 17.1 GB/s835.6195 ms约 30.2 GB/s从有效带宽看uncoalesced 的带宽利用率明显低得多。尤其是stride32时有效带宽只有约9.6 GB/s而 coalesced 约为245 GB/s两者差距非常大。10. 和矩阵、图像、深度学习有什么关系Memory Coalescing 在很多场景中都非常重要。10.1 矩阵访问如果矩阵按行主序存储A[row*Ncol]那么同一个 warp 中线程访问连续col通常比较快A[row][0], A[row][1], A[row][2], ...但如果按列访问A[0][col], A[1][col], A[2][col], ...在行主序内存中就是跨步访问可能变慢。10.2 图像处理图像通常也是二维数组。如果线程按行连续处理像素pixel[y][x], pixel[y][x1], pixel[y][x2]访问通常较好。如果线程跨列、跨通道、跨行访问可能产生非合并访问。10.3 深度学习 Tensor Layout深度学习中常见 layoutNCHW NHWC不同 layout 会影响某些 kernel 中线程访问是否连续。所以在高性能推理框架中layout 选择和 memory coalescing 密切相关。11. 实践优化建议最重要的工程建议是11.1 让threadIdx.x对应连续内存在 CUDA kernel 中尽量让threadIdx.x 0,1,2,3...对应访问data[base 0] data[base 1] data[base 2] data[base 3]而不是data[base 0 * stride] data[base 1 * stride] data[base 2 * stride]11.2 二维数组优先保证行方向连续如果数据是 C/C 行主序data[row*widthcol]那么通常应该让threadIdx.x 对应 col这样同一个 warp 的线程更容易访问连续地址。11.3 对非连续访问进行重排如果算法天然需要跨步访问可以考虑1. 改变数据布局 2. 使用 shared memory 做 tile 重排 3. 让读取阶段 coalesced计算阶段在 shared memory 中调整访问模式 4. 合并多个小访问减少随机访问矩阵转置就是经典例子Global 内存读写尽量 coalesced 中间用 Shared Memory 处理转置本课结论CUDA kernel 的性能不仅取决于计算量也强烈依赖 Global 内存访问模式。实验结果显示stride8 时uncoalesced 约慢 8 倍 stride16 时uncoalesced 约慢 14.6 倍 stride32 时uncoalesced 约慢 25.5 倍这说明warp 内线程访问连续地址 → 内存事务少 → 带宽利用率高 → kernel 快 warp 内线程访问分散地址 → 内存事务多 → 带宽利用率低 → kernel 慢一句话总结Memory Coalescing 的本质是让一个 warp 内的线程尽量访问连续 Global 内存地址从而减少内存事务、提高带宽利用率在 CUDA 优化中访存模式往往比算术指令本身更决定性能。