CUDA 编程模型

前面几卷讲的是硬件（SM、Tensor Core、内存层次），这一章把硬件变成代码。CUDA 的编程模型用三个概念映射硬件：grid/block/thread 映射并行，五层内存空间映射内存层次，kernel 是一次启动的执行单位。理解这三组映射，是读得懂高性能内核（FlashAttention、CUTLASS）的前提。

三层并行：grid → block → thread

一个 CUDA kernel 的并行结构是三层嵌套：

// 每个 thread 算一个输出元素（朴素写法）
__global__ void add(float* C, const float* A, const float* B, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;  // 全局索引
    if (i < n) C[i] = A[i] + B[i];
}
// 启动：grid 里有多少个 block，每个 block 里有多少个 thread
add<<<ceil(n/256), 256>>>(C, A, B, n);

层	硬件映射	说明
thread	单个线程（寄存器私有）	你写的逻辑单位
block	一个 SM 上的常驻单元	block 内共享共享内存，可同步
grid	全 GPU	block 分发到各 SM

三个映射要点：

block 是资源分配单位：一个 block 的线程只能待在一个 SM 上（共享内存与同步的边界）；block 太小浪费 SM，太大受限于每 SM 常驻上限
warp 是执行单位（硬件把 block 内线程按 32 个切成 warp，见GPU 架构）：分支发散与合并访存都在 warp 粒度发生
grid 是吞吐单位：block 数要远大于 SM 数（H100 132 个 SM，通常启动几千个 block），SM 才有调度的余地

五层内存空间

CUDA 显式暴露了硬件的内存层次，内核代码要自己选择每份数据放哪：

空间	范围	生命周期	用途
寄存器	thread 私有	kernel 内	局部变量，最快
共享内存	block 内共享	kernel 内	tile 暂存，手动管理
L2	全 GPU	硬件管理	自动缓存
全局内存（HBM）	全 GPU	跨 kernel	数据的家
常量/纹理内存	只读缓存	—	只读数据的缓存优化

共享内存是高性能内核的核心：手动把数据分块搬进共享内存，反复使用——这就是分块（tiling），GEMM 与 FlashAttention 的全部秘密：

HBM →（整块搬运）→ 共享内存（tile） →（反复读）→ 寄存器 → Tensor Core

每一次"从 HBM 读进 tile"被 tile 内的所有计算复用——算术强度就是这样被抬高的。

同步与异步

两种同步原语对应两种边界：

block 内同步：__syncthreads()——block 内线程栅栏，tile 加载后同步再计算。block 之间没有同步原语（只有 cooperative groups 的特殊机制），这是内核设计的硬约束
流（stream）异步：同一流内顺序执行，不同流并发。cudaMemcpyAsync + 多流是"搬运与计算重叠"的原始手段；Hopper 的 TMA 把这个模式做进硬件（异步 tile 搬运，见GPU 架构）

现代内核的标准结构就是多级流水：TMA 搬下一块 → warp 计算当前块 → 两者重叠，用同步原语（mbarrier）协调。

一个 tile 化 GEMM 的骨架

把前面的概念拼起来，一个高性能矩阵乘的完整结构（CUTLASS 同款思路的简化版）：

__global__ void gemm_tiled(float* C, const float* A, const float* B) {
    __shared__ float As[TILE][TILE], Bs[TILE][TILE];   // 共享内存 tile

    for (int k = 0; k < K; k += TILE) {
        // 1. 协作加载：block 内所有线程把 A、B 的 tile 搬进共享内存
        load_tile(As, A, k); load_tile(Bs, B, k);
        __syncthreads();                                // 2. 栅栏同步

        // 3. 寄存器累加：每个线程算自己负责的输出块
        for (int i = 0; i < TILE; i++)
            acc += As[ty][i] * Bs[i][tx];               // 实际用 Tensor Core
        __syncthreads();
    }
    C[...] = acc;                                       // 4. 写回 HBM
}

四步结构（加载 → 同步 → 计算 → 写回）就是所有高性能内核的模板，差异只在：tile 多大、Tensor Core 怎么喂、加载怎么与计算重叠。

写内核 vs 用内核

工程上 90% 的情况不需要手写 CUDA：

选择	场景
直接用库（cuBLAS/CUTLASS）	标准算子（矩阵乘、卷积），见算子库
torch.compile / XLA	标准模型结构，编译器自动融合，见下一章
Triton	自定义算子但不想管线程，见下一节
手写 CUDA	新算子、极限性能、特殊硬件特性

手写 CUDA 的真实成本：写对容易，写快极难（合并访存、占用率、Tensor Core 喂数据、异步流水，每一层都是坑）。先确认瓶颈在算子上（用 Roofline 分析），再考虑手写。

本章要点

grid/block/thread 三层并行映射 GPU 硬件：thread 是逻辑单位，block 是资源单位（SM 常驻），grid 是吞吐单位。
五层内存空间显式暴露硬件层次；共享内存 + 分块复用是高性能内核的核心手法。
block 内同步（__syncthreads()）、跨 block 靠流与异步；Hopper 的 TMA 把异步搬运做进硬件。
"加载 → 同步 → 计算 → 写回"的 tile 化模板是所有 GEMM/FlashAttention 的骨架。
写对容易写快难：先 Roofline 分析瓶颈，再决定要不要手写。