写CUDA的人大多经历过这种场景：内核跑通了，结果也对，top命令一开，GPU利用率50%上下，再往上调就不知从哪下手。改block大小、换循环顺序、加restrict，试了半天，要么没变化，要么更慢。我在这种阶段浪费过大量时间，后来才意识到问题不在“试得不够多”，而在于没有把内核在硬件上的真实行为量化出来。NVIDIA Nsight Compute就是干这个的Kernel级性能分析工具，今天这篇入门，我按自己实际踩过的路径来讲，尽量把那些文档里没写透的细节一起说清楚。
Nsight Compute能解开三件我之前一直靠猜的事：内核到底卡在访存还是计算上、SM内部哪些流水线空闲、Warp的stall到底花在等什么。本文面向准备开始用工具优化CUDA程序的人，包括已经会写内核但对Profile结果发怵的同学，以及想用命令行做自动化基准的开发者。
1. 先搞清楚Nsight Compute到底解决了什么问题
1.1 传统的定时和采样为什么不够用
很多初学者优化GPU程序，只会用cudaEvent测整体时间，或者在CPU侧用std::chrono掐表。这两种办法只能告诉你“这段内核跑了多少毫秒”，但完全无法解释“为什么是这么多毫秒”。同样一个内核，换一台GPU，时间变了，你依然不知道为什么。
后来有人用nvprof读硬件计数器，这比纯计时进了一步。但nvprof提供的指标粒度仍然是设备级的，比如总内存吞吐量、总SM利用率。问题是现代GPU上有几十个SM，每个SM内部又有多个子单元——Load/Store单元、整数/浮点单元、张量核心、共享内存、L1缓存等等。内核在一个时钟周期内可能同时发生多个事件：某些Warp在等着访存返回，某些Warp在算数，另一些在发射指令。设备级聚合计数把这些信息全混在一起，自然看不出瓶颈在哪。
Nsight Compute和它是两代思路。Nsight Compute直接在每个SM内部做指令级和Warp级的采样与计数，能告诉你是哪一行源码、哪一条SASS指令占用了最多的时间，以及该指令在流水线中等待了多久。也就是说，它把优化对象从“整个内核”下钻到了“一条指令”，这才叫Kernel级分析。
1.2 和Nsight Systems的分工：别用错工具
我第一次接触到Nsight家族时也有困惑：NVIDIA同时提供了Nsight Systems和Nsight Compute，名字相似，界面上有重叠，到底用哪个？后来形成的判断很简单：
- 当你不知道瓶颈在哪一个阶段、想全局看CPU与GPU整体时间线时，用Nsight Systems；
- 当你已经锁定某个内核，想把它在SM里的行为拆碎、逐指令分析时，用Nsight Compute。
我做应用优化时通常是两层配合：先跑Nsight Systems看整个应用里哪个内核耗时最多，再针对这个热点内核跑Nsight Compute。前者回答“哪里慢”，后者回答“为什么慢”。如果只用Nsight Compute，可能会在一个无关紧要的内核上浪费半天；如果只用Nsight Systems，热点内核内部的结构性问题永远看不到。
2. 环境准备与首次Profile实操
2.1 安装与版本选择
Nsight Compute的安装方式有几种：全栈开发时装的NVIDIA HPC SDK里自带了一套，但版本跟着SDK走，更新慢；我更推荐直接单独下载Nsight Compute独立安装包，这样可以第一时间用上新架构的支持和新增指标。安装包解压后有两个核心可执行文件：命令行版ncu和图形版ncu-ui。
版本上有一点必须提醒：Nsight Compute对GPU架构和CUDA版本有对应要求。条件允许的话，尽量用比CUDA Toolkit略新一点的Nsight Compute版本，因为它通常保留了对更多旧架构的后向兼容，同时对新的Ada、Hopper架构支持更完整。换新GPU时第一件事就是查工具版本的架构支持列表，否则Profile出来一堆解释不了的警告。
注意：Nsight Compute的分析会附加到GPU设备上，和其他同时运行的GPU任务会发生资源争夺。做Profile前关闭不必要的图形界面、浏览器硬件加速等，避免计数器读数不稳定。
2.2 编译选项和第一个采样
先从一个非常常见的saxpy内核开始。下面这段代码看上去没什么好说的，但刚好适合用来解释Nsight Compute报告里的基础概念：
__global__ void saxpy(int n, float a, float *x, float *y) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) { y[i] = a * x[i] + y[i]; } }编译时务必打开调试信息和行号映射。我用nvcc习惯加三个选项：
nvcc -arch=sm_80 -lineinfo -O3 -o saxpy saxpy.cu-lineinfo让Nsight Compute能把SASS指令关联回你写的.cu源码行，这是Source视图能不能用起来的前提。不加这个选项，后面报告里的源代码关联就是空的，分析就变成了纯汇编研究。
第一次采集，最省事的方式是用GUI：启动ncu-ui，新建一个Profile项目，指定可执行文件路径和工作目录，点击运行即可。默认会把应用里所有的内核都收集一遍。假如上面的saxpy内核只运行一次，采集十几秒内就能结束。
命令行版本也很快上手：
ncu -f -o saxpy_prof --set full --launch-count 1 ./saxpy-f是覆盖输出文件，-o saxpy_prof会把后面的数据写成ncu格式的报告文件，--set full表示采集全部指标集合，--launch-count 1表示只测第一次内核启动。第一次跑应用推荐加上--launch-count 1，因为对于分析来说，单次启动足够，还能避免重复数据文件过大。
2.3 读懂Profile采集过程里的状态输出
运行ncu时，它并不安静。终端会打出当前在重放哪一个内核、采了多少Warp、是否有警告。下面这段是我实际跑过的输出片段，为了简化做了删节：
==PROF== Connected to process 9527 ==PROF== Profiling "saxpy" from process 9527 ==PROF== Profiling "saxpy" ... 8/8 ................ 100% ==PROF== Collecting data from kernel "saxpy" (ID 0) ==PROF== Report: saxpy_prof.ncu-rep如果看到==WARNING==开头的行，别急着跳过。常见的警告包括“Clock rate was not locked”和“Sample rate may have been reduced due to high overhead”。前者说明测试期间GPU频率可能有波动，后者说明采样开启时间不足。出现这类警告时，优化结论要留有余地，最好重启干净环境再复测。
命令行采集完毕后，用下面的命令回看结果：
ncu --import saxpy_prof.ncu-rep --page details这样可以直接在终端里翻看关键指标，不必开GUI。不过做深度分析时我还是更推荐打开ncu-ui，交互式勾选指标，视觉效果强得多。
3. 核心指标与瓶颈判断方法
3.1 把GPU想象成一条有多条流水线的加工厂
理解Nsight Compute报告，最好先把GPU内部运行方式做个类比。比如把SM想象成一条加工厂生产线，Warp就是一批待加工的零件。生产线上有多个工位：负责从内存取料的访存工位、负责算术运算的计算工位、负责在Warp之间切换的调度工位。任何一个工位堵了，整线都会慢下来，但你站在车间外只能看到整体产量下降了，看不到究竟是哪个工位堵的。
Nsight Compute干的事情就是在每个工位旁边装一个传感器，告诉你这台机器每个工位的利用率、排队时间、物料等待时间。它报告里的核心指标，本质上就分成三类：吞吐量、占用率、Stall原因。这三类对应着工厂里的产能利用率、工人数量、停工原因。
3.2 必看的4个关键指标
打开报告后，不用眉毛胡子一把抓，先把下面这几个指标记住。
| 指标名 | 含义 | 我判断时的经验阈值 |
|---|---|---|
| Achieved Occupancy | 实际占用的Warp槽位与理论最大槽位的比例 | 低于25%要警惕调度不足 |
| SM Throughput | 各SM单元平均活跃周期占比 | 常看执行单元利用率 |
| Memory Throughput | 各类内存层级（L1/L2/DRAM）的吞吐量 | 和SM Throughput对比找主瓶颈 |
| Warp Stall Reasons | Warp等待的具体原因分布 | 看占比最大的Stall类型 |
以saxpy为例，它的访存和计算都很简单，报告里的SM Throughput和Memory Throughput通常都不高。因为每个线程只做一次乘加、两次读和一次写，数据量远小于带宽上限，所以你会看到Memory Throughput比较高，比如70%到85%，而Compute Throughput只有20%上下。这种形态说明内核是访存密集型的典型例子。
3.3 用Speed of Light视图判断主瓶颈
Nsight Compute的Speed of Light视图是我最常用的一页。它把理论峰值和实际值做对比，用柱状图画出SM吞吐、内存吞吐、L1/L2缓存吞吐。理想情况是至少有一项接近满负荷，说明硬件主要资源在干活。如果所有柱状图都在50%以下，那大概率是延迟受限或工作负载太小，空转时间太多。
有一次我优化一个自定义算子，测试时看到Compute和Memory都只有30%左右，一开始以为工具坏了。后来查Stall原因，发现Long Scoreboard占比超过一半，说明所有Warp都在等访存数据返回，算数单元自然闲着。普通的吞吐量视图没法直接看出这个问题，但Stall采样一眼就暴露了。所以我的建议是：看Metrics视图之前，先看Warp State和Stall采样分布。
Stall原因里几个高频关键词需要区分：
Long Scoreboard：Warp在等待全局或局部内存数据返回，说明访存延迟没被隐藏；Wait：等待固定延迟流水线或内存操作完成；Short Scoreboard：主要是共享内存、常量缓存等短延迟等待；MIO Throttle：访存指令队列满了，吞吐遇到瓶颈；Math Pipe Throttle：计算指令队列满了，计算吞吐是瓶颈。
如果一个内核的Stall主要是MIO Throttle，说明它已经在压榨内存流水线；如果是Long Scoreboard，说明问题更可能在“延迟隐藏不足”，这时单纯堆带宽意义不大，反而应该提高Occupancy或者增加每个线程的并行度。
3.4 Occupancy高不等于性能好
关于占用率，我翻过一个大跟头。早期我以为Occupancy越高越好，于是把所有内核都调成高Occupancy，结果性能不升反降。原因是Occupancy只是Warp槽位的使用比例，不等于每个Warp都高效工作。高Occupancy但每条指令都在同一个内存地址附近打转，或者所有Warp都卡在同一条内存等待上，整体照样慢。
Nsight Compute把理论Occupancy和实际Achieved Occupancy分开显示，这一点很重要。理论Occupancy只和内核的寄存器和共享内存用量有关；实际Occupancy还要受到调度器、指令延迟、同步行为的影响。比如一个只有8个Warp的小内核，即使寄存器再省，它也不可能达到高Occupancy，因为GPU有64个Warp槽位，就喂了8个进去。这时候问题就成了“工作负载本身太小”，而不是内核写错了。
条件允许时，我会用--launch-count多次测量加--set basic来对比不同block size下的Achieved Occupancy变化，找出分水岭在哪。
4. 实战：完整走一遍Nsight Compute分析流程
4.1 准备一个受控的测试程序
为了把分析流程讲完整，我们准备一个稍微复杂点的小程序：矩阵乘法。这个例子能让访存、计算、共享内存多个因素都参与进来，也方便观察不同优化带来的指标变化。
#define TILE 16 __global__ void matmul_tiled(float *A, float *B, float *C, int N) { __shared__ float As[TILE][TILE]; __shared__ float Bs[TILE][TILE]; int bx = blockIdx.x, by = blockIdx.y; int tx = threadIdx.x, ty = threadIdx.y; int row = by * TILE + ty; int col = bx * TILE + tx; float sum = 0.0f; for (int k = 0; k < N / TILE; ++k) { As[ty][tx] = A[row * N + (k * TILE + tx)]; Bs[ty][tx] = B[(k * TILE + ty) * N + col]; __syncthreads(); for (int i = 0; i < TILE; ++i) { sum += As[ty][i] * Bs[i][tx]; } __syncthreads(); } C[row * N + col] = sum; }启动配置里，我用blockDim(16,16)，网格按矩阵大小计算。这个版本比naive版本多了共享内存复用，是分析共享内存负载和bank conflict的合适载体。
4.2 第一次采集：先看宏观再下钻
运行下面的命令行采集命令，把该内核作为目标单独抓出来：
ncu -f -o matmul_basic --set full --kernel-name matmul_tiled --launch-count 3 ./matmul-kernel-name支持正则表达式，内核多的时候很有用。--launch-count 3表示采三次该内核的启动，Nsight Compute默认会自动重放内核来完成所需的多轮计数，这里的三次是为了让报告里的平均数据更稳定。
打开报告后，我首先看Speed of Light。一个没有被优化的矩阵乘法，在合理尺寸下通常表现为Compute Throughput在30%到50%，Memory Throughput也不高，两条都没到顶，这是典型的“并行度不足+共享内存带宽浪费”混合状态。此时如果直接看具体指标：sm__throughput.avg.pct和gpu__compute_memory_throughput.avg.pct，就能确认整体资源利用率。
接着看Launch Statistics里的Block Size和Waves Per CTA。所谓Wave，可以理解为一轮同时能启动的CTA波次。如果Wave数量是小数，意味着最后一个Wave只有部分CTA在跑，SM资源一定会有空窗。做优化时可以适当减小Block Size，让Wave数变成整数，比如4个Wave整数比3.7个Wave利用率高不少。
4.3 直达瓶颈：Source视图和SASS指令映射
当宏观数据定位到Compute或者Memory之后，我习惯切到Source视图，按处理时间排序查看热点指令。这里能直观看到哪些行代码吃掉了最多时钟周期。以矩阵乘法为例，内层循环那一行sum += As[ty][i] * Bs[i][tx]通常会是热点。点开对应的SASS指令，笔记本上能看到取共享内存的指令和浮点乘加指令。如果共享内存指令的Stall数据里Short Scoreboard占大头，说明线程在等共享内存数据返回；如果Bank Conflict相关指标很高，就得考虑改数据布局或padding。
查共享内存bank冲突时，用下面这条指标最直接：
ncu --metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum,l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum ./matmul数字如果不是0，说明同一个Warp里多个线程访问了同一个bank下的不同地址，共享内存访问次数成倍增加。这个指标我可以反复强调：很多开发者在共享内存优化上花了大力气，结果瓶颈其实在bank conflict，白白浪费了改数据布局的时间。
4.4 用Profile数据驱动一次小优化
以我上面给的矩阵乘法为例，定义TILE为16，共享内存数组大小是16x16个float。每行16个float正好铺满16个bank，同一行的不同列访问时，As[ty][i]这种模式在所有线程里是按不同bank访问的，没冲突。但Bs[i][tx]这一项，当内层循环i变化时，同行不同列线程访问不同行的同一列，可能产生冲突。解决办法很经典：给数组第二维加padding，把它从float Bs[TILE][TILE]改成float Bs[TILE][TILE+1]。这样一行跨越17个float，下一行首地址偏移到不同bank，冲突幅度立刻下降。
改完之后重新采集同一次基准，比较前后的bank_conflicts和内核耗时。Nsight Compute报告支持直接排版对比：在GUI里导入两份报告，选中Speed of Light和Source，可以把两个版本的指标并列看。用这种方式做验证，比单纯看运行时间要扎实得多，因为你能看到到底是共享内存访问次数变少了，还是只是GPU频率变高了。
5. 常见问题与排查技巧实录
5.1 Profile结果不稳定：频率和温度在捣乱
刚用Nsight Compute时，我最常遇到的怪现象是两次采集之间内核时间差异居然达到20%。排查到最后，发现GPU的Boost频率受温度影响很大。跑完一次满负载，GPU温度升高，下一次频率下降，时间自然就长了。Nsight Compute虽然支持锁定时钟，但需要在Profile前额外设置。
这里要区分GPU架构。在大部分桌面和数据中心GPU上，可以通过下面的命令锁定频率：
nvidia-smi -lgc 1500跑完恢复默认：
nvidia-smi -rgc对于不支持nvidia-smi -lgc的平台，也不用强求，把报告里的温度、SM时钟频率记下来，只在相同频率条件下的两组结果之间比较，结论仍然有效。
注意：锁定时钟是在GPU全局生效的，测完务必恢复默认。如果忘记恢复，后续其他显存密集任务的表现都会被影响。
5.2 采集不到某些内核：命名和重放机制
Nsight Compute不是每个内核都能顺利分析。它依靠在GPU上重放内核来采集不同计数器，如果内核本身是非确定性的，比如使用随机数生成之后修改控制流，重放时的行为可能和第一次不同，报告数据就失真了。这种场景下尽量减少外部随机性，把随机种子固定下来再采集。
还有一类情况是内核太短。一个内核运行时间只有几微秒时，Nsight Compute需要大量重复采样才能拿到稳定数据，很可能出现“Skipped”或“Insufficient work”提示。这时可以增大数据规模来延长内核时间，或者利用--launch-count配合多次启动的累积效应。
5.3 指标太多看不懂：先做减法
打开完整报告，界面上可能同时列出几十个指标，我第一次看时直接傻眼。后来养成一个固定的阅读顺序，分享给大家：
明确定位问题的步骤：
- 先看
Speed of Light，确定是内存受限、计算受限还是延迟受限的世界； - 再看
Warp State和Stall Sampling，确定Warp到底卡在哪一类事件； - 最后看
Source视图，落到具体指令和行号； - 如果涉及共享内存，才去查bank conflict和共享内存分配细节。
这个顺序能避开大部分噪音。指标永远是为问题服务的，而不是指标本身越全越好。
5.4 动态并行和MIG环境要注意什么
涉及动态并行时，子内核会从设备端启动，Nsight Compute也能捕获，但父内核和子内核的关联需要在报告里仔细定位。MIG（多实例GPU）环境下，每个实例的硬件计数器范围被隔离，使用工具前先确认实例的显存和SM划分符合你的预期，否则计数读到的可能是整个物理GPU的聚合结果，和预期完全不符。
5.5 测量开销对短内核的影响
Nsight Compute是出了名的“重”工具。--set full模式下，它对某些内核对性能的测量扰动可能非常明显。做严格基准时，我一般会同时记录gpu__time_duration.avg和wall_clock，如果两者差距超过正常范围，说明采样流程本身就已经拖慢了应用。此时减小指标集合，改用--set basic，再复核关键时间。
优化是否有效的最终判断标准仍然是应用在真实数据上的运行时间。Nsight Compute用来解释原因、指导方向，最终确认还是要回到一次干净的实战运行。
6. 融入日常开发流程的一些习惯
6.1 把Profile变成提交前的固定步骤
我在团队里推动过一个很简单的规则：涉及CUDA内核的改动，提交前必须附上一份Nsight Compute的Speed of Light截图或者CLI导出的关键指标。这样做有三个好处：
- 防止“看起来更快但说不清为什么”的优化进入代码库；
- 让后续接手的人能在同一组指标上做对比；
- 尽早发现那些只对特定GPU架构有效的优化，避免在异型集群上翻车。
实际操作上，把命令行采集命令写进一个脚本，固定指标集和输出命名规则，这样每次跑完都能直接和上一次结果比对。脚本化的好处还在于可以批量测试不同blockDim或不同编译选项下的表现。
6.2 和Nsight Systems配合的整体视图
Nsight Compute的粒度足够细，但视角太微观。一个应用可能有几十个内核，如果每个内核都做全量分析，时间成本很高。我的做法是先用Nsight Systems跑一次完整应用，拿到每个内核的耗时占比，只对Top 3的热点内核用Nsight Compute做下钻。这样既不会错过全局瓶颈，又控制住了每次优化周期的时间。
有一点经验：Nsight Compute报告里的Duration和Nsight Systems里的同一内核时长会有差异，这是采样机制不同造成的正常现象，不要惊慌。对比时只比较同一工具内部的数据，跨工具绝对数值没有意义。
6.3 关于架构知识的积累
Nsight Compute再强大，也不会替你决定怎么改代码。同一个Stall原因，在安培架构和霍普架构上的对策可能完全不同。我自己的体会是：在使用工具的过程中，逐渐去理解GPU的调度器、执行流水线、内存层级，比熟记任何一组指标都重要。工具只是给这些硬件行为配上了更精确的温度计，最终做判断的还是你对硬件模型的理解。
如果你刚接触这个工具，建议把GPU Blackboard里那些术语翻一翻，理解Scoreboard到底是什么状态跟踪机制，Warp Slot和Scheduler之间怎么配合。基础概念补齐后，再看报告里的指标，很多结论几乎是显而易见的。
7. 选择指标集与命令行自动化的更多细节
7.1 核心指标集如何取舍
Nsight Compute的--set参数决定了一次分析采集多少指标。它们的耗时和干扰是递增的。实际使用中我会按场景选择：
| 集合 | 特点 | 适用场景 |
|---|---|---|
--set basic | 指标不多，开销小 | 快速验证内核是否正常、估算时间 |
--set full | 全指标，数据量大 | 深度分析单个热点内核 |
--set detailed | 更多硬件计数器组合 | 研究特定架构行为 |
| 自定义指标列表 | 用--metrics指定 | 日常回归、CI环境 |
不建议每次都用--set full。一方面它内部会把内核重放很多遍，运行时间明显变长；另一方面大量指标会干扰注意力。我在日常基准中固定使用一个短线指标清单，例如：
ncu --metrics gpu__time_duration.sum,sm__throughput.avg.pct,dram__throughput.avg.pct,sm__warps_active.avg.pct.per_cycle_active ./app这样一条命令跑完，就能快速回到终端看到内核时间、SM利用率、内存利用率、活跃Warp比例四个关键值。符合需求时再上--set full补全细节。
7.2 批处理与报告命名
自动化分析时，报告文件名和输出文件路径要规范。我的脚本一般长这样：
ncu -f -o prof_$(date +%Y%m%d)_${kernel} --set basic --kernel-name ${kernel} --launch-count 3 ./app把内核名、日期记进文件名，方便后续横向对比。CI环境里建议加上--target-processes all或精确匹配进程名，避免误采集到其他辅助进程。
7.3 动态查看：--page参数
命令行模式下，--page决定了打印输出的组织方式。最常用的是--page details和--page speedoftlight。前者输出分节的详细信息，后者只输出Speed of Light概要。我习惯用前者，因为关键指标都聚在一页，省去开GUI的时间。
如果报告已经保存在.ncu-rep里，还可以用--import反复读取不同页面，不用重新跑应用。这个特性对调整分析视角特别方便，也方便团队内部共享报告文件，让没有GPU环境的人也能查看分析结果。
8. 最后再提几个最常被忽略的实操细节
第一个细节是关于重放的。Nsight Compute在采集某些指标时需要多次重放内核，以保证计数器不会溢出。如果你的内核内部有依赖外部状态的逻辑，重放结果可能不一致。所以尽量在Profile版本里把输入数据固定成静态文件，避免随机噪声。
第二个细节是Warm-up。首次调用CUDA API会触发上下文初始化和模块加载，第一次内核启动时间通常偏长。我在命令行里会先让应用弱跑一个预热内核，再用--launch-count 3采集后两次的结果，这样报告里的数据更贴近真实稳态。
第三个细节是数据类型的匹配。分析整数密集型和浮点密集型内核时，sm__inst_executed_pipe_fma和sm__inst_executed_pipe_alu是不同的计数器。不要只盯着总体的SM Throughput，而是去匹配具体执行单元。张量核心的计算要用sm__pipe_tensor_op_{hmma,imma}.avg.pct这类指标，普通浮点指令在张量核心指标里一点都不会体现，容易造成“编辑器说计算几乎空闲，但实际已经在用张量核心全速跑”的误会。
把这几条养成习惯后，Nsight Compute才真正融入了我的日常循环：写内核、跑Profile、看瓶颈、改代码、再跑Profile。它不再是偶尔拿出来看看的玩具，而是定位性能问题的第一反应工具。希望这篇入门详解，能帮你省掉我当年独自摸索时浪费掉的那些时间。
