2. GPU硬件架构与线程模型映射

理解CUDA编程，最关键的一步是建立“软件线程”与“硬件执行单元”之间的映射直觉。很多初学者写出能跑的核函数，却说不清为什么某个写法快、某个写法慢，根源往往在于没有把线程层次和GPU物理结构对应起来。本节从硬件侧出发，逐层落到软件侧，最后给出可复用的索引计算方法。

一、从SM到warp：硬件如何“消化”线程

一块NVIDIA GPU由若干流多处理器（Streaming Multiprocessor，SM）组成。以常见的Ampere架构为例，一个SM内部包含多个CUDA核心、若干warp调度器、寄存器文件、共享内存/L1缓存以及特殊功能单元。SM是真正执行指令的物理单元，也是资源分配的基本单位：一个线程块（block）被整体调度到某一个SM上，块内线程不会跨SM拆分。

SM内部以warp为调度粒度。一个warp固定包含32个线程，线程编号连续。硬件为每个warp分配一个程序计数器，同一warp内的32个线程在同一时刻执行同一条指令，这就是SIMT（单指令多线程）模型。当warp内线程因条件分支走向不同路径时，会发生分支发散：不同路径被串行执行，未走该路径的线程被掩码屏蔽，效率下降。因此，让同一warp内线程尽量走相同分支，是核函数编写的重要原则。

SM通过warp调度器在多个常驻warp之间快速切换来隐藏延迟。当某个warp因等待全局内存数据而停顿时，调度器立刻切换到另一个就绪warp继续发射指令。这解释了为什么“线程数要足够多”：足够的常驻warp才能填满流水线。但常驻warp数量受寄存器、共享内存等资源限制，并非越多越好。

二、三级层次与四个内置变量

CUDA把并行组织成网格—线程块—线程三级。一次核函数启动（<<<grid, block>>>）产生一个网格，网格由若干线程块构成，每个块由若干线程构成。硬件侧对应关系是：网格整体提交给GPU，块被分配到SM，块内线程被划分为warp。

核函数内部可直接读取四个内置变量，它们都是dim3类型（含x、y、z三个分量）：

• gridDim：网格在各维度上的块数，即启动配置中的grid参数；
• blockIdx：当前线程块在网格中的坐标，从0开始；
• blockDim：每个线程块在各维度上的线程数，即启动配置中的block参数；
• threadIdx：当前线程在块内的坐标，从0开始。

注意区分“维度”和“大小”：blockIdx.x是坐标，gridDim.x是该维度的块总数。未使用的维度默认为1，因此一维启动时y、z分量恒为0或1，不影响计算。

三、线程索引计算：从一维到三维

线程索引的本质是把多维坐标“压平”成线性编号，以便访问一维数组。约定以x维变化最快（行优先），这与C语言数组的行优先存储一致。

一维网格最简单，全局索引为：

int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) { /* 处理 a[i] */ }

这里的if (i < n)是边界保护。当数据量n不是blockDim.x整数倍时，最后一个块中会有线程索引越界，必须屏蔽，否则会读写非法内存。

二维网格常用于图像、矩阵。设矩阵宽为width，则：

int col = blockIdx.x * blockDim.x + threadIdx.x;  // 列，变化最快
int row = blockIdx.y * blockDim.y + threadIdx.y;  // 行
if (row < height && col < width) {
    int idx = row * width + col;   // 行优先压平
    /* 处理 data[idx] */
}

启动配置示例：处理 1920×1080 的图像，取 block 为 (32, 8)，则 grid 为 (ceil(1920/32), ceil(1080/8)) = (60, 135)。块内共256线程，是32的整数倍，warp划分整齐。

三维网格用于体数据、三维仿真等。压平公式按x最快、y次之、z最慢：

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
if (x < nx && y < ny && z < nz) {
    int idx = (z * ny + y) * nx + x;
}

四、映射关系带来的实践约束

理解了映射，就能解释几条常见工程规则。第一，块大小取32的整数倍：若blockDim.x不是32的倍数，最后一个warp会有线程空闲，浪费发射槽。第二，块不宜过小或过大：过小则每块占用一个SM的调度资源却并行度不足；过大则受寄存器与共享内存限制，SM能常驻的块数减少，反而降低延迟隐藏能力。常见取值如128、256、512。第三，二维块应让x维连续，因为同一warp内threadIdx.x连续，访问相邻内存可形成合并访问（coalescing），显著提升全局内存带宽利用率。

最后强调一个易错点：blockIdx与threadIdx都是块内/网格内的相对坐标，不能直接当作全局编号使用；必须乘以对应的blockDim或gridDim再相加。把映射关系记牢，后续学习共享内存、同步与性能优化时，才能准确判断每个线程在硬件上究竟落在哪个warp、哪个SM，以及为什么某种访问模式会快或慢。