在第 15 章中，我们剖析了 GPU 的物理硬件——CUDA 核心、Tensor 核心、流式多处理器（SM）以及 SIMT 执行模型。但硬件不会自己工作，需要一套软件抽象把数万个计算核心组织成可编程的并行计算资源。这套抽象就是 CUDA 编程模型。
对于大模型开发者，你未必每天手写 .cu 文件，但无论是调优 DeepSpeed 的分布式策略、分析推理框架的 Kernel 耗时，还是理解 FlashAttention 为什么快，都必须先建立对 CUDA 线程层级的直觉。本节把这套层级拆成可理解的地图。
一、Host 与 Device：两个世界
CUDA 程序天然运行在异构计算环境中：
- Host（主机端）：CPU 及其内存（DRAM）。负责逻辑控制、数据准备、启动 Kernel。
- Device（设备端）：GPU 及其显存（HBM/GDDR）。负责大规模并行计算。
一个典型的执行流程是：Host 把数据从内存拷贝到显存 → 在 Device 上启动并行函数（Kernel）→ 计算完成后把结果拷回内存。现代大模型训练框架（PyTorch、DeepSpeed）已经把这层拷贝封装得很深，但当你看到 cudaMemcpyAsync 或 pin_memory 相关日志时，本质上就是在处理 Host 与 Device 之间的数据边界。
二、三级层级结构：Thread → Block → Grid
CUDA 用一套三级层级来组织并行任务。你可以把它想象成“工厂 → 车间 → 工人”的映射：
| 层级 | 英文 | 作用 | 类比 |
|------|------|------|------|
| 线程 | Thread | 最基本的执行单元，执行一条串行指令流 | 一个工人 |
| 线程块 | Block | 由若干线程组成，内部可协作、可同步 | 一个车间班组 |
| 网格 | Grid | 由若干 Block 组成，完成一次 Kernel 启动的全部工作 | 整座工厂的一次生产任务 |
1. 线程（Thread）：并行的原子
线程是 CUDA 中最细粒度的执行实体。每个线程有自己的：
- 私有寄存器（Register）：存局部变量，速度最快，容量极少；
- 局部内存（Local Memory）：寄存器溢出时的兜底，实际在显存中，慢；
- 程序计数器与执行状态：在 SIMT 模型中，一个 Warp（32 个线程）共享指令分发单元，但每个线程维护自己的执行上下文（15.4 节提到的 SIMT 核心特征）。
实用认知：大模型中的每一个 Token 的矩阵元素乘法、每一个注意力头的计算，最终都会被拆解成数以万计的 Thread 在硬件上同时执行。
2. 线程块（Block）：协作与共享的边界
Block 是一组可以互相协作的线程集合，关键特征：
- 共享内存（Shared Memory）：Block 内所有线程可访问的一块高速片上存储（通常每 Block 几十 KB）。这是 CUDA 性能优化的兵家必争之地——FlashAttention（17.3 节）的核心优化之一，就是把注意力计算中的重复访存从慢速 HBM 搬到快速的 Shared Memory 中进行复用。
- 线程同步：通过
__syncthreads()实现 Block 内全屏障同步。Block 之间无法直接同步，也不能假设执行顺序。 - 执行单元归属：一个 Block 只能被调度到一个 SM（流式多处理器）上运行，且一个 SM 可同时驻留多个 Block。
实用认知：Block 尺寸（通常取 128、256、512、1024）直接影响 Occupancy（占用率）——即 SM 上实际活跃的 Warp 比例。Block 太小会导致 SM 资源闲置；太大则可能因寄存器/Shared Memory 不足，反而让占用率下降。推理引擎（如 TensorRT、vLLM）的 Kernel 调优大量工作就是在反复试验 Block 尺寸。
3. 网格（Grid）：Kernel 的总览
Grid 是一次 Kernel 启动所涉及的全部 Block 的集合。它定义了这次并行任务的总工作量。
Grid 与 Block 都可以是一维、二维或三维的（dim3 类型）。这种多维设计最初是为了方便图像处理（x, y 像素坐标），在大模型中则常用于映射矩阵乘法的 (M, N) 维度或批处理 (batch, head, seq_len) 维度。
三、硬件映射：软件层级如何落到物理 GPU 上
这是从“写代码”到“调性能”的关键一跳。CUDA 的线程层级并非纯软件概念，它与第 15 章的硬件架构有严格的映射关系：
Grid (Kernel)
└── Block 0 ──→ 调度到 SM 0 ──→ 拆分为 Warp 0, Warp 1, Warp 2...
└── Block 1 ──→ 调度到 SM 1 ──→ 拆分为 Warp 0, Warp 1...
└── Block 2 ──→ 调度到 SM 0 (如果 SM 0 还有余量)
...- Warp（线程束）：硬件真正的调度单位，固定 32 个线程。一个 Block 里的线程按 32 个一组被 SM 的 Warp Scheduler 发射执行。
- 分支发散（Branch Divergence）：同 Warp 内的 32 个线程必须共享同一条指令流。如果代码出现
if (threadIdx.x % 2)这样的分支，导致一半线程走 A 路径、一半走 B 路径，硬件会串行化执行两条路径，屏蔽不参与的线程。这在注意力掩码（Attention Mask）处理中是常见的性能陷阱。 - SM 资源竞争：每个 SM 的寄存器文件和 Shared Memory 总量固定。Block 如果申请了过多资源，SM 能同时驻留的 Block/Warp 数量就会减少，直接拉低占用率。
实用认知：当你在大模型训练日志中看到 “Low GPU utilization” 但显存未满时，往往不是因为算力不够，而是因为 Kernel 的 Block 配置不合理，导致 Warp 无法充分填充 SM，或者频繁等待显存（Global Memory）而空闲。
四、线程定位：如何知道自己该算哪一部分
CUDA 为每个线程提供内置变量来确定自己的全局坐标：
| 内置变量 | 含义 | 示例 |
|----------|------|------|
| threadIdx | 当前线程在 Block 内的索引 (x, y, z) | threadIdx.x |
| blockIdx | 当前 Block 在 Grid 内的索引 (x, y, z) | blockIdx.x |
| blockDim | 每个 Block 的维度大小 | blockDim.x = 256 |
| gridDim | 整个 Grid 的维度大小 | gridDim.x = 4096 |
全局线程索引计算（一维情况）：
int global_tid = blockIdx.x * blockDim.x + threadIdx.x;在大模型矩阵运算中，这个 global_tid 通常被映射到输出矩阵的某个坐标 (row, col)，每个线程负责计算结果矩阵的一个或一组元素。例如，在计算 Query × Key^T 的注意力分数时，一个线程块可能负责处理一个注意力头中的一块局部注意力矩阵。
五、存储层级与线程层级的伴生关系
CUDA 的线程层级与存储层级（15.5 节）是伴生绑定的。理解这种绑定，是看懂后续 FlashAttention、PagedAttention 等优化的前提：
| 存储层级 | 速度 | 可见范围 | 与线程层级的关系 |
|----------|------|----------|------------------|
| 寄存器 | 极快（1 cycle） | 单个 Thread 私有 | 每个线程独享 |
| Shared Memory | 快（~20 cycles） | 同 Block 内共享 | Block 是共享内存的边界 |
| L2 Cache | 中等 | 全 GPU 共享 | 跨 Block 自动缓存 |
| Global Memory (HBM) | 慢（~400 cycles） | 全 GPU 共享 | Grid 级别，所有 Thread 可访问 |
大模型场景的典型数据流：
- 权重矩阵和输入激活从 HBM（Global Memory）加载；
- 按 Block 粒度切分到 SM，利用 Shared Memory 做数据复用（如 Q、K 向量的分块加载）；
- 每个 Thread 用寄存器保存自己的局部累加值；
- 计算结果写回 Shared Memory 汇总，最终一次性写回 HBM。
如果跳过 Shared Memory 这层优化，每个线程每次都去 HBM 读数据，GPU 的算力会被显存带宽完全饿死——这正是大模型推理中常见的 Memory-Bound 场景（16.2 节）。
六、给大模型开发者的实用认知
1. 你通常不需要手写 Kernel，但必须能读懂瓶颈
PyTorch、Triton、CUDA 生态已经封装了 torch.matmul、F.scaled_dot_product_attention 等高层 API。但当训练/推理性能不达标时，性能分析工具（Nsight Compute、PyTorch Profiler）会给出 Warp 级、Block 级的统计。不懂线程层级，你就看不懂 “Shared Memory Bank Conflict” 或 “Warp Stalls due to Execution Dependency” 这类报告。
2. 推理框架的性能差异本质是 Kernel 调度差异
vLLM 的 PagedAttention、FlashAttention 的 Kernel Fusion，核心 tricks 都是在 Block/线程层级重新组织内存访问模式，减少 HBM 读写。它们没有改变算法数学，只是改变了 CUDA 的并行执行策略。
3. 批量大小（Batch Size）与 Grid 规模的联动
推理时增大 Batch Size，本质上是增大了 Grid 中的总线程数，让 GPU 有更多 Warp 可供调度，从而掩盖显存延迟。但当 Batch 大到超过 L2 Cache 容量时，又会退化为带宽瓶颈。这个“甜蜜点”的寻找，需要结合具体模型的 Block 配置和 GPU 的 SM 数量来判断。
七、小结
CUDA 编程模型的核心是一套层级化的并行抽象：
- Thread：执行的原子，拥有私有寄存器；
- Block：协作的班组，共享 Shared Memory，受限于 SM 的物理资源；
- Grid：一次 Kernel 任务的总览，所有 Block 独立调度。
这三层结构通过 Warp（32 线程束）映射到 GPU 的 SM 硬件上。理解这一映射，你就理解了为什么大模型优化不仅是“算法问题”，更是“如何把矩阵切分成合适大小的 Block，让数万个线程在 SM 上高效运转”的工程问题。
在 17.2 节中，我们将进一步拆解 CUDA 的软件栈层级——驱动、Runtime、cuDNN、NCCL 之间的关系，以及在大模型分布式训练中，这些组件是如何协同工作的。