CUDA Samples实战指南:从零跑通你的第一个GPU示例
2026/9/18 3:04:20 网站建设 项目流程

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_Introduction9_CUDA_Tile共10个层级,由浅入深覆盖基础、技巧、特性、库、领域应用。
  • 真实可运行:每个示例都有CMakeLists.txt,配好CUDA Toolkit就能编译运行,不是伪代码。
  • 覆盖主流场景:向量加法、矩阵乘法、直方图、FFT、图像滤波、科学计算、深度学习算子都有对应样例。
  • 双语言支持:C++示例之外,python/目录提供基于cuda.core的Python示例,门槛更低。
  • 可交叉编译cmake/toolchains/里带Tegra等嵌入式平台的工具链文件,能编译到板子上跑。

为什么值得学?因为它的示例结构直接对应CUDA编程的真实工作流:分配显存、拷数据、启动内核、校验结果。把前三个目录的示例读一遍,你就掌握了这个套路。

🚀 三步跑起来

环境搭好后,三步就能编译出可运行的程序。

环境配置一步到位

  1. 准备硬件与驱动:一台带支持CUDA的NVIDIA GPU的机器,装好匹配的显卡驱动。
  2. 安装CUDA Toolkit:从NVIDIA官网下载对应平台的版本,当前示例库配套13.3版本。装完验证:nvcc --version能看到版本号即可。
  3. 安装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_Introduction1_Utilities2_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 Graphs3_CUDA_Features/simpleCudaGraphsjacobiCudaGraphs演示把多次内核启动打包成一张图提交,降低CPU调度开销。
  • 张量核心cudaTensorCoreGemmbf16TensorCoreGemm展示了用mma指令直接调用Tensor Core,是写自定义GEMM算子的起点。
  • 纹理与表面内存simpleTexturebindlessTexture讲解2D过滤采样,做图像处理绕不开。
  • 多GPU与P2PsimpleMultiGPUstreamOrderedAllocationP2P演示跨卡通信,做分布式推理必读。
  • 动态并行与CDPcdpQuadtree等示例展示在GPU上从内核里再启动内核的递归场景。
  • Python侧python/2_CoreConcepts/parallelReductiontmaTensorMap用Python写同样内容,方便快速验证想法。

⚡ 怎么跑得更快

这些优化清单来自示例里的通用做法,写自己的程序时直接对照。

  • 合并内存访问:相邻线程访问相邻地址,让显存事务一次拉满。vectorAdd默认就是按这个模式写的。
  • 共享内存分块:matrixMul的分块策略,减少全局内存往返次数。
  • 用shuffle和硬件归约指令:reduction示例里__shfl_down_sync__reduce_add_sync比共享内存方案更快。
  • 异步拷贝与流重叠simpleStreamsstreamingCopyComputeOverlap演示传输和计算并行,避免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),仅供参考

需要专业的网站建设服务?

联系我们获取免费的网站建设咨询和方案报价,让我们帮助您实现业务目标

立即咨询