CUDA 算子工程：手写 FlashAttention v2 之路

第 3 章 CUDA 编程模型七层

作者 杨艺韬 · 5,224 字 · 发布于 · 更新于

第 2 章讲的是硬件长什么样，这一章讲的是 CUDA 允许你怎么去指挥它。 两者不是一一对应的：每一层抽象都对应一种"协作的粒度"，而不是一块具体的硅。

3.1 为什么需要七层抽象

第 2 章我们看到了 Hopper 硬件的精确结构——132 SM、每 SM 4 Sub-Core、每 Sub-Core 32 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 的快速通信（__shfl_sync / __shfl_xor_sync / __shfl_down_sync）。
  • 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 在 Block 层分 tile、在 Warp 层做矩阵乘（Hopper 上一个 warp-group 协作发出 WGMMA）；分布式 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. 标量局部变量通常就是寄存器：不必担心它们的"开销"；但寄存器用量超出上限会溢出（spill）到 local memory，按运行时下标访问的局部数组也会被放进 local memory，那才有真实代价（第 4 章）。
  3. 避免线程间的"个性化分支"：如第 1 章讲的，warp 内两条分支各占一半线程时，这段代码的吞吐直接减半。

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 模式）
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 的优势不在"单条指令延迟更低"——SHFL 走的是和 SMEM 相近的数据通路，延迟量级是可比的。它赢在别处：

经 SMEM 交换一次:  一条 store + 一条 load + 一次 __syncthreads, 还占掉 SMEM 容量
Warp shuffle:      一条 SHFL 指令, 不占 SMEM 容量, 不需要 block 级同步, 也没有 bank conflict

这是为什么 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 的值都非零
int all_nonzero = __all_sync(0xFFFFFFFF, my_val != 0);

// 2) any: 是否有任意 lane 的值非零
int 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>
#include <cooperative_groups/reduce.h>   // cg::reduce / cg::plus 在这个头里，漏了编译不过
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 (CC 9.0 硬件上限)
SM 上 SMEM 总:                  228 KB
SM 上寄存器总:                  256 KB (65536 个 32-bit)

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

  1. Occupancy（占用率）：希望 SM 上活跃 warp 数多。block 太大时资源按整块分配，零头容易浪费；block 太小又会先撞上每 SM 32 个 block 的上限（32 线程的 block 最多凑出 32 个 warp，只有 50% occupancy，至少 64 线程才可能满 64 warp）。
  2. 协作效率：block 太小（比如 32）时 block 级协作退化成 warp 内协作，SMEM 协作能力用不上。
  3. SMEM 容量：如果一个 block 用了 80 KB SMEM（加上 CUDA 每 block 预留的 1 KB 即 81 KB），228 KB 的 SM 只能放 2 个 block，occupancy 受限。
  4. 寄存器压力：每线程寄存器用量越多，能 active 的线程越少。

常用的 block size 经验值：

场景 推荐 Block Size 理由
Element-wise / 简单 reduction 256 / 512 高 occupancy，不需要太多 SMEM
LayerNorm / Softmax 256 / 512 / 1024 一个 block 处理一行/一个 token
GEMM (Tiled) 128 / 256 tile 内一组 warp 协作
FA2 Forward 128 (4 warps × 32)；hdim 192/256 部分配置用 256 4 个 warp 沿行分摊一个 query block（csrc/flash_attn/src/flash_fwd_launch_template.h）

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——注意每 block 的上限是 227 KB（每 SM 的 228 KB 减去 CUDA 预留的 1 KB），写 228 KB 会直接失败：

cudaFuncSetAttribute(my_kernel,
    cudaFuncAttributeMaxDynamicSharedMemorySize,
    227 * 1024);   // Hopper 每 block 上限；用 prop.sharedMemPerBlockOptin 查更稳

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

3.4.3 __syncthreads() 的代价

__syncthreads() 让 block 内所有线程等到同一个点。屏障指令本身不贵，真正的代价是等待：block 里最慢的那个 warp 到达之前，其余 warp 都停在屏障前不能发射后续指令。

一个常见的低效写法：单 buffer，每轮两次同步，加载和计算完全串行：

// 低效：两次 sync 都省不掉（第二次防止下一轮覆盖还在被读的 smem），但加载与计算串行
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;
    float v = arr[i * BLOCK_SIZE + tid];  // 先发出下一块的全局加载（落到寄存器）
    /* compute on smem[cur] */             // 加载在途时算当前 buffer
    smem[next][tid] = v;                   // 算完再写进 next buffer
    __syncthreads();
}

/* compute on smem[(N / BLOCK_SIZE - 1) % 2] */

注意循环体里的顺序。如果直接写 smem[next][tid] = arr[...]，写 SMEM 的指令要等全局加载返回才能发射，这个 warp 的计算排在它后面，谈不上重叠（nvcc 13.4 编到 sm_90a，SASS 里 STS 紧跟在 LDG 之后、先于计算用的 LDS/FFMA，不同 CUDA 版本可能不同）。先读进寄存器、算完再写 SMEM，正是 CUTLASS include/cutlass/gemm/threadblock/mma_pipelined.h 的寄存器中转做法。

这是 Tiled GEMM 的核心模式。第 11 章会反复用到，并换成 Ampere 起的 cp.async，由硬件把数据异步搬进 SMEM，连寄存器中转都省了。

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 启动时指定)
//    （这种方式下 kernel 不能再带 __cluster_dims__）
cudaLaunchConfig_t config = {};
config.gridDim = grid;               // 各维须是 cluster 对应维的整数倍
config.blockDim = block;
config.dynamicSmemBytes = smem_size;
config.stream = stream;

cudaLaunchAttribute attrs[1];        // 必须活到 cudaLaunchKernelEx 调用之后
attrs[0].id = cudaLaunchAttributeClusterDimension;
attrs[0].val.clusterDim.x = 2;
attrs[0].val.clusterDim.y = 2;
attrs[0].val.clusterDim.z = 1;
config.attrs = attrs;
config.numAttrs = 1;

cudaLaunchKernelEx(&config, my_kernel, args...);

CUDA 保证可移植的 cluster 上限是 8 个 Block；H100 上还可以 opt-in（cudaFuncAttributeNonPortableClusterSizeAllowed）到 16，但这会让 kernel 绑死在特定硬件上。

3.5.2 Cluster 内 API

Cluster 内的 Block 通过 Cooperative Groups 的 cluster_group 协作：

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

    unsigned cluster_size = cluster.num_blocks();  // 总 block 数（dim_blocks() 返回的是 dim3）
    unsigned 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。对映射后地址的读写经 GPC 内专用的 SM-to-SM Network 完成，数据不必绕道 L2，延迟因此显著更低——这正是 DSMEM 的全部价值所在（本专栏没有条件实测具体周期数，§4.1 那张延迟表给的是各层的量级）。

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
    //    注意 blockIdx 仍然是 grid 全局编号（cluster 不会把它重新编号），
    //    所以全局下标照常写；cluster 内的位置才用 block_rank()。
    int gid = blockIdx.x * blockDim.x + tid;
    /* ... block-level reduce 到 smem[0] ... */
    __syncthreads();

    // 2. **先等所有 block 都把 partial sum 写进自己的 smem[0]**。
    //    这一步不能省：cluster.sync() 之前去读 peer 的 SMEM，读到的是未定义值。
    cluster.sync();

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

    // 4. 退出前再同步一次：只要还有 block 在读我的 SMEM，我就不能退出。
    cluster.sync();
}

两次 cluster.sync() 一次都不能少：前一次保证被读的数据已经写好，后一次保证被读的 SMEM 在读完之前不会随着 block 退出而失效。漏掉任何一次都是竞态——而且在小 cluster 上大概率"跑起来是对的"，等换个规模才炸。

这种"Cluster reduce"省掉的是第二次 kernel launch（单次 launch 开销在微秒量级）以及中间数组的一次 HBM 往返。数据量越小、kernel 本身越短，这部分固定开销占比越高，收益也越明显；数据量一大，它就淹没在 HBM 带宽里了。

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

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

3.6.1 Grid 间通信的几种手段

  1. 全局 atomic：原子操作在 L2 上执行，同一地址争用激烈时会被串行化，但简单。
  2. 多次 kernel launch：第一个 kernel 算 partial 结果，第二个 kernel 汇总。每次 launch 在 host 端有几微秒开销。
  3. CUDA Graph：把多个 kernel 串成一个 graph，整体提交，省去重复 launch 开销。
  4. Cooperative Launch + grid.sync()：少数情况下可以用 grid-level 同步，但要求所有 block 同时驻留，grid 大小不能超过 SM 数 × 每 SM 可驻留 block 数。

LLM 算子里最常见的是 多次 kernel launch——比如 reduce 大数组：先 launch 一个 kernel 算每 block 的 partial，再 launch 一个 kernel 汇总。对 prefill 这类单个 kernel 就要跑几百微秒以上的场景，launch 开销可以忽略；但 decode 阶段每个 kernel 只跑几十微秒，几百次 launch 的开销与有效工作同一量级，这时就要靠 CUDA Graph（附录 A）。

3.6.2 Grid 大小的选择

理论上 grid_size 可以很大（gridDim.x ≤ 2^31 − 1, 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 数 × 每 SM 驻留 block 数]
    S2B[每个 Block 用 grid-stride loop 循环处理多个 tile]
    S2C[Block 数固定不变]
  end

One Block per Output 是默认的写法，简单直观。Persistent Kernel 是 Hopper 上 CUTLASS、FA3 等高性能 kernel 普遍采用的写法（第 18 章详细讲），优势是：

  1. 摊薄 Block 派发开销：persistent kernel 用一个 Block 处理多个输出 tile，省下的是反复的 block 派发、资源分配与回收，以及尾部的调度抖动（第 18 章解释为什么不能按"block 数 × 单 block 开销"串行相加）。
  2. 更好的负载均衡：每个 Block 用 atomic counter 抢下一个 tile，自然均衡。
  3. 数据复用在 L2：SMEM 只是几个 stage 的环形缓冲，每个 tile 仍要重新载入；可调度的是 tile 的处理顺序，让相邻 tile 共用的数据尽量命中 L2（§18.2.2）。

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 -> 每 block 的 SMEM Optin 上限 (227 KB Hopper：
//                              每 SM 228 KB，CUDA 每 block 预留 1 KB)
// 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 会详细讲。

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

DGX H100 (8 张 H100):
  GPU 0   GPU 1   ...   GPU 7
    │ ╲     │ ╲           │ ╲      每张 GPU 的 18 条 NVLink 分别连到
    ▼  ▼    ▼  ▼          ▼  ▼     4 颗 NVSwitch 芯片（每颗都连着全部 8 张卡）
  [NVSwitch 0] [NVSwitch 1] [NVSwitch 2] [NVSwitch 3]
  每张 GPU 到 NVSwitch 结构 900 GB/s 双向 (全互联, 不随对端是谁而变)

NVLink 4 在 H100 SXM5 上每条链路单向 25 GB/s，每张有 18 条链路 = 450 GB/s 单向 / 900 GB/s 双向。按双向口径比，NVLink 4 的 900 GB/s 是 PCIe Gen5 x16（128 GB/s 双向）的 ~7 倍；但仍然远低于 H100 SXM5 HBM3 的 3350 GB/s。所以张量并行（TP）的扩展常受 NVLink 带宽制约——TP 每层都要做跨卡归约，卡数一多，通信很容易成为瓶颈。

跨 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 等，其中 CUDA Graph 与 stream 附录 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[H100 SXM5 · 132 SM + L2 + HBM]
    GPC3[GPC · 全片每 GPC 18 SM]
    SM3[SM]
    SubCore3[Sub-Core · 32 FP32 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 逐个派发给 SM，SM 只看得到 block。
  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 与 cluster 层，第 11 章 Tiled GEMM 用到 block 层 SMEM 与 __syncthreads，第 17 章用 TMA 与 warp specialization 把 FA2 写到 SOTA。

第 4 章我们会换一个最关键的视角——内存层级。GPU 有 register / SMEM / L1 / L2 / HBM 五层内存，层与层之间的延迟和带宽差距从几倍到几十倍。LLM 算子优化的核心几乎全部围绕"怎么让数据多在 SMEM 和寄存器里待着，少往 HBM 跑"展开。读完第 4 章，读者会有一份按量级标注各层延迟与带宽的"内存代价表"，可以放在桌面，写每一行 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 工作。