2. GPU硬件架构与CUDA执行模型

要写出真正高效的CUDA程序，必须先理解程序在硬件上“长什么样”。CUDA的执行模型并不是软件凭空抽象出来的，它几乎一一对应着GPU的物理结构：线程对应执行单元，线程块对应SM（Streaming Multiprocessor，流多处理器），网格对应整个GPU。本节从SM、warp调度、线程层次、内存层次到计算能力版本，把这条“软件到硬件”的映射关系讲清楚。

SM：GPU的基本计算单元

一块GPU由若干SM组成，SM是真正执行指令的物理单元。以较常见的架构为例，一个SM内部包含：多个CUDA核心（用于整数/浮点运算）、若干专用功能单元（如SFU用于超越函数、Tensor Core用于矩阵乘加）、寄存器文件、共享内存/L1缓存、warp调度器与分派单元。不同架构的SM内部配比差异很大，这正是“计算能力版本”要区分的原因。

关键点在于：线程块被整体调度到一个SM上执行，块内线程共享该SM的共享内存和寄存器资源。因此一个SM能同时驻留多少个线程块，取决于寄存器、共享内存、线程数等资源上限。如果某个块占用资源过多，SM就无法容纳更多块，占用率（occupancy）随之下降。

warp：调度的最小单位

CUDA把线程按32个一组打包成warp。warp是硬件调度和指令发射的基本单位，而不是单个线程。一个warp内的32个线程在同一时刻执行同一条指令，但操作各自的数据——这就是SIMT（单指令多线程）模型。

由此产生两个重要后果。第一，分支发散：如果warp内线程走不同分支，硬件会串行执行各分支路径，只有走该路径的线程有效，其余线程被屏蔽，效率成倍下降。第二，合并访存：如果一个warp的32个线程访问连续且对齐的全局内存地址，硬件可合并为少量内存事务；若地址分散，则事务数激增，带宽利用率骤降。

当某个warp因等待内存或依赖而停顿时，warp调度器会切换到另一个就绪warp继续发射指令，用“多warp隐藏延迟”。这就是为什么足够的占用率对性能至关重要——它给调度器提供了可切换的warp池。

线程层次结构：Grid → Block → Warp → Thread

CUDA用四级层次组织线程：网格（Grid）→ 线程块（Block）→ warp → 线程（Thread）。程序员在核函数中用内置变量定位线程：blockIdx（块索引）、threadIdx（块内线程索引）、blockDim（块维度）、gridDim（网格维度）。全局线程编号常写成：

int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) { /* 处理第 idx 个元素 */ }

块内线程可以通过共享内存和 __syncthreads() 协作，但不同块之间无法直接同步或共享共享内存。因此算法设计要尽量把需要频繁通信的数据放在同一个块内处理。块大小通常取32的整数倍（如128、256、512），避免出现“半个warp”浪费执行槽位。

内存层次与访问代价

GPU内存层次按延迟从低到高大致为：寄存器 → 共享内存/L1 → L2 → 全局内存 → 主机内存。寄存器访问几乎无延迟，但数量有限；共享内存由块内线程共享，是块内数据复用的关键；全局内存容量大但延迟高，必须靠大量并发warp来隐藏。

优化访存的核心原则是合并访问与数据复用。例如矩阵乘法中，把子矩阵块先加载到共享内存，再在块内反复使用，可把全局内存访问量降低一个数量级。此外，常量内存适合所有线程读取同一地址的场景，纹理内存适合有空间局部性的访问。

计算能力版本：决定可用特性与资源上限

计算能力（Compute Capability，CC）用主次版本号表示，如7.5、8.6、9.0。它决定了SM内部结构、每SM最大线程数、寄存器数量、共享内存大小、支持的指令集（如Tensor Core、异步拷贝、分布式共享内存等）。编译时用 -arch=sm_86 指定目标架构。

表中数值为常见配置示例，具体以官方文档为准。理解这些上限，才能判断一个核函数是否受寄存器或共享内存限制而无法达到理想占用率。

把本节串起来看：你写的每个线程最终落在某个SM的某个warp槽位上执行；块决定共享内存与同步边界；warp决定分支与访存效率；计算能力决定这些资源的数量上限。后续所有优化技巧，本质上都是在这套映射关系里做权衡。