Triton 教程

主题

11. 性能建模与 Roofline 分析 ​

写代码前预知上限，写完后定位瓶颈——这是高级 GPU 工程师与新手的分水岭。本章把性能分析从"猜测 + benchmark"升级到"建模 + 验证"。你将学会在写一行代码之前算出"这个 kernel 最多能跑多快"，并在跑完之后用 Nsight Compute 的指标定位到 cycle 级别的瓶颈。

本章内容概览 ​

  • 11.1 性能建模的目标
  • 11.2 Roofline Model 详解
  • 11.3 Arithmetic Intensity 计算实战
  • 11.4 Occupancy 分析
  • 11.5 Nsight Compute 关键指标
  • 11.6 用 Nsight Compute 生成 Roofline Chart
  • 11.7 带宽利用率与 TFLOPS 利用率
  • 11.8 案例研究：用 Roofline 分析教程示例 kernel
  • 11.9 本章小结

11.1 性能建模的目标 ​

性能建模回答两个问题：

时机问题方法
写代码 前这个算子最多能跑多快？值不值得花时间手写？Roofline + Arithmetic Intensity 估算
写代码 后跑出来的速度离上限差多远？瓶颈在哪？Nsight Compute 实测 + SOL / occupancy 分析

为什么要"先建模"

很多新手会一头扎进 kernel 写作，结果发现：

  • 改了 100 行代码，性能涨 2%（因为算子本身就是 memory-bound，再怎么优化也跑不过 PyTorch）
  • 或者改了 5 行代码，性能涨 10×（因为之前根本不知道有 tensor core 的潜力可以挖）

写之前算 Roofline，能让你在 5 分钟内做出"该不该手写这个算子"的判断。

11.1.1 一个反例：为什么向量加法不值得手写 ​

算子: vector_add (fp16, N=1M elements)
FLOP: N (一次加法)
IO:   3N × 2 B = 6N bytes
AI:   1/6 ≈ 0.17 FLOP/byte

A100 ridge point (BF16 TC): 156 FLOP/byte
→ AI 远小于 ridge → memory-bound
→ 性能上限 = 2.0 TB/s × 0.17 = 340 GFLOPS
→ PyTorch 已经能跑 ~330 GFLOPS（97% 上限）
→ 手写没有任何空间，时间投入比 1:0

反之 fused LayerNorm + Residual：

算子: fused_layernorm_residual (fp16, B×D=8192×4096)
原始 PyTorch: 4 个独立 kernel × 各自往返 HBM ≈ 4× IO
融合后:       1 个 kernel × 1× IO
→ 即使 AI 不变，IO 减 4×，时间 ≈ 减 75%
→ 这就是 Triton 的甜点：fused kernel 的性能上限取决于 IO 量，IO 量取决于你怎么融合

11.2 Roofline Model 详解 ​

11.2.1 核心公式 ​

attainable_perf(FLOPS) = min( π,  β × I )
                            ↑       ↑
                  peak compute   peak BW × arithmetic intensity
  • π = 硬件峰值算力（FLOPS / 秒）
  • β = 峰值 HBM 带宽（bytes / 秒）
  • I = arithmetic intensity = 总 FLOP / DRAM bytes（FLOP / byte）

Ridge point（拐点）= π / β，单位 FLOP/byte。kernel 的 I 落在 ridge 左侧即 memory-bound，右侧即 compute-bound。

11.2.2 主流 GPU 的 π、β、Ridge ​

系统Peak compute（TFLOPS）HBM BW（TB/s）Ridge（FLOP/byte）
A100 80GB SXM, BF16 TC3122.0156
H100 SXM, BF16 TC9893.35295
H100 SXM, FP8 TC19793.35592
H200 SXM, BF16 TC9894.8206
B200, BF16 TC22508.0281
B200, FP8 TC45008.0562
B200, FP4 TC90008.01125

来源：Modal GPU Glossary、NVIDIA H100 数据手册。

算力 peak 的"口径陷阱"

  • A100 FP16/BF16 Tensor Core 峰值 312 TFLOPS 是 dense；带 2:4 sparsity 翻倍到 624。
  • H100 SXM5 BF16/FP16 dense 是 989 TFLOPS，sparsity 翻倍到 1979。
  • 不同来源经常混用 dense / sparsity 数字。做 roofline 分析前先统一稀疏假设，否则结论差 2×。
  • 表中均用 dense 口径。

11.2.3 几何意义：文字版 Roofline 图 ​

log-log 坐标，横轴 AI、纵轴可达 FLOP/s：

log(FLOP/s)
    ▲
    │                            ┌──────────────  ← π (compute ceiling)
    │                           ╱
    │                          ╱   <─── compute-bound 区
    │                         ╱
    │                        ╱
    │              <──┐     ╱
    │   memory-bound  │    ╱
    │           区    │   ╱     斜率 = β
    │                 │  ╱      (memory ceiling)
    │                 │ ╱
    │                 │╱
    │                 ●  ridge point = π / β
    │                 │
    └─────────────────┴────────────────────────────►  log(AI, FLOP/byte)
                     156 (A100 BF16)
                     295 (H100 BF16)
                     1125 (B200 FP4)

怎么读：

  • 找到你的 kernel 的 AI（横轴）
  • 沿垂直线往上走，撞到 roofline 的位置 = 性能上限
  • AI < ridge：撞到斜线（memory-bound），上限 = β × AI
  • AI > ridge：撞到平顶线（compute-bound），上限 = π

11.2.4 怎么拿到硬件 π / β ​

方法 1：官方数据手册 ​

最方便但最理想化。比如 NVIDIA H100 SXM5 数据手册。

方法 2：实际频率查询 ​

bash
nvidia-smi --query-gpu=memory.bandwidth_utilization,clocks.mem --format=csv
nvidia-smi --query-gpu=clocks.current.sm --format=csv,noheader

注意：GPU 在低负载时会降频。测性能前要先做 1 秒的 warmup 把频率拉到 boost。

方法 3：Nsight Compute 实测 peak ​

bash
ncu --set=@roofline python my_kernel.py

它会跑一组 micro-benchmark 自动得到"机器实测峰值"。这个数字通常比官方 spec 低 10~15%，因为是 sustained 而非 instantaneous peak。生产环境建议用这个数。

11.3 Arithmetic Intensity 计算实战 ​

11.3.1 通用公式 ​

AI = total_FLOP / total_DRAM_bytes

注意几个常见误算：

  1. DRAM bytes ≠ load + store：要包括所有真正访问 HBM 的字节，不包括 cache 命中的部分。
  2. FLOP 按指令计：FMA = 2 FLOP（一次乘 + 一次加）。
  3. 混合精度算子：累加器是 fp32 但输入是 fp16 时，FLOP 不变，bytes 按各自精度算。

11.3.2 Elementwise ​

FLOP = ops_per_element × N
bytes = (n_loads + n_stores) × element_size × N
AI = ops_per_element / ((n_loads + n_stores) × element_size)

举例：

算子opsloadsstoresdtypeAI
out = x + y (fp32)1214 B1/12 = 0.083
out = x + y (bf16)1212 B1/6 = 0.167
out = x*x + y*y (fp16)3 (mul+mul+add)212 B3/6 = 0.5
out = ReLU(x*y+z) (fp16)3 (FMA算2+ReLU算1)312 B3/8 = 0.375

结论：elementwise 永远是 memory-bound。AI 上限大约 0.5~1.0 FLOP/byte，远低于任何 GPU 的 ridge。

11.3.3 GEMM ​

公式：

GEMM C = A × B，shape A=(M,K), B=(K,N), C=(M,N)
FLOP  = 2 × M × N × K           (每个 C 元素 K 次 FMA × 2)
bytes = (M×K + K×N + M×N) × dtype_bytes    (读 A 读 B 写 C)
AI    = 2MNK / ((MK + KN + MN) × dtype)

当 M = N = K = D（方阵）：

AI = 2D³ / (3D² × dtype) = (2D) / (3 × dtype)

关键洞察：GEMM 的 AI 与边长 D 线性相关——大矩阵 compute-bound，小矩阵 memory-bound。

实例：fp16 GEMM，M=N=K=4096 ​

FLOP  = 2 × 4096³ = 1.374e11
bytes = 3 × 4096² × 2 = 9.83e7
AI    = 1.374e11 / 9.83e7 = 1398 FLOP/byte
  • A100 ridge = 156，4096 远大于 ridge → 深度 compute-bound
  • 性能上限 = π = 312 TFLOPS
  • 实测 cuBLAS / 优化好的 Triton matmul 在 D=4096 时跑 ~280 TFLOPS = 90% 上限

不同 D 下的 AI 速查表（fp16） ​

D (=M=N=K)AI (FLOP/byte)A100 上限
642142 GFLOPS (memory-bound)
1284386 GFLOPS (memory-bound)
25685170 GFLOPS (memory-bound)
512170267 TFLOPS (刚跨过 ridge，仍偏 memory)
1024341312 TFLOPS (compute-bound, 但 occupancy 可能不足)
2048683312 TFLOPS (compute-bound)
40961365312 TFLOPS (深度 compute-bound)

为什么 batched 小矩阵要拼成大矩阵

看上表：D=128 时性能上限只有 86 GFLOPS = 0.027% 的峰值算力。这就是为什么 grouped GEMM / batched GEMM 要把多个小矩阵拼成大矩阵——本质是把 AI 拉到 ridge 右边。

11.3.4 FlashAttention ​

朴素 attention 的 AI 分析（不算 softmax 临时存储）：

Q, K, V shape: (B, H, S, D)
S = sequence length, D = head dim

attention scores: Q × K^T → (B, H, S, S)
softmax(scores)
output:           scores × V → (B, H, S, D)

朴素 FLOP  ≈ 4 × B × H × S² × D
朴素 bytes ≈ 2 × (3 × B × H × S × D + B × H × S²)
            ↑    ↑                     ↑
          fp16   Q,K,V                 S×S attention 矩阵 (写出 + 读回)
朴素 AI ≈ D / 2  (当 S ≫ D 时)

FlashAttention 的关键改进：把 S×S attention 矩阵保留在 SRAM 中，不往返 HBM。这样：

FA FLOP  ≈ 4 × B × H × S² × D     (相同)
FA bytes ≈ 2 × 3 × B × H × S × D   (只读 Q, K, V，不写 attention 矩阵)
FA AI   ≈ (4 × S × D) / (6 × D) × S / S = 2S/3 → 实际接近 D
        (经过 tile 优化后，准确值约 D/2 ~ D)
算子D=64, S=2048 AID=128, S=2048 AI
朴素 attention32 (memory-bound)64 (memory-bound)
FlashAttention~64 (memory-bound, 但更高)~128 (memory-bound, 接近 ridge)

FA 仍然是 memory-bound

FlashAttention 把 AI 从 32 提到 128 已经很惊人了，但 A100 上 D=128 时 AI=128 仍然小于 ridge 156——所以 FA 在 A100 上仍是 memory-bound，性能上限 = β × AI = 2.0 × 128 ≈ 256 TFLOPS（而非 312 peak）。这也是为什么 H100 上 FA-3 通过 warp specialization 把"加载 K/V"和"算 P*V"真正并行起来，进一步压榨内存带宽。

11.3.5 常见 kernel 的 AI 速查表 ​

KernelFLOPDRAM BytesI (FLOP/byte)
Vector add (fp32, N elt)N12N0.083
Vector add (bf16)N6N0.167
LayerNorm (fp16, B×D)~8BD4BD2
Softmax (fp16, B×N)~5BN4BN1.25
RMSNorm (fp16)~4BD4BD1
GEMM (fp16, M=N=K=D)2D³6D²D/3
FlashAttn fwd (D=128, S seq, bf16)4·S²·D6SD~D/1.5
Conv2d (fp16, common)dependsdependstyp. 100~300

→ vector add / softmax / layernorm / RMSNorm 永远 memory-bound，优化目标是带宽利用率； → 大 GEMM 在 K 长时 compute-bound，瓶颈是 Tensor Core 利用率； → FlashAttention 在 D=64~128、长 seq 时仍是 memory-bound（在 attention matrix 保留进 SRAM 后才接近 ridge）。

11.4 Occupancy 分析 ​

11.4.1 定义 ​

achieved_occupancy = active_warps_per_SM / max_warps_per_SM

max_warps_per_SM 在 A100/H100 上都是 64（compute capability 8.0/9.0）。

11.4.2 三大约束（取下限） ​

每个 SM 能同时跑多少 warp，由三个独立约束的最小值决定：

约束 1：寄存器数 ​

每 SM 65536 个 32-bit register
每个 thread 最多 255 register

例：thread 用 128 reg，block 256 thread → 一个 block 用 128 × 256 = 32768 register → 一个 SM 最多挂 65536 / 32768 = 2 个 block = 2 × 256 / 32 = 16 warp，理论 occupancy = 16/64 = 25%。

如果 thread 用 32 reg：每 block 8192 reg → 8 block / SM → 64 warp / SM = 100% occupancy。

约束 2：Shared Memory ​

GPU单 SM SMEM 物理大小单 block 最大可配
A100192 KB164 KB
H100256 KB227 KB（CUDA 保留 1 KB）

例：block 用 100 KB SMEM → A100 上一个 SM 最多 1 block，H100 最多 2 block。

约束 3：Block / Warp 总数 ​

GPU每 SM 最大 block每 SM 最大 warp
A100 / H1003264

block 太小（如 32 thread = 1 warp）会被 block 数撞顶：32 block × 1 warp = 32 warp = 50% occupancy（哪怕寄存器和 SMEM 都不卡）。

11.4.3 Triton 视角 ​

Triton 参数影响哪个约束
BLOCK_SIZE_M × BLOCK_SIZE_N决定累加器大小 → 寄存器
num_warps× 32 = block size → 影响 block 数和总 thread 数
num_stages每多一级 SMEM buffer → 线性增加 SMEM
BLOCK_SIZE_K影响 K-buffer 大小 → SMEM

11.4.4 一个完整算例 ​

配置：A100、fp16、BLOCK_M=BLOCK_N=128、BLOCK_K=32、num_stages=3、num_warps=4

SMEM 占用：

每级 K-buffer = A tile + B tile
            = 128×32×2 B + 32×128×2 B
            = 8 KB + 8 KB = 16 KB
num_stages = 3 → 总 SMEM = 3 × 16 KB = 48 KB
A100 单 block 上限 164 KB → 不卡 → 1 SM 上 SMEM 维度可以挂 164/48 ≈ 3 block

寄存器占用（假设 ptxas 报 192 reg/thread）：

每 block thread 数 = num_warps × 32 = 4 × 32 = 128 thread
每 block reg 数  = 192 × 128 = 24576
每 SM 最多 block = 65536 / 24576 = 2 block （向下取整）

综合：

SMEM 维度：3 block/SM
Reg 维度： 2 block/SM
Block 数：32 block/SM
→ 取最小 = 2 block/SM

总 warp = 2 block × 4 warp = 8 warp
Occupancy = 8 / 64 = 12.5%

怎么拉高？

  • 降到 num_warps=8：block 内 thread 翻倍到 256，每 block 64 KB SMEM、49152 reg → 仍 1 block/SM，总 warp 8 → 还是 12.5%
  • 降 BLOCK_K=16：SMEM 减半到 24 KB，但 reg 不变 → 仍 2 block，仍 12.5%
  • 加 maxnreg=128 限制 ptxas：thread 用 128 reg → block 用 16384 reg → 4 block/SM × 4 warp = 16 warp = 25%
  • 降 BLOCK_M/N 到 64：累加器矩阵小 4 倍 → reg 估算 ~96 → 4 block/SM × 4 warp = 25%

11.4.5 从 Triton 读取实际寄存器数 ​

python
# 方式 1：metadata
compiled = my_kernel[(grid,)](*args)
print(compiled.metadata)
# {'num_warps': 4, 'num_stages': 3, 'shared': 49152, ...}
print(compiled.n_regs)       # 每 thread 寄存器数
print(compiled.n_spills)     # spill 字节数（>0 警报）

# 方式 2：ptxas 日志
# TRITON_DUMP_PTXAS_LOG=1 python ... 看输出

11.4.6 高 occupancy ≠ 高性能 ​

Volker Volkov 那篇经典论文（Better Performance at Lower Occupancy）证明：25% 占用 + 充分 ILP（每个 thread 持有更多寄存器以 batch 更多独立操作）能跑赢 100% 占用。

Triton 上 GEMM 在 25~50% occupancy 区间最常达到峰值——因为 wgmma/tensor core 的延迟主要靠 K 轴软件流水线隐藏，而非线程级并发。

不要迷信 occupancy

看到 achieved_occupancy = 25% 不要立刻去调。先看：

  1. SM SOL / Memory SOL 是否已经 > 80%（说明算力或带宽已饱和）
  2. Warp stall reason（如果是 long_scoreboard 多 → 真的是 latency 没被掩盖，加 occupancy 才有用）

如果 SOL 已经很高，低 occupancy 不是问题。

11.5 Nsight Compute 关键指标 ​

Nsight Compute (NCU) 是 NVIDIA 官方 GPU profiler。它的指标命名约定：

<unit>__<counter>.<rollup>.<sub_metric>
  • unit：sm、smsp（SM sub-partition）、dram、l1tex、lts（L2）
  • rollup：sum、avg、max
  • sub_metric：pct_of_peak_sustained_active、per_second 等

11.5.1 Memory 类（找带宽瓶颈） ​

指标含义
dram__bytes_read.sumHBM 读字节数 = sectors_read × 32 B
dram__bytes_write.sumHBM 写字节数
dram__throughput.avg.pct_of_peak_sustained_elapsedDRAM 带宽利用率（"SOL FB"）
dram__cycles_elapsed.avg.per_secondDRAM 实际频率
l1tex__throughput.avg.pct_of_peak_sustained_activeL1/Tex cache 单元利用率
l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sumL1 全局 load sector 数
lts__t_sectors_hit_rate.pctL2 命中率（关键：低命中率 → 多次跑 DRAM）
lts__throughput.avg.pct_of_peak_sustained_elapsedL2 吞吐
smsp__sass_average_data_bytes_per_sector_mem_global.pct每 sector 平均有效字节占比（< 100% 说明未合并）

11.5.2 Compute 类（找算力瓶颈） ​

指标含义
sm__sass_thread_inst_executed_op_fadd_pred_on.sumFP32 add 指令数
sm__sass_thread_inst_executed_op_ffma_pred_on.sumFP32 FMA（×2 算 FLOP）
sm__sass_thread_inst_executed_op_hfma_pred_on.sumFP16 标量 FMA
sm__pipe_tensor_cycles_active.avg.pct_of_peak_sustained_activeTensor Core 利用率
sm__inst_executed_pipe_tensor.sumTensor Core 发射的 mma 指令数
sm__throughput.avg.pct_of_peak_sustained_elapsedSM 综合吞吐（"SM SOL"）

Tensor Core FLOP 计算

Nsight Compute 自带的 roofline 只统计 fp32/fp64 scalar；要画 BF16/FP8/Tensor Core roofline 必须自定义 section file。Tensor Core 的 FLOP 要按指令 shape 加权：

  • mma.m16n8k16.f16：一次完成 16×8×16×2 = 4096 FLOP
  • wgmma.m64n128k16：一次 64×128×16×2 = 262144 FLOP

参考 NVIDIA 论坛 Nsight Compute-Roofline chart。

11.5.3 Occupancy 类 ​

指标含义
sm__warps_active.avg.pct_of_peak_sustained_activeAchieved Occupancy
launch__waves_per_multiprocessor一个 grid 在 SM 上跑几"波"（< 1 即 tail effect 严重）
launch__block_size每 block thread 数
launch__registers_per_thread每 thread 寄存器数
launch__shared_mem_per_block_static + _dynamic每 block SMEM 字节数
Theoretical Occupancy由 block_size、registers、SMEM 决定的上限（NCU 自动算）

11.5.4 Warp Stall Reasons（找瓶颈类型） ​

smsp__warps_issue_stalled_* 系列告诉你 warp 为什么不能发指令：

指标何时是瓶颈修复方向
short_scoreboard等本地 MIO（SMEM、ldmatrix）检查 SMEM bank conflict / 减小 SMEM 占用
long_scoreboard等 global load / L1 miss加 num_stages、提升 occupancy
barrier等 __syncthreads减少同步、缩小 block
membar等 memory fence（fp atomic）减少原子操作
no_instructionI-cache miss / 大 kernel拆分 kernel、减少展开
wait等 ALU pipelineILP 不够，让 thread 持有更多寄存器

11.5.5 Speed of Light (SOL) ​

SOL = 各类资源相对其物理峰值的利用率。NCU 报告里两个核心数字：

  • SM SOL % = sm__throughput.avg.pct_of_peak_sustained_elapsed
  • Memory SOL % = gpu__compute_memory_throughput.avg.pct_of_peak_sustained_elapsed

两者由子单元中最高利用率那个决定。判定规则：

现象诊断优化方向
SM SOL > 80%Compute-bound进一步优化要降 FLOP 或换更快 ALU（如 fp16 → fp8）
Mem SOL > 80%Memory-bound进一步优化要降 byte 或 fuse kernel
都低 (< 60%)Latency-bound看 stall reason；常见 occupancy 太低或 warp scheduler 选不到 ready warp
SM ≈ Mem ≈ 50%均衡难再榨油，看是否有微观调度优化

SOL 是性能分析的"第一眼"

看一个 kernel 的 NCU 报告时，永远先看 SOL 那两个数。它直接告诉你瓶颈大类，免得在错的方向上花时间。

11.6 用 Nsight Compute 生成 Roofline Chart ​

11.6.1 安装与权限 ​

bash
# 安装（CUDA Toolkit 自带）
which ncu

# 解锁 NVIDIA 性能计数器权限（一次性，需 root）
sudo modprobe -r nvidia_uvm nvidia
sudo modprobe nvidia NVreg_RestrictProfilingToAdminUsers=0
sudo modprobe nvidia_uvm

# 或永久写入
echo 'options nvidia NVreg_RestrictProfilingToAdminUsers=0' \
    | sudo tee /etc/modprobe.d/nvidia-profiling.conf

不解锁权限会报 ERR_NVGPUCTRPERM。

11.6.2 最小命令 ​

bash
ncu --set=full \
    --section=SpeedOfLight_RooflineChart \
    --section=SpeedOfLight \
    --section=ComputeWorkloadAnalysis \
    --section=MemoryWorkloadAnalysis \
    -o my_profile \
    python my_triton_script.py

输出 my_profile.ncu-rep，用 ncu-ui my_profile.ncu-rep 打开看 GUI。

11.6.3 只跑指定 kernel ​

autotune 期间一个 Python 脚本可能调几百次 kernel，全部 profile 太慢。过滤方法：

bash
ncu --set=full \
    --kernel-name regex:matmul_kernel \
    --launch-count 1 \
    --launch-skip 5 \
    -o matmul_profile \
    python my_triton_script.py
  • --kernel-name regex:...：用正则匹配 kernel 名
  • --launch-skip 5：跳过前 5 次调用（让 autotune 完成 + warmup）
  • --launch-count 1：只 profile 1 次

11.6.4 Triton kernel 实战配方 ​

Red Hat / Spheron 推荐的标准流程：

bash
# Step A: 让 Triton 把每个 autotune config 都跑一次并 dump 最优
TRITON_PRINT_AUTOTUNING=1 \
TRITON_KERNEL_DUMP=1 \
python infer.py

# Step B: 选定最优 config，固定下来
# @triton.autotune(configs=[best_only])  # 关掉其他 config

# Step C: NCU 只跑这一个
ncu --replay-mode kernel \
    --set full \
    --kernel-name flash_attn_fwd \
    -o fa_profile \
    python infer.py

# Step D: GUI 看
#   - Speed of Light → Roofline 标签：单点图
#   - Memory Workload Analysis：各级 cache 流量
#   - Source View：哪一行 SASS 是热点

11.6.5 Roofline chart 怎么看 ​

打开 .ncu-rep 后切到 Speed of Light → Roofline 标签：

  • 横轴 = Arithmetic Intensity（FLOP/byte），分 fp32 / fp64 双轴
  • 纵轴 = 实测 GFLOPS
  • 斜线 = β × I（HBM、L2 各一条）
  • 平顶线 = π
  • 你的 kernel = 单点；hover 看实际 BW 利用率与 SOL%

判断 kernel 在哪个 bound：

  • 单点贴近 HBM 斜线 → memory-bound（HBM 带宽吃满了）
  • 单点贴近 L2 斜线（高于 HBM 斜线） → L2-bound（数据从 HBM 来得慢，但 L2 已经撑爆）
  • 单点贴近平顶线 → compute-bound
  • 单点远低于所有 roofline → latency-bound / occupancy 问题

11.6.6 自定义 Roofline（BF16 / FP8 Tensor Core） ​

NCU 默认 roofline 只算 fp32/fp64。要画 Tensor Core 的：

  1. 复制 NCU SDK 里 SpeedOfLight_RooflineChart.section
  2. 修改 Y 轴公式：把 sm__pipe_tensor_cycles_active.sum × per-instruction FLOP 加进去
  3. 修改 X 轴分母：把 dram__bytes.sum 加进去
  4. 用 --section-folder 指向你的自定义 section

参考 NCU 文档 Custom Sections。

11.7 带宽利用率与 TFLOPS 利用率 ​

11.7.1 公式 ​

BW utilization    = achieved_bytes / (peak_BW × kernel_time)
FLOPS utilization = achieved_FLOPS / peak_FLOPS
                  = achieved_FLOPS / π

11.7.2 从测量到 FLOPS ​

python
# 用 triton.testing.do_bench 测时
ms = triton.testing.do_bench(lambda: kernel[(grid,)](*args))

# 算理论 FLOP
total_flop = 2 * M * N * K   # GEMM 例

# 实际 TFLOPS
achieved_TFLOPS = total_flop / (ms * 1e-3) / 1e12

# A100 BF16 peak = 312 TFLOPS
util = achieved_TFLOPS / 312
print(f"FLOPS utilization = {util * 100:.1f}%")

11.7.3 从测量到 BW ​

python
# 算理论搬运字节数
total_bytes = (M * K + K * N + M * N) * 2   # fp16

# 实际 GB/s
achieved_GBps = total_bytes / (ms * 1e-3) / 1e9

# A100 80GB peak BW = 2039 GB/s
util = achieved_GBps / 2039
print(f"BW utilization = {util * 100:.1f}%")

11.7.4 各种 kernel 的"好"利用率标准 ​

算子类型应该看什么利用率"好"的标准
Vector add / elementwiseBW utilization> 90% peak BW = 优秀
LayerNorm / SoftmaxBW utilization> 85% peak BW = 优秀
Reduction / sum / meanBW utilization> 80% peak BW = 优秀
GEMM (D ≥ 1024)FLOPS utilization> 85% peak FLOPS = 优秀（cuBLAS 水平）
GEMM (D ≤ 256)BW utilization看 AI 在 ridge 哪侧
FlashAttention (S ≥ 1024, D=128)BW utilization> 70% peak BW = 优秀
Conv2dFLOPS utilization> 70% peak FLOPS = 优秀

注意分母对齐

计算 achieved_bytes 时要算"最低限度需要搬运的字节"（即理论 IO 量），而不是 dram__bytes_read.sum + dram__bytes_write.sum（后者包含 cache 回写、未合并浪费等"额外"流量）。

如果你算"理论 BW"用前者，但跟 NCU 上的 dram__throughput 比，会发现 NCU 显示 95% 但你自己算只有 60%——这是分母不一致。两个数都对，但意义不同：

  • 算法层 BW 利用率（用理论 IO 量）：告诉你"算法本身离硬件极限多远"
  • 硬件层 BW 利用率（NCU 报告）：告诉你"硬件忙不忙"

调优时关注前者；判断是否打满硬件用后者。

11.8 案例研究：用 Roofline 分析教程示例 kernel ​

把前 9 章的 4 个 example kernel 放进 roofline 框架走一遍。GPU 假设 = A100 80GB SXM。

11.8.1 Case 1：Vector Add (Chapter 4) ​

配置: fp32, N = 64M elements
FLOP  = N = 6.7e7
bytes = 3N × 4 = 8.05e8
AI    = 0.083 FLOP/byte
项目值
Roofline 上限 (β × AI)2039 × 0.083 = 169 GFLOPS
实测~165 GFLOPS（97% 上限）
实测带宽~1980 GB/s（97% peak BW）
结论memory-bound 区，已饱和，无优化空间

11.8.2 Case 2：Softmax (Chapter 7 fused) ​

配置: fp16, B=8192, N=4096
FLOP  ≈ 5 × B × N = 1.67e8
bytes = 4 × B × N = 1.34e8
AI    = 1.25 FLOP/byte
项目值
Roofline 上限 (β × AI)2039 × 1.25 = 2549 GFLOPS
朴素 PyTorch (多 kernel)~150 GFLOPS（5.9% 上限）
Triton fused (online softmax)~2400 GFLOPS（94% 上限）
实测带宽 (fused)~1920 GB/s
结论memory-bound，但融合能省 4× IO，性能涨 16×

诊断逻辑：

  1. AI = 1.25 << ridge (156) → memory-bound
  2. 优化方向：减少 IO（融合）、不要做"提升算力"
  3. 融合后 BW utilization 已 94% → 没空间

11.8.3 Case 3：Matmul D=4096 (Chapter 5) ​

配置: fp16, M=N=K=4096
FLOP  = 2 × 4096³ = 1.37e11
bytes = 3 × 4096² × 2 = 9.83e7
AI    = 1398 FLOP/byte
项目值
Roofline 上限 (compute)π = 312 TFLOPS
Triton (无 grouped ordering)220 TFLOPS (71% 上限)
Triton (grouped ordering)245 TFLOPS (79% 上限)
Triton (autotuned)280 TFLOPS (90% 上限)
cuBLAS290 TFLOPS (93% 上限)
结论compute-bound，瓶颈是 Tensor Core 利用率 + L2 cache 复用

诊断逻辑：

  1. AI = 1398 >> ridge (156) → compute-bound
  2. NCU 看：SM SOL > 80%，Memory SOL < 50% → compute 主导
  3. 优化方向：grouped ordering（L2 复用）、num_stages 调优、autotune
  4. 不要做：再加 SRAM、再融合 kernel——对 compute-bound 无效

11.8.4 Case 4：FlashAttention fwd (Chapter 8) ​

配置: fp16, B=16, H=16, S=2048, D=128
FLOP  ≈ 4 × B × H × S² × D = 5.49e11
bytes ≈ 2 × 3 × B × H × S × D = 6.29e6 (假设 K/V/Q 各读一遍)
AI    ≈ 87.3 FLOP/byte  (经过 tile, 实际更高一点)
项目值
Roofline 上限 (β × AI)2039 × 87 = 177 TFLOPS
Roofline 上限 (compute)312 TFLOPS
实际 min(BW, compute)177 TFLOPS (memory-bound)
Triton FA-2 实测~150 TFLOPS (85% 上限)
结论memory-bound，FA 的精妙在于把 AI 从 ~D/4 提到 ~D（朴素 attention 的 4×）

诊断逻辑：

  1. AI = 87 < ridge (156) → 仍 memory-bound（H100 上 AI 接近 ridge）
  2. NCU 看：Memory SOL > 80%，SM SOL ≈ 50%
  3. 优化方向：进一步减少 HBM 流量（H100 上用 warp specialization 让 K/V 加载与 P*V 计算真并行）
  4. A100 上空间已经不大；H100 上有 FA-3 的 20~30% 提升空间

11.8.5 四个案例对比总结 ​

KernelAIBound优化策略当前实现 vs 上限
Vector add0.083memory没招97%
Softmax (fused)1.25memory融合94%
Matmul D=40961398computegrouped/autotune90%
FA fwd D=12887memorytiling85%

洞察：

  • 三个 memory-bound kernel 都跑到 85~97% 上限——这就是 Triton 的杀手锏：把带宽榨到极致。
  • 唯一 compute-bound 的 matmul 也跑到 90%（接近 cuBLAS）。
  • 没有任何 kernel 跑到 100% 上限，永远留 5~15% 是硬件实际可持续 vs 理论 peak 的差距。

11.9 性能建模工作流（一图总结） ​

                  ┌─────────────────────────────┐
                  │ 1. 写 kernel 之前           │
                  │    算 AI                    │
                  │    画 roofline 找上限       │
                  └──────────────┬──────────────┘
                                 ▼
                  ┌─────────────────────────────┐
                  │ AI < ridge?                 │
                  └──────────────┬──────────────┘
              memory-bound       │       compute-bound
                  ▼                              ▼
        ┌─────────────────┐          ┌─────────────────┐
        │ 目标：BW 90%+    │          │ 目标：FLOPS 85%+ │
        │ 优化：融合、     │          │ 优化：Tensor Core│
        │       减 IO、    │          │       autotune、 │
        │       调 stages  │          │       grouped    │
        └────────┬─────────┘          └────────┬─────────┘
                 ▼                              ▼
        ┌──────────────────────────────────────────────┐
        │ 2. 写完 kernel                               │
        │    do_bench 测 ms                            │
        │    算 BW util / FLOPS util                   │
        └──────────────────┬───────────────────────────┘
                           ▼
        ┌──────────────────────────────────────────────┐
        │ utilization < target?                        │
        └──────────────────┬───────────────────────────┘
                           ▼ Yes
        ┌──────────────────────────────────────────────┐
        │ 3. Nsight Compute 深挖                       │
        │    SOL 数字 → memory/compute/latency bound   │
        │    Stall reason → 找根因                     │
        │    Roofline chart → 验证理论分析             │
        └──────────────────────────────────────────────┘

本章小结 ​

  • Roofline 模型 把性能上限简化成 min(π, β × AI)，让你写代码前 5 分钟就能预知上限。
  • Arithmetic Intensity 是建模核心，elementwise (AI < 1) 永远 memory-bound，大 GEMM (AI > ridge) 才 compute-bound。
  • A100 ridge = 156、H100 ridge = 295、B200 ridge = 281（BF16）——越新的卡 ridge 越右，越难打满算力。
  • Occupancy 三大约束：registers / SMEM / block 数。Triton 上 GEMM 在 25~50% occupancy 区间最常达到峰值——高 occupancy ≠ 高性能。
  • Nsight Compute 优先级：先看 SM SOL 和 Memory SOL → 找 bound 类别 → 再看 warp stall reason → 找根因。
  • 典型 Triton kernel 利用率标准：memory-bound 90%+ BW util、compute-bound 85%+ FLOPS util 即可叫优秀。

至此"进阶篇"前两章（编译器原理、性能建模）讲完了。下一章我们把视角推到最新硬件——Hopper 的 TMA / WGMMA / 集群、Blackwell 的 tcgen05 / TMEM——看看 Triton 如何把这些"几年前还要写一千行 CUDA"的新特性变成 30 行 Python。

思考题 ​

  1. Roofline 决策题：你接到任务"用 Triton 加速这个 PyTorch 训练 step"，里面包含 4 个算子：

    • (a) gelu(x)：fp16，shape (8192, 4096)
    • (b) linear(x, W)：fp16，x 是 (8192, 4096)，W 是 (4096, 4096)
    • (c) linear(y, V)：fp16，y 是 (8192, 16384)，V 是 (16384, 4096)
    • (d) add(out1, out2)：fp16，两个 (8192, 4096) tensor 相加

    请对每个算子：(1) 算 AI；(2) 判断 A100 上是 memory/compute bound；(3) 估算性能上限；(4) 给出"是否值得手写 Triton 加速"的判断（与朴素 PyTorch 对比）。

  2. Occupancy 调参：你的 fp16 matmul kernel 在 A100 上配置 BLOCK_M=BLOCK_N=128, BLOCK_K=64, num_warps=8, num_stages=3，ptxas 报 Used 168 registers, 0 bytes spill stores。请算出：

    • (a) 每 block 寄存器数与 SMEM 字节数
    • (b) 三大约束分别允许多少 block/SM
    • (c) achieved occupancy
    • (d) 如果想拉到 50% occupancy，给出至少 2 种调整方案以及各自的潜在代价

  3. NCU 实战诊断：你写的 fused RMSNorm kernel 处理 fp16 (4096, 4096) 输入耗时 0.18 ms，理论 IO 量 67 MB，A100 80GB peak BW = 2039 GB/s。NCU 报告显示：

    • SM SOL = 22%
    • Memory SOL = 55%
    • achieved_occupancy = 28%
    • smsp__warps_issue_stalled_long_scoreboard 占总 stall 的 62%
    • lts__t_sectors_hit_rate.pct = 78%

    请：(1) 算实测 BW 利用率（用理论 IO）；(2) 判断瓶颈类型；(3) 解读 NCU 数据组合反映了什么问题；(4) 给出至少 3 个具体的优化方向（不要泛泛说"提升 occupancy"，要指明改哪个参数为什么有效）。

基于 MIT 协议发布

Copyright © 2026 Triton 中文教程