如果你已经会写一点CUDA那么这套官方给出的CUDA Samples 13.3一定是绕不开的参考代码如果你刚接触GPU开发我更建议把它当成一套带注释的实战教材来读。这个版本在目录结构、示例覆盖面和构建体验上都比较成熟200多个示例覆盖了从入门核函数到多卡通信、CUDA Graph、统一内存、图形互操作等几乎所有主题。我这次做GPU开发时把整套源码反复对照过好几遍也踩了不少驱动、编译和性能验证层面的坑所以打算把架构全景、源码分层逻辑、经典示例的拆解思路以及怎么把这套Samples真正变成自己的工程能力一次性整理清楚。1. 拿到13.3源码包后我建议你先忘掉“跑通Demo”这件事很多人在GitHub上下载了NVIDIA/cuda-samples仓库或者在CUDA Toolkit安装目录里找到samples文件夹然后选中一个示例make一下看到屏幕上滚出一堆数字就算完事。这种做法不能说错但基本浪费了这套源码80%的价值。CUDA Samples的价值不在“能跑”而在“为什么这么写”“换一种写法性能会差多少”“哪些代码可以直接搬进生产项目”。1.1 这套Samples到底装了什么从目录树看全貌先看获取方式。第一是直接克隆GitHub上的nvidia/cuda-samples仓库切到对应Release标签比如13.3第二是安装CUDA Toolkit时自带的samples目录Linux下通常在/usr/local/cuda/samplesWindows下在CUDA安装目录下的Samples文件夹里。两条路径内容一致区别只是GitHub仓库更新更及时。解压或克隆下来之后第一眼看到的是一串带编号的顶层目录。这个编号不是随便排的它既是难度梯度也是应用领域的划分目录名称主要内容适合人群0_SimpleSimple最基础的CUDA概念示例如vectorAdd、matrixMul、reduce刚入门的新手1_UtilitiesUtilities设备查询、带宽测试、P2P通信等工具类程序所有开发者尤其做环境验证时2_GraphicsGraphicsOpenGL/DirectX互操作、图形渲染相关图形学、可视化方向3_ImagingImaging图像处理算法如boxFilter、bilateralFilter、convolutionSeparable图像处理、CV方向4_FinanceFinance金融衍生品定价如BlackScholes、BinomialOptions量化计算方向5_SimulationsSimulations物理仿真如FDTD3d、SmokeParticles、fluidsGL科学计算、物理模拟方向6_AdvancedAdvanced更底层的进阶主题如CUDA Graph、stream、multiGPU、cooperativeGroups有经验的性能优化开发者7_CUDALibrariesCUDALibraries调用cuBLAS、cuFFT、cuRAND、NPP、thrust等官方库的示例准备用库而不是自己写kernel的开发者这个表格看一眼就明白Samples不是单纯的教学代码它还承担着“官方API效果展示”和“软硬件能力验证”的功能。读懂这个目录结构比你多跑几个示例重要得多。1.2 评测环境与跑通的第一个sample我这次是在一台Ubuntu 22.04系统上做的评测CUDA Toolkit版本12.8显卡分别试过Ada架构的消费卡和一块Ampere架构的数据中心卡。Ubuntu下编译Samples非常简单以最经典的0_Simple/matrixMul为例cd 0_Simple/matrixMul make clean make ./matrixMul正常运行会先打印设备信息然后是“Matrix Multiplication”的测试结果最后是一行Performance相关的数字。这里有个细节如果make报错说找不到nvcc先确认nvcc -V有没有输出没有就检查/usr/local/cuda/bin有没有加进PATH或者CUDA_HOME环境变量是不是指向了正确位置。Windows平台则简单很多直接用Visual Studio打开Samples_vs2022.sln或者对应版本的解决方案文件选中某个示例设置为启动项目编译即可。但我必须强调一点跑通这个sample只是热身。真正有价值的是下一步——打开它的.cu源文件把每一行和CUDA编程模型对照起来读。2. 架构全景CUDA Samples 13.3的分层逻辑不是随手排的如果你把整个仓库按目录拉平对着源码逐个看会发现NVIDIA的分层思路实际上很清晰前几个目录解决“这个功能怎么做”后几个目录解决“这个功能怎么做好”。理解这种双轨逻辑是读源码的门槛。2.1 从0_Simple到7_CUDALibraries每一层的门槛与价值0_Simple这一层的核心是“一个示例对应一个知识点”而且尽量不引入额外概念。比如vectorAdd只讲grid/block/thread三级组织和简单的数据拷贝matrixMul讲共享内存和线程协作reduce讲跨线程跨block归约。对新手来说这一层的代码量小、注释少而精稍微有点C语言基础就能跟上。1_Utilities这一层我建议所有人优先跑一遍。deviceQuery会列出当前机器的设备属性包括计算能力Compute Capability、共享内存每block大小、最大线程数、常量内存大小等bandwidthTest测试设备到主机的PCIe/系统内存带宽p2pBandwidthLatencyTest专门验证多卡之间的P2P通信带宽和延迟。一个常见场景是拿到一台新服务器先用deviceQuery确认驱动和CUDA是否装好再跑bandwidthTest确认数据传输带宽是否正常。这个习惯能帮你省掉大量环境排查时间。6_Advanced这一层的门槛明显抬高但它是整个Samples里最值得反复咀嚼的部分。比如cudaGraphs演示如何用CUDA Graph把多个kernel依赖关系固化成一张图减少调度开销streamOrderedAllocation演示如何用流序内存分配解决多stream间的内存依赖cooperativeGroups讲解如何用协作组做更细粒度的线程同步。这些主题都有对应的官方文档但文档写得再好也不如一个能编译、能运行、能改参的示例来得直接。7_CUDALibraries则换了一条路不教你写kernel而是教你怎么用官方库。对绝大多数业务开发来说能用cuBLAS、cuFFT、thrust解决的事情根本没必要自己造轮子这一层就是“官方库使用说明书”。2.2 泛型算法、嵌入式与可视化样例的位置与取舍细看目录会发现几个泛型算法示例被分散在不同的层级0_Simple/reduce是基础归约6_Advanced里还有更复杂的归约变体3_Imaging里有卷积、滤波、直方图这些图像算法但它们的优化思路同样适用于通用数据处理。NVIDIA这么安排的用意很明显一方面让初学者在简单目录里就能接触核心思想另一方面在进阶目录里给出满足生产性能的版本。比如二维卷积simple版可能直接用全局内存加原子操作imaging版会刻意优化共享内存访问模式两者对比阅读你对“优化到底优化了什么”的认识会特别清晰。此外仓库里还有一批面向Jetson等嵌入式平台的示例通常以tegra或EGLStream命名。这类示例对纯服务器开发的人暂时用不上但对做边缘计算、机器人、自动驾驶的朋友就是宝贵参考。3. 源码级拆解三个经典sample读透CUDA核心机制接下来进入重头戏。我不打算把200多个示例全部过一遍那样既琐碎也没意义。挑三个我反复研究过、也是面试和实际项目中最高频出现的经典示例matrixMul、threadFenceReduction、histogram。把这三个读透你对CUDA的线程同步、共享内存、原子操作、归约设计这几个核心机制的理解会上一个台阶。3.1 matrixMul共享内存tiling与线程同步的教科书matrixMul对应的完整源码在0_Simple/matrixMul目录文件名叫matrixMul.cu。它实现的是经典的C A * B矩阵乘法朴素思路很直观每个线程负责计算输出矩阵的一个元素于是需要读取A的一整行和B的一整列。问题在于全局内存的访问开销很大一个线程读一遍整块矩阵要反复从显存搬运数据。优化的核心是共享内存分块也就是tiling。基本流程是定义一个BLOCK_SIZE比如16x16每个block负责计算输出矩阵中BLOCK_SIZE x BLOCK_SIZE的一个子块。计算这个子块时先协作把A的对应列切片和B的对应行切片加载到共享内存里然后__syncthreads()等待所有数据就位再继续乘加。下一次循环再把下一块数据加载进来。关键代码结构大致是__global__ void MatrixMulKernel(float* C, const float* A, const float* B, int wA, int wB) { __shared__ float As[BLOCK_SIZE][BLOCK_SIZE]; __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE]; int bx blockIdx.x, by blockIdx.y; int tx threadIdx.x, ty threadIdx.y; int aBegin wA * BLOCK_SIZE * by; int aEnd aBegin wA - 1; int aStep BLOCK_SIZE; int bBegin BLOCK_SIZE * bx; int bStep BLOCK_SIZE * wB; float Csub 0.0f; for (int a aBegin, b bBegin; a aEnd; a aStep, b bStep) { As[ty][tx] A[a wA * ty tx]; Bs[ty][tx] B[b wB * ty tx]; __syncthreads(); #pragma unroll for (int k 0; k BLOCK_SIZE; k) { Csub As[ty][k] * Bs[k][tx]; } __syncthreads(); } int c wB * BLOCK_SIZE * by BLOCK_SIZE * bx; C[c wB * ty tx] Csub; }读这段代码时有一个细节特别值得注意循环体内为什么在加载数据之后和计算完成之后各有一个__syncthreads()第一个同步是为了确保As和Bs里的所有数据都写完了才允许任何线程去读第二个同步是为了防止计算较快的线程提前进入下一轮循环把共享内存里的旧数据覆盖掉。少了任何一个结果都是偶发错误这种错误极其恶心因为它不是必现而是跟线程调度顺序有关。共享内存tiling为什么能提速从访存量来看朴素版本每个线程访问A一行中所有元素同一个A矩阵元素会被多个线程重复读取tiling之后每个元素在共享内存里只加载一次然后被block内多个线程复用。全局内存访问量降了一个量级性能差距自然明显。实测下来对较大矩阵这个tiling版本相比朴素版本能获得数倍到数量级的提升具体倍数跟矩阵规模相关。3.2 threadFenceReduction一次kernel里跨block归约的fence语义这个示例提供了很多开发者在刚写多block协作时犯迷糊的解法。通常归约需要把所有block的中间结果合到一起最简单的方案是启动两次kernel第一次每个block做局部归约第二次用单个block把所有局部结果收尾。threadFenceReduction却只启动一次kernel就完成了跨block归约关键靠两样东西共享内存里的block私有归约结果、全局内存里的计数器和__threadfence()。它的逻辑大致是每个block先把自己的最终结果写到全局内存数组的对应位置然后执行__threadfence()确保这条写入对其他block可见接着一个线程对全局计数器做原子加一得到返回值。如果返回值等于gridDim.x减一说明当前block是最后一个到达的block由它读取其他block写好的所有局部结果算出最终结果写回。这里的核心是__threadfence()的语义它保证调用线程在fence之前的所有全局内存写入对另一个线程在观察到某个同步标志之后执行的读取操作是可见的。没有这个fence最后一个block去读其他block结果时可能读到的是没来得及刷新到显存的数据。很多刚接触的人会问能不能用__syncthreads()做跨block同步不能。__syncthreads()只同步block内线程它没有任何机制保证一个block能看到另一个block的数据。threadFenceReduction里的计数器加原子操作实际上是在构建一个轻量的全局软件屏障。我从实际开发里得到的经验是这种单kernel归约写法能省一次kernel启动开销在kernel启动延迟敏感的短计算任务里有用但它的可扩展性和调试性不如两阶段kernel。真正做生产项目时我通常优先考虑两阶段方案只有在对延迟要求极其苛刻、并且block数量可控的场景下才用软件屏障。Samples的价值就在于它把这种“偏方”和标准做法摆在同一个仓库里让你在选择时心里有数。3.3 histogram原子操作与局部直方图的取舍histogram是图像处理和数据分析里极其常见的算子但它在GPU上优化起来有不少门道。样例位于3_Imaging/histogram展示了几种不同实现。最朴素的实现是每个线程处理一个像素直接对全局内存的直方图数组做atomicAdd。功能上完全正确性能却很糟因为所有线程在同时争抢少数几个计数器全局原子操作在显存级别上的冲突会让吞吐量崩塌。优化思路是对直方图做“私有化”privatization每个block维护自己的一份直方图放在共享内存里线程优先更新共享内存内的计数器因为共享内存的原子操作比全局内存快得多。等到整个block处理完数据再做一次合并把block的局部直方图加到全局结果里。合并过程可以用一个线程顺序做也可以用多个线程配合原子操作并行做。这个示例还揭示了两个权衡。第一是共享内存容量bin数量越多私有直方图占用的共享内存越大。假设直方图有256个bin每个bin用32位整数每block就需要1KB共享内存看起来不多但如果bin数量到4096或更高或者你想同时维护多个通道开销就很可观。第二是归并成本block数越多合并阶段的全局原子操作还是越多所以block粒度需要根据数据规模和硬件资源调。我后来在做自定义算子时把histogram里的私有化思想扩展到了很多“先统计、后归并”的算法里这类模式几乎可以成为GPU归约类问题的通用模板。4. 从Samples移植到工程错误处理、基准确认与可复用模式读源码和写工程之间还有一道鸿沟示例代码为了突出重点常会省略一些工程细节但它也留下了大量值得直接复用的骨架。下面这几个东西是我从Samples里拷贝到生产项目里最频繁的。4.1 CUDA_SAFE_CALL这类宏是Samples留给工程的第一笔财富CUDA API有一个容易踩的坑大量API是异步的kernel启动之后错误并不会立刻返回而是在某个同步点才会冒出来。如果你每次API调用后都不检查返回值遇到“莫名其妙的脏数据”时排查成本极高。Samples里经常出现的checkCudaErrors宏本质上是一个封装的错误检查器#define CUDA_CHECK(call) \ do { \ cudaError_t err (call); \ if (err ! cudaSuccess) { \ fprintf(stderr, CUDA error %s:%d: %s\n, __FILE__, __LINE__, \ cudaGetErrorString(err)); \ exit(EXIT_FAILURE); \ } \ } while (0)生产环境里你甚至可以把exit换成更友好的错误上抛或饥饿信号处理。但最核心的思想是一致的每一次cudaMalloc、cudaMemcpy、cudaLaunchKernel之后都要做检查这样才能把问题定位到具体调用点。这个习惯是真的能救命我调试过的最难缠问题里有相当一部分最后都归因于某次kernel启动后没有检查异步错误。4.2 每个sample都带CPU baseline性能验证才有据可依细读Samples时你会发现它几乎给每个优化示例都配了一个CPU参考实现。比如matrixMul在GPU kernel之外还写了一个CPU端的朴素乘法用于校验结果是否正确nbody示例里也保留了一个CPU端计算引力的参考函数。这个设计非常值得抄进自己的工程里。做GPU性能优化时只测GPU kernel的运行时间是不够的。你得知道三件事一、结果对不对二、比CPU快了多少三、比“朴素但正确”的GPU版本快了多少。没有CPU baseline你无法判断加速比没有朴素GPU版本对照你无法判断优化手段的实际收益。我现在的习惯是每个优化版本至少保留一个对应的正确性校验函数和一个朴素版本跑性能时同时输出三组数据。Samples自带的这套“参考实现计时结果比对”模板几乎是现成的测试框架。另外计时本身也有讲究。Samples里普遍使用cudaEvent来记录GPU时间这是因为GPU kernel的启动是异步的用clock()或者chrono测到的经常包含CPU端排队等待时间。正确做法是创建一个cudaEvent在kernel前后分别cudaEventRecord然后cudaEventSynchronize取cudaEventElapsedTime。Samples里到处是这样测时的例子照抄就行。4.3 直接能从Samples拷走的几个代码模式和习惯把整个仓库当“代码模板库”能直接拿走的东西比我预期的多deviceQuery初始化后自动获取设备属性、选择计算能力最高设备的模式适合放在任何GPU项目的启动阶段。checkCudaErrors错误处理宏上面已经讲了。CPU参考实现与GPU结果比对的模式适合作为单元测试的雏形。用cudaEvent计时的模板比用CPU时间戳可靠。矩阵乘法中通过宏控制BLOCK_SIZE、通过编译参数切换kernel版本的思路适合做性能调参实验。每个示例独立目录、自带Makefile、一条make就能跑的工程组织方式我把它迁移到了自己的实验仓库里效果很好。5. 编译、调试与性能核验13.3在真实环境中的操作链路拿到Samples只是第一步跑起来、改起来、验证起来才有价值。这一节记录我在真实环境里编译、调试、以及用性能工具核验示例时的一些操作链路和踩坑记录。5.1 跨平台构建make体系、VS工程与架构参数Linux/macOS下Samples的构建依赖一套统一的make体系顶层有Makefile子目录里也有简单的Makefile它们会调用common.mk里的公共逻辑。日常用到的基本命令make # 编译当前目录的示例输出到 bin/linux/release/ make dbg1 # 编译debug版本 make clean # 清理 make TARGET_ARCHx86_64 # 显式指定架构如果你是交叉编译比如在x86主机上构建Jetson设备用的示例可以指定类似make TARGET_ARCHaarch64这个过程中最容易踩的坑是“架构不匹配”你有一块Ada显卡但make默认生成的目标代码里没有包含适合你显卡的计算能力。Samples的Makefile一般会通过cuda_arch变量控制编译目标架构比如make cuda_archsm_89如果不确定自己的显卡计算能力先跑deviceQuery看输出或者用nvidia-smi --query-gpucompute_cap --formatcsv。算力不对后果有两种一种是你拿到的是一个JIT编译出来的PTX性能打折扣另一种是驱动太老不认识新的SASS直接报加载失败。Windows端相对轻松打开解决方案文件之后在工程属性里指定目标计算架构、CUDA Toolkit路径即可。但有几个sample依赖OpenGL、GLFW等图形库可能需要额外装了依赖才能编译过这类sample通常会在顶层README里标注编译失败时优先看依赖说明。5.2 Nsight Compute把sample的性能假设核验一遍源码里的注释说得再怎么天花乱坠都不如自己跑一次性能分析工具来得直接。CUDA Toolkit自带的Nsight Compute命令行工具叫ncu是我评测Samples时用得最多的工具。以matrixMul为例编译好之后执行ncu --set full ./matrixMul完成后会给出详细报告核心指标包括占用率Achieved Occupancy、共享内存吞吐、DRAM吞吐、指令数、分支发散、bank conflict等。我第一次跑的时候最关注的就是Shared Memory Bank Conflicts这一项。Samples默认的矩阵读取方式基本不会出现bank conflict但如果你把As[ty][k] * Bs[k][tx]这个访问模式改成按行优先存储的行向量操作就很容易制造出冲突。这里有一个通用经验不要只看一个指标。实际工作中我最常犯的错误是用“更高占用率”作为唯一目标结果发现shared memory或者寄存器压力被推高局部性反而变差kernel耗时并没有下降。NVIDIA在Nsight Compute里给出的memory workload分析和occupancy关联提示能帮你快速理解指标之间的耦合关系。Samples里每个优化示例都能作为这种“改参数、看指标”实验的起点。还有一个小技巧用ncu分析某个advanced目录下的示例之前先确认有没有其他进程占用GPU因为现代GPU是分时复用的被其他任务占用会导致测出来的指标漂移。跑分析时用CUDA_VISIBLE_DEVICES把进程限定到某一块空闲卡上能减少干扰。5.3 通用GPU运行时的坑显存、驱动版本与运行参数不同版本的Samples对CUDA运行时版本有最低要求。如果你的驱动版本较老编译可能能过但运行时会报symbol lookup error之类的错比如找不到cudaGraphExecKernelNodeSetParams这类新API。遇到这种情况第一选择是升级驱动第二选择是查看Samples的README确认它需要的CUDA版本。Samples 13.3整体比较新建议搭配较新的CUDA Toolkit使用。显存问题是另一个高频坑。有些示例比如粒子系统、体积渲染默认分辨率或数据规模会吃掉大量显存。如果机器上有多个进程占用了显存运行时会报内存不足。我习惯先跑nvidia-smi看看每块卡的显存和利用率再决定是否设置CUDA_VISIBLE_DEVICES。另外多卡机器的默认设备是0号卡但0号卡不一定是最适合跑计算的卡。生产机器上常有集显或低端卡占着0号位置直接用deviceQuery确认计算能力最强的设备ID然后在运行前设置环境变量指定export CUDA_VISIBLE_DEVICES1这个操作能避免你在不适配的设备上白跑很久。6. 把Samples当成起点而不是终点我的阅读路径与踩坑记录最后这部分是经验之谈。我在把Samples吃透的过程中踩了不少坑也总结出了一条比较高效的阅读路径。如果你准备开始读这套源码可以参考一下。6.1 适合不同目标人群的阅读顺序如果你是纯新手我的建议是这样走先跑0_Simple/vectorAdd搞懂grid/block/thread三级结构和cudaMemcpy的数据搬运逻辑再跑0_Simple/matrixMul理解共享内存和线程同步接着跑1_Utilities/deviceQuery和bandwidthTest了解你手头硬件的真实参数之后可以看3_Imaging/boxFilter或者5_Simulations/SmokeParticles体验一下稍微完整的算法在GPU上怎么组织。如果你已经在写CUDA但觉得性能上不去直接跳到6_Advanced。重点看cudaGraphs、streamOrderedAllocation、cooperativeGroups、concurrentKernels这几个目录。它们解决的是“kernel启动开销怎么降”“多个kernel怎么并行”“线程之间更灵活的协作方式”这些恰恰是生产性能瓶颈最常出现的地方。如果你做的是深度学习推理或数值计算优先看7_CUDALibraries里的cuBLAS、cuFFT、thrust示例。了解官方库的调用范式、内存布局要求比手写kernel更容易落地。6.2 我踩过的三个坑也顺手写在这里第一个坑是盲目改BLOCK_SIZE。我曾经把matrixMul里的BLOCK_SIZE从16改成32想当然地认为更大的block能提高并行度。实测发现性能反而下降用Nsight Compute一测才明白占用率受共享内存容量和寄存器数量限制block变大之后每个SM能驻留的block数减少最终占用率不升反降。Samples的默认值往往是针对大多数显卡调试过的平衡点改之前要带着假设去改而不是“感觉”去改。第二个坑是漏了__threadfence()。我基于threadFenceReduction自己写一个跨block归约时一开始只用了原子操作加一作为计数漏掉了fence导致偶尔计算结果错误。这种问题特别难复现调了一天才定位到。后来我的习惯是涉及跨block可见性时先画清楚“谁在何时写、谁在何时读”然后再决定同步原语绝不靠猜。第三个坑是私有直方图忘清零。在仿照histogram写一个自定义统计算子时我把每个block的私有直方图放到了全局内存里方便多个kernel复用结果忘了在每次kernel调用前清零第二次调用时结果直接叠加上一次的数据。示例里其实有类似的初始化流程我因为没细读这部分踩了坑。这提醒我从Samples里抄代码时不仅要抄核心逻辑还要把初始化、清理、边界条件这些“配角”一起抄过来。6.3 一个可以长期复用的源码阅读习惯GPU实验簿最后分享一个我保持了很久的习惯为每个想深入研究的sample建一个独立实验分支一次只改一个变量并用Nsight Compute记录固定的一组指标。比如对matrixMul我记录Duration、Achieved Occupancy、Shared Memory Bank Conflicts、DRAM Throughput、L1/L2 Cache Hit Rate。改完一个参数运行一遍把数据填进一个markdown表格。时间久了你会建立起非常扎实的“参数-性能数据”直觉。这个方法的核心价值在于它逼着你用数据而不是感觉做判断。Samples提供的就是一套干得很干净的实验素材你只需要在此基础上架一把刻度尺。长期积累下来这套“实验簿”会成为你自己的性能优化手册比任何官方文档都更适合你手头的硬件。
CUDA Samples 13.3实战解析：从共享内存到性能优化，解锁GPU开发全流程
如果你已经会写一点CUDA那么这套官方给出的CUDA Samples 13.3一定是绕不开的参考代码如果你刚接触GPU开发我更建议把它当成一套带注释的实战教材来读。这个版本在目录结构、示例覆盖面和构建体验上都比较成熟200多个示例覆盖了从入门核函数到多卡通信、CUDA Graph、统一内存、图形互操作等几乎所有主题。我这次做GPU开发时把整套源码反复对照过好几遍也踩了不少驱动、编译和性能验证层面的坑所以打算把架构全景、源码分层逻辑、经典示例的拆解思路以及怎么把这套Samples真正变成自己的工程能力一次性整理清楚。1. 拿到13.3源码包后我建议你先忘掉“跑通Demo”这件事很多人在GitHub上下载了NVIDIA/cuda-samples仓库或者在CUDA Toolkit安装目录里找到samples文件夹然后选中一个示例make一下看到屏幕上滚出一堆数字就算完事。这种做法不能说错但基本浪费了这套源码80%的价值。CUDA Samples的价值不在“能跑”而在“为什么这么写”“换一种写法性能会差多少”“哪些代码可以直接搬进生产项目”。1.1 这套Samples到底装了什么从目录树看全貌先看获取方式。第一是直接克隆GitHub上的nvidia/cuda-samples仓库切到对应Release标签比如13.3第二是安装CUDA Toolkit时自带的samples目录Linux下通常在/usr/local/cuda/samplesWindows下在CUDA安装目录下的Samples文件夹里。两条路径内容一致区别只是GitHub仓库更新更及时。解压或克隆下来之后第一眼看到的是一串带编号的顶层目录。这个编号不是随便排的它既是难度梯度也是应用领域的划分目录名称主要内容适合人群0_SimpleSimple最基础的CUDA概念示例如vectorAdd、matrixMul、reduce刚入门的新手1_UtilitiesUtilities设备查询、带宽测试、P2P通信等工具类程序所有开发者尤其做环境验证时2_GraphicsGraphicsOpenGL/DirectX互操作、图形渲染相关图形学、可视化方向3_ImagingImaging图像处理算法如boxFilter、bilateralFilter、convolutionSeparable图像处理、CV方向4_FinanceFinance金融衍生品定价如BlackScholes、BinomialOptions量化计算方向5_SimulationsSimulations物理仿真如FDTD3d、SmokeParticles、fluidsGL科学计算、物理模拟方向6_AdvancedAdvanced更底层的进阶主题如CUDA Graph、stream、multiGPU、cooperativeGroups有经验的性能优化开发者7_CUDALibrariesCUDALibraries调用cuBLAS、cuFFT、cuRAND、NPP、thrust等官方库的示例准备用库而不是自己写kernel的开发者这个表格看一眼就明白Samples不是单纯的教学代码它还承担着“官方API效果展示”和“软硬件能力验证”的功能。读懂这个目录结构比你多跑几个示例重要得多。1.2 评测环境与跑通的第一个sample我这次是在一台Ubuntu 22.04系统上做的评测CUDA Toolkit版本12.8显卡分别试过Ada架构的消费卡和一块Ampere架构的数据中心卡。Ubuntu下编译Samples非常简单以最经典的0_Simple/matrixMul为例cd 0_Simple/matrixMul make clean make ./matrixMul正常运行会先打印设备信息然后是“Matrix Multiplication”的测试结果最后是一行Performance相关的数字。这里有个细节如果make报错说找不到nvcc先确认nvcc -V有没有输出没有就检查/usr/local/cuda/bin有没有加进PATH或者CUDA_HOME环境变量是不是指向了正确位置。Windows平台则简单很多直接用Visual Studio打开Samples_vs2022.sln或者对应版本的解决方案文件选中某个示例设置为启动项目编译即可。但我必须强调一点跑通这个sample只是热身。真正有价值的是下一步——打开它的.cu源文件把每一行和CUDA编程模型对照起来读。2. 架构全景CUDA Samples 13.3的分层逻辑不是随手排的如果你把整个仓库按目录拉平对着源码逐个看会发现NVIDIA的分层思路实际上很清晰前几个目录解决“这个功能怎么做”后几个目录解决“这个功能怎么做好”。理解这种双轨逻辑是读源码的门槛。2.1 从0_Simple到7_CUDALibraries每一层的门槛与价值0_Simple这一层的核心是“一个示例对应一个知识点”而且尽量不引入额外概念。比如vectorAdd只讲grid/block/thread三级组织和简单的数据拷贝matrixMul讲共享内存和线程协作reduce讲跨线程跨block归约。对新手来说这一层的代码量小、注释少而精稍微有点C语言基础就能跟上。1_Utilities这一层我建议所有人优先跑一遍。deviceQuery会列出当前机器的设备属性包括计算能力Compute Capability、共享内存每block大小、最大线程数、常量内存大小等bandwidthTest测试设备到主机的PCIe/系统内存带宽p2pBandwidthLatencyTest专门验证多卡之间的P2P通信带宽和延迟。一个常见场景是拿到一台新服务器先用deviceQuery确认驱动和CUDA是否装好再跑bandwidthTest确认数据传输带宽是否正常。这个习惯能帮你省掉大量环境排查时间。6_Advanced这一层的门槛明显抬高但它是整个Samples里最值得反复咀嚼的部分。比如cudaGraphs演示如何用CUDA Graph把多个kernel依赖关系固化成一张图减少调度开销streamOrderedAllocation演示如何用流序内存分配解决多stream间的内存依赖cooperativeGroups讲解如何用协作组做更细粒度的线程同步。这些主题都有对应的官方文档但文档写得再好也不如一个能编译、能运行、能改参的示例来得直接。7_CUDALibraries则换了一条路不教你写kernel而是教你怎么用官方库。对绝大多数业务开发来说能用cuBLAS、cuFFT、thrust解决的事情根本没必要自己造轮子这一层就是“官方库使用说明书”。2.2 泛型算法、嵌入式与可视化样例的位置与取舍细看目录会发现几个泛型算法示例被分散在不同的层级0_Simple/reduce是基础归约6_Advanced里还有更复杂的归约变体3_Imaging里有卷积、滤波、直方图这些图像算法但它们的优化思路同样适用于通用数据处理。NVIDIA这么安排的用意很明显一方面让初学者在简单目录里就能接触核心思想另一方面在进阶目录里给出满足生产性能的版本。比如二维卷积simple版可能直接用全局内存加原子操作imaging版会刻意优化共享内存访问模式两者对比阅读你对“优化到底优化了什么”的认识会特别清晰。此外仓库里还有一批面向Jetson等嵌入式平台的示例通常以tegra或EGLStream命名。这类示例对纯服务器开发的人暂时用不上但对做边缘计算、机器人、自动驾驶的朋友就是宝贵参考。3. 源码级拆解三个经典sample读透CUDA核心机制接下来进入重头戏。我不打算把200多个示例全部过一遍那样既琐碎也没意义。挑三个我反复研究过、也是面试和实际项目中最高频出现的经典示例matrixMul、threadFenceReduction、histogram。把这三个读透你对CUDA的线程同步、共享内存、原子操作、归约设计这几个核心机制的理解会上一个台阶。3.1 matrixMul共享内存tiling与线程同步的教科书matrixMul对应的完整源码在0_Simple/matrixMul目录文件名叫matrixMul.cu。它实现的是经典的C A * B矩阵乘法朴素思路很直观每个线程负责计算输出矩阵的一个元素于是需要读取A的一整行和B的一整列。问题在于全局内存的访问开销很大一个线程读一遍整块矩阵要反复从显存搬运数据。优化的核心是共享内存分块也就是tiling。基本流程是定义一个BLOCK_SIZE比如16x16每个block负责计算输出矩阵中BLOCK_SIZE x BLOCK_SIZE的一个子块。计算这个子块时先协作把A的对应列切片和B的对应行切片加载到共享内存里然后__syncthreads()等待所有数据就位再继续乘加。下一次循环再把下一块数据加载进来。关键代码结构大致是__global__ void MatrixMulKernel(float* C, const float* A, const float* B, int wA, int wB) { __shared__ float As[BLOCK_SIZE][BLOCK_SIZE]; __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE]; int bx blockIdx.x, by blockIdx.y; int tx threadIdx.x, ty threadIdx.y; int aBegin wA * BLOCK_SIZE * by; int aEnd aBegin wA - 1; int aStep BLOCK_SIZE; int bBegin BLOCK_SIZE * bx; int bStep BLOCK_SIZE * wB; float Csub 0.0f; for (int a aBegin, b bBegin; a aEnd; a aStep, b bStep) { As[ty][tx] A[a wA * ty tx]; Bs[ty][tx] B[b wB * ty tx]; __syncthreads(); #pragma unroll for (int k 0; k BLOCK_SIZE; k) { Csub As[ty][k] * Bs[k][tx]; } __syncthreads(); } int c wB * BLOCK_SIZE * by BLOCK_SIZE * bx; C[c wB * ty tx] Csub; }读这段代码时有一个细节特别值得注意循环体内为什么在加载数据之后和计算完成之后各有一个__syncthreads()第一个同步是为了确保As和Bs里的所有数据都写完了才允许任何线程去读第二个同步是为了防止计算较快的线程提前进入下一轮循环把共享内存里的旧数据覆盖掉。少了任何一个结果都是偶发错误这种错误极其恶心因为它不是必现而是跟线程调度顺序有关。共享内存tiling为什么能提速从访存量来看朴素版本每个线程访问A一行中所有元素同一个A矩阵元素会被多个线程重复读取tiling之后每个元素在共享内存里只加载一次然后被block内多个线程复用。全局内存访问量降了一个量级性能差距自然明显。实测下来对较大矩阵这个tiling版本相比朴素版本能获得数倍到数量级的提升具体倍数跟矩阵规模相关。3.2 threadFenceReduction一次kernel里跨block归约的fence语义这个示例提供了很多开发者在刚写多block协作时犯迷糊的解法。通常归约需要把所有block的中间结果合到一起最简单的方案是启动两次kernel第一次每个block做局部归约第二次用单个block把所有局部结果收尾。threadFenceReduction却只启动一次kernel就完成了跨block归约关键靠两样东西共享内存里的block私有归约结果、全局内存里的计数器和__threadfence()。它的逻辑大致是每个block先把自己的最终结果写到全局内存数组的对应位置然后执行__threadfence()确保这条写入对其他block可见接着一个线程对全局计数器做原子加一得到返回值。如果返回值等于gridDim.x减一说明当前block是最后一个到达的block由它读取其他block写好的所有局部结果算出最终结果写回。这里的核心是__threadfence()的语义它保证调用线程在fence之前的所有全局内存写入对另一个线程在观察到某个同步标志之后执行的读取操作是可见的。没有这个fence最后一个block去读其他block结果时可能读到的是没来得及刷新到显存的数据。很多刚接触的人会问能不能用__syncthreads()做跨block同步不能。__syncthreads()只同步block内线程它没有任何机制保证一个block能看到另一个block的数据。threadFenceReduction里的计数器加原子操作实际上是在构建一个轻量的全局软件屏障。我从实际开发里得到的经验是这种单kernel归约写法能省一次kernel启动开销在kernel启动延迟敏感的短计算任务里有用但它的可扩展性和调试性不如两阶段kernel。真正做生产项目时我通常优先考虑两阶段方案只有在对延迟要求极其苛刻、并且block数量可控的场景下才用软件屏障。Samples的价值就在于它把这种“偏方”和标准做法摆在同一个仓库里让你在选择时心里有数。3.3 histogram原子操作与局部直方图的取舍histogram是图像处理和数据分析里极其常见的算子但它在GPU上优化起来有不少门道。样例位于3_Imaging/histogram展示了几种不同实现。最朴素的实现是每个线程处理一个像素直接对全局内存的直方图数组做atomicAdd。功能上完全正确性能却很糟因为所有线程在同时争抢少数几个计数器全局原子操作在显存级别上的冲突会让吞吐量崩塌。优化思路是对直方图做“私有化”privatization每个block维护自己的一份直方图放在共享内存里线程优先更新共享内存内的计数器因为共享内存的原子操作比全局内存快得多。等到整个block处理完数据再做一次合并把block的局部直方图加到全局结果里。合并过程可以用一个线程顺序做也可以用多个线程配合原子操作并行做。这个示例还揭示了两个权衡。第一是共享内存容量bin数量越多私有直方图占用的共享内存越大。假设直方图有256个bin每个bin用32位整数每block就需要1KB共享内存看起来不多但如果bin数量到4096或更高或者你想同时维护多个通道开销就很可观。第二是归并成本block数越多合并阶段的全局原子操作还是越多所以block粒度需要根据数据规模和硬件资源调。我后来在做自定义算子时把histogram里的私有化思想扩展到了很多“先统计、后归并”的算法里这类模式几乎可以成为GPU归约类问题的通用模板。4. 从Samples移植到工程错误处理、基准确认与可复用模式读源码和写工程之间还有一道鸿沟示例代码为了突出重点常会省略一些工程细节但它也留下了大量值得直接复用的骨架。下面这几个东西是我从Samples里拷贝到生产项目里最频繁的。4.1 CUDA_SAFE_CALL这类宏是Samples留给工程的第一笔财富CUDA API有一个容易踩的坑大量API是异步的kernel启动之后错误并不会立刻返回而是在某个同步点才会冒出来。如果你每次API调用后都不检查返回值遇到“莫名其妙的脏数据”时排查成本极高。Samples里经常出现的checkCudaErrors宏本质上是一个封装的错误检查器#define CUDA_CHECK(call) \ do { \ cudaError_t err (call); \ if (err ! cudaSuccess) { \ fprintf(stderr, CUDA error %s:%d: %s\n, __FILE__, __LINE__, \ cudaGetErrorString(err)); \ exit(EXIT_FAILURE); \ } \ } while (0)生产环境里你甚至可以把exit换成更友好的错误上抛或饥饿信号处理。但最核心的思想是一致的每一次cudaMalloc、cudaMemcpy、cudaLaunchKernel之后都要做检查这样才能把问题定位到具体调用点。这个习惯是真的能救命我调试过的最难缠问题里有相当一部分最后都归因于某次kernel启动后没有检查异步错误。4.2 每个sample都带CPU baseline性能验证才有据可依细读Samples时你会发现它几乎给每个优化示例都配了一个CPU参考实现。比如matrixMul在GPU kernel之外还写了一个CPU端的朴素乘法用于校验结果是否正确nbody示例里也保留了一个CPU端计算引力的参考函数。这个设计非常值得抄进自己的工程里。做GPU性能优化时只测GPU kernel的运行时间是不够的。你得知道三件事一、结果对不对二、比CPU快了多少三、比“朴素但正确”的GPU版本快了多少。没有CPU baseline你无法判断加速比没有朴素GPU版本对照你无法判断优化手段的实际收益。我现在的习惯是每个优化版本至少保留一个对应的正确性校验函数和一个朴素版本跑性能时同时输出三组数据。Samples自带的这套“参考实现计时结果比对”模板几乎是现成的测试框架。另外计时本身也有讲究。Samples里普遍使用cudaEvent来记录GPU时间这是因为GPU kernel的启动是异步的用clock()或者chrono测到的经常包含CPU端排队等待时间。正确做法是创建一个cudaEvent在kernel前后分别cudaEventRecord然后cudaEventSynchronize取cudaEventElapsedTime。Samples里到处是这样测时的例子照抄就行。4.3 直接能从Samples拷走的几个代码模式和习惯把整个仓库当“代码模板库”能直接拿走的东西比我预期的多deviceQuery初始化后自动获取设备属性、选择计算能力最高设备的模式适合放在任何GPU项目的启动阶段。checkCudaErrors错误处理宏上面已经讲了。CPU参考实现与GPU结果比对的模式适合作为单元测试的雏形。用cudaEvent计时的模板比用CPU时间戳可靠。矩阵乘法中通过宏控制BLOCK_SIZE、通过编译参数切换kernel版本的思路适合做性能调参实验。每个示例独立目录、自带Makefile、一条make就能跑的工程组织方式我把它迁移到了自己的实验仓库里效果很好。5. 编译、调试与性能核验13.3在真实环境中的操作链路拿到Samples只是第一步跑起来、改起来、验证起来才有价值。这一节记录我在真实环境里编译、调试、以及用性能工具核验示例时的一些操作链路和踩坑记录。5.1 跨平台构建make体系、VS工程与架构参数Linux/macOS下Samples的构建依赖一套统一的make体系顶层有Makefile子目录里也有简单的Makefile它们会调用common.mk里的公共逻辑。日常用到的基本命令make # 编译当前目录的示例输出到 bin/linux/release/ make dbg1 # 编译debug版本 make clean # 清理 make TARGET_ARCHx86_64 # 显式指定架构如果你是交叉编译比如在x86主机上构建Jetson设备用的示例可以指定类似make TARGET_ARCHaarch64这个过程中最容易踩的坑是“架构不匹配”你有一块Ada显卡但make默认生成的目标代码里没有包含适合你显卡的计算能力。Samples的Makefile一般会通过cuda_arch变量控制编译目标架构比如make cuda_archsm_89如果不确定自己的显卡计算能力先跑deviceQuery看输出或者用nvidia-smi --query-gpucompute_cap --formatcsv。算力不对后果有两种一种是你拿到的是一个JIT编译出来的PTX性能打折扣另一种是驱动太老不认识新的SASS直接报加载失败。Windows端相对轻松打开解决方案文件之后在工程属性里指定目标计算架构、CUDA Toolkit路径即可。但有几个sample依赖OpenGL、GLFW等图形库可能需要额外装了依赖才能编译过这类sample通常会在顶层README里标注编译失败时优先看依赖说明。5.2 Nsight Compute把sample的性能假设核验一遍源码里的注释说得再怎么天花乱坠都不如自己跑一次性能分析工具来得直接。CUDA Toolkit自带的Nsight Compute命令行工具叫ncu是我评测Samples时用得最多的工具。以matrixMul为例编译好之后执行ncu --set full ./matrixMul完成后会给出详细报告核心指标包括占用率Achieved Occupancy、共享内存吞吐、DRAM吞吐、指令数、分支发散、bank conflict等。我第一次跑的时候最关注的就是Shared Memory Bank Conflicts这一项。Samples默认的矩阵读取方式基本不会出现bank conflict但如果你把As[ty][k] * Bs[k][tx]这个访问模式改成按行优先存储的行向量操作就很容易制造出冲突。这里有一个通用经验不要只看一个指标。实际工作中我最常犯的错误是用“更高占用率”作为唯一目标结果发现shared memory或者寄存器压力被推高局部性反而变差kernel耗时并没有下降。NVIDIA在Nsight Compute里给出的memory workload分析和occupancy关联提示能帮你快速理解指标之间的耦合关系。Samples里每个优化示例都能作为这种“改参数、看指标”实验的起点。还有一个小技巧用ncu分析某个advanced目录下的示例之前先确认有没有其他进程占用GPU因为现代GPU是分时复用的被其他任务占用会导致测出来的指标漂移。跑分析时用CUDA_VISIBLE_DEVICES把进程限定到某一块空闲卡上能减少干扰。5.3 通用GPU运行时的坑显存、驱动版本与运行参数不同版本的Samples对CUDA运行时版本有最低要求。如果你的驱动版本较老编译可能能过但运行时会报symbol lookup error之类的错比如找不到cudaGraphExecKernelNodeSetParams这类新API。遇到这种情况第一选择是升级驱动第二选择是查看Samples的README确认它需要的CUDA版本。Samples 13.3整体比较新建议搭配较新的CUDA Toolkit使用。显存问题是另一个高频坑。有些示例比如粒子系统、体积渲染默认分辨率或数据规模会吃掉大量显存。如果机器上有多个进程占用了显存运行时会报内存不足。我习惯先跑nvidia-smi看看每块卡的显存和利用率再决定是否设置CUDA_VISIBLE_DEVICES。另外多卡机器的默认设备是0号卡但0号卡不一定是最适合跑计算的卡。生产机器上常有集显或低端卡占着0号位置直接用deviceQuery确认计算能力最强的设备ID然后在运行前设置环境变量指定export CUDA_VISIBLE_DEVICES1这个操作能避免你在不适配的设备上白跑很久。6. 把Samples当成起点而不是终点我的阅读路径与踩坑记录最后这部分是经验之谈。我在把Samples吃透的过程中踩了不少坑也总结出了一条比较高效的阅读路径。如果你准备开始读这套源码可以参考一下。6.1 适合不同目标人群的阅读顺序如果你是纯新手我的建议是这样走先跑0_Simple/vectorAdd搞懂grid/block/thread三级结构和cudaMemcpy的数据搬运逻辑再跑0_Simple/matrixMul理解共享内存和线程同步接着跑1_Utilities/deviceQuery和bandwidthTest了解你手头硬件的真实参数之后可以看3_Imaging/boxFilter或者5_Simulations/SmokeParticles体验一下稍微完整的算法在GPU上怎么组织。如果你已经在写CUDA但觉得性能上不去直接跳到6_Advanced。重点看cudaGraphs、streamOrderedAllocation、cooperativeGroups、concurrentKernels这几个目录。它们解决的是“kernel启动开销怎么降”“多个kernel怎么并行”“线程之间更灵活的协作方式”这些恰恰是生产性能瓶颈最常出现的地方。如果你做的是深度学习推理或数值计算优先看7_CUDALibraries里的cuBLAS、cuFFT、thrust示例。了解官方库的调用范式、内存布局要求比手写kernel更容易落地。6.2 我踩过的三个坑也顺手写在这里第一个坑是盲目改BLOCK_SIZE。我曾经把matrixMul里的BLOCK_SIZE从16改成32想当然地认为更大的block能提高并行度。实测发现性能反而下降用Nsight Compute一测才明白占用率受共享内存容量和寄存器数量限制block变大之后每个SM能驻留的block数减少最终占用率不升反降。Samples的默认值往往是针对大多数显卡调试过的平衡点改之前要带着假设去改而不是“感觉”去改。第二个坑是漏了__threadfence()。我基于threadFenceReduction自己写一个跨block归约时一开始只用了原子操作加一作为计数漏掉了fence导致偶尔计算结果错误。这种问题特别难复现调了一天才定位到。后来我的习惯是涉及跨block可见性时先画清楚“谁在何时写、谁在何时读”然后再决定同步原语绝不靠猜。第三个坑是私有直方图忘清零。在仿照histogram写一个自定义统计算子时我把每个block的私有直方图放到了全局内存里方便多个kernel复用结果忘了在每次kernel调用前清零第二次调用时结果直接叠加上一次的数据。示例里其实有类似的初始化流程我因为没细读这部分踩了坑。这提醒我从Samples里抄代码时不仅要抄核心逻辑还要把初始化、清理、边界条件这些“配角”一起抄过来。6.3 一个可以长期复用的源码阅读习惯GPU实验簿最后分享一个我保持了很久的习惯为每个想深入研究的sample建一个独立实验分支一次只改一个变量并用Nsight Compute记录固定的一组指标。比如对matrixMul我记录Duration、Achieved Occupancy、Shared Memory Bank Conflicts、DRAM Throughput、L1/L2 Cache Hit Rate。改完一个参数运行一遍把数据填进一个markdown表格。时间久了你会建立起非常扎实的“参数-性能数据”直觉。这个方法的核心价值在于它逼着你用数据而不是感觉做判断。Samples提供的就是一套干得很干净的实验素材你只需要在此基础上架一把刻度尺。长期积累下来这套“实验簿”会成为你自己的性能优化手册比任何官方文档都更适合你手头的硬件。