CUDA Samples 实战走查从跑通第一个 kernel 到看懂性能差距【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples你写过你好世界级别的向量加法看到屏幕上蹦出Test PASSED就收工了。但换个稍复杂的 kernel编译能过、运行不报错速度却只有理论值的零头——瓶颈在哪全靠自己猜。NVIDIA 官方的 CUDA Samples 就是给这种状态准备的一个按学习路径组织的 GPU 并行编程示例代码库从最基础的向量加法到张量核心、CUDA Graphs、CUDA Tile 这些新特性每个示例都带 README 说明涉及哪些 API。这篇走查不逐个目录念菜单挑最能打透的两个示例带你读代码。️ 十分钟跑通第一个 CUDA 示例前置条件只有一个装好 CUDA Toolkit比如 13.x 版本。装完先确认一件事nvcc --version能打印版本号就够了别的不用查。然后拉代码、构建git clone https://gitcode.com/GitHub_Trending/cu/cuda-samples.git cd cuda-samples mkdir build cd build cmake .. make -j$(nproc)仓库本身就是个 CMake 工程cmake ..会自动找到 CUDA 工具链。全量编译所有示例要几分钟你完全不用等它跑完——make vectorAdd -j单独构建这一个目标半分钟就出来。然后在build/bin/0_Introduction下执行./vectorAdd最后看到Test PASSED链路就算通了编译工具链、驱动、GPU、示例代码四样都没问题。把 vectorAdd 读成 GPU 编程的完整骨架 cpp/0_Introduction/vectorAdd/vectorAdd.cu 总共不到 200 行但把 GPU 计算的全流程走了一遍主机内存分配 →cudaMalloc设备内存 →cudaMemcpy拷入 → 启动 kernel → 拷回结果 → 逐元素校验。读懂它后面所有示例都是在这套骨架上加料。kernel 本体只有三行int i blockDim.x * blockIdx.x threadIdx.x; if (i numElements) { C[i] A[i] B[i] 0.0f; }第一行是 CUDA 编程里出现频率最高的公式每个线程算出自己的全局编号。第二行的边界检查不是摆设——元素数通常不是线程块大小的整数倍多出来的线程不做保护就会越界写内存。值得动手调的参数有两个。一个是每块线程数代码里是 256int threadsPerBlock 256; int blocksPerGrid (numElements threadsPerBlock - 1) / threadsPerBlock;第二行是标准的向上取整算块数写法值得抄进你以后的代码里。把 256 改成 512 或 128 再跑结果不变但启动配置变了——这正是新手理解网格维度怎么算最快的方式。同目录里还有几个变体值得扫一眼vectorAddDrv用 Driver API 实现同样的事vectorAddMMAP演示跨进程映射显存vectorAdd_nvrtc在运行时才编译 PTX。同一个加法三种编程模型对比着读比单看一个收获大。矩阵转置一行代码差距百倍速度差距如果说 vectorAdd 教你怎么跑cpp/6_Performance/transpose/transpose.cu 教你为什么慢。转置看起来毫无并行技巧含量就是out[y][x] in[x][y]但朴素 kernel 跑起来比优化版慢一到两个数量级原因全在内存访问模式上。朴素版本里一个线程块 32x32 的区域写输出时相邻线程的写入地址相隔 1024 个 float矩阵宽度。一行里连续 32 个线程实际只踩满了一个缓存行剩下 31 行全是浪费。优化版的做法是把 32x32 的 tile 先搬进共享内存__shared__ float tile[TILE_DIM][TILE_DIM 1];加载时按行读连续对齐缓存友好写出时按列写也连续对齐。注意数组第二维是TILE_DIM 1多出的那列是故意留的用来打散共享内存 bank 冲突——转置场景里不加这一列不同线程会撞进同一个 bank。这个文件里排着约六个 kerneltransposeNaive、带共享内存的版本、带 block swizzle 的版本各自测时输出。把几个数字并排看一遍你对缓存对齐四个字的理解会比读十页文档深。它所在的 cpp/6_Performance/ 目录还有cudaGraphsPerfScaling、alignedTypes等几个专门对性能做实验的示例。踩坑实录三个新手最常碰壁的地方第一次cmake ..就报错九成是nvcc不在 PATH 里。Toolkit 装完后bin目录默认不一定进环境变量配好之后重开终端再试——旧 shell 里的 PATH 不会自己刷新很多人在这一步空耗一小时。能编译、一运行就CUDA_ERROR_NO_BINARY_FOR_GPU是架构没对上。代码是按某个 compute capability 编译的你的 GPU 比它新或旧。重新cmake时带上目标架构比如-DCMAKE_CUDA_ARCHITECTURES89或者用native让工具链自动探测。第三个坑比较隐蔽部分示例编译不出来但其实没报错——仓库里的 CMake 对 FreeImage、FreeGLUT、Vulkan 这类可选依赖是探测式的缺库就静默跳过对应示例。所以我明明装了为什么 bin 目录里没有那个程序多半是某个可选库没装翻一下build阶段的 configure 输出能找到线索。另外两个实用开关想上 cuda-gdb 做设备端断点configure 时加-DENABLE_CUDA_DEBUGTrue代价是编译器优化被关掉别拿它测性能想批量回归验证仓库根目录有 run_tests.py配test_args.json可以一次跑完所有已构建的示例。⚡ 什么时候该开始抠性能数据量还停在十万级的时候别优化——拷贝时间占大头瓶颈在cudaMemcpy而不是 kernel。真正该动手的信号是kernel 耗时稳定了但数字远低于硬件峰值带宽。这时候看 transpose 里的测时套路就行它用 CUDA Event 记录、重复NUM_REPS100 次取平均把一次启动的噪声平均掉。你测自己的 kernel 时也照这个来单次测量的数字没有参考价值。工具链这边构建层有 CMake 开关比如前文提到的ENABLE_CUDA_DEBUG分析层直接用 Nsight Systems 抓 timelinevectorAdd 这种小例子里你会直观看到 kernel 只有几微秒前面那两段拷贝和 launch 开销反而更长——小数据量下GPU 计算很快是错觉开销全在搬运和启动上。想直接看库级性能差距的话cpp/4_CUDA_Libraries/nvJPEG 这类示例会同时给出 GPU 解码结果和 CPU 参考对比按你的方向挑路走图像处理从 cpp/5_Domain_Specific/ 下手。起点建议SobelFilter边缘检测全链路和bilateralFilter保边平滑图像数据文件都放在示例目录里改参数跑一遍就能看到效果差异。科学计算看 cpp/4_CUDA_Libraries/conjugateGradient同一个共轭梯度求解器有普通版、UM 统一内存版、多设备版三套实现同一算法三种写法对比着读想理解数值变换本身cpp/2_Concepts_and_Techniques/dct8x8/ 把 8x8 DCT 拆成多个 kernel 逐级实现。AI 推理与矩阵计算重点扫 cpp/3_CUDA_Features/ 里的张量核心系列——cudaTensorCoreGemm、bf16TensorCoreGemm以及最新的 cpp/9_CUDA_Tile/ 目录tileMatmul、tileBmm这是面向新一代 tile 编程模型的官方示例。Python 用户python/ 目录从1_GettingStarted/vectorAdd到2_CoreConcepts/reduction一条线走完核心概念3_FrameworkInterop/里还有直接写自定义 PyTorch kernel 的示例。今天就能做的一件事打开 transpose 源码把TILE_DIM从 32 改成 64重新编译跑一遍对比输出里各 kernel 的耗时——你会亲眼看到 tile 大小和 bank 冲突之间的博弈这比任何教程段落都直接。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考