CUDA两种块维度配置下内存访问速度差异原因问询

这个问题其实戳中了CUDA全局内存访问最核心的一个点——合并访问的优先级远高于单个线程的内存遍历步长，咱们拆解开来聊：

首先得明确两个关键前提：

• CUDA的全局内存是按行优先存储的（也就是A[i][j]的内存地址是A + i*N + j）；
• CUDA是SIMT（单指令多线程）架构，线程块内的所有线程会同步执行相同的指令，它们的内存访问是同时发起的，而全局内存的性能几乎完全取决于这些并发访问是否能被硬件合并成更少的内存事务。

情况1：block(1,32)，每个线程遍历行（stride 1） #

你设置的线程块是1×32的，意味着每个线程块里有32个线程（对应threadIdx.y从0到31），每个线程负责遍历一整行的元素。

当这些线程同时开始访问各自行的第一个元素时，它们的内存地址是：
(blockIdx.y*32 + 0)*N, (blockIdx.y*32 +1)*N, ..., (blockIdx.y*32 +31)*N

这些地址之间的间隔是N个元素（如果是float类型，就是4*N字节），完全是分散的！这种情况下，硬件无法合并访问，只能为每个线程单独发起一个内存事务，带宽利用率极低——哪怕单个线程后续是连续的stride 1访问，也救不了整体的内存效率。

情况2：block(32,1)，每个线程遍历列（stride N） #

线程块是32×1的，32个线程对应threadIdx.x从0到31，每个线程负责遍历一整列的元素。

当这些线程同时开始访问各自列的第一个元素时，它们的内存地址是：
0*N + (blockIdx.x*32 +0), 0*N + (blockIdx.x*32 +1), ..., 0*N + (blockIdx.x*32 +31)

这些地址是连续的！硬件可以把这32个访问合并成1个内存事务（刚好对应128字节的内存 burst，符合CUDA的内存访问对齐要求），带宽利用率接近100%。

虽然单个线程后续访问的是stride N的内存，但因为线程块的并发访问是合并的，整体内存吞吐量比情况1高得多，所以速度更快。

核心结论 #

• 不要只看单个线程的内存遍历步长，线程块内的并发访问是否连续合并才是全局内存性能的关键；
• CUDA的合并访问是按threadIdx.x维度对齐的，所以通常推荐线程块的x维度设为32的倍数（比如32、64），这样更容易实现合并访问；
• SIMT架构下，线程块内的线程同步执行，它们的内存访问行为是绑定在一起的，这一点和CPU的多线程完全不同。

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