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

B200 · BLACKWELL GB100 · 特性专题

TMA 增强 · + Cluster 16

★ Blackwell 搬运主线 · device 端改 tensormap · gather4/scatter4 · nonportable cluster 8→16

01

是什么

DEFINITION

Hopper TMA 的全部能力（1D–5D bulk 搬运 / multicast / mbarrier 完成计数）在 Blackwell 原样继承，叠加三件新东西[2]：

① tensormap.replace——tensor map（拷贝描述符）可以在 kernel 内原地改写：global 基址、维度、stride、boxDim 都能换，不再必须回 host 重编码。② tile::gather4 / scatter4——按索引数组一次搬 4 行非连续行（MoE 专家行聚合 / embedding 查表场景的原生指令）。③ cluster nonportable 上限 8→16——16 个 CTA 的 DSMEM 共享域（portable 上限仍为 8）。

02

为什么（vs Hopper TMA）

RATIONALE · VS SM_90 TMA
维度Hopper (sm_90)Blackwell (sm_100)
换 tensor map回 host 重新 cuTensorMapEncodeTiledtensormap.replace 在 kernel 内原地改[2]
非连续行搬运软件循环逐行发 TMA / gather 索引展开tile::gather4 / scatter4 按索引数组搬 4 行
cluster 上限8 portable / 16 nonportable同 8/16；但 multicast mask 16-bit 与 16 CTA 对齐，DSMEM 域实用价值放大
变长序列 / 动态 shape每变一次 shape 一次 host 往返kernel 内热更新，LLM decode 变 batch 零往返
典型受益算子规则 tile GEMM / attention+ MoE 行聚合、embedding 查表、变长 batch

一句话：Hopper TMA 解决「规则 tile 搬得快」，Blackwell 补上「不规则访问与动态描述符」——这两块以前只能退回逐线程访存或 host 干预。

03

关键数字

KEY NUMBERS · CONSTRAINTS
项值 / 约束注
replace 可改字段global_address / globalDim / globalStrides / boxDim / elementStridestensormap.replace.tile.*.global.b1024
replace 可见性需 fence proxy（tensormap.cp_fenceproxy）改完必须 fence 再发 cp.async
gather4 粒度一次 4 行（tile 第 0 维），行号来自索引数组索引可为 smem / global，行内连续
scatter4 对称4 行写回不连续目标行与 gather4 同族，方向相反
cluster 维度portable ≤ 8 · nonportable ≤ 1616 需 cudaFuncAttributeNonPortableClusterSizeAllowed
multicast mask16-bit（每位 1 CTA）恰好覆盖 16-CTA cluster
Hopper 遗产约束boxDim ≤ 256/维 · stride 16B 倍数 · 单线程发射全部沿用，见 H200 TMA 特性页
04

可编译示例

COMPILABLE EXAMPLE · VERBATIM

① Kernel 内热更新 tensor map（免 host 往返）

// 逐字对齐 PTX ISA tensormap.replace 语义（sm_100+）：
if (tid == 0) {
  // 换基址（变 batch / 换 buffer 时）；dims/strides/boxDim 同族指令
  asm volatile("tensormap.replace.tile.global_address.global.b1024.b64"
                " [%0], %1;"
    :: "l"(&tmap), "l"(new_ptr));
  // 改完必须 fence proxy，之后的 cp.async.bulk.tensor 才能看到新值
  asm volatile("tensormap.cp_fenceproxy.global.shared::cta.b1024.b64"
                " [%0], [%1], %2;"
    :: "l"(&tmap), "l"(smem_tmap_copy), "r"(size));
}
__syncthreads();

② gather4：按索引数组一次搬 4 行（MoE 行聚合）

// idx_smem: 4 个行号（uint64）；每行 = boxDim[0] 连续元素
if (elect_one_sync()) {
  asm volatile(
    "cp.async.bulk.tensor.2d.tile::gather4"
    ".shared::cluster.global.mbarrier::complete_tx::bytes"
    " [%0], [%1, {%2}], [%3], [%4];"
    :: "r"(smem_dst), "l"(&tmap), "r"(col0),
       "r"(idx_smem), "r"(mbar) : "memory");
}
// scatter4 对称：.tile::scatter4 写回 4 条不连续目标行

③ 16-CTA nonportable cluster（DSMEM 域翻倍）

// host：显式 opt-in 才能用 16
cudaFuncSetAttribute(kernel,
    cudaFuncAttributeNonPortableClusterSizeAllowed, 1);
cudaLaunchKernelEx(&cfg, kernel, args);  // cfg.clusterDim = {16,1,1}
// device：16 CTA 的 DSMEM 邻居直访（同 Hopper 语义）
cg::cluster_group c = cg::this_cluster();
int* nb = c.map_shared_rank(smem, 15);  // rank 0..15
05

注意事项 / 陷阱

PITFALLS · CHECKLIST
CAUTION · 陷阱清单
  • replace 后必须 fence proxy（tensormap.cp_fenceproxy），且要跨线程可见（__syncthreads）——漏 fence 后续 TMA 读到旧 map，静默错数据。
  • tensor map 本体仍是 host 编码、64B 对齐、__grid_constant__ 传入；replace 只改字段不重建结构。
  • gather4 固定 4 行粒度：行数不是 4 的倍数需 padding 索引；索引数组元素为行号（坐标第 0 维），行内仍要求连续+对齐约束。
  • 16-CTA cluster 是 nonportable：不保证在所有 sm_100 设备可调度（资源不足直接 launch 失败），fallback 逻辑要自己写。
  • multicast + 16 CTA 时 mask 16 位正好用满；跨 cluster multicast 依旧不存在。
  • gather/scatter 走 sm_100/sm_103 专属：编译目标写成 sm_90 会静默编不过或运行期非法，检查 -arch。
06

数据源

SOURCES