1. 从一次调优翻车说起为什么需要 Roofline 模型去年帮一个团队排查推理性能问题他们的场景是把一个视觉模型部署到 GPU 服务器上做批量图片处理。硬件配置不差单卡显存 80GB理论带宽标称 2TB/s 级别FP16 算力也是几百 TFLOPS 的量级。但实测吞吐量只有理论峰值的不到 15%加卡也不见线性提升团队一度怀疑是框架版本问题折腾了两周换驱动、换容器镜像、换推理后端收效甚微。后来我把他们的 kernel 逐个拉出来算了一下计算强度Arithmetic Intensity也就是每读取一字节数据能完成多少次浮点运算再对照硬件的带宽和算力画了张图问题一目了然绝大多数算子落在图的左侧斜坡区域也就是显存带宽受限区算力单元大部分时间在等数据搬运。换句话说瓶颈根本不在算力而在数据供给。这个判断方法就是Roofline 性能模型最朴素的应用。Roofline 模型不是某个具体工具也不是某个框架的插件它是一种性能分析的心智框架。它用一张二维图把“硬件能给你多少”和“你的程序实际需要多少”放在同一个坐标系里对比让你在动手优化之前先搞清楚到底该往哪个方向使劲。对于做 GPU 计算、kernel 算子开发、深度学习推理部署、科学计算加速的人来说这套方法能省下大量盲目试错的时间。这篇文章我会从模型的基本原理讲起把计算强度、带宽墙、算力屋顶这几个核心概念拆开揉碎然后给出完整的实操流程——怎么测带宽、怎么算算力、怎么采集 kernel 的实际数据、怎么画图、怎么读图。最后会分享几个我在实际项目中踩过的坑和总结出来的判断经验。不管你是刚接触 GPU 性能分析的新手还是已经做过一些调优但总觉得缺个系统方法的老手应该都能从中拿到能直接用的东西。2. Roofline 模型的核心原理拆解2.1 一张图看懂三个关键概念Roofline 模型的图形结构其实非常简单横轴是计算强度纵轴是性能通常用 FLOPS 表示。图上有两条线构成一个“屋顶”形状第一条是带宽斜线。它从原点出发斜率为硬件的峰值显存带宽。这条线表达的含义是当你的计算强度很低时性能上限完全由带宽决定性能 计算强度 × 带宽。比如带宽是 1000 GB/s计算强度是 2 FLOP/Byte那性能上限就是 2000 GFLOPS跟你的算力有多强没关系。第二条是算力水平线。它是一条水平线高度等于硬件的峰值算力。当计算强度足够高时数据搬运不再是瓶颈性能上限由算力决定。两条线的交点对应的计算强度就是这台硬件的拐点Ridge Point。拐点左侧是带宽受限区右侧是算力受限区。拐点的计算强度 峰值算力 / 峰值带宽。举个例子某 GPU 峰值 FP16 算力 300 TFLOPS峰值带宽 2 TB/s那拐点就是 300e12 / 2e12 150 FLOP/Byte。也就是说你的 kernel 计算强度必须超过 150才有可能碰到算力屋顶。注意这里的算力和带宽都要用实际可达到的值而不是规格表上的理论峰值。实际峰值受频率、功耗墙、ECC 开销等因素影响通常会打七八折甚至更多。2.2 计算强度到底怎么算计算强度是 Roofline 模型的灵魂参数定义是总浮点运算次数除以总访存字节数。听起来简单但实际算的时候有几个容易搞错的地方。首先浮点运算次数要区分是 FMAFused Multiply-Add还是单独的乘加。一次 FMA 算两次浮点运算一次乘法一次加法很多硬件文档里说的算力就是按 FMA 计算的。如果你数指令的时候把一条 FMA 当成一次运算算出来的计算强度会偏小一半。其次访存字节数要算的是实际发生的数据传输量不是数据规模。比如一个矩阵乘法 C A × BA 是 M×KB 是 K×NC 是 M×N浮点运算次数是 2×M×N×K。但访存字节数取决于你的分块策略理想情况下每个数据只从显存读一次那就是 (M×K K×N M×N) × 4 字节FP32。但实际实现中由于分块和缓存的存在访存模式会复杂得多。我一般用这个公式做快速估算计算强度 浮点运算次数 / (输入字节数 输出字节数)对于逐元素操作比如 ReLU、加法每个元素读一次写一次计算强度极低通常不到 1。对于矩阵乘法计算强度随矩阵规模增大而增大因为运算次数是 O(N³) 而访存是 O(N²)。对于卷积计算强度取决于卷积核大小、通道数、特征图尺寸通常比逐元素操作高但比大矩阵乘法低。2.3 带宽墙和算力屋顶的物理含义带宽墙的本质是数据供给速度跟不上计算消耗速度。GPU 的算力单元SM 里的 CUDA Core、Tensor Core数量有限但速度极快而显存带宽受限于物理接口和内存颗粒频率增长相对缓慢。过去十几年GPU 算力的增长速度远快于带宽增长速度导致拐点不断右移越来越多的算子落入带宽受限区。算力屋顶也不是铁板一块。不同精度FP32、FP16、INT8、FP8的算力差异巨大Tensor Core 和普通 CUDA Core 的算力也差一个数量级。所以画 Roofline 图的时候往往要画多条水平线分别对应不同精度和不同计算单元的峰值。还有一个容易被忽略的点L2 缓存的影响。Roofline 模型最初是针对 DRAM 带宽的但现代 GPU 的 L2 缓存容量和带宽都相当可观。如果一个 kernel 的数据能大部分命中 L2那实际可用的带宽远高于 DRAM 带宽Roofline 图上的斜线斜率会变大。所以严格来说应该针对不同的存储层级画多条斜线。3. 实操前的准备工作摸清硬件底细3.1 实测显存带宽的三种方法规格表上的带宽是理论值实际能跑出多少需要自己测。我常用三种方法从简单到精确依次是方法一用现成工具快速测。NVIDIA 平台可以用bandwidthTestCUDA Samples 里自带AMD 平台可以用rocm-bandwidth-test。这些工具会跑一组不同大小的数据传输给出设备到主机、主机到设备、设备内拷贝的带宽。优点是快缺点是测的是纯拷贝带宽跟实际 kernel 的访存模式有差异。方法二写一个简单的拷贝 kernel 自己测。核心思路是分配两个大数组反复做dst[i] src[i]统计总传输字节数和耗时。关键是要让数组足够大超过 L2 缓存容量循环次数足够多取稳定后的平均值。下面是一个 CUDA 示例__global__ void copyKernel(float* dst, const float* src, size_t n) { size_t idx blockIdx.x * blockDim.x threadIdx.x; size_t stride gridDim.x * blockDim.x; for (size_t i idx; i n; i stride) { dst[i] src[i]; } } // 调用时 grid 大小设为 SM 数量的若干倍block 设为 256 或 512 // 传输字节数 n * sizeof(float) * 2读一次写一次 // 带宽 传输字节数 / 耗时方法三用 STREAM 风格的 Triad 测试。做a[i] b[i] scalar * c[i]访存模式更接近真实计算场景。这种方法测出来的带宽通常比纯拷贝略低但更有参考价值。我实测下来的经验是同一张卡纯拷贝带宽通常能达到规格值的 85% 到 92%Triad 带宽会再低 5% 到 10%。做 Roofline 分析时我倾向于用 Triad 测出来的值作为带宽上限这样更保守也更接近实际。3.2 峰值算力的获取与验证峰值算力可以从硬件规格文档里查但要注意区分精度和计算单元。以 NVIDIA A100 为例精度CUDA Core 算力Tensor Core 算力FP3219.5 TFLOPS-FP1678 TFLOPS312 TFLOPSINT8-624 TOPS如果你用的是 Tensor Core 做矩阵乘法那参考的应该是 312 TFLOPS 而不是 78。如果 kernel 里混合了 Tensor Core 和普通 CUDA Core 操作那要按实际占比加权。验证峰值算力可以用一个计算密集型的微基准比如大矩阵乘法。用 cuBLAS 跑一个 8192×8192 的 FP16 矩阵乘法看实际能达到多少 TFLOPS。如果只能跑到峰值的 60% 到 70%那说明要么矩阵规模不够大要么库的配置没调好。我一般会跑几组不同规模取最大值作为实际可达峰值。3.3 确定拐点位置拐点 峰值算力 / 峰值带宽。用实测值算出来的拐点才是真正有参考意义的。比如实测 FP16 Tensor Core 算力 280 TFLOPS实测带宽 1.8 TB/s那拐点就是 280e12 / 1.8e12 ≈ 155 FLOP/Byte。这个数字告诉你计算强度低于 155 的 kernel理论上都是带宽受限的。你可以拿这个标准去筛你项目里的所有 kernel快速定位哪些有优化空间。提示不同精度的拐点不同。FP32 算力低拐点也低可能只有 20 到 30 FLOP/Byte。所以一个 kernel 在 FP32 下是算力受限换成 FP16 后可能就变成带宽受限了。精度选择本身就会改变瓶颈位置。4. 采集 kernel 实际数据并绘制 Roofline 图4.1 用性能分析工具抓取关键指标要画 Roofline 图你需要每个 kernel 的两个核心数据浮点运算次数和访存字节数。这两个数据可以从性能计数器中获取。NVIDIA 平台用 Nsight Computencu最方便。跑一条命令ncu --metrics sm__sass_thread_inst_executed_op_fadd_pred_on.sum,sm__sass_thread_inst_executed_op_fmul_pred_on.sum,sm__sass_thread_inst_executed_op_ffma_pred_on.sum,dram__bytes_read.sum,dram__bytes_write.sum ./your_program这里分别统计了 FADD、FMUL、FFMA 指令数以及 DRAM 读写字节数。浮点运算次数 FADD FMUL FFMA × 2。访存字节数 读字节 写字节。计算强度 浮点运算次数 / 访存字节数。AMD 平台用rocprof思路类似指标名称不同但含义对应。Intel GPU 可以用unitrace或 VTune。如果你用的是 PyTorch 这类框架不想深入到 kernel 级别也可以用torch.profiler拿到每个算子的耗时和访存量估算。但精度不如直接抓硬件计数器适合做粗筛。4.2 数据整理与计算强度计算抓完数据后我一般整理成这样的表格Kernel 名称浮点运算次数访存字节数计算强度 (FLOP/Byte)实测性能 (GFLOPS)elementwise_add2.1e68.4e60.25120conv2d_3x31.8e92.4e87.5850gemm_10242.1e91.2e71754200softmax5.6e61.1e70.5195计算强度那一列就是横轴坐标。实测性能那一列就是纵轴坐标。把每个 kernel 当成一个点画到图上再叠加带宽斜线和算力水平线Roofline 图就出来了。实测性能的计算方式是浮点运算次数 / kernel 耗时。注意单位统一FLOP 对应秒TFLOPS 对应 10¹² FLOP/s。4.3 用 Python 快速画图我习惯用 matplotlib 画 Roofline 图灵活且容易定制。下面是一段可以直接用的代码import matplotlib.pyplot as plt import numpy as np # 硬件参数用实测值 peak_bw 1800 # GB/s peak_flops 280000 # GFLOPS (FP16 Tensor Core) # 计算强度范围 ai np.logspace(-2, 3, 500) # 带宽斜线 bw_line ai * peak_bw # 算力水平线 flops_line np.full_like(ai, peak_flops) # 实际性能上限取两条线的较小值 roofline np.minimum(bw_line, flops_line) # 绘制 plt.figure(figsize(10, 6)) plt.loglog(ai, roofline, k-, linewidth2, labelRoofline) plt.loglog(ai, bw_line, b--, alpha0.5, labelBandwidth Limit) plt.axhline(ypeak_flops, colorr, linestyle--, alpha0.5, labelCompute Limit) # 标注拐点 ridge_ai peak_flops / peak_bw plt.axvline(xridge_ai, colorgray, linestyle:, alpha0.7) plt.annotate(fRidge Point\nAI{ridge_ai:.1f}, xy(ridge_ai, peak_flops), xytext(ridge_ai*2, peak_flops*0.3), arrowpropsdict(arrowstyle-, colorgray)) # 画 kernel 数据点 kernels { elementwise_add: (0.25, 120), conv2d_3x3: (7.5, 850), gemm_1024: (175, 4200), softmax: (0.51, 95), } for name, (x, y) in kernels.items(): plt.scatter(x, y, s80, zorder5) plt.annotate(name, (x, y), textcoordsoffset points, xytext(8, 5), fontsize9) plt.xlabel(Arithmetic Intensity (FLOP/Byte), fontsize12) plt.ylabel(Performance (GFLOPS), fontsize12) plt.title(Roofline Model, fontsize14) plt.legend(loclower right) plt.grid(True, whichboth, alpha0.3) plt.tight_layout() plt.savefig(roofline.png, dpi150) plt.show()这段代码画出来的图能直观看到每个 kernel 离屋顶有多远。离得越远优化空间越大。4.4 读图的三个层次拿到 Roofline 图后怎么读很关键。我一般分三个层次来看第一层看位置。kernel 点在斜线附近说明带宽受限在水平线附近说明算力受限离两条线都很远说明还有别的瓶颈比如延迟、同步、分支发散。第二层看距离。点到屋顶的垂直距离就是性能差距。差距大不代表一定能优化到屋顶但至少说明有空间。我一般把差距超过 2 倍的 kernel 列为重点优化对象。第三层看趋势。把同一类 kernel 的不同实现画在一起能看出哪种策略更接近屋顶。比如不同分块大小的矩阵乘法计算强度不同位置也不同一眼就能看出最优分块方向。5. 基于 Roofline 的优化策略与实战案例5.1 带宽受限 kernel 的优化手段带宽受限的 kernel优化方向就一个减少数据搬运。具体手段有提高数据复用率。最典型的是分块Tiling。把大矩阵切成小块让每个小块能塞进共享内存或寄存器块内数据反复使用减少对显存的访问。矩阵乘法的计算强度能从 O(1) 提升到 O(√N)就是这个道理。合并访存。确保相邻线程访问相邻内存地址这样一次内存事务能取回更多有效数据。如果访存模式是跨步的或者随机的实际带宽利用率会大幅下降。我见过一个 kernel改成合并访存后带宽利用率从 30% 提升到 85%性能直接翻倍。使用低精度数据类型。FP16 比 FP32 少一半字节INT8 再少一半。在精度允许的前提下换低精度能直接减少访存量等效于提高了带宽。这也是为什么现在推理场景大量使用 FP16 和 INT8。利用 L2 缓存。现代 GPU 的 L2 缓存有几十 MB如果数据能大部分命中 L2实际带宽远高于 DRAM。可以通过调整数据布局、控制 kernel 执行顺序来提高 L2 命中率。5.2 算力受限 kernel 的优化手段算力受限的 kernel优化方向是提高计算效率使用 Tensor Core。如果精度允许把矩阵乘法、卷积这类操作映射到 Tensor Core 上算力能提升一个数量级。但要注意数据布局要符合 Tensor Core 的要求否则会有额外的转换开销。减少指令开销。算力受限不代表所有指令都在做有效计算。地址计算、循环控制、分支判断这些指令也占用发射槽。通过循环展开、减少分支、使用更高效的指令序列能提高有效算力占比。提高占用率。如果 SM 里的 warp 数量太少算力单元会因为等待而空闲。通过调整 block 大小、减少寄存器使用、优化共享内存分配可以提高占用率让更多 warp 并行执行掩盖延迟。指令级并行。在单个线程内让多条独立指令交错执行能更好地利用流水线。比如展开循环后多条 FMA 可以背靠背发射减少依赖停顿。5.3 一个真实案例从带宽受限到接近屋顶回到开头那个视觉模型推理的例子。用 Roofline 分析后发现主要瓶颈是几个逐元素操作和归一化层计算强度都不到 1远在带宽受限区。优化措施第一步把多个逐元素操作融合成一个 kernel。原来 ReLU、加法、乘法是三个独立 kernel每个都要读写一遍显存。融合后只读写一遍访存量降到原来的三分之一。第二步把归一化层的统计计算和归一化操作合并避免中间结果写回显存。第三步对卷积层调整分块策略把计算强度从 5 左右提升到 20 以上虽然还是带宽受限但性能提升了近 3 倍。整体优化后端到端吞吐量从原来的不到 15% 峰值提升到 55% 左右。剩下的差距主要来自一些无法融合的算子和框架调度开销。实操心得融合 kernel 是带宽受限场景下性价比最高的优化手段。但融合不是越多越好融合太多会导致寄存器压力大、占用率下降反而可能变慢。我一般会做几组不同融合程度的对比测试找平衡点。6. 常见问题与排查技巧实录6.1 数据采集阶段的典型坑问题一性能计数器数据不准。有些 kernel 执行时间太短计数器采样可能不准确。解决办法是让 kernel 重复执行多次取平均值。或者用--launch-count和--launch-skip参数控制采样范围。问题二访存字节数统计的是 DRAM 还是 L2。不同指标含义不同。dram__bytes是 DRAM 级别的lts__t_bytes是 L2 级别的。做 Roofline 分析时如果 kernel 数据能命中 L2应该用 L2 带宽作为上限否则会低估性能潜力。问题三浮点运算次数统计遗漏。有些运算可能被编译器优化成整数运算或者位运算计数器抓不到。比如除以 2 可能被优化成移位。这种情况需要看 SASS 汇编确认实际指令。6.2 读图时的常见误判误判一把所有 kernel 都往屋顶上靠。有些 kernel 天生就是带宽受限的比如逐元素操作计算强度就是上不去。硬要优化只会浪费时间。正确的做法是接受它的带宽受限本质转而优化访存效率。误判二忽略延迟和同步开销。Roofline 模型只考虑带宽和算力不考虑延迟。如果一个 kernel 的瓶颈是内存延迟或者线程同步那它在图上可能离屋顶很远但优化方向不是提高计算强度而是增加并行度或减少同步。误判三用理论峰值画图。理论峰值和实际可达值差距很大用理论值画图会导致所有 kernel 看起来都离屋顶很远失去参考意义。一定要用实测值。6.3 常见问题速查表现象可能原因排查方法解决方向kernel 在斜线下方很远访存效率低检查访存模式是否合并优化数据布局合并访存kernel 在水平线下方很远算力利用率低检查占用率和指令效率提高占用率减少指令开销计算强度算出来异常高浮点运算次数统计有误核对 SASS 指令确认 FMA 计数方式计算强度算出来异常低访存字节数统计偏大检查是否统计了 L2 流量区分 DRAM 和 L2 指标优化后性能不升反降融合过度导致占用率下降对比优化前后占用率减少融合程度找平衡点同一 kernel 多次测量差异大频率波动或缓存影响固定频率预热后测量多次测量取稳定值6.4 几个独家避坑技巧技巧一先粗后细。不要一上来就抓所有 kernel 的详细数据。先用框架自带的 profiler 看哪些 kernel 耗时占比高只对 Top 10 的 kernel 做详细分析。大部分性能问题集中在少数几个 kernel 上。技巧二建立基线。优化前先记录当前的计算强度和性能优化后再测一次对比变化。没有基线你无法判断优化是否有效。技巧三关注拐点变化。换硬件或者换精度后拐点会变。原来算力受限的 kernel 可能变成带宽受限优化方向也要跟着变。每次换平台都要重新算拐点。技巧四不要迷信屋顶。Roofline 模型给出的是理论上限实际能达到 70% 到 80% 就已经很不错了。追求 100% 屋顶是不现实的也不必要。技巧五结合其他模型。Roofline 只考虑带宽和算力不考虑延迟、占用率、指令吞吐等因素。实际分析时要结合占用率分析、指令分析、内存分析等多个维度才能全面定位瓶颈。7. 我在实际项目中的几点体会Roofline 模型最大的价值不是那张图而是它强迫你在优化之前先想清楚一个问题这个 kernel 到底受限于什么我见过太多团队一遇到性能问题就加卡、换框架、调参数折腾很久才发现瓶颈根本不在他们以为的地方。花半个小时算一下计算强度画一张 Roofline 图往往比盲目试错一周更有效。另一个体会是计算强度是可以设计的。同样的计算逻辑不同的实现方式计算强度可能差一个数量级。融合 kernel、调整分块、改变数据布局这些手段本质上都是在提高计算强度让 kernel 从带宽受限区往算力受限区移动。所以在写 kernel 的时候脑子里要有一根弦我这么写计算强度大概是多少离拐点还有多远最后分享一个小技巧如果你手头没有性能分析工具或者环境不允许跑 profiler可以用纸笔估算的方式做快速判断。数一下 kernel 里每个线程做了多少次浮点运算读写了多少字节除一下就是计算强度。虽然粗糙但足以判断大方向。我经常在 review 代码的时候就这么干几分钟就能看出一个 kernel 有没有明显的性能问题。这个模型后续还可以往几个方向扩展一是加入 L2 缓存层级画多条斜线二是加入延迟约束画出延迟受限区三是针对特定算子比如注意力机制做定制化的 Roofline 分析。但核心思想不变先搞清楚瓶颈在哪再动手优化。
