我第一次拿到 Nsight Compute 报告时是有点发懵的。满屏的英文缩写、百分比和直方图一眼扫过去全是 Duration、Memory Throughput、Achieved Occupancy、Warp Stall每个数都像结论又都说不出结论是什么。后来调过几个算子、踩过一些坑才慢慢摸清它想告诉我什么。Nsight Compute日常命令行就是ncu是 NVIDIA 官方的 kernel 级性能剖析器它能精确回答一个 CUDA kernel 在 GPU 内部每个环节到底消耗了多少时间、多少带宽、多少指令。这篇文章的目的很单纯把报告里最常见的指标讲明白它们是什么含义、怎么读、如何据此判断内核该往哪个方向优化。无论你是刚接触 CUDA 性能分析的开发者还是已经跑过几次 ncu 但对结果半懂不懂的选手这篇内容应该都有参考价值。需要说明的是性能分析这种事指标永远是服务于定位瓶颈的。我会尽量用实际调优时的观察方式来解释每个指标而不是单纯给一份指标字典。你拿到的报告长什么样、界面叫什么名字不同版本会有细节差异但核心指标的含义是稳定的理解了底层逻辑换显卡、换版本都不慌。1. SOL 视图的两个吞吐指标先判断内核类型再谈优化1.1 SOL 模型的判断逻辑瓶颈由最长的那块木板决定打开 Nsight Compute 的 CPU 报告第一个值得看的页面通常是 SOLSpeed of Light光速视图。它只做两件事测出这个 kernel 实际的计算吞吐和内存吞吐然后分别除以对应硬件的理论峰值得到两个百分比。Compute (SM) ThroughputSM 上算术、逻辑、访存指令实际完成的吞吐量占理论峰值的比例。这里的计算吞吐统计的是指令发射和计算管线利用率。Memory Throughput实际访存流量通常是 L2 到 DRAM 或者 SM 到 L2与硬件峰值带宽的比例不同架构统计口径略有差异但含义都是内存系统被用了多少。SOL 视图最核心的结论是SOL %这一行。它的取值逻辑是max(Memory Throughput, Compute Throughput)也就是取两个吞吐里更高的那个百分比作为内核效率参考。这个取法常让新人困惑为什么不取平均值或两个都要看因为 GPU 里计算和访存是并行管线你的 kernel 不可能同时既算得慢又等内存等得慢最终决定运行时间的一定是更吃紧的那一侧。就像一根水管被中间最窄的一段卡住两头的粗细都不能代表整体流量。实际调优里我先看这两百分比就知道方向。假设某 kernel 算出来Compute Throughput 92%、而Memory Throughput 18%那它就是个纯计算密集型内核优化重点应该放在减少指令、提高 ILP指令级并行、用更便宜的数学指令上反过来如果Memory Throughput 88%、计算吞吐只有 25%再折腾减少 FMA 指令就是浪费功夫真正该做的是改善访存合并、提高缓存命中、压带宽占用。这里还有一个很容易被忽视的点SOL 里的吞吐百分比是相对于该 GPU 的理论峰值不是相对你的程序理想状态。所以一个 kernel 跑到Memory Throughput 80%已经算是相当接近带宽天花板没必要再追求数值上超过 90%因为任何合法的访问模式都会带一些地址开销和读写混合损耗强行压榨那 10% 往往性价比很低。我记得有一次把读写分离、去掉多余的非对齐拷贝从 82% 提到 87%再往后怎么优化都纹丝不动后来看了 DRAM 层的 write/read 切换开销才明白那 13% 是内存控制器物理层面的固有成本。1.2 别忽略 Duration、SM 利用率与后面的层面SOL 页面里Duration是内核总耗时这个数字本身没有太多秘密但它和时钟频率强相关。GPU 跑 boost 频率还是基础频率Duration 会差 15% 甚至更多。所以做多次比较时最好记录 GPU 当时的实际核心频率和显存频率不然你优化完发现耗时降了 10%可能有一半是今天散热好、Boost 频率上去了这个后面讲实测环境时会展开。紧接着要看的是SM Efficiency有些版本叫SM Active。这个指标反映的是整个执行期间所有 SM 里有百分之多少的时间处于忙碌状态。如果你的 kernel 只把数据分布在少部分 SM 上跑比如一个很小的 block 数只占满了一颗 GPU 二十个 SM 里的四个那么即使单个 SM 跑得飞快SM Efficiency也会很低整体耗时被直接拉长。这种情况下问题不在指令也不在访存而在并行规模本身不够需要增大 block/grid 维度、让更多 SM 参与工作或者考虑把多个小 kernel 合并成大 kernel 减少启动空隙。有人会遇到一种奇怪现象Memory Throughput和Compute Throughput加在一起超过 100%但SOL %却是 90%。这不是工具算错了而是 SOL 的 Compute 吞吐往往只统计了部分硬件执行单元如 FMA/ALU而内存吞吐统计的是另一套独立硬件单元两者在物理上可以同时达到各自峰值的若干比例理论上限并不是 100% 的互斥分配。理解这一点就行不用纠结。从实操角度第一个排查看完 SOL基本就能把内核归成三类纯计算瓶颈、纯内存瓶颈、以及两边都不高但耗时长的延迟敏感型。前两类方向明确第三类说明内核可能受依赖等待、低占用率或同步开销限制接下来就看占用率和 Warp 状态了。2. 占用率指标理论值、实测值与越高越好的误区2.1 Theoretical Occupancy 是怎么算出来的占用率Occupancy衡量的是 SM 里的 Warp Slot 有多少比例被真正放进了活跃 Warp。理论占用率Theoretical Occupancy是根据内核的静态资源配置和硬件限制算出来的理论上最多能住下多少 Warp跟你程序实际跑起来的表现无关。它由几个硬限制共同决定取最小值每 SM 最大线程数常见架构是 2048 个线程也就是 64 个 Warp。每 SM 最大 Block 数一般在 16 或 32老架构可能更低。每线程寄存器数量如果内核每线程用 64 个寄存器SM 上寄存器文件总量又是 65536 个那么能放的线程就是 1024占用率直接腰斩到 50%。每 Block 共享内存用量类似共享内存总量除以每 block 的字节数得出 block 数量限制。举个例子你感受下计算方式。假设显卡每 SM 支持 2048 线程内核设置每 block 512 线程、每线程 40 个寄存器、每 block 共享内存 20KBSM 总寄存器 65536、总共享内存 100KB。线程限制是 2048/512 4 block寄存器限制是 65536/(512×40) 3.2向下取到 3 block共享内存限制是 100/20 5 block。三者取最小最多 3 个 block 同时驻留也就是 1536 线程理论占用率 75%。你如果能从每线程 40 个寄存器压到 32 个寄存器限制变成 4 block理论占用率就能到 100%。在 Nsight Compute 的 Occupancy 页面它会直接把每项限制、当前内核的资源用量、以及占用率随 Block 大小变化的曲线图全部列出来。如果你发现理论占用率只有 25%去这个页面看一眼通常一眼就能看出是寄存器还是共享内存在卡脖子然后针对性优化寄存器太紧就调__launch_bounds__或-maxrregcount共享内存太紧就砍共享内存或者分块调度block 太小就调大 block。2.2 Achieved Occupancy 与理论值之间的落差说明了什么理论占用率是静态上限实际运行时的占用率叫 Achieved Occupancy也就是在全部执行周期里平均下来每个周期处于 Active 状态的 Warp 数占硬件 Warp Slot 的比例。理论上限 100% 不代表实测一定接近 100%两者之间的落差是很有价值的信息。常见的落差原因有三种。第一是同步屏障kernel 里大量使用__syncthreads()一旦执行快的 Warp 已经到达屏障等别人它就从 Active 变成 Stalled实测占用率自然往下掉第二是资源退出空档Block 算完退出、新 Block 还没完全调度进来的那段时间里 Warp Slot 是空的第三是分支发散和依赖等待Warp 内在等待长延迟访存结果时如果隐藏不了调度器没法持续发射新指令活动 Warp 虽然在 Slot 里但实际处于等待状态这会影响 Warp State 页面的等待数据也会让有效工作比例降低。我见过不少Theoretical Occupancy 100%但Achieved Occupancy 50%的 kernel原因是全局内存访问依赖链太长每个线程 load 完立刻要用中间又没有其他独立 Warp 来填充等待周期。这种时候光增加理论占用率没用因为调度器还没有足够多的可发射Warp 来填满洞。真正有效的动作是提高并行度更多 Warp 数、减少长时间 stall改善局部性、预取、向量化或者削减同步频率。2.3 实战里占用率到底要调到多少这是我最早进入误区的地方一度以为占用率必须拉到 100%于是拼命压寄存器、缩共享内存结果性能反而更差。原因是过度限制寄存器会导致寄存器溢出到 local memory访存量暴增而把每 block 线程数调过大又可能引入不必要的共享内存 bank conflict。占用率和性能之间是曲线关系不是直线正比关系。对于访存密集型 kernel达到 50%~75% 的占用率通常就足够把带宽延迟藏住了再往上收益很小对于计算密集且指令依赖短的 kernel高占用率意义也不大反而因为 L1 缓存被更多 Warp 瓜分而伤害命中率。比较稳妥的做法是先用ncu --set full或直接看默认 SOL 报告里的 Achieving Occupancy如果它在 90% 以上就别再把时间花在抠寄存器上如果低于 60%同时 SOL 的 Memory/Compute 吞吐都不高说明延迟没藏住才值得去提升占用率。这个阈值不是绝对的但它帮我把精力投入真正影响性能的地方。3. Warp 状态和 Stall 指标把时间花在哪一秒说清楚3.1 Stall 采样的基本读法Nsight Compute 的 Warp State / Scheduler Statistics 页面是定位延迟问题的核心。它做的事情可以这么理解GPU 运行期间硬件会随机采样正在活跃的 Warp 和正处于停滞状态的 Warp记录此时它们停在哪里、为什么停。采样量足够大之后工具把停滞原因按比例统计出来就能告诉你这段时间里 Warp 们到底为什么等着。页面里通常有几类关键数字。Active Warps Per Scheduler表示每个调度器平均有多少活跃 Warp越多说明并行填充越好Issued Warp Per Scheduler表示每个调度器每周期成功发射的指令数最大就是 1Warp Cycles Per Issued Instruction是每条发射指令摊到的 Warp Cycles这个值越大说明每个 Warp 发射下一条指令的间隔越长等待比重就越大。接下来是一长串Stall Reason例如 Barrier、Long Scoreboard、Short Scoreboard、Wait、Math Pipe Throttle、MIO Throttle、LG Throttle、Branch Resolving、No Instruction、Not Selected、Sleeping 等。这些指标本质是采样占比不是精确的周期计数但对判断瓶颈方向非常可靠。下面逐个讲。3.2 常见 Stall 原因逐项释义与优化方向Long Scoreboard等待全局内存返回。这个是最常见的 Stall 原因之一说明 Warp 发出的 global load/store 尚未从较远的内存层级L2 或 DRAM回来。它本身不代表访问效率低只代表延迟没藏住。优化思路和访存优化一致合并访问、提高 L1/L2 命中、让访存和计算尽量错开、减少 load 次数、使用向量化加载比如 float4。Short Scoreboard等待共享内存、常量命中、纹理等短延迟操作的返回。相比 Long Scoreboard它的延迟小得多但如果大量 Warp 都在排队访问共享内存同样会形成明显等待。重点排查共享内存 bank conflict、是否把高频数据塞进了常量缓存。Wait固定延迟的流水线依赖通常是指令结果还没算完就被下一指令使用比如连续的 FMA 依赖链。优化方向是重建计算顺序、增加独立表达式、让互相依赖的操作间隔更远。Barrier等待__syncthreads()之类的同步。Barrier 占比高说明线程负载不均衡或同步点过多检查循环里有没有不必要的同步或者条件分支导致一部分线程需要等另一部分完成大量计算。Math Pipe Throttle数学流水线吞吐不够指令都排到了同一个数学单元比如 SFU/64 位运算上。某些特殊函数sin、cos、exp、log、除法吞吐很低大量使用会让指令堵在管道入口。优先用更便宜的近似指令或混合精度同时也要看它是否配合高占用率刷爆了某条专用指令通道。MIO ThrottleMIOMemory Input/Output指令队列满了。MIO 管道负责共享内存、常量缓存、一些全局内存与特殊指令的处理如果访存指令太密集队列就满了后续 Warp 进不来。常见原因是共享内存访问和全局访问混合轰炸、地址计算复杂导致访存指令过多。LG ThrottleLGLoad/Global指令队列满了通常是大量全局 load/store 指令在短时间内集中发射。这和 Long Scoreboard 的区别在于一个是在入口排队、一个是在等结果返回。遇到这种情况可以尝试减少访存条数、批量处理独立访存、用更宽的向量加载。Branch Resolving分支目标解析耗费的等待。如果代码里大量复杂的循环转移、switch、嵌套分支分支解析器会忙不过来。把明显能提级循环的条件判断拿到循环外、减少发散分支通常有帮助。No InstructionWarp 处于活跃状态但当前没有可执行的指令常见于指令缓存I-Cache未命中后取指延迟或者极短的 kernel 里取指令跟不上执行。把被频繁调用的内联函数收缩、减少超大函数体能降低这个比例。Not SelectedWarp 可以执行但调度器每个周期只能挑一个 Warp 发射其他可运行的 Warp 只能排队。这其实不完全是坏消息它说明并行度足够只是调度器繁忙。如果这个比例极高且其他 Stall 都不高可以试着降低一点占用率或减少每条指令的依赖反而能减少调度碰撞。SleepingWarp 在等待独立线程调度等机制一般出现在有independent thread scheduling或异常处理代码里常规 kernel 里占比很低。3.3 不要被单一 Stall 原因带偏场景化分析看 Stall 原因时最容易犯的错是哪个高就优化哪个。实际上很多 stall 是连锁反应。举个例子我在调一个归约 kernel 时发现 Long Scoreboard 占比 45%理所当然去优化全局访存合并改成向量加载后 Long Scoreboard 降到 30%但出现了新的高占比 MIO Throttle因为 float4 加载让 LSU/MIO 队列更拥挤。最后同时减少了访存指令总数、又不把所有 load 堆在一个循环开头才把两个指标都压下来。正确姿势是结合上下文看如果 Achieved Occupancy 很低同时 Long Scoreboard 很高优先加占用率藏延迟如果占用率已经很高但 Long Scoreboard 依旧高说明单 Warp 访问的延迟真的过长去看这是不是非合并访存、L2 未命中多如果 Not Selected 占比很高说明并行度饱和了再堆占用率意义不大。把这些指标当成互相补充的证据链而不是孤立的红灯。4. 从 L1 到 DRAM内存层级指标关键含义与典型访问模式4.1 Nsight Compute 报告里的内存指标家族内存相关指标分布在 Memory Workload 页面和 SOL 下方的小节里主要围绕三层缓存L1包含在 SM 内的统一缓存、L2片上多 SM 共享的最后一级缓存、DRAM显存。常见指标大概有这些DRAM Throughput实际访问显存的流量相对峰值比例。它最接近内存瓶颈的直接证据。Memory Throughput在部分架构下统计的是 L2 侧的吞吐包含了 L1 miss 走向 L2 的流量理解它时要知道它不等同于 DRAM 带宽。L1 Cache Hit Rate和L2 Cache Hit Rate全局访问命中率。命中的 Load 不需要走到 DRAM但即便 L1 命中也有 Short Scoreboard 延迟L2 命中则有 Long Scoreboard只是路径比 DRAM 短。Sectors/Request平均每个访存请求访问了多少个 32 字节扇区。这个数很关键它可以暴露访存合并程度。理想情况下一次 warp 的 32 个线程访问连续地址只需要 4 个扇区128B cache line 由 4 个 sectors 组成如果每个线程访问的是步长为 8 字节的地址需要的扇区数可能翻倍甚至十几倍带宽利用率直线下降。Local Memory/Register Spill寄存器溢出后临时数据落到 local memory本质是全局内存的量溢出越多隐式访存越多Performance 损失往往很明显。Shared Memory Bank Conflicts共享内存同一时钟周期内多个地址落在同一 bank 导致串行化报告里会给出冲突次数和百分比。4.2 全局内存访问效率sectors 与 cache hit rate判断一段全局访存代码写得好不好最直接的指标是Sectors/Request和 L1/L2 命中率的组合。假设一个 Warp 内每个线程读一个 float416 字节32 线程一共需要 512 字节连续数据正好是 4 个 128B 缓存行、16 个 sectorsSectors/Request 应该是 16。如果每个线程按 stride 间隔读 16 字节那么同样 32 线程分布到十几条缓存行上Sectors/Request 会膨胀到几十甚至上百实际从 DRAM 搬运的数据量是实际需求的好几倍DRAM Throughput会迅速飙高但 L1/L2 命中率反而一塌糊涂。这种模式一旦被你在报告里用高 DRAM Throughput 低 L2 命中率 高 Sectors/Request锁定基本就是非合并访问回去把数据布局改成结构体数组SoA或者按连续 tile 加载即可。还有一种隐藏较深的情况读多写少或者写多读少的 kernelDRAM Throughput 可能没有打满但 DRAM 侧的读写切换开销非常大表现为 Duration 长、Memory Throughput 又不到峰值一半。此时报告里 Memory Workload 的读写比例数据会给你提示优化方向往往是分批次把读和写拆开、增加中间缓冲减少 DRAM 总线来回切换。4.3 共享内存 bank conflict 与 local memory 泄漏共享内存是 GPU 上唯一可以手工控制低延迟存储的地方但它有 bank 结构32 个 bank每 bank 4 字节一个周期最多同时服务 32 个不同 bank 的地址。如果同一 warp 里两条地址落在同一个 bank就会被分成两拍甚至更多拍执行。Nsight Compute 的报告里会直接给出Shared Memory Bank Conflicts数值比如某个内核的共享内存访问冲突率 200%意思是平均每条共享内存指令多付了两倍时间。排查时记住一点常见冲突来自按threadIdx.x直接索引共享数组的列式访问或者 stride 恰好是 bank 数量的整数倍。解决办法是给数组加 padding比如__shared__ float s[8][32 1]让行末多出的 1 个元素把 bank 对齐关系打破。这个改动很小但效果立竿见影。Local memory 溢出则是另一类容易被忽略的问题。在 Nsight Compute 的Memory Workload里看到 local 相关指标不断增长或者汇编里有STL/LDL指令说明寄存器不够用数据被挤到 local memory。local memory 物理上落在全局内存上虽然先经过 L1但每次都带来可观的延迟和带宽开销。处理方法降低每个线程寄存器占用来提升驻留 Warp 数或者反直觉地增加__launch_bounds__里的线程数让编译器更激进压寄存器也可以直接用float4等向量类型让编译器更好地利用寄存器减少溢出。4.4 通过内存指标反推代码问题把几个指标组合起来能做到望闻问切。我下面列一个平时判读用的对应表你可以拿自己的报告对照指标组合现象大概率代码问题优先处理动作DRAM Throughput 高 L2 Hit Rate 低 Sectors/Request 高非合并全局访问、随机散布读改 SoA、向量加载、连续 tileL2 Hit Rate 高 DRAM Throughput 低 性能仍差数据量不适合 L2 缓存、或 L2 抖动分块计算提高局部性、调整 block 顺序Shared Memory Bank Conflict 高共享内存索引 bank 冲突数组 padding 或改变存储布局Local/LDL/STL 指令量大寄存器溢出压寄存器、调 launch_bounds、向量化Memory Throughput 中低 Long Scoreboard 高访存延迟未隐藏并行度不足提高占用率、增强独立并行DRAM Throughput 低但读写切换频繁读写交错太碎读写分组、设计中间缓冲这表不是万能药但足够把方向引到正确位置。尤其注意单纯看L1 Hit Rate 高不一定就是好事它只说明数据在 L1 有复用如果数据根本不需要复用那高命中率可能只是把本该直接走的流量绕了一圈反而不如直接绕过 L1 的ld.global.cg指令。5. 从报告到改代码的完整排查路数5.1 典型优化流程先判型再定位N 轮的调优有点像病理检查顺序很重要乱翻指标只会浪费时间。我自己整理出来的流程是这样的跑一次完整 profile先看 SOL 页面。判断是 Compute-bound、Memory-bound 还是 Latency-bound。Latency-bound 的标准是两边吞吐都不超过 50%且 Duration 很长。第二看占用率。如果 Theoretical Occupancy 低于 60%翻 Occupancy 页面找出限制资源如果 Achieved Occupancy 远低于 Theoretical去 Warp State 里看等待原因。第三看 Warp State / Stall Reason。结合刚才的类型判断选占比前三的 stall 原因对应查找代码里的相关语句。第四看 Memory Workload。验证访存模式重点查 Sectors/Request、Cache Hit Rate、Bank Conflict、Local Memory。改代码后重新 profile和 baseline 对比。一定要保留第一次的 ncu 报告不然改完说不清是代码生效还是频率波动。这套顺序背后的逻辑是粗粒度到细粒度的金字塔先知道卡在计算还是内存再看并行度够不够然后落到具体指令/访存模式最后验证改动。跳步的结果就是瞎猜——有人一看 Long Scoreboard 高就加#pragma unroll结果瓶颈在共享内存 bank conflictUnroll 再狠也白搭。5.2 命令行实操如何稳定地拿到一份可靠报告命令行下建议用这样的基本姿势# 最简单的全量分析生成当前 kernel的报告入口 ncu --set full -o my_report ./my_app # 只看特定 kernel跳过前几次冷启动统计后三次 ncu -k my_kernel --launch-skip 2 --launch-count 3 -o my_report ./my_app # 只看 SOL 相关指标快速判断类型 ncu --section SpeedOfLight_RooflineChart -k my_kernel ./my_app # 导出指标到CSV方便自己汇总对比 ncu --metrics gpu__time_duration.avg,gpu__compute_memory_throughput.avg.pct_of_peak_sustained_elapsed \ --csv -k my_kernel ./my_app有几个参数是我平时必加的--launch-skip和--launch-count。GPU kernel 往往有很多次 launch前几次可能是热身也可能受状态影响只分析目标 kernel 并跳过前几轮能拿到更稳定的数据。-o导出报告文件也很必要ncu 生成的是 sqlite 格式的 .ncu-rep 文件之后可以用ncu --import重新打开全部细节比只看控制台输出强得多。另一个值得关注的参数是--target-processes all。如果你的程序有主进程拉起子进程、或者接了其他多进程框架不加这个参数可能只 profile 到主进程而漏掉真正跑 kernel 的子进程。大多数单进程程序用默认模式就行但遇到奇怪的空报告时排查一下这个参数不吃亏。5.3 我在实际 profile 中踩过的一些坑第一次拿到一份特别差的带宽数据我把罪魁祸首归咎于访存代码折腾两天优化毫无起色。后来发现同一台机器上另一个显卡任务在跑把 DRAM 带宽抢走了一大半我的 kernel 再优化也是给别人做陪跑。从那以后我做两步跑 profile 前先确认独占 GPUnvidia-smi看有没有其他进程占用必要时把 GPU 切成 exclusive 模式隔离出来的数字才算数。动态频率也是个大坑。GPU 核心频率会受功耗和温度影响同样一段代码冬天冷启动时 Boost 频率高跑得快夏天连续跑多次频率可能更低。对比两组 profile 结果之前先看报告里的SM Frequency和Memory Frequency。如果两次频率差超过 5%直接比较 Duration 就不公平最好把指标换算成周期数或者看吞吐百分比以规避频率波动带来的噪声。小 kernel 的 profile 结果尤其要小心。一个只跑 5 微秒的 kernel启动和收尾的开销占比巨大采样样本不足Stall Reason 的统计稳定性差数值波动也大。对这种 kernel要么把--launch-count调到多次要么承认它不值得精细剖析优先用 Nsight Systems 从系统层面看整体启动开销而不是死磕 ncu 的单 kernel 指标。内核代码改完之后比较报告时不要只盯一个数字变没变。我习惯把改前改后两个 .ncu-rep 文件放到 Nsight Compute 的对比视图里同时看 SOL、占用率、Stall 三项确认瓶颈是不是真正转移了。很多时候一个优化只是把瓶颈从 A 点推到 B 点如果你只看到 A 点指标变好而 B 点指标明显恶化性能可能完全没提升甚至倒退。最后补几句实在话用 Nsight Compute 的时间越长越觉得它最厉害的地方不是给出一个性能分而是能逼着你去理解 GPU 的执行模型。指标本身没有魔法每一个数字背后都是硬件在某一层做了什么、等了多久。拿到一份报告先别急着开优化开关把每个高亮指标当成一个问题为什么它高它和旁边那个指标是什么关系如果我把它降下来会不会把别的指标顶上去我个人的实操体会是把 baseline 报告存好每次改动只动一个变量跑完就看那一两个关键指标其他指标的变化只做参考。这样一天下来你能快速建立代码写法到指标表现的对应直觉以后再看到类似报告基本不用细查就能猜出代码大概长什么样。这个积累过程比任何现成的优化清单都值钱。如果你拿到的第一份报告看不懂从 SOL 页面的两个百分比和 Warp State 页面的 stall 原因开始对照本文逐个过一遍很快就能入门了。