B200// TCGEN05
← B200 速查主页 / 特性卡片墙

B200 · BLACKWELL GB100 · 特性专题

tcgen05 · 五代 Tensor Core + TMEM

★ Blackwell 计算主线 · wgmma 的继任者：单线程发射 · 累加器驻 Tensor Memory

01

是什么

DEFINITION

tcgen05.* 是第五代 Tensor Core 的 PTX 指令族，取代 Hopper 的 wgmma（wgmma 仅 sm_90a 目标，sm_100 已移除）。三个根本变化：① 单线程发射——不再需要 128 线程 warpgroup 协作；② 累加器（和可选的 A 操作数）驻留在 TMEM——一块每 SM 256KB、独立于寄存器文件的矩阵存储；③ cta_group::2——两个 CTA 在两个 SM 上协同发射一个 2 倍宽的 MMA。

执行模型仍是异步：单线程发射 tcgen05.mma，用 tcgen05.commit 把「这批 MMA 完成」绑定到 mbarrier，消费侧等 mbarrier 再 tcgen05.ld 读回 TMEM。

02

为什么（vs Hopper wgmma）

RATIONALE · VS WGMMA (SM_90)
维度Hopper wgmma (cc 9.0)Blackwell tcgen05 (cc 10.0)
发射主体warpgroup = 128 线程协作1 个线程（elect_one）发 1 条指令
累加器位置寄存器文件（每 warp 一大块）TMEM 256KB/SM，不占寄存器
寄存器压力大 tile 时 accum 吃满 RF，挤占占用度RF 只留地址/epilogue——占用度解耦
最大 tile (atom)m64n256k16（单 warpgroup）128×256×16（cta_group::2 跨两 SM）
单指令延迟32–128 cyc，随 tile 线性涨[6]~11 cyc 恒定（11.0–11.4，tile 几乎不敏感）[6]
低精度FP8 / INT8+ FP6 / FP4(NVFP4) / block-scale（MX）
指令编码立即数字段32-bit idesc 指令描述符（SMEM 中，运行期可改）

结果：GEMM 主循环里「寄存器喂 MMA」的瓶颈被整体移走——累加器在 TMEM 里增量累积，寄存器只干指针和 epilogue。实测单指令 ~11 cyc 且不随 tile 变大而劣化[6]，这是 Blackwell 上大 tile 更划算的直接原因。

03

关键数字

KEY NUMBERS · CONSTRAINTS
项值 / 约束注
TMEM 容量256 KB/SM = 128 lanes × 512 col × 32-bit[2]与 RF 等大；全片 38MB
tcgen05.alloc 粒度列数 2 的幂：32 / 64 / … / 512地址按 32 列对齐；dealloc 必须配对
TMEM 带宽读 ~16 TB/s · 写 ~8 TB/s（每 SM）[6]与 L1/SMEM 带宽叠加，不互占
cta_group::2 tile128×256×16 atom（2 CTA / 2 SM）两 CTA 必须 cluster launch 且 rank 相邻
操作数位置D 累加器：TMEM；A：TMEM 或 SMEM；B：SMEMA 在 TMEM 可省一次 SMEM 读
kind 路径.kind::f16 / .kind::tf32 / .kind::f8f6f4 / .kind::mxf8f6f4mxf8f6f4 = block-scale（MX/NVFP4）
idesc32-bit 指令描述符（SMEM 常驻）dtype/M/N/transpose 等运行期可改
tcgen05.mma 延迟~11 cyc（FP16 11.2 / FP8 11.8 / FP4 12.6）[6]吞吐口诀：一次 mma 占 ~11cyc 流水位
读回通道tcgen05.ld / .st（32x32b 等 lane 布局）读后必须 tcgen05.wait::ld 才可用
04

可编译示例

COMPILABLE EXAMPLE · VERBATIM

模式逐字对齐 PTX ISA 与 CUTLASS SM100 gemm_sm100.hpp / Colfax 教程的 canonical 流程：alloc → mma → commit → wait → ld → dealloc。

① 单线程分配 TMEM（64 列）

// 分配结果（TMEM 起始列地址）由硬件写进 smem 变量
__shared__ uint32_t taddr;
if (threadIdx.x == 0) {
  asm volatile("tcgen05.alloc.cta_group::1.sync.aligned.shared::cta.b32 [%0], 64;"
    :: "r"((uint32_t)__cvta_generic_to_shared(&taddr)));
}
__syncthreads();  // 全 CTA 可见 taddr；列数必须是 2 的幂 (32–512)

② 单线程发射异步 MMA（A/B 在 SMEM，D 在 TMEM）

// idesc: 32-bit 指令描述符，编码 dtype / M / N / transpose / saturate…
// a_desc/b_desc: SMEM 矩阵描述符（swizzle 与布局必须匹配 TMA 写入布局）
if (threadIdx.x == 0) {
  asm volatile("tcgen05.mma.cta_group::1.kind::f16"
                " [%0], %1, %2, %3, %4;"
    :: "r"(d_tmem), "r"(a_desc), "r"(b_desc),
       "r"(idesc), "r"(scale_c));  // scale_c=0 覆盖, =1 累加
  // 完成通知绑到 mbarrier（一批发多条 mma 共用一次 commit）
  asm volatile("tcgen05.commit.cta_group::1.mbarrier::arrive::one"
                ".shared::cluster.b64 [%0];"
    :: "r"(mbar_smem));
}

③ 读回 TMEM → 寄存器（epilogue 前半）

// 等 mbarrier 翻相后，按 32x32b lane 布局读 64 列 × 每线程若干元素
asm volatile("tcgen05.ld.sync.aligned.32x32b.x64.b32 "
              "{%0,…,%63}, [%1];"
  : "=r"(d0), /* …62 个寄存器… */ "=r"(d63)
  : "r"(d_tmem));
asm volatile("tcgen05.wait::ld.sync.aligned;");  // 数据就绪屏障

// ④ 用完释放（列数与 alloc 一致）
if (threadIdx.x == 0) {
  asm volatile("tcgen05.dealloc.cta_group::1.sync.aligned.b32 %0, 64;"
    :: "r"(taddr));
}
05

注意事项 / 陷阱

PITFALLS · CHECKLIST
CAUTION · 陷阱清单
  • 发射与 commit 都只允许单线程（或 cta_group::2 下两 CTA 各一线程）；warpgroup 集体发射是 wgmma 语义，别照搬。
  • TMEM 列数必须 2 的幂、≥32，dealloc 的列数必须与 alloc 完全一致——泄漏 TMEM 等于泄漏显存，后续 CTA alloc 直接失败。
  • TMEM 是新的占用度维度：512 列总量，一个 m128n256 累加器就占 128 列，并发 CTA 数被 TMEM 卡死时降寄存器没用。
  • tcgen05.ld 之后必须 tcgen05.wait::ld 才能消费数据；.st 同理有 wait::st。
  • cta_group::2 要求 cluster launch 且两 CTA 的 SM 相邻；只开 1 个 CTA 用 ::2 是 UB。
  • A 用 TMEM 传参时（A 也在 TMEM），写入 A 与读 A 的 mma 之间同样要 mbarrier 同步——TMEM 没有 fence-free 的跨线程可见性保证。
  • idesc 的 dtype/M/N 与 a_desc/b_desc 不一致不报编译错（运行期描述符），错配 = 静默数据错乱。
06

数据源

SOURCES