CuTe C++ 学习路线

基于当前仓库提交 0b55a2f6（CUTLASS 4.8），检查日期：2026-10-21。

范围按官方 CuTe Quickstart 的定义：examples/cute 下全部 16 个 .cu 示例，结合 C++ 文档、CMake 配置和相关 C++ 单元测试。Python DSL 不在本路线内。其他编号的 CUTLASS 示例只作为后续应用入口，不计入这 16 个。

建议主线：Layout → Tensor → Copy → Predication → 普通 GEMM → 流水线 → Tensor Core → 按架构进阶。

先假设你已经会写基本 CUDA kernel，并能阅读 C++ 模板、auto 和 constexpr。基础主线安排 4–5 周，每周约 8–12 小时；Hopper、SM100、Rubin 再按目标选择约 3–5 周。是否进入下一阶段，以能完成练习为准。

全部示例与推荐顺序

下面的编号是学习顺序，不是仓库目录顺序。sgemm_sm80.cu 分两遍阅读，但文件只计一次。

顺序	示例	主要新增概念	练习产出
1	tiled_copy.cu	tiled_divide、local_partition、线程布局/值布局、Copy_Atom、TiledCopy	说明线程 0、1、31、32 分别复制哪些元素
2	tiled_copy_if.cu	identity tensor、坐标谓词、copy_if	完成不能被 tile 整除的矩阵复制
3	sgemm_1.cu	CTA 分块、copy/compute 的不同线程划分、SMEM、寄存器累加、epilogue	从零复写一个对齐尺寸的普通 FMA GEMM
4	sgemm_2.cu	TiledCopy/TiledMMA 分工、UniversalFMA、GMEM 预取到寄存器、SMEM padding	与 sgemm_1 对照，逐项解释新增抽象
5	sgemm_sm70.cu	GMEM→RMEM→SMEM 与 SMEM→RMEM→FMA 两层软件流水	画出两个 K tile 和内部 K block 的时序
6	sgemm_sm80.cu	cp.async、多 stage SMEM；half 分支再引入 mma.sync、LDSM、swizzle、retile	先解释异步流水线，再追踪一个 half 元素进入 MMA 的全过程
7	hopper/wgmma_sm90.cu	保留 cp.async，引入 warpgroup MMA、SMEM descriptor、异步 MMA 等待	对照 SM80，指出操作数和 accumulator 的位置
8	hopper/wgmma_tma_sm90.cu	TMA tensor、tma_partition、transaction barrier、PipelineState	解释同一 SMEM stage 何时可读、何时可重写
9	blackwell/01_mma_sm100.cu	tcgen05.mma、TMEM 分配与释放、make_tmem_copy；输入先用普通 cooperative copy	画出 SMEM→MMA→TMEM→RMEM→GMEM 数据流
10	blackwell/02_mma_tma_sm100.cu	在 01 基础上用 TMA 加载 A/B	对照 01，解释 TMA barrier 与 MMA barrier 各等待什么
11	blackwell/03_mma_tma_multicast_sm100.cu	cluster、TMA multicast、mask、跨 CTA 同步	推导 A 沿 N、B 沿 M 广播的接收 CTA 集合
12	blackwell/04_mma_tma_2sm_sm100.cu	2SM MMA/TMA、leader/peer CTA、双 SM TMEM 分配	解释谁发出 MMA、谁发出 TMA、transaction bytes 为什么改变
13	blackwell/05_mma_tma_epi_sm100.cu	epilogue 分块、TMA load C/store D、store fence/commit/wait	解释 RMEM→SMEM→GMEM 输出路径中的每次同步
14	rubin/rubin_fp8.cu	当前分支 SM107 FP8 MMA、1CTA/2CTA、B matrix buffer reuse	对照有/无 B reuse 时的 MMA 和 tile 形状
15	rubin/rubin_fp8_blockscaled.cu	SFA/SFB scale tensor、SMEM→TMEM 的 UTCCP、带 scale 的 MMA；也包含 FP8/FP4 混合输入	单独画 scale 数据流，并与 A/B tile 对齐
16	rubin/rubin_fp4_blockscaled.cu	FP4 subbyte 存储、scale vector size、多种 scale 类型	区分逻辑元素个数、物理字节数、scale 个数

阶段 0：先把坐标映射打牢

预计 6–8 小时。先读 01_layout.md，再读 03_tensor.md。02_layout_algebra.md 先掌握 coalesce、composition、divide 和 product 的用途，遇到具体分块时再回看推导。

掌握顺序：

1. Int<N> 与运行时整数；Shape、Stride、嵌套 tuple。
2. Layout 是从逻辑坐标到索引的映射；区分 rank、depth、size、cosize。
3. Tensor 的 engine/iterator 与 Layout；区分内存视图和实际拥有存储的 Tensor。
4. _ 切片、local_tile、local_partition；变换的是视图还是数据。
5. composition(A,B) 的 A(B(coord))；divide 如何拆出 tile 与剩余坐标。
6. product 如何组合布局，尤其是线程布局和值布局。

练习：手算 Layout<Shape<_4,_3>,Stride<_1,_8>> 的完整地址表，说明为什么 size=12、cosize=20。再切出一列，说明为什么不需要复制数据。用 print、print_layout 核对结果。

验收：给出一个两层嵌套的 shape/stride，你能先手算坐标对应的地址，再用小程序确认。暂时不用穷尽 inverse、complement、nullspace 的实现细节。

补充阅读采用“小案例查证”的方式：composition.cpp、logical_divide.cpp、logical_product.cpp、coalesce.cpp。这些测试帮助理解性质，不必一开始阅读全部测试组合。

阶段 1：把 Copy 单独学清楚

预计 4–6 小时，学习表中的 1。

tiled_copy.cu 中先读简单的 copy_kernel，再读 copy_kernel_vectorized；当前 main 默认启动的是后者。将“CTA 负责一个 tile”和“线程负责 tile 内哪些值”分开推导。

重点追踪：tensor_S → tiled_tensor_S → tile_S → thr_tile_S → fragment。

先固定 Copy Atom，分别修改线程布局和值布局；记录每个线程访问哪些地址，再判断合并访存。最后比较单元素与 128-bit copy，检查地址、连续性和对齐条件。

验收：能说明 partition_S/partition_D 如何协调同一次 copy 的源和目的，以及线程布局和值布局如何共同决定访问模式。不要把“每线程有 4 个值”直接等同于“必然生成一条 128-bit 指令”。

阶段 2：先学会处理边界

预计 2–4 小时，学习表中的 2，配合 0y_predication.md。

tiled_copy_if.cu 的默认 shape 是 (528,300)，tile 是 (128,64)。它将 identity tensor 变成坐标谓词，再对数据和谓词使用一致的 tile/线程划分。

先做标量 copy 的任意尺寸边界处理，再看向量化版本。向量化还受连续性和对齐条件约束，不能随意换成奇数尺寸后保留原 Copy Atom。

验收：能从某个边界线程的坐标算出哪些值有效；知道补零、跳过读取、跳过写回是不同的处理。后面给 GEMM 加 K 尾块时，要使无效乘法输入为零；不能只屏蔽 C 的写回。

阶段 3：从最简单的 GEMM 建立全局图

预计 6–8 小时，学习表中的 3，配合 0x_gemm_tutorial.md 的 sgemm_1 部分和 04_algorithms.md。

按以下顺序阅读，不必从第一行模板声明开始：

1. gemm_nt/gemm_tn：固定一个配置，明确 shape、stride、CTA tile 和线程数。
2. mA/mB/mC → gA/gB/gC：整矩阵到 CTA 视图。
3. tAgA/tAsA 与 tCsA/tCsB/tCgC：复制划分和计算划分。
4. tCrC 与 K 循环：共享内存复用、寄存器累加和同步。
5. axpby：写回完整的 alpha*AB + beta*C。

特别注意：CuTe GEMM 中 B 通常用 (N,K) 表示，cute::gemm 的坐标约定不能直接套用日常 (K,N) 的形状记忆。

第一遍固定对齐尺寸。独立写一个小版本，对照 CPU 或 CUDA reference 检查数值。之后再把阶段 2 的方法用于 M/N/K 边界。

验收：能画出 A/B 从 GMEM 到 SMEM、每线程读取到累加器的路径，并说明两次 __syncthreads 分别保护什么。

阶段 4：理解 Atom/Tiled 与软件流水线

预计 8–10 小时，学习表中的 4、5。

先对照 sgemm_1 → sgemm_2。前者用线程 Layout 进行划分；后者显式使用 TiledCopy 与 TiledMMA，并预取下一 tile 到寄存器。此时 MMA Atom 是 UniversalFMA。

再对照 sgemm_2 → sgemm_sm70。后者显式管理 SMEM→RMEM 的分批加载与计算，使 K tile 内部也形成流水。

当前 sgemm_sm70.cu 的计算仍是 UniversalFMA。它的教学重点是适合该代架构的软件流水线，文件名并不意味着它演示 Volta Tensor Core 指令。

练习：画至少两个 K tile 的时序，标注 GMEM→RMEM、RMEM→SMEM、SMEM→RMEM、FMA。对每个临时寄存器和 SMEM tile，记录何时产生、何时最后被读取、何时允许覆盖。

验收：能解释为什么 TiledMMA 不必对应硬件 Tensor Core；能说明预取发生在什么位置，以及为什么不会覆盖还在使用的数据。

阶段 5：把 SM80 拆成两课

预计 10–14 小时，学习表中的 6，配合 0t_mma_atom.md。

第一课读使用 UniversalFMA 的泛型 NT/TN 配置，专注 SM80_CP_ASYNC_*、cp_async_fence、cp_async_wait 和 PIPE 维。当前示例使用 3 个 SMEM stage。通过表格追踪 smem_pipe_read、smem_pipe_write 与 K tile 索引。

第二课读 half TN 重载：SM80_16x8x16_F16F16F16F16_TN、SM75_U32x4_LDSM_N、Swizzle、make_tiled_copy_A/B、retile_D。同一个文件此时才进入硬件 Tensor Core 路径。当前 main 使用 half 类型，泛型 float 路径需要在练习副本中调整类型后实例化。

练习：先证明 copy 视图与 MMA fragment 描述的是同一批逻辑元素，再分析 LDSM 的 lane/value 映射；追踪一个元素从 GMEM、swizzled SMEM、寄存器直到 MMA 的位置。retile_D 调整寄存器 Tensor 的视图，本身不搬运数据。

验收：能独立解释“为什么等这个 async group”“为什么 SMEM 还不能覆盖”“为什么 Copy 和 MMA 的视图不同却能对接”。然后再改 stage 数和 tile 大小，记录耗时、SMEM 用量与寄存器变化。

相关测试按需要阅读：ampere/tiled_cp_async.cu、ampere/ldsm.cu、ampere/cooperative_gemm.cu。测试中的某些 Atom 有独立硬件要求，不能把整个测试文件视为所有新 GPU 都可运行的集合。

阶段 6：Hopper，先换计算指令，再换搬运方式

预计 8–12 小时，学习表中的 7、8。先完成 SM80 的核心概念，再读 0z_tma_tensors.md。

wgmma_sm90 保留 cp.async 搬运，并引入 WGMMA。对照 SM80 观察：这里的 SS MMA A/B fragment 是指向 SMEM 的 descriptor 视图，不能继续一律当成普通寄存器数组；WGMMA 由 warpgroup 协同参与。

wgmma_tma_sm90 再引入 TMA，重点是 host descriptor 构造、TMA 坐标 tensor、tma_partition、预期传输字节数，以及 producer/consumer barrier 的 index 和 phase。

验收：对每一段 SMEM 说明生产者、消费者、完成信号与复用条件。不要把“发出了异步指令”直接等同于“已经获得理想的计算/访存重叠”；教程中的等待策略相对直观，后续才研究 warp specialization。

阶段 7：SM100，严格按 01→05 增加能力

预计 12–18 小时，学习表中的 9–13。

这一组是显式设计的递进教程，保留原顺序即可。第一遍把重点放在每个相邻文件的差异，而不是重复阅读所有打印和验证代码。

重点变化有三条：accumulator 从寄存器进入 TMEM；MMA 的划分要理解 CTA/cluster 层级；multicast 和 2SM 需要在跨 CTA 的数据使用结束后才能复用 SMEM。每一课只新增表中对应的能力。

验收：读完 05 后，能画出 A/B 主循环与 C/D epilogue 两条数据流，并解释 barrier、phase、transaction bytes、multicast mask、TMEM 生命周期和 TMA store fence 的作用。

阶段 8：Rubin 与低精度缩放

预计 8–12 小时，学习表中的 14–16。它们不是基础 CuTe 的前置要求。

先看非 blockscaled FP8，理解当前分支 SM107_MMA_*、1CTA/2CTA 和 B buffer reuse 的配置关系。再读 FP8 blockscaled：把 SFA/SFB 当成与 A/B 并行的数据通路，追踪 GMEM→SMEM→TMEM、UTCCP，以及带 scale 的 MMA 输入。最后读 FP4 的打包存储、scale vector size 和 scale 类型。

验收：能区分普通 FP8、blockscaled FP8、FP4 的数据/scale 布局；能正确计算逻辑元素、存储字节和 scale 的关系。精度实验用小矩阵，并与显式反量化后的 reference 比较。

结合当前机器选择实践路线

当前机器检测到 RTX 5060 Ti，compute capability 12.0（SM120），NVCC 13.2。

内容	当前机器的安排
Layout/Tensor 小实验	可以在 host 上进行
两个 Copy 与四个基础 GEMM 文件	优先动手；按本机架构编译，使用满足示例约束的配置
Hopper WGMMA 两例	源码学习；实测需要匹配的 Hopper GPU 与编译目标
SM100 01–05	源码学习；需要匹配的 SM100 GPU，不适用于 SM120 GeForce
Rubin 三例	源码学习；本分支要求 SM107 GPU，且相关宏要求 CUDA 13.4+

因此，本机实践主线是 阶段 0–5 → 自写一个有边界处理和 reference 验证的 GEMM → SM120 应用分支。Hopper、SM100、Rubin 保留为架构阅读路线。

SM120 应用入口可以选 79a NVFP4 GEMM 和 87a FP8 blockwise GEMM。它们属于 CUTLASS 集成示例，主要帮助将 CuTe 概念对应到 collective/builder/mainloop；适合基础完成后阅读。

编译与运行入口

基础阶段可以直接用 NVCC 编译单个示例。以当前机器为例：

cd /home/x/workspace/cutlass

nvcc -std=c++17 -O2 --expt-relaxed-constexpr -arch=sm_120 \
  -I include -I tools/util/include \
  examples/cute/tutorial/tiled_copy_if.cu -o /tmp/cute_cpp_copy_if
/tmp/cute_cpp_copy_if

nvcc -std=c++17 -O2 --expt-relaxed-constexpr -arch=sm_120 \
  -I include -I tools/util/include \
  examples/cute/tutorial/sgemm_sm80.cu -o /tmp/cute_cpp_sm80
/tmp/cute_cpp_sm80 256 256 256 T N

本次已用 NVCC 13.2、-arch=sm_120 编译并在 RTX 5060 Ti 上执行这两个源码入口：tiled_copy_if 的普通和向量化复制均输出 Success.；SM80 half TN 示例正常运行并输出耗时。后者原程序没有做 reference 数值比较，因此此次仅确认编译与运行，不据此声明数值验证通过。

也可使用仓库 CMake，配置 CUTLASS_NVCC_ARCHS=120 并只构建需要的 target。基础 target 是 cute_tutorial_tiled_copy、cute_tutorial_tiled_copy_if、cute_tutorial_sgemm_1、cute_tutorial_sgemm_2、cute_tutorial_sgemm_sm70、cute_tutorial_sgemm_sm80。SM100 教程 target 在当前 CMake 中以 CUTLASS_NVCC_ARCHS 包含 100a 为条件。

源码使用中需要记住：

1. sgemm_sm80 的 main 默认 half、NT，但 half NT 重载中有 assert(false && "Not implemented")。用上面的 T N 参数走已实现的 half TN 路径。
2. sgemm_1/2/sm70/sm80 的教学主循环默认要求尺寸适配 tile。grid 使用 ceil-div 不代表 kernel 自动具备完整边界保护。
3. 基础 GEMM 程序主要展示运行和测速，不能把打印 GFLOP/s 当成数值正确性的证明；自写版本要加 reference 比较，half 累加要采用合适的误差标准。
4. SM100、SM120 都叫 Blackwell，但使用的 MMA 能力不同；架构编号不能简单按数值大小理解为指令兼容关系。

每个示例统一采用的学习记录

每次只记录以下五项，便于之后比较不同架构：

1. 形状：关键 Tensor 每一维是什么，变换前后对应哪些逻辑元素。
2. 位置：数据实际在 GMEM、SMEM、寄存器还是 TMEM；对象是存储、视图还是 descriptor。
3. 归属：哪些 CTA、warp、lane 负责哪些元素或发出哪些指令。
4. 时序：何时产生、何时可读、何时允许覆盖；由哪个同步操作保证。
5. 验证：手算/打印映射、reference 数值、资源用量和耗时各证明什么。

第一轮目标是能解释并修改；第二轮独立复写；第三轮再优化并分析生成的指令。基础主线完成后，再追 include/cute/atom 的布局构造和 include/cute/arch 的指令封装，会更容易将模板与实际数据流联系起来。