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 编程里出现频率最高的公式：每个线程算出自己的全局编号。第二行的边界检查不是摆设——元素数通常不是线程块大小的整数倍，多出来的线程不做保护就会越界写内存。
值得动手调的参数有两个。一个是每块线程数，代码里是 256：
int 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。
这个文件里排着约六个 kernel：transposeNaive、带共享内存的版本、带 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_ARCHITECTURES=89，或者用native让工具链自动探测。
第三个坑比较隐蔽：部分示例编译不出来，但其实没报错——仓库里的 CMake 对 FreeImage、FreeGLUT、Vulkan 这类可选依赖是探测式的，缺库就静默跳过对应示例。所以"我明明装了为什么 bin 目录里没有那个程序"，多半是某个可选库没装，翻一下build阶段的 configure 输出能找到线索。
另外两个实用开关：想上 cuda-gdb 做设备端断点，configure 时加-DENABLE_CUDA_DEBUG=True（代价是编译器优化被关掉，别拿它测性能）；想批量回归验证，仓库根目录有 run_tests.py，配test_args.json可以一次跑完所有已构建的示例。
⚡ 什么时候该开始抠性能
数据量还停在十万级的时候，别优化——拷贝时间占大头，瓶颈在cudaMemcpy而不是 kernel。真正该动手的信号是：kernel 耗时稳定了，但数字远低于硬件峰值带宽。这时候看 transpose 里的测时套路就行：它用 CUDA Event 记录、重复NUM_REPS（100 次）取平均，把一次启动的噪声平均掉。你测自己的 kernel 时也照这个来，单次测量的数字没有参考价值。
工具链这边，构建层有 CMake 开关（比如前文提到的ENABLE_CUDA_DEBUG），分析层直接用 Nsight Systems 抓 timeline：vectorAdd 这种小例子里你会直观看到 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），仅供参考
