搞懂 GPU 推理的两堵墙带宽瓶颈 vs 计算瓶颈以及怎么判断你撞了哪堵面向对象刚上手大模型推理优化的同学大学生竞赛 / 课程项目都适用你将学到为什么「GPU 看着一直很忙」却不代表没优化空间怎么用 Roofline 和 profiler 判断瓶颈在带宽还是算力这套分析怎么落到 LLM 推理prefill / decode和实际赛题上。本文只讲原理与思路不贴具体工程代码重在让你建立正确的「判断框架」。目录搞懂 GPU 推理的两堵墙带宽瓶颈 vs 计算瓶颈以及怎么判断你撞了哪堵0. 从一个很常见的困惑说起1. GPU 的两种「忙」算力 vs 访存带宽2. Roofline 模型一张图判断瓶颈3. 为什么「GPU 总有 kernel」看不出瓶颈4. 怎么判断 / 测试一套可操作的方法论5. 把框架套到 LLM 推理prefill 和 decode 是两堵不同的墙5.1 Prefill处理输入 prompt—— 偏计算墙5.2 Decode逐 token 生成—— 典型带宽墙5.3 吞吐怎么由这两段决定6. 对应到赛题优化点到底在哪7. 四个层次的优化各自在打哪堵墙原理版8. 在 HIP / 国产加速卡DCU 等上怎么落地8.1 术语 / 工具对照CUDA → HIP / ROCm8.2 在 HIP 上做 Roofline / 瓶颈判断8.3 HIP 编程里影响优化的关键差异8.4 国产 DCU海光 gfx9 系部署注意9. 推荐参考GitHub / 资料10. 小结0. 从一个很常见的困惑说起有同学问我「我跑 torch profile 看推理时 GPU 上几乎总有 kernel 在运行这不就说明 GPU 一直在干活、问题和调度关系不大吗是不是没什么可优化的了」这个困惑非常典型而且结论恰恰相反。GPU「一直在跑 kernel」≠「GPU 在高效干活」。很多时候那些 kernel 正趴在那里等内存——计算单元空转等着数据从显存搬过来。要看懂这件事你得先理解 GPU 上的两堵墙计算墙和带宽墙。1. GPU 的两种「忙」算力 vs 访存带宽一块 GPU 干活本质上消耗两种资源算力compute每秒能做多少次浮点运算单位 FLOP/s比如几十 TFLOP/s访存带宽memory bandwidth每秒能从显存HBM/GDDR搬多少字节单位 GB/s比如 1~3 TB/s。任何一个 kernelGPU 上的一段并行计算跑起来要么卡在算力上要么卡在带宽上计算瓶颈compute-bound数据早就搬到了但运算太多算不过来 → 卡算力带宽瓶颈memory-bound运算很少但要搬的数据太多搬不过来 → 卡带宽计算单元在等数据。判断一个 kernel 属于哪种看一个关键指标——算术强度Arithmetic IntensityI 这个 kernel 做的浮点运算次数 (FLOPs) 它要读写的字节数 (Bytes) I \frac{\text{这个 kernel 做的浮点运算次数 (FLOPs)}}{\text{它要读写的字节数 (Bytes)}}I它要读写的字节数(Bytes)这个kernel做的浮点运算次数(FLOPs)​直观理解搬一个字节进来能顺带做多少次运算。I II大 → 搬一次算很多次 → 算力先撑不住 → compute-boundI II小 → 搬一堆数据只算几下 → 带宽先撑不住 → memory-bound。2. Roofline 模型一张图判断瓶颈把上面两堵墙画成一张图就是经典的Roofline 模型。一个 kernel 实际能达到的性能上限是P achievable min ⁡ ( P peak , B × I ) P_{\text{achievable}} \min\big(\,P_{\text{peak}},\; B \times I\,\big)Pachievable​min(Ppeak​,B×I)其中P peak P_{\text{peak}}Ppeak​是 GPU 峰值算力FLOP/sB BB是峰值带宽Bytes/sI II是算术强度。当B × I P peak B \times I P_{\text{peak}}B×IPpeak​即I II小性能被带宽限制 → memory-bound处在 roofline 的「斜坡」上当B × I ≥ P peak B \times I \ge P_{\text{peak}}B×I≥Ppeak​即I II大性能被算力限制 → compute-bound处在 roofline 的「平顶」上。两者的分界点叫脊点ridge pointI ridge P peak B I_{\text{ridge}} \frac{P_{\text{peak}}}{B}Iridge​BPpeak​​ 一句话记住I I ridge I I_{\text{ridge}}IIridge​就是带宽墙I I ridge I I_{\text{ridge}}IIridge​就是计算墙。性能 P ▲ │ ________________ ← 平顶compute-bound卡算力 P_peak │ / │ / ← 斜坡memory-bound卡带宽斜率 B │ / └────────●─────────────────▶ 算术强度 I I_ridge为什么这张图重要因为它告诉你优化的天花板在哪如果你在斜坡上memory-bound再怎么优化算法/减少计算都没用——你得减少数据搬运或提高带宽利用率如果你在平顶上compute-bound优化访存没用——你得减少计算量或用更高吞吐的算力单元如 Tensor Core。搞错了墙方向就全错了。3. 为什么「GPU 总有 kernel」看不出瓶颈回到开头的困惑。profiler 的 timeline 上「一直有 kernel」反映的是GPU 占用率occupancy/utilization高但占用率高 ≠ 高效一个 memory-bound 的 kernel计算单元大部分时间在stall停顿等 HBM 返回数据。从 timeline 看这个 kernel「在运行」占着 GPU但实际有效算力利用率可能只有个位数百分比。也就是说timeline 的「满」只能说明「没空闲」不能说明「跑得快」。⚠️关键认知判断「还有没有优化空间」不要看「GPU 忙不忙 / 有没有 kernel」要看带宽利用率或算力利用率——实测吞吐距离硬件峰值还有多远。所以那位同学的 timeline 分析没错但解读错了层次timeline 能告诉你「谁在跑、跑多久」但要定位瓶颈得进一步看每个 kernel 的memory throughput / compute throughputprofiler 里有这些指标算出它离 roofline 哪条边更近。4. 怎么判断 / 测试一套可操作的方法论下面是定位瓶颈的标准流程用伪代码描述工具用 NVIDIA Nsight Compute、AMD rocprof或框架自带 profiler 都行流程 1把 kernel 按类别归并找大头 ───────────────────────────────────── 导出 profiler trace含每个 GPU kernel 的耗时 按算子类别聚合GEMM / attention / norm / elementwise / ... 对每类总耗时、调用次数、占总时间百分比 → 找出占大头的那 1~2 类优化它们性价比最高流程 2对大头 kernel 做 Roofline 判断 ───────────────────────────────────── 对目标 kernel 估算 FLOPs这步做了多少次乘加 估算 Bytes读写了多少字节输入/权重/输出 I FLOPs / Bytes # 算术强度 I_ridge P_peak / B # 硬件脊点 if I I_ridge: 瓶颈 memory-bound带宽墙 理论下界时间 Bytes / B # 搬完数据至少要这么久 else: 瓶颈 compute-bound计算墙 理论下界时间 FLOPs / P_peak流程 3算「利用率」判断还有多少肉 ───────────────────────────────────── 实测时间 profiler 测得的 kernel 耗时 利用率 理论下界时间 / 实测时间 # 越接近 1 越到墙 if 利用率 0.7: kernel 没喂满硬件有优化空间 # 改访存模式 / 调度 / 融合 elif 利用率 0.9: 已贴近硬件墙这条路到顶了 # 换思路别死磕技巧很多 profiler如 Nsight Compute直接给你Memory Throughput %和Compute (SM) Throughput %两个百分比——哪个高就卡哪个。这比自己估 FLOPs/Bytes 更快但自己会估能帮你建立物理直觉。5. 把框架套到 LLM 推理prefill 和 decode 是两堵不同的墙大模型自回归推理分两个阶段它们撞的墙完全不同这是理解 LLM 推理优化的核心5.1 Prefill处理输入 prompt—— 偏计算墙一次性处理整段输入比如几千个 token内部是大矩阵乘GEMM和注意力。token 多 → 矩阵大 → 算术强度高 → 通常compute-bound。它决定TTFTTime To First Token首字延迟注意力是O ( n 2 ) O(n^2)O(n2)长输入时往往是 prefill 的最大头。5.2 Decode逐 token 生成—— 典型带宽墙生成阶段一次只算1 个 token。但即使只算 1 个 token也要把整个模型的权重从显存读一遍。于是I decode 少量运算 读整个模型权重的字节 ≈ 非常小 I_{\text{decode}} \frac{\text{少量运算}}{\text{读整个模型权重的字节}} \;\approx\; \text{非常小}Idecode​读整个模型权重的字节少量运算​≈非常小算术强度极低 →铁板钉钉的 memory-bound。decode 的速度几乎完全由「把权重从显存搬一遍要多久」决定T decode per token ≳ 每 token 需读取的权重字节数 B HBM T_{\text{decode per token}} \;\gtrsim\; \frac{\text{每 token 需读取的权重字节数}}{B_{\text{HBM}}}Tdecode per token​≳BHBM​每token需读取的权重字节数​这就是TPOTTime Per Output Token每个输出 token 的耗时的理论下界。划重点decode 是 memory-bound瓶颈是权重读取带宽不是算力。在 decode 上拼命优化「计算」是白费力气——计算单元本来就在等内存。5.3 吞吐怎么由这两段决定单条请求生成N out N_{\text{out}}Nout​个 token 的总时间约为T total ≈ TTFT ( N out − 1 ) × TPOT T_{\text{total}} \approx \text{TTFT} (N_{\text{out}} - 1)\times \text{TPOT}Ttotal​≈TTFT(Nout​−1)×TPOT输出吞吐Throughput ≈ N out TTFT N out × TPOT \text{Throughput} \approx \frac{N_{\text{out}}}{\text{TTFT} N_{\text{out}}\times \text{TPOT}}Throughput≈TTFTNout​×TPOTNout​​当输出很长N out N_{\text{out}}Nout​大时N out × TPOT N_{\text{out}}\times\text{TPOT}Nout​×TPOT主导分母于是Throughput → N out 大 1 TPOT \text{Throughput} \xrightarrow{N_{\text{out}}\text{ 大}} \frac{1}{\text{TPOT}}ThroughputNout​大​TPOT1​输出越长吞吐越由 TPOT也就是 decode 带宽说了算。6. 对应到赛题优化点到底在哪很多大模型推理赛题的评测设定是单条请求并发为 1、固定且较长的输出长度在时延 SLA 约束下比吞吐。把第 5 节的结论套进去吞吐 ≈ 1 / TPOT输出长 → decode 主导→核心是降 TPOT其次是降长输入档的 TTFT。TPOT 是 memory-bound→ 要盯的是decode 阶段权重读取的带宽利用率你用的计算后端 / 注意力实现有没有把 HBM 喂满库的默认 kernel 在「逐 token」这种又瘦又长的矩阵-向量乘形状上往往不是最优的。调度参数基本无效像「关闭 chunked prefill」「关闭 prefix caching」这类调度优化是为多请求 batch / 重复前缀设计的。单请求、无重复前缀的场景下它们基本是杯水车薪——方向不在调度。判断到没到顶的标尺用第 4 节的方法算出 decode 的「带宽利用率」。利用率离峰值还远 → 有肉已贴近 HBM 峰值 → 这条路到顶得换思路或承认硬件天花板。用伪代码把「赛题瓶颈定位」串起来赛题瓶颈定位 ───────────────────────────────────── 1. 确认评测口径并发数、输出长度 2. if 输出长 and 并发1: 主瓶颈 TPOTdecode, memory-bound 次瓶颈 长输入档的 TTFTprefill, compute-bound 3. 测 decode每 token 权重读字节 / HBM 峰值带宽 TPOT 地板 4. 带宽利用率 TPOT地板 / 实测TPOT 5. if 利用率低: 优化 decode 的权重读 kernel喂满带宽 else: 撞带宽墙承认或换非带宽路线7. 四个层次的优化各自在打哪堵墙原理版业界常把推理优化分成四层。理解了「两堵墙」就能看清每层在干嘛模型层权重 / 显存减少要搬的数据量——量化权重用更少 bit、压缩。本质是降带宽墙的压力。⚠️ 注意很多赛题禁止持久化量化/剪枝且「运行时不改变读取字节数的伪量化」对 memory-bound 没用。框架层KV-cache / 调度管理 KV cache、batch 调度。对多请求收益大对单请求长输出KV 读在 decode 里占比可能很小收益有限。算子层tuning / 融合让单个或多个 kernel 更高效——算子融合减少中间结果往返显存降带宽墙、kernel tuning 提升带宽/算力利用率。memory-bound 场景下融合和访存优化是主战场。硬件层Tensor Core / 低精度用专用矩阵单元和低精度提算力——主要帮compute-bound如 prefill 的 GEMM。对 memory-bound 的 decode低精度计算帮不上瓶颈在访存不在算力但低精度存储减字节能帮——前提是合规。一句话先用 Roofline 判断每个热点撞的是哪堵墙再去对应的层找手段。拿打 compute 的招去治 memory-bound必然无效。8. 在 HIP / 国产加速卡DCU 等上怎么落地前面的「两堵墙 Roofline profile」方法论是硬件无关的——带宽墙的物理本质在 NVIDIA 和 AMD / 海光上一模一样。但工具链和术语不同国产 DCU 走的是ROCm / HIP体系HIP 是对标 CUDA 的编程模型hipcc编译绝大多数 CUDA 代码可一键hipify移植。给一份对照和注意点。8.1 术语 / 工具对照CUDA → HIP / ROCm能力 CUDA (NVIDIA) HIP/ROCm (AMD / 海光 DCU) ---------------- ------------------ --------------------------- 监控 nvidia-smi rocm-smi / hy-smi(海光) kernel profiler Nsight Compute rocprof / rocprofiler / omniperf 设备 / 架构信息 deviceQuery rocminfo 编程模型 CUDA C HIP C (hipcc 编译) 线程束 warp 32 线程 wavefront 64 线程 (gfx9 wave64) shared memory __shared__ LDS (Local Data Share) BLAS 库 cuBLAS/cuBLASLt rocBLAS / hipBLASLt 注意力 / DNN cuDNN MIOpen 低精度矩阵单元 Tensor Core Matrix Core (gfx9 的 MFMA/mmac 类指令)8.2 在 HIP 上做 Roofline / 瓶颈判断方法论和第 4 节完全一样只是换工具抓 kernel 指标用rocprof导出 trace或用omniperfAMD 的性能分析套件能直接出 Roofline 图、给 memory/compute throughput对应 Nsight Compute。峰值带宽B BB怎么拿一是查卡的 HBM 规格二是写一个纯读带宽探针——一个「只把数据读进来、几乎不算」的 kernel测它达成的 GB/s这就是该卡实际可持续的读带宽上限比理论规格更可信用来当 Roofline 斜坡的真实斜率。判断逻辑不变实测带宽 / 峰值带宽 利用率越接近 1 越到墙。HIP 上判断「decode 到没到带宽墙」 ───────────────────────────────────── 探针带宽 跑「纯读 kernel」测得的可持续 GB/s # 该卡真实上限 理论下界 每 token 权重字节 / 探针带宽 利用率 理论下界 / 实测 TPOT if 利用率 ≈ 1: 已撞 HBM 墙访存侧到顶 else: kernel 没喂满带宽还有肉8.3 HIP 编程里影响优化的关键差异wavefront 64gfx9跨线程规约、__shfl类操作的宽度是 64 不是 32搬移 CUDA kernel 时这点最容易错。访存合并 / 对齐用宽 load如 128-bit 的dwordx4保证 coalesced对齐到 128B流式只读数据可用 non-temporal load 绕开 cache。double-buffer / 软件流水memory-bound kernel 提带宽利用率的核心手段——用 LDS 或寄存器做双缓冲预取下一块数据与当前计算重叠把 global load 延迟藏起来这点 CUDA/HIP 通用。inline asm 与矩阵单元HIP 也支持 inline 汇编调矩阵核指令但不同代、不同厂商的指令助记符可能不同移植时编译器内建函数__builtin_*在某些定制工具链里可能未声明需用对应汇编指令这是国产卡移植的常见坑。后端生态成熟度很多为 NVIDIA 高度优化的 kernel如某些 flash attention 实现、库的特定路径在 HIP 上可能落到Triton 通用实现或需重编译性能未必最优——要按你的实际形状 A/B 实测选后端别默认照搬。8.4 国产 DCU海光 gfx9 系部署注意容器必须映射设备节点--device/dev/kfd --device/dev/dri否则框架看不到卡。架构 overrideHSA_OVERRIDE_GFX_VERSION要和实际 gfx 架构号对应rocminfo查填错算子会崩。库的形状敏感性rocBLAS/hipBLASLt在不同 batch / 矩阵形状上表现可能和 NVIDIA 很不一样——尤其推理decode 那种 M1 的瘦长矩阵-向量乘库的默认 kernel 选择不一定最优务必 A/B。新算子库的适配面一些为较新架构优化的算子库如 AITER在老 gfx 架构上可能反而更慢按实测开关别盲目开。一句话换到 HIP / 国产卡方法论Roofline 判墙、盯带宽利用率一个字不用改改的只是工具rocprof / omniperf、要注意 wave64 与库后端选择。「带宽墙」是 HBM 的物理属性decode 的 memory-bound 不会因为换成 HIP 就消失——这也是为什么国产卡上做 LLM 推理优化主战场同样是「把权重读带宽喂满」。9. 推荐参考GitHub / 资料vLLMgithub.com/vllm-project/vllm—— 主流推理框架PagedAttention、连续批处理、CUDA Graph读它的 attention/线性层实现能学到工程化的瓶颈处理。FlashAttentiongithub.com/Dao-AILab/flash-attention—— 注意力的访存优化经典tiling online softmax 减少 HBM 往返理解「算子融合降带宽墙」的范本。flash-linear-attention (FLA)github.com/fla-org/flash-linear-attention—— 线性注意力 / 状态空间模型的高效 kernel新架构Mamba/GatedDeltaNet 等的递归 decode 优化。NVIDIA Nsight Compute官方 kernel 级 profiler直接给 Memory/Compute Throughput判断瓶颈的首选工具AMD 对应rocprofiler/rocprof。Roofline 模型原文Williams et al.,Roofline: An Insightful Visual Performance Model(CACM 2009) —— 想吃透 roofline 看这篇。mit-han-lab系列github.com/mit-han-lab量化、稀疏、KV cache 等推理优化工作的集合竞赛找思路的好去处。ROCm / HIPgithub.com/ROCm/ROCm、github.com/ROCm/HIPAMD / 国产卡的开源计算栈CUDA→HIP 移植看 HIP 文档与hipify工具。omniperfgithub.com/ROCm/omniperfROCm 上的性能分析 Roofline 工具对应 NVIDIA 的 Nsight Compute国产卡判瓶颈用它。Tritongithub.com/triton-lang/triton跨 NVIDIA / AMD 的 kernel 语言很多框架的算子用它写HIP 后端也能跑是看「融合 kernel 怎么写」的好入口。10. 小结GPU 有两堵墙计算墙compute-bound和带宽墙memory-bound由算术强度I II与脊点I ridge P peak / B I_{\text{ridge}}P_{\text{peak}}/BIridge​Ppeak​/B决定。「GPU 一直有 kernel 在跑」不能判断瓶颈——要看带宽 / 算力利用率timeline 的「满」只是「没空闲」。用Roofline profiler三步走归并热点 → 判断撞哪堵墙 → 算利用率看还有多少肉。LLM 推理prefill 偏计算墙影响 TTFTdecode 是带宽墙影响 TPOT输出越长吞吐越由1/TPOT决定。单请求长输出的赛题主攻 decode 的权重读带宽利用率调度参数基本无效用带宽利用率判断有没有到硬件天花板。觉得有用点个赞收藏⭐。看懂「两堵墙」你做推理优化就不会再「拿错锤子砸错钉子」了。