CUTLASS / CuTe 主线
8 课现代 GPU 模板库与 Layout 代数
CUTLASS / CuTe 核心主线
本主线深入 NVIDIA CUTLASS 3.x 与核心抽象 CuTe，学习 GPU 硬件层级演进、Tensor Core 算术强度、CuTe Layout 代数与分层 GEMM 实现。
学习目标
理解 Volta → Ampere → Hopper → Blackwell 的 Tensor Core 演进；
掌握 CuTe Layout 核心代数：Shape、Stride、坐标变换与切片；
掌握 Tiled Copy 与 Tiled MMA 的构建与使用；
理解异步数据加载（
cp.async）、Shared Memory 多级流水与 Epilogue 融合。
课程目录
| 课程名称 | 模块归属 |
|---|---|
| 第 1 课：GPU 层级、Tensor Core 与算术强度 | CUTLASS / CuTe 主线 |
| 第 2 课：CuTe Layout：Shape、Stride 与坐标映射 | CUTLASS / CuTe 主线 |
| 第 3 课：CuTe Tensor、视图与线程分区 | CUTLASS / CuTe 主线 |
| 第 4 课：Tiled Copy、向量化与 Predication | CUTLASS / CuTe 主线 |
| 第 5 课：MMA Atom、Tiled MMA 与三级 Tiling | CUTLASS / CuTe 主线 |
| 第 6 课：Shared Memory 布局与异步流水 | CUTLASS / CuTe 主线 |
| 第 7 课：混合精度、Accumulator 与 Epilogue | CUTLASS / CuTe 主线 |
| 第 8 课：Kernel 组合、Dispatch 与 Profiling | CUTLASS / CuTe 主线 |
核心架构图解
CUTLASS 教程：从手写 GEMM 到 CuTe / Tensor Core Kernel
学习目标：看懂 CUTLASS 的层级抽象，能把 ThreadblockShape / WarpShape / InstructionShape 和数据搬运、pipeline、epilogue 对上
一、先把 GEMM 三重循环映射到 GPU 层级
数学 GEMM
C[M,N] = A[M,K] x B[K,N]
C[M,N] = A[M,K] x B[K,N]
CTA / Threadblock Tile
ThreadblockShape{kM,kN,kK}
一个 block 负责一个 C tile
ThreadblockShape{kM,kN,kK}
一个 block 负责一个 C tile
Warp Tile
多个 warp 分摊 CTA tile
从 smem 读到寄存器
多个 warp 分摊 CTA tile
从 smem 读到寄存器
MMA Atom
mma.sync / wgmma
Tensor Core 指令
mma.sync / wgmma
Tensor Core 指令
Epilogue
alpha/beta
bias/activation/store
alpha/beta
bias/activation/store
二、CUTLASS 2.x / 传统 GEMM API 怎么读
device::Gemm
host 侧入口，语义接近 cuBLAS
host 侧入口，语义接近 cuBLAS
threadblock::MmaPipelined
global -> shared -> register
K 维 mainloop
global -> shared -> register
K 维 mainloop
warp::Mma
warp 级 outer-product 累加
warp 级 outer-product 累加
Iterator / Layout
描述 tile 访问、边界 predication、访存布局
描述 tile 访问、边界 predication、访存布局
读法：先定位 device API，再追 mainloop，最后看 iterator 如何把矩阵 tile 化。面试重点是 tile 层级、共享内存复用、bank conflict、occupancy 与 pipeline stage。
三、CUTLASS 3.x / CuTe 怎么理解
Collective Mainloop
GmemTiledCopy + SmemLayout
TiledMma / pipeline
GmemTiledCopy + SmemLayout
TiledMma / pipeline
Collective Epilogue
accumulator -> output tile
融合 scale/bias/activation
accumulator -> output tile
融合 scale/bias/activation
CuTe Layout Algebra
Shape / Stride / Layout
把索引映射写成类型
Shape / Stride / Layout
把索引映射写成类型
Atom / Tiled MMA Copy
MMA 原子操作 + 分块组合
Hopper/Blackwell 可接 TMA/WGMMA
MMA 原子操作 + 分块组合
Hopper/Blackwell 可接 TMA/WGMMA
读法：CuTe 的核心不是新数学，而是用 Layout/Tile/Atom 精确表达数据在哪、怎么搬、由哪个 MMA 原子消费。先学 layout，再学 tensor，再看 GEMM tutorial。
四、面试可讲的优化主线
Global Memory
coalesced load
A/B tile 从 gmem 搬入 smem
coalesced load
A/B tile 从 gmem 搬入 smem
Shared Memory
数据复用
bank conflict free layout
数据复用
bank conflict free layout
Register Accumulator
每线程/warp 保持 C fragment
提升 ILP
每线程/warp 保持 C fragment
提升 ILP
Pipeline Stage
cp.async / TMA 预取
隐藏 gmem latency
cp.async / TMA 预取
隐藏 gmem latency
Epilogue Fusion
store 前融合
减少额外 kernel
store 前融合
减少额外 kernel
必会问题 1
ThreadblockShape 变大为什么可能更快，也可能 occupancy 下降
ThreadblockShape 变大为什么可能更快，也可能 occupancy 下降
必会问题 2
Warp tile / instruction tile 分别对应什么硬件层级
Warp tile / instruction tile 分别对应什么硬件层级
必会问题 3
为什么需要 epilogue，bias/activation 放哪里融合
为什么需要 epilogue，bias/activation 放哪里融合
练习路线
naive GEMM -> tiled GEMM -> WMMA -> CUTLASS profiler -> CuTe layout
naive GEMM -> tiled GEMM -> WMMA -> CUTLASS profiler -> CuTe layout
CUTLASS 教程：GEMM 到 CuTe 原理 · Excalidraw 源文件
ldmatrix 指令与 CuTe TV Layout 映射
核心：搞清 (Thread, Value) → (Row, Col) 的映射，以及 ldmatrix 为什么能把 SMEM 数据直接放成 mma.sync 期望的寄存器格式。
Part 1：ldmatrix.m8n8.x1，一个 8×8 fp16 矩阵
谁提供地址
T0–T7 各提供一行的 SMEM 地址（每行 16 B = 8 个 fp16）。
T8–T31 的地址在 .x1 中不使用，但 32 个线程都必须执行这条
.sync.aligned 指令。谁拿到数据
row = lane / 4
col_start = (lane % 4) * 2
reg (u32) = { M[row][col_start],
M[row][col_start+1] }
每 4 个线程分一行，每线程拿 2 个相邻 fp16。
| 线程组 | lane%4 = 0 | 1 | 2 | 3 |
|---|---|---|---|---|
| T0–T3 | T0: (0,0) (0,1) | T1: (0,2) (0,3) | T2: (0,4) (0,5) | T3: (0,6) (0,7) |
| T4–T7 | T4: (1,0) (1,1) | T5: (1,2) (1,3) | T6: (1,4) (1,5) | T7: (1,6) (1,7) |
| T8–T11 | T8: (2,0) (2,1) | T9: (2,2) (2,3) | T10: (2,4) (2,5) | T11: (2,6) (2,7) |
| … | … | … | … | … |
| T28–T31 | T28: (7,0) (7,1) | T29: (7,2) (7,3) | T30: (7,4) (7,5) | T31: (7,6) (7,7) |
Part 2：ldmatrix.x4，一次加载 16×16 的 A 操作数
| 矩阵 | 地址来自 | 覆盖区域 | 写入寄存器 | 对应 mma A 片段 |
|---|---|---|---|---|
| Mat 0 | T0–T7 | M 0–7，K 0–7 | d0 | a0, a1 |
| Mat 1 | T8–T15 | M 8–15，K 0–7 | d1 | a2, a3 |
| Mat 2 | T16–T23 | M 0–7，K 8–15 | d2 | a4, a5 |
| Mat 3 | T24–T31 | M 8–15，K 8–15 | d3 | a6, a7 |
地址按这个顺序给，d0–d3 就正好是 mma.m16n8k16 的 A 片段顺序。线程 t 的 4 个寄存器：
g = t / 4, c = (t % 4) * 2
d0 = {A[g][c], A[g][c+1]} d1 = {A[g+8][c], A[g+8][c+1]}
d2 = {A[g][c+8], A[g][c+9]} d3 = {A[g+8][c+8], A[g+8][c+9]}
例 t = 5: g = 1, c = 2 → d0 = A[1][2..3], d1 = A[9][2..3],
d2 = A[1][10..11], d3 = A[9][10..11]
Part 3：CuTe 里的 SM80_16x8x16 ALayout
ALayout: (T32, V8) → (M16, K16)，按列主序 idx = m + 16·k
Shape <Shape <_4, _8>, Shape <_2, _2, _2>>
Stride<Stride<_32,_1>, Stride<_16,_8,_128>>
| 模式 | 大小 | stride | 含义 |
|---|---|---|---|
| 线程 t % 4 | 4 | 32 | K 方向 +2 列 |
| 线程 t / 4 | 8 | 1 | M 方向 +1 行 |
| 值 v0 | 2 | 16 | K 方向 +1 列（同一寄存器里的两个 fp16） |
| 值 v1 | 2 | 8 | M 方向 +8 行（下半块） |
| 值 v2 | 2 | 128 | K 方向 +8 列（右半块） |
Thread 0 持有 (0,0) (0,1) | (8,0) (8,1) | (0,8) (0,9) | (8,8) (8,9)，与 Part 2 的 d0–d3 一致。
make_tiled_copy_A(Copy_Atom<LDSM>{}, tiled_mma) 会按 MMA 的 ALayout 反推每个线程该拷哪些元素，所以 ldmatrix 的输出可以直接喂给 mma.sync。K-major 的 A 用 SM75_U32x4_LDSM_N（ldmatrix.x4）；M-major 的 tile 才用 SM75_U16x8_LDSM_T（.x4.trans）。Part 4：全流程 GMEM → SMEM → Register → MMA
Global Memory
A[M×K]、B[N×K]
cp.async：
SM80_CP_ASYNC_CACHEGLOBAL<uint128_t>Shared Memory
SmemLayoutAtom +
Swizzle<3,3,3>把行号 XOR 进 16 B 块索引，ldmatrix 读的 8 行落在不同 bank，无冲突。
Register
ldmatrix 写成 MMA 格式：A 每线程 4 个 u32（8 fp16），B 每线程 2 个 u32（4 fp16）。
Accumulator
mma.sync.m16n8k16：C/D 每线程 4 个 f32，(T32, V4) → (M16, N8)。
using MMA = TiledMMA<MMA_Atom<SM80_16x8x16_F32F16F16F32_TN>>;
auto copy_A = make_tiled_copy_A(Copy_Atom<SM75_U32x4_LDSM_N, half_t>{}, mma);
auto copy_B = make_tiled_copy_B(Copy_Atom<SM75_U32x4_LDSM_N, half_t>{}, mma);
cute::copy(copy_A, tCsA, tCrA_view); // 发射 ldmatrix
cute::gemm(mma, tCrA, tCrB, tCrC); // 发射 mma.sync
ldmatrix TV Layout 详解 · Excalidraw 源文件