第 13 课：DeepSeek DeepGEMM 核心架构与 Hopper FP8 算子拆解

本课目标与完成标准

学完后你应能：

• 深入推演 DeepSeek-V3 核心算子库 DeepGEMM 的设计动机，说清其相比传统 CUTLASS 3.x / cuBLASLt 的架构取舍；
• 掌握 Hopper 架构（SM90）的核心指令级原语：Tensor Memory Accelerator (TMA)、Warp-Group MMA (wgmma.mma_async) 与 mbarrier 硬件事务屏障；
• 理清 Fine-grained Block-wise FP8 量化（$1 \times 128$ 细粒度分块缩放）的数学原理，以及在主循环内把部分和乘缩放系数并提升到 FP32 累加的实现方式；
• 掌握 MoE 场景下 Contiguous Grouped GEMM 与 Masked Grouped GEMM 的调度机制及其与 DeepEP 通信库的协同。

通过必须同时满足：

• 独立补齐本课唯一的代码填空题（Hopper WGMMA + 细粒度反量化分块模拟内核），并通过给定测试断言；
• 三个问答题均说明因果链与硬件微架构状态迁移，而不是只报术语；
• 能指出至少一个正确性边界和一个性能瓶颈取舍；
• 总分不低于 8/10，且没有一票否决级概念错误。

前置关系

• 课程前置：第 10 课：共享内存 Tiled GEMM、CUDA 异步拷贝机制、计算机体系结构（SM/Shared Memory/Tensor Core）。
• 本课在路线中的作用：作为算子核心主线的现代化高阶实战课，将经典的 Shared Memory Tiling 跨越升级到现代 Hopper (SM90) / Blackwell (SM100) 的无寄存器搬运异步流水线与 MoE 定制 Grouped GEMM。

核心心智模型

1. 它是什么，解决什么问题

DeepGEMM 是 DeepSeek 专为大语言模型（特别是 DeepSeek-V3 / V2.5 / R1 等超大规模 MoE 模型）研发的开源 FP8 Tensor Core GEMM 算子库。

(1) 为什么现有开源库（CUTLASS 3.x / cuBLASLt）不够用？

1. CUTLASS 3.x 模板层次深、编译耗时长：
   • 传统 CUTLASS 依赖极深的模板嵌套与代数表达（Layout、Tile、Iterator、Collective Builder），仅编译一组 GEMM 内核通常需要数分钟乃至数十分钟；
   • 二次开发、定制特殊访存或做微架构调优时，需要理解大量模板推导，调试门槛高。
2. DeepSeek-V3 细粒度分块量化（$1 \times 128$ Block-wise FP8）缺少原生支持（如图 13-1** 所示）**：

DeepGEMM 的细粒度 FP8 缩放

激活 A：1 × 128 分组

• 每个 token 沿 K 每 128 个元素一个缩放
• $\alpha = \max|A_{\text{group}}| / 448$（E4M3 最大值）
• 离群值只影响所在的 128 个元素

权重 B：128 × 128 分块

• 每个 $128 \times 128$ 块一个缩放 $\beta$
• 离线量化，缩放存为 FP32
• 分块与 Tensor Core 的 K 块对齐

主循环内提升累加

• 每个 128 宽 K 块：WGMMA 算 FP8 部分和
• 在 CUDA Core 上乘 $\alpha \beta$，累加到 FP32
• Epilogue 只做类型转换和写回

$C_{ij} = \sum_g \alpha_i^{(g)}\, \beta^{(g,\, \lfloor j/128 \rfloor)} \sum_{k=0}^{127} A^{\text{fp8}}_{i,\,128g+k}\, B^{\text{fp8}}_{128g+k,\,j}$

图 13-1：激活 1×128、权重 128×128 分块缩放，部分和在主循环内提升到 FP32

• 常规 TensorRT-LLM 或 cuBLASLt 主要面向 Tensor-wise（整张量单标量缩放） 或 Channel-wise（逐通道缩放）；
• DeepSeek-V3 为了在 FP8 下获得接近 BF16 的精度，采用 Block-wise Fine-grained Scaling：激活按 $1 \times 128$ 分块（每个 token 在 $K$ 维度每 128 个元素一个缩放系数），权重按 $128 \times 128$ 分块（每 128 个输出通道 × 128 个 $K$ 元素一个缩放系数），缩放系数为 FP32（较新的 DeepGEMM 版本也支持 UE8M0）。由于缩放系数沿 $K$ 维变化，反量化不能推迟到 Epilogue 统一完成，必须在主循环里每个 128 宽的 $K$ 块之后进行，cuBLASLt 当时不能原生融合这种做法。

3. MoE 解码（Decoding）与 CUDA Graph 的静态分配冲突：
   • 在推理 Decode 阶段，各 Expert 分配到的 Token 数量极小且每步动态变化；
   • 若使用传统变长 Grouped GEMM，CPU 必须每步重新统计各专家 Token 偏移量并下发 Kernel 参数，打破 CUDA Graph 捕获机制，带来沉重的 CPU Launch 延迟。

(2) DeepGEMM 的设计要点

• KISS 原则：不使用 CUTLASS 的多层类结构，核心内核代码量小且自包含，通过轻量级 JIT（即时编译）模块在运行时按形状编译，免去安装时的预编译开销；
• 直接使用 Hopper 硬件原语：利用 SM90 的 TMA（张量内存加速器）与 WGMMA 指令，官方在 H800 上报告最高约 1550 TFLOPS（H800 稠密 FP8 理论峰值约 1979 TFLOPS，约为峰值的 78%）；
• 双模 MoE Grouped GEMM 支持（如图 13-3** 所示）**：

DeepGEMM MoE 分组矩阵乘执行模式

Contiguous 连续布局 (训练与 Prefill)

动态 Token 按专家拼接
(M 维长度不规则)

M 维按 Tile 对齐
(M-alignment)

K-Grouped / M-Grouped
分组矩阵乘执行

Masked 掩码布局 (推理 Decode + CUDA Graph)

固定最大容量静态显存
(Zero-Allocation 缓冲)

每个专家的有效行数
(masked_m 张量)

接收 DeepEP 通信数据
(零 CPU 开销直传)

Masked GEMM 算子
仅计算有效活跃 Token

图 13-3：DeepGEMM 面向 MoE 的 Contiguous 与 Masked 两大布局设计对比

• Contiguous：用于训练前向/反向与推理 Prefill，连续内存拼接，算力利用率最大化；
• Masked：用于推理 Decode，配合 DeepEP 低延迟内核，以每个专家的有效行数 masked_m 标记有效 Token，形状固定，兼容 CUDA Graph。

---

2. Hopper (SM90) 硬件微架构机制剖析（如图 13-2 所示）

DeepGEMM Hopper (SM90) 异步生产者-消费者流水线

生产者：1 个 TMA Warp Group (128 线程)

TMA 异步批量读取 Matrix A
(cp.async.bulk.tensor)

TMA 异步批量读取 Matrix B
(cp.async.bulk.tensor)

TMA 异步加载 1x128 缩放因子 Scale Factors

mbarrier.arrive.expect_tx
(向硬件事务屏障注册字节数)

片上共享内存 (SMEM 多级循环缓冲区: kNumStages = 3~4)

SMem A Tile [BLOCK_M, 128] FP8

SMem B Tile [BLOCK_N, 128] FP8

SMem Scale Factors FP32

消费者：Math Warp Groups (128/256 线程)

mbarrier.try_wait.parity
(非阻塞硬件轮询事务完成)

wgmma.mma_async
(Tensor Core 直接从 SMEM 读入 A/B, 累加至 FP32 寄存器)

每个 128 宽 K 块：
FP32 Acc += D·SF_A·SF_B（CUDA Core）
Epilogue 转 BF16 写回

硬件异步直通写入SS 模式零寄存器开销

图 13-2：DeepGEMM 在 Hopper 架构下的 TMA 硬件直连与 WGMMA 异步多级流水线架构

要看懂 DeepGEMM 的源码，需要理解 Hopper 架构相对于 Ampere 架构的三个主要变化：

(1) TMA (Tensor Memory Accelerator)：异步硬件批量搬运

• Ampere 痛点：在 Ampere（SM80）中，虽然引入了 cp.async，但每个线程仍需显式计算地址偏移，执行指令发射。
• Hopper 创新：TMA 是片上独立的专用硬件单元。线程只需配置一个 cute::TmaDescriptor（包含张量基地址、形状、跨度、分块尺寸及 Swizzle 模式），随后仅需由单个线程发起一条 cp.async.bulk.tensor 指令，硬件直接将多维张量瓦片从全局显存（Global Memory）搬运至共享内存（Shared Memory），全程完全不占用 SM 通用寄存器和 ALU 算力。

(2) WGMMA (wgmma.mma_async)：Warp-Group 级异步矩阵乘

• Ampere 痛点：mma.sync 是 Warp 级同步指令，线程必须先从 SMEM 把数据显式 ldmatrix 加载到通用寄存器（Register File），消耗大量寄存器配额，极易导致寄存器溢出（Spilling）而降低 Occupancy。
• Hopper 创新：WGMMA 提升至 Warp-Group（4 个 Warp，共 128 个线程） 协同粒度。更关键的是，它支持 SS（Shared-Shared）寻址模式：Tensor Core 直接从 Shared Memory 读取矩阵 A 与矩阵 B 的切片执行点积，中间完全跳过通用寄存器。仅输出累加矩阵 $D$ 保存在 FP32 寄存器中，极大节省了片上寄存器开销。

(3) 生产者-消费者解耦（Producer-Consumer Splitting）与 mbarrier

DeepGEMM 将一个 ThreadBlock（例如 256 或 384 线程）进行角色解耦：

• 生产者（TMA Warp-Group，128 线程）：负责调度 TMA 硬件发起下一阶段的 A、B 及 Scale Factor 数据预取，并在硬件事务屏障（ClusterTransactionBarrier）上注册预期传输字节数（expect_tx）；
• 消费者（Math Warp-Groups，128 或 256 线程）：通过 mbarrier.try_wait.parity 轮询硬件屏障。一旦 TMA 数据到位，立即执行 wgmma.mma_async，用多级流水掩盖访存延迟（$k\text{NumStages} = 3 \sim 4$）。读完一个 stage 后，消费者还要在该 stage 的 empty barrier 上到达，生产者等到它之后才能复用这块 SMEM。

---

3. DeepGEMM 核心代码架构逐行拆解

以下为简化示意（非源码原文），结构参考 DeepGEMM 源码 sm90_fp8_gemm_1d1d.cuh 等文件，只保留关键骨架：

// 简化示意（非源码原文）：DeepGEMM SM90 FP8 GEMM 主循环的生产者-消费者骨架
// 函数名为示意，真实 API 与参数以 DeepGEMM / CUTLASS 源码为准
template <uint32_t BLOCK_M, uint32_t BLOCK_N, uint32_t BLOCK_K, uint32_t kNumStages>
__device__ void gemm_loop_sketch(uint32_t k_tiles /* TMA 描述符、SMEM 指针等省略 */) {
    // full_bar[s]：TMA 写完 stage s 后完成（生产者 -> 消费者）
    // empty_bar[s]：消费者读完 stage s 后到达（消费者 -> 生产者），防止 TMA 覆盖尚未读完的数据（WAR）
    __shared__ cutlass::arch::ClusterTransactionBarrier full_bar[kNumStages];
    __shared__ cutlass::arch::ClusterBarrier            empty_bar[kNumStages];
    // 由一个线程初始化：full_bar[s].init(1); empty_bar[s].init(消费者到达者数量); 之后 fence + __syncthreads()

    const bool is_producer = (threadIdx.x / 128 == 0);

    #pragma unroll 1
    for (uint32_t k = 0; k < k_tiles; ++k) {
        const uint32_t s      = k % kNumStages;
        const uint32_t parity = (k / kNumStages) & 1;   // 每绕环形缓冲区一圈翻转一次

        if (is_producer) {
            if (threadIdx.x == 0) {
                // 第一圈缓冲区本来就是空的；之后要等消费者释放上一圈的同一 stage
                if (k >= kNumStages) empty_bar[s].wait(parity ^ 1);
                full_bar[s].arrive_and_expect_tx((BLOCK_M + BLOCK_N) * BLOCK_K * sizeof(__nv_fp8_e4m3) /* + SF 字节 */);
                // cp.async.bulk.tensor：TMA 写完后由硬件扣减 full_bar[s] 的事务字节数
                tma_load_2d(&tensor_map_a, &full_bar[s], smem_a[s], k * BLOCK_K, m_block_idx * BLOCK_M);
                tma_load_2d(&tensor_map_b, &full_bar[s], smem_b[s], k * BLOCK_K, n_block_idx * BLOCK_N);
            }
        } else {
            full_bar[s].wait(parity);                    // 等待本 stage 的 TMA 数据到位
            warpgroup_arrive();                          // wgmma.fence.sync.aligned
            wgmma_ss(desc_a(s), desc_b(s), accum, /*scale_d=*/false);  // 若干条 wgmma.mma_async（SS 模式）
            warpgroup_commit_batch();                    // wgmma.commit_group.sync.aligned
            warpgroup_wait<0>();                         // wgmma.wait_group.sync.aligned 0：SMEM 已读完，accum 可用
            empty_bar[s].arrive();                       // 通知生产者 stage s 可以复用
            // 主循环内提升累加（CUDA Core，FP32）：final_accum += accum * sf_a * sf_b
            promote_with_scales(final_accum, accum, sf_a(s), sf_b(k));
        }
    }
    // Epilogue：final_accum 转 BF16 写回 GMEM
}

---

实践任务：唯一代码填空题

练习说明

在不支持真实 Hopper 物理 GPU 的模拟环境或静态代码审查中，我们需要深刻理解其分块细粒度反量化逻辑与算子数值等价性。

请实现一个细粒度分块量化 GEMM 模拟器（Block-wise FP8 GEMM：A 按 1x128、B 按 128x128 分块缩放）：

• 输入矩阵 $A \in \mathbb{R}^{M \times K}$ 与矩阵 $B \in \mathbb{R}^{N \times K}$ 采用 FP8 E4M3 数值格式（在练习中用低精度浮点模拟）；
• $K$ 维度按 $128$ 元素切分成若干 Block。A 的每一行在每个 $K$ 块上有一个缩放因子，$S_A \in \mathbb{R}^{M \times (K/128)}$；B 按 $128 \times 128$ 分块，$S_B \in \mathbb{R}^{(N/128) \times (K/128)}$；
• 真实数学计算：对于分块 $b \in [0, K/128 - 1]$，子块点积结果需乘上对应的缩放标量 $S_A[i, b] \times S_B[\lfloor j/128 \rfloor, b]$ 后再累加；
• 请补齐缺失的循环步进与反量化数学逻辑。

只能修改 TODO/______ 位置，不得删除断言或放宽通过条件。

import numpy as np

def fine_grained_block_gemm_simulation(A_fp8, B_fp8, S_A, S_B, block_k=128, block_n=128):
    """
    模拟 DeepGEMM 的细粒度 Block-wise FP8 GEMM 计算流（A: 1x128，B: 128x128）
    A_fp8: [M, K]  (E4M3 量化整型或低精度模拟)
    B_fp8: [N, K]  (E4M3 量化整型或低精度模拟，注意 B 为 Transposed 存储)
    S_A:   [M, K // block_k]  (FP32 缩放因子，每行每个 K 块一个)
    S_B:   [N // block_n, K // block_k]  (FP32 缩放因子，每 128x128 块一个)
    return: D [M, N] FP32 结果
    """
    M, K = A_fp8.shape
    N, K_b = B_fp8.shape
    assert K == K_b, "K dimension mismatch"
    assert K % block_k == 0, "K must be divisible by block_k"
    assert N % block_n == 0, "N must be divisible by block_n"
    num_blocks = K // block_k
    n_block_of_col = np.arange(N) // block_n  # 第 j 列属于 B 的第几个 128 行块

    D = np.zeros((M, N), dtype=np.float32)

    # 模拟 Hopper 内核外层循环：逐个 Block_K 推进并进行细粒度反量化累加
    for b in range(num_blocks):
        k_start = b * block_k
        k_end = ______

        # 提取当前 128-channel 瓦片
        A_tile = A_fp8[:, k_start:k_end]  # [M, block_k]
        B_tile = B_fp8[:, k_start:k_end]  # [N, block_k]

        # 1. 执行片上 Tile 点积累加: Tile_Acc = A_tile @ B_tile.T  -> [M, N]
        tile_acc = np.matmul(A_tile, B_tile.T)

        # 2. 提取当前 Block 的缩放因子外积: scale_tile[i, j] = S_A[i, b] * S_B[j // block_n, b] -> [M, N]
        scale_tile = ______

        # 3. 将反量化后的局部瓦片累加至全局输出矩阵 D
        D += ______

    return D

# 测试用例验证
np.random.seed(42)
M, N, K = 128, 256, 256
block_k = 128
block_n = 128

# 构造模拟输入
A_sim = np.random.randint(-10, 10, size=(M, K)).astype(np.float32)
B_sim = np.random.randint(-10, 10, size=(N, K)).astype(np.float32)
S_A = np.random.uniform(0.01, 0.05, size=(M, K // block_k)).astype(np.float32)
S_B = np.random.uniform(0.01, 0.05, size=(N // block_n, K // block_k)).astype(np.float32)

# 正确标准数学对照
ref_D = np.zeros((M, N), dtype=np.float32)
for b in range(K // block_k):
    A_sub = A_sim[:, b*block_k : (b+1)*block_k] * S_A[:, b:b+1]
    B_sub = B_sim[:, b*block_k : (b+1)*block_k] * S_B[np.arange(N) // block_n, b][:, None]
    ref_D += np.matmul(A_sub, B_sub.T)

out_D = fine_grained_block_gemm_simulation(A_sim, B_sim, S_A, S_B, block_k=128)
assert np.allclose(out_D, ref_D, atol=1e-5), f"Max difference: {np.max(np.abs(out_D - ref_D))}"
print("DeepGEMM 细粒度分块量化模拟测试通过。")

检查方法

运行脚本验证断言通过；核对 num_blocks 累加与数学基准的一致性。

Q1

为什么 DeepSeek-V3 选择 $1 \times 128$ 的 Block-wise FP8 量化，而不是直接沿用整个 Tensor 共享一个 Scale（Tensor-wise）？这对 Hopper WGMMA 硬件执行提出了什么挑战？

你的答案：

Q2

在 Hopper 架构中，WGMMA 指令的 SS 模式（A、B 两个操作数都从 Shared Memory 读取）与 Ampere 架构下的 ldmatrix + mma.sync 相比，消除了哪些微架构开销？为什么能大幅提升寄存器利用率和算子性能？

你的答案：

Q3

在 MoE 推理的 Decoding 阶段，DeepGEMM 为什么专门设计了 m_grouped_fp8_gemm_nt_masked 算子？为什么变长 Contiguous Grouped GEMM 无法与 CUDA Graph 协同工作？

你的答案：

评分与通过规则

• 代码 4 分：正确填补 k_end、scale_tile 与 D 累加逻辑并通过断言；
• Q1～Q3 各 2 分：必须从算子精度下溢、寄存器/访存层次、CUDA Graph 静态内存模型等本质因果链回答；
• 一票否决：无法讲清 WGMMA 与 TMA 的分工，或混淆 Tensor-wise 与 Block-wise 量化。

Tip💡 点击展开参考答案与详细代码实现

参考答案（仅 answer 分支）

def fine_grained_block_gemm_simulation(A_fp8, B_fp8, S_A, S_B, block_k=128, block_n=128):
    M, K = A_fp8.shape
    N, K_b = B_fp8.shape
    assert K == K_b, "K dimension mismatch"
    assert K % block_k == 0, "K must be divisible by block_k"
    assert N % block_n == 0, "N must be divisible by block_n"
    num_blocks = K // block_k
    n_block_of_col = np.arange(N) // block_n  # 第 j 列属于 B 的第几个 128 行块

    D = np.zeros((M, N), dtype=np.float32)

    for b in range(num_blocks):
        k_start = b * block_k
        k_end = (b + 1) * block_k

        # 提取当前 128-channel 瓦片
        A_tile = A_fp8[:, k_start:k_end]  # [M, block_k]
        B_tile = B_fp8[:, k_start:k_end]  # [N, block_k]

        # 1. 执行片上 Tile 点积累加
        tile_acc = np.matmul(A_tile, B_tile.T)

        # 2. 提取当前 Block 的缩放因子外积 (利用广播生成 [M, N])
        scale_tile = S_A[:, b:b+1] * S_B[n_block_of_col, b][None, :]

        # 3. 反量化并累加
        D += tile_acc * scale_tile

    return D

# 执行验证
out_D = fine_grained_block_gemm_simulation(A_sim, B_sim, S_A, S_B, block_k=128)
assert np.allclose(out_D, ref_D, atol=1e-5)
print("答案执行验证通过。")

Q1 参考答案

1. 数值精度与动态范围（Dynamic Range）： FP8（E4M3 格式）的动态范围很小（共 8 位：4 位指数、3 位尾数，最大绝对值 448，精度明显低于 BF16/FP16）。大模型各激活通道间存在严重的离群值（Outliers）。若采用 Tensor-wise 单一缩放因子，为了容纳极大离群值，缩放系数必须非常大，导致大多数正常数值下溢为 0，精度严重下降。细粒度分块量化把一个缩放系数的作用范围限制在激活的 $1 \times 128$ 块或权重的 $128 \times 128$ 块内，离群值只影响所在小块，局部信噪比更高。

2. 对 Hopper WGMMA 的挑战： WGMMA 只能在一段 $K$ 上连续累加，不能在累加过程中给不同 $K$ 块乘不同的缩放系数；另外 H800 上 FP8 WGMMA 的累加只保留约 14 位有效精度，$K$ 很大时全部在 Tensor Core 内累加会损失精度。DeepGEMM 把 $K$ 维分块大小（BLOCK_K）设为 $128$，与缩放分块对齐：每算完一个 128 宽的 $K$ 块，消费者线程把 WGMMA 的部分和乘以对应的 Scale_A * Scale_B，再在 CUDA Core 上累加到单独的 FP32 寄存器（DeepSeek-V3 报告中的 promotion 做法）。反量化因此发生在主循环内，Epilogue 只负责类型转换和写回。

Q2 参考答案

1. 消除寄存器往返搬运（Zero Register Allocation for Operands）： Ampere 的流程是：Global Memory -> (cp.async) -> SMEM -> (ldmatrix) -> Registers -> (mma.sync) -> Registers。A 和 B 矩阵必须同时常驻在寄存器堆中，每个线程需要几十甚至上百个寄存器存放输入操作数。 Hopper WGMMA 的 SS 模式是：Global Memory -> (TMA) -> SMEM -> (wgmma.mma_async) -> Registers(仅存放累加器 D)。矩阵 A 和 B 的输入操作数永远停留在 SMEM，Tensor Core 直接通过硬件总线并发抓取 SMEM 数据进行 FMA。

2. 性能收益： A/B 操作数不再占用寄存器，省下的寄存器用于更大的累加器 Tile（以及 FP32 提升累加用的寄存器），提高每个 CTA 的计算密度。DeepGEMM 通常每个 SM 只驻留 1 个 CTA，受 SMEM（多级 A/B 缓冲）限制，收益并不来自 Occupancy 提升。同时省去了 ldmatrix 指令的发射开销，减轻 Warp Scheduler 的发射压力；生产者 warp group 还可以通过 setmaxnreg 让出寄存器给消费者。

Q3 参考答案

1. CUDA Graph 核心不变性原理： CUDA Graph 依赖“静态执行图”消除 CPU Launch 开销。在 Graph 捕获期间，所有 Kernel 的 Launch 参数、输入输出指针、Grid/Block 维度必须完全固化，执行期不得发起任何主机端动态内存分配或 CPU-GPU 间的同步（如根据各专家实际 Token 偏移重算元数据）。

2. Masked 布局解决动态偏移： 在 Decode 阶段，每个 GPU 每步分配给各个 Expert 的 Token 数完全动态不可预知。m_grouped_fp8_gemm_nt_masked 使用形状固定的输入 [num_experts, max_m, K]，每个专家预留最大 Token 容量。DeepEP 低延迟 dispatch 把收到的 Token 写入各专家的固定区域，并给出每个专家实际收到的 Token 数（对应 DeepGEMM 的 masked_m 张量，形状 [num_experts]）。Kernel 启动参数固定，运行时在 GPU 上读取 masked_m[e]，只计算第 $e$ 个专家前 masked_m[e] 行所在的 M 方向 Tile，跳过其余 Tile。这样不需要 CPU 重算偏移，也不打断 CUDA Graph 回放。

面试高频追问与优化扩展

• 追问 1：TMA 的 Multicast 广播特性在 GEMM 中如何使用？
  • 回答要点：Hopper 引入了 Thread Block Cluster（可移植上限为 8 个 CTA）。当 Cluster 内多个 CTA 计算相同 $M$ 行块（或相同 $N$ 列块）的输出瓦片时，它们需要的 A（或 B）瓦片相同。TMA 支持 Cluster Multicast：一次 TMA 读取可同时写入 Cluster 内多个 CTA 的 Shared Memory，减少对 L2 的重复读取。DeepGEMM 通常使用 2 个 CTA 的 Multicast。
• 追问 2：DeepGEMM 的 JIT 用 NVCC 还是 NVRTC？各有什么取舍？
  • 回答要点：DeepGEMM 在运行时按 GEMM 形状、Block 大小等参数生成并编译内核，编译产物会缓存，后续调用直接加载。早期版本只调用 NVCC；后来加入了 NVRTC 选项（设置 DG_JIT_USE_NVRTC=1），官方说明编译速度最高约提升 10 倍，但某些形状上性能可能略低，因此默认仍使用 NVCC。取舍是首次编译耗时与内核性能之间的平衡。

源码参考与延伸阅读

• DeepSeek DeepGEMM 官方开源仓库
• NVIDIA PTX ISA 8.5: WGMMA & TMA Specifications
• DeepSeek-V3 Technical Report: Multi-head Latent Attention and DeepSeekMoE FP8 Training

← 上一课
第 12 课：ArgMax 对归约与确定性 下一课 →
第 14 课：DeepSeek DeepEP …