CUDA中按列相加为何比按行相加更快？技术疑问求解

这是个非常好的问题，刚好戳中了CUDA内存访问优化里最核心的点——内存合并访问（Memory Coalescing），咱们一步步拆解你的疑惑：

先理清两个核函数的内存访问模式 #

首先得明确：CUDA全局内存的效率，核心取决于同一个Warp（32个线程为一组）内的线程，是否在访问连续的内存地址，而不是单个线程内部的访问是否连续。

1. th_single_row_add的问题：线程间访问完全分散 #

你的行求和核函数里，每个线程负责一整行的元素。假设我们用常见的256线程/Block配置，那：

• 线程0访问的是数组的第0行：a[0], a[1], ..., a[999]（连续的1000个元素）
• 线程1访问的是第1行：a[1000], a[1001], ..., a[1999]
• ...
• 线程31（同一个Warp的最后一个线程）访问的是第31行：a[31000], ..., a[31999]

看起来单个线程是连续访问，但同一个Warp里的32个线程，它们的起始地址间隔是1000个float（也就是4000字节），完全不连续。CUDA的全局内存是按事务（Transaction）读取的，这种分散的访问会导致每个线程都要发起独立的内存事务，内存带宽利用率极低——相当于明明可以一次拉取一大块连续内存，结果却要分32次拉取零散的小块，自然慢。

2. th_single_col_add的优势：线程间访问完美合并 #

而列求和核函数里，每个线程负责一整列的元素：

• 线程0访问的是第0列：a[0], a[1000], a[2000], ..., a[999000]（单个线程内是间隔访问）
• 线程1访问的是第1列：a[1], a[1001], ..., a[999001]
• ...
• 线程31访问的是第31列：a[31], a[1031], ..., a[999031]

重点来了：同一个Warp里的32个线程，第一次访问的地址是连续的32个float（0到31），刚好能合并成1个内存事务（对应128字节，符合CUDA的事务粒度）；第二次循环时，它们访问的地址是1000到1031，同样是连续的，又能合并成1个事务……以此类推，每次循环的内存访问都能完美合并，内存带宽利用率拉满，速度自然快很多。

你的误解点在哪里？ #

你之前以为“单个线程内索引间隔大速度慢”，但CUDA的内存优化逻辑是优先保证线程组（Warp）的访问连续性，而非单个线程的访问连续性。单个线程内的间隔访问，只要线程组整体是连续的，依然能高效利用内存带宽；反过来，单个线程连续访问但线程组访问分散，反而会严重浪费带宽。

额外建议：更高效的元素级并行 #

其实对于矩阵加法这种操作，最理想的核函数是让每个线程处理一个元素，这样天生就能保证完美的内存合并访问：

__global__ void th_elementwise_add(float* a, float* b, float* c, int n) {
    int global_idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (global_idx >= n * n) return;
    c[global_idx] = a[global_idx] + b[global_idx];
}

你可以测试一下这个核函数，速度会比你的列求和版本还要快不少——因为它没有循环开销，每个线程只做一次加法，内存访问完全合并。

内容的提问来源于stack exchange，提问作者Егор Лебедев