一次“转置”，半张算力：DeepSeek SwapAB 漂流史——从 H20 社区 PR 到 NVIDIA 官方命名

摘要SwapAB 是近两年开源 AI 基础设施领域最漂亮的微观样本之一在矩阵乘指令里交换 A/B 操作数把硬件最僵硬的维度分配给模型里最静态的张量。它的旅程横跨两代架构与六方生态2025 年 9 月社区贡献者 Wangzheee 在 DeepGEMM PR #192 中为 Hopper 写下wgmma(desc_b, desc_a, …)H20 上最高 66.7%一年半后NVIDIA 在官方技术博客中将同一技巧命名为 SwapsMmaAb披露 Blackwell 上稀疏 MLA 吞吐提升 47.03%。本文沿这条上游化路径完整拆解 SwapAB 的机制、代价与边界并以五组真实数字为标本讨论 kernel 级收益从社区 PR 流向集成方日志、厂商博客与官方 README 的过程中口径是如何变形的。0. 引子一个社区 PR 的上游之旅2026 年 9 月 26 日NVIDIA 发布 TensorRT-LLM 技术博客第 15 期《Optimizing DeepSeek-V3.2 on NVIDIA Blackwell GPUs》。其中一节以SwapsMmaAb为名专门介绍了一个矩阵乘优化技巧在矩阵乘中交换 A、B 操作数以提升张量核利用率。追根溯源这个技巧最初既不出自 NVIDIA也不出自 DeepSeek 的正式员工——它来自 2025 年 9 月社区贡献者 Wangzheee 向 DeepGEMM 提交的 PR #192最初跑在 Hopper 架构的 H20 上。一年半、两代架构、六方生态社区、DeepSeek、NVIDIA、vLLM、FlashInfer、昇腾一个几十行代码的技巧完成了从社区 PR 到芯片厂商官方命名的上游化之旅。这件事值得拆开写因为它同时具备三个特殊性技术上是深的它触及张量核指令最底层的形状约束且横跨 Hopper 与 Blackwell 两代架构传播上是广的它横跨社区贡献者、DeepSeek、NVIDIA、vLLM、FlashInfer、华为昇腾六方生态叙事上是有戏剧性的最初的手笔并非出自 DeepSeek 内部而是一位社区贡献者最终却被芯片厂商的官方博客收编命名——开源史上上游化的罕见案例。先把可核实的东西钉在墙上第 4 节将用它解剖五组真实数字数字口径硬件来源52.5% ~ 66.7%masked grouped FP8 GEMMm32 档两行H20HopperDeepGEMM PR #192 自带基准~40%vLLM 集成 DeepGEMM 的 GEMM kernel 提升#19820 初次集成报 ~40%适配 DeepGEMM breaking 更新上游 PR #112后以 #20087 合入特定 batch size 下 ~40%B200vLLM 贡献者 Wentao Ye 开发日志yewentao256.github.io初版曾因 breaking 变更经历 revert/重提社区记录 #2004947.03%FP8 稀疏 MLA 吞吐提升GB300Blackwell UltraNVIDIA blog1532.64%–64.20%各项 kernel fusion 的端到端收益区间同上NVIDIA blog151.5%DeepGEMM Linear 对端到端吞吐的提升BlackwellvLLM v0.10.2 release notes不看清张量核的行数之笼就理解不了这些数字为什么会相差五倍——下面先从硅片层面开始。1. 两代张量核同一个行数之笼1.1 指令形状被硅片焊死的 M 维从 Volta 到 Blackwell张量核的宏观趋势是单条指令的块头越来越大而 M 维输出矩阵的行始终被硬件数据通路焊死Volta / AmpereSM70/SM80mma.m16n8k8/16M 16。块头小浪费有限HopperSM90wgmma.mma_asyncM 64固定N 可在 {8, 16, …, 256} 中选取但为编译期立即数BlackwellSM100tcgen05.mmaUMMA单 CTA 发射时M ∈ {64, 128}两个 CTA 组成 cluster 走cta_group::2时M ∈ {128, 256}N 为 8 的倍数且最大 256关键在于——支持运行时改写K 由数据类型决定如 FP8/FP6/FP4 混合精度的kind::mxf8f6f4每次 UMMA 吃 K32。Blackwell 的累加器也搬了家不在寄存器文件而在 SM 私有的张量内存TMEM128 行 × 512 列 × 32bit。TMEM 的 128 行与数据通路一一对应第 m 行结果由第 m 条数据通路持有——这就是 M 刚性的物理来源。铁律任何自然行数远小于指令 M 的 GEMM都要在 M 维上空转。图 1 把这行铁律画成了两张脸谱同一个小 token 数的 GEMM摆在 M 维利用率上限 25%swap 到权重侧再落到 N 维则接近打满。后文所有收益都可以回溯到这两张脸谱之差。图 1M 维之笼swap 前后 M/N 维的分配对比本文自制示意图1.2 历史纵深为什么 SwapAB 在 Hopper 萌芽、在 Blackwell 绽放M 的笼子不是 Blackwell 才有为什么 swap 技巧到这一代才出圈把时间线拉长看答案在 N 维身上。Hopper 上社区已经做出了 swap——PR #192 的原文写得明白Swap A B:WGMMA::wgmma(desc_b, desc_a, accum, k)即把权重描述符塞进 wgmma 的 M64 侧、激活塞进 N 侧。它有效但天花板清晰可见wgmma 的 N 是编译期参数expert 尾块不足一个 tile 时只能在软件层补齐动态收缩到 valid_m做不到。这个痛处当时就写在 PR 自己的 TODO 里——LowLatency 场景下 BLOCK_M/BLOCK_N 如何选择仍待探索一句话编译期 N 让尾块问题无解作者只能把它留给未来。Blackwell 把 N 变成了运行时可改写的指令描述符字段desc.n_dim_ umma_n 3再配合 2-CTA 的 TMA multicast 语义swap 才从一个补救动作升级成换维 动态结算 搬运减半的完整三连招。硬件刚性的形状没变变的是刚性的可编程性——这是 SwapAB 故事的隐藏主线。2. 病灶当 token 比行数还少1.2 节给出了一半答案Blackwell 的可编程刚性让 swap 的收益倍增另一半同样关键——真正把矛盾推到嗓子眼的是 DeepSeek 的负载形态。撞墙的位置有两处MoE 解码V3/V4 的 MoE 层有 256 个 expert、top-k8。decode 时 batch 的 token 数本就不大经 router 稀释后单个 expert 摊到的 token 常常只有几十。Grouped GEMM 的每个子 GEMM 以 token 数为行数离 64/128 差着数量级稀疏 MLA 张量并行V3.2 的 DSA 保留 128 个 Q 头的设计TP 切分后每 rank 只剩 128/TP 个头TP4 时 32注意力 GEMM 又一次以小行数撞 M 维。写成算术自然 M32、指令 M128利用率上限 32/128 25%——权重照读、电照耗四分之三的张量核空转。这种空转格外贵原因藏在负载的算术强度里decode 是低算术强度负载本就贴着带宽墙跑Blackwell 的 roofline ridge 点比 Hopper 更靠右同样的 decode 负载在新硬件上离计算上限更远、贴带宽更近图 2——也就是说带宽的账一分省不下来计算的账却在大把白付。图 2decode 的 roofline 位置带宽受限且 ridge 点右移来源thesoftwarefrontier.com补充一句避免误读上面的带宽受限是 decode 小批量下的典型形态不是算子的宿命——batch 足够大时 MoE 同样转入计算受限长上下文 prefill 里的 MLA 也会变重带宽。真实的分界是算术强度不是算子名字。3. SwapAB把最僵硬的维度分给最静态的张量3.1 起源一位社区工程师和一张 H20 基准表先纠正传播中的一个理所当然的误会swapAB 的第一作者不是 DeepSeek 员工。2025 年 9 月社区贡献者Wangzheee向 DeepGEMM 提交 PR #192《support swapAB for m_grouped_fp8_gemm_nt_masked》——目标是 DeepSeek-V3/R1 MoE 解码用的 masked grouped FP8 GEMM测试卡是H20Hopper。开关方式朴素得可爱export ENABLE_SWAPAB1。PR 描述里附了完整基准后文第 4 节将逐行解剖这张表。此刻先看规律M 越小收益越大m32 时两组配置分别提升 66.7% 与 52.5%m96 回落到 26%~47%m160 仍有 56%~61%图 3。一条清晰的画像swapAB 是小 M 专科医生M 越大越不需要它。这个画像在第 3.5 节还会回来收账。这张表还自带一个坦诚的细节作者在 TODO 里写明如何为 LowLatency 场景选 BLOCK_M/BLOCK_N 仍待探索并计划把 H100/H800 的支持补上。它是一份社区 PR 的原始面貌不是厂商白皮书——这种颗粒度的开放是整个故事的地基。图 3PR #192 自带基准可视化。测试条件H20num_groups16对齐 Mn/k 为权重维度数据取自 PR 原文本文制图3.2 手法Blackwell 版把 swap 写进了静态断言先看可验证的一手证据。PR #192 的 diff 里swap 的核心动作只有两行Hopper/WGMMA 路径读者打开 PR 即可自行核对# PR #192 diff 节选Hopper /WGMMA 路径 exportENABLE_SWAPAB1 # 运行期开关默认关闭 WGMMA::wgmma(desc_b, desc_a, accum, k); # 描述符对调权重上 M64 侧激活上 N 侧一行环境变量加一行描述符对调就是 swapAB 的初始形态。传统摆法里 A 激活行 token、B 权重行 输出特征——直觉来自训练batch 维天然是行。SwapAB 摆法A 权重B 激活。权重的特征维1536/4096/7168……去对齐 MMA 的 M 硬约束token 数落到灵活的 N 维。到 DeepSeek-V4 的 MegaMoE SM100 kernel 里这条规则已固化// deep_gemm/include/deep_gemm/impls/sm100_fp8_fp4_mega_moe.cuh // MMA 配置固定 swap A/BA 为权重维、B 为激活维2-CTA UMMAK-major constexpruint32_t LAYOUT_AD_M 128;// 单 CTA 的 TMEM 布局高度 constexpruint32_t UMMA_M LAYOUT_AD_M *2;// 2562-CTA 发射后的有效高度 constexpruint32_t UMMA_N BLOCK_M;// Swap 后N 维 expert token 块 constexpruint32_t UMMA_K 32;// mxf8f6f4 单条 UMMA 的 K 宽 constexpruint32_t LOAD_BLOCK_M BLOCK_M /2;// A 做 multicast 后单 CTA 只读一半 DG_STATIC_ASSERT(BLOCK_M %160,Invalid block M); DG_STATIC_ASSERT(BLOCK_N LAYOUT_AD_M,Invalid block N);// 权重维必须 128 DG_STATIC_ASSERT(BLOCK_K 128,Invalid block K);代码出处说明以上及 3.3 节的代码片段转引自第三方源码解析langcopilot.com 对 DeepSeek-V4 MegaMoE 的系列分析尚未在公开 DeepGEMM 仓库的一手代码中逐字核对——截至发稿无法在公开仓库中检索到对应文件或因版本迭代或该文件未公开。该 kernel 的存在性有独立旁证vLLM PR #53040 描述将 V4 的 replicated FP8 shared expert 融合进 DeepGEMM 的 persistent SM100 MegaMoE kernelSGLang 的 release notes 亦提及DeepSeek V4 的 DeepGEMM 与 FlashMLA kernel含 MegaMoE。引用时请以思路级证据而非仓库级证据对待。同样需要说明版本关系避免读者混淆PR #192 属于 V3/R1 时代的 Hopper kernel上面这段代码来自V4 时代的 Blackwell MegaMoE kernel——同一思路在两代架构上的两代实现。V4 已随 TRT-LLM 的 DeepSeek V4 Pro QA 测试用例与大量第三方部署公开存在。3.3 第二招N 维动态收缩到 valid_mSwap 之后的意外之喜N 维不仅容忍小尺寸还能运行时改写。expert 尾块不足BLOCK_M时直接改指令描述符// mma/sm100.cuhUMMA_N 以 8 为编码单位写入指令描述符 voidupdate_instr_desc_with_umma_n(UMMA::InstrDescriptorBlockScaled desc, constuint32_t umma_n){ desc.n_dim_ umma_n 3;// 有效 token 数 valid_m其余由 TMA OOB 补零忽略 }// K 循环内发射注意 fma 的 a/b 描述符与 SFA/SFB 槽位全部对调 ptx::SM100_MMA_MXF8F6F4_2x1SM_SS::fma( b_desc, a_desc,// ← SwapABb_desc 在前 accum_stage_idx * UMMA_N, k_block_idx 0|| k 0,// 首条 UMMA 用 scale::0 顺便清零累加器 runtime_instr_desc, kTmemStartColOfSFB, kTmemStartColOfSFA);尾块不再补到 128 再算——算多少写多少空转税从按 128 收变成按 8 的粒度收。这一招是 Blackwell 独有的Hopper 的 wgmma 里 N 是编译期立即数想做也做不到1.2 节那行 TODO 就是它的墓志铭。3.4 第三招TMA multicast 顺势白嫖2-CTA UMMA 的硬件约定操作数 A 由两个 CTA各持 M 的一半操作数 B 由两个 CTA共享完整 N。Swap 前 A激活两个 CTA 各读不同的 token 行Swap 后 A权重两边读同一块权重——权重加载天然走 TMA multicasttma::copyBLOCK_K, LOAD_BLOCK_N, kSwizzleBMode, b_dtype_t( tensor_map_b_ptr, full_barriers[stage_idx], smem_b[stage_idx], k_idx, n_idx,2);// ★ num_tma_multicast 2一次加载广播两半集群权重流量减半累加器按列 tokens落在 TMEM、由 epilogue 按列取数。对权重远大于激活的 MoE这一页是带宽账本上实打实的进账。图 4swap 后的 2-CTA 数据流共享权重、TMA multicast、TMEM 按列存 token本文自制示意图3.5 代价与边界五笔账SwapAB 不是开关一拨就万事大吉这笔账至少有五笔布局锁死BLOCK_N必须等于 128DG_STATIC_ASSERT权重在 SMEM 的 swizzle 模式、TMA 描述符的构建都随之定死kernel 的通用性让位给特化大 M 时收益消失甚至倒贴PR #192 自己的数据表明 m1024 的对齐场景仅 15%——M 本来就满时 swap 只是增加约束。因此 adaptive kernel selection 是这套技巧不可分割的一半而且阈值随硬件而变TRT-LLM PR #4430 的 host 端选择逻辑在 H2078 个 SM上是 m_per_expert 64 启用 swap在 H100/H200 上则收紧到 32——切换线同时取决于问题形状与 GPU 的 SM 数量特化的组合爆炸每增加一个 (数据类型 × layout × swap 开关 × tile 形状) 组合JIT 就多编译一个 kernel 变体。DeepGEMM 在 PR #112 之后以低 CPU 开销的 JIT 编译缓存应对但面对 serving 端千变万化的 shape 流变体数量与编译耗时的平衡仍是部署中的现实约束TMEM 列占用swap 后累加器按列存放、列数tokens。N 维过大的 batch 会吞掉 TMEM 的 512 列配额挤压流水线的双缓冲余量全链路重构成本TMA→SMEM→UTCCP→UMMA→TMEM→epilogue 每一环的地址推导都要按列即 token重写epilogue 里SM100_TMEM_LOAD_16dp256b1x的列偏移链这也是模仿者多、做对者少的原因。4. 数字的口径同一个技巧五组数字SwapAB 的收益数字散落在一年半的传播链上几乎每一组都带着不同的口径字段。把它们并排放在一起比任何单点数字都更能说明问题。第一组PR #192 的 H20 基准2025-09kernel 级Hopper。前文已经见过这张图这里把六行数据完整列出H20num_groups16对齐 M配置OriginalSwapAB提升m32, n4096, k7168117 TFLOPS195 TFLOPS66.7%表内最高m32, n7168, k2048118 TFLOPS180 TFLOPS52.5%m96, n4096, k716818126647.0%m96, n7168, k204818823625.5%m160, n4096, k716816326361.3%m160, n7168, k204815924855.9%规律依旧M 越小收益越大而一旦 M 对齐m1024收益立刻回落到 15%——这组数字天生属于小 M 场景图 3见 3.1 节。第二组vLLM 集成实测2025-06 起B200特定 batch size。vLLM 贡献者 Wentao Ye 的开发日志记录#19820 初次集成报 ~40% GEMM perf improvementDeepGEMM SM100 重构上游 PR #112带来 breaking 变更初版集成随之 revert 并重提最终以 #20087 合入——表述也谨慎地改为 特定 batch size 下~40%。同一段旅程的后续 PR 还报告了 ~15% 的端到端提升。限定词本身就是信息一个 batch size 敏感的技巧离开它的主场数字立刻缩水。第三组NVIDIA blog152026-09-26GB300厂商口径。FP8 稀疏 MLA 吞吐 47.03%Top-K 索引 7.41×kernel fusion 端到端 32.64%~64.20%。这是全链条中口径最完整的一组厂商博客不需要取悦任何一方因此同时给出了吞吐、端到端与单项数据。第四组vLLM E2Ev0.10.21.5%。这组数字提醒kernel 收益 × 负载占比 ≠ 端到端收益。摊薄是非线性的随 batch size 剧烈变化、存在木桶效应GEMM 提速后瓶颈立即转移到访存与调度、kernel 收益本身还受 L2 命中与带宽余量约束。1.5% 是全部约束共同作用后的诚实结果。第五组DeepGEMM README对基线 1.1×–1.2×。官方 README 的保守标注与 PR 里的 66.7% 并存同一事实的两种文体——面向潜在用户的 README 需要数字长期站得住面向评审的 PR 只需要数字当下站得住图 6。五组数字五套口径硬件不同H20 / B200 / GB300、层级不同kernel / 端到端、来源立场不同社区 PR / 集成方日志 / 厂商博客 / 官方 README。图 5 以灰线分列。这组对照提供了本文最想留下的读数纪律读任何提速 X%之前先问三个字段——什么硬件什么层级谁说的图 5同一个技巧五组数字口径分列本文自制。判读三字段什么硬件 · 什么层级 · 谁说的图 6DeepGEMM 官方 README 基准表对基线的保守标注来源DeepGEMM README经 MarkTechPost 转载5. 生态回响一个社区 PR 的上游漂流SwapAB 的传播路径是观察开源基础设施博弈的最佳样本2025-02-13DeepGEMM 开源2025-07-20PR #112 完成 SM90/SM100 双架构重构SM100 支持四种 layout 与 UE8M0 packed 缩放2025-09PR #192Wangzheee为 Hopper masked grouped GEMM 加入 swapAB——本文主角登场身份是社区贡献者而非 DeepSeek 员工2025-12-24TRT-LLM PR #4430 合并据 PR 页面记录将同一思路引入 NVIDIA 官方栈FlashInfer 亦提供 swapAB 风格 kernelvLLM 路线图明确提及整合2026-02-13vLLM 官方博客《DeepSeek-V3.2 on GB300: Performance Breakthrough》发布实测数据图 82026-09-26NVIDIA blog15 把该技巧命名为SwapsMmaAb单列一节披露稀疏 MLA 47.03%、Top-K 7.41×、kernel fusion 最高 64.20% E2E2026-04-27DeepGEMM issue #320社区工程师公开追问 mega kernel 与 w4a8 grouped GEMM 中 swapab 的收益来源——技巧的工程讨论已经日常化这是传播广泛可以坐实的形态2026-09-30DeepGEMM-Ascend 发布FlashMLA 仓库同日新闻宣布昇腾 950 NPU 版稀疏算子prefill/decode 410/360 TFLOPS标称 95%/83% 峰值并附技术报告——换维对齐硬件约束的思路完成跨硬件验证。图 7SwapAB 的漂流史红Hopper 起源绿Blackwell 官方正名本文自制横向看竞争格局同一技巧在生态里有三种实现姿态DeepGEMM 走 JIT 路线——按 shape 现场编译、启发式选择是否 swap3.5 节的组合爆炸账极限性能最高、编译有开销FlashInfer 走预编译库路线——swapAB 风格 kernel 开箱即用牺牲一点峰值换确定性CUTLASS 在 SM100 collective builder 的模板参数层面支持 A/B 角色互换UMMA atom 不区分语义只认 layout是重剑无锋的通用解。vLLM/SGLang 在调度层按问题形状在三家之间选型——技巧的扩散终点是它变成所有人的默认选项。纵向看 DeepSeek 自身的技术版图还有一条常被忽略的路线分野DeepGEMM/FlashMLA 代表的 PTX 级手写 kernel与 DeepSeek 开源的 tile 级 DSL TileLang 代表的生产力路线。SwapAB 能否在 DSL 中表达这个问题其实已有初步答案SGLang PR #15763 的 TileLang 实现明确标注 inspired by TRT-LLM #4430 and DeepGEMM #192在 RTX 4090 上对比 Triton 报 5.77×、在 H20 上对比 DeepGEMM 报 1.42×数字与硬件组合据 PR 标注。真正的问题随之升级DSL 表达后的性能天花板能否在更多硬件上追平手写内核这决定了手写派 vs 编译派的分界线往哪边移。再看为什么是 DeepSeek。先交代论证性质以下不是有访谈或内部文件背书的结论而是本文基于公开事实一个社区 PR 的走向提出的解释框架——姑且称之为广场模型——供读者检验。它要回答三个问题。其一DeepSeek 与 Meta/Google 的基础设施开源有何本质区别一个可能的解释是Meta 开源模型但 kernel 栈如 FBGEMM深绑推荐场景、与模型发布脱节Google 的 XLA 是编译器通用路线不为自家模型的独特架构做手工特化。而 DeepSeek 让模型与 kernel 作为一套东西同时开源DSA、细粒度 MoE 的每一个非常规设计几周内就会在 DeepGEMM/FlashMLA 里长出对应的算子——如果广场模型成立这正是广场有人的原因。其二这个广场收得到租金吗本文的假说是租金的来源是负载的独特性——只要下一代架构继续由 DeepSeek 定义非常规适配的先发权就留在广场中央反之若某代模型退回常规 GQA 稠密架构租金即刻归零。这是该策略的内在风险而非外部威胁也可以反过来读作 DeepSeek 坚持架构创新的动机之一。其三Wangzheee 为什么把 PR 交给 DeepGEMM 而非 CUTLASS/FlashInfer作者未留下动机自述但 PR 文本本身提供了推断依据他动手的 API 是 DeepGEMM 的 masked grouped GEMM、测试卡是 H20——前者是 DeepSeek MoE 推理的原生 API后者是其主力部署硬件二者本就配套PR 合入后数周内就能经由 vLLM/SGLang 的依赖进入真实生产集群而对 CUTLASS漫长的评审周期足以让技巧错过一代模型的部署窗口。若此推断成立贡献者用脚投票投的是哪条路径最短地抵达算力账单。6. 定位软硬协同设计的第三层SwapAB 值得写进教科书因为它示范了 co-design 的第三个层次第一层模型架构迁就硬件——MLA 把 KV 压成 576 维潜向量本质是给 HBM 带宽松绑第二层编译与调度软化硬件——persistent kernel、CUDA Graph、Wave 调度把碎片负载拼成规整流第三层改写指令语义层的题面——不搬数据、不改算法、不动精度只在指令描述符里交换 A/B 的含义让硬件约束从绊脚石变成顺风车。PR #192 → SwapsMmaAb 的漂流恰好是第三层的两次登台Hopper 上的版本换维证明了题面可改Blackwell 上的版本换维 动态结算 multicast证明了题面可被持续改写。它的出现有两个前提其一新指令把 GEMM 的题面参数M/N 形状、操作数角色、缩放因子槽位暴露得足够细其二开源 kernel 社区把读 PTX 写 kernel变成了模型公司与社区工程师的标准能力。推广一步凡自然行数小于硬件指令 M 的负载——MoE decode、稀疏 MLA、投机解码的 draft 验证、小 batch LoRA serving——都存在同构机会。反过来若下一代硬件把 M 维也做灵活、或把 multicast 做进指令语义这个技巧就会自然消亡。好技巧的价值正在于它精确瞄准了某代硬件的刚性。7. Rubin 前夜结语2026 年 4 月的 MLPerf Inference v6.0 上GB300 NVL72 跑 DeepSeek-R1 交出了 8064 token/s/GPU较 v5.1 提升 2.77 倍。几个月前vLLM 官方博客的 GB300 实测2026-02-13则给这种乐观浇了一盆清醒的冷水DeepSeek-R1NVFP4EP2prefill 峰值达 22476 TGS较 H200 提升约 8 倍而 V3.2 的稀疏注意力在现行栈里尚未兑现纸面优势——indexer 路径使单层 kernel 耗时达到 MLA 的 2.7 倍TTFT 反而比 R1 差约 55%官方小标题直言 Still Way To Go图 8。kernel 级技巧的合力再强也还补不上架构新原语的适配欠账——这是 SwapAB 故事最诚实的注脚也是一切 kernel 级乐观最需要的对冲。图 8vLLM 官方 GB300 实测R1 与 V3.2 的性能对比来源vLLM 官方博客2026-02-13V3.2 差距分析见 Still Way To Go 一节而 NVIDIA 官方博客为 R1 创纪录推理性能配的宣传图从另一个角度记录了 Blackwell 的绝对水位B200 NVL8 将 DeepSeek-R1 671B 的单用户输出速度推到 253 token/sH200 NVL8 为 54图 9。SwapAB 们不会出现在任何一张宣传图的标题里但它们藏在这些柱子的每一个像素里。图 9B200 NVL8 上 DeepSeek-R1 671B 单用户输出速度纪录来源NVIDIA 官方开发者博客于是可以给 SwapAB 一个总体评价了。它指向的变革是真实且深刻的当社区工程师在 Hopper 上写下第一行wgmma(desc_b, desc_a, …)当 DeepSeek 把这条思路写进 Blackwell kernel 的静态断言当芯片厂商为它命名并写进官方博客——AI 基础设施的竞争就进入了纳米级时代胜负不再只由 FLOPS 和显存带宽决定也由谁更懂得在 128 行的 TMEM 里安置一个不满 32 行的张量决定。SwapAB 终将在某代硬件上消亡——所有瞄准刚性的技巧都会。它留下的不是技巧而是一种提问方式在宣布算力提升之前先数清楚这 128 行里究竟坐着多少 token。创作不易禁止抄袭转载请附上原文标题及其链接。参考资料DeepGEMM PR #192《support swapAB for m_grouped_fp8_gemm_nt_masked》作者 Wangzheee2025-09——全部 H20 基准数字与 swap 的 WGMMA 实现export ENABLE_SWAPAB1、WGMMA::wgmma(desc_b, desc_a, accum, k)github.com/deepseek-ai/DeepGEMM/pull/192NVIDIA TensorRT-LLM Tech Blog 15:Optimizing DeepSeek-V3.2 on NVIDIA Blackwell GPUs2026-09-26——SwapsMmaAb、FP8 稀疏 MLA 47.03%、Top-K 7.41×、kernel fusion 32.64%–64.20% E2E。TRT-LLM PR #44302025-12-24 合并据 PR 页面记录——swapAB 引入官方栈其 host 端选择逻辑H2078 SM于 m_per_expert64 启用H100/H200 于 32 启用阈值细节转引自社区核查刊出前建议对照 PR diff 复核。DeepGEMM 仓库 README、News 与 Issues——2025-07-20 PR #112SM90/SM100 全量重构、issue #3202026-04-27社区对 mega kernel 与 w4a8 grouped GEMM 中 swapab 收益的追问、2026-09-10 #432、2026-09-30 #462 与 DeepGEMM-Ascend。vLLM 官方博客《DeepSeek-V3.2 on GB300: Performance Breakthrough》2026-02-13DaoCloud vLLM team——图 8 实测数据R1 prefill 22476 TGS、V3.2 7360 TGS、DSA 单层 kernel 耗时为 MLA 的 2.7 倍、R1 TTFT 较 V3.2 低约 55%vLLM v0.10.2 Release Notes——DeepGEMM Linear 1.5% E2E。B200 GEMM ~40% 出处vLLM 贡献者 Wentao Ye 开发日志yewentao256.github.io / wentao.site——#19820初次集成~40%与 #20087适配上游 PR #112 breaking 更新后合入特定 batch size 下 ~40%初版集成曾随 breaking 变更 revert/重提社区记录 #20049。vLLM PR #53040——将 V4 的 replicated FP8 shared expert 融合进 DeepGEMM 的 persistent SM100 MegaMoE kernelMegaMoE kernel 存在性的旁证SGLang release notes——DeepSeek V4 的 DeepGEMM 与 FlashMLA kernel含 MegaMoE。DeepSeek-V4 MegaMoE 源码解析langcopilot.com2026-05-15——3.2/3.3 节代码片段的转引来源未经一手仓库逐字核对见正文出处说明。FlashMLA 仓库 README2026-09-30 更新——B200 稀疏 MLA 1350/1024 TFLOPS昇腾 950 NPU 版 410/360 TFLOPS据仓库新闻表述。thesoftwarefrontier.com《The Blackwell Migration Question》——图 2 roofline 推导。SGLang PR #15763——TileLang 版 swapAB 实现标注 inspired by TRT-LLM #4430 and DeepGEMM #192RTX 4090 上 5.77× over Triton、H20 上 1.42× over DeepGEMM据 PR 标注MarkTechPost 对 DeepGEMM 的报道2025-02-25——图 7 官方 README 基准表转载NVIDIA 开发者博客《NVIDIA Blackwell 实现世界纪录的 DeepSeek-R1 推理性能》——图 9MLPerf Inference v6.0 GB300 NVL72 8064 token/s/GPU2026-04ithome 报道。