## 一、先从一个"跑偏了"的优化说起

2026 年 3 月，Tri Dao 和 Colfax、Meta、NVIDIA 的人一起发了 FlashAttention-4（arXiv:2603.05451，3 月 5 日）。论文里有一句话很值得琢磨（来源：FlashAttention-4 论文及 Colfax Research 解读博客，2026-03）：

> 一个朴素的观点可能认为，GEMM 的速度完全决定了 attention 内核的性能，其他部分至少在一阶近似下可以忽略。但在 B200 上做"喂料与速度"分析，结论恰恰相反：主要瓶颈不在张量核心做 MMA 有多快，而在于（a）前向计算时 softmax 指数运算用的 SFU 单元，（b）反向计算时的共享内存流量。

翻译成人话：**你以为 GPU 在忙着算矩阵，其实它有一半时间在排队等指数运算、等共享内存倒腾数据。**

为什么会这样？因为硬件是"偏科"长大的。从 Hopper H100 到 Blackwell B200，BF16 张量核心吞吐从 1 PFLOPS 涨到 2.25 PFLOPS，涨了 2.25 倍；但 **SFU 数量没变，共享内存带宽也没变**。作者给这种现象起了个名字，叫**非对称硬件扩展（asymmetric hardware scaling）**——算力单元一路狂奔，其他配套设施原地踏步。

这件事的意义超出 attention 本身。它意味着：**你从 H100 换到 B200，代码一行不改，内核大概率不会快 2.25 倍，因为卡住你的那个东西根本没升级。**

## 二、Roofline：一张图分清"算不完"和"等数据"

判断卡在哪，最趁手的工具是 **Roofline 模型**。它把硬件的两条性能上限画在一张图上：一条水平线是峰值算力，一条斜线是峰值内存带宽，两条线组成一个"屋顶"。你的内核是坐标上的一个点，落在斜线下面就是内存受限（等数据），贴到水平线就是计算受限（算不完）（来源：*AI Systems Performance Engineering* 第 6 章 6.13 节）。

![fig12-roofline.png](https://developer.qcloudimg.com/http-save/yehe-11479707/12033061331f2fc0baea9ce83015d392.png)


图里两条线的交点叫**脊点（ridge point）**，对应的横坐标就是"算术强度"的门槛。算术强度的定义很朴素：**每从显存搬 1 字节数据，能干多少浮点运算（FLOPs/byte）**。

原书举了个特别有说服力的例子：一个内核读两个 FP32 浮点数（8 字节），做一次加法（1 FLOP），写回一个 FP32（4 字节）。一共搬了 12 字节，干了 1 个 FLOP，算术强度是 **0.083 FLOPs/byte**。而一块 Blackwell 级 GPU 的脊点在 **10 FLOPs/byte** 左右（约 80 TFLOPS ÷ 8 TB/s）。

差多少？**约 120 倍**——原书写"100 倍以上 / 两个数量级"，120 是精确算出来的。

这个内核的算术强度比门槛低了两个数量级，它连算术逻辑单元的牙缝都塞不满。向量加法、激活函数、LayerNorm、逐元素乘——所有"读一遍写一遍"的操作，算术强度都在 0.1 这个量级，全是内存受限。反过来，一个大尺寸的矩阵乘法，数据搬进来要在片上被复用成百上千次，算术强度能轻松过百，那就是计算受限。

Nsight Compute 里怎么分辨？原书写得很直接：**内存受限的内核，DRAM 带宽利用率非常高，同时算力类指标（比如 ALU 利用率）很低**，说明 warp 大部分时间卡在访存上；而 warp stall reason 里会显示大量的内存依赖停顿。

大白话版：算术强度就是"出一趟差能办多少事"。你跑一趟 HBM 要几百个时钟周期，如果只为了加两个数就回来，那时间全耗在路上了；聪明做法是搬一批数据到片上，在片上把能算的都算完再回去。

## 三、内存层次：一步慢，处处慢

既然关键是把数据留在片上，就得知道片上到底有多快。Blackwell 的内存层次是这样的（来源：*AI Systems Performance Engineering* 第 6 章表 6-5）：

| 层级         | 容量                              | 延迟             | 带宽（约）            |
| ---------- | ------------------------------- | -------------- | ---------------- |
| 寄存器        | 每 SM 64K 个 32 位寄存器（每线程最多 255 个） | 单周期，基本免费       | 每 SM 数十 TB/s     |
| TMEM       | 每 SM 256 KB（张量核心专用）             | ~10 周期         | TB/s 级，直连张量核心    |
| 共享内存 / L1  | 每 SM 228 KB（可用 227 KB）          | ~20-30 周期      | TB/s（无 bank 冲突时） |
| 常量缓存       | 每 SM ~8 KB                      | ~1 周期（warp 广播） | TB/s 级           |
| L2         | 全 GPU 126 MB                    | ~200 周期        | 多 TB/s 聚合        |
| 全局内存 HBM3e | B200 最多 180 GB（B300 约 288 GB）   | 数百 → 1000+ 周期  | ~8 TB/s          |

![fig12-mem-hierarchy.png](https://developer.qcloudimg.com/http-save/yehe-11479707/c5760c600b17f7911f40a87bc1fba95a.png)

看最后两行的落差：**从共享内存到 HBM，延迟从几十周期跳到几百上千周期，差了一个到两个数量级；带宽侧则是另一种故事——共享内存每 SM 是 TB/s 级的吞吐，而 HBM 那 8 TB/s 是整张卡共享的封顶值，几千个线程来抢。**

这里有个隐蔽的坑：**寄存器溢出（register spilling）**。每个线程最多 255 个寄存器，如果你的内核局部变量太多、或者融合得太狠，编译器会把溢出的变量丢进"本地内存"——而本地内存就 backing 在 DRAM 上，延迟直接回到数百周期。你以为自己在用最快的存储，其实早就掉回最慢那一层了。这也是为什么 Nsight Compute 会专门追踪 spill 指标。

还有个 Blackwell 的新成员值得认识：**TMEM**。它是每 SM 256 KB 的专用片上内存，直连张量核心，延迟约 10 周期，不在 CUDA C++ 里直接指针寻址，而是由 TMA（Tensor Memory Accelerator）硬件搬运数据。你可以把它理解成"给张量核心开的专属备料台"：TMA 负责把数据从全局内存搬到共享内存，张量核心指令（tcgen05）再在共享内存和 TMEM 之间隐式倒腾，全程不回 HBM，人不用手动管。

顺带一提，B200 看起来是一块 GPU，实际是两个受光刻限制的裸片通过 10 TB/s 的片间互连拼起来的，每边挂 4 个 HBM3e 堆栈，共 8 个。地址空间对开发者是统一的，看到性能 profile 时注意一下"跨裸片访问"可能有差异就行。

## 四、合并访问：一个 stride 把带宽打到 25%

知道了层次，下一步是"怎么搬才不亏"。这里有个原书第 7 章的实测案例，数字挺吓人。

现代 GPU 的缓存行是 **128 字节**，由 **4 个 32 字节扇区**组成。如果一个 warp 的 32 个线程访问连续且 128 字节对齐的地址，硬件会合并成最少的 4 个扇区，一次拿完。但如果线程用跨步索引（比如 stride=2），或者起始地址没对齐，请求就会跨缓存行，扇区数膨胀。

原书实测：跨步版本的每个 warp 请求被拆成平均 **8 个独立的 32 字节扇区**，是理想值的两倍。结果是 **Nsight Compute 只报告约 25% 的持续峰值 DRAM 吞吐**——四分之三的带宽白扔了。打个比方：32 个人打饭，挨着拿一次端走；隔一个拿一次，同样的饭得跑八趟，多出来的七趟全在排队。

对应的诊断指标也有明确阈值（来源：同书第 7 章）：

- **Global Memory Load Efficiency** 明显偏低
- **DRAM 扇区读取计数**偏高
- **平均每请求扇区数（sectors per request）超过 4.0**
- DRAM 吞吐百分比远低于峰值

在 PyTorch 里这个坑很好踩：一个带跨步索引的 `torch.index_select`，底层 gather 内核就会发出非合并加载——32 个线程访问的地址相隔 `stride * 4` 字节，硬件没法合并，只能发 32 次独立加载。

修法通常很简单：让每个 warp 的 32 个线程读连续元素，用 `idx = blockIdx.x * blockDim.x + threadIdx.x` 这种线性索引；或者把数据布局从 AoS（结构体数组）换成 SoA（数组结构体），让线程 i 永远访问元素 i。改完再看 Nsight Compute，效率指标会立刻回升。

同类的坑还有**共享内存 bank 冲突**——共享内存像 32 个并行收银台，多个线程同时挤同一个台，请求就被串行化，TB/s 级的吞吐说没就没。原书给的规避办法是填充（padding）或交叉（swizzle）数据。这两类问题都属于"代码看起来没错，但带宽只剩零头"。

## 五、提算术强度的三板斧 + 一个偏方

判断清楚、访问也合并了，接下来才是真正提性能。手段就三类，都是围绕"每搬一个字节多干点活"。

![fig12-arith-intensity.png](https://developer.qcloudimg.com/http-save/yehe-11479707/a388147348f383a0e8cfeb90555c2c02.png)


**第一斧：分块复用（tiling）。** 把 A、B 的小子矩阵搬进共享内存，让每个字节在片上被反复用。极端一点算：一个 32×32 的块加载到共享内存后，每块贡献 1024 次乘加（32×32），相当于平均每个元素被复用 **32 次**。相比每次都从 DRAM 取数，算术强度直接乘上复用因子。Blackwell 上还有个更狠的招：**线程块集群 + TMA 多播**，2×2 的四个 CTA 协作加载同一个块，全局内存流量最多能降到 1/4。Blackwell 支持到 16 个 CTA 的集群，但要注意——集群越大占用率越低，开之前先测（来源：同书第 9 章）。

**第二斧：内核融合（fusion）。** 把多个操作合成一个内核，中间结果留在寄存器或共享内存里，不写回 HBM。原书给了个直观对比：

- 不融合：4 FLOPS 对应 36 字节的加载与存储（中间结果来回倒腾）
- 融合后：4 FLOPS 只对应 12 字节

同样的活，内存流量少了三分之二，算术强度是原来的 3 倍。这也解释了为什么 TorchInductor 之类的编译器那么热衷于自动融合逐元素操作——收益太直接了。不过融合不是无脑堆，寄存器压力上去会触发溢出，反而吃亏，所以原书提醒"一定要分析融合后的内核"。

**第三斧：降精度。** 这是最便宜的杠杆。同一个 128 字节事务，能装 32 个 FP32、64 个 FP16、128 个 FP8 或 256 个 FP4。把 FP32 换成 FP16，每字节干的工作量**立刻翻倍**。换个出差的比喻：之前一趟车装 32 个 FP32，现在换装 64 个 FP16——路费一样，每趟办的活翻倍。降精度之所以最划算，就因为它几乎不动代码结构。Blackwell 甚至引入了硬件解压缩——权重可以用比 FP4 更激进的方案压着存在 HBM 里，加载时硬件现场解压，等于凭空多出一部分有效带宽。

**偏方：重计算（recomputation）。** 内存带宽是瓶颈时，宁可重新算，也别存了再读。比如 `x*x` 这种廉价表达式，在寄存器里重算一遍通常比从显存把之前算好的 `x*x` 读回来快。放大到模型层面就是**激活检查点（activation checkpointing）**：反向传播要用到的激活值不存了，到时候重算，用一点额外 FLOPS 换掉大量的显存占用和访存流量。推理引擎也常用这招省显存、给更大 batch 腾空间（来源：原书第 9.5 节）。

**我的实践排序**：先做降精度（改一行 `torch.set_float32_matmul_precision('high')` 就能看到效果），再做融合（TorchInductor 自动），最后才动分块（要重写调度）。前三板斧吃下大部分收益，分块属于拼最后那 10%。改了之后记得复测——可能动了精度，数值稳定性要先观察。

## 六、Rubin 来了，墙更高了

讲到这里，你可能会想：新一代硬件总该把带宽提上来了吧？提了，但没提够。

看 NVIDIA 官方数字（来源：NVIDIA 技术博客 *Inside NVIDIA Rubin GPU Architecture*，2026）：Rubin 用 HBM4，接口宽度相对 HBM3e 翻倍，单 GPU 提供**最多 288 GB 容量、最高 22 TB/s 带宽**，是 Blackwell 的 **2.8 倍**；同时第三代 Transformer Engine 提供**最高 50 PFLOPS 的 NVFP4 推理算力**。

| 维度             | Blackwell B200        | Rubin              | 倍数                              |
| -------------- | --------------------- | ------------------ | ------------------------------- |
| HBM 带宽         | ~8 TB/s（HBM3e）        | 最高 22 TB/s（HBM4）   | 2.8×                            |
| 显存容量           | 180 GB（B300 约 288 GB） | 288 GB             | 1.6×                            |
| NVFP4 推理算力     | ~10 PFLOPS（dense）     | 50 PFLOPS（稀疏）      | dense 对 dense 约 3.5×，稀疏口径标称约 5× |
| NVLink（双向/GPU） | 1.8 TB/s              | 3.6 TB/s（NVLink 6） | 2×                              |

（来源：NVIDIA 官方技术博客，2026；不同口径下数字会有差异）

留意 NVIDIA 给 Rubin 标的 50 PFLOPS 是 **NVFP4 稀疏口径**；B200 对应数字是 dense FP4，dense 对 dense 的实际倍数约 3.5×。但不论按哪个口径看，**算力的涨幅都明显大于带宽**——这意味着同一份代码搬到 Rubin 上，掉进内存受限区的概率反而更高：想喂饱那些张量核心，每个字节得干比以前多一截的活。

## 七、2026 的四种破墙姿势

既然墙在加高，2026 年最前沿的内核优化本质都是跟它较劲。FlashAttention-4 论文实际上列了四招，最关键的第一招是**重设计的全异步 MMA 流水线**——利用 Blackwell 完全异步的 MMA 指令和更大 tile，让 tensor core 与 softmax 指数运算、内存搬运三件事真正并行起来。这是后面三招能见效的基础。

第二到第四招是：

1. **软件模拟指数函数**：既然 SFU 是瓶颈、而 FMA 单元充裕，那就用多项式近似（Cody-Waite + 三次多项式）在通用算力单元上算 exp，只覆盖 10-25% 的指数运算元素，把稀缺 SFU 解放出来。
2. **条件 softmax 重缩放**：online softmax 通常每次运行最大值变化都要重缩放，FA4 只在变化大到威胁数值稳定时才做，据 Tri Dao 在 Hot Chips 的分享，**重缩放操作减少约 10 倍**。
3. **2-CTA MMA + TMEM 存中间结果**：反向传播把中间结果放 TMEM 缓解共享内存流量，再用 Blackwell 的 2-CTA MMA 模式让一个 UMMA 跨 CTA 对执行，进一步减少共享内存流量，同时把原子归约砍半。

结果：B200 BF16 下前向峰值 **1613 TFLOPS**，硬件利用率 **71%**，比 cuDNN 9.13 快最高 **1.3 倍**，比 Triton 快 **2.7 倍**（来源：FA4 论文 arXiv:2603.05451，2026-03-05）。另外 FA4 整个用 CuTe-DSL（Python 内嵌的 DSL，属于 CUTLASS）实现，编译时间比传统 C++ 模板快 **20-30 倍**——内核开发从写 C++ 模板变成写 Python，迭代速度上来才有空间反复试错（比如 FA4 一连试了多种 softmax 重缩放策略，普通工程节奏根本撑不住）。论文同时提到 NVIDIA 后续已把部分 FA4 技术吸收进新版 cuDNN。

再讲个"土法炼钢"的真实案例，能让你对内存层次有多敏感。DeepSeek 的 DeepEP 专家并行库里用了这么一条定制 PTX 指令：

```
ld.global.nc.l1::no_allocate.l2::256b
```

意思是：发起非一致的全局加载，**不在 L1 分配**（`l1::no_allocate`），而是以 **256 字节块直接灌进 L2**（`l2::256b`）。为什么？因为专家并行的 all-to-all 通信内核每次调度都只读一遍大块连续缓冲区，如果这些加载走 L1，会把 SM 正在频繁使用的缓存行挤出去。绕开 L1 直取 L2，就保护了热数据。

严格地说，这几个修饰符本身在 NVIDIA PTX ISA 文档里能找到，但 DeepSeek 把它们组合起来精确控制 H800 受出口限制版本的缓存行为，被 DeepEP README 标为"未定义/文档外用法"。风险也很实在：不保证跨代稳定，所以 DeepEP 里专门留了构建开关 `DISABLE_AGGRESSIVE_PTX_INSTRS=1`，出问题可以一键关掉（来源：*AI Systems Performance Engineering* 第 9 章）。

这故事说明两件事：一是内存层次的行为值得抠到指令级；二是非常规手段要配逃生开关。

## 八、给你的排查清单

把这篇和上一篇的四步体检接起来，落到单卡内核上，顺序大概是这样：

1. **先算算术强度**：你的内核每搬 1 字节干多少活？对照脊点（Blackwell 约 10 FLOPs/byte），先判断是在"等数据"区还是"算不完"区。方向错了，后面全白做。
2. **看内存层次有没有用起来**：Nsight Compute 里查寄存器溢出、共享内存命中率、L2 命中率。溢出到 local memory 是最容易被忽略的性能杀手。
3. **查合并访问**：sectors per request 是不是超过 4.0？Global Memory Load Efficiency 低不低？PyTorch 侧重点查 `index_select` 之类的 gather。
4. **上三板斧**：降精度 → 内核融合 → 分块复用。优先级按这个顺序，因为前两个几乎不动代码结构，分块要重写调度。
5. **带宽还是紧就考虑重计算**：廉价表达式重算、激活检查点，用算力换带宽。
6. **改一处测一处**：每次改动后用同一个 ncu 命令复测，看内核在 Roofline 上是往右移了（强度提高）还是往上移了（更接近屋顶）。

---


2026 年 3 月，Tri Dao 和 Colfax、Meta、NVIDIA 的人一起发了 FlashAttention-4（arXiv:2603.05451，3 月 5 日）。论文里有一句话很值得琢磨（来源：FlashAttention-4 论文及 Colfax Research 解读博客，2026-03）：

AI Infra 系列 · 第六章 GPU 性能工程（二）：B200 算力涨 2.25 倍，运数据的路一寸没宽，你的内核其实在等数据

这篇就讲清楚：怎么用 Roofline 判断你是"算不完"还是"等数据"，Blackwell 的内存层次差多少倍，一个 stride 怎么把带宽打到 25%，以及 2026 年这堵墙为什么还在加高。

运维开发工程师

硬件

人工智能

架构设计

GPU内核性能优化常被误认为由矩阵运算速度决定，但FlashAttention-4等前沿研究揭示，真正瓶颈在SFU指数运算与共享内存带宽。本文详解Roofline模型、内存层次、合并访问等关键诊断工具，提供降精度、内核融合、分块复用等实操提效策略，并前瞻Rubin架构带宽增长滞后于算力的挑战，助你精准定位性能墙。

代码解释

编译器

DeepSeek

Python

2026上云采购 | AI焕新·智启新局

tione

4核4G3M云服务器 新用户低至38元/年！

文章

问答

视频

教程

学习中心

腾讯云实验室

直播

竞赛

腾讯云代码分析专区

腾讯iOA零信任安全管理系统专区

腾讯云架构师技术同盟交流圈

腾讯云数据库专区

腾讯云智能顾问专区

腾讯云原生专区

腾讯混元专区

腾讯云TCE专区

腾讯云Lighthouse专区

腾讯云HAI专区

腾讯云Edgeone专区

腾讯云存储专区

腾讯云智能专区

腾讯轻联专区 

腾讯云开发专区

TAPD专区

腾讯轻量云游戏服专区

EdgeOne AI 安全实战专区

腾讯云最具价值专家

腾讯云架构师技术同盟

腾讯云创作之星

腾讯云开发者先锋

腾讯云代码助手

云原生构建

TAPD 敏捷项目管理

Cloud Studio

SDK中心

API中心

命令行工具

涵盖代码开发、场景应用、自动测试全流程，助你从零构建专属AI助手

一站式MCP教程库，解锁AI应用新玩法

聚焦“写作效率、视觉美观与运行性能”三方面进行全面升级，为您提供更高效、稳定的创作环境

社区富文本&Markdown编辑器全新改版上线，欢迎大家体验!

诚挚邀请您参与本次调研，分享您的真实使用感受与建议。您的反馈至关重要，感谢您的支持与参与！

社区新版编辑器体验调研

AI Infra 系列 · 第六章 GPU 性能工程（二）：B200 算力涨 2.25 倍，运数据的路一寸没宽，你的内核其实在等数据-腾讯云开发者社区-腾讯云

AI Infra 系列 · 第六章 GPU 性能工程（二）：B200 算力涨 2.25 倍，运数据的路一寸没宽，你的内核其实在等数据

AI Infra 系列 · 第六章 GPU 性能工程（二）：B200 算力涨 2.25 倍，运数据的路一寸没宽，你的内核其实在等数据

社区

活动

圈层

关于

腾讯云开发者

热门产品

热门推荐

更多推荐