连续线程需访问连续数据吗？CUDA索引方式性能对比分析

测试代码与场景 #

我们对比两种矩阵加法的线程索引实现（基于Grace Hopper GPU，采用16×16线程块）：

__global__ void matrixAdd1(int *d_a, int *d_b, int *d_c, int width, int height) {
  // 推荐版本：连续线程访问连续内存
  int row = blockIdx.y * blockDim.y + threadIdx.y;
  int col = blockIdx.x * blockDim.x + threadIdx.x;
  // 替代版本：连续线程跨步访问内存
  // int col = blockIdx.y * blockDim.y + threadIdx.y;
  // int row = blockIdx.x * blockDim.x + threadIdx.x;

  if (row < height && col < width) {
    int idx = row * width + col;
    d_c[idx] = d_a[idx] + d_b[idx];
    // 修改后重复加法操作以放大内存开销
    // for(int i=0; i<100; i++) d_c[idx] += d_a[idx];
  }
}

测试场景：百万级元素矩阵，初始测试性能差异微小，修改代码重复加法操作后，性能出现3倍差距。两类索引方式的性能指标对比如下：

推荐索引方式（正确）性能指标 #

----------------------- ----------- ------------
    Metric Name             Metric Unit Metric Value
    ----------------------- ----------- ------------
    DRAM Frequency                  Ghz         2.62
    SM Frequency                    Ghz         1.51
    Elapsed Cycles                cycle      272,148
    Memory Throughput                 %        69.18
    DRAM Throughput                   %        40.12
    Duration                         us       179.58
    L1/TEX Cache Throughput           %        66.13
    L2 Cache Throughput               %        87.64
    SM Active Cycles              cycle   263,722.43
    Compute (SM) Throughput           %        48.22
    ----------------------- ----------- ------------

替代索引方式（错误）性能指标 #

----------------------- ----------- ------------
    Metric Name             Metric Unit Metric Value
    ----------------------- ----------- ------------
    DRAM Frequency                  Ghz         2.62
    SM Frequency                    Ghz         1.53
    Elapsed Cycles                cycle      743,618
    Memory Throughput                 %        66.57
    DRAM Throughput                   %        14.89
    Duration                         us       486.11
    L1/TEX Cache Throughput           %        60.71
    L2 Cache Throughput               %        66.57
    SM Active Cycles              cycle   739,352.92
    Compute (SM) Throughput           %        17.57
    ----------------------- ----------- ------------

问题解答 #

1. 两种索引方式本应存在性能差异吗？ #

必然存在差异，核心原因是内存访问的合并效率不同：

• 推荐方式中，同一warp内的连续线程（threadIdx.x递增）对应矩阵的连续列，计算出的线性索引idx = row*width + col是连续的，符合GPU内存系统的合并访问要求，单个内存事务就能满足warp内所有线程的请求，DRAM带宽利用率高（测试中达40.12%）。
• 替代方式中，连续线程对应矩阵的连续行，线性索引idx = row*width + col会以width为步长跨步访问内存。这种情况下，warp内每个线程的内存请求无法合并，需要发起大量独立的内存事务，DRAM带宽利用率骤降（仅14.89%），最终导致执行时间大幅增加。

初始测试差异微小，是因为百万级数组可能部分或全部被L2缓存覆盖，缓存的高命中率掩盖了DRAM访问模式的劣势；当重复加法操作后，内存访问的总开销占比提升，跨步访问的低效性被放大，性能差距就显现出来。

2. 如何验证索引方式的优劣？ #

可以通过以下几种方式验证：

• 性能分析工具监控核心指标：使用Nsight Compute查看内存合并率（Memory Coalescing）、内存事务数（Memory Transactions）、DRAM吞吐量。合并率越高、事务数越少，索引方式越优。
• 跨数据规模测试：在小数据（缓存可覆盖）、大数据（超出缓存容量）场景下分别测试。小数据场景可能差异不明显，大数据场景能直接体现内存访问模式的影响。
• 理论带宽利用率计算：根据线程访问模式，计算理想带宽利用率。比如连续访问时，32个线程访问连续32个int（128字节），刚好匹配Grace Hopper的内存事务粒度，利用率接近100%；跨步访问时，若width远大于事务粒度，利用率会低至1/width级别，和实际测试结果对比即可验证。

3. 连续线程是否需要访问连续数据？ #

是的，这是CUDA编程中实现高效内存访问的核心原则：
GPU的SM以warp为单位调度线程，同一warp内的线程会同时发起内存请求。当连续线程访问连续的对齐内存地址时，内存控制器会将这些请求合并为少量事务，最大化带宽利用率；反之，连续线程跨步访问非连续地址，会导致大量无法合并的独立请求，不仅浪费DRAM带宽，还会增加内存访问延迟，显著降低kernel性能。

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