Memory Coalescing内存合并访问是 GPU 编程中针对Global Memory全局内存最基础且最重要的访存性能优化机制。它的核心思想是当一个 Warp32 个线程同时发起全局内存读写时如果这 32 个线程请求的内存地址是连续且对齐的GPU 硬件内存控制器会将这 32 次独立的内存请求“合并”为 1 次或极少数次大带宽的突发传输Memory Transaction。一、 为什么需要 Memory CoalescingGPU 的板载显存HBM / GDDR属于高延迟硬件。显存控制器并不是以单个 Byte 为单位传输数据的而是以Cache Line / 内存段通常为 32 字节、64 字节或 128 字节为基本单位进行突发读取。理想合并情况Coalesced AccessWarp 中的 Thread 0 读地址0, Thread 1 读地址4, …, Thread 31 读地址124连续读 32 个float共 128 字节。结果内存控制器发现这 128 字节恰好落在一个连续对齐的 128B Cache Line 内仅用1 次显存事务Transaction就满足了 32 个线程的全部需求。显存带宽利用率达到100%。未合并情况Uncoalesced / Strided Access如果线程间存在跨步Stride如 Thread 0 读0, Thread 1 读128, …, Thread 31 读3968。结果32 个线程请求的数据分散在 32 个不同的 Cache Line 中硬件不得不发起32 次独立传输。虽然每次传输取回了 128 字节但线程实际上只用了其中的 4 字节有效显存带宽利用率降至3.125%1/32严重挂起计算核心。二、 合并与未合并模式对比访问模式线程与地址对应关系 (ithreadIdx.xi \text{threadIdx.x}ithreadIdx.x)硬件显存事务数 (Transactions)带宽利用率完全合并 (Coalesced)AddressBasei×4\text{Address} \text{Base} i \times 4AddressBasei×41 次(128-Byte)100%乱序合并 (Permuted)地址在同一个 128B 块内但线程顺序随机1 次(128-Byte)100%现代 GPU 可重排跨步访问 (Strided)AddressBasei×Stride×4\text{Address} \text{Base} i \times \text{Stride} \times 4AddressBasei×Stride×4多倍增长最大 32 次极低随着 Stride 增大急剧下降完全随机 (Random)乱序指针/散列查表最大 32 次最低三、 经典非合并场景与解决策略1. 结构体数组AoS变 数组结构体SoA痛点使用 C/C 传统的 AoSArray of Structures模式例如定义struct Point { float x, y, z; } points[N];。当 Warp 中的线程并发读取points[i].x时相邻线程读取的地址中间隔着y和zStride 3导致无法合并。解法重构成 SoAStructure of Arrays将数据按属性拆分为独立连续数组float x[N], y[N], z[N];让线程依次读取x[i]。2. 矩阵按列读取Column-wise Read痛点行主序存储Row-Major的M×NM \times NM×N矩阵如果线程按列方向去读取数据如A[i * N j]中iii随线程变化相邻线程之间的内存间隔为NNN触发严重跨步。解法Shared Memory 中组转Staging。先将 Global Memory 中的矩阵块按行合并加载Coalesced Copy到 Shared Memory 中在 Shared Memory 中做转置Transpose再供线程读取。3. 向量化加载Vectorized Load / 128-bit优化现代 GPU 支持单个线程一次读取 64-bitfloat2或 128-bitfloat4/LDG.128指令。通过让少量的线程一次性搬运大块连续内存进一步减少硬件传输指令数最大化硬件并发数。四、 在现代 DSLTileLang / Triton中的体现在 TileLang 或 Triton 等高级算子语言中开发者通常不需要手写线程级别的索引但底层的内存合并原理依然起决定性作用在 TileLang 中写T.copy(A_global[row, col:col64], A_shared)时编译器在生成 CUDA / PTX 代码时会自动将连续切片col:col64映射为合并的连续访存指令。如果在 TileLang/Triton 中对 Global Memory 做了不连续的步长切片或复杂 Gather 操作底层同样会退化为未合并访存降低 Memory Throughput。