2. 为什么聚焦到 Kernel 级：先搞清楚工具边界
1.1 项目概述
1.2 场景适配说明
2.1 从定位差异看必要性
3.5 跑一趟 profile 的命令行实录
4.2 定位一行代码或一条指令
5.3 几个经常被问到的“差不多先生”问题
四.1 章节原始问题
该章节存在原始编号问题（#### 1.1 / ### 1.2），不符合 H2/H3 编号规范，需重排后统一输出。 ## 1. 项目概述
1.1 核心需求解析
这标题看着像“又是一个调参工具”，实际用过你就知道，Nsight Compute 这类 Kernel 级性能分析工具，跟 nvidia-smi、Nsight Systems 是完全不同层面的东西。一句话说清定位：nvidia-smi 告诉你 GPU 现在忙不忙，Nsight Systems 告诉你 kernel 之间怎么排队、哪个阶段在等待，而 Nsight Compute 直接钻进单个 kernel 内部，看清楚每一行 SASS 指令到底为什么跑得慢。
这个工具解决的是 CUDA 优化里最扎心的问题：你已经写完了 kernel，跑起来结果也对，但就是觉得不够快。卡在“感觉能优化但不知道从哪下手”的阶段。Nsight Compute 的价值就是给你一份“kernel 体检报告”，把计算吞吐、访存命中、warp 调度、指令发射、寄存器压榨、shared memory bank conflict、尾巴效应这些全部量化，甚至能定位到是哪一条 SASS 指令造成了 stall。
适合谁来学？刚进 CUDA 优化门的人、已经调过一阵子但只会试探性改参数的人、被老板逼着交出性能报告的人。这篇文章我会从工具选型逻辑讲起，把 profile 流程、报告解读、指标定位、常见踩坑串成一条可复用的路径，保证你读完能自己跑一轮完整的 kernel 分析。
1.2 环境准备与版本选择
开始实操前先把环境说清楚。Nsight Compute 有两种形态：独立 GUI 版，以及作为 CUDA Toolkit 一部分安装的命令行版ncu。如果你装了 CUDA Toolkit，默认命令行工具就在bin目录下。GUI 则在 NVIDIA Nsight Compute 独立安装包里，也可以从 Visual Studio / CodeXL 插件里调起来。
版本选型这块有个实际经验——尽量让 Nsight Compute 版本不低于当前驱动所支持的 CUDA 运行版本。版本太旧的话，profile 新架构（比如 Hopper、Ada Lovelace）上的新指标时经常会出现“Unknown metric”或者干脆拒绝生成报告。我自己遇到过 CUDA 12.0 配了旧版 NCU，profile 一个用了async内存拷贝的 kernel 时直接崩掉，换到配套版本的 NCU 就好了。所以第一步建议：ncu --version看一眼版本，再对照一下驱动的 CUDA 版本，兼容性有问题就先升级工具。
提示：CLI 工具
ncu才是日常定点分析的主力，GUI 更适合做整盘报告浏览。下面的流程以命令行新开为主，GUI 作为补充确认手段。
2. 为什么聚焦到 Kernel 级：先搞清楚工具边界
2.1 从定位差异看必要性
很多人在做 CUDA 性能分析时，习惯性先从宏观层面看整体耗时，然后直接拍脑袋优化某个循环。但 GPU 加速的瓶颈往往不是“整体耗时”能说明白的——两个 kernel 总耗时差不多，一个是因为访存等待，另一个是因为指令依赖，你拿着总耗时表根本分不出来。这正是 Nsight Compute 存在的意义：它在 kernel 执行层面做插桩和计数器采集，把每一个 scheduler 每一拍干了什么摊开给你看。
它的采集方式主要依赖硬件性能计数器（硬件 counter）和软件插桩的结合。硬件 counter 是 GPU 微架构本身记录的事件，比如 issued instructions、executed instructions、L1 hit/miss、DRAM read/write、warp stall 原因等。NCU 通过对这些 counter 做多轮 pass 采样、汇聚和分析，最终生成结构化的报告。这个过程跟其他 profiling 工具最大的区别就是“粒度”：Nsight Systems 的时间线只到 kernel 边界，NCU 则深入到指令流。
所以选型逻辑很清晰：如果要做“kernel 之间怎么排布、CPU 与 GPU 之间怎么 overlapped”这种宏观瓶颈，用 Nsight Systems；如果定位到了某一个 kernel 内部执行效率低，立刻切到 Nsight Compute。两个工具分工明确，混着用容易把问题问模糊。我自己在实战中更习惯先用 Systems 做一轮整体扫描，确认哪几个 kernel 占比最大，再用 NCU 做深度定点——省时、省力、定位准。
2.2 它到底采集了什么数据
其实你只要抓住硬件 counter 这个核心，NCU 的一切输出都顺理成章。CUDA 的 GPU 内部有大量可编程的性能计数器，分布在 GPU 子引擎（比如 SASS 指令分发单元、L1/L2 缓存控制器、DRAM 控制器、内存管理单元）上。NCU 通过 CUpti（CUDA Performance Tools Interface）和底层驱动按特定顺序开关这些 counter，采集指定 kernel 上的计数。
采集完的数据按层次组织起来，会生成一份报告。报告里有几个大类信息：
Speed of Light（SOL）：衡量 kernel 执行时 GPU 各子系统的利用率百分比。最常用的两个就是 Compute（SM 里的 ALU 和 FMA 流水线等）和 Memory（访存管线、L2/DRAM 吞吐）。SOL 是一眼判断瓶颈方向的入口。
主导性能限制指标：在 SOL 之外，报告会告诉你是“计算吞吐受限”还是“访存带宽受限”——这对应了经典 Roofline 模型里的两类瓶颈：compute bound 还是 memory bound。
Warp Stall 原因分布：每个 warp 在流水线中停顿（stall）的原因归类。Clang 的统计里最常见的有
Long Scoreboard（等待 memory 数据）、Wait（固定延迟指令）、Short Scoreboard、Math Pipe Throttle、MIO Throttle、LG Throttle、Branch Resolving、Barrier、Membar、Imiss等，每一个原因对应一类架构行为。Occupancy：提交到 SM 上的 warp 数占最大 warp 数的比例。Occupancy 低会导致无法隐藏延迟、而 Occupancy 过高则可能导致寄存器溢出（spill）或 shared memory 不足，所以这不完全是越高越好。
内存访问 / 缓存命中统计：Hit rate、sector efficiency、bank conflict 等细粒度访存指标。
数据汇总之后，NCU 还可以进一步通过 Source Counters 和 Instruction Counters 把性能映射回 CUDA C++ 源码行或 PTX/SASS 指令上，这也是第 4 部分定位“哪一行有问题”的基础。
2.3 与 Roofline 模型的关系
Roofline 是 HPC 领域做 kernel 分析时绕不开的抽象模型，它把硬件性能画成一张“天花板”图：横轴是算术强度（FLOP/Byte），纵轴是可达的 FLOPs。一个 kernel 如果在 Roofline 图里就碰到天花板，说明它在那个硬件限制下已经接近极值；如果离得很远，说明还有优化空间。
NCU 的 SOL 页面里面就内嵌了这个思想，不过实现略不同——它直接对比“当前 kernel 实际达到的 compute / memory throughput”与“各子系统的硬件峰值”。你不用自己去算 FLOP 数，只要看数字就能知道瓶颈到底在哪侧。
实操时一个很有用的参考是：如果 compute throughput 已经达到峰值 85% 以上，就别再死磕计算指令那一侧了，反过来看访存、occupancy、调度之类的维度。反之如果 compute 利用率很低但 memory 很高，那目标就明确了——减少数据搬移、提升缓存命中或者改数据布局。
3. 第一次完整跑通分析流程
3.1 四种 profile 模式怎么选
正式上手之前，先把 NCU 的驱动模式搞清楚，不然你会在命令行上栽跟头。常用参数是--set和--section。
默认模式（Default）：不指定任何 section 时，NCU 会跑一组基础集合。包含了一些核心 section，例如 Compute Workload Analysis、Memory Workload Analysis、Scheduler Statistics、Warp State Statistics 等，可以快速对 kernel 做体检，但指标不完整。
Profile 集合（--set）：快速选中成组配置。比如
--set full会开启几乎所有 section，数据量大但耗时长，适合最终深挖；--set default就是默认集；--set roofline专门生成 Roofline 分析的指标集。用ncu --set full加--list-sets可以看有哪些内置集合。自定义 section（--section）：可以指定单个 section 开启，比如
--section ComputeWorkloadAnalysis --section MemoryWorkloadAnalysis。这种方式针对性最强，profile 时间最短。范围过滤（--kernel-id）：把分析限定在某个 kernel（或按名字正则匹配），避免 profile 一个巨型程序里几百个 kernel 只为了看其中一个。
我的建议：第一轮跑 default 或 roofline 集合先拉大方向，再用自定义 section 深挖细节。直接全量 section 一来耗时翻倍，二来面对几十个指标页很容易迷失重点。
3.2 基础命令行：单 kernel 定点分析
先把命令给出来，注释也写在旁边，方便你直接抄：
# 1. 列出程序里的 kernel 名字，确认要分析的目标 ncu --list-kernels ./my_app arg1 arg2 # 2. 对指定 kernel（按名字数字索引）跑默认分析 ncu --kernel-id 1 --set default ./my_app arg1 arg2 # 3. 只做 roofline 分析 ncu --set roofline ./my_app arg1 arg2 # 4. 保存报告到文件，后面可以反复看 ncu --section ComputeWorkloadAnalysis --section MemoryWorkloadAnalysis \ --export report_file.ncu-rep ./my_app arg1 arg2初次使用建议先用--list-kernels过一遍，知道程序里到底有几个 kernel、名字带什么后缀（比如__$sm_这种模板后缀）。然后按 kernel-id 分析，不然 NCU 默认会把每一 kernel 都跑一遍相关指标，耗时会翻很多倍。
有一点特别提醒：NCU 依赖 PC Sampling 或插桩来写内存数据，profile 本身会显著改变 kernel 的执行行为。你在 profile 模式下看到的绝对耗时往往比正常运行要慢，但这不影响相对量化（比如吞吐百分比、stall 分布），这些指标是硬件计数，基本可信。所以不要拿 profile 时间直接做最终性能报告，要拿 profile 看方向和趋势，拿独立计时做最终验收。
3.3 用 GUI 浏览报告
命令行产出.ncu-rep文件后，一般用 GUI 打开更直观。命令：
ncu --gui report.ncu-rep如果只是终端环境没有 GUI，也可以用ncu --import report.ncu-rep --page details这类命令导出文本摘要。GUI 里最常用的两个视图：
Summary Page：展示 SOL 条形图、occupancy 数据、内核基本信息。适合在最上层确认瓶颈方向，红色标记的项目就是离硬件峰值最远的维度。
Details Page：左侧是 section 树，右侧是具体数据表。这里可以看到 Solver 的详细指标，比如
sm__throughput.avg.pct_of_peak_sustained_elapsed、gpu__compute_memory_throughput.avg.pct_of_peak_sustained_elapsed、sm__warps_active.avg.pct_of_peak_sustained_active等原始指标名，一般建议直接对照 Source 视图看源码映射。
首次打开 GUI 可能会被一堆 NVLINK、NVLink 相关的心跳信息吓到，忽略即可，重点是 SOL. 另外，高版本 GUI 的 Source View 支持直接点击源码行显示 sampled warp stall 和 memory access 位置，这个对定位缓存对齐问题特别有用。
3.4 从 CLI 输出快速读懂报告
为了不用每次都进 GUI，我习惯在命令行跑完直接看关键指标。举例：
ncu --set default ./my_app arg1 arg2输出的典型摘要长这样（加了一点格式，实际 CLI 的排版会稍密集）：
==PROF== Profiling "kernel_name" from process 12345 Duration : 145.86 us Compute Throughput : 63.62% Memory Throughput : 12.45% Occupancy : 42.84% Achieved Occupancy : 40.10% WARP Stall Reasons (% of total stalls) Long Scoreboard : 78.1% Wait : 6.2% Math Pipe Throttle: 5.1% ...看到Compute Throughput 63.62%同时Memory Throughput 12.45%、Long Scoreboard 78.1%，就基本可以断定 kernel 是典型访存延迟受限（memory latency bound），不是带宽受限。对策重点放在合并访存、增加并行度、减少无谓依赖上，而不是去改计算指令顺序。反过来，若 Compute 95%、Memory 20%、stall 集中在Math Pipe Throttle，那就是纯计算 bound，策略变成降低指令数、提高 ILP 或换更高效的数学指令。
3.5 跑一趟 profile 的命令行实录
用一个我手头常见的模拟场景来演示整个流程。假设有一个图像处理类的 CUDA 程序，主计算 kernel 叫execute_filter_pipeline，运行参数是输入长度和迭代次数。
# 第一步：列出 kernel ncu --list-kernels ./image_demo 4096 8 # 输出： # [0] execute_filter_pipeline # [1] execute_postprocess_kernel # 第二步：对主 kernel 跑默认分析 ncu --kernel-id 0 --set default ./image_demo 4096 8 # 输出摘要中计算吞吐约 27%，内存吞吐 76%，occupancy 38%，stall 集中在 Long Scoreboard这一步已经基本确认是 memory bound。接下来做 roofline 确认：
ncu --kernel-id 0 --set roofline ./image_demo 4096 8 # 打开 .ncu-rep 看 roofline 图，kernel 的点靠近 memory 端的线条到这里，优化方向就清晰了：核心是访存优化，而不是盲目开大 loop 或用更多指令并行。后面把报告导出来慢慢拆：
ncu --kernel-id 0 --section MemoryWorkloadAnalysis --section WarpStateStatistics \ --export ./report_filter.ncu-rep ./image_demo 4096 8 ncu --gui report_filter.ncu-rep我会在 GUI 里重点看 Memory Workload 的 “Memory Access Pattern” 和 Warp State 的 stall breakdown，确认是否存在非合并访问、通过 shared memory 做数据重排的必要性，以及是否可以降低寄存器压力提高 occupancy。
4. 核心指标深度解读：从数字到优化动作
4.1 SOL 百分比是“方向”，不是“结论”
SOL 页面会给一行Compute (SM) [%]和Memory [%]，很多刚入门的同学看到 compute 60%、memory 70% 就觉得两头都还行，然后不晓得该优化哪里。其实这两个数字要结合 stall 来读。
比如Memory Workload Analysis里的Memory Throughput偏高、但Compute Workload Analysis偏低，同时Long Scoreboard又占了大部分 stall，这就是典型的“带宽接近用尽、但计算单元浪费了”的状态。而Long Scoreboard本身说明等待 L2/DRAM 数据返回的事件非常多。真正优化动作是，把访存模式从“每线程取几个离散元素”改成“合并取连续数组切片”，必要时用__ldg或__ldcs显式走 read-only 或 streaming cache。
反过来，如果Compute (SM)已经到 85% 以上、Long Scoreboard也没那么突出，但Occupancy偏低，说明调度隐藏已接近极限，再调 occupancy 提升空间有限——因为 SM 里的执行单元已经忙不过来了。这种局面要换个思路，从算法层面减少指令数量（常数折叠、查表等），或者做 kernel 融合减少反复启动和边界开销。
4.2 Warp Stall 分布是“归因”，不是“惩罚”
Warp 是指一组 32 个并行线程，GPU 执行时按 warp 调度。当一个 warp 因为某种原因不能继续发射指令，就发生了 stall。NCU 会统计每一个 warp stall 的原因，百分比最大的就是最需要关注的瓶颈。
常见 stall 原因和对策可以整理成一张速查表：
| Stall 原因 | 典型含义 | 优化动作 |
|---|---|---|
| Long Scoreboard | 等待全局/本地内存数据返回 | 合并访问、增加每线程访存个数、提升 ILP |
| Short Scoreboard | 等待共享内存或常量内存数据 | 检查 bank conflict、使用__constant__配合广播 |
| Wait | 固定延迟指令（如整数乘、浮点加）还未完成 | 调整指令顺序、展开循环提升指令级并行 |
| Math Pipe Throttle | 数学流水线（SFU 等）过载 | 减少超越函数（sin/cos/exp）调用，必要时查表 |
| MIO Throttle | 访存指令队列（MIO）堵塞 | 减少访存指令条数，把分散访问合并成大块传输 |
| LG Throttle | L1/L2 请求过载 | 降低访存密度，利用 shared memory 或纹理缓存 |
| Branch Resolving | 分支指令等待解析 | 减少 divergent branch，尽量用算术逻辑替代分支 |
| Barrier | 等待 __syncthreads() 或协作组同步 | 优化同步点、减少同步频率，考虑用__syncwarp替代 |
这里有实操技巧：不要只看 stall 占比，还要结合stall_cycles看绝对周期数。比如某个 warp 有 20% 的 stall 原因是 Barrier，但它的绝对 stall 周期数很小，那整体影响就不大；反之 Long Scoreboard 占比 40%、绝对周期数巨大，那才是动刀的位置。
4.3 Occupancy 并不是“越高越好”
很多教程把 Occupancy 说成“尽可能高”，实际经验要打折扣。高 Occupancy 让更多 warp 驻留，能掩盖访存延迟；但代价是每个 warp 可用寄存器变少，很容易发生 spill；同时在访存带宽已经打满的场景下，更高的 occupancy 不会带来额外吞吐，反而可能增加 cache thrash。
举个例子：某个 kernel 每个线程需要 40 个寄存器才能避免 spill，但 occupancy 拉满时每个线程只能分到 24 个寄存器。这时候超过一半的变量会被压到 local memory，反而造成额外的访存流量，性能和预期相反。所以Occupancy页面里，NCU 其实会标出Achieved Occupancy（运行时的实际占用率）和Theoretical Occupancy（静态分析上限），两者差距大就说明存在启动边界效应、尾部效应或同步阻塞。
真正合理的做法是：看sm__warps_active.avg.pct_of_peak_sustained_active这个指标，同时对照smsp__average_warp_latency_per_inst_issued.ratio。如果 latency 高但 occupancy 已经充裕，那隐藏延迟已经是够用的；如果 latency 高且 occupancy 被非调度原因卡住（比如 shared memory 用太多），那才值得通过__launch_bounds__或者cudaFuncAttributeMaxDynamicSharedMemorySize调整。
4.4 内存访问模式：Miss、Sector 与 Bank Conflict
这一小节属于 NCU 报告里信息量最大、但也是新手最容易忽略的。Memory Workload Analysis里通常有两块核心数据：
Cache hit rates / miss rates：L1 text cache、L2 命中率。如果 L1 hit rate 极低而访问模式又是连续数组，优先怀疑访问步长不对或者没有用
__ldg。Memory Access Pattern：显示 warp 内 32 个线程访问的地址分布。这个图最大价值是判断 global load/store 是否合并。合并的意思是同一 warp 的 32 个线程访问的地址落在一个或少数几个 32 字节 sector 内，最好就在同一 128 字节 cache line。
如果sector efficiency很低（比如低于 50%），说明发生了大量跨 sector 的非合并访问——即使没有 visible bank conflict，全局内存带宽也已经浪费了。常见原因是数组按“列”访问但没有转置，或者 struct-of-arrays（SoA）误写成了 array-of-structs（AoS）。调整数据结构后 sector efficiency 能从 30% 一下拉到 90% 以上，kernel 性能可以获得级数级提升。
同样地，shared memory 的bank conflict也是重点。每个 bank 每个周期只能服务一个地址，如果同一 warp 内多个线程访问同一个 bank 的不同地址，就会造成 N 路冲突，实际吞吐降为 1/N。NCU 会直接在 source view 里把 conflict 地址标红，你可以看是访问shared_data[threadIdx.x * 128 + i]这类模式还是出现了跨 bank 步长问题。经典解决法是把数组 padding 一下（比如从[32]扩展到[33]），错开 bank 编号。
4.5 指令统计：你看到的循环可能不是你以为的循环
指令层面的分析在Instruction Statistics里能看到：
Executed Instruction Mix：各类型指令（FMA、INT、LDG、STS 等）的比例。Issued Warp Instructions与Executed Warp Instructions的区别，用于判断指令发射效率。Replay：某些指令（比如 shared memory bank conflict）会被硬件重放，导致实际执行的指令数大于代码里写出的指令数。
举例：一段本来只有一条 STS 的代码，若遇到 4-way bank conflict，会被重放 4 次，对应Instruction Replay Overhead明显上升。这个指标能帮你定位“代码看着没问题为什么慢”的隐性问题。减少 replay 的方式又回到了 bank padding 和访问模式设计。
5. 基于分析结果的优化路径与验证闭环
5.1 先按瓶颈类型选择优化策略
读完报告后，不要直接就去改代码，先把瓶颈类型归类，按我下面这个决策路径走省力气：
Memory bound：优先看 Memory Workload Analysis 的 hit rate、sector efficiency、访问模式。优化动作优先级：合并访存 → 提升 L2 命中 → 使用只读缓存/纹理路径 → 必要情况下用 shared memory 做数据分块重排 → 进一步减少重复 global 访问。
Latency bound（stall 以 Long Scoreboard 为主）：提高每线程 ILP（立即数并行性）、增加占用率、调整循环展开因子，或者把独立指令穿插进访存指令之间。典型策略是“一次多取、分批算”，让多个独立访存在飞行中重叠。
Compute bound：减少冗余指令（常数折叠、查找表）、谨慎使用 SFU 指令（
sinf等）、提高 FMA 比例（FFMA 每条能算乘加而我们经常拆成两行）、检查精度需求后按需用--use_fast_math。Occupancy 受限：寄存器限制
__launch_bounds__、减少 shared memory 使用、调整 block size 让资源利用率落到整数区间。注意：如果 kernel 本身访存并发很低，提高 occupancy 往往没有收益。
把策略先写下来，再动手，后面每一步改动就可以用 NCU 复测验证。不要“边改边猜”。
5.2 一个具体案例：仿真的访存优化
我拿一个模拟的“矩阵按行更新” kernel 来演示优化过程。假设核心代码长这样（示意）：
__global__ void update_by_row(const float* input, float* output, int width) { int tx = threadIdx.x; int row = blockIdx.x; for (int i = 0; i < width; i += blockDim.x) { int col = tx + i; float val = input[row * width + col]; output[row * width + col] = val * 2.0f + 1.0f; } }实际运行中这个 kernel 的 NCU 报告显示Memory Throughput高了、Compute Throughput很低，Long Scoreboard占比巨大，sector efficiency 大概 60%。为什么会这样？
这里行方向访问本质上是连续的，问题出在 block 划分上：width和blockDim.x若不能整除，尾部 warp 会访问到不完整的 sector。sector efficiency 下降。改进思路有几种：判断width是否是 32 的整数倍，若不是则用更细粒度的 mask；或者把内层循环改成“每个线程连续访问多个元素”（strided 和 block-linear 混用），让每个 warp 完整覆盖连续 sector。
改成 block-linear 后 sector efficiency 提升到 95% 以上，Long Scoreboard明显下降，核整体耗时减少了约 40%。这个实例的核心启示：报告里的 sector efficiency 直接对应优化动作，而不是给你一堆形容词。
5.3 验证闭环：每次改动后跑同一份指标
优化不是改一版就完事，我养成良好的习惯是——在每次改动之后，固定跑同一份 NCU section 集合，记录 SOL、occupancy、top 3 stall 原因、sector efficiency 这四个核心指标的数值变化。用表格对比，比如：
| 改动点 | Compute | Memory | Occupancy | Top Stall | 耗时 |
|---|---|---|---|---|---|
| 原始版本 | 27% | 76% | 38% | Long Scoreboard | 145 us |
| 改成合并访问 | 34% | 64% | 42% | Long Scoreboard | 95 us |
| 加上launch_bounds | 41% | 55% | 55% | Math Pipe | 82 us |
| 使用 shared memory 缓冲 | 46% | 48% | 52% | Wait | 75 us |
看到没有，Memory数字一直在下降，但耗时也在降——因为瓶颈在访存，memory throughput 低反而说明浪费少了。这跟“越高越好”的直觉正好相反。所以读报告、验证优化，都要带着“瓶颈方向”这个锚点，不要追着某个百分比跑。
6. 实用技巧：加速分析、压榨报告信息量
6.1 用 section 组合有效控制 profile 耗时
默认 full section 的 profile 时长经常是正常执行的好几倍，尤其是在大 kernel 上，每多一个 section 就多一轮 counter 采集。对策有两个：
- 用
--set default+ 少量自定义 section 组合，限定分析范围。如果只是确认访存瓶颈，用--section MemoryWorkloadAnalysis就够。 - 用
--kernel-id只分析目标 kernel，必要时配合--launch-count 1、--launch-skip跳过前几次空跑（通常第一次运行存在缓存冷启动，跳过更稳）。
另外，ncu支持--target-processes all指定多进程程序的全部子进程，但会显著增加分析和 filter 复杂度，除非必要建议只针对主进程。
6.2 文本导出与脚本化分析
GUI 看报告体验好，但做自动化巡检、性能趋势对比时，CLI 的文本导出更方便。常用命令：
ncu --export full_report.ncu-rep ./app arg # 之后可以这样提取纯文本： ncu --import full_report.ncu-rep --page details > details.txt ncu --import full_report.ncu-rep --page summary > summary.txt除此之外还可以用--csv导出 CSV 格式，之后用脚本处理多轮测试的指标对比。这让“跑 N 组参数对比占用/命中率”这类工作可以被批量执行，而不是每次人工点上 GUI 点五次。
6.3 结合 Source View 定位到具体代码行
Nsight Compute 最“香”的功能之一就是 Source View。开启方式一般是 GUI 的 Details 页里选择 Source Counters 标签，或者在 CLI 用--set full并开启 source-level 信息。
做过几次优化后的经验是：当Long Scoreboard是最主要 stall 原因时，直接在 Source View 的 load/store 指令行上盯着看。NCU 会标出对应每一行代码的 stall 热度，有时你会惊奇发现，热点并不在“你以为的计算行”，而是一行不起眼的指针解引用或结构体拷贝。这时候可以针对性地手动内联、拆分访存或加#pragma unroll 4。
7. 实操中的常见问题与排查技巧
7.1 问题一：profile 结果显示所有指标全是 0 或根本不生成报告
最常见原因是权限或硬件计数器被占用。默认情况下需要足够权限来打开硬件计数器。排查步骤：确认用户是否在video组、尝试以sudo或者修改 NVreg 相关参数；另一个原因是多个 profile 工具同时跑，或驱动处于只读模式，可以重启后重试。
另一个原因是 kernel 在 profile 过程中启动失败、或者 CUDA error 没有被检查。代码里很多cudaGetLastError被吞掉了，profile 模式下这种 error 会被放大。先给关键步骤加cudaGetLastError检查，排除代码异常。
7.2 问题二：报告里的 Duration 和我用 event 计时差很多
这是正常现象，不是 bug。NCU 要完成指令采样、计数、读写，会拉长执行时间；另外它在 profile 时也可能关闭一些自动调优（如 launch bounds 动态调整、persistent L2 设置）。以 NCU 的 Duration 看相对变化即可，绝对数值参考独立计时器。
7.3 问题三：经常被问到的几个“差不多先生”问题
“我的 kernel Compute 利用率 95%，是不是就没优化空间了？” —— 不一定，要看 stall 分布。95% 的 compute throughput 但
Long Scoreboard依然很高，说明计算管线在等待数据时虽然整体利用率高，但还有可压榨的空间。可以做更高并行度的数据预取或者优化访存模式来让计算和访存重叠更充分。“Occupancy 从 40% 提到 80%，性能是不是就该翻倍？” —— 并不直接。如果没有足够访存并行度或指令级并行，提高 occupancy 只是让更多 warp 等待，对延迟掩盖效果有限。验证方式就是跑一遍
Scheduler Statistics看eligible warps per cycle是否远低于最大值。“NCU 说我的 bank conflict 是 8-way，为什么改动数组 padding 后就不 conflict 了？” —— 这是 padding 改变了地址与 bank 的映射关系，把同一 bank 的冲突请求分散到多个 bank 上。这个现象也说明，bank conflict 和数组 stride 有强耦合，报告里的 warning 其实给了你改数据布局的依据。
7.4 技巧：先减变量，再看指标，最后改结构
可能是我被问得最多的建议：不要同时改一大堆东西再跑 profile。变量的相互影响会让指标完全失去归因能力。正确做法是一次只动一个变量，然后看同一组指标。比如先只做访存合并，确认Long Scoreboard下降了；再加__launch_bounds__，看 occupancy 变化；最后再动 shared memory。这样每一步的指标变化都能对应到一个可解释的优化动作。
从个人的使用习惯来说，Nsight Compute 跟其他 NVIDIA 工具不同：它不是“看一下时间线”那种感知层面的工具，而是需要你带着问题去查的显微镜。它不会直接告诉你“你应该把某个循环展开 4 次”，但能告诉你“这一行 load 指令制造了多少 stall 周期”“这段代码 bank conflict 多严重”“哪些存储器请求根本没利用满带宽”。拿到这些数据后，优化就变成了一道道明确的选择题，而不是玄学。
如果你正准备拿手上第一个 CUDA kernel 开刀，我的建议是：先跑一轮ncu --set default，重点看Compute Throughput、Memory Throughput和Warp Stall三个数字，然后按照第 5 节的策略路径定方向，最后再用自定义 section 去做专项确认。等你跑过两三次完整流程，就会建立起“指标-硬件行为-优化动作”的对应直觉，到那时候 NCU 就不再是报告生成器，而是你优化过程里最顺手的那把卡尺。
