如何合并CUDA内核中与块大小一致的分块全局内存读取

核心问题分析 #

你当前的L2扇区读写比1.09~1.12已经属于极高合并度的水平，仅存的少量额外访问几乎都来自分块首地址未对齐到GPU缓存行边界，你之前把基地址存入共享内存再偏移的操作属于冗余操作，反而会额外增加共享内存访问开销，对合并访问没有任何帮助。

优化方案 #

1. 优先做分块地址对齐（收益最高）

CUDA的全局内存合并访问要求连续访问序列的起始地址对齐到缓存扇区边界：

• 你使用的float2为8字节，32线程的Warp单次连续读取32个float2的总大小为256字节，对应安培及更新架构的2个128字节L2扇区
• 只要你将所有分块的首地址对齐到256字节边界（即32个float2的大小），即可实现100%合并访问，扇区读写比降到1.0
• 对齐实现方法：
  • 分配全局数组时直接按256字节对齐（cudaMalloc默认分配的内存已经对齐到256字节，无需额外处理）
  • 每个分块之间增加填充字节，直到分块首地址满足256字节对齐要求，业务逻辑处理时跳过填充位即可

2. 移除冗余的共享内存基地址存储

块内所有线程访问的分块基地址是统一的，直接将基地址存储在寄存器中即可，直接使用global_ptr[block_base + threadIdx.x]的寻址模式，CUDA的地址生成单元会自动识别连续访问模式，触发最优合并策略，不需要绕一层共享内存读取基地址。

3. 可选：用更大的向量加载进一步提升吞吐量

如果你的业务逻辑允许，可改用float4类型批量加载，每个线程一次加载2个连续的float2，512个元素仅需256次内存请求，内存吞吐量可提升近一倍：

// 假设分块首地址已经对齐到16字节（float4的对齐要求）
float4 val4 = *reinterpret_cast<const float4*>(global_ptr + block_base + threadIdx.x * 2);
float2 val0 = make_float2(val4.x, val4.y);
float2 val1 = make_float2(val4.z, val4.w);

4. 针对归约阶段单处未合并的排查

如果仅第二处读取出现未合并，优先检查该读取的地址偏移计算逻辑，大概率是循环迭代时的偏移量没有保证Warp内连续，可在Nsight Compute的Memory Workload Analysis面板中查看对应指令的访问地址分布，确认是否存在跨扇区的零散访问。

补充说明 #

你当前的合并度已经非常高，如果全局内存访问不是你kernel的核心瓶颈，不需要过度优化，上述对齐操作即可覆盖所有可优化的空间。

内容的提问来源于stack exchange，提问作者tripatheea