第二十课：如何使用 Nsight Compute 分析 CUDA Kernel

导语：这一课进入实际性能诊断：如何从 Nsight Compute 的报告中判断 Kernel 是计算受限、带宽受限，还是因为延迟与调度停顿而变慢。

Nsight Compute 是面向单个 CUDA Kernel 的性能分析器，可以查看：

Kernel 执行时间
SM 与显存吞吐率
Warp 调度状态
缓存和内存事务
Occupancy
流水线利用率
Warp Stall Reason
源码与 SASS 对应关系
Roofline

完整程序中哪个 Kernel 最耗时，通常先通过时间线工具定位；确定关键 Kernel 后，再使用 Nsight Compute 深入分析。当前 Nsight Compute 同时提供图形界面 ncu-ui 和命令行工具 ncu。

一、Nsight Systems 和 Nsight Compute 的分工

两者解决的问题不同。

Nsight Systems

观察整个应用程序：

CPU 在做什么
GPU 在做什么
Kernel 何时启动
数据传输是否与计算重叠
不同 Stream 是否并发
哪个 Kernel 占用时间最长

Nsight Compute

深入某个 Kernel：

为什么这个 Kernel 慢
计算单元是否饱和
显存是否饱和
Warp 为什么无法发射
访存是否合并
Shared Memory 是否冲突
寄存器是否限制 Occupancy

正确流程通常是：

Nsight Systems
定位关键 Kernel
        ↓
Nsight Compute
分析关键 Kernel 内部瓶颈

NVIDIA 当前工具流程也支持从系统时间线选中 Kernel，再转入 Nsight Compute 进行单 Kernel 分析。

二、分析前如何编译程序

建议使用优化版本，并保留源码行号：

nvcc -O3 -lineinfo kernel.cu -o app

其中：

-O3
启用优化

-lineinfo
保留源码与 GPU 指令之间的行号对应

不要使用下面的版本进行正式性能测试：

nvcc -G kernel.cu -o app

-G 面向设备端调试，会抑制大量优化，生成的性能通常不能代表正式版本。

如果还想查看编译资源，可以使用：

nvcc -O3 -lineinfo \
    -Xptxas=-v \
    kernel.cu -o app

重点查看：

Used registers
spill stores
spill loads
shared memory
stack frame

三、最基本的 ncu 命令

直接分析程序：

ncu ./app

只分析名称中包含 vector_add 的 Kernel：

ncu \
  -k regex:vector_add \
  ./app

只分析匹配到的第一次启动：

ncu \
  -k regex:vector_add \
  -c 1 \
  ./app

其中：

-k / --kernel-name
过滤 Kernel 名称

-c / --launch-count
限制分析的 Kernel 启动次数

当前命令行工具支持精确名称或 regex: 正则表达式过滤，-c 只统计已经匹配过滤条件的 Kernel。

保存分析报告

ncu \
  -k regex:vector_add \
  -c 1 \
  -o vector_add_report \
  ./app

然后使用图形界面打开：

ncu-ui vector_add_report.ncu-rep

也可以直接：

ncu \
  --open-in-ui \
  -k regex:vector_add \
  -c 1 \
  ./app

ncu 可以把分析结果打印到终端或保存为报告文件，ncu-ui 可以打开该报告进行图形化查看。

四、Set 和 Section 是什么

GPU 可收集的性能计数器非常多，不可能每次都低成本地全部读取。

Nsight Compute 将相关指标组织成 Section，例如：

SpeedOfLight
LaunchStats
Occupancy
SchedulerStats
WarpStateStats
MemoryWorkloadAnalysis
ComputeWorkloadAnalysis
InstructionStats
SourceCounters

多个 Section 又组成 Set。

默认情况下使用较轻量的：

ncu --set basic ./app

收集完整信息：

ncu --set full ./app

查看当前版本支持的 Set：

ncu --list-sets

查看支持的 Section：

ncu --list-sections

当前工具默认收集 basic Set，而 --set full 会收集完整 Section；Section 越多，分析开销通常越大。

只选择需要的 Section

例如：

ncu \
  --section SpeedOfLight \
  --section LaunchStats \
  --section Occupancy \
  --section SchedulerStats \
  --section WarpStateStats \
  --section MemoryWorkloadAnalysis \
  --section SourceCounters \
  -k regex:my_kernel \
  -c 1 \
  -o my_report \
  ./app

这比对程序中的所有 Kernel 使用：

--set full

更可控。

五、为什么分析过程会重复执行 Kernel

GPU 同时可用的硬件计数器数量有限。

如果请求的指标太多，Nsight Compute 可能把指标拆成多个 Pass：

Pass 0：收集一组指标
Pass 1：重新执行 Kernel，收集另一组指标
Pass 2：再次执行
……

为了保持每次执行输入一致，工具可能需要保存和恢复 Kernel 能访问的内存。

因此你可能看到：

Profiling kernel - 46 passes

这并不表示原程序执行了46遍，而是分析器为收集指标进行了 Replay。请求更多 Section、分析更多 Kernel，都会显著提高分析成本。

这会带来三个注意事项：

不要用完整分析运行时间代替程序正常运行时间；
被分析 Kernel 最好具有确定性；
尽量只分析一个有代表性的 Kernel 实例。

六、正确的性能分析顺序

不要一打开报告就直接看 Warp Stall。

推荐顺序是：

1. Kernel 是否真的耗时
2. Grid 是否足够大
3. 计算或内存吞吐率是否接近上限
4. Roofline 属于哪个区域
5. 调度器是否能够持续发射
6. Warp 在等待什么
7. 哪一行源码造成这些等待

因为同一种 Stall，在不同上下文中含义不同。

例如：

Long Scoreboard 很高

如果调度器每周期仍能充分发射其他 Warp，这个等待可能已经被完全隐藏，并不是首要瓶颈。

NVIDIA 文档也明确建议：只有当调度器无法持续发射指令时，才重点分析 Stall Reason。

七、第一步：看 Kernel Duration

最直接的指标是：

Duration

它表示该次 Kernel 执行所需时间。

但单独看 Duration 没有意义，必须有比较对象：

优化前：120 μs
优化后：80 μs

或者与理论需求比较：

处理数据量：1 GB
时间：1 ms
有效带宽约：1 TB/s

分析时应确认：

输入规模一致；
GPU 型号一致；
编译选项一致；
数据内容具有代表性；
Kernel 已经预热；
没有调试编译；
多次运行结果稳定。

八、Launch Statistics：先看任务规模和资源

Launch Statistics 通常包含：

Grid Size
Block Size
Threads
Registers Per Thread
Static Shared Memory
Dynamic Shared Memory
Waves Per SM

这一部分回答：

Kernel 是否向 GPU 提供了足够多的工作？

LaunchStats 汇总 Grid、Block 和资源需求；这些启动配置会直接影响设备利用率。

1. Grid 太小

假设 GPU 有120个 SM，Kernel 只有30个 Block：

30 个 SM 有任务
90 个 SM 空闲

此时即使每个 Block 内部效率很高，整块 GPU 也无法充分利用。

报告可能表现为：

SM Throughput 低
Memory Throughput 低
Achieved Occupancy 可能不低
但 Grid Size 很小

问题不在单个 SM，而在全局任务数量不足。

2. Waves Per SM

可以把一轮能够同时占据所有 SM 的 Block 称为一个 Wave。

例如：

GPU 有 100 个 SM
每个 SM 同时放 2 个 Block

一轮容量约为：

100\times2=200\text{ Blocks}

如果 Grid 只有100个 Block：

Waves < 1

部分容量没有被使用。

如果 Grid 为450个 Block：

第一轮：200
第二轮：200
第三轮：50

最后一轮会产生尾部效应。

3. 资源是否过多

Launch Statistics 还应检查：

每线程寄存器是否过多
每 Block Shared Memory 是否过大
Block 是否过大

这些资源可能限制一个 SM 能同时驻留的 Block 和 Warp 数量。

九、第二步：看 Speed Of Light

SpeedOfLight 提供 Compute 和 Memory 两类高层吞吐率。

可以理解为：

Compute Throughput
计算资源达到理论持续峰值的百分比

Memory Throughput
相关内存资源达到理论持续峰值的百分比

该 Section 会进一步显示不同子单元中利用率最高的贡献者。

四种基本情况

情况一：Memory 高，Compute 低

Compute Throughput：20%
Memory Throughput：90%

大概率是内存受限。

常见于：

数组复制；
向量加法；
Elementwise 算子；
简单激活函数；
低复用的稀疏访问。

优化方向：

减少数据读写
融合 Kernel
提高数据复用
降低数据位宽
改善合并访存

情况二：Compute 高，Memory 低

Compute Throughput：90%
Memory Throughput：30%

大概率是计算管线受限。

常见于：

高复用矩阵乘法；
Tensor Core GEMM；
大量特殊函数；
密集卷积。

优化方向：

使用更高吞吐率数据类型
使用 Tensor Core
减少无效计算
平衡执行管线
提高指令级并行

情况三：Compute 和 Memory 都低

Compute Throughput：25%
Memory Throughput：20%

这并不表示程序没有瓶颈。

更可能是：

Grid 太小；
Warp 大量等待依赖；
Occupancy 太低；
访存延迟没有隐藏；
分支发散；
同步等待；
指令获取问题；
任务负载不均衡。

此时要继续看：

Scheduler Statistics
Warp State Statistics
Launch Statistics

情况四：Compute 和 Memory 都较高

说明 Kernel 同时对计算和内存系统产生较大压力。

这可能是：

已经接近良好平衡；
某个共享资源成为瓶颈；
不同阶段分别受不同资源限制。

需要继续查看子流水线和 Memory Chart，而不能仅根据两个总百分比下结论。

十、Speed Of Light 不等于“有效计算比例”

例如：

Compute Throughput = 90%

不一定表示90%的运算都有用。

如果程序执行大量冗余 FMA，仍可能把计算单元跑满。

同样：

Memory Throughput = 90%

可能包括大量未合并、重复或无效的数据流量。

所以还要问：

是否做了正确且必要的计算？
是否搬运了必要的数据？

Profiler 只能告诉你硬件在忙什么，不能自动判断算法工作是否有价值。

十一、第三步：看 Roofline

Roofline 将三个量放到同一张图中：

\text{compute performance}

\text{memory bandwidth}

\text{arithmetic intensity} = \frac{\text{FLOPs}}{\text{Bytes}}

图中：

纵轴：FLOP/s
横轴：FLOPs/Byte
斜线：内存带宽上限
水平线：计算峰值上限
交点：Ridge Point
圆点：当前 Kernel

Nsight Compute 的 Roofline 用斜线表示 Memory Bandwidth Boundary、水平线表示 Peak Performance Boundary，二者交点是 Ridge Point。

Memory Bound 区域

如果点位于 Ridge Point 左侧：

算术强度低
每搬运大量数据只做少量计算

性能上限近似为：

P\leq B\times AI

优化重点应是增加算术强度：

Shared Memory Tiling；
寄存器复用；
Kernel Fusion；
减少中间结果写回；
降低精度；
避免重复加载。

Compute Bound 区域

如果点位于 Ridge Point 右侧：

算术强度足够高
计算峰值成为主要上限

优化重点包括：

Tensor Core；
更低精度；
FMA/MMA 利用率；
减少分支和无效指令；
平衡计算流水线；
增加 ILP。

点距离 Roofline 很远

如果当前点距离对应上限很远：

既没有跑满带宽
也没有跑满算力

通常存在其他限制：

延迟；
并行度不足；
指令依赖；
不规则访存；
同步；
管线吞吐率；
任务负载不均衡。

Roofline 中点所在区域指出主要理论限制，而点到边界的距离表示潜在改进空间。

十二、第四步：分析 Memory Workload

MemoryWorkloadAnalysis 会分析：

Global Memory；
Local Memory；
Shared Memory；
L1/TEX；
L2；
Device Memory；
System Memory；
Peer Memory；
各级请求、吞吐率和命中率。

Memory Chart 中的连线可以显示指令数、请求数、数据传输量和链路峰值利用率；具体可见单元取决于 GPU 架构。

内存限制的三种不同形式

Nsight Compute 将内存瓶颈进一步区分为：

Mem Busy
某个内存硬件单元长时间处于忙碌状态

Max Bandwidth
某条数据通路的带宽接近极限

Mem Pipes Busy
发射内存指令的管线吞吐率接近极限

也就是说，“内存受限”并不一定只是 HBM 带宽跑满，也可能是 Load/Store 指令过多，或某一级缓存、端口和管线饱和。

十三、分析 Global Memory

重点检查：

DRAM Throughput
L2 Throughput
L1/TEX Throughput
Cache Hit Rate
Requests
Sectors
Bytes
Load/Store Instructions

需要回答以下问题。

1. 显存带宽是否接近上限

如果：

Device Memory Throughput 很高

说明片外显存带宽可能已经成为瓶颈。

此时继续提高 Occupancy，通常不会显著提高性能，因为显存已经无法提供更多数据。

2. 请求数据和实际传输数据是否差距过大

假设一个 Warp 真正需要：

128\text{ Bytes}

但由于地址分散，实际传输了：

1024\text{ Bytes}

说明大量带宽被无效数据占用。

常见原因：

Warp 地址不连续；
访问步长过大；
结构体布局不合理；
地址未对齐；
随机 Gather/Scatter。

优化方向是改善合并访存，而不是简单增加缓存。

3. L1/L2 命中率是否合理

高命中率通常有利，但不能单独评价性能。

例如：

L2 Hit Rate 很高
L2 Throughput 也达到极限

Kernel 可能受 L2 带宽限制。

又例如：

L1 Hit Rate 低
DRAM Throughput 也低

可能是随机小事务，无法产生足够内存并行性。

所以应同时观察：

命中率
传输量
吞吐率
请求数量
Warp 等待情况

十四、检查 Local Memory 和寄存器溢出

Local Memory 常由以下情况产生：

动态索引数组；
大型局部结构；
寄存器不足导致 Spill。

Local Memory 逻辑上属于单线程，但物理上位于设备内存地址空间，其性能特性接近 Global Memory，并经过缓存。

如果报告中出现大量：

Local Load
Local Store

应检查：

-Xptxas=-v

是否报告：

spill stores
spill loads

还应查看源码中是否存在：

float array[large_size];
float x = array[dynamic_index];

优化方向可能是：

减少同时活跃变量；
减少过度循环展开；
拆分 Kernel；
固定小数组索引；
调整每线程输出数量；
不要过度限制寄存器。

十五、检查 Shared Memory

Shared Memory 分析重点包括：

Load/Store 数量
Bank Conflict
Wavefronts
吞吐率
Short Scoreboard
MIO Throttle

理想情况下，一个 Warp 的请求能够以较少的 Wavefront 完成。

如果存在 Bank Conflict，同一个 Warp 的访问可能被拆成多次处理：

一次逻辑 Shared Memory 指令
              ↓
多个物理 Wavefront

优化方向：

改变 Shared Memory 布局；
增加 Padding；
调整线程到数据的映射；
使用广播模式；
使用更宽、数量更少的 Load；
避免同一 Warp 多线程访问同一 Bank 的不同地址。

十六、第五步：看 Scheduler Statistics

每个 Warp Scheduler 维护一组 Warp。

报告中常见四个概念：

Theoretical Warps
理论上由启动配置允许的 Warp 数

Active Warps
已经驻留在调度器上的 Warp

Eligible Warps
当前依赖已经满足、可以发射的 Warp

Issued Warps
本周期真正被选择发射的 Warp

如果某周期没有 Eligible Warp，调度器就无法发射指令，这说明延迟没有被充分隐藏。

情况一：Active Warps 很少

可能原因：

Occupancy 低；
Grid 太小；
寄存器使用过多；
Shared Memory 使用过多；
Block 太大；
Block 数量上限。

此时先解决资源和启动配置问题。

情况二：Active Warps 很多，Eligible Warps 很少

说明 Warp 虽然已经驻留，但大部分都在等待：

等待 Global Memory
等待 Shared Memory
等待前一条指令
等待 Barrier
等待执行管线

这时 Occupancy 继续提高未必有效，因为大量 Warp 可能都在等待相同资源。

需要查看 Warp Stall Reason。

情况三：Eligible Warps 很多，Issued 达到上限

说明调度器一直有候选 Warp 可以发射。

此时调度侧通常比较健康，瓶颈更可能是：

某个执行流水线饱和；
显存带宽饱和；
指令吞吐上限；
算法执行了过多工作。

情况四：Eligible Warps 很多，但并非全部都被选中

这通常是正常现象。

调度器每周期发射能力有限，即使很多 Warp 都 Ready，也只能选其中部分。

对应的：

Not Selected

并不一定是坏 Stall。大量 Not Selected 通常表示候选 Warp 足够多，延迟隐藏能力较好。

十七、Warp Stall Reason 到底怎么看

Warp Stall Reason 表示：

某个 Warp 为什么无法发射下一条指令。

但它并不自动等于 Kernel 的性能损失。

假设 Warp 0 在等待显存，但调度器立即发射 Warp 1：

Warp 0：Long Scoreboard
Warp 1：Issued

GPU 没有空闲，因此 Warp 0 的等待被隐藏了。

所以第一原则是：

先看调度器有没有空闲发射槽
再看 Stall Reason

Warp Stall 数据由调度器状态采样获得；官方建议仅在调度器不能持续发射时，将 Stall 作为主要优化目标。

十八、Long Scoreboard

Long Scoreboard

通常表示 Warp 正在等待来自 L1TEX 路径的数据依赖，包括：

Global Memory；
Local Memory；
Texture；
Surface。

典型代码：

float x = A[index];
float y = x * 2.0f;

第二条指令必须等待 A[index] 返回。

如果 Long Scoreboard 很高，同时 Eligible Warp 很少，说明长延迟访存没有被隐藏。

优化方向：

合并访存；
提高数据局部性；
提高 L1/L2 命中；
Shared Memory Tiling；
增加独立计算；
软件预取；
增加适量 Warp；
降低 Local Memory Spill；
减少指针追踪式访问。

十九、Short Scoreboard

Short Scoreboard

通常表示等待 MIO 路径相关依赖，常见来源包括：

Shared Memory；
特殊数学指令；
动态分支；
某些 Warp 级操作。

如果 Shared Memory 使用较多，应首先检查：

Bank Conflict
Wavefronts

高 Short Scoreboard 经常与 Shared Memory 冲突或连续依赖有关。

二十、LG Throttle

LG Throttle

表示本地或全局内存指令的相关队列压力过大。

常见原因：

Global Load/Store 指令过于密集；
Local Memory Spill；
频繁的小宽度访存；
算术指令过少，访存指令连续堆积。

优化方向：

减少冗余 Load/Store
合并多个窄访问
减少 Spill
让计算与访存交错
提高数据复用

官方说明中，LG Throttle 对应 Local/Global Memory 指令队列拥塞。

二十一、MIO Throttle

MIO Throttle

表示 Memory Input/Output 指令队列压力较大。

可能来自：

Shared Memory 指令；
特殊数学指令；
动态分支；
某些同步或 Warp 指令。

如果由 Shared Memory 引起，可以考虑：

减少 Shared Memory 指令数量；
使用更宽的访问；
增加寄存器复用；
检查 Bank Conflict。

二十二、Math Pipe Throttle

Math Pipe Throttle

表示多个 Warp 都想使用同一类数学执行流水线，导致该管线供不应求。

例如：

大量 FP32 FMA
大量 FP64
大量整数乘法
大量特殊函数

如果某一条管线利用率很高，而其他管线较低，说明指令组合过于集中。

优化方向可能是：

使用 Tensor Core；
降低精度；
减少该类运算；
调整算法；
利用其他执行管线；
提高 Warp 数以隐藏流水延迟。

二十三、Wait

Wait

表示等待固定延迟的执行依赖。

例如长依赖链：

sum = sum * a + b;
sum = sum * c + d;
sum = sum * e + f;

每一条都依赖前一条的 sum。

优化方向：

float sum0 = ...;
float sum1 = ...;
float sum2 = ...;
float sum3 = ...;

使用多个独立累加器，最后再归并，可以增加 ILP。

还可以：

适度展开循环；
重新排列独立指令；
使用延迟更低的操作；
增加适量可执行 Warp。

官方将 Wait 描述为等待固定延迟执行依赖，并建议通过更多活跃 Warp、代码重构或循环展开隐藏延迟。

二十四、Barrier

Barrier

表示 Warp 在等待同一个 Block 中的其他 Warp 到达 CTA Barrier，例如：

__syncthreads();

高 Barrier Stall 可能说明：

同步过于频繁；
同步前工作量不均衡；
分支导致不同 Warp 到达时间差异大；
Block 太大；
部分 Warp 执行长路径。

优化方向：

减少不必要的 __syncthreads()；
平衡各 Warp 工作量；
使用 Warp 级原语；
缩小同步范围；
使用双缓冲和异步 Pipeline；
把大 Block 拆成更独立的小工作组。

二十五、Memory Barrier

Membar

表示 Warp 在等待内存屏障以及相关未完成内存操作。

它与普通 CTA Barrier 不完全相同：

Barrier
等待线程或 Warp 到达同步点

Membar
等待内存操作满足指定可见性和顺序

如果 Membar 很高，应检查：

是否存在不必要的 Fence；
前面的 Store 是否过多；
Store 访问是否合并；
原子和同步设计是否可以改进。

二十六、Branch Resolving 与 No Instructions

Branch Resolving

Warp 等待分支目标和 PC 更新。

可能说明：

分支数量多；
动态控制流复杂；
间接跳转；
严重分支发散。

No Instructions

Warp 没有可供发射的已获取指令，可能来自：

指令缓存 Miss；
代码体积过大；
大量跳转；
Kernel 太短且工作量不足。

过度循环展开和超大融合 Kernel 有时会增加指令缓存压力。

二十七、Occupancy 应该怎么看

Occupancy 定义为：

\text{Occupancy} = \frac{\text{Active Warps}} {\text{Maximum Warps}}

报告通常会给出：

Theoretical Occupancy
理论上根据资源计算的上限

Achieved Occupancy
运行时实际平均活跃 Warp 比例

Occupancy 主要决定隐藏延迟的候选 Warp 数量。高 Occupancy 不保证高性能，但过低 Occupancy 会降低延迟隐藏能力；理论值与实际值差距很大时，还可能存在负载不均衡。

Occupancy 限制来源

报告通常会指出限制来自：

Registers
Shared Memory
Block Size
Warp Limit
Block Limit
Barrier Resources
Cluster Resources

例如：

Theoretical Occupancy：25%
Register Limit：主要限制

说明每线程寄存器使用过多。

但不要立即限制寄存器，因为：

寄存器减少
    ↓
Occupancy 提高
    ↓
可能产生 Spill
    ↓
性能反而下降

什么时候应提高 Occupancy

以下组合表明可能需要提高：

Occupancy 低
Eligible Warps 低
调度器频繁无法发射
Long Scoreboard 或 Wait 很高
计算和内存吞吐率都低

可以尝试：

减少寄存器；
减少 Shared Memory；
调整 Block 大小；
减少每线程工作量；
增加 Grid；
缩小 Tile。

什么时候不必提高 Occupancy

如果：

Occupancy 只有 50%
Eligible Warps 足够
调度器基本每周期发射
Tensor Core 接近峰值

则50%已经足够。

此时强行提高 Occupancy，可能会破坏寄存器分块和数据复用。

二十八、Compute Workload Analysis

这一部分显示不同执行流水线利用率，例如：

FP32
FP64
Integer
Tensor
Load/Store
Special Function

ComputeWorkloadAnalysis 会报告 IPC 和不同 SM 执行流水线利用率；某条管线利用率很高时，它可能限制整体性能。

单一流水线饱和

例如：

FP32 Pipe：95%
Tensor Pipe：0%
其他 Pipe：较低

说明 Kernel 主要依赖 FP32。

如果算法是矩阵乘法，可以考虑：

FP16/BF16/TF32；
Tensor Core；
库函数；
MMA 指令。

SFU 饱和

如果程序大量使用：

sin()
cos()
exp()
log()
rsqrt()

可能由特殊函数相关管线限制。

即使总体 Compute Throughput 没到100%，某条稀缺流水线也可能已经满载。

INT 管线压力

地址计算、循环变量和索引运算通常会使用整数管线。

某些复杂张量索引 Kernel 可能算术运算很少，但地址计算非常多。

优化方向包括：

简化索引公式；
预计算常量；
使用连续布局；
减少除法和取模；
合并循环；
使用指针递增。

二十九、Instruction Statistics

InstructionStats 显示实际执行的 SASS 指令类型和数量，可以判断：

FMA 是否生成；
Tensor Core 指令是否出现；
Load/Store 是否过多；
分支是否过多；
Local Memory 指令是否出现；
特殊函数是否占比过高。

该 Section 面向实际执行的低级 SASS 指令；指令类型过于集中，通常意味着少数流水线承担了主要压力。

例如一个矩阵乘法本应使用 Tensor Core，但报告中只有大量：

FFMA

没有相应 MMA 指令，说明 Tensor Core 路径可能没有启用。

三十、如何定位到具体源码

使用：

nvcc -O3 -lineinfo ...

后，Nsight Compute 可以把部分指标对应到：

CUDA C++ 源码
PTX
SASS

Source 页面中的 Hot Spot 表可以把高 Stall、内存访问或执行指标定位到具体代码位置。

例如：

sum += A[index] * B[index];

可能对应：

IMAD    地址计算
LDG     加载 A
LDG     加载 B
FFMA    乘加

如果这行显示大量 Long Scoreboard，真正产生等待的通常是其中的 LDG，而不是源代码表面上的乘法。

一行源码可能对应多条 SASS

例如：

C[i] = A[i] + B[i];

可能包含：

计算 A 地址
计算 B 地址
加载 A
加载 B
执行加法
计算 C 地址
写入 C

所以必须展开 SASS 观察哪条指令承担指标。

三十一、诊断案例一：向量加法

Kernel：

C[i] = A[i] + B[i];

假设报告为：

Compute Throughput：8%
Memory Throughput：88%
DRAM Throughput：很高
Global Memory 合并良好
Eligible Warps：足够

结论：

显存带宽受限

这时以下优化通常作用不大：

增加更多计算线程；
强行提高 Occupancy；
展开加法；
减少一两条整数指令。

更有效的方向：

与上下游算子融合
避免中间数组 C
使用 FP16
一次完成更多 Elementwise 操作
减少总读写次数

三十二、诊断案例二：非合并访存

Kernel：

x = A[threadIdx.x * stride];

报告可能是：

Memory Throughput 不高
Long Scoreboard 高
Eligible Warps 很少
实际传输数据远大于有效数据
L2/DRAM 请求数量大

虽然 Kernel 在等待内存，但显存带宽没有达到峰值。

这说明它不是理想的“带宽饱和”，而是：

访存效率差导致的内存延迟受限

优化方向：

重新组织数据；
让 Warp 访问连续地址；
使用转置或预处理；
Shared Memory 重排；
将 AoS 改成 SoA；
提高请求合并率。

三十三、诊断案例三：长依赖链

Kernel：

for (int k = 0; k < K; ++k) {
    sum = sum * a[k] + b[k];
}

报告可能是：

Compute Throughput 不高
Memory Throughput 不高
Wait 高
Eligible Warps 低
Occupancy 尚可

原因：

每次循环都依赖上一轮 sum

可以使用多个累加器：

float sum0 = 0.0f;
float sum1 = 0.0f;
float sum2 = 0.0f;
float sum3 = 0.0f;

for (int k = 0; k < K; k += 4) {
    sum0 += ...;
    sum1 += ...;
    sum2 += ...;
    sum3 += ...;
}

float sum = sum0 + sum1 + sum2 + sum3;

这样增加 ILP。

但也要检查寄存器是否增加过多。

三十四、诊断案例四：Shared Memory 冲突

报告表现为：

Short Scoreboard 高
Shared Memory Wavefronts 多
Bank Conflict 明显
Memory Throughput 总体不一定高

可能的代码：

value = tile[threadIdx.x][fixed_column];

如果地址映射使多个线程落入同一 Bank，就会串行处理。

可能改成：

__shared__ float tile[32][33];

通过 Padding 改变 Bank 映射。

优化后应重新观察：

Wavefronts 是否减少
Short Scoreboard 是否下降
Duration 是否下降

三十五、诊断案例五：同步等待

报告表现为：

Barrier 高
不同 Warp 到达同步点时间差异大
Eligible Warps 低

例如：

if (threadIdx.x < complex_count) {
    // 不同线程工作量差异很大
}

__syncthreads();

先完成的 Warp 必须等待慢 Warp。

优化方向：

让每个 Warp 工作量更均匀；
把不规则任务重新排序；
使用更细粒度任务；
减少 Block 大小；
使用 Warp 级同步；
消除不必要的全 Block Barrier。

三十六、诊断案例六：任务规模太小

报告表现为：

Compute Throughput 低
Memory Throughput 低
Grid 很小
Waves Per SM 很低
单次 Kernel Duration 很短

这时瓶颈可能不是 Kernel 内部效率，而是：

SM 覆盖不足；
Kernel 启动开销；
任务切分太细。

优化方向：

合并多个小任务
Kernel Fusion
批量处理
Persistent Kernel
CUDA Graph

Kernel 启动和 CPU—GPU 时间线问题应结合 Nsight Systems 判断，而不是只看单 Kernel 内部指标。

三十七、一个实用的诊断决策树

可以按照下面的顺序。

三十七、一个实用的诊断决策树

图示：三十七、一个实用的诊断决策树

三十八、不要孤立地看 Stall 百分比

假设报告为：

Long Scoreboard：40%
Not Selected：35%
Wait：10%

不能简单得出：

40% 的执行时间浪费在内存上

因为 Warp Stall 指标通常是调度状态或采样结果，不等于整个 SM 空闲时间。

正确问题是：

调度器是否还有其他 Eligible Warp？
执行单元是否仍然在工作？
该 Stall 是否位于关键源码位置？

三十九、不要孤立地看 Cache Hit Rate

例如：

L2 Hit Rate：95%

不一定代表性能好。

可能是程序反复读取同一批数据，但读取次数本身远超必要数量。

又例如：

L2 Hit Rate：20%

对于只读取一次的大型流式数组，低命中率可能完全正常。

应根据算法预期判断：

数据本来是否应该复用？
是否值得留在 Cache？

四十、不要把 100% Occupancy 当作目标

性能目标是：

\text{shorter execution time}

而不是：

\text{higher Occupancy}

高性能 GEMM 可能使用大量寄存器保存输出 Tile，Occupancy 不高，但 Tensor Core 利用率极高。

相反，一个简单 Kernel 可能有100% Occupancy，却因为随机访存而很慢。

四十一、分析结果需要与源码修改形成闭环

推荐每次只修改一个主要因素：

版本 A：原始实现
版本 B：改善合并访存
版本 C：加入 Shared Memory
版本 D：调整 Tile
版本 E：增加寄存器分块

然后比较：

Duration
Speed Of Light
Roofline
Memory Traffic
Eligible Warps
Stall Reason
Registers
Occupancy

Nsight Compute 支持将某次结果设为 Baseline，并在报告中比较当前结果与基线差异。

四十二、推荐的最小分析命令

第一步，轻量分析：

ncu \
  --set basic \
  -k regex:my_kernel \
  -c 1 \
  ./app

第二步，针对性收集：

ncu \
  --section SpeedOfLight \
  --section LaunchStats \
  --section Occupancy \
  --section SchedulerStats \
  --section WarpStateStats \
  --section MemoryWorkloadAnalysis \
  --section ComputeWorkloadAnalysis \
  --section SourceCounters \
  -k regex:my_kernel \
  -c 1 \
  -o my_kernel_analysis \
  ./app

第三步，打开报告：

ncu-ui my_kernel_analysis.ncu-rep

第四步，在确认确实需要时再使用：

ncu --set full ...

四十三、本课核心结论

第一，Nsight Compute 用于深入分析单个 CUDA Kernel，而不是首先观察整个程序时间线。

第二，正式性能分析应使用优化编译，并加入 -lineinfo 建立源码与指令对应。

第三，应通过 Kernel 名称和启动次数过滤，只分析有代表性的关键实例。

第四，收集大量指标可能触发 Kernel Replay，因此分析运行时间不能直接作为正常执行时间。

第五，性能分析应先看：

Duration
Launch Statistics
Speed Of Light
Roofline

然后再深入 Warp Stall。

第六，Compute Throughput 高通常表示计算资源压力较大，Memory Throughput 高通常表示内存系统压力较大。

第七，计算和内存吞吐率都低时，应检查任务规模、Eligible Warp、依赖和同步。

第八，Roofline 通过算术强度区分 Memory Bound 和 Compute Bound。

第九，Memory Bound 需要继续区分：

带宽真正饱和
访存事务低效
访存延迟未隐藏
内存指令管线饱和

第十，Scheduler Statistics 中最关键的关系是：

Active Warp
    ↓
Eligible Warp
    ↓
Issued Warp

第十一，Long Scoreboard 通常与 L1TEX 路径的长延迟数据依赖有关。

第十二，Short Scoreboard 常与 Shared Memory、特殊函数或 MIO 操作有关。

第十三，Barrier 表示 Block 内 Warp 同步等待，Wait 常表示固定延迟执行依赖。

第十四，Not Selected 通常不是坏事，它可能说明调度器有足够多 Eligible Warp。

第十五，Occupancy 只表示驻留 Warp 比例，不直接等于硬件利用率或性能。

第十六，最终优化必须以 Duration 是否下降为判断标准，并同时验证数值正确性。

下一课将进行完整的实战分析：

第二十一课：从一个低效 Reduction Kernel 开始逐步优化

包括：

原子累加版本
相邻线程归约
分支发散问题
Shared Memory Bank Conflict
树形归约
Warp Shuffle
每线程处理多个元素
向量化加载
多级归约
如何使用 Nsight Compute 验证每一步优化

一、Nsight Systems 和 Nsight Compute 的分工

Nsight Systems

Nsight Compute

二、分析前如何编译程序

三、最基本的 ncu 命令

保存分析报告

四、Set 和 Section 是什么

只选择需要的 Section

五、为什么分析过程会重复执行 Kernel

六、正确的性能分析顺序

七、第一步：看 Kernel Duration

八、Launch Statistics：先看任务规模和资源

1. Grid 太小

2. Waves Per SM

3. 资源是否过多

九、第二步：看 Speed Of Light

四种基本情况

情况一：Memory 高，Compute 低

情况二：Compute 高，Memory 低

情况三：Compute 和 Memory 都低

情况四：Compute 和 Memory 都较高

十、Speed Of Light 不等于“有效计算比例”

十一、第三步：看 Roofline

Memory Bound 区域

Compute Bound 区域

点距离 Roofline 很远

十二、第四步：分析 Memory Workload

内存限制的三种不同形式

十三、分析 Global Memory

1. 显存带宽是否接近上限

2. 请求数据和实际传输数据是否差距过大

3. L1/L2 命中率是否合理

十四、检查 Local Memory 和寄存器溢出

十五、检查 Shared Memory

十六、第五步：看 Scheduler Statistics

情况一：Active Warps 很少

情况二：Active Warps 很多，Eligible Warps 很少

情况三：Eligible Warps 很多，Issued 达到上限

情况四：Eligible Warps 很多，但并非全部都被选中

十七、Warp Stall Reason 到底怎么看

十八、Long Scoreboard

十九、Short Scoreboard

二十、LG Throttle

二十一、MIO Throttle

二十二、Math Pipe Throttle

二十三、Wait

二十四、Barrier

二十五、Memory Barrier

二十六、Branch Resolving 与 No Instructions

Branch Resolving

No Instructions

二十七、Occupancy 应该怎么看

Occupancy 限制来源

什么时候应提高 Occupancy

什么时候不必提高 Occupancy

二十八、Compute Workload Analysis

单一流水线饱和

SFU 饱和

INT 管线压力

二十九、Instruction Statistics

三十、如何定位到具体源码

一行源码可能对应多条 SASS

三十一、诊断案例一：向量加法

三十二、诊断案例二：非合并访存

三十三、诊断案例三：长依赖链

三十四、诊断案例四：Shared Memory 冲突

三十五、诊断案例五：同步等待

三十六、诊断案例六：任务规模太小

三十七、一个实用的诊断决策树

三十八、不要孤立地看 Stall 百分比

三十九、不要孤立地看 Cache Hit Rate

四十、不要把 100% Occupancy 当作目标

四十一、分析结果需要与源码修改形成闭环

四十二、推荐的最小分析命令

四十三、本课核心结论

第二十一课：从一个低效 Reduction Kernel 开始逐步优化

相关文章

第七课：GPU 是算力受限，还是带宽受限

第二十七课：GPU 架构学习总结与统一分析框架

第二十六课：如何从零设计一个简化的 SIMT GPU