搞GPU开发这些年，被问得最多的一句话就是：“我的程序在GPU上跑得很慢，到底慢在哪？”CPU上有perf、vtune、gprof一堆工具可以追，一换到GPU，很多人就抓瞎了——显存占用看着正常，程序也不报错，可性能就是上不去。这时候真正缺的，就是一套系统性的GPU profiling方法。
GPU profiling，简单说就是把跑在GPU上的kernel、显存访问、线程调度、计算单元利用率这些环节全部量化出来，搞清楚瓶颈到底卡在哪一步。这篇文章不是教科书式的名词解释，而是我从实际项目里攒下来的一套流程：怎么选工具、怎么读指标、怎么从数据反推代码问题、怎么排查那些看起来跟profiling无关的GPU异常。适合三类人看：写CUDA/OpenCL kernel的异构计算开发者，做深度学习训练和推理部署的同学，以及搞GPU驱动、运行时栈的底层工程师。不论你在哪个层次，这套思路都能直接拿来用。
1. GPU Profiling 到底在测什么：先搞懂GPU的底层执行模型
1.1 kernel、grid、block、thread 与 warp 的真实关系
很多人一上来就打开profiler看数字，结果连数字代表什么都说不清。想读懂profiling报告，必须先理解GPU的执行模型。
CUDA编程模型里，host端启动一个kernel，这个kernel会以grid的形式发射到设备端。grid由若干block组成，block又由若干thread组成。thread是逻辑上的最小执行单元，但硬件真正调度的时候，根本不是一条一条thread来跑的。NVIDIA这边，32个连续编号的thread组成一个warp，这才是硬件调度的基本单位。一条指令发出去，整个warp的32个线程同时执行。AMD那边对应的概念叫wavefront，一般是64个线程一组。
这里就要说到热词里那个“cooperative thread array”（CTA）。在CUDA里，CTA其实指的就是线程块thread block，是一组能够在shared memory和同步原语上互相协作的线程。到了Cooperative Groups编程模型出来之后，CTA又有了更严格的含义：它指的是一组必须同时驻留在GPU上、可以跨block同步的线程块集合。如果这些block不能同时驻留，那同步操作就会直接死锁。所以warp和CTA是两层概念：warp是硬件层面的固定执行粒度，CTA是软件层面为了协作和同步而组织的逻辑分组。profiler里的occupancy、warp stall这些指标，全都建立在理解这两层概念的基础上。
1.2 profiling必须关注的关键指标
工具输出的名词一大堆，但核心指标就几个，我把它们按重要性排个序：
| 指标 | 含义 | 瓶颈指向 |
|---|---|---|
| Occupancy（占用率） | 活跃warp数占SM最大可容纳warp数的比例 | 寄存器过多、block太小、shared memory超限 |
| SM Active Warps | 每个SM上同时活跃的warp数量 | 并行度是否足够 |
| Memory Throughput | 实际访存带宽占峰值带宽的百分比 | 是否memory bound、访存是否合并 |
| L1/L2 Cache Hit Rate | 各级缓存命中率 | 访存模式好坏 |
| Warp Stall Reasons | warp卡住的原因分布 | long scoreboard、barrier、wait等 |
| Instruction Mix / Pipe Utilization | 各类指令占比和计算流水线利用率 | 指令选择是否合理、是否compute bound |
这些指标不是孤立的。比如occupancy高不代表性能好，如果因为访存不合并导致所有线程都在等数据，那SM上warp再多也是白等。反过来，如果ALU流水线利用率已经到90%以上，那考虑优化访存也意义不大，这就是典型的compute bound，该考虑算法层级的修改了。
1.3 什么信号出现时，应该启动profiling
根据我的经验，下面这些情况出现任何一个，都值得做一轮完整profiling：
- kernel执行时间在总耗时里占比很高，但加速比远低于理论值。
- GPU利用率看着不低，但训练或推理吞吐量死活上不去。
- 相同代码在不同GPU上表现差异巨大，想搞清楚原因。
- 程序偶发卡顿、显存报错、驱动重置，需要确认是否因为资源耗尽或访存越界。
还有一种情况，就是新接手别人的代码，什么都不懂，先跑一遍profiler建立基线数据。这一步特别重要，后面优化有没有效果，全靠基线的对照。
2. 工具选型解析：不同阶段用不同武器
2.1 主流GPU profiling工具对比
很多人一提到GPU profiling就只想到Nsight Compute，这是个大误区。工具选型取决于你想回答什么问题。我常用的工具分成几个层次：
| 工具 | 定位 | 适用场景 | 典型输出 |
|---|---|---|---|
| nvidia-smi | 硬件状态监控 | 快速健康检查：温度、显存、功耗、ECC | 命令行表格 |
| Nsight Systems | 系统级时间线分析 | CPU-GPU交互、kernel启动、内存拷贝、API开销 | 时间线视图 |
| Nsight Compute | kernel级微架构分析 | 单kernel内部的占用率、访存、指令、stall原因 | SOL图、Stall分析 |
| NVPROF | 旧版命令行profiler | 脚本化批量采集、老环境兼容 | 文本报告 |
| rocprof | AMD平台profiling | MI系列GPU上的kernel分析 | 文本/表格 |
| VTune | Intel平台分析 | Intel集成显卡和Arc显卡 | 时间线 |
2.2 我的实际选型逻辑
先说结论：先系统级，后kernel级；先健康检查，后微架构分析。
我见过太多人拿到一个性能问题，直接打开Nsight Compute抓单个kernel，折腾半天发现瓶颈根本不在kernel内部——可能是cudaMemcpy阻塞了整个pipeline，也可能是kernel启动太频繁导致launch overhead过大。Nsight Compute再精细也回答不了这种全局问题。
正确的打开方式是这样：先跑一遍nvidia-smi确认GPU没有硬件异常，包括温度、功耗、显存占用、ECC错误计数。没有问题就用Nsight Systems抓全局时间线，看看CPU和GPU的流水。kernel时间占比高说明GPU侧确实是热点，占比低就去查API调用、内存拷贝、CPU侧逻辑。确认热点kernel之后，再用Nsight Compute做单kernel深挖。
顺带提一句，chrome://gpu这种页面也能算入门级profiling，它能看到浏览器是否启用了GPU加速、用了哪块GPU、支持的加速特性有哪些。Windows下遇到“GPU not support acceleration”这种问题，第一反应就应该是打开这个页面看状态。
3. 实操全过程：定位一个真实kernel的性能瓶颈
3.1 环境准备与GPU状态确认
profiling开始之前，先把环境确认一遍，这一步能省掉后面太多排查时间。我用一套固定的前置检查命令：
nvidia-smi nvcc --version nvidia-smi -q -d ECCnvidia-smi看驱动版本、GPU型号、显存、当前利用率；nvcc确认CUDA工具链版本；ECC查询重点关注有没有显存错误。如果ECC错误计数一直在涨，那后面profiling数据可能都是脏的，因为硬件在反复重试和纠错。
需要注意，nsys和ncu的版本最好跟CUDA主版本匹配。比如CUDA 12.x配Nsight Compute 2024.x，不要让工具版本和驱动版本差距太大，否则采样经常失败，报一堆看不懂的错。
3.2 全局时间线采集：Nsight Systems先行
前置检查没问题，先做全局采集：
nsys profile --stats=true -o app_profile ./my_app这条命令会跑一遍程序，输出所有API调用、kernel启动、内存拷贝的时间线。跑完看几个关键点：
- kernel总耗时占比：如果kernel只占20%，剩下80%在cudaMemcpy，那改kernel算法可能收益不大，优先考虑用pinned memory或者异步拷贝。
- kernel启动数量：如果每秒启动成千上万个kernel，每个kernel只干一点点活，那launch overhead就是瓶颈，合并kernel或者用CUDA Graphs能大幅改善。
- CPU和GPU之间的间隔：如果GPU经常空闲等CPU提交任务，说明host端逻辑卡住了。
这一步的目的是圈定问题范围，然后才轮到微观分析。
3.3 微观分析：Nsight Compute深挖热点kernel
全局定位到热点kernel之后，跑单kernel采集：
ncu --set full -o kernel_profile -k my_kernel_name ./my_app-k参数指定要分析的kernel名字，避免采集全部kernel导致时间太长。跑完之后用Nsight Compute的界面打开kernel_profile.ncu-rep文件。
我一般按下面这个顺序读报告：
第一眼看Speed of Light（SOL）图，它会给出两个关键百分比：Compute (SM) Throughput和Memory Throughput。哪个接近100%，就说明瓶颈在哪一侧。如果Memory Throughput 95%，那就别纠结指令优化了，专心解决访存问题。
第二眼看Occupancy。这里要对比Theoretical Occupancy和Achieved Occupancy。如果理论值不高，看右侧的占用率限制因素：Registers Per Thread、Shared Memory Per Block、Block Size。举例来说，每个线程用了64个寄存器，导致一个SM只能驻留一半的block，这种情况可以通过__launch_bounds__限制寄存器数量来提升占用率。
第三眼看Warp State里的Stall Reasons。这个指标直接告诉你warp在等什么：
- Long Scoreboard：等全局内存数据返回，说明访存延迟没被隐藏，优先看访存是否合并。
- Barrier：等block内同步，说明线程负载不均或者同步太频繁。
- Wait：等固定延迟的指令完成，比如整数除法这类慢指令。
- Not Selected：warp可以执行但没被调度器选中，说明并行度太高而执行资源不够，不是大问题。
还有一个小技巧，Nsight Compute的Source view可以关联到具体代码行。看到某个循环对应的汇编指令吞吐异常，基本就能锁定问题代码。
3.4 从profiling结果反推代码优化：一个访存合并实例
拿我之前调过的一个矩阵转置kernel举例。原始代码长这样：
__global__ void transpose_naive(const float* src, float* dst, int width, int height) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (x < width && y < height) { dst[x * height + y] = src[y * width + x]; } }Nsight Compute的报告显示Memory Throughput接近90%，而Compute Throughput只有20%，典型的memory bound。再看Stall Reasons，Long Scoreboard占比超过60%。这就说明问题几乎可以确定是访存不合并：相邻线程访问src矩阵时，列方向跨度是width个float，每个线程读的地址差得很远，一个warp的32次访问完全无法合并成少数几次cache line传输。
解决方案是分块转置，用shared memory做数据交换：
#define TILE_SIZE 32 __global__ void transpose_tiled(const float* src, float* dst, int width, int height) { __shared__ float tile[TILE_SIZE][TILE_SIZE + 1]; int x = blockIdx.x * TILE_SIZE + threadIdx.x; int y = blockIdx.y * TILE_SIZE + threadIdx.y; if (x < width && y < height) { tile[threadIdx.y][threadIdx.x] = src[y * width + x]; } __syncthreads(); x = blockIdx.y * TILE_SIZE + threadIdx.x; y = blockIdx.x * TILE_SIZE + threadIdx.y; if (x < height && y < width) { dst[y * height + x] = tile[threadIdx.x][threadIdx.y]; } }注意tile声明成[TILE_SIZE][TILE_SIZE + 1]，加的那个1是为了避免bank conflict。修改之后再跑一次ncu，Long Scoreboard占比从60%掉到15%左右，Memory Throughput虽然还是高，但整体kernel时间缩短了差不多4倍。这个过程就是profiling的标准循环：定位瓶颈、分析原因、修改代码、重新采集、对比数据。
3.5 时钟锁定：保证测量结果可复现
还有一个值得单独说的点：profiling时GPU频率会动态变化，同样的代码跑两次，结果可能差20%。如果要做严谨的对比实验，建议锁频：
sudo nvidia-smi -lgc 1500 # 用完解锁 sudo nvidia-smi -rgc锁频能消除DVFS带来的波动，但锁到过高频率有风险，笔记本GPU散热跟不上还会触发降频甚至崩溃。我一般锁到该GPU boost clock的80%左右，既稳定又安全。桌面端GPU锁频相对随意，笔记本上操作要格外小心。
4. 常见GPU异常与Profiling现场排查实录
4.1 底层硬件错误大合集：代码43、Xid 79、GPU Crash Dump
这些错误我在不同机器上都遇到过，每一个都会让profiling无法进行，必须先处理。
Windows下最常见的“NVIDIA GeForce RTX 4060 Laptop GPU”设备管理器报代码43，系统提示“Windows 已停止此设备，因为其报告了问题”。看了一眼事件查看器，通常伴随几个显示驱动相关的警告。排查顺序是：先更新或回滚驱动排除驱动问题；再看是否近期超频导致显存不稳；最后检查是不是笔记本双显卡切换导致独显被禁用。如果是在虚拟机里透传GPU，代码43几乎是常态，需要确认宿主机和客户机驱动都匹配。
Xid 79: GPU has fallen off the bus，这类错误字面意思是GPU从PCIe总线上掉线了。可能原因包括：供电不足、PCIe插槽接触不良、显卡被物理拔出、驱动崩溃后重置失败。遇到这个，先看dmesg里有没有反复出现，如果频繁出现，建议更换PCIe插槽或电源。笔记本用户要留意是不是用了低功耗电源适配器，满载时供电跟不上很容易复现。
GPU crash dump triggered这种提示，在Linux上通常伴随Xid错误一起出现。Crash dump本身是驱动在异常发生后做的现场保存机制，生成的文件大小动辄几百MB。排查重点是确认崩溃前有没有跑什么重负载程序，以及ECC错误计数。
nvidia-smi -q -d ECC | grep -A 5 "ECC"如果ECC错误持续增加，基本可以判断显存有硬件隐患，继续profiling数据已经没有参考价值。
4.2 上层软件问题：GPU加速不可用、ComfyUI单GPU模式
有些问题和硬件无关，纯粹是软件配置。比如Chrome底部提示“GPU not support acceleration”，chrome://gpu页面里的WebGL和硬件加速全红。检查思路：驱动版本太老；显卡被驱动设置禁用了硬件加速；在远程桌面会话里运行，GPU被虚拟适配器接管。前者更新驱动即可，后者要把Chrome的硬件加速关掉或改用物理会话。
ComfyUI在Windows上报“On Windows we are currently forcing single GPU mode”，这是因为ComfyUI为了规避NVIDIA驱动在多GPU环境下的一些已知问题，强制单GPU运行。如果你确实有多卡需求，需要手动设置CUDA_VISIBLE_DEVICES，并确认两张卡都支持当前计算能力要求，有些新卡比如RTX 5070 Laptop的sm_120在旧版CUDA下会直接不兼容，这时候要用支持该架构的新版CUDA才行。
4.3 多卡、容器与GPU虚拟化场景下的Profiling限制
在k8s里调用GPU时，节点上要装好NVIDIA device plugin，Pod里申请nvidia.com/gpu资源，容器内才看得到GPU。但这里有个常见坑：容器里跑nvidia-smi有时能显示，Nsight Compute却采不到完整指标。原因在于容器隔离了硬件计数器，或者device plugin没有把必要的设备节点都挂载进去。解决方案是给容器加privileged权限或挂载/proc/driver/nvidia和/dev/nvidia-caps路径，具体看运行时的要求。
HAMI这类GPU虚拟化方案会把物理GPU切分成多个vGPU，profiling时尤其要留意：你看到的SM占用率可能只是你那个虚拟实例的视图，不代表物理卡真实状态。多个租户同时跑，硬件计数器互相干扰，采集出来的数据稳定性很差。云平台上的GPU配额预冻结机制也是类似逻辑，配额不够会自动冻结队列，这时候不是代码问题，是资源管理问题，找平台申请配额扩容或者错峰跑才是正解。
4.4 Profiling工具本身的坑：采样开销与权限问题
Nsight Compute为了方便分析，会在GPU上做指令重放和数据采集，这会让kernel执行时间翻好几倍。所以绝对不要用profiling模式下的耗时跟正常模式的耗时对比，这是新手最容易犯的错误。正常性能数据应该在不加profiler的普通模式下测，profiler只用来取指标。
另一个坑是权限。Windows上跑ncu经常遇到“Failed to initialize NVML”或“Permission denied”，在终端里没有用管理员身份运行。Linux上如果当前用户不在video组，也要sudo。否则采集到的数据不完整，有些计数器直接读不到。
还有一点，Profiling多卡程序时，如果多个进程同时往同一块GPU提交任务，计数器会串。我遇到过采集数据波动特别大的情况，最后发现是同一个节点上另一个师兄在跑训练任务。后来养成习惯：做profiling之前，先nvidia-smi确认GPU空闲，不然不如不跑。
5. 两年GPU profiling实战后的一些个人体会
最后分享一点自己的经验，不写总结，就说几个我在实际项目中踩出来的原则。
第一，profiling一定是个循环过程，不是跑一次就完事。我的习惯是先建立基线数据，然后一次只改一个变量，改完重新采集，再跟基线对比。不要同时调block大小、寄存器数、访存方式，否则数据出来根本分不清是哪个改动起的效果。每次profiling生成的文件保留下来，标注好日期和改动内容，这个习惯救过我很多次。
第二，动手优化之前，先用SOL图确认是compute bound还是memory bound。很多人一上来就增加block数、加大并行度，结果因为每个线程寄存器占用过高，occupancy反而掉下来，性能一点没提升。方向错了，越努力越糟糕。
第三，做GPU驱动或运行时层面的人，不能只依赖应用层的profiler。驱动开发场景下，kernel崩溃、设备丢失、中断风暴这些问题，光看Nsight数据是不够的，要结合dmesg、Xid错误、firmware日志一起看。profiling工具是起点，不是终点。
还有一个小建议：把profiling做成常态化工作。每次提交代码前跑一次基线，合入之后跑一次对比。看起来多花了一点时间，但那些“明明没改什么性能突然掉了一半”的诡异问题，基本都能靠这个机制第一时间发现。我自己搭了一个简单的脚本，一键跑nsys和ncu，输出格式化报告。你觉得有必要也可以这样搭一套，成本不高，收益长期看非常明显。
