Windows 下 CUDA 性能分析实战从 nvcc 编译排障到 Nsight Systems、Nsight Compute 定位 Kernel本文面向刚开始分析 CUDA 程序性能的开发者记录一套在 Windows PowerShell 下可复用的工作流先解决nvcc cl.exe NVTX 3编译问题再用 Nsight Systems 找到耗时 Kernel最后用 Nsight Compute 分析 Occupancy、缓存、全局内存和源码热点。代码来源与文章说明本文使用的 CUDA Profiling 示例代码来自 GitHub 开源课程仓库Infatoshi/cuda-course具体章节为作者 / 仓库维护者GitHub 用户Infatoshi课程仓库Infatoshi/cuda-course示例代码章节05_Writing_your_First_Kernels/03 Profiling原章节提供了00 nvtx_matmul.cu、01_naive_matmul.cu和02_tiled_matmul.cu等课程示例。本文是在 Windows、RTX 2080 Ti 和 CUDA 13.3 环境中对这些示例进行编译、运行和性能分析后整理的学习记录重点补充Windows PowerShell 下nvcc找不到cl.exe的排查过程NVTX 3 与旧版nvToolsExt.lib链接参数的差异Nsight Systems、Nsight Compute 的实际采集步骤本机实测指标的解释、结论边界和常见误区。课程代码版权归原作者及相应贡献者所有。本文中的环境排障、性能数据和分析结论来自个人学习实验如需使用或修改课程代码请以原仓库说明和实际许可状态为准。一、为什么需要两个性能分析工具分析 CUDA 程序时一开始最容易混淆的问题是Nsight Systems 和 Nsight Compute 到底有什么区别最实用的判断方式是确认程序能够正确运行 ↓ Nsight Systems程序的时间花在哪里哪个 Kernel 值得分析 ↓ Nsight Compute这个 Kernel 为什么快或慢 ↓ 提出优化假设修改代码再重新测量工具观察层级主要回答的问题Nsight Systemsnsys系统级、进程级时间线CPU API、GPU Kernel、内存复制如何排列哪个阶段最耗时Nsight Computencu单个 CUDA KernelKernel 的吞吐率、Occupancy、缓存、DRAM、指令和源码热点如何如果教学程序中只有一个已知 Kernel可以直接使用 Compute面对包含大量 Kernel、数据复制和 CPU 调度的真实项目通常应该先看 Systems。二、实验环境与示例本次实测环境如下项目环境操作系统与终端Windows PowerShellGPUNVIDIA GeForce RTX 2080 Ti显存22528 MiBCompute Capability7.5NVIDIA 驱动610.62CUDA Toolkit /nvcc13.3 / V13.3.73Nsight Systems2026.1.3.425Nsight Compute2026.2.1.0Visual StudioVisual Studio 2022示例包含三个 CUDA 程序程序用途00 nvtx_matmul.cu用 NVTX 标记内存分配、H2D、Kernel、D2H、释放等阶段01_naive_matmul.cu每个线程计算一个输出元素的朴素矩阵乘02_tiled_matmul.cu使用两个16×16共享内存 tile 的分块矩阵乘矩阵规模为1024×1024Kernel 配置为Grid (64, 64, 1) 4096 blocks Block (16, 16, 1) 256 threads/block 总线程数 4096 × 256 1,048,576 输出元素数 1024 × 1024 1,048,576在这个规模下恰好由一个线程计算一个输出元素。三、第一步解决 Windows 下的 CUDA 编译问题3.1nvcc fatal: Cannot find compiler cl.exe in PATH在普通 PowerShell 中直接编译nvcc-O3-lineinfo-o 00_nvtx_matmul.exe00 nvtx_matmul.cu-lnvToolsExt出现错误nvcc fatal : Cannot find compiler cl.exe in PATH这个错误不代表 CUDA Toolkit 没装好。Windows 上的nvcc需要 MSVC 作为 Host Compiler而普通 PowerShell 通常没有加载 Visual Studio C 开发环境。可以先使用vswhere.exe查找 Visual Studio$vswhere${env:ProgramFiles(x86)}\Microsoft Visual Studio\Installer\vswhere.exe$vswhere-latest-products*-requires Microsoft.VisualStudio.Component.VC.Tools.x86.x64-property installationPath查到安装位置后在当前 PowerShell 中加载开发环境Set-ExecutionPolicy-ScopeProcess-ExecutionPolicy Bypass Visual Studio安装目录\Common7\Tools\Launch-VsDevShell.ps1-Arch amd64-HostArch amd64这里使用-ScopeProcess只会临时修改当前 PowerShell 进程的执行策略关闭窗口后自动失效不会修改系统级策略。如果重新打开普通 PowerShell需要再次加载 Visual Studio 开发环境。3.2 NVTX 3 不需要链接旧的nvToolsExt.lib加载 Visual Studio 环境后继续使用-lnvToolsExt链接器报错LINK : fatal error LNK1181: 无法打开输入文件“nvToolsExt.lib”本例源码使用#includenvtx3/nvToolsExt.h在当前 NVTX 3 用法下功能由头文件提供不需要再链接旧的nvToolsExt.lib。删除-lnvToolsExt之后编译成功。3.3 最终使用的编译命令nvcc-O3-lineinfo-archsm_75-Xcompiler/utf-8-o 00_nvtx_matmul.exe00 nvtx_matmul.cunvcc-O3-lineinfo-archsm_75-Xcompiler/utf-8-o 01_naive_matmul.exe01_naive_matmul.cunvcc-O3-lineinfo-archsm_75-Xcompiler/utf-8-o 02_tiled_matmul.exe02_tiled_matmul.cu参数作用使用原因-O3开启编译优化更接近实际性能分析场景-lineinfo保留源码行号关联信息让 Compute 的 Source 页面关联 CUDA C 源码-archsm_75生成 Turingsm_75代码与 RTX 2080 Ti 的 CC 7.5 对应-Xcompiler/utf-8向 MSVC 传递 UTF-8 选项避免代码页相关警告-o指定输出文件名便于区分程序和报告版本编译完成后不要只看有没有生成.exe还应检查退出码并实际运行$LASTEXITCODEGet-Item.\01_naive_matmul.exe.\01_naive_matmul.exe$LASTEXITCODE如果示例本身没有printf运行后直接返回提示符且退出码为0属于正常现象。四、第二步用 Nsight Systems 查看全局时间线4.1 采集 CUDA 与 NVTX 数据nsys profile--tracecuda,nvtx--statstrue--force-overwritetrue-o nsys_00_nvtx.\00_nvtx_matmul.exe参数含义参数作用profile启动一次性能采集--tracecuda,nvtx只跟踪 CUDA 和 NVTX减少无关内容--statstrue采集后直接输出统计摘要--force-overwritetrue允许覆盖同名报告-o nsys_00_nvtx指定报告名前缀生成nsys_00_nvtx.nsys-rep nsys_00_nvtx.sqlite4.2 NVTX 能解决什么问题NVTX 可以把程序中的不同业务阶段标记出来。例如Matrix Multiplication ├── Memory Allocation ├── Memory Copy H2D ├── Kernel Execution ├── Memory Copy D2H └── Memory Deallocation本次单次采集数据NVTX Range时间Matrix Multiplication706.715755 msMemory Allocation696.761166 msMemory Copy H2D4.138120 msKernel Execution3.237010 msMemory Copy D2H1.632196 msMemory Deallocation0.914174 ms需要注意Matrix Multiplication是包含其他子范围的外层范围。NVTX Summary 对嵌套范围可能重复统计不能把各行百分比简单相加。4.3 第一次cudaMalloc很慢不代表每次都很慢三次cudaMalloc的统计如下指标时间Average85.645864 msMedian0.128668 msMin0.122189 msMax256.686735 ms最大值远高于另外两次说明第一次 CUDA 调用还包含上下文、驱动和运行时初始化成本。因此不能得出每次 cudaMalloc 都需要约 85.6 ms更准确的说法是首次调用显著拉高了平均值。进行稳定基准测试前应先预热 CUDA 上下文和目标 Kernel。4.4 不要混淆三种“执行时间”本次报告中同时出现完整 Matrix Multiplication NVTX 范围706.716 ms CPU 侧 Kernel Execution NVTX 范围 3.237 ms GPU 上 matrixMulKernel 实际执行 2.386 ms它们分别代表完整业务阶段包括分配、复制、计算和释放CPU 发起 Kernel、同步等待和少量调用开销Kernel 真正在 GPU 上执行的时间。在 Systems 时间线上经常能看到CPUcudaLaunchKernel → cudaDeviceSynchronize 等待 ↕ GPU matrixMulKernel 执行CPU 轨道上的cudaDeviceSynchronize很长表示 CPU 正在等待此前提交的 GPU 工作并不是 GPU 在执行一个名为cudaDeviceSynchronize的 Kernel。4.5 读取已有报告不重新运行程序也可以从已有.nsys-rep中读取 Kernel 汇总nsys stats--report cuda_gpu_kern_sum.\nsys_01_naive.nsys-rep nsys stats--report cuda_gpu_kern_sum.\nsys_02_tiled.nsys-rep如果出现NOTICE: Existing SQLite export found表示nsys正在复用已有 SQLite 导出不是错误。如果.nsys-rep已经更新可增加--force-exporttrue强制刷新 SQLite。五、第三步用 Nsight Compute 深入分析 Kernel5.1 先采集 Basic 报告ncu--setbasic--launch-count 1--force-overwrite--export ncu_01_naive.\01_naive_matmul.exe ncu--setbasic--launch-count 1--force-overwrite--export ncu_02_tiled.\02_tiled_matmul.exe打开 GUIncu-ui.\ncu_01_naive.ncu-repBasic 报告建议先看GPU Speed Of Light ThroughputLaunch StatisticsOccupancyDuration朴素矩阵乘的单次 Basic 数据指标实测值Duration2.09 msCompute (SM) Throughput61.70%Memory Throughput61.70%L1/TEX Cache Throughput92.56%L2 Cache Throughput13.06%DRAM Throughput1.21%Registers/Thread52Theoretical Occupancy100%Achieved Occupancy97.71%这里有三个非常容易误判的指标误区一Memory Throughput等于显存带宽不是。本例中Memory Throughput为61.70%但真正对应显存活动的DRAM Throughput只有1.21%。误区二Occupancy 100% 说明 Kernel 已经最优不是。Occupancy 只说明 SM 上活跃 warp 的数量不直接代表指令效率、访存效率或算法效率。误区三某个百分比变高就一定更好百分比必须结合绝对耗时、请求数量和具体硬件单元理解。单独看一个比率很容易得出错误结论。5.2 使用 Baseline 比较两个版本在 Nsight Compute GUI 中将朴素版本加入Baselines再打开分块版本可以直接看到绝对值和相对变化。本次 Basic 单次对比指标朴素版分块版相对变化Duration2.09 ms1.32 ms-36.81%Compute Throughput61.70%73.20%18.64%DRAM Throughput1.21%2.84%134.29%Registers/Thread5239-25.00%Static Shared Memory/Block02.05 KB从 0 增加Achieved Occupancy97.71%97.78%0.07%两个版本的 Occupancy 基本相同因此不能把加速归因于 Occupancy 提升。DRAM Throughput相对增加超过 100%也不表示传输了更多数据。它受到运行时间和活动集中程度影响必须继续查看绝对字节数和请求数量。5.3 Detailed 采集为什么很慢ncu--setdetailed--launch-count 1--force-overwrite--export ncu_01_naive_detailed.\01_naive_matmul.exe ncu--setdetailed--launch-count 1--force-overwrite--export ncu_02_tiled_detailed.\02_tiled_matmul.exeDetailed 模式需要读取更多硬件计数器。GPU 无法在一次运行中同时提供所有指标因此工具可能重放 Kernel 多次。所以即使 Kernel 本身只执行12 ms整条ncu命令运行很久也属于正常现象不能把工具采集总时间当成 Kernel 耗时。5.4 Raw 页面用于查精确计数器当已经知道需要什么指标时可以在 Raw 页面搜索sass__inst_executed_global_loads分块版本显示4,194,304 (-93.75%)Raw 对比通常显示当前值并在括号中给出相对 Baseline 的变化。反推得到朴素版 Global Load 指令67,108,864 分块版 Global Load 指令 4,194,304 减少 93.75%即降至 1/16其他计数器进一步说明了数据从全局内存复用转移到了共享内存指标朴素版分块版sass__inst_executed_global_loads67,108,8644,194,304sass__inst_executed_shared_loads041,943,040dram__bytes_read.sum8.98 MB9.93 MBdram__bytes_write.sum5.00 MB5.09 MB如果 Baseline 为0相对变化可能显示为inf%。这表示从零增长无法计算有限百分比不是报告损坏。5.5 Source 页面把热点定位到源码编译时保留-lineinfo后可在 Source 页面使用Left ViewSource Right ViewSASS Navigate ByInstructions Executed分块 Kernel 的核心代码为sharedA[ty][tx]A[row*Ktile*TILE_SIZEtx];sharedB[ty][tx]B[(tile*TILE_SIZEty)*Ncol];__syncthreads();sumsharedA[ty][k]*sharedB[k][tx];__syncthreads();C[row*Ncol]sum;乘加行对应约58.51%的Instructions Executed因为它位于 tile 循环和 tile 内乘加循环中。但需要注意Instructions Executed表示源码行关联的已执行指令占比不是精确的时间占比。六、为什么共享内存分块能减少 16 倍全局加载朴素版本每个 block 的源码级全局加载次数256 threads × 1024 loop iterations × 2 loads(A/B) 524,288 global loads/block分块版本有 64 个 tile每个 tile 由线程合作加载 A、B 各 256 个元素64 tiles × (256 A loads 256 B loads) 32,768 global loads/block两者相除524,288 / 32,768 16这与 Nsight Compute Raw 页面观察到的全局加载指令减少 16 倍一致。两个16×16FP32 tile 占用的共享内存为sharedA 16 × 16 × 4 bytes 1024 bytes sharedB 16 × 16 × 4 bytes 1024 bytes 合计 2048 bytes ≈ 2.05 KB/block分块版用少量全局加载将数据放入共享内存再由同一 block 内的线程反复复用。代价是增加共享内存访问和__syncthreads()但本次单次测量中收益大于成本。一个看似反常的现象是分块版 L1 命中率从约87%降到约2%Kernel 却更快。原因是朴素版发出了大量重复的全局内存请求其中很多被 L1/L2 命中分块版通过共享内存复用直接消除了大部分全局请求。请求总量减少以后L1 命中率下降并不代表性能变差。这给我的经验是先看绝对耗时和请求总量再看命中率不能孤立地追求某个比率变高。七、为什么 Systems 和 Compute 的耗时结果不同本次单次采样结果工具/采集朴素版分块版单次加速比Nsight Systems2.389605 ms1.996393 ms1.20×Nsight Compute Basic2.09 ms1.32 ms1.58×Nsight Compute Detailed2.05 ms1.30 ms1.58×三个结果方向一致分块版更快但具体幅度不同。可能影响因素包括两个工具的采集目标和开销不同Compute 为采集计数器可能重放 Kernel每份报告来自独立进程和独立运行GPU 频率、温度和缓存状态存在变化本次没有预热每个 Kernel 只有一个有效样本。因此当前能够下的结论是已验证单次采样分块版本耗时更低全局加载指令降至朴素版本的1/16。但不能把1.20×或1.58×写成稳定、通用的加速比。要获得可靠性能结论还需要预热、多次运行并报告 Median、Min、Max 和标准差。八、一套可复用的 CUDA Profiling 清单8.1 编译nvcc-O3-lineinfo-archsm_75-Xcompiler/utf-8-o app.exe app.cu架构参数与目标 GPU 匹配。使用-lineinfo支持源码关联。检查编译和运行退出码。检查 CUDA API、Kernel launch 和同步错误。8.2 Systems 全局定位nsys profile--tracecuda,nvtx--statstrue--force-overwritetrue-o app_systems.\app.exe nsys stats--report cuda_gpu_kern_sum.\app_systems.nsys-rep查看 CPU API、GPU Kernel、Memcpy 和 NVTX 的先后关系。区分业务范围、CPU 等待时间与 GPU Kernel 时间。找到总耗时最大或调用次数最多的 Kernel。不把第一次 CUDA 初始化误判为稳定 API 成本。8.3 Compute 深入分析ncu--setbasic--launch-count 1--force-overwrite--export app_basic.\app.exe ncu--setdetailed--launch-count 1--force-overwrite--export app_detailed.\app.exe先看 Duration再看吞吐率和 Occupancy。不把Memory Throughput直接等同于 DRAM 带宽。不把高 Occupancy 直接等同于高性能。使用 Baseline 保持对照关系。需要精确数量时进入 Raw 页面。使用 Source 页面将热点关联回 CUDA C 源码。九、总结这次实践形成了一个比较清晰的 CUDA 性能分析顺序解决编译环境 ↓ 确认程序正确运行 ↓ 用 NVTX 标记业务阶段 ↓ 用 Nsight Systems 找到关键 Kernel ↓ 用 Nsight Compute 查看 Basic 和 Baseline ↓ 用 Detailed、Raw、Source 验证优化原因 ↓ 保持相同条件重新测量最重要的经验不是记住某个指标而是避免孤立解读指标高 Occupancy 不等于高性能高命中率不等于请求数量合理CPU 同步时间不等于 GPU Kernel 时间第一次 CUDA 调用不能代表稳定运行成本单次采样只能说明本次观察不能证明稳定加速。参考资料Infatoshicuda-courseInfatoshiCUDA Profiling 示例05_Writing_your_First_Kernels/03 Profiling