这张图基本可以概括我今天想聊的全部内容一个最朴素的向量加法 kernel从能跑到能看懂它为什么快、为什么慢中间隔着一次 Nsight Compute 的性能剖析。作为 CUDA 编程系列第三篇这篇文章不再停留在“怎么写 kernel”的阶段而是把重点放在“怎么看性能”。我会从向量加法这个最简案例出发先理清 CUDA 执行模型里几个必须建立的概念再完整过一遍从代码编写、编译到运行的实操流程最后用 Nsight Compute 对 kernel 做一次真实分析逐步拆解 occupancy、带宽利用率和 stall 这些指标。不管你是刚开始接触 CUDA 的学生还是从 PyTorch 等框架侧切进来、想了解底层到底发生什么的开发者这篇文章都能给你一条能直接落地的路径。1. 运行模型与向量加法的关键设计细节向量加法在 CUDA 里相当于编程语言里的 Hello World但很多人学完只是看了个热闹写出的代码能跑却说不清为什么 grid 和 block 要这么设为什么数据要先拷贝到显存再拷回来。这些问题不解决后面看 Nsight Compute 的分析结果就会一头雾水所以先把关键概念压一压。1.1 从线程层级开始理解 GPU 的“人海战术”CPU 擅长处理复杂分支和单线程性能但 GPU 的思路完全不同它用成千上万个简单线程同时执行同一条指令靠“量”取胜。CUDA 把这套并行模型抽象成三层grid一次 kernel 启动对应一个 grid可以理解为一整个任务。blockgrid 内部由多个 block 组成每个 block 内部有一个共享内存空间block 里的线程可以相互协作。thread真正执行指令的单元每个线程通过内置变量 threadIdx、blockIdx、blockDim 来定位自己处理的数据范围。向量加法里最自然的做法是一个线程负责计算一个输出元素。假设有 N 个元素我们要启动 N 个线程。但 GPU 硬件不可能管理无限多个独立的“调度单元”它真正调度的是 warp——通常 32 个线程为一组。所以如果 block 大小是 128一个 block 内部就包含 4 个 warp硬件以 warp 为单位分发指令。这句话非常重要因为后面分析 Nsight Compute 的很多指标比如 warp stall、occupancy都离不开这个基础认知。1.2 内存模型与数据拷贝为什么不能直接操作“普通内存”另一类必须建立的概念是内存。GPU 有自己独立的显存device memoryCPU 侧的数据必须先拷贝到显存kernel 执行完再把结果拷贝回来。这不是 CUDA 设计得麻烦而是物理架构决定的PCIe 总线带宽和显存带宽不是一个量级数据搬移成本往往比计算本身还高。CUDA 用 cudaMalloc 在显存上分配空间用 cudaMemcpy 在主机端和设备端之间搬运数据最后一个 cudaFree 释放显存。很多刚上手的人会忽略 cudaMemcpy 的同步语义——它是阻塞的也就是说在数据拷贝完成之前CPU 会等在那里。这一点看起来不起眼却会导致整个程序的时间统计出现偏差后面我会专门讲怎么处理。1.3 为什么向量加法是个“带宽受限”问题向量加法的计算强度极低每个元素只需要一次浮点加法却要读两个数、写一个数。换句话说计算单元大部分时间在等数据从显存传过来而不是在计算。这种问题被归为 memory-bound性能好坏主要取决于你能否把显存带宽用满而不是能不能把 ALU 用满。理解这一点后你就会明白为什么优化策略和计算密集型的矩阵乘法完全不同矩阵乘要多用共享内存和寄存器重用数据而向量加法首先要减少无效的数据搬运、让访问尽量对齐合并。Nsight Compute 的核心任务之一就是帮你判断当前 kernel 是计算受限还是访存受限然后才有对症下药的优化方向。2. 向量加法从零实现与第一次性能感知理论说清楚了接下来真正动手写代码。我不会只贴一份能跑的代码而是把每一行背后为什么这么写、有哪些容易踩的坑讲明白。2.1 完整代码实现#include stdio.h #include cuda_runtime.h __global__ void vectorAdd(const float *a, const float *b, float *c, int n) { int i blockIdx.x * blockDim.x threadIdx.x; if (i n) { c[i] a[i] b[i]; } } int main() { int n 1 20; size_t bytes n * sizeof(float); float *h_a (float *)malloc(bytes); float *h_b (float *)malloc(bytes); float *h_c (float *)malloc(bytes); for (int i 0; i n; i) { h_a[i] 1.0f; h_b[i] 2.0f; } float *d_a, *d_b, *d_c; cudaMalloc(d_a, bytes); cudaMalloc(d_b, bytes); cudaMalloc(d_c, bytes); cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice); cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice); int blockSize 256; int gridSize (n blockSize - 1) / blockSize; vectorAddgridSize, blockSize(d_a, d_b, d_c, n); cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost); int errors 0; for (int i 0; i n; i) { if (h_c[i] ! 3.0f) errors; } printf(errors: %d\n, errors); cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); free(h_a); free(h_b); free(h_c); return 0; }这一段代码覆盖了完整流程分配主机内存 - 初始化数据 - 分配显存 - 数据拷贝到设备 - 启动 kernel - 数据拷回主机 - 校验结果 - 释放资源。核心那行索引计算blockIdx.x * blockDim.x threadIdx.x是整个 CUDA 编程里最常见的代码必须一眼看懂blockIdx.x 表示当前是第几个 blockblockDim.x 表示一个 block 里有几个线程threadIdx.x 表示当前线程在 block 里的编号。三者一乘一加就把所有线程映射到数组下标上。2.2 边界检查与 gridSize 计算我见过很多初学者写 kernel 时不做边界检查结果数组长度不是 blockSize 整数倍时就会有线程越界读写显存程序可能出现无法解释的随机错误。正确做法是启动足够多的线程比如 N120、blockSize256 时gridSize (N blockSize - 1) / blockSize。这个计算公式的本质是向上取整如果 N 不能被 blockSize 整除最后一个 block 会有部分线程超出边界。kernel 内部那行if (i n)就是拦截这些越界线程的防线。2.3 头文件、编译脚本与依赖你会遇到的第一批坑代码本身不难但第一次编译的人很容易卡在环境配置上。这里基于我实际使用经验总结几个要点编译命令用nvcc -o vector_add vector_add.cunvcc 是 CUDA toolkit 自带的编译器驱动它会把 .cu 文件里的主机代码交给 gcc/g设备代码编译成 cubin/PTX。如果你在 Windows 上用cl.exe的环境一定注意 nvcc 需要能和 Visual Studio 的编译器配合。经常有人装好 CUDA 后一编译就报找不到 cl.exe一般是没在“x64 Native Tools Command Prompt”里操作。Ubuntu 22.04 下装 CUDA 时建议用 runfile 或者 deb 包装完记得把/usr/local/cuda/bin加到 PATH/usr/local/cuda/lib64加到 LD_LIBRARY_PATH。很多人跑nvcc --version有输出但一运行程序就报libcudart.so: cannot open shared object file基本就是 LD_LIBRARY_PATH 没配好。在 WSL2 里跑 CUDA 比较省心Windows 侧装好 NVIDIA 驱动后WSL2 里直接装 CUDA toolkit 就能用。但要注意WSL2 里不支持 NVIDIA 厂商驱动只支持微软提供的驱动映射功能受限但日常学习足够。编译运行完后你应该会看到errors: 0。到这里代码已经跑通了但性能如何完全不清楚。下一步就用 Nsight Compute 来量化。2.4 一个常见的坑运行时 API 与驱动 API 辨析很多人在排查环境时会看到两个概念runtime API如 cudaMalloc、cudaMemcpy和 driver API如 cuInit、cuCtxCreate。前者写业务代码用得多后者负责更底层的上下文管理。Nsight Compute 本质上是基于 driver API 和 NVIDIA 硬件性能计数器做采样的工具。如果机器上 CUDA 版本和驱动版本不匹配或者 Nsight Compute 版本落后于驱动很容易出现“工具打不开”、“分析失败”的问题。另外nvidia-smi里显示的 CUDA Version 是驱动支持的最大版本不代表当前环境实际用的是哪个 toolkit。很多热词里有人问“3060 怎么确定安装了 CUDA”正确做法是先跑nvcc --version确认 toolkit 版本再跑nvidia-smi确认驱动版本两个都存在且能编译运行 sample才算真正装好。只装了显卡驱动但没有安装 CUDA toolkitnvcc是找不到的。3. Nsight Compute 性能分析实操Nsight Compute 是 NVIDIA 官方推出的 kernel 级性能分析工具取代了早期的 Nsight Visual Studio Edition 里的一部分功能专门用于 Profiling CUDA kernel。它能给出 occupancy、带宽利用率、指令占比、warp stall 原因等极其详细的数据是优化 CUDA 程序的重要依据。3.1 基本分析流程首先得保证 Nsight Compute 已安装。CUDA toolkit 完整安装默认会带上 Nsight Compute装完后可以直接在命令行用ncu访问它。分析单个 kernel 最简单的方式是在可执行文件后面直接加 profile 参数ncu --set full ./vector_add这条命令会对程序里每个 kernel 做一次完整分析。由于 vector_add 只有一个 kernel所以输出不会太长。第一次跑的时候你大概率会看到一大堆整齐排列的表格重点先看几个关键行Durationkernel 真正在 GPU 上执行的时间单位是纳秒或微秒。注意它不包含数据拷贝时间。Compute (SM) ThroughputSM流处理器的计算资源利用率。数值越低说明计算越闲。Memory Throughput显存带宽利用率。向量加法这个值通常远高于计算利用率这正好验证了它是个 memory-bound 问题。Achieved Occupancy实际达到的占用率指每个 SM 上活跃 warp 数占最大可容纳 warp 数的比例。有时候你配置的 block 数不少但受限寄存器数或共享内存实际 occupancy 并不高。我自己第一次跑 vectorAdd 时结果很典型Compute (SM) Throughput 只有百分之十几Memory Throughput 却到了百分之七八十Achieved Occupancy 在 80% 以上。看到这几个数基本可以下结论kernel 本身没问题瓶颈确实在显存访问。3.2 如何读懂 Data Movement 信息--set full会输出很多细节其中有一个 Section 叫“Data Movement”或“Memory Workload Analysis”它会显示全局内存读请求、写请求以及访问模式。对向量加法而言最理想的情况是每个线程访问的数据地址连续且对齐这样 warp 内 32 个线程访问的 128 字节可以合并成少数几个 sector 访问传输效率最高。如果你把 kernel 改成c[i n/2]或c[i * 4]就能看到所谓的“非合并访问”uncoalesced access出现性能立刻下降。这算是验证内存访问模式对性能影响最快的手段。3.3 重点指标解读occupancy 与 stall 是两个不同维度对于刚接触性能分析的人最容易混淆的是 occupancy 和 stall。Occupancy 说的是“台子上同时有多少活人在干活”。占用率高不代表每个人效率高只代表可调度的 warp 数量充足能把显存延迟掩盖掉一部分。Stall 说的是“当前 warp 为什么停下来等”。比如等待全局内存数据返回long scoreboard、等待固定功能单元wait、等待下一轮指令not selected等等。高质量的分析流程是先看 occupancy 够不够再看 stall 主要是什么原因。如果 occupancy 很高但仍有大把 stall 耗费在访存等待上说明问题出在内存延迟本身优化的思路可以是提升访存局部性、使用向量化访问、或者用共享内存预处理。如果 occupancy 本身很低那首先要查是不是寄存器分配太多或者 block 尺寸不合理把活跃 warp 数顶上去再说。一个很容易犯的错误是一上来就追求 100% occupancy。比如我的向量加法如果改用 128 block、每线程一个元素occupancy 反而受 block 数量限制。实际优化时只要达到 50% 以上且能把带宽用满对 memory-bound 的 kernel 来说就足够了。3.4 查看源级关联与汇编级执行另一件值得提的事Nsight Compute 的 Source 和 SASS 视图。它能精确到某一行 C 代码对应哪些 SASS 指令以及这些指令的 issue 情况。很多人以为优化是玄学真到了 SASS 层面你会发现编译器已经把很简单的循环展开成了向量指令或者反过来说编译器并没有按你想的那样把中间计算缓存到寄存器里而是反复访问了全局内存。拿向量加法来说你a[i] b[i]这里很简短但编译器是否生成LDG.E这种 128 位向量访存指令直接决定性能上限。分析 SASS 的意义在于核实编译器做了什么而不是站在 C 代码层面想当然。这个能力对理解复杂 kernel 更加重要向量加法只是一个起点。4. 常见问题与排查技巧实录这部分内容来自我日常折腾环境和性能分析过程中踩过的坑也结合大家在常见问题里的搜索需求整理成一个个快速定位的清单。4.1 性能计数器打开失败ERR_NVGPUCTRPERM刚在 Linux 下跑ncu的同学非常容易碰到类似ERR_NVGPUCTRPERM: The user does not have permission to access NVIDIA GPU Performance Counters的报错。原因是 Linux 默认禁止普通用户访问性能计数器需要 root 权限或者给当前用户配置权限。处理办法其实不复杂基本思路是让当前用户加入系统允许访问性能计数器的组。不同发行版组名不同常见的是wheel或adm也有的是直接修改nvidia模块参数sudo usermod -a -G wheel $USER改完需要注销重新登录。如果还不行去 NVIDIA 开发者论坛搜一下有没有针对你内核对性能计数器加锁的情况。4.2 数据拷贝时间为什么也被算进“程序运行时间”写性能测试的时候很多人会在 kernel 启动前后用clock()或std::chrono计时却把 cudaMemcpy 的时间也算进去。这会让结果看起来特别慢因为 PCIe 搬运本身就有几百微秒级别的开销。如果只是想测 kernel 本身应该用 CUDA eventcudaEvent_t start, stop; cudaEventCreate(start); cudaEventCreate(stop); cudaEventRecord(start); vectorAddgridSize, blockSize(d_a, d_b, d_c, n); cudaEventRecord(stop); cudaEventSynchronize(stop); float ms 0.0f; cudaEventElapsedTime(ms, start, stop); printf(kernel time: %f ms\n, ms);Event 和之前讲的内存拷贝不同它是异步的需要 cudaEventSynchronize 等 GPU 执行完成后再取时间差。这个点虽然小却是很多人统计时间时出错的原因。4.3 核函数启动失败却不报错默认情况下kernel 启动是异步的CPU 不会立刻感知 GPU 上的错误。常见表现是程序不崩、后续 cudaMemcpy 卡住然后莫名其妙报一个cudaErrorIllegalAddress或者cudaErrorInvalidValue。排查手段就一条每次 kernel 启动后检查错误码cudaError_t err cudaGetLastError(); if (err ! cudaSuccess) { printf(CUDA error: %s\n, cudaGetErrorString(err)); }这属于写 CUDA 代码最基本的卫生习惯强烈建议每个 kernel 后面都跟一个。4.4 与 cuBLAS 相关的经典报错很多人搜过cublas_status_execution_failed这个错误这里也顺带说几句。这个报错高发于调用 cuBLAS API 时输入输出指针不是显存地址、维度参数错误或者设备上下文被前一个 kernel 的错误污染。它作为通用执行失败码属于“下游结果”真正的根因往往在上面几行。我的排查习惯是先检查所有传入 cuBLAS 的指针是否显存地址再看有没有先前的 kernel 已经报错最后核对 lda 这类矩阵leading dimension参数是否合法。很多时候把前面的错误清掉这个报错自然消失。4.5 在 WSL2 或 Windows 下使用 Nsight Compute 的经验Windows 和 WSL2 下对 Nsight Compute 的支持已经有明显改善但仍有几个坑值得注意WSL2 上跑ncu驱动需要支持 GPU Performance Counters API部分数据集驱动初版可能有兼容问题建议先更新到较新驱动。Windows 原生环境下ncu 从命令行运行没问题但如果想图形界面 GUI则需要 Nsight Compute 的独立版本。命令行操作其实够用也更容易自动化。如果在 C4D 渲染时遇到there is no cuda device which is selected类似的提示这其实是渲染器对 CUDA 设备的枚举出错了常见原因是显卡驱动和渲染器版本不兼容或者显卡处于持久模式未被正确识别。解决方向上先更新驱动确认显存能被识别再考虑重装渲染器程序。4.6 快速确认一张卡到底能不能用 CUDA搜索里经常有 “3060 怎么确定安装了 CUDA” 这类问题。我的建议是三步验证nvidia-smi能看到显卡且右上角 CUDA Version 不是 None。nvcc --version能看到 toolkit 版本。编译并运行一个 sample比如deviceQuery或我们这篇文章的向量加法能正确算出来。三步都过环境就没问题。如果确实需要同一台机器上多个 CUDA 版本可以用环境变量切换PATH和LD_LIBRARY_PATH或者用 conda 在虚拟环境里指定 cudatoolkit这个方案对研究和开发来说都更省心。5. 写在最后的个人体会从向量加法到 Nsight Compute看起来只是加了一个 profiling 步骤其实是两套完全不同的思维模式。前者要求你理解 CUDA 的线程模型和内存模型后者要求你把性能数据当成另一门“语言”去解读。我个人在实际调试 kernel 的时候最大的心得体会是不要拿着任何单一指标下结论。occupancy 高不代表性能好Memory Throughput 满了也不代表一定到极限SASS 里指令少了也不一定快。真正的分析流程是不断组合多个指标在大方向上判断瓶颈类型在小方向上定位具体指令和访存模式。比如今天这个向量加法如果只看计算资源你可能会觉得 GPU 很闲但一旦结合 Memory Throughput你会立刻意识到瓶颈在显存带宽上。后面如果要继续深入可以从两个方向走一是把 vectorAdd 从全局内存版本改为 shared memory 分块版本做一个完整的数据复用优化对比二是换一个计算密集型的例子比如矩阵乘法看计算受限问题里如何用 tiling 和指令级并行榨干 SM。不管走哪条路N s ight Compute 都会是你最常用的分析工具之一如何读懂它的报告远比多写几个 kernel 更重要。