企业官网建设流程全解析

提到GPU profiling，可能很多人第一反应是打开任务管理器看一眼GPU占用率，看到100%就以为"已经跑满了"，看到20%就抱怨"显卡没发挥作用"。我做了几年GPU相关的性能分析，可以负责任地说：这两个判断大概率都是错的。GPU profiling从来不是看一个占用率数字那么简单，它要回答的问题是——你的程序在GPU上到底把时间花在了哪里、把资源浪费在了哪里，以及为什么明明硬件在忙，业务上却不出活儿。

这篇文章是我基于自己的实际项目经验整理的一份GPU profiling实战笔记。内容覆盖从底层执行模型、常用工具链，到真实瓶颈定位、驱动层异常排查，再到优化动作落地的完整链路。适合正在做CUDA/CUDA内核开发、PyTorch深度学习训练调优、GPU集群运维的工程师，也适合刚接触异构计算、想搞明白"GPU利用率为什么上不去"的初学者。

1. 先厘清概念：GPU profiling到底在"剖"什么

1.1 profiling不是monitoring，别把两个概念混着用

我经常看到有人把GPU profiling和GPU监控混为一谈。监控（monitoring）是持续采集GPU的状态——利用率、显存、温度、功耗、时钟频率，典型工具是nvidia-smi、nvidia-smi dmon、DCGM，它回答的是"现在GPU在干什么"；而profiling是对某一次执行过程做深入的采样和插桩分析——kernel耗时、启动开销、访存吞吐、warp调度情况、寄存器压力，典型工具是Nsight Systems、Nsight Compute、PyTorch Profiler，它回答的是"刚才那次运行，时间去哪了，为什么这么慢"。

这两个层次都很重要，但排障顺序应该是先做monitoring发现问题，再做profiling定位根因。我接手过一个训练脚本，nvidia-smi看GPU利用率只有18%，所有人第一反应是"GPU没喂饱，数据加载有瓶颈"。结果一profile发现，问题根本不在数据管线，而在一个自定义kernel的启动配置——block数量设置得太小，GPU的几十个SM只激活了零星几个，其余都在空转。这就是典型的"只看监控不看profile"会误判的场景。

1.2 profiling要关注的四个维度

我习惯把GPU profiling拆成四个维度来思考，每个维度背后对应一类不同的瓶颈：

很多时候瓶颈不是单一维度，而是多个维度叠加。比如一个kernel既访存密集又启动过于频繁，那单纯优化访存模式只能改善一部分，还得考虑把多个小kernel合并成一个。

1.3 什么时候该做GPU profiling

我总结了三类典型触发场景：

• 性能不达预期：模型训练速度比理论上慢很多，infra延迟高，GPU利用率异常，这时需要对训练循环和核心kernel做profile。
• 资源异常：显存OOM但排查不出谁占了显存，或者GPU温度/功耗异常升高，需要profile定位是哪个算子或哪个批次的显存峰值。
• 环境变更后的回归验证：换了驱动、换了CUDA版本、从单卡扩展到多卡、换了GPU型号之后，哪怕程序能跑，也要做一次基线profile，确认性能没有劣化。

2. 底层执行模型：看懂GPU怎么干活，才能看懂profile结果

2.1 从host到device：一次kernel launch的全流程

GPU profiling的所有指标，本质上都是在度量这条链路上每一环的开销。一个CUDA kernel从CPU发起执行，要经历这么几步：

1. 应用通过CUDA Runtime API提交kernel launch请求。
2. 驱动把kernel、参数、相关的CUDA context信息打包，写入GPU命令缓冲区。
3. 命令缓冲区通过PCIe/系统总线传给GPU前端。
4. GPU前端解析命令，把kernel调度到各个SM上。
5. SM上的调度器把线程块分派到执行单元，按warp逐条发射指令。
6. kernel执行结束，结果写回显存，CPU侧通过同步操作感知完成。

很多人忽略的一点是：kernel launch本身有一笔固定开销，通常是几微秒到十几微秒。当你的kernel只有几十微秒的执行时间时，launch开销就占了总耗时的大头。Nsight Systems的时间线里能清清楚楚看到CPU发指令和GPU执行之间的空隙，这就是很多"小kernel扎堆"程序慢的根本原因。

2.2 warp和cooperative thread array（CTA）到底什么关系

我在网上看到不少人在搜"cooperative thread array（CTA）和warp是什么关系"，这确实是理解GPU profiling绕不开的底层概念。直接给结论：

• thread：你写代码时的最小执行单元，每个线程跑同一段kernel函数。
• warp：硬件真正执行时的调度单位，由32个线程组成。GPU的SIMT架构决定了一个warp里的线程在同一时刻执行同一条指令（可以有分支分歧，但本质上是一条指令喂给32个lane）。
• CTA（cooperative thread array）：也就是常说的线程块（thread block），是一组可以互相协作的线程集合，它们可以在block内做同步（__syncthreads()）、通过共享内存交换数据。一个CTA通常包含多个warp，比如256线程的CTA就是8个warp。

用一句话概括：CTA是程序员视角的逻辑并发结构，warp是硬件视角的调度执行结构。CTA多大由你定，warp多大是NVIDIA固定死的（32线程）。

这个区分对profiling很重要。因为很多profile工具报告的是"每SM活跃warp数""warp占用率"这类指标，你要能心算出来：一个256线程的CTA占用8个warp，如果SM能容纳的并发warp上限是64，那8个warp只占了12.5%的"理论并发容量"——哪怕kernel执行时ALU一直在运算，这种低占用率的结构也会导致延迟无法被隐藏。

2.3 occupancy越高一定越好吗：算一笔账

所谓occupancy（占用率），指一个SM上活跃warp数与理论最大warp数的比值。很多人调到高occupancy就觉得优化到位了，这是个常见误区。

occupancy的真正作用是隐藏访存延迟。当一个warp因为读显存被阻塞时，SM调度器切换到另一个warp执行，让访存和计算重叠。低occupancy时，warp阻塞就是真的空闲；高occupancy时，访存延迟能被其他warp的计算掩盖。

但高occupancy本身不直接等于高性能。假设一个kernel是纯计算密集（所有数据都在寄存器里），没有访存等待，那计算单元始终满载，多高的occupancy都无所谓，反而是低occupancy还能减少调度开销。反过来，一个访存密集的kernel，如果occupancy太低，所有warp都在等数据，SM的ALU干瞪眼，这时提高occupancy就非常有效。

我做过一个例子：一个kernel每个线程用64个寄存器，导致一个SM上最多同时驻留1024个线程（20世纪NVIDIA的经典上限是每个SM 2048线程，寄存器文件大小固定），occupancy只有50%。我把寄存器用量压到40个之后，occupancy到了80%，虽然每线程的"局部计算能力理论峰值"略微下降，但访存延迟被更好地隐藏了，整体kernel时间反而降了35%。这就是为什么profile工具里要同时看"寄存器压力"和"occupancy"两个指标，不能只看一个。

2.4 内存层次：profile里那些带宽数字意味着什么

GPU的内存/缓存层次对profiling结果的影响，经常被低估。以NVIDIA A100为例，一个SM有：

• 寄存器文件（最大256KB）
• 一级缓存和共享内存（加起来192KB，可配置比例）
• L2缓存（40MB，所有SM共享）

不同级别的带宽差异巨大。我用一个类比帮助理解：寄存器相当于你的大脑工作记忆，取用几乎零延迟；共享内存相当于手边的笔记本，几十个周期能翻到；L1/L2缓存相当于书架上的资料，需要起身查；显存相当于仓库，每次取货要走一个很长的通道。

profile工具里常见的指标是"global memory throughput""L1/TEX hit rate""L2 hit rate"。当L2命中率高，说明数据被复用了；当global memory throughput接近峰值，说明这个kernel是访存受限的。如果你看到一个kernel算力利用率低、访存带宽利用率极高，那基本可以判定为memory-bound，优化方向是减少无效访存、提高数据复用，而不是堆更多计算。

3. 我的常用profiling工具栈与安装避坑

3.1 NVIDIA工具族谱：Nsight Systems和Nsight Compute的分工

NVIDIA官方的profiling工具迭代了好几代。早年的nvprof在CUDA 10之后基本废弃了，新项目直接用Nsight系列。很多人第一次接触这俩工具会混淆，我帮你理清分工：

• Nsight Systems（命令行是nsys）：做全流程时间线的分析。它能看到CPU上的API调用、CUDA kernel的启动和结束、内存拷贝、CUDA graphs、多流之间的重叠情况。回答的问题是"哪个阶段耗时最长""kernel有没有和拷贝重叠""CPU有没有拖GPU后腿"。这是你拿到一个程序后第一个该跑的工具。
• Nsight Compute（命令行是ncu）：针对单个kernel做深度分析。它会对kernel做多次重放（replay），测量SM内部的各类硬件计数器——计算吞吐、访存吞吐、warp状态、寄存器分配、branch divergence、shared memory bank conflict等。回答的问题是"这个kernel为什么这么慢""瓶颈在计算还是访存"。这是精确定位kernel内部瓶颈的第二站。

我建议的统一流程是：先用nsys看全局，找出最耗时的3~5个kernel，再用ncu去逐个分析这几个kernel内部到底发生了什么。直接拿ncu扫全程序会非常慢，因为每个kernel都要重放很多次。

常用的两个命令模板：

# 全局时间线分析，跟踪CUDA和NVML nsys profile --trace=cuda,nvtx,osrt --output=my_profile -o report python train.py # 对指定kernel做深度分析，输出到report.ncu-rep ncu --set full --kernel-name regex:kernel_name --launch-count 3 --export report ./my_app

3.2 PyTorch训练怎么profile：torch.profiler实操

做深度学习训练的GPU profiling，我一般不会直接上Nsight，先用PyTorch自带的torch.profiler快速定位热点算子。它的好处是能直接把算子名、CPU耗时、CUDA耗时、显存占用对应起来，免去从CUDA kernel名字反推PyTorch操作的心智负担。

一段最常用模板：

from torch.profiler import profile, ProfilerActivity with profile(activities=[ProfilerActivity.CPU, ProfilerActivity.CUDA]) as prof: for step in range(10): train_one_step() prof.step() # 每个step打一个标记 # 按CUDA时间排序，输出前20个算子 print(prof.key_averages().table(sort_by="cuda_time_total", row_limit=20))

看输出表的时候有个关键心法：先看cuda_time_total前几名是不是符合你的预期。如果你以为瓶颈是Attention，结果排第一的是DataLoader到GPU的传输，那问题在IO管线而不是模型。我排查过很多"训练慢"的案例，第一热点经常是aten::to或cudaMemcpy，这种拷贝瓶颈根本不该靠优化kernel解决，而是要在数据管线上动手（prefetch、持久化worker、直接加载到GPU端）。

torch.profiler还有一个很实用的with_stack=True参数，能把算子对应的python调用栈打出来，定位是哪一行代码触发了这个算子。在复杂的训练脚本里找"到底谁的显存/耗时暴涨"，这个参数能省很多时间。

3.3 集群和云GPU场景：DCGM + Prometheus持续画像

单机排障可以交互式跑nsys，但上了集群、云GPU，你需要的是持续采集的profiling/监控数据。NVIDIA官方的DCGM（Data Center GPU Manager）是目前最成熟的方案。

我推荐的最小可用组合是：

• dcgm服务采集所有GPU的利用率、显存、SM占用、PCIe吞吐、温度、功耗、uncorrected errors。
• dcgm-exporter把指标暴露成Prometheus格式。
• Prometheus + Grafana负责存储和可视化。

这套东西的价值在于，你可以拉出时间维度上的画像。比如训练任务每天凌晨某个时段GPU利用率骤降，配合时间线一看，发现是另一个定时任务抢占CPU导致数据加载变慢——这种跨组件的问题，单靠一次profile根本看不出来，必须靠持续监控数据。

很多用Kubernetes的朋友会在上面叠加GPU虚拟化或调度配额（比如热搜里有人提到k8s调用GPU、GPU资源分配、根组织云原生的GPU配额），这里我提醒一句：虚拟化/配额层的profiling和裸金属完全不同。如果GPU被多个任务分时共享，单个任务看到的nvidia-smi利用率是别人的，profile时间线也会被其他任务插队。这时候要先确认有没有MIG、时间片切分一类的机制，再看结果才有意义。

3.4 Intel GPU和AMD GPU的profiling工具

并非所有人都在NVIDIA生态里。热搜词里出现"Intel UHD Graphics""昇腾系列GPU"，说明现在异构设备的谱系已经很宽了。

Intel GPU带核显的机器做profiling，可以用Intel自家的VTune Profiler，它的GPU analysis能看EU（Execution Unit）利用率、访问带宽、shader占用。如果是跑oneAPI程序，Level Zero接口下还有ze_tracer这类工具。AMD方面对应的是ROCm生态，工具叫rocsolver/rocprof和OmniTrace，能对标Nsight Compute。国产的昇腾系列用的是华为自己的Ascend Profiling工具链，API风格也类似。

我的个人经验是：非NVIDIA平台的profiling工具成熟度普遍落后一截，遇到问题先看硬件厂商官方文档，别指望社区经验太丰富。日常调试可以先从官方提供的nvidia-smi替代命令开始（AMD有rocm-smi，Intel有xpu-smi），做第一层可用性判断，再决定要不要上重型工具。

4. 实测案例：从"GPU利用率20%"到定位真正的瓶颈

4.1 案例一：数据加载瓶颈，GPU利用率忽高忽低

这是一个非常典型的PyTorch训练场景。某个训练脚本跑起来，nvidia-smi里的GPU利用率像心电图一样在0%和95%之间来回跳。第一反应是"数据加载太慢"，但不要凭直觉下结论，我用Nsight Systems快速验证。

nsys时间线出来后的画面是：GPU上有kernel在跑，但每个kernel结束之后，GPU有一段长长的空白，CPU侧的cudaMemcpyAsync调用和DataLoader迭代的耗时占了绝大多数。这说明GPU不是没活干，而是喂数据的CPU管线把GPU饿着了。

解决办法按优先级排列：

1. 把DataLoader的num_workers调到4~8，让数据在后台线程预处理。
2. 开pin_memory=True，让page-locked内存的H2D拷贝走快路径。
3. 用prefetch_factor提前多取几个batch。
4. 极端情况下把数据预处理挪到GPU上（比如解码、归一化直接用CUDA tensor）。

改完之后再用nsys看same时间线，GPU利用率的锯齿明显消失了。这个案例的教训是：profile的value不在于告诉你"GPU利用率低"，而在于告诉你低的原因在哪一层。

4.2 案例二：小kernel扎堆，启动开销吃掉了执行时间

另一个我常见的"反直觉"现象：GPU利用率很高，但程序还是很慢。表面上看显存带宽也高、SM也忙，但总执行时间就是下不来。

用nsys看时间线，你会看到一串密密麻麻的kernel，每个只有几微秒。比如一个推理程序里，每个batch切成了几十个很小的elementwise算子。这些kernel本身计算量极小，但因为每个都走一遍"CPU提交->命令传输->GPU调度"的流程，launch overhead反而占了60%以上的耗时。

这种case的解法是kernel融合。把一连串elementwise操作合并成一个kernel，或者用PyTorch的torch.compile或TensorRT做自动化图优化，减少kernel数量。CUDA Graph是另一个利器（下面章节详述）。当你看到时间线上有大量"短命kernel"时，别纠结优化单个kernel的寄存器使用率，那是捡芝麻；先解决kernel数量和调度开销，那才是西瓜。

4.3 案例三：一个kernel的occupancy和访存冲突，用Nsight Compute定位

还有一种情况是：全局时间线没问题，数据加载没问题，kernl数量也不多，但整个程序就是慢。这种就要用ncu进kernel内部看。

我优化过一个自定义的矩阵归约kernel，现象是ncu报告里：

• Compute Workload Analysis显示SM吞吐只有峰值的63%。
• Memory Workload Analysis显示DRAM吞吐接近95%。
• Warp State Stats显示大量warp在long scoreboard状态（等显存数据）。

这三个数字合起来，答案非常明确：这是个memory-bound kernel，SM算力没满是因为在等数据，不是不会算。优化方向不是堆指令，而是减少每个线程读显存的次数。我做了两件事：一是让每个线程循环加载多个元素到寄存器（grid-stride loop），做一次访存多算几次；二是用shared memory做block内的数据复用，让同一block的线程尽量从shared memory拿邻居的数据。

改完之后的ncu报告，DRAM吞吐从95%降到60%，但SM吞吐反而升了，总耗时降了40%。这就是典型的"访存换计算"收益远高于"死磕计算指令"。

5. 驱动层异常排查：当profiling撞上硬件故障

做GPU profiling不可能只面对正常的程序。我遇到过好几次profile到一半，驱动崩了或者设备掉了，然后整个排查方向从性能分析转向硬件/驱动故障诊断。这一节把那些高频问题一次说清楚。

5.1 dmesg里的XID错误：以XID 79为例

热搜词里有XID 79，这是NVIDIA Linux驱动里的经典错误，完整信息一般是GPU has fallen off the bus。我第一次遇到时以为是PCIe插槽接触不良，后来发现原因五花八门。

排查XID 79的完整链路我整理为：

1. 先抓证据：执行dmesg -T | grep -i xid，找到完整的XID条目和它周围几行日志，看有没有伴随NVRM字样、温度告警、电源事件。
2. 确认供电和散热：用nvidia-smi -q -d POWER看当前功耗和最大功耗，用nvidia-smi -q -d TEMPERATURE看温度。XID 79在笔记本上经常出现在高负载+电池供电的场景，因为功耗墙和瞬时电流波动导致GPU瞬间掉链子。
3. 检查PCIe链路状态：用lspci -vvv看GPU对应PCIe设备的状态，确认link width和link speed是否正常（比如应该16x的变成8x或1x，说明链路在降速）。温度过高也会触发PCIe链路降速。
4. 尝试降频复现：用nvidia-smi -lgc 700锁定一个较低的GPU时钟，如果复现频率大幅降低，基本可以坐实是供电/散热不稳问题。
5. 驱动回滚对比：XID 79也跟驱动bug有关，虽然不常见。保留一个旧版本驱动做A/B测试，注意先nvidia-smi -pm 1或者卸载干净再装。

我印象最深的教训是：在移动工作站上，XID 79常常不是硬件问题，而是电源管理策略问题。Windows下混合显卡本子在电池模式切换到核显输出时，独显供电被切掉，Linux下如果没有正确配置nvidia-persistenced也会出现类似现象。最简单的验证方式：插上电源适配器、设置高性能模式，再跑一次同样的负载。

5.2 Windows下"错误代码43"的排查

Windows设备管理器里NVIDIA显卡报"错误代码43"，是另一种高频问题，热搜词里也出现了"英伟达gpu错误代码43"。这个错误意味着Windows已停止该设备，原因是设备报告了问题。对独显笔记本用户来说尤其常见。

我的排查顺序是：

1. 先排除驱动问题：到官网下载对应型号的最新驱动，用DDU（Display Driver Uninstaller）在安全模式下彻底卸载旧驱动，再全新安装。很多43是驱动残留冲突导致的。
2. 查供电和显卡切换：笔记本上如果BIOS里有独显直连/混合模式选项，切换到独显直连试试，排除核显切换导致的设备状态异常。
3. 检查是否被节能策略禁用：设备管理器里右键显卡，属性->电源管理，看看有没有"允许计算机关闭此设备以节约电源"被勾选，把钩去掉。
4. 查硬件接触：如果台式机，重新插拔显卡和供电线；如果是笔记本，尽量送修之前确认不是散热问题——过热后芯片虚焊也会出现43。

经验上有个规律：43如果在刚开机时就出现，多半是驱动或固件问题；如果用一会儿之后出现，往往伴随高热或供电不稳。结合任务管理器里的GPU温度趋势判断，方向会更准。

5.3 PyTorch装完却用不了GPU：验证链路

热搜里有"pytorch安装教程gpu"和"验证paddle gpu是否验证成功"，这类问题本质上是环境验证链路不完整。我每次配新环境，都会按这个顺序验证：

# 1. 驱动层：nvidia-smi能看到GPU和驱动版本 nvidia-smi # 2. CUDA运行时层：nvcc版本（注意它与驱动支持的CUDA版本可以不同） nvcc --version # 3. PyTorch层：能否识别GPU python -c "import torch; print(torch.cuda.is_available(), torch.cuda.device_count())"

常见坑有三个：

• pip把PyTorch装成了CPU版本：这是最大的坑。直接用pip install torch在某些源镜像上会是CPU版，torch.cuda.is_available()返回False。正确做法是去PyTorch官网选择对应的CUDA版本安装命令，比如pip install torch torchvision torchaudio --index-url https://download.pytorch.org/whl/cu121（版本号要匹配你的CUDA驱动兼容性）。判断方法：python -c "import torch; print(torch.version.cuda)"，如果不是类似11.8/12.1这样的数字，说明装错了。
• 驱动支持版本低于CUDA要求：比如新PyTorch要CUDA 12.x，但你的老卡驱动只支持CUDA 11.4。这时要么升级驱动，要么装旧版PyTorch。日志往往不会说清楚，torch.cuda.is_available()直接就False。
• WSL/容器里宿主机驱动没打通：WSL2里跑GPU，Windows侧要有对应的GPU驱动，容器里则要确保--gpus all或用NVIDIA Container Toolkit挂载了驱动。

排查的时候记住一条主线：从上往下逐层验证——驱动层能看到GPU吗？CUDA runtime能初始化吗？深度学习框架版本匹配吗？每一层的失败都往上找原因，很少需要直接怀疑硬件。

6. profiling只是开始：把分析结果变成可落地的优化动作

6.1 从profile指标到优化决策的映射表

工具跑完，报告看了，下一步自然是"动手改代码"。我总结了一份从profile观测到优化动作的快速映射表，每次拿到新的report就先对号入座：

这张表的背后逻辑很简单：profile给你的不是"答案"，是"线索"。每个指标对应一条优化路径，你在实际项目中只需要按图索骥。

6.2 几个我验证过的实用优化动作

针对上面映射表里的高价值动作，扩展一下具体落地方案。

H2D/D2H拷贝最小化。GPU profiling里最容易被忽视的开销就是CPU与GPU之间的数据拷贝。很多人习惯每个step把tensor.cpu()拿回来打印一下loss，多几次这种操作，拷贝就占了训练时间的1/3。优化思路是：批量同步、异步拷贝、让拷贝和非依赖计算在不同CUDA stream上重叠。nsys时间线里如果看到cudaMemcpy和kernel是串行的，那说明stream使用有优化空间。

CUDA Graphs是"小kernel扎堆"的终极解法。这个技术把一连串kernel launch在host侧录制成一个graph对象，然后一次性提交执行，免去每次launch的驱动开销。PyTorch里可以这样粗粒度地体验：

from torch.cuda.graphs import make_graphed_callables graphed_fn = make_graphed_callables(my_model, sample_args)

实测在大量小kernel的模型上，整体速度提升10%到30%很正常。但注意CUDA Graph在动态shape、有数据依赖分支的场景下会有约束，不适合无脑套用。

混合精度不是"降精度"，是"提速手段"。FP16/FP32混合精度训练能显著降低访存带宽压力，对memory-bound的模型尤其有效。但要小心梯度下溢和精度损失，PyTorch的torch.cuda.amp已经做了loss scaling，遇到不收敛再逐渐关掉。用ncu对比一下开启混合精度前后DRAM吞吐的变化，你会直观看到为什么它能提速。

kernel融合的正确姿势。融合不是把所有逻辑堆到一个kernel里就行。我见过有人把不相干的逻辑硬融进去，结果寄存器爆炸、occupancy暴跌，反而更慢。正确的判断标准是：要融合的操作之间有数据依赖（一个的输出是另一个的输入），且融合后每个线程的工作负载适中、不会导致资源超卖。用ncu比较融合前后kernel的SM吞吐和耗时，用数据说话。

6.3 优化闭环：改完一定要再profile一次

我犯过的最大错误是——改完代码觉得"肯定更快了"，然后就没再跑profile，直接上线。结果往往是被隐藏的问题在另一个维度爆发：数据加载快了，但GPU在等kernel同步；kernel快了，但显存碎片化导致OOM。

我的习惯是：每次优化改动之后，必须跑一次和基线完全相同的profile命令，对比关键指标。基线profile文件保留在项目目录里（比如baseline_nsys.sqlite、baseline_ncu.ncu-rep），改动后用同样的命令生成新报告，两者用Nsight的diff功能对比，或者自己盯core指标——总耗时、Top-5 kernel耗时、DRAM吞吐、occupancy。

这种"优化-复测"的循环，比任何一套优化口诀都重要。因为GPU程序的性能受硬件、驱动、数据shape、批大小影响太大，没有基线对比，你很难判断一次改动到底是变好了还是变坏了。

7. 写在最后：一点个人的实操体会

用了这么多年GPU profiling工具，我最大的体会是：工具只是镜子，照出问题的从来不是工具本身，而是你对程序执行模型的判断力。nvidia-smi看一眼就能判断"利用率高就是快"当然简单，但它会让你漏掉最核心的问题——GPU忙什么才是关键。

如果让我给刚开始接触GPU profiling的朋友一个建议，那就是：从最小的单kernel开始建立基线。写一个简单的矩阵乘法或者归约kernel，分别用nsys和ncu跑一遍，把每个指标的含义对应到自己写的代码上。这个过程花不了半天，但它建立的心智模型，会让之后遇到"GPU利用率低""XID 79""PyTorch识别不到GPU"这些问题时，第一反应是"查哪一层、看哪个工具、用什么逻辑排除"，而不是去抖音搜玄学教程。这条路我自己就是这么走过来的，确实值得你走一遍。