L2.3 经典 CUDA 编程精要与进阶

三维坐标　layer: L2（数据与算子）｜level: Senior｜pillar: 编程与编译

上一篇我们站在「编程与编译」支柱的山脚，知道一行 a @ b 会被编译成某个底层 kernel。本文要钻进这个 kernel 内部：为什么同样一段矩阵乘法，手写得好与坏能差出 10 倍以上。答案不在算法，而在「如何把数据搬到离计算单元最近的地方、并让上千个线程协同复用它」——这正是 CUDA 编程的本质，也是所有高性能算子（FlashAttention、cuBLAS、CUTLASS）的共同地基。

学习目标

• 前置知识：读过 L1.1（SM 内部结构、warp 调度与延迟掩盖、occupancy 概念）与 L2 编程支柱开篇（知道一行 a @ b 会被编译成 kernel）；会写基础 C/C++ 与 Python；知道「GPU 靠海量线程并发掩盖访存延迟」即可。无需任何 CUDA 编程经验。
• 学完产出：① 能把 CUDA 的软件三级层次（Thread / Block / Grid，外加 Warp）一一映射到硬件实体（CUDA Core / SM / GPU），并据此解释 gridDim/blockDim 该怎么配；② 能讲清 Shared Memory「程序员显式管理的 L1」为什么是高性能 kernel 的胜负手，并用「HBM 访问从 O(N) 降到 O(N/TILE)」量化 tiled GEMM 的加速来源；③ 能定位并消除 bank conflict——说清「32 Bank、Bank = (地址/4) % 32、列访问触发 32 路冲突」的因果链，并用 +1 padding 一招化解；④ 能讲清从「逐线程搬运 → cp.async 异步双缓冲 → Hopper TMA 单线程发起批量搬运」的演进主线，理解每代硬件都在把「搬数据」从计算线程上卸载；⑤ 亲手实现 naive 与 tiled 两版 GEMM 并实测加速比，把「访存优化」从概念变成手里的真实数字。
• 阅读姿势：盯住一条主线——「优化 CUDA kernel 的全部功夫，都花在『让数据离计算单元更近、并被尽可能多的线程复用』上」。无论是 Shared Memory tiling、bank conflict 规避，还是 cp.async/TMA 异步流水，本质都在解决同一堵墙：算力增长远快于带宽，绝大多数 kernel 卡在「等数据」而非「算不过来」。带着这把尺子去读每一节的图、表与代码。

背景与现状

CUDA（Compute Unified Device Architecture） 是 NVIDIA 提供的并行编程模型，它的核心思想只有一句话：用海量轻量级线程掩盖内存延迟。CPU 靠几个强核 + 巨大缓存来「躲」延迟，GPU 反其道而行——它有成千上万个线程，当一批线程在等显存数据时，调度器立刻切换到另一批就绪线程继续算，用并发把访存延迟「藏」在计算背后。

理解这一点，才能理解 CUDA 编程的全部矛盾来源：算力（FLOPS）增长远快于带宽（GB/s）增长。一块 H100 的 FP16 算力接近 1000 TFLOPS，而 HBM3 带宽约 3.35 TB/s——算力与带宽之比（即 计算强度 / arithmetic intensity 的临界值）高达数百 FLOP/Byte。这意味着：绝大多数朴素 kernel 都卡在「等数据」上，而非「算不过来」。优化 CUDA kernel，本质上就是在做一件事：减少对慢速全局显存（HBM）的访问次数，最大化对快速片上存储（Shared Memory / 寄存器）的数据复用。

从业界视角看，CUDA 编程能力的演进可以概括为三个阶段：

• 手写 kernel 时代（2007–2016）：开发者直接用 CUDA C 写 kernel，性能强依赖个人对硬件的理解。
• 模板库时代（2017–2021）：CUTLASS（CUDA Templates for Linear Algebra Subroutines）把 tiling、双缓冲等模式抽象成 C++ 模板，cuBLAS 在内部大量复用，普通人难以手写超越。
• DSL / 编译器时代（2021 至今）：Triton、MLIR 等让算法工程师用 Python 级抽象写出接近手写性能的 kernel——但它们生成的代码本质仍是本文讲的这套 tiling + shared memory + 双缓冲范式。

业界信号：FlashAttention 之所以能颠覆 Attention 计算，核心不是新算法，而是把 Softmax 重组成可以在 Shared Memory 上分块计算（tiling）的形式，从而避免把巨大的 N×N 注意力矩阵写回 HBM。这再次印证：高性能算子的胜负手是「访存」，不是「算法」。本文讲透的 tiled GEMM，就是理解 FlashAttention 的前置必修课。

原理与架构

要写出快的 kernel，必须把软件抽象（Thread/Block/Grid）与硬件实体（SM / Shared Memory / 寄存器）一一对应起来。

2.1 Thread / Block / Grid：软件层次到 SM 的映射

CUDA 用三级层次组织线程，每一级都精确映射到一个硬件概念：

关键映射规则（决定你的 kernel 配置）：

软件概念	硬件实体	工程含义
Grid	整块 GPU 上的所有 SM	gridDim 应 ≥ SM 数量，否则部分 SM 闲置
Block	整块驻留在单个 SM 上（不会跨 SM）	同 Block 内才能用 Shared Memory 协作与 __syncthreads() 同步
Warp（32 线程）	Warp Scheduler 的调度单位	Block 大小应为 32 的倍数；同 warp 内分支发散（branch divergence）会串行化
Thread	占用寄存器	寄存器用量过高会降低 occupancy（占用率），削弱延迟掩盖能力

核心洞察：一个 SM 能同时驻留多个 Block，前提是它们的「寄存器 + Shared Memory」总需求不超过 SM 的物理上限。Occupancy = 活跃 warp 数 / SM 理论最大 warp 数——它衡量「调度器手里有多少就绪 warp 可用来掩盖延迟」。这就是为什么 Shared Memory 和寄存器用得越省，往往能驻留越多 Block、藏住越多延迟。

2.2 Shared Memory：程序员显式管理的 L1 缓存

Shared Memory 是位于 SM 片上、由程序员显式分配和管理的高速存储（延迟约为 HBM 的 1/20～1/30）。它和 L1 Cache 共享同一块物理 SRAM，区别在于：L1 由硬件自动管理，Shared Memory 由你用 __shared__ 关键字手动控制。

它的杀手级用途是 数据复用：把一块数据从 HBM 搬进 Shared Memory 一次，让 Block 内成百个线程反复读取它，从而把对 HBM 的访问次数降低一个数量级。tiled GEMM（见动手实践一节）就是这个思想的教科书级应用。

2.3 Bank Conflict：32 个 Bank 与串行化惩罚

Shared Memory 在物理上被切分为 32 个 Bank（恰好对应一个 warp 的 32 个线程），每个 Bank 每周期只能服务一次访问。理想情况：一个 warp 的 32 个线程访问 32 个不同 Bank → 一周期并行完成。Bank Conflict：若多个线程访问同一个 Bank 的不同地址，硬件被迫串行化这些访问，N 路冲突 = N 倍延迟。

经典冲突场景：在 GEMM 中按列访问一个 [32][32] 的 Shared Memory 数组。地址 tile[threadIdx.x][k] 当 threadIdx.x 在 warp 内变化时，相邻线程访问的元素相隔 32 个 float（即 32 * 4 = 128 字节）——由于 Bank = (地址/4) % 32，所有线程恰好落到同一个 Bank，触发 32 路冲突。

免除技巧——Padding（补齐）：把数组声明从 __shared__ float tile[32][32]; 改为 __shared__ float tile[32][33];（多补一列）。此时列内相邻元素的地址相隔 33 个 float，33 % 32 = 1，强制错开到不同 Bank，冲突消失。代价仅是每行多 4 字节的 Shared Memory 占用，性价比极高。

2.4 异步内存拷贝与 Double Buffering 流水线

朴素 tiled kernel 的执行节奏是「串行」的：搬一块数据 → 同步 → 计算 → 搬下一块……计算单元在搬数据时是闲置的。Double Buffering（双缓冲） 用两块 Shared Memory 缓冲区交替，让「计算当前块」与「预取下一块」并行重叠，从而掩盖访存延迟。

从 Ampere（SM80）起，CUDA 提供 cp.async（异步拷贝） 指令：它能直接把数据从 HBM 搬进 Shared Memory 而不经过寄存器，且不阻塞计算线程——发起拷贝后线程继续算，到需要数据时再 cp.async.wait_group 同步。这是实现高效双缓冲的硬件基石。

2.5 先进架构硬核指令：异步管道与 Hopper TMA

• Ampere 异步管道（cuda::pipeline）：CUDA 11 把 cp.async 封装为标准库 cuda::pipeline，用 producer_acquire / commit / consumer_wait 显式描述多级流水线，把双缓冲推广到**多级（multi-stage）**预取。
• Hopper TMA（Tensor Memory Accelerator）：H100（SM90）引入的专用硬件单元。此前 cp.async 仍需每个线程参与地址计算与发起；TMA 让单个线程发起一次描述符（descriptor）调用，由专用引擎批量异步搬运一整块多维 tile（含自动地址计算与边界处理），彻底把搬运逻辑从计算线程剥离。配合 TBC（Thread Block Cluster） 与分布式 Shared Memory，让多个 Block 的 Shared Memory 互相直接访问——这是新一代 GEMM/Attention kernel（如 FlashAttention-3、cuBLAS Hopper 路径）极致性能的来源。

演进主线一句话：手动逐线程搬运 →（Ampere）cp.async 异步、绕过寄存器 →（Hopper）TMA 单线程发起、专用引擎批量多维搬运。每一代硬件都在把「搬数据」这件苦活从宝贵的计算线程上卸载下去。

动手实践：手写 tiled GEMM 对比 naive

实验目标：亲手实现 naive GEMM 与 tiled GEMM（Shared Memory 分块） 两个 kernel，在同一矩阵规模下测速，实测加速比，并解释加速来源 = 「HBM 访问次数从 O(N) 降到 O(N/TILE)」。产出物：一份两版 kernel 的 GFLOPS 对比与加速比数字。

3.1 路径选择

• 路径 A — NVIDIA GPU（推荐）：本地有 NVIDIA 显卡 + 已装 CUDA Toolkit（含 nvcc），或用 Google Colab 免费 T4 GPU / RunPod / Lambda 租 GPU。这是唯一能跑真 CUDA kernel 的路径。
• 路径 B — Mac / 无 GPU：无法运行 .cu，改用 NumPy 分块矩阵乘法演示 tiling 的数据复用原理（见 3.5），并按 3.6 在 Colab 上跑真 GPU 版。

3.2 完整 CUDA 代码 gemm.cu

#include <cstdio>
#include <cuda_runtime.h>

#define TILE 16          // 每个 Block 处理 16x16 的输出 tile
#define N    1024         // 方阵边长（C = A * B，均为 N x N）

// ---------- Kernel 1：Naive GEMM（每个输出元素直读全局显存）----------
__global__ void gemm_naive(const float* A, const float* B, float* C, int n) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < n && col < n) {
        float acc = 0.0f;
        for (int k = 0; k < n; ++k)
            acc += A[row * n + k] * B[k * n + col];  // 每次乘加都读两次 HBM
        C[row * n + col] = acc;
    }
}

// ---------- Kernel 2：Tiled GEMM（Shared Memory 分块复用）----------
// 注意 +1 列 padding 以免除 bank conflict
__global__ void gemm_tiled(const float* A, const float* B, float* C, int n) {
    __shared__ float As[TILE][TILE + 1];
    __shared__ float Bs[TILE][TILE + 1];

    int ty = threadIdx.y, tx = threadIdx.x;
    int row = blockIdx.y * TILE + ty;
    int col = blockIdx.x * TILE + tx;
    float acc = 0.0f;

    // 沿 K 维分块：每次把 A、B 的一个 tile 搬进 Shared Memory 复用
    for (int t = 0; t < n / TILE; ++t) {
        As[ty][tx] = A[row * n + (t * TILE + tx)];
        Bs[ty][tx] = B[(t * TILE + ty) * n + col];
        __syncthreads();                       // 等全 Block 搬完

        #pragma unroll
        for (int k = 0; k < TILE; ++k)         // 这 TILE 次乘加全部命中 Shared Memory
            acc += As[ty][k] * Bs[k][tx];
        __syncthreads();                       // 等全 Block 算完再搬下一块
    }
    if (row < n && col < n) C[row * n + col] = acc;
}

// ---------- 计时辅助 ----------
float time_kernel(void (*launch)(const float*, const float*, float*, int),
                  const float* dA, const float* dB, float* dC, int n,
                  dim3 grid, dim3 block) {
    cudaEvent_t s, e; cudaEventCreate(&s); cudaEventCreate(&e);
    // 预热
    launch<<<grid, block>>>(dA, dB, dC, n); cudaDeviceSynchronize();
    cudaEventRecord(s);
    for (int i = 0; i < 20; ++i) launch<<<grid, block>>>(dA, dB, dC, n);
    cudaEventRecord(e); cudaEventSynchronize(e);
    float ms = 0; cudaEventElapsedTime(&ms, s, e);
    cudaEventDestroy(s); cudaEventDestroy(e);
    return ms / 20.0f;   // 单次平均毫秒
}

int main() {
    size_t bytes = (size_t)N * N * sizeof(float);
    float *hA = (float*)malloc(bytes), *hB = (float*)malloc(bytes);
    for (int i = 0; i < N * N; ++i) { hA[i] = 1.0f; hB[i] = 2.0f; }

    float *dA, *dB, *dC;
    cudaMalloc(&dA, bytes); cudaMalloc(&dB, bytes); cudaMalloc(&dC, bytes);
    cudaMemcpy(dA, hA, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(dB, hB, bytes, cudaMemcpyHostToDevice);

    dim3 block(TILE, TILE);
    dim3 grid(N / TILE, N / TILE);

    float t_naive = time_kernel(gemm_naive, dA, dB, dC, N, grid, block);
    float t_tiled = time_kernel(gemm_tiled, dA, dB, dC, N, grid, block);

    double gflop = 2.0 * N * N * N / 1e9;   // 2*N^3 次浮点运算
    printf("Naive : %.3f ms  | %.1f GFLOPS\n", t_naive, gflop / (t_naive / 1e3));
    printf("Tiled : %.3f ms  | %.1f GFLOPS\n", t_tiled, gflop / (t_tiled / 1e3));
    printf("Speedup (naive/tiled) = %.2fx\n", t_naive / t_tiled);

    free(hA); free(hB); cudaFree(dA); cudaFree(dB); cudaFree(dC);
    return 0;
}

3.3 编译与运行（NVIDIA GPU 路径）

# 查看 GPU 与 nvcc 是否就绪
nvidia-smi
nvcc --version

# 编译：-arch 按你的卡填（T4=sm_75, A100=sm_80, H100=sm_90）
nvcc -O3 -arch=sm_75 gemm.cu -o gemm

# 运行
./gemm

典型输出（数值随卡而异，T4 上约）：

Naive : 6.842 ms  | 313.8 GFLOPS
Tiled : 1.271 ms  | 1689.6 GFLOPS
Speedup (naive/tiled) = 5.38x

3.4 加速比从哪来？（核心解释）

朴素版每个输出元素都从 HBM 读 2N 个 float，全 kernel 总 HBM 读 ≈ 2N³。tiled 版每个 tile 数据被 Block 内 TILE 个线程复用，HBM 访问次数降到约 2N³ / TILE——TILE=16 时理论上访存量降到 1/16。实测加速比通常是 4～10 倍（受访存非全瓶颈、occupancy 等影响），这就是「数据复用」带来的真金白银。想进一步逼近 cuBLAS，下一步是寄存器分块（每线程算多个输出）+ 双缓冲（cp.async）。

3.5 无 GPU 替代：NumPy 分块演示 tiling 复用原理

无 GPU 时，用 NumPy 分块矩阵乘法演示「同一块数据被多次复用」的核心思想（CPU 上不会有 GPU 那样的加速，目的是理解原理而非测速）：

import numpy as np

N, TILE = 512, 64
A = np.ones((N, N), dtype=np.float32)
B = np.full((N, N), 2.0, dtype=np.float32)
C = np.zeros((N, N), dtype=np.float32)

# 模拟 GPU tiled GEMM 的三层分块循环：搬一块 -> 复用一块
for i0 in range(0, N, TILE):
    for j0 in range(0, N, TILE):
        acc = np.zeros((TILE, TILE), dtype=np.float32)
        for k0 in range(0, N, TILE):
            a_tile = A[i0:i0+TILE, k0:k0+TILE]   # 对应搬进 Shared Memory 的 As
            b_tile = B[k0:k0+TILE, j0:j0+TILE]   # 对应 Bs
            acc += a_tile @ b_tile               # 这一块数据被 TILE^2 次乘加复用
        C[i0:i0+TILE, j0:j0+TILE] = acc

assert np.allclose(C, A @ B)
print("分块结果与直接 matmul 一致；a_tile/b_tile 即对应 GPU 的 Shared Memory 复用单元")

python tiled_numpy.py

对应关系：a_tile / b_tile ↔ Shared Memory 的 As / Bs；最内层 k 循环 ↔ kernel 里命中 Shared Memory 的乘加；「搬一块、复用一块」的结构与 CUDA 完全同构。

3.6 在 Google Colab 免费 GPU 上跑真 CUDA 版

# Colab 新建 notebook → 菜单「代码执行程序 / Runtime」→「更改运行时类型」→ 选 GPU(T4)
# 在 cell 中执行（! 开头为 shell 命令）：
!nvidia-smi
%%writefile gemm.cu
# ……把 3.2 的代码粘进来……
!nvcc -O3 -arch=sm_75 gemm.cu -o gemm && ./gemm

踩坑预警 (Gotchas)

• 不 cudaDeviceSynchronize 就计时 = 测了个寂寞：kernel 异步下发，必须用 CUDA Event 包裹并 cudaEventSynchronize 才测到真耗时；首次调用含上下文初始化，务必预热。
• -arch 填错跑不起来或退化：架构号要匹配你的卡（T4=sm_75 / A100=sm_80 / H100=sm_90）。填低了用不上新指令，填高了可能直接报 no kernel image is available。
• __syncthreads() 放在分支里会死锁：它要求 Block 内所有线程都到达；写在 if (row<n) 内部、部分线程跳过，会 hang。同步必须在所有线程都执行的路径上。
• Bank Conflict 静默吃性能：去掉 Shared Memory 数组的 +1 padding，kernel 仍能跑出正确结果，但慢一截且无报错。用 ncu --metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum ./gemm 才能看见。
• N 不是 TILE 整数倍会越界：示例为简化假设 N % TILE == 0；通用代码须对 tile 边界做 if 判断并对越界元素填 0。
• gflop 用 int 算 2*N*N*N 溢出：N=1024 时 2*N³ ≈ 2.1e9 超过 int32，务必用 double（代码中已用 2.0）。

深入思考

下面三题每题先给题干，再用 <details> 折叠一份图文并茂的参考答案。建议先合上答案自己想 3 分钟，再展开对照。

思考题 1：复用上限的代价权衡

把 TILE 从 16 增大到 32，理论上每块数据被复用更多次、HBM 访存量再降一半，但实测加速比往往不升反降。结合 2.1 节 occupancy 机制（Shared Memory / 寄存器占用 → 一个 SM 能驻留几个 Block），解释「为什么『分块越大越好』是个陷阱」，并说明你会用什么实验流程找到最优 TILE。

展开参考答案（含 TILE-occupancy 权衡图 + 算一遍）

结论：TILE 同时控制两件事——「数据复用率」与「每 Block 的 Shared Memory 占用」。前者随 TILE 增大而改善（访存量降），后者却随 TILE² 平方膨胀，挤占 SM 上有限的 Shared Memory，使能并发驻留的 Block 数骤减、occupancy 跌破能掩盖延迟的临界点。当「并行度损失」压过「访存量收益」，加速比就掉头向下。最优 TILE 是这条凹曲线的顶点，只能实测出来。

用具体数字算一遍（数量级估算，假设 SM 可用 Shared Memory 约 100 KB，FP32）：

1. tiled kernel 每 Block 占两块 tile：As + Bs ≈ 2 × TILE² × 4 字节。
2. TILE=16：每 Block ≈ 2 × 256 × 4 = 2 KB → 单 SM 理论可驻留约 100 / 2 = 50 个 Block（实际还受 warp/寄存器上限约束，但 Shared Memory 远不是瓶颈）→ occupancy 充足。
3. TILE=32：每 Block ≈ 2 × 1024 × 4 = 8 KB → 单 SM 仅可驻留约 100 / 8 ≈ 12 个 Block，且每 Block 线程数从 16×16=256 升到 32×32=1024，触顶「每 SM 最多 1024～2048 线程」的硬上限，常驻 Block 进一步被压到 1～2 个。
4. 后果：访存量虽降一半，但常驻 warp 数大跌；一旦不足以填满 HBM 的数百周期延迟，SM 就开始空转等数据——这正是 2.1 节「Shared Memory 用得越省，越能藏住延迟」的反面教材。

怎么找最优 TILE：把 TILE 做成编译期参数，对 {8, 16, 32} 各编译一版，在目标卡上实测 GFLOPS 扫描；同时用 ncu 看 Achieved Occupancy 与 Shared Memory Per Block，找到「occupancy 尚未塌、访存量已显著下降」的拐点。通常 TILE=16/32 在多数卡上接近甜点，但唯一可信的是你卡上的实测曲线，而非纸面访存量——这与 3.4 节「实测加速比 4～10 倍受 occupancy 等影响」完全呼应。

思考题 2：从 tiled GEMM 到 FlashAttention

FlashAttention 把 Softmax-Attention 重写成可以在 Shared Memory 上分块计算、从不把 N×N 注意力矩阵写回 HBM 的形式。请用本文「tiled GEMM 把 HBM 访问降到 1/TILE」（2.2 节 + 3.4 节）的思路，解释 FlashAttention 的加速本质，并指出它比普通 tiled GEMM 多解决了什么难题（提示：Softmax 的全局归约 / online softmax）。

展开参考答案（含 FlashAttention 分块数据流图 + 对比表）

结论：FlashAttention 的加速和 tiled GEMM 同根同源——都是「把数据搬进 Shared Memory 复用、避免对 HBM 的反复读写」。但它比 tiled GEMM 难一个数量级：GEMM 的每个输出块只是简单累加，而 Attention 中间要做 Softmax，Softmax 需要对整行做全局归约（求 max、求和），与「只看一个 tile 的局部」天然冲突。FlashAttention 用 online softmax 把全局归约改写成「逐块增量更新」，从而在不物化 N×N 矩阵的前提下算出正确结果。

和 tiled GEMM 的同与异（对比表）：

维度	tiled GEMM	FlashAttention
核心手法	tile 搬进 Shared Memory，Block 内复用	同左：Q/K/V 分块进 Shared Memory，流式复用
HBM 访存收益	降到约 1/TILE	不物化 N×N 分数矩阵，访存量降一个数量级
中间运算	纯乘加累加，块间独立、直接相加	含 Softmax，需对整行做全局 max/sum 归约
多解决的难题	——	online softmax：把全局归约改写成逐块增量更新

多解决的那个难题，具体是什么：Softmax 要 exp(xᵢ − max) / Σexp(·)，分母依赖整行所有元素——而分块计算时你一次只看到一个 K 块，看不到全行。online softmax 的巧思是：每处理一个新 K 块，就更新「running max m」与「running sum l」，并对已经累积的输出 O 乘一个 rescale 因子做校正，使得「分块逐步累加」的结果与「先看全行再算」数学完全等价。这样既享受了 tiled GEMM 的访存红利，又绕开了「Softmax 需要全局视野」的死结——这正是 3.4 节结尾「想进一步逼近，下一步是寄存器分块 + 双缓冲」之上，把同一套访存哲学推向 SOTA 的范例（详见 Dao-AILab/flash-attention 的 csrc/）。

思考题 3：硬件代际迁移的红利不是免费的

你的 tiled GEMM 在 Ampere（A100）上已接近访存瓶颈。迁到 Hopper（H100）后，仅重编译 -arch=sm_90 几乎没有提升。结合 2.4 节（cp.async / 双缓冲）与 2.5 节（TMA / Thread Block Cluster），说明要吃到 H100 的红利，kernel 的「搬运」与「协作」逻辑需要怎样重写？为什么说「新硬件的算力红利不是免费的」？

展开参考答案（含搬运逻辑代际重写图 + 对比表）

结论：H100 的红利藏在「专用搬运引擎 TMA」与「跨 Block 协作 TBC」里，而这两者都是新的硬件指令/单元——旧 kernel 的「每个线程自己算地址、自己发起 cp.async」既用不上 TMA 的批量异步搬运，也享受不到 Block 间分布式 Shared Memory。编译器无法自动把「逐线程搬运范式」改写成「单线程发描述符 + 专用引擎搬运」范式，所以只重编译几乎无效——你必须按新指令重写搬运与协作逻辑。这就是「算力红利不是免费的」。

搬运与协作要怎么重写（对比表）：

环节	Ampere 旧写法	Hopper 重写后
谁算地址	每个线程自己算 tile 内偏移	单线程填一次 TMA 描述符，引擎自动算地址与边界
搬运指令	逐线程 cp.async 发起	cuda::memcpy_async + TMA，专用引擎批量搬多维 tile
流水深度	双缓冲（2 级）	多级 cuda::pipeline（multi-stage）深度预取
协作范围	仅 Block 内 Shared Memory	TBC 跨 Block 分布式 Shared Memory，互相直接读写
计算线程负担	仍要兼顾搬运	搬运彻底剥离，计算线程专注 MMA

为什么红利不免费：编译器能做指令选择、寄存器分配，但改不了你的算法结构与数据搬运范式。-arch=sm_90 只是允许生成 SM90 指令，它不会把你「逐线程 cp.async + Block 内双缓冲」的代码自动重组成「TMA 描述符 + 多级流水 + TBC 协作」——这是程序员（或 CUTLASS/Triton 这类库）才能做的结构性改写。这正呼应 2.5 节那句演进主线：每代硬件都在把「搬数据」从计算线程上卸载，但要享受这份卸载，必须主动改用新的搬运/协作原语。所以工业界宁可依赖 cuBLAS/CUTLASS——把「贴着每代硬件重写 kernel」这件苦活交给专门维护的库，而不是每换一代卡就手写一遍（这也是 L2.4 Triton 试图用编译器自动化的目标）。

延伸阅读

1. 核心 Paper / 官方文档

• CUDA C++ Programming Guide（NVIDIA 官方）— 必读 §Memory Hierarchy、§Shared Memory、§Asynchronous Data Copies 三章。
• FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness（2022）— 理解「IO-aware」如何把本文的 tiling 思想推向 SOTA。
• CUTLASS 官方文档与 GTC「CUTLASS: Python API, Enhancements, and More」 — 工业级 tiling / 多级双缓冲的范式参考。

2. 相关高 Star 仓库与源码必读路径

• NVIDIA/cutlass — 看 include/cutlass/gemm/ 下的分层 tiling 与 pipeline 抽象，是手写 GEMM 的工业级天花板。
• Dao-AILab/flash-attention — 看 csrc/ 下 kernel 如何用 Shared Memory tiling + online softmax 避免 HBM 物化。
• siboehm/SGEMM_CUDA — 一步步从 naive 优化到接近 cuBLAS 的经典教学仓库，与本文实验完美衔接。

3. 优质博客 / 视频 / 课程

• Simon Boehm「How to Optimize a CUDA Matmul Kernel for cuBLAS-like Performance」博客 — 本文实验的进阶续集，逐版优化讲到寄存器分块与双缓冲。
• NVIDIA GTC「CUDA Programming Model for Hopper」/「Programming the Hopper TMA」技术讲座。
• 课程：UIUC ECE408 / Stanford CS149（Parallel Computing）建立 SIMT 与访存优化的系统观。

---

下一篇 → L2.4 Triton 编译器源码级实战：当你已亲手体会到手写 CUDA 的繁琐与威力，下一步看 Triton 如何用 Python 级抽象，自动生成本文这套 tiling + 双缓冲范式的 kernel。