小泉水1. 合并访问的核心原理与价值
在GPU编程中，全局内存访问效率直接决定了计算性能的上限。很多开发者会遇到一个典型现象：精心设计的CUDA内核，理论计算强度很高，但实际运行效率却远低于预期。通过Nsight工具分析，常常会发现"Memory Throughput Utilization"指标低得惊人——这正是非合并访问（Uncoalesced Access）的典型表现。
1.1 从硬件视角理解合并访问
现代GPU的显存子系统采用高带宽设计，但实现峰值带宽需要满足特定访问模式。以NVIDIA Ampere架构为例，每个内存控制器对应256-bit的显存接口，最小传输单元为32字节（对应L1 cache line大小）。当线程束（Warp）中的32个线程发出内存请求时，硬件会将这些请求组合成内存事务（Memory Transaction）。
合并访问的本质，是将线程束内所有线程的内存访问请求合并为最少次数的内存事务。理想情况下，32个线程访问连续的128字节对齐内存块（32线程×4字节访问），仅需1次128字节的内存事务。而非合并访问最坏情况下，每个线程产生独立的内存事务，总共需要32次32字节事务——带宽利用率直接下降为1/32。
关键理解：合并访问不是软件层面的"优化选项"，而是硬件设计强制的效率门槛。不符合合并条件的访问模式会直接触发硬件层面的带宽浪费。
1.2 线程束与合并访问的关系
CUDA执行模型以线程束为基本调度单位，这决定了合并访问的判断粒度：
- 每个线程束32个线程共享相同的程序计数器
- 内存请求在线程束级别进行合并判断
- 合并效果仅在线程束内部评估，不同线程束之间无合并关系
以常见的float类型数据为例，合并访问要求：
- 线程束内线程访问的地址必须连续
- 所有访问地址必须落在对齐的128字节边界内（32个float×4字节）
- 访问顺序最好与线程ID顺序一致（threadIdx.x升序）
1.3 合并访问的数学表达
用数学语言描述，对于线程束内线程i（0≤i<32），其访问的全局内存地址应满足：
code复制address[i] = base_address + (threadIdx.x + blockIdx.x * blockDim.x) * element_size
其