CUDA Samples实战指南:从零跑通你的第一个GPU示例
【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples
NVIDIA CUDA Samples是NVIDIA官方维护的CUDA示例代码库,配套CUDA Toolkit 13.3,用C++和Python演示GPU编程的完整套路。它适合没写过CUDA、想系统入门并行计算的新手,也适合想查某个API怎么用、怎么优化的开发者。本文带你三步跑起来,并规划好接下来的学习路径。
🧩 这到底是干嘛的
上一节说了项目背景,这里先讲清楚它能帮你解决什么问题。
你可以把CUDA Samples理解成一个"带答案的习题册":每道题都是一个可运行的程序,题面写在README里,解题步骤就是源码。它适合三类人:刚学GPU编程的新手、需要查API用法的工程师、做性能优化的老手。
核心能力有五条:
- 分层教学:从
0_Introduction到9_CUDA_Tile共10个层级,由浅入深覆盖基础、技巧、特性、库、领域应用。 - 真实可运行:每个示例都有
CMakeLists.txt,配好CUDA Toolkit就能编译运行,不是伪代码。 - 覆盖主流场景:向量加法、矩阵乘法、直方图、FFT、图像滤波、科学计算、深度学习算子都有对应样例。
- 双语言支持:C++示例之外,
python/目录提供基于cuda.core的Python示例,门槛更低。 - 可交叉编译:
cmake/toolchains/里带Tegra等嵌入式平台的工具链文件,能编译到板子上跑。
为什么值得学?因为它的示例结构直接对应CUDA编程的真实工作流:分配显存、拷数据、启动内核、校验结果。把前三个目录的示例读一遍,你就掌握了这个套路。
🚀 三步跑起来
环境搭好后,三步就能编译出可运行的程序。
环境配置一步到位
- 准备硬件与驱动:一台带支持CUDA的NVIDIA GPU的机器,装好匹配的显卡驱动。
- 安装CUDA Toolkit:从NVIDIA官网下载对应平台的版本,当前示例库配套13.3版本。装完验证:
nvcc --version能看到版本号即可。 - 安装CMake:需要3.20及以上版本,Linux下执行
sudo apt install cmake。
构建命令三连
进入仓库根目录,执行下面三条命令:
mkdir build && cd build cmake .. make -j$(nproc)第二条会扫描所有CMakeLists.txt并检查CUDA环境,第三条并行编译全部示例。Windows下用VS的"x64 Native Tools Command Prompt"执行同样的步骤即可。
验证版本的方法
编译成功后,示例会被安装到build/bin/x64/linux/release(默认路径按架构和系统自动组织)。直接运行一个:
./bin/x64/linux/release/cpp/0_Introduction/vectorAdd/vectorAdd看到[Vector addition of 50000 elements]且输出Test PASSED,说明工具链、驱动、编译、运行全部打通。如果只想构建单个示例,也可以进对应子目录单独执行cmake,速度更快。
提示:1_Utilities里的deviceQuery是官方推荐的第二站,它能打印你GPU的显存、计算能力和架构编号,是排查环境问题的首选工具。
📂 目录布局速览
跑通之后,你大概率想知道"接下来看什么"。目录结构就是学习地图。
- cpp/0_Introduction/:入门层,vectorAdd、matrixMul、simplePrintf等,先把基本流程跑熟。
- cpp/1_Utilities/:工具层,deviceQuery、topologyQuery,查硬件信息用。
- cpp/2_Concepts_and_Techniques/:概念技巧层,直方图、reduction归约、扫描、排序网络等核心算法。
- cpp/3_CUDA_Features/:特性层,CUDA Graphs、张量核心(Tensor Core,GPU里专门加速矩阵乘的硬件单元)、PTX JIT等。
- cpp/4_CUDA_Libraries/:库层,cuBLAS、cuFFT、NPP、nvJPEG等NVIDIA库的用法。
- cpp/5_Domain_Specific/:领域应用层,图像处理、金融期权定价、粒子系统、流体模拟。
- cpp/6_Performance/:性能层,转置优化、统一内存性能、CUDA Graphs扩展。
- cpp/7_libNVVM/与cpp/8_Platform_Specific/:NVVM中间表示示例,以及Tegra嵌入式平台专属示例。
- cpp/9_CUDA_Tile/:Tile编程模型的新特性示例。
- python/:Python版示例,结构与C++侧对齐,按GettingStarted、CoreConcepts等分组。
- Common/:所有C++示例共享的头文件与工具代码,如helper_cuda.h。
推荐阅读顺序:0_Introduction→1_Utilities→2_Concepts_and_Techniques(重点看vectorAdd、matrixMul、reduction、histogram)→ 按你的兴趣方向选3~6。
上图是2_Concepts_and_Techniques/dct8x8示例附带的DCT基函数图。DCT(离散余弦变换)是JPEG压缩的核心算法,这个示例演示了用GPU并行做DCT的完整过程,可以作为"算法上GPU"的模板参考。
📖 示例精讲
知道看哪里了,下面挑三个最典型的示例,讲透它们的思路。
vectorAdd:GPU编程的最小完整流程
解决什么问题:把两个向量逐元素相加,C[i] = A[i] + B[i]。看似简单,但它包含了GPU程序的完整生命周期。
核心思路:主机malloc分配数组 →cudaMalloc在显存分配对应空间 →cudaMemcpy把数据传上去 → 启动内核(每个线程算一个元素,线程编号i = blockDim.x * blockIdx.x + threadIdx.x)→ 结果拷回主机 → CPU上校验正确性。源码在 vectorAdd.cu,内核本身不到10行。
能迁移到什么场景:任何"逐元素独立"的计算都能套用,比如图像调色、数组过滤、张量逐点激活函数。你以后写的所有CUDA程序,骨架都和它一样。
matrixMul:共享内存的第一个实战
解决什么问题:矩阵乘法,线性代数的基本操作,也是深度学习的主力运算。
核心思路:这个示例刻意不追求极致性能,而是为了"讲清楚"。每个线程块把要用的A、B小块先搬到共享内存(GPU上类似CPU寄存器缓存的高速片上内存),再反复读取,大幅减少显存访问次数。同一目录还附带cuBLAS版本,对比手写内核和官方库的差距。
能迁移到什么场景:卷积、注意力计算、图上的矩阵运算。读懂它的分块(tiling)策略,你就掌握了CUDA性能优化的第一块基石。
reduction:归约模式与优化手段全演示
解决什么问题:把一个大数组求和,这是统计、采样、注意力softmax里到处都有的操作。
核心思路:朴素写法是每轮把数组长度减半。示例展示了四级优化:用共享内存组织块内归约、用__shfl_down_sync指令让线程间直接交换数据(省掉共享内存往返)、用__reduce_add_sync硬件归约指令、用cooperative groups的reduce写法。源码见 reduction_kernel.cu。
能迁移到什么场景:求最大最小值、概率统计、树形结构汇总。它是学习"同一个算法如何用不同指令实现"的最佳样本。
🔭 再往深一点
基础通了之后,这些方向可以按兴趣深挖。
- CUDA Graphs:
3_CUDA_Features/simpleCudaGraphs和jacobiCudaGraphs演示把多次内核启动打包成一张图提交,降低CPU调度开销。 - 张量核心:
cudaTensorCoreGemm、bf16TensorCoreGemm展示了用mma指令直接调用Tensor Core,是写自定义GEMM算子的起点。 - 纹理与表面内存:
simpleTexture、bindlessTexture讲解2D过滤采样,做图像处理绕不开。 - 多GPU与P2P:
simpleMultiGPU、streamOrderedAllocationP2P演示跨卡通信,做分布式推理必读。 - 动态并行与CDP:
cdpQuadtree等示例展示在GPU上从内核里再启动内核的递归场景。 - Python侧:
python/2_CoreConcepts/parallelReduction和tmaTensorMap用Python写同样内容,方便快速验证想法。
⚡ 怎么跑得更快
这些优化清单来自示例里的通用做法,写自己的程序时直接对照。
- 合并内存访问:相邻线程访问相邻地址,让显存事务一次拉满。vectorAdd默认就是按这个模式写的。
- 共享内存分块:matrixMul的分块策略,减少全局内存往返次数。
- 用shuffle和硬件归约指令:reduction示例里
__shfl_down_sync和__reduce_add_sync比共享内存方案更快。 - 异步拷贝与流重叠:
simpleStreams、streamingCopyComputeOverlap演示传输和计算并行,避免GPU空转。 - 用CUDA Graphs减少启动开销:小内核反复调用时,打包成图提交,CPU开销显著下降。
- 选对数据类型:fp16、bf16在Tensor Core上吞吐远高于fp32,见
fp16ScalarProduct和bf16 GEMM示例。 - 先测量再优化:用Nsight Systems看时间线、cuda-gdb查内核(构建时加
-DENABLE_CUDA_DEBUG=True开启),别凭感觉调参。
🛠 出错了怎么办
编译和运行CUDA示例最常见的坑,按现象找原因。
- 现象:
cmake ..报找不到CUDA Toolkit。- 原因:nvcc不在PATH,或Toolkit没装全。
- 解法:先跑
nvcc --version确认;不行就设置CMAKE_CUDA_COMPILER指向具体路径,如cmake -DCMAKE_CUDA_COMPILER=/usr/local/cuda/bin/nvcc ..。
- 现象:运行时提示没有可用的CUDA设备。
- 原因:驱动版本低于CUDA要求,或机器上根本没有NVIDIA GPU。
- 解法:更新驱动到与Toolkit匹配的最低版本;先跑
deviceQuery确认GPU能被识别。
- 现象:
make阶段部分示例编译失败,报Vulkan、libdrm缺文件。- 原因:平台专属依赖(如simpleGL、simpleVulkan)没装。
- 解法:这些示例可跳过,README里给出注释掉对应
add_subdirectory或用make --keep-going忽略错误继续构建其余示例的方法。
- 现象:程序跑起来了但结果校验失败。
- 原因:GPU架构编号不在编译目标列表里,内核未真正执行。
- 解法:查看根目录CMakeLists.txt中的
CMAKE_CUDA_ARCHITECTURES,确认包含你GPU对应的编号(deviceQuery里可查)。
- 现象:显存不足(out of memory)退出。
- 原因:示例默认数据规模超过显存。
- 解法:改小程序开头的
numElements或矩阵维度常量;大任务改用分块或流式处理。
🎯 学完能做什么
按下面两条路线走,效率最高。
初学者路线(2~4周):第1周跑通vectorAdd并逐行读懂,跑一遍deviceQuery认识自己的GPU;第2周精读matrixMul和reduction,把共享内存、归约两个模式吃透;第3周做histogram和scan,体会"共享内存+原子操作+同步"的组合拳;第4周挑一个5_Domain_Specific里的完整应用(如SobelFilter或BlackScholes)独立编译运行,模仿其结构写自己的小项目。
进阶方向:性能方向看6_Performance加Nsight Systems做profiling;深度学习方向看张量核心GEMM系列和python/3_FrameworkInterop里PyTorch、TensorFlow的自定义内核;系统方向看多GPU、IPC、CUDA Graphs;嵌入式方向用cmake/toolchains/里的工具链把示例交叉编译到Tegra。
上图是5_Domain_Specific/bilateralFilter示例的处理结果,双边滤波能在平滑噪声的同时保住边缘,这类示例直接对应工业界的图像管线需求。
CUDA Samples把GPU编程从零到优化的每一步都变成了可运行、可对照的代码。建议你今晚就做一件事:装好环境,把vectorAdd跑起来,明天开始逐行读它的源码。两周后,你就能独立写出并优化自己的第一个CUDA程序了。
【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples
创作声明:本文部分内容由AI辅助生成(AIGC),仅供参考