Roofline 实战:算术强度、层级带宽与瓶颈定位(含 AXPY/GEMM 最小复现)
标题
Roofline 实战:算术强度、层级带宽与瓶颈定位(含 AXPY/GEMM 最小复现)
关键词
Roofline、Arithmetic Intensity(FLOPs/Byte)、Memory-bound、Compute-bound、层级 Roofline(HBM/L2/L1)、Nsight Compute、AMD Omniperf、Intel Advisor、FMA=2 FLOPs、Tensor Core
摘要
一张 Roofline 图就够把性能问题说清:横轴是算术强度,纵轴是实测性能;点贴斜线是带宽受限,贴水平线是计算受限。本文用可复现的 AXPY(带宽型)与 GEMM(计算型)做对照,给出统一的 FLOPs/Bytes 计数方法(含 FMA 与不同精度通道),并用层级 Roofline 把“被卡在哪一层内存”落到实处。配套三套主流工具的最小流程(Nsight Compute / Omniperf / Advisor),读图-验证-改进一气呵成,避免空谈。
目录
- 方法论速览:两条上限线与一个点——算术强度如何落到工程数据
- 峰值与带宽的获取:理论参数 vs 可达值(拷贝基准与计数器)
- 层级 Roofline:HBM/L2/L1 的多斜线与同一内核的多点位
- 计数规则与精度通道:FMA=2、FP32/TF32/BF16/FP16/FP8 与张量核心
- 最小实验 A(AXPY):合并访问、步长与对齐,让带宽型内核“上抬不走样”
- 最小实验 B(GEMM):尺寸与精度对水平屋檐的逼近;cuBLAS/rocBLAS 快速到位
- 读图与决策:带宽还是算力?三步判定与常见误判(尺寸效应、统计口误、层级混淆)
- 实操清单:采集命令、关键指标表、基线脚本与对比模板
1) 方法论速览:两条上限线 + 一个点(先把“图”刻在脑子里)
Roofline 只做两件事:给出一条计算上限(横向屋檐,芯片在某种精度下的可达 FLOPs/s)和一条带宽上限(斜向屋檐,可达带宽 × 算术强度 AI),然后把你的内核落成一个点。
- 横轴:算术强度 AI = FLOPs / Bytes(做了多少运算/搬了多少字节)。
- 纵轴:实测性能(FLOPs/s)。
判读极其直接:点贴斜线=带宽受限,贴水平线=计算受限。这一模型来自 Williams/Patterson 的经典工作,是至今最实用的性能“第一诊断”。(ACM Digital Library)
落点最快路径:任选一款厂商工具直接画图——NVIDIA Nsight Compute 的 Roofline 视图、AMD Omniperf 的 Roofline 报告、Intel Advisor 的 GPU Roofline。跑一遍就能看到你的点到底被哪条“屋檐”压住。(NVIDIA Docs)
2) 上限别抄参数表:峰值 FLOPs 和“可达带宽”怎么拿
计算上限从来不是一条线就完事——不同精度/计算路径各有一条(FP32、TF32、BF16、FP16、FP8…;是否走张量/矩阵核心)。计数规则通常默认1 次 FMA = 2 FLOPs,这在 Roofline 文献与主流工具里是一致的;你要对照自己使用的精度看对应的“水平屋檐”。(ACM Digital Library)
带宽上限更不能只看 HBM/GDDR 标称值;真正有参考意义的是你的可达带宽(实际测出来的)。做法很朴素:
- 用工具内置的度量/基准去测“拷入/拷出/内核流式访问”的吞吐,把这条斜线画在图上再谈点位;
- Nsight Compute 的 Memory Workload、Omniperf 的分析面板都能给出你这台机器的可达数;Intel Advisor 也支持 GPU Roofline 并给出相应带宽视图。(NVIDIA Docs)
小结:水平线看精度与计算路径,斜线看你机器上真实能喂到的数据速率——两条线都要“量”,别只搬规格书。
3) 层级 Roofline:HBM/L2/L1 各有一条斜屋檐(同一内核会出现多个点)
现代 GPU 是分层存储:HBM/DRAM → L2 → L1/纹理 → 片上寄存/共享存储。同一个内核,有的阶段主要吃 L2,有的阶段不得不回 HBM,因此在“不同内存层”的 Roofline 上会出现不同点位:
- HBM 斜线最“平”,最容易被它卡住;
- 到了 L2/L1,斜率更陡,点就没那么容易贴线。
Nsight Compute 在 UI 里直接提供层级 Roofline与内存分析;AMD 的 Omniperf 会跑一组加速器特定微基准来估算层级屋檐(MI200+ 支持),并生成两套 Roofline(FP32/64 与 I8/FP16 分开,读图更干净);Intel Advisor 的 GPU Roofline 也支持在不同层之间切换。(NVIDIA Docs)
实操建议(两步就够)
- 先看 Roofline 点贴的是 HBM 斜线 还是 L2/L1 斜线;
- 再对照 Memory/Cache 面板的命中率与事务统计,锁定“卡你的是哪一层”。
定位清楚“层”,后续该右移(提 AI/复用)还是该上抬(提指令/并行度)就一目了然。
4) 算清 FLOPs 与 Bytes:别混了层,也别忘了合并访问
FLOPs 怎么算
- 统一用 FMA=2 FLOPs;
- AXPY(
y=a*x+y)每元素 ≈ 2 FLOPs; N×NGEMM ≈ 2N³。
不同精度让“水平屋檐”高度不同,工具会把你的点放到对应精度图层里。(ACM Digital Library)
Bytes 怎么算
- 以“从某一层实际搬了多少”为准:同一内核在 HBM 视图 与 L2 视图 的 Bytes 不同,AI 也会不同;
- 这就是为什么要在工具里确认“这一屏 Roofline 是哪一层”的原因(Nsight Compute 的 Metric Details、Advisor 的层级切换都能明确口径;Omniperf 会给出分层估计)。(NVIDIA Developer Forums)
一个常被忽略的硬伤:合并访问
AI 相同,合并的全局访存(同一 warp 的连续对齐访问合并成更少事务)能让点沿斜线上抬;非合并/跨步访问会把点压下去。这在 CUDA 最佳实践与近月的官方博文里都被标为“高优先级”。你的第一步优化,往往就是把访问排整齐。(NVIDIA Docs)
5) 最小实验 A:AXPY(带宽型)——把“搬运”排整齐,点就沿斜线上抬
要点先说透
AXPY:y = a*x + y。单精度下,每个元素 ≈ 2 FLOPs,但要读 x(4B)+ 读 y(4B)+ 写回 y(4B)≈ 12B,AI≈2/12≈0.167。天生贴带宽斜线。想提速,不是多算,而是把同一 warp 的访问做成连续对齐的合并访问(coalescing)。
两版可复现实验(对齐 vs. 跨步)
// axpy.cu 编译: nvcc -O3 axpy.cu -o axpy
#include <cuda_runtime.h>
__global__ void axpy_good(const float* __restrict__ x, float* __restrict__ y, float a, int n){
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) y[i] = a * x[i] + y[i]; // 连续对齐,合并事务
}
__global__ void axpy_bad(const float* __restrict__ x, float* __restrict__ y, float a, int n, int S){
int i = blockIdx.x * blockDim.x + threadIdx.x;
int j = i * S; // 人为破坏合并:跨步 S
if (j < n) y[j] = a * x[j] + y[j];
}
int main(){
const int N = 1<<26; // ~6.7e7,约 256MB 级数据
const int S = 8; // 坏访问:跨步 8
float *x,*y; cudaMalloc(&x,N*4); cudaMalloc(&y,N*4);
dim3 blk(256), grid((N+blk.x-1)/blk.x);
axpy_good<<<grid,blk>>>(x,y,2.0f,N); cudaDeviceSynchronize(); // 预热
axpy_good<<<grid,blk>>>(x,y,2.0f,N); cudaDeviceSynchronize(); // 好访问
axpy_bad <<<grid,blk>>>(x,y,2.0f,N,S); cudaDeviceSynchronize(); // 坏访问
cudaFree(x); cudaFree(y); return 0;
}
怎么量、看什么
- NVIDIA:
ncu --set full ./axpy→ Roofline 视图 + Memory Workload。两颗点都会贴斜线,但 good 版本更高(更接近可达带宽);DRAM/L2 吞吐明显更高,Global Load/Store 事务更少。 - AMD:
omniperf profile -n run1 -- ./axpy→omniperf analyze -p ./workloads/run1 --gui。Roofline 同理,good 点更高。
扩展 1(略右移)
在块内把相邻元素读入共享存储/寄存器做小批操作(例如两次 AXPY 合并一次写回),减少往返次数,AI 会略向右,点随之沿斜线上抬。
结论
AXPY 的 AI 基本由算法决定;想变快,先把访问排整齐。访问模式决定你能不能贴近“可达带宽”,这就是带宽型内核的核心手感。
6) 最小实验 B:GEMM(计算型)——让点顶到水平屋檐
要点先说透
C=A×B(N×N):FLOPs≈2N³,字节量≈12N²(单精度),AI≈N/6。N 足够大时,点天然在右侧,倾向计算受限。最短路径:直接用成熟库(cuBLAS/rocBLAS),先把点顶到水平屋檐附近,再讨论细节。
可复现实验(NVIDIA 版)
// gemm_cublas.cu 编译: nvcc -O3 gemm_cublas.cu -lcublas -o gemm
#include <cublas_v2.h>
#include <cuda_runtime.h>
int main(){
const int N=4096; float *A,*B,*C;
cudaMalloc(&A,1ll*N*N*4); cudaMalloc(&B,1ll*N*N*4); cudaMalloc(&C,1ll*N*N*4);
cublasHandle_t h; cublasCreate(&h);
const float alpha=1.f, beta=0.f;
// 预热
cublasSgemm(h, CUBLAS_OP_N,CUBLAS_OP_N, N,N,N, &alpha, A,N, B,N, &beta, C,N);
cudaDeviceSynchronize();
// 正式
cublasSgemm(h, CUBLAS_OP_N,CUBLAS_OP_N, N,N,N, &alpha, A,N, B,N, &beta, C,N);
cudaDeviceSynchronize();
cublasDestroy(h); cudaFree(A); cudaFree(B); cudaFree(C);
}
AMD 对应:用 rocblas_sgemm 跑同规模;用 Omniperf 看 Roofline。
怎么量、看什么
- Roofline:N=4096 时,点应很靠右并逼近水平线。把 N 降到 512/1024,点会“掉下来”,甚至回到斜线附近(复用不够、调度/启动成本变显眼)。
- 精度试验:把
cublasSgemm换cublasGemmEx(如 BF16/FP16/TF32),点应进一步上抬,贴近更高的水平屋檐(矩阵/张量核心通道)。
结论
GEMM 的第一原则是尺寸与精度。先让点贴到水平线,再考虑 tile、寄存器压力等细活;太小的矩阵会直接把你“拉回斜线阵营”。
7) 读图与决策:带宽还是算力?三步就够
第一步:看位置
- 贴斜线 → 带宽型。先查访问模式(连续、对齐、事务大小)、是否存在跨步与反复回写。
- 贴水平线 → 计算型。先查尺寸、并行度、是否走到矩阵/张量核心、是否被小核启动开销稀释。
第二步:看“哪一层”在卡你
- 打开层级 Roofline(HBM/L2/L1)。
- 如果点在 HBM 斜线附近:优先做数据复用/合并访问;
- 如果点在 L2/L1 斜线附近:多数是块内数据组织不佳或共享存储冲突,先把块内重排、转置、对齐做干净。
第三步:用一个小改动验证
- 自认带宽受限?把 AXPY 的步长 8→1,点应沿斜线上抬;
- 自认算力受限?把 GEMM 的 N 拉大或启用低精度
GemmEx,点应靠近对应精度的水平线。
判读要能被一个改动证伪或证实,这一步能帮你避免在错误方向上耗时间。
8) 实操清单:命令、指标、对比模板(拎包就能复现)
采集命令
-
NVIDIA:
ncu --set full --target-processes all ./your_app- 打开 Floating-Point Operations Roofline、Memory Workload、Speed-of-Light。
-
AMD:
omniperf profile -n run1 -- ./your_appomniperf analyze -p ./workloads/run1 --gui(Roofline 在 GUI 内)。
-
Intel:Advisor 的 GPU Roofline Insights(GUI/CLI 任选),切换 HBM/L2/L1 视图。
关键观察点
- Roofline:点靠斜线还是水平线;是否随“一个改动”出现预期的上抬/右移。
- 内存面板:DRAM/L2 吞吐、合并访问相关统计(事务大小、对齐/效率)、缓存命中。
- 计算面板:实际 FLOPs(含 FMA 记 2)、是否启用矩阵/张量核心通道、活跃并行度。
对比模板(建议你固定下来复用)
-
带宽型内核:
- A/B 两版:步长=1 vs 步长=8;
- 记录:Roofline 点位、DRAM/L2 吞吐、事务数;
- 预期:AI 基本不变,步长=1 的点沿斜线上抬。
-
计算型内核:
- A/B 两版:N=4096 vs N=1024;或 FP32 vs BF16/FP16(GemmEx);
- 记录:Roofline 点位、FLOPs/s、张量核心是否生效;
- 预期:大 N / 低精度版更贴近水平线。
到此为止
你已经用两颗点把“带宽 vs 计算”的分界跑实了,也学会了在三套工具里复现、验证和对比。接下来可以把这一套方法复制到你手头的热点算子上:先落点,再决策,不盲调。
个人简介
作者简介:全栈研发,具备端到端系统落地能力,专注人工智能领域。
个人主页:观熵
个人邮箱:privatexxxx@163.com
座右铭:愿科技之光,不止照亮智能,也照亮人心!
专栏导航
观熵系列专栏导航:
具身智能:具身智能
国产 NPU × Android 推理优化:本专栏系统解析 Android 平台国产 AI 芯片实战路径,涵盖 NPU×NNAPI 接入、异构调度、模型缓存、推理精度、动态加载与多模型并发等关键技术,聚焦工程可落地的推理优化策略,适用于边缘 AI 开发者与系统架构师。
DeepSeek国内各行业私有化部署系列:国产大模型私有化部署解决方案
智能终端Ai探索与创新实践:深入探索 智能终端系统的硬件生态和前沿 AI 能力的深度融合!本专栏聚焦 Transformer、大模型、多模态等最新 AI 技术在 智能终端的应用,结合丰富的实战案例和性能优化策略,助力 智能终端开发者掌握国产旗舰 AI 引擎的核心技术,解锁创新应用场景。
企业级 SaaS 架构与工程实战全流程:系统性掌握从零构建、架构演进、业务模型、部署运维、安全治理到产品商业化的全流程实战能力
GitHub开源项目实战:分享GitHub上优秀开源项目,探讨实战应用与优化策略。
大模型高阶优化技术专题
AI前沿探索:从大模型进化、多模态交互、AIGC内容生成,到AI在行业中的落地应用,我们将深入剖析最前沿的AI技术,分享实用的开发经验,并探讨AI未来的发展趋势
AI开源框架实战:面向 AI 工程师的大模型框架实战指南,覆盖训练、推理、部署与评估的全链路最佳实践
计算机视觉:聚焦计算机视觉前沿技术,涵盖图像识别、目标检测、自动驾驶、医疗影像等领域的最新进展和应用案例
国产大模型部署实战:持续更新的国产开源大模型部署实战教程,覆盖从 模型选型 → 环境配置 → 本地推理 → API封装 → 高性能部署 → 多模型管理 的完整全流程
Agentic AI架构实战全流程:一站式掌握 Agentic AI 架构构建核心路径:从协议到调度,从推理到执行,完整复刻企业级多智能体系统落地方案!
云原生应用托管与大模型融合实战指南
智能数据挖掘工程实践
Kubernetes × AI工程实战
TensorFlow 全栈实战:从建模到部署:覆盖模型构建、训练优化、跨平台部署与工程交付,帮助开发者掌握从原型到上线的完整 AI 开发流程
PyTorch 全栈实战专栏: PyTorch 框架的全栈实战应用,涵盖从模型训练、优化、部署到维护的完整流程
深入理解 TensorRT:深入解析 TensorRT 的核心机制与部署实践,助力构建高性能 AI 推理系统
Megatron-LM 实战笔记:聚焦于 Megatron-LM 框架的实战应用,涵盖从预训练、微调到部署的全流程
AI Agent:系统学习并亲手构建一个完整的 AI Agent 系统,从基础理论、算法实战、框架应用,到私有部署、多端集成
DeepSeek 实战与解析:聚焦 DeepSeek 系列模型原理解析与实战应用,涵盖部署、推理、微调与多场景集成,助你高效上手国产大模型
端侧大模型:聚焦大模型在移动设备上的部署与优化,探索端侧智能的实现路径
行业大模型 · 数据全流程指南:大模型预训练数据的设计、采集、清洗与合规治理,聚焦行业场景,从需求定义到数据闭环,帮助您构建专属的智能数据基座
机器人研发全栈进阶指南:从ROS到AI智能控制:机器人系统架构、感知建图、路径规划、控制系统、AI智能决策、系统集成等核心能力模块
人工智能下的网络安全:通过实战案例和系统化方法,帮助开发者和安全工程师识别风险、构建防御机制,确保 AI 系统的稳定与安全
智能 DevOps 工厂:AI 驱动的持续交付实践:构建以 AI 为核心的智能 DevOps 平台,涵盖从 CI/CD 流水线、AIOps、MLOps 到 DevSecOps 的全流程实践。
C++学习笔记?:聚焦于现代 C++ 编程的核心概念与实践,涵盖 STL 源码剖析、内存管理、模板元编程等关键技术
AI × Quant 系统化落地实战:从数据、策略到实盘,打造全栈智能量化交易系统
大模型运营专家的Prompt修炼之路:本专栏聚焦开发 / 测试人员的实际转型路径,基于 OpenAI、DeepSeek、抖音等真实资料,拆解 从入门到专业落地的关键主题,涵盖 Prompt 编写范式、结构输出控制、模型行为评估、系统接入与 DevOps 管理。每一篇都不讲概念空话,只做实战经验沉淀,让你一步步成为真正的模型运营专家。
🌟 如果本文对你有帮助,欢迎三连支持!
👍 点个赞,给我一些反馈动力
⭐ 收藏起来,方便之后复习查阅
🔔 关注我,后续还有更多实战内容持续更新
更多推荐

所有评论(0)