关卡阅读材料 // 04
张量计算
把逐元素的标量乘加重新表达成 Warp 协作的矩阵 Tile，并在输入、乘法与累加之间选择合适的精度。
建议边读边写下答案；时长包含代码推演与练习。
进入互动关卡→- 01导读与目标05 MIN
- 02心智模型08 MIN
- 03概念深挖15 MIN
- 04代码推演12 MIN
- 05历史生态08 MIN
- 06练习复盘12 MIN
优先使用成熟的加速库与工具，在真实应用中选择算法、数据类型和可验证的性能路径。
关卡解释 Tensor Core、Warp Tile 与混合精度为什么成立；实验把手写标量 GEMM 与 cuBLAS 基线并列，避免把专用硬件当作自动加速。
- 课程 00–01
- 基础矩阵乘法
- 理解 FP16 / FP32 表示差异
- cuBLAS
- Nsight Compute
- CUTLASS Profiler（可选）
导读与目标
先记住这一句
专用计算单元只有在程序把工作表达成它认识的形状时才会生效。
完成后你应该能够
- 01
区分标量线程矩阵乘法与 Warp 级矩阵 Tile。
- 02
解释 WMMA Fragment 的协作所有权和不透明 Lane 映射。
- 03
分析低精度输入与高精度累加的数值交换。
- 04
判断应使用库、模板库还是手写 WMMA 微内核。
建立心智模型
- 01Tile
从大矩阵中取出固定形状的小块，作为一次 Warp 级矩阵操作的单位。
- 02Fragment
Warp 共同拥有矩阵片段；元素如何分布到 32 个 Lane 是不透明的实现细节。
- 03累加
较低精度输入可以减少计算与搬运成本，但较高精度累加常用于保护求和结果。
概念深挖
WMMA 是 Warp 级操作
load、mma 与 store 必须由整个 Warp 一致参与，不能把它理解成某个线程独占的矩阵指令。
混合精度是路径设计
它不等于把所有变量改成 FP16；输入、累加、主权重与损失缩放承担不同职责。
库是生态的一部分
cuBLAS、cuDNN 与 CUTLASS 把 Tile、布局和代际差异封装成更稳定的接口，通常比手写微内核更适合生产。
从输出元素并行到矩阵操作并行
传统 kernel 可以让每个线程负责一个 C[row][col]，但线程内部仍沿 K 维做标量循环。这已经有线程并行，却没有把工作表达成 Tensor Core 支持的矩阵乘加 Shape。WMMA 进一步让一个 Warp 共同拥有并计算一个 Tile。
大 GEMM 仍需多层 Tiling：线程块负责较大 Tile，Warp 负责子 Tile，数据从全局内存进入 Shared Memory，再进入 Fragment。Tensor Core 只加速核心 MMA；加载、布局、边界和写回仍决定完整 kernel 的效率。
停下来想一想256 个线程各算一个输出，为什么仍不等于使用 Tensor Core？＋
参考答案它们可能只执行普通标量 FMA。必须使用支持的库或 WMMA/MMA 路径，把操作、类型与 Shape 映射到 Tensor Core。
Fragment 的 Lane 映射故意不透明
wmma::fragment 表示一个 Warp 分布式持有的矩阵片段。每个 Lane 实际保存哪些元素由架构和实现决定，文档并不保证稳定映射。程序应通过 load_matrix_sync、mma_sync 与 store_matrix_sync 操作 Fragment。
Warp 中所有参与线程必须使用一致的模板参数和控制流到达 WMMA 操作。若部分 Lane 走不同分支，可能导致挂起或错误。需要逐元素修改时，应只使用文档明确支持的 Fragment 成员语义，或先写回普通存储。
停下来想一想能否假设 Lane 0 永远持有矩阵左上角元素？＋
参考答案不能。Fragment 元素到 Lane 的映射是不透明的，可能随架构变化。
混合精度保护的是误差路径
低精度乘数降低存储与计算成本，但累加会把许多项相加，小量可能在大部分和附近被舍入。使用 FP32 Accumulator 可以降低这种累计损失，却不能恢复输入量化前已经丢失的信息。
训练还可能需要 FP32 主权重、Loss Scaling 和特定算子保持高精度。生产代码通常应先使用 cuBLAS、cuDNN 或框架自动混合精度，再用 CUTLASS 或 WMMA 处理真正需要定制的路径，并用数值基准比较结果。
停下来想一想FP16 输入 + FP32 累加能否保证与全 FP32 完全相同？＋
参考答案不能。输入在转成 FP16 时可能已经量化；FP32 累加只保护后续求和的一部分误差。
把概念放回代码
01wmma::fragment<matrix_a, 16, 16, 16, half, row_major> a;02wmma::fragment<matrix_b, 16, 16, 16, half, col_major> b;03wmma::fragment<accumulator, 16, 16, 16, float> c;04wmma::load_matrix_sync(a, A, lda);05wmma::load_matrix_sync(b, B, ldb);06wmma::mma_sync(c, a, b, c);代码逐段推演
- 01DECLARE类型与 Shape 是接口的一部分
Matrix A/B 的布局、输入类型与 16×16×16 Shape 必须是目标架构和接口支持的组合。
- 02LOAD整个 Warp 一致加载
基址、Leading Dimension 和对齐必须满足要求；load_matrix_sync 把存储布局转换为不透明 Fragment。
- 03MMA一次表达 D = A×B + C
mma_sync 由 Warp 协作执行；Accumulator 类型决定求和路径，但不改变已经量化的输入。
- 04STORE把 Fragment 返回普通布局
store_matrix_sync 将分布式状态写回行主序或列主序存储，之后才能按普通矩阵访问。
知识图谱
补充参考 · 不要求一次读完，也不计入核心 60 分钟
核心术语
- GEMM
- 通用矩阵乘法 C = αAB + βC，是深度学习和科学计算中大量算子的计算核心。
- Tile
- 把大矩阵分解成适合线程组、片上存储和矩阵指令处理的小块。
- Tensor Core
- 面向小矩阵乘加的专用执行单元，需要合适的数据类型、布局、Shape 与软件路径。
- Accumulator
- 累加乘积的中间精度；它可以高于输入或输出精度，以控制数值误差。
- Mixed Precision
- 在同一计算中组合不同数据类型，平衡吞吐、容量、带宽与数值稳定性。
带数字推演
- 01算运算量
2 × 1024³ ≈ 21.5 亿次浮点运算。
- 02算最少字节
A 与 B 各约 2 MiB，C 约 4 MiB，理想一次性流量合计约 8 MiB。
- 03算强度
约 21.5 亿 FLOP ÷ 8 MiB ≈ 256 FLOP/byte。
诊断手册
调用 GEMM 却没有矩阵指令
- 先检查
- 数据类型、Compute Type、布局/对齐、Shape 和库选择的算法
- 需要的证据
- Profiler 中实际指令路径与库日志，而不是 API 名称
Tensor 路径更快但结果偏差大
- 先检查
- 输入量化范围、累加类型、归约长度与误差判断方式
- 需要的证据
- 绝对/相对误差分布和高精度参考结果
小矩阵吞吐很差
- 先检查
- 启动开销、Batch 数量、边界 Tile 与并行度
- 需要的证据
- 按 Shape 扫描的延迟、Batch GEMM 对照和 Occupancy
从硬件到生态
- 01CUDA Core 标量/向量 FMA
矩阵乘法由线程显式完成乘加和 Tile。
布局、Shared Memory 与寄存器复用决定效率。 - 02Volta Tensor Core
Warp 级矩阵乘加进入专用硬件路径。
软件栈需要表达矩阵 Shape、数据类型与累加语义。 - 03库 + 编译器 + Transformer Engine
cuBLAS、CUTLASS 和框架根据硬件选择 Tile 与低精度 Recipe。
优化逐渐从手写指令转向选择库、验证数值并分析实际路径。
硬件与生态坐标
Volta 首次在 SM 中加入 Tensor Core，CUDA 9 通过 WMMA 暴露 Warp 级矩阵乘加。之后的软件栈持续把更多数据类型、布局和模型算子映射到张量计算路径。CUDA Core 并未因此消失。
练习与复盘
在 WMMA 中，谁拥有一个 16×16×16 的矩阵 Tile？
- A. 一个线程
- B. 一个 Warp 协作拥有
- C. 整个 CPU 进程
查看答案＋
WMMA 的 load、mma 与 store 是 Warp 级协作操作，Fragment 在 Lane 之间的分布不应被程序假定。
每线程一个输出、每 Warp 一个 Tile、每 Block 一个大 Tile，分别是什么层次？
提示＋
区分标量所有权、Warp 协作与 Block 分工。
参考答案＋
每线程一个输出是标量并行；每 Warp 一个 Tile 对应 WMMA/MMA；每 Block 大 Tile 是更上层的分块，内部通常包含多个 Warp Tile。
为什么 4096 上反复累加 0.5 时，FP16 Accumulator 可能看不到变化？
提示＋
比较 4096 附近的 FP16 间隔。
参考答案＋
该量级的 FP16 相邻可表示值间隔大于 0.5，小增量在每次舍入时消失；FP32 累加有更细的尾数分辨率。
只有 threadIdx.x % 32 < 16 的 Lane 调用 mma_sync。问题在哪里？
提示＋
WMMA 的协作范围。
参考答案＋
同一 Warp 的 Lane 没有一致参与 WMMA，行为未定义。应让整个 Warp 在一致控制流中执行。
标准大矩阵乘法应先手写 WMMA 还是调用 cuBLAS？
提示＋
比较成熟库与定制成本。
参考答案＋
通常先用 cuBLAS并验证算法、布局和精度；只有库无法满足融合、特殊布局或研究目标时，再考虑 CUTLASS/WMMA。
可选实践实验
不计入核心 60 分钟 · 需要相应 CUDA / GPU 环境
实验目标
从标量 GEMM 到库与 Tensor 路径
建立正确且可重复的 GEMM 基线，比较标量 Kernel、cuBLAS 与支持的 Tensor Core 数据类型。建议步骤
- 01
生成确定性矩阵并用 CPU 或高精度库结果作为参考，分别定义绝对误差与相对误差阈值。
- 02
运行课程中的标量 Kernel，记录 Shape、布局、Leading Dimension 与端到端时间。
- 03
用 cuBLAS 实现同一操作，显式记录输入、累加和输出数据类型。
- 04
用 Nsight Compute 检查矩阵指令、内存与占用指标，确认实际路径，而不是从 API 名称猜测。
提交三种实现的误差、时间和 Profiler 路径证据，并说明为什么生产代码通常先选库。
使用 CUTLASS Profiler 比较两个 Tile/Stage 配置；只在 Shape、类型和正确性相同的条件下比较。