B200 的 2CTA 怎样调优：从矩阵分块到流水线

Blackwell 的 2CTA MMA 让两个 SM 上的一对线程块协作计算一个矩阵块。它可以减少重复搬运、降低每个 CTA 的输入存储需求，并支持更大的计算块；同时也引入配对调度、跨 CTA 同步和更大的累加结果。

优化的关键，是用数据复用和流水线收益覆盖这些成本。本文以 B200 上的 FP16／BF16 GEMM 为例，先算清一组分块的资源账，再讨论 tile、TMA、流水线、TMEM 和调度怎样配合。

一、先明确：2CTA 协作的是什么

Cluster：把多个 CTA 组成一个协作组

CTA 就是 CUDA 的 Block。Thread Block Cluster（线程块集群，本文简称 Cluster）是 Grid 与 Block 之间的一层可选分组：把若干 Block 组织在一起，允许它们在一次 kernel 执行中通信和同步。 这项能力从 Hopper 的计算能力 9.0 开始提供，B200 继续支持。

例如，一个 Grid 有 8 个 Block，每个 Cluster 包含 2 个 Block，就形成 4 个 Cluster：

Grid：一次 kernel 启动，共 8 个 Block
├─ Cluster 0：Block 0、1
├─ Cluster 1：Block 2、3
├─ Cluster 2：Block 4、5
└─ Cluster 3：Block 6、7

Cluster 是程序定义的协作组，SM 和 GPC 是硬件资源。 GPC（Graphics Processing Cluster）包含多个 SM；同一 Cluster 的 Block 保证被共同调度到同一个 GPC 中执行，每个 Block 仍由一个 SM 执行。共同调度保证这些块能在执行期间协作，不要求所有线程逐条指令同步推进，也不表示整个 Grid 的所有 Cluster 同时运行。CUDA 的 Thread Block Clusters

Cluster 提供哪些功能

能力	程序能做什么
共同调度	同组 Block 同时具备协作条件，支持跨块数据交接
跨 Block 同步	用 cluster.sync() 等机制等待组内线程完成准备，再进入下一阶段
分布式共享内存（DSMEM）	读取、写入或原子更新其他块的共享内存，分摊输入存储、交换结果
为硬件协作指定范围	在支持的架构上使用 TMA multicast 广播输入，或用 2CTA MMA 配对计算

其中，DSMEM 把组内各块的共享内存暴露为可跨块访问的地址空间。每块仍分配自己的 SMEM；访问远端 SMEM 需要遵守地址映射、同步和生命周期规则。它也不会让不同 Cluster 自动共享片上数据。分布式共享内存的使用规则

Cluster 的大小可以是一维、二维或三维。CUDA 的可移植上限通常为 8 个 Block，实际配置还受设备和 MIG 分区限制；B200 在显式启用非可移植配置后可支持 16 个。更大的 Cluster 会增加共同驻留的资源要求，应结合活跃 Cluster 数评估收益。B200 的 Cluster 大小与限制

一个最小例子：读取同组另一块的共享内存

下面每个 Cluster 固定包含两个 Block。每块把自己的 Grid 编号写入本地共享内存，再读取同组另一块写入的编号：

#include <cooperative_groups.h>
namespace cg = cooperative_groups;

__global__ void __cluster_dims__(2, 1, 1) cluster_peer(int* out) {
    cg::cluster_group cluster = cg::this_cluster();
    __shared__ int marker;

    if (threadIdx.x == 0) marker = blockIdx.x;
    cluster.sync();  // 同组各块已启动，并完成 marker 初始化

    if (threadIdx.x == 0) {
        int peer_rank = 1 - static_cast<int>(cluster.block_rank());
        int* peer_marker = cluster.map_shared_rank(&marker, peer_rank);
        out[blockIdx.x] = *peer_marker;
    }

    cluster.sync();  // 等所有远端读取完成，各块才能退出
}

// d_out 已分配，可容纳 4 个 int；显存管理、错误检查省略。
cluster_peer<<<4, 128>>>(d_out);  // 4 个 Block，形成 2 个 Cluster
// 完成并复制回主机后，结果应为 [1, 0, 3, 2]。

这里 <<<4,128>>> 的第一个参数仍是 Block 总数，不是 Cluster 数；Grid 各维的 Block 数必须是对应 Cluster 维度的整数倍。block_rank() 是块在本 Cluster 内的编号，blockIdx.x 是块在 Grid 内的编号。两次 cluster.sync() 分别保证数据已准备好、读取结束前对方不会退出；所有线程都要参与，__syncthreads() 只能同步本块，不能替代它。这个例子只演示 Cluster 通信，没有执行矩阵运算。

2CTA：在 Cluster 内配对做矩阵乘加

2CTA MMA，也称 2-SM MMA，指同一 Cluster 中一对 CTA 协作完成矩阵乘加。 每个 CTA 仍由自己的 SM 执行，两个 CTA 保留各自的共享内存和 Tensor Memory（TMEM，保存矩阵累加结果的专用片上存储）。

本文讨论的运算是 D = A × B，矩阵维度分别为 A[M,K]、B[K,N]、D[M,N]。一个 tile 是从大矩阵中切出的小块；沿 K 方向分段累加，最终得到一块 D。

同一个 Cluster：包含一对 CTA

CTA 0 → SM 0本地 SMEM：部分 A、B本地 TMEM：部分累加结果

CTA 1 → SM 1本地 SMEM：部分 A、B本地 TMEM：部分累加结果

一对 CTA 协作完成一次 2CTA MMA
tcgen05.mma.cta_group::2

图 1：两个 CTA 在同一个 GPC 内的两个 SM 上协作。SM 编号为示意；Cluster 规定协作范围，各块的线程和存储资源仍有边界。

PTX 用 .cta_group::2 指定配对协作。只把 Cluster 大小设为 2，不会自动把普通矩阵运算变成 2CTA MMA。实际实现还需要匹配的矩阵指令、数据布局、TMEM 分配与同步。PTX 的发射粒度与 CTA Pair

Cluster 可以容纳多对 CTA；2CTA 是一条 MMA 的协作粒度，Cluster 大小是更外层的组织参数。 下文先用只含一对 CTA 的 Cluster，避免把配对计算和更大范围的数据广播混在一起。

二、算一笔账：为什么可能更快

这一例子里，A 的总量没有减少；减少的是两个 CTA 对 B 的重复存储和加载。 矩阵元素没有压缩，计算结果也没有少算。先把矩阵乘法拆开，再看数据放在哪里。

1. 相同的输出，需要哪些输入

设这一轮要完成 256 × 256 的输出更新，沿 K 方向先处理 64 个元素。A 分块为 256 × 64，B 分块为 64 × 256。

将 A 按行切成上下两半，将 B 按列切成左右两半：

A × B → 本轮输出更新

A256 行64 列

B：64 行 × 256 列 →

B左64 × 128

B右64 × 128

A上128 × 64

A下128 × 64

A上 × B左输出左上

A上 × B右输出右上

A下 × B左输出左下

A下 × B右输出右下

输出：256 行 × 256 列；每格 128 × 128

图 1a：左侧 A 按行分，上方 B 按列分；沿 A 的行、B 的列找到交点，就是对应的输出分块。图中尺寸均为“行数 × 列数”，这一轮沿 K 方向处理 64 个元素。

例如，输出上半的某一行，既要乘 B 的左 128 列，也要乘右 128 列。因此，负责上半输出的计算需要完整 B，负责下半输出的计算也需要完整 B。 它们分别使用不同的 A，但共用同一份 B。

这里的 64 是一次搬入 SMEM 的 K 分段，不代表一条底层 MMA 指令就覆盖 K=64；实现可通过多条 MMA 累加完成。若整个矩阵的 K 更长，再重复处理后续分段。

2. 两个独立 1CTA：各自存完整 B

先用两个普通的 1CTA MMA 任务，各算一半输出。每个任务的矩阵指令使用本 CTA 准备的输入：

CTA 0 的 SMEM：A上 + B左 + B右 → 输出上半
CTA 1 的 SMEM：A下 + B左 + B右 → 输出下半

A上、A下各存一次；B左、B右却都存了两次。两个独立 CTA 即使同时运行在两个 SM 上，也不会仅因为数据相同，就自动把另一块的 SMEM 作为自己的 1CTA MMA 输入。

3. 2CTA：B 分开存，计算时协同使用

换成配对的 2CTA MMA 后，输入这样放置：

CTA 0 的 SMEM：A上 + B左
CTA 1 的 SMEM：A下 + B右

存储按这个方式分摊，计算仍要覆盖四个乘积：

CTA 0 的 SMEM：A上、B左
CTA 1 的 SMEM：A下、B右

输出归属	左 128 列	右 128 列
CTA 0 上 128 行	A上 × B左B 在本地	A上 × B右B 来自对端
CTA 1 下 128 行	A下 × B左B 来自对端	A下 × B右B 在本地

图 1b：四格构成完整的 256 × 256 输出更新。每格为 128 × 128；行表示输出在两块 TMEM 中的归属，格内标明所需 B 位于本地还是对端。示意不表示四条独立指令或固定执行顺序。

CTA 0 对应的计算需要 A上、B左和 B右，CTA 1 对应的计算需要 A下、B左和 B右。2CTA MMA 允许硬件使用分布在配对 CTA 共享内存中的 B，两个 SM 因而能共同使用完整 B，而无须在每块 SMEM 中都常驻一份完整副本。 PTX 对 B 的位置明确包含当前 CTA 和 peer CTA。PTX 的 tcgen05 矩阵操作数

实际流程是：两个 CTA 按布局准备各自输入，完成必要同步，再由配对中选出的一个线程发起 2CTA MMA；硬件完成跨 SM 的配合，结果按布局累加到两块各自的 TMEM。CTA 0 负责的输出覆盖左、右两半，CTA 1 也一样；只计算图中的两个对角格，会漏掉一半输出。

这与官方教程使用的 A 按行分、B 按列分、输出按行分的布局一致。其他 MMA 形状和布局需按对应规则推导。NVIDIA 的 2CTA 分块教程

4. 省了多少：一份 B，而不是一半 A 加一半 B

A、B 都使用 2 字节的 FP16／BF16，1 KiB = 1024 B。完成相同的 256 × 256 × 64 工作，单轮输入存储如下：

单轮输入存储	两个独立的 1CTA，各算 128 × 256	一对 2CTA，共算 256 × 256
每个 CTA 的 A	128 × 64 × 2 B = 16 KiB	16 KiB
每个 CTA 的 B	64 × 256 × 2 B = 32 KiB	64 × 128 × 2 B = 16 KiB
每个 CTA 的 A＋B	48 KiB	32 KiB
两个 CTA 合计	96 KiB	64 KiB

按两个 CTA 合计，更容易看出差别：

两个独立 1CTA：一份 A（32 KiB）＋ 两份 B（64 KiB）＝ 96 KiB
一对 2CTA：   一份 A（32 KiB）＋ 一份 B（32 KiB）＝ 64 KiB

所以，A 总量保持 32 KiB，B 的 SMEM 副本从两份变一份；合计省下 32 KiB。 每个 CTA 的输入缓冲由 48 KiB 降到 32 KiB。省下的空间可以用于增加流水线 stage，或为 epilogue 留出空间。Stage 是循环缓冲区的一份独立槽位。

两边仍完成同样的乘加；2CTA 增加了跨 SM 协作和同步成本。是否更快，取决于减少重复加载、缓解存储压力和改善流水线带来的收益，能否覆盖这些成本。

这里统计的是算法要求加载到 SMEM 的数据量。独立 CTA 对 B 的重复读取可能命中 L2，因此不能据此断言 HBM 流量也下降三分之一，更不能把它直接换算成加速比。

同一轮执行 2 × 256 × 256 × 64 = 8,388,608 次浮点运算。只除以输入加载字节数，计算强度从约 85.3 提高到 128 FLOP/B。这是忽略输出、缓存和同步的模型，作用是解释复用收益。

硬件总量没有翻倍：两个独立的 1CTA 本来就能使用两个 SM。2CTA 的优势要从减少冗余、指令形状和执行效率中获得。

5. 能自己分片，为什么还要用 2CTA 指令

分片由程序决定，矩阵乘加由执行单元完成。自己用 CUDA Core 写循环和 FMA，也能算出同一结果；tcgen05.mma.cta_group::2 则让两个 SM 的 Tensor Core 按规定布局协同计算，直接使用分布在两块中的输入，并将结果累加到各自的 TMEM。对于支持的数据类型和大块稠密矩阵，这条专用路径通常有更高的乘加吞吐。

输入复用也有不同层次：TMA multicast 可以把同一输入送给多个 CTA，但每个接收者仍保存自己的副本；前例的 2CTA MMA 则允许 B 的左右半片分放两块，计算时协同使用。普通 1CTA MMA 需要在本地备齐当前运算所需的 B，也可以分多轮加载较小子块来节省空间。2CTA 的输入分布与计算方式

因此，2CTA 的价值包括 Tensor Core 跨 SM 协同、支持更大的 MMA 形状，以及减少输入副本。软件分片仍然需要做，专门指令负责高效执行这些分片之间的乘加。

三、先选 tile，再调整流水线

M、N 决定复用，也决定有多少工作可以调度

增大输出 tile，通常能让一次输入读取服务更多乘加；代价是累加结果增大、独立输出任务减少。2CTA 配对后，一个任务又需要两个 SM，不能只看任务数减少了多少。

例如把 128 × 256 的 1CTA tile 换成 256 × 256 的 2CTA tile，任务数大致减半，但每个任务使用的 SM 数也加倍。是否改善整卡利用率，还要看实际驻留资源、任务数量与最后一批工作的分布。

边界同样重要。若 M 只有 128，却使用覆盖 256 行的 2CTA tile，在补零并计算完整 tile 的实现里，M 方向有效工作比例只有 50%。为整齐的大矩阵选出的配置，可能在小 batch 或 MoE 的小专家矩阵上吃亏。

下面这个小计算器可在 CPU 上估算输出覆盖率与输入缓冲。它只做预算，不判断某个后端是否支持该配置：

from math import ceil

M, N, K = 384, 4096, 4096
BM, BN, BK = 256, 256, 64  # 一对 CTA 的 tile
stages = 3
input_bytes = 2

pair_tiles = ceil(M / BM) * ceil(N / BN)
coverage = M * N / (pair_tiles * BM * BN)
smem_per_cta = stages * ((BM // 2) * BK + BK * (BN // 2)) * input_bytes

print("输出 tile 数:", pair_tiles)                 # 32
print("M/N 覆盖率:", f"{coverage:.1%}")            # 75.0%
print("每 CTA 输入缓冲:", smem_per_cta // 1024, "KiB")  # 96 KiB
print("K 分段数:", ceil(K / BK))                   # 64

覆盖率是几何指标，不是实际 Tensor Core 利用率。若实现能跳过部分工作，或者还有 K 尾块，实际情况需要另算。

K_tile 与 stage 数要一起看

固定 M、N 分块时，增大 K_tile 会同时增加单轮运算量和输入字节数，不会凭空提高前面那个输入计算强度。它可能减少循环和同步次数，却也让每个 stage 更大，挤掉可驻留任务或预取空间。

在上面的 2CTA 布局里，K_tile 从 64 增至 128，每个 CTA 的输入缓冲从每 stage 32 KiB 增至 64 KiB。3 stage 就从 96 KiB 增至 192 KiB，尚未计入屏障、padding、缩放因子和 epilogue。

因此应搜索 tile × stage 数的组合，先排除资源超限和边界浪费严重的候选，再实测。B200 的共享内存有每 SM、每 Block 的限制，Cluster 内也不能把两个 SM 的容量当作单个 CTA 可任意使用的一大块内存。Cluster 驻留情况应使用 cudaOccupancyMaxActiveClusters 等接口计算。Blackwell 资源与 Cluster occupancy

四、TMA：配对共享与 multicast 分开调

TMA（Tensor Memory Accelerator）负责异步搬运矩阵块。2CTA 调优时，需要区分两件事：

机制	数据怎样使用	主要调优点
一对 CTA 的 MMA 协作	两块存放不同输入片段，共同参与一次 MMA	分片、描述符、配对完成通知
TMA multicast	将一份输入送到多个需要它的 CTA	接收掩码、广播范围、屏障记账

对于前面的单对 CTA，B 左右各存一半，已经能支持配对计算；无需为了“两个 CTA 都要完整 B”再各复制一整份。若 Cluster 包含多个输出任务，沿 N 排列的任务可以复用 A，沿 M 排列的任务可以复用 B，再考虑 multicast。

优化时先确定每个接收者需要哪一片，再由布局生成 multicast mask。广播范围更大，也会增加接收存储和同步范围。 只向会消费数据的 CTA 广播；不要把普通 1CTA 的 TMA atom、掩码和期望字节数直接复制到 2CTA 实现。

SMEM 的 swizzle 是地址重排，用于配合访存和矩阵指令。应先沿用库为该 MMA 提供的布局，再根据瓶颈改动；手工转置或压缩缓冲时，需要同时校验 TMA 描述符与 MMA 看到的布局。官方 CuTe 教程的 TMA 分片、掩码与 SMEM 布局

五、流水线：让搬运、MMA、结果处理重叠

2CTA 的输入更省空间，一个直接用途是多放几份输入缓冲：当前 K 分段做 MMA，后续 K 分段由 TMA 预取。

角色	阶段 1	阶段 2	阶段 3
TMA 生产者	加载 K 段 0	加载 K 段 1	加载 K 段 2
MMA 发起者	等待输入	累加 K 段 0	累加 K 段 1
Epilogue	处理上一输出 tile 的完整累加结果

图 2：TMA 与 MMA 处理当前输出 tile 的各个 K 分段；epilogue 同时处理上一 tile。假设上一 tile 已完成全部 K 累加，且结果缓冲独立。时间段仅为示意，不按比例。

一个线程发起 MMA，为什么还分配一个 Warp

“一个线程发起”是指令要求，“分配一个 Warp”是程序的分工方式。 一个 Warp 有 32 个线程；tcgen05.mma.cta_group::2 由整对 CTA 中选出的一个线程发起，Tensor Core 随后异步执行矩阵乘加。专职 MMA Warp 通常负责持续运行“等输入 → 发起 MMA → 更新流水线状态 → 下一轮”的控制循环，发起时选出一个线程即可。PTX 的发起粒度

Warp Specialization 将这些控制工作与搬运、结果处理分给不同 Warp，让各角色独立推进。角色也可以合并：同一个 Warp 在不同阶段承担多种工作，具体由流水线设计决定。JAX 的 Blackwell Warp Specialization 教程

专职 Warp 的成本与收益

专职 Warp 会占用 32 个线程名额、一个驻留 Warp 槽位及相应寄存器。例如每块 8 个 Warp，其中 1 个专职控制，便用了 1/8 的线程名额；这不等于损失 1/8 的计算吞吐。等待中的 Warp 不会独占 32 个 CUDA Core，执行单元可以服务其他就绪 Warp。

更可能划算	更可能不划算
K 循环较长，少量控制指令就能持续驱动大量 Tensor Core 计算	小矩阵、短 K，流水线很快收尾
搬运、MMA、epilogue 有足够工作重叠，分工后减少互相等待	额外 Warp 和寄存器挤掉驻留任务，其他角色又经常空等

收益来自让 Tensor Core 更少断粮，同时让其他 Warp 处理 softmax、类型转换等工作。调优时比较整个 kernel 的耗时，再决定拆开角色还是合并角色。

输入缓冲怎样交接

用状态机看一份输入缓冲，比记住几条 barrier 指令更容易：

EMPTY：前一轮 MMA 已不再读取这份输入
  ↓ 生产者取得槽位，设置正确的 TMA 完成条件
FILLING：TMA 正在写入
  ↓ 相关输入全部到达
FULL：允许发起使用这份输入的 MMA
  ↓ MMA 已完成对这份输入的使用
EMPTY：允许下一轮覆盖

发起异步 MMA，不等于它已经读完 SMEM；发起 TMA，也不等于数据已可用。 提前释放输入缓冲可能覆盖仍在被读取的数据。循环复用屏障时，还需要正确维护 stage 索引和 phase，区分本轮完成与上一轮完成。PTX 的 tcgen05 内存顺序与异步完成规则

Stage 数增加后，先看 MMA 等输入的停顿是否减少。若等待没改善、驻留任务却减少，继续加深流水线通常没有收益。K 很短的任务，还可能来不及进入稳态就已经收尾。

六、两块的同步与 TMEM 生命周期要一起设计

2CTA 中有几个容易混淆的粒度：

操作	需要关注的参与者
发起 .mma.cta_group::2	一对 CTA 中选出的一个线程；库封装可能要求一个 Warp 共同进入调用
.alloc/.dealloc.cta_group::2	两个 CTA 各有一个 Warp，配对执行分配／释放
__syncthreads()	只同步当前 CTA 的线程
Cluster barrier／异步完成通知	按规则协调多个 CTA，或等待对应异步操作

这里的分工是：一个线程代表整对 CTA 发起 MMA，两个 CTA 各派一个 Warp，共同完成 TMEM 的分配和释放。tcgen05 的 Issue Granularity

同步负责交接缓冲：MMA 用完输入后复用 SMEM，epilogue 读完结果后复用 TMEM；两块结束相互访问后，再释放资源。CUTLASS 的 pipeline 和 TMEM allocator 可以管理这些交接，调优时再用细粒度完成通知减少整组等待。

七、TMEM 与 epilogue：计算变快以后，瓶颈会移动

MMA 算结果，累加缓冲存结果，epilogue 处理结果

累加缓冲保存一个输出 tile 已经算出的部分和。 沿 K 每处理一段，MMA 就把新的乘积加进去；全部 K 段完成后，缓冲中才是这一块完整的矩阵乘法结果。在本文的 B200 实现里，它位于 TMEM。

Epilogue（尾处理）读取累加结果，把它变成最终输出。 最简单的是转换类型并写回，也可以融合 bias、激活和缩放。用一个输出元素表示，两阶段的逻辑如下：

# 逻辑示意：实际 GPU 指令一次处理矩阵分块
acc = 0.0
for k in range(K):
    acc += A[i, k] * B[k, j]       # 主循环：持续累加

D[i, j] = to_fp16(acc + bias[j])   # epilogue：处理并输出

① MMA读取 A、B 分块沿 K 逐段乘加

→

② 累加缓冲 · TMEM部分和 → 完整结果对应一个输出 tile

→

③ Epilogue读到寄存器中处理转换／激活／归约

→

④ 输出到显存直接存储或经 SMEM＋TMA

累加槽位 0MMA 正在累加 tile p

累加槽位 1Epilogue 正在读取 tile p−1

图 3：上方为一块输出的数据流；下方为双累加缓冲的一种重叠状态。两槽轮换使用，输入 A、B 的 SMEM 缓冲另行管理。图中阶段和箭头不表示耗时比例。

MMA 与累加缓冲之间会反复发生读改写，直到 K 累加完成。Epilogue 随后把 TMEM 中的结果分批取到寄存器，做相应运算并输出；2CTA 中两个 CTA 分别处理自己拥有的输出部分。tcgen05 的累加与输出数据流

Epilogue 的工作通常分给谁

在 Warp Specialization 实现中，由一组 epilogue Warp 分担输出元素。 线程按布局各取一批结果，进行转换和计算，再协同写回；它们与 MMA 发起 Warp 是不同角色。这样的分工让 Tensor Core 做下一块乘加时，普通算术和访存管线可以处理上一块结果。

例如，CUTLASS 的一种 SM100 持久化 GEMM 将 Warp 0 分给 MMA，Warp 1 负责调度，Warp 2 搬运输入，Warp 3 加载 epilogue 需要的额外输入，Warp 4～7 执行 epilogue。Warp 数由 kernel 设计决定，这个配置是实例，不是 B200 的固定分配。官方 Warp 分工示例

简单 kernel 也可以由完成前一阶段工作的线程继续做 epilogue。是否单独安排一组 Warp，取决于能否获得有用的重叠，以及为此增加的线程和缓冲开销。

工作轻或重，先看依赖范围

to_fp16：边取结果，边转换和写回

FP32 转 FP16 是逐元素工作，每个输出可以独立处理。通常直接放在 GEMM 的 epilogue：epilogue Warp 各自取出一批累加结果，向量化转换并写回，不必为类型转换再启动一次 kernel。

这时更值得关注的是 TMEM 读取、输出访存和子块布局。运算指令虽少，整块输出的数据搬运仍有成本；给更多 Warp 也要看这些管线是否还有余量。

Gate／SwiGLU：让同一批线程拿到成对结果

如果 gate 指的是 y = up × SiLU(gate)，每个位置只需要一对投影结果，没有整行归约。可以把 up 和 gate 投影安排在同一个 GEMM 的相邻或交错列块中，再由 epilogue 线程配对取值、计算并输出。

它比单纯转换类型多了激活运算和输入配对，仍适合做融合 epilogue。关键是让两路结果在当前处理范围内就绪，避免为了配对而增加大量跨块交换。cuDNN 的 SM100 GEMM＋SwiGLU 实现就采用了列块配对。官方 SwiGLU 融合说明

Gate 若指 MoE 的专家选择网络，则还包含下面的 router 工作，组织方式会不同。

Softmax：先汇总一行，再得到每个输出

Softmax 需要计算行最大值、指数和行内总和，再归一化。它的关键依赖是同一行的其他元素。短行可以交给一个 Warp；更宽的行可由多个 Warp 分担，再合并局部最大值和总和。

如果一整行就在当前 CTA 的输出范围内，可以融合在本地处理。若一行横跨多个输出 tile，就需要合并这些 tile 的统计量：可以先由 GEMM 写出 logits，再交给专门的 softmax kernel；也可以为整个融合算法设计分块归约或在线更新。

Attention 是后一种方式的典型场景：QKᵀ → softmax → PV 可以组织成多阶段融合 kernel，softmax 在中间参与流水线，而不只是最后一次写回前的尾处理。CUTLASS v4.8.0 的一个 Blackwell FMHA 配置给两组 softmax 各分配 4 个 Warp，另设 correction、MMA、load 和最终 epilogue 角色；更多线程投入归约与归一化，MMA 仍由专门角色发起。官方 FMHA 分工源码、FlashAttention-2 的分块与工作划分

MoE router：把选专家与 token 整理分阶段安排

Router 通常先通过投影得到每个 token 的专家分数，再做评分、top-k 选择和路由权重计算。后面还要统计各专家的 token 数、安排输出位置，并把 token 整理成专家 GEMM 需要的布局。

一种常见组织是 router GEMM → 融合评分与 top-k 的路由 kernel → token dispatch → 专家 GEMM。选择阶段可以按 token 分配 Warp 或 CTA；批量计数、前缀和和重排则按各自的数据依赖安排。TensorRT-LLM 提供了将 softmax、top-k 和权重归一化合成一个路由操作的实现。官方路由融合示例

若专家数较少、完整分数行已在当前处理范围内，也可以进一步融合投影后的选择工作。选择融合边界时，需要比较省下的启动和显存往返，能否覆盖额外的寄存器、同步与驻留开销。Router 本身在小 batch 下还可能主要受启动开销影响，并不总是一个“大计算量”操作。

2CTA 的“两块”，与双缓冲的“两份”

2CTA 决定一块输出怎样由两个 CTA 合作完成；双累加缓冲决定能同时保存几块输出的累加结果。 前者是计算分工，后者是存储槽位数，可以分别选择。

本文的 2CTA MMA 把累加结果保存在 TMEM。高性能持久化 kernel 常把它与多累加槽位一起设计，让前后输出 tile 的处理重叠；简单实现也可以只用一个槽位。CUTLASS 的教程就分别说明了单 stage 和双 stage 的累加布局，1CTA 同样可以使用双缓冲。官方累加缓冲布局

继续使用 256 × 256 的输出 tile：CTA 0 拥有上半结果，CTA 1 拥有下半结果。两块各存一半，合起来才是一份完整结果；双缓冲再为下一块输出准备一组这样的空间。

同一对 CTA，管理两组结果槽位

累加槽位	CTA 0 · SM 0	CTA 1 · SM 1
槽位 0 输出 tile P	P 的上半128 × 256	P 的下半128 × 256
槽位 1 输出 tile Q	Q 的上半128 × 256	Q 的下半128 × 256

横看一行：两块共同拥有一个 tile 的结果。
纵看一列：一个 CTA 的 TMEM 中放两份不同 tile 的结果。

图 4：单缓冲只需要槽位 0；双缓冲使用两个槽位。增加的是两块各自的结果存储空间，CTA 数量始终为 2。尺寸表示 FP32 累加元素的行数 × 列数。

两个 CTA 仍然共同计算 P，再共同计算 Q。双缓冲并没有变成“CTA 0 单独算 P、CTA 1 单独算 Q”，也不需要再增加两个 CTA。2CTA 的每块 TMEM 分片

单缓冲和双缓冲，执行怎样不同

假设同一对 CTA 连续处理 P、Q、R 三块输出。单缓冲要先把 P 从 TMEM 读出，才能覆盖同一个槽位计算 Q；双缓冲则可以一边读 P，一边在另一槽位计算 Q。

单累加缓冲：MMA 与结果读取轮流使用槽位 0

时间 →	1	2	3	4	5	6
MMA	P₀	等	Q₀	等	R₀	等
读 TMEM	—	读 P₀	—	读 Q₀	—	读 R₀

双累加缓冲：读上一块时，在另一槽位计算下一块

时间 →	1	2	3	4	5	6
MMA	P₀	Q₁	R₀	—	—	—
读 TMEM	—	读 P₀	读 Q₁	读 R₀	—	—

P／Q／R 是输出 tile，下标 ₀／₁ 是累加槽位。
时段 2：两块共同计算 Q₁，同时各自读取 P₀ 的对应部分。

图 5：依赖关系示意，假设每块 MMA 和结果读取各占一格，且资源允许两阶段重叠。省略输入搬运及读出后的转换、激活和写回；格数不代表 B200 实测耗时或加速比。

双缓冲的槽位按 0 → 1 → 0 → 1 轮换：P 放槽位 0，Q 放槽位 1；P 读完后，R 复用槽位 0。MMA 与 epilogue 通过完成通知交接这些槽位。

这里的“读完”是指对应结果最后一次 TMEM 读取完成。单缓冲也能让后续 MMA 与寄存器中的转换、激活或写回重叠；双缓冲进一步把“从 TMEM 读出上一块”的阶段也与下一块 MMA 重叠。 槽位释放不必一律等到整个 epilogue 结束。CUTLASS 的累加缓冲交接实现

多放一份缓冲的成本与收益

前面的每个 CTA 拥有 128 × 256 个 FP32 累加元素，一份占 128 × 256 × 4 = 128 KiB。单缓冲时，两块合计 256 KiB；双缓冲时，每块需要 256 KiB，两块合计 512 KiB，分别位于两个 SM 上。

按这个布局，每份占 256 个 TMEM 列，双份占 512 列，达到文档给出的单 CTA TMEM 地址空间列数上限。更大的输出 tile 或额外 TMEM 用途，就要一起调整布局和槽位数。TMEM 的容量与分配规则

输入 stage 的 SMEM、累加缓冲的 TMEM，以及 epilogue 自己使用的寄存器和 SMEM，要分别计入资源预算。前文 2CTA 省下的是 B 的 SMEM 副本，不能直接拿来充当第二份 TMEM 累加缓冲。

同一对 CTA 连续处理许多输出 tile 时，双缓冲有持续的重叠机会；若只处理一块输出，或结果读取已很快，多占一份 TMEM 的收益就可能有限。选择时应观察 MMA 等待空槽位的时间，以及增加缓冲后的整体吞吐。

若 epilogue 的计算或归约持续拖慢流水线，可以在支持的配置中增加处理 Warp、调整子块或重新划分阶段；若瓶颈在输出带宽，增加 Warp 通常帮不了多少。加深缓冲可以隐藏延迟，却无法长期弥补 epilogue 的吞吐不足。CUTLASS 的 StageCountAutoCarveout 会为 epilogue 的 SMEM 留出空间，再选择主循环 stage，融合操作变化后也要重新配合调整。主循环与 epilogue 的资源配置

八、Cluster、L2 与调度：让足够多的配对持续工作

先用一个 CTA pair 跑通，再比较更大的 Cluster。扩大 Cluster 可以增加 multicast 复用，却也提高共同驻留的要求。更大 Cluster 的数据复用收益，要和可同时工作的 Cluster 数一起比较。

输出 tile 的遍历顺序同样影响 L2。让相邻任务在一段时间内复用相同 A 或 B 分块，可能减少显存访问；应测试 M 优先、N 优先或分组遍历，而不是假定一种顺序适合所有矩阵长宽比。CUTLASS 的 tile rasterization

持久化 kernel 让一个 worker 连续处理多个输出 tile，可以摊薄初始化成本。Blackwell 的 CLC 则允许完成工作的 worker 获取尚未启动的工作。用于 2CTA 时，两个 CTA 必须获取一致的任务坐标并同步推进；CLC 按 Cluster 粒度领取工作，不能让同一对 CTA 各自拿一个无关 tile。CUTLASS 的 CLC 调度模型

CLC 改善的是任务分配，无法补救过大的 tile、过少的总工作量或已经饱和的显存带宽。小 M/N、短 K 的问题应保留 1CTA 候选；M/N 小但 K 很长时，还可独立评估 Split-K 等增加任务并行度的方式，同时计入归约成本。

九、用官方示例核对实现

下面固定 CUTLASS v4.8.0 的 CuTe 教程 04。它展示 2CTA MMA、TMA 分片、multicast、屏障和 TMEM；适合核对机制，尚不是用于比较峰值性能的完整 benchmark。运行需要 B200、兼容驱动及支持 SM100 的 CUDA 工具链，例如 CUDA 13.0。

git clone --branch v4.8.0 --depth 1 https://github.com/NVIDIA/cutlass.git cutlass-2cta
cd cutlass-2cta

该版本教程默认的 Cluster 是 4 × 4 × 1。为先观察单个 CTA pair，在 examples/cute/tutorial/blackwell/04_mma_tma_2sm_sm100.cu 中，将 cluster_shape 的定义改为：

// 仅改变 Cluster 组织；教程已经选用了 2SM MMA 和对应的 TMA／同步。
auto cluster_shape = make_shape(Int<2>{}, Int<1>{}, Int<1>{});

然后构建和运行：

cmake -S . -B build-2cta \
  -DCUTLASS_NVCC_ARCHS=100a \
  -DCUTLASS_ENABLE_TESTS=OFF
cmake --build build-2cta --target cute_tutorial_04_mma_tma_2sm_sm100 -j 4
./build-2cta/examples/cute/tutorial/blackwell/cute_tutorial_04_mma_tma_2sm_sm100 512 1024 256

教程按位置接收 M、N、K，当前分块要求维度可整除，不支持任意越界 tile。先确认输出 Execution is successful.；出现 Test is waived 表示没有执行目标实现，不能算验证通过。目标名与行为可核对版本化源码及CMake 定义。

之后再迁移到带多级流水线和计时的 GEMM 示例。公平的 1CTA／2CTA 对比，应允许两边各自选择合适 tile、stage 和 epilogue；仅把一个底层指令替换掉，既可能不合法，也无法代表两种实现的最佳表现。

十、按瓶颈调参，用实测决定保留哪项改动

建议按这个顺序收敛：先验证布局与同步，再搜索 tile 和 stage，接着调整 epilogue，最后测试更大 Cluster、multicast、遍历顺序和 CLC。每轮只改变一类因素，保留上一版作为对照。

观察到的现象	优先核对	值得尝试的改动
MMA 经常等输入	TMA 完成条件、缓存、可预取深度	调整 stage、减少重复加载、调整布局
大 tile 快，小矩阵慢	边界覆盖率、独立任务数、初始化成本	缩小 tile，保留 1CTA 分派
stage 增加后变慢	每 CTA SMEM、活跃 Cluster 数	减少 stage 或 K_tile
MMA 等累加缓冲	epilogue 时长、TMEM 生命周期	调整结果缓冲或 epilogue 子块
SM 利用不均，收尾很长	总任务数、Cluster 大小、任务分配	缩小 Cluster，测试持久化／CLC
带宽已满，算力未满	实际 DRAM/L2 字节数、缓存命中	减少数据移动，先测复用收益

表中的对应关系是诊断假设，不是仅凭一条指标就能成立的结论。Nsight Systems 用于看时间线和 kernel 间关系，Nsight Compute 用于看资源与吞吐瓶颈；具体指标名称随工具版本变化。

计时前完成分配与初始化，分开 warm-up 和测量，多次运行并报告中位耗时与波动。用 CUDA Event 或 GPU 时间线测设备耗时，不能用 CPU 提交时间替代。记录 dtype、布局、M/N/K、GPU、工具链、Cluster、tile、stage 和融合操作。NVIDIA 的 GEMM 测量方法

还要分别测试缓存已热与轮换输入缓冲的情况。反复乘同一组小矩阵，测到的可能主要是 L2 复用；真实模型若不断换权重或输入，结论可能不同。最后用完整算子或模型负载确认收益，尤其关注小 batch、MoE 专家大小分布和融合成本。

本文的字节数、覆盖率和运算量均为明确假设下的计算示例，未提供 B200 实测加速比。源码入口固定为 CUTLASS v4.8.0，技术资料核对日期为 2026 年 10 月 25 日。