第 3 章 CUDA 编程模型七层

"The mark of a great architect is not how complicated they can make a system, but how cleanly they can layer it so each layer can be reasoned about independently." ——一句被反复引用的工程箴言

3.1 为什么需要七层抽象

第 2 章我们看到了 Hopper 硬件的精确结构——132 SM、每 SM 4 Sub-Core、每 Sub-Core 16 FP32 cores 等等。但程序员不会直接编程到这些硬件单元上。CUDA 在硬件之上铺了一层编程模型抽象，把硬件的复杂性藏起来，给程序员一个相对干净的并行 API。

这一层抽象一共有七级：

flowchart TB
  Host[Host · CPU 端]
  Device[Device · GPU 整卡]
  Grid[Grid · 一次 kernel launch 的所有线程]
  Cluster[Cluster · 一组协作的 Block · Hopper 新增]
  Block[Block · 一组共享 SMEM 的线程]
  Warp[Warp · 32 线程的硬件调度组]
  Thread[Thread · 单线程标量代码]

  Host -->|cudaMalloc, kernel<<<...>>>| Device
  Device -->|kernel 启动| Grid
  Grid -->|可选| Cluster
  Cluster --> Block
  Block --> Warp
  Warp --> Thread

为什么要这么多层？答案是：每一层都对应一种"协作的粒度"。

• Thread 层：你写普通的标量代码（局部变量、算术、判断、循环）。
• Warp 层：32 线程同步执行，可用 warp shuffle 做无 SMEM 的快速通信（warp_shfl_xor / warp_reduce）。
• Block 层：上千线程共享 SMEM，可用 __syncthreads 同步。
• Cluster 层（Hopper 新增）：多 Block 之间通过分布式共享内存协作。
• Grid 层：一次 kernel launch 的所有线程。Grid 内不同 Block 默认无法直接通信，只能通过 atomic 或多次 launch。
• Device 层：整张 GPU 卡，host 通过 device API（cudaMalloc、cudaMemcpy）操作。
• Host 层：CPU 上的代码，负责数据准备、kernel launch、与其他系统组件交互。

不同算子在不同层次上做协作。比如 LayerNorm 主要在 Block 层协作（一个 token 的所有维度由一个 Block 处理）；GEMM 主要在 Warp 层协作（一个 warp-group 处理一个输出 tile）；分布式 GEMM 跨 Device 协作（多卡 NCCL 通信）。

下面把每一层从下往上展开。

3.2 Thread 层：你写的代码看起来是单线程的

CUDA C++ 写 kernel 时，函数体内的代码看起来就是一段普通的 C++ 单线程代码：

__global__ void simple_kernel(float* arr, int N) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid < N) {
        arr[tid] = arr[tid] * 2.0f + 1.0f;
    }
}

这段代码每个线程独立执行一遍。每个线程通过 threadIdx、blockIdx、blockDim、gridDim 这几个内置变量知道"自己是谁"。每个线程有：

• 私有寄存器：局部变量（如上面的 tid）默认存在寄存器里。
• 私有的 PC（程序计数器）：每个线程在 Volta 之后有独立的 PC（Independent Thread Scheduling）。
• 共享的 SMEM 与 HBM：和其他线程协作的载体。

线程层的关键性能直觉：

1. 不要假设线程间有任何顺序：线程 0 不一定比线程 1 先执行任何指令。如果你需要顺序，必须用同步原语。
2. 局部变量是寄存器：不要担心局部变量"开销"——它们都是寄存器，零开销。
3. 避免线程间的"个性化分支"：如第 1 章讲的，warp 内 divergence 会让性能减半。

3.3 Warp 层：32 线程的同步乐队

每 32 个连续线程组成一个 warp。warp 是 SIMT 模型的真正调度单位——硬件给一个 warp 发同一条指令。

CUDA 在 warp 层暴露的关键 API：

3.3.1 Warp Shuffle：寄存器间的零开销通信

__shfl_* 系列指令允许一个 warp 内 32 个线程直接交换寄存器值，不经过 SMEM：

// 把 lane 0 的值广播到 warp 内所有 32 个 lane
int val = my_value;
val = __shfl_sync(0xFFFFFFFF, val, 0);

// 把每个 lane 的值传给它的"对偶 lane"（XOR 模式）
int val = my_value;
val = __shfl_xor_sync(0xFFFFFFFF, val, 16);

// 树形归约
for (int offset = 16; offset > 0; offset >>= 1) {
    val += __shfl_xor_sync(0xFFFFFFFF, val, offset);
}
// 此时所有 lane 都有 warp 内 32 个值的和

第一个参数 0xFFFFFFFF 是 mask，表示 32 个 lane 全部参与（如果某些 lane 因为 divergence 不活跃，需要相应调整 mask）。

Warp shuffle 比 SMEM 快得多：

访问 SMEM (无 bank conflict):  20-30 cycles
Warp shuffle:                  4-6 cycles

这是为什么 reduce 算法在 warp 内一定要用 shuffle 而不是 SMEM。第 5 章会有完整的 warp-level reduce 实现。

3.3.2 Warp Vote：32 票快速决策

__all_sync / __any_sync / __ballot_sync 让 32 个 lane 快速投票：

int my_val = ...;
// 1) all: 是否所有 lane 的值都非零
unsigned all_nonzero = __all_sync(0xFFFFFFFF, my_val != 0);

// 2) any: 是否有任意 lane 的值非零
unsigned any_nonzero = __any_sync(0xFFFFFFFF, my_val != 0);

// 3) ballot: 收集每个 lane 的布尔值到一个 32-bit 整数
unsigned mask = __ballot_sync(0xFFFFFFFF, my_val > 0);
// mask 的第 i 位 = (lane i 的 my_val > 0)
int count = __popc(mask);  // 多少个 lane 满足条件

这些指令在以下场景非常有用：

• 稀疏 attention：只对满足条件的 token 做计算，先用 ballot 收集 mask。
• 动态 batch：判断当前 warp 是否有任何线程还有活要做。
• 变长输入：每个 lane 处理不同长度的序列，用 ballot 决定提前退出哪些 lane。

3.3.3 Cooperative Groups：现代化的 warp API

CUDA 9 引入的 Cooperative Groups 库提供了更现代的 warp 协作 API：

#include <cooperative_groups.h>
namespace cg = cooperative_groups;

__global__ void kernel() {
    auto block = cg::this_thread_block();
    auto warp = cg::tiled_partition<32>(block);

    int val = ...;
    int sum = cg::reduce(warp, val, cg::plus<int>());

    if (warp.thread_rank() == 0) {
        // 此时 sum 是这个 warp 32 个值的和
    }
}

Cooperative Groups 比裸的 __shfl_* 更安全（自动处理 mask）、更可读，但本质做的事是一样的。新代码推荐用 CG，老代码常见 __shfl_*。

3.4 Block 层：SMEM 协作的舞台

一个 Block 是一组 1-1024 个线程的集合，关键约束：

1. 同一个 Block 的所有线程必须调度到同一个 SM 上。这是硬件约束，不是软件选择。
2. Block 内的线程可以共享 SMEM：声明 __shared__ 变量即可。
3. Block 内可以用 __syncthreads() 同步：让所有线程等到同一点再继续。

Block 是 CUDA 编程中最常用的协作粒度。一个典型的 Block-level reduce：

template <int BLOCK_SIZE>
__global__ void block_reduce(const float* arr, float* out, int N) {
    __shared__ float smem[BLOCK_SIZE];
    int tid = threadIdx.x;
    int gid = blockIdx.x * blockDim.x + tid;

    // 1. 每线程读一个值到 SMEM
    smem[tid] = (gid < N) ? arr[gid] : 0.0f;
    __syncthreads();

    // 2. block 内分层归约
    for (int s = BLOCK_SIZE / 2; s > 0; s >>= 1) {
        if (tid < s) smem[tid] += smem[tid + s];
        __syncthreads();
    }

    // 3. 写出 partial sum
    if (tid == 0) out[blockIdx.x] = smem[0];
}

Block 大小（block size，又叫 CTA size）的选择是个值得讨论的问题。

3.4.1 怎么选 Block Size

Hopper 的硬件约束：

Block 大小上限:                 1024 线程 (32 warps)
SM 上同时活跃线程数上限:        2048 线程 (64 warps)
SM 上同时活跃 Block 数上限:     32 (软件配置, Hopper 默认 32)
SM 上 SMEM 总:                  228 KB
SM 上寄存器总:                  256 KB (65536 个 32-bit)

选 Block Size 时要平衡多个因素：

1. Occupancy（占用率）：希望 SM 上活跃 warp 数多 → block 不要太大（小 block 让 SM 上能容纳更多 block，提供更多 warp 备选）。
2. 协作效率：block 太小（比如 32）时无法做 block-level reduce，浪费 SMEM 协作能力。
3. SMEM 容量：如果一个 block 用了 80 KB SMEM，那一个 SM 只能放 2-3 个 block，occupancy 受限。
4. 寄存器压力：每线程寄存器用量越多，能 active 的线程越少。

常用的 block size 经验值：

Block size 没有"最优值"——必须在具体 kernel 上 profile。但 128 / 256 / 512 是出现频率最高的三个。

3.4.2 SMEM 的两种声明方式

// 1. 静态 SMEM (大小编译期已知)
__shared__ float smem_static[1024];

// 2. 动态 SMEM (大小运行时确定)
extern __shared__ float smem_dynamic[];

// host 端 launch:
int dynamic_smem_size = 1024 * 4;  // 1024 floats = 4 KB
my_kernel<<<grid, block, dynamic_smem_size>>>(...);

动态 SMEM 适合 tile 大小可调的算子（比如 GEMM 不同输入需要不同 tile size）。Hopper 上动态 SMEM 上限默认 48 KB，超过需要用 cudaFuncSetAttribute opt-in 到 228 KB：

cudaFuncSetAttribute(my_kernel,
    cudaFuncAttributeMaxDynamicSharedMemorySize,
    228 * 1024);

第 11 章手写 GEMM 时会用到这一点。

3.4.3 __syncthreads() 的代价

__syncthreads() 让 block 内所有线程等到同一个点。它的开销在 Hopper 上是 15-30 cycles——不算贵，但也不便宜。

一个常见的反模式：在循环内频繁同步：

// 反例：每次循环都 sync
for (int i = 0; i < N; i += BLOCK_SIZE) {
    smem[tid] = arr[i + tid];
    __syncthreads();
    /* compute on smem */
    __syncthreads();
}

更好的写法是用软件流水线（software pipelining）让计算和加载重叠：

// 加载 buffer 0
smem[0][tid] = arr[0 * BLOCK_SIZE + tid];
__syncthreads();

for (int i = 1; i < N / BLOCK_SIZE; ++i) {
    int cur = (i - 1) % 2;
    int next = i % 2;
    // 加载 next buffer 同时算 cur buffer
    smem[next][tid] = arr[i * BLOCK_SIZE + tid];
    /* compute on smem[cur] */
    __syncthreads();
}

这是 Tiled GEMM 的核心模式。第 11 章会反复用到。

3.5 Cluster 层：Hopper 引入的 Block 协作

Cluster 是 Hopper 引入的新一层抽象。Pre-Hopper 代码可以完全忽略它（兼容性默认 cluster_size=1），但 LLM 算子在 Hopper 上越来越多地利用它。

3.5.1 启用 Cluster

CUDA 提供两种方式声明 Cluster 大小：

// 1. 静态声明 (编译期固定)
__global__ void __cluster_dims__(2, 2, 1) my_kernel(...) {
    // 这个 kernel 的 cluster 是 2×2×1 = 4 个 block
}

// 2. 动态配置 (host 启动时指定)
cudaLaunchConfig_t config = {
    .gridDim = grid,
    .blockDim = block,
    .dynamicSmemBytes = smem_size,
    .stream = stream,
    .attrs = (cudaLaunchAttribute[1]){
        {.id = cudaLaunchAttributeClusterDimension,
         .val.clusterDim = {2, 2, 1}}
    },
    .numAttrs = 1,
};
cudaLaunchKernelEx(&config, my_kernel, args...);

Cluster 大小最大 16（推荐 ≤ 8，因为 GPC 内 SM 数有限）。

3.5.2 Cluster 内 API

Cluster 内的 Block 通过几个新内置变量协作：

__global__ void cluster_kernel() {
    cooperative_groups::cluster_group cluster =
        cooperative_groups::this_cluster();

    int cluster_size = cluster.dim_blocks();      // 总 block 数
    int block_rank   = cluster.block_rank();      // 当前 block 在 cluster 内的 ID

    // 关键: 跨 block 访问 SMEM
    extern __shared__ float smem[];
    float* peer_smem = cluster.map_shared_rank(smem, target_block_rank);
    // 现在 peer_smem 指向另一个 block 的 SMEM, 可直接读写
}

map_shared_rank 是 Cluster 的核心：它把当前 block 的 SMEM 指针重映射到另一个 block 的 SMEM。映射通过专用的 SM-to-SM Network 完成，延迟比走 L2 低很多（约 100 cycles vs 300+ cycles）。

3.5.3 Cluster 的同步

Cluster 内的 block 同步用 cluster.sync()：

__global__ void cluster_reduce(float* arr, float* out, int N) {
    auto cluster = cooperative_groups::this_cluster();
    extern __shared__ float smem[];
    int tid = threadIdx.x;

    // 1. 每 block 算自己的 partial sum
    int gid = (cluster.block_rank() * blockDim.x + tid) +
              cluster.dim_blocks() * blockDim.x * blockIdx.x;
    /* ... block-level reduce 到 smem[0] ... */
    __syncthreads();

    // 2. 跨 block 收集 (block 0 收集所有 partial sum)
    if (cluster.block_rank() == 0 && tid == 0) {
        float total = smem[0];
        for (int i = 1; i < cluster.dim_blocks(); ++i) {
            float* peer = cluster.map_shared_rank(smem, i);
            total += peer[0];
        }
        out[blockIdx.x] = total;
    }
    cluster.sync();
}

这种"Cluster reduce"比传统两阶段（先 block reduce，再 launch 第二个 kernel 跨 block reduce）快 2-3x，因为省去了第二次 kernel launch 的 ~5μs 开销，对小数据特别有用。

3.6 Grid 层：一次 launch 的全部线程

一次 kernel<<<grid, block>>> 启动的所有 block 组成一个 grid。Grid 之间的 block 是同伴关系——它们看到一样的全局内存，但默认无法直接通信。

3.6.1 Grid 间通信的几种手段

1. 全局 atomic：成本最高（HBM atomic 几百 cycle），但简单。
2. 多次 kernel launch：第一个 kernel 算 partial 结果，第二个 kernel 汇总。每次 launch ~5μs 开销。
3. CUDA Graph：把多个 kernel 串成一个 graph，整体提交，省去重复 launch 开销。
4. Cooperative Launch + grid.sync()：少数情况下可以用 grid-level 同步，但要求所有 block 同时活跃，对 occupancy 限制大。

LLM 算子里最常见的是 多次 kernel launch——比如 reduce 大数组：先 launch 一个 kernel 算每 block 的 partial，再 launch 一个 kernel 汇总。一次推理的几百次 launch 加在一起也只占 1-2ms，可以接受。

3.6.2 Grid 大小的选择

理论上 grid_size 可以很大（gridDim.x ≤ 2^31, y/z ≤ 65535）。实践中两种典型策略：

flowchart LR
  subgraph Strategy1 [策略 1: One Block per Output]
    S1A[每个 Block 处理一个输出 tile]
    S1B[grid_size = N_outputs]
    S1C[Block 数过多时硬件自动队列]
  end
  subgraph Strategy2 [策略 2: Persistent Kernel]
    S2A[grid_size = SM 数 × occupancy]
    S2B[每个 Block 用 grid-stride loop 循环处理多个 tile]
    S2C[Block 数固定不变]
  end

One Block per Output 是默认的写法，简单直观。Persistent Kernel 是 Hopper 上推荐的写法（第 18 章详细讲），优势是：

1. 避免 Block launch 开销：Block 启动有 ~100-200 cycle 的开销，persistent kernel 用一个 Block 处理多个输出 tile，摊薄。
2. 更好的负载均衡：每个 Block 用 atomic counter 抢下一个 tile，自然均衡。
3. 更大的 tile cache：Block 内 SMEM 在不同 tile 间可以复用（如 KV cache）。

3.7 Device 层：一张卡

Device 层的 API 是 host 端调用的，关键的有：

// 设备查询
int device_count;
cudaGetDeviceCount(&device_count);

cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);
// prop.multiProcessorCount  -> SM 数 (132 for H100 SXM5)
// prop.sharedMemPerBlock    -> SMEM 默认上限 (48 KB)
// prop.sharedMemPerBlockOptin -> SMEM Optin 上限 (228 KB Hopper)
// prop.regsPerBlock         -> 寄存器/block 上限
// prop.warpSize             -> 32

// 设备选择
cudaSetDevice(0);

// 内存管理
float* d_ptr;
cudaMalloc(&d_ptr, N * sizeof(float));
cudaMemcpy(d_ptr, h_ptr, N * sizeof(float), cudaMemcpyHostToDevice);
cudaFree(d_ptr);

// Stream 与 Event (异步)
cudaStream_t stream;
cudaStreamCreate(&stream);
my_kernel<<<grid, block, 0, stream>>>(...);

cudaEvent_t event;
cudaEventCreate(&event);
cudaEventRecord(event, stream);
cudaEventSynchronize(event);

LLM 推理中最常用的 device API 是 stream（异步执行）和 graph（批量提交）。附录 A 会详细讲。

3.7.1 多 GPU 与 NVLink

Hopper 系统中常见的多 GPU 拓扑：

DGX H100 (8 张 H100):
  GPU 0 ─── NVSwitch ─── GPU 1
   │                        │
   ├─── NVSwitch ─── GPU 2  ├─── NVSwitch ─── GPU 3
   │                        │
   ...
  900 GB/s 全互联带宽 (任意 GPU 对都是 900 GB/s)

NVLink 4 在 H100 上每条链路 25 GB/s，每张 H100 有 18 条链路 = 450 GB/s 单向 / 900 GB/s 双向。这个带宽是 PCIe Gen5 (~64 GB/s 单向) 的 14x，但仍然远低于 HBM3 (3350 GB/s)。这就是为什么张量并行 (TP) 通常受限于 NVLink 带宽——大模型推理时跨卡通信是关键瓶颈。

跨 GPU 的 CUDA API：

// 显式 P2P 拷贝
cudaMemcpyPeer(dst_ptr, dst_dev, src_ptr, src_dev, size);

// 或 NCCL (推荐)
ncclAllReduce(send_buf, recv_buf, count, ncclFloat, ncclSum, comm, stream);

本书是单 GPU 视角，多 GPU 通信的细节请见《vLLM 内核探秘》第 14 章。

3.8 Host 层：把所有东西串起来

Host 端的代码负责：

int main() {
    // 1. 准备数据
    float* h_a = (float*)malloc(N * sizeof(float));
    init_data(h_a, N);

    // 2. 上 GPU
    float* d_a;
    cudaMalloc(&d_a, N * sizeof(float));
    cudaMemcpy(d_a, h_a, N * sizeof(float), cudaMemcpyHostToDevice);

    // 3. Launch kernel
    int block_size = 256;
    int grid_size = (N + block_size - 1) / block_size;
    my_kernel<<<grid_size, block_size>>>(d_a, N);

    // 4. 同步等待结果
    cudaDeviceSynchronize();

    // 5. 取回结果
    cudaMemcpy(h_a, d_a, N * sizeof(float), cudaMemcpyDeviceToHost);

    // 6. 清理
    cudaFree(d_a);
    free(h_a);
    return 0;
}

绝大多数 host 代码长这个样。但 LLM 推理服务里，host 端比这复杂得多——CUDA Graph、多 stream pipeline、page-locked memory、cudaMallocAsync 等，附录 A 会专门讲。

3.9 七层抽象与硬件的映射

最后把七层抽象与第 2 章的硬件结构对应起来：

flowchart TB
  subgraph Software [编程模型抽象]
    Host3[Host CPU]
    Device3[Device · 整 GPU]
    Grid3[Grid]
    Cluster3[Cluster]
    Block3[Block · 1024 thread]
    Warp3[Warp · 32 thread]
    Thread3[Thread]
  end
  subgraph Hardware [硬件实体]
    CPU3[CPU + System Memory]
    GPU3[GH100 Die · 132 SM + L2 + HBM]
    GPC3[GPC · 18 SM]
    SM3[SM]
    SubCore3[Sub-Core · 16 cores]
    Reg3[Register file slot]
  end
  Host3 --- CPU3
  Device3 --- GPU3
  Grid3 -.-> GPU3
  Cluster3 -.-> GPC3
  Block3 -.-> SM3
  Warp3 -.-> SubCore3
  Thread3 -.-> Reg3

注意几个映射的特点：

1. 不是一一对应：一个 SM 可以同时跑多个 block，所以 Block ↔ SM 是 N:1。
2. 硬件无感知 Grid：硬件不知道"grid"是什么概念，它只调度 block。Grid 是软件抽象。
3. Cluster ↔ GPC 是新的对应：Cluster 内的 block 必须在同一 GPC 内，这是 Hopper 才有的硬件约束。

这个映射理解清楚后，性能调优就有方向了：

• 想要 warp-level 高效？关注 warp shuffle、warp vote、ballot。
• 想要 block-level 高效？关注 SMEM 利用、__syncthreads 频率、bank conflict。
• 想要 cluster-level 高效？关注 distributed SMEM 拓扑、cluster_sync 频率。
• 想要 grid-level 高效？关注 occupancy、kernel launch 次数、persistent kernel 模式。

3.10 这一章的小结与下一章

这一章从软件角度梳理了 CUDA 的七层编程模型抽象：

1. Thread：你写的标量代码。
2. Warp：32 线程的硬件调度组，有 shuffle / vote / ballot 等专属 API。
3. Block：SMEM 共享、__syncthreads 同步的舞台。
4. Cluster（Hopper 新增）：通过分布式 SMEM 跨 block 协作。
5. Grid：一次 launch 的全部 block。
6. Device：一张卡。
7. Host：CPU 端代码。

每一层都有自己的协作 API 和性能直觉。后续章节会反复回到这七层——比如第 5 章 Reduce 用到 warp 和 block 层，第 11 章 Tiled GEMM 用到 block 层 SMEM 与 __syncthreads，第 17 章 FA2 SOTA 用到 cluster 层与 warp specialization。

第 4 章我们会换一个最关键的视角——内存层级。GPU 有 register / SMEM / L1 / L2 / HBM 五层内存，每一层的延迟和带宽差几十倍。LLM 算子优化的核心几乎全部围绕"怎么让数据多在 SMEM 和寄存器里待着，少往 HBM 跑"展开。读完第 4 章，读者会有一份精确到 cycle 的"内存代价表"，可以放在桌面，写每一行 kernel 代码时心里都有数。

本章动手练习：

1. 写一个 kernel，每个 block 输出自己的 (blockIdx, threadIdx, warp_rank, lane_rank)，验证 32 线程组成 warp 的规律。
2. 用 Cooperative Groups 写一个 warp-level reduce，对比第 5 章会给的 __shfl_xor_sync 版本，看哪种更易读。
3. 在 H100 上写一个最简单的 cluster kernel：cluster_size=2，让 block 0 写 SMEM，block 1 通过 map_shared_rank 读 block 0 的 SMEM，验证 distributed SMEM 工作。