简介《CUDA BY EXAMPLE》配套源码包是围绕GPU高性能编程经典著作整理的全套示例实现适合刚接触CUDA C/C的开发者、科研人员及高校学生对照学习。压缩包共57个文件包含36个.cu源文件、9个.h头文件、4个.lib库和4个.dll动态库等整体仅443KB轻量便携、解压即用。代码覆盖书中第3章至第11章及附录A的核心示例涉及线程块与网格组织、全局内存与共享内存管理、原子操作、多数据流并发、多GPU协作、OpenGL互操作以及性能分析等关键知识点每个示例与书中章节一一对应可帮助读者在真实可运行代码中理解并行计算的设计与调试方法。包内另附README、许可说明及自我介绍等文档方便快速确认环境与使用方式。目前已有410人学习浏览是深入掌握CUDA编程并提升GPU计算效率的实用参考。 《CUDA BY EXAMPLE》中文版叫《GPU高性能编程CUDA实战》我前前后后刷了三遍书里配套的所有例子源码直到现在还在我硬盘上躺着。这篇东西不是书评而是把我整理源码、编译运行、踩坑排查的全过程记录一遍希望能帮到同样拿这本书入门GPU高性能编程的人。这本书是Jason Sanders和Edward Kandrot写的成书时间很早对应的CUDA版本大概在1.x到2.x但读下来你会发现一个很扎心的事实CUDA的生态已经膨胀到难以置信可核心编程模型几乎没变依然是线程层次、内存模型、kernel启动、共享内存、原子操作这些老底子。所以你完全不用担心“书太老代码过时”。书里所有例子源码只要能跑起来你对GPU编程的理解会比看十篇教程都扎实。1. 这本书和配套源码为什么现在还要折腾它1.1 老书不老CUDA编程模型二十年没变很多人一看到“CUDA By Example”这种老书就下意识觉得没必要看这恰恰是误区。现在很多做深度学习、AI推理的人用PyTorch、TensorFlowGPU的调度被框架封装得严严实实甚至连__global__都不用写。可一旦你遇到显存优化、自定义算子、多流并发、模型推理延迟这类问题不懂底层线程模型就是抓瞎。书里的例子源码走的是“最朴素路线”从最简单的hello_world开始到矢量求和、Julia集、点积、波纹、直方图、多GPU协同每一个例子都只围绕一个知识点展开。没有花哨的库没有工程化包袱反而非常适合用来理解GPU到底是怎么把数据拆开、并行、再合并的。我推荐这本书的另一个原因是它的“例子导向”风格。第3章的hello_world会教你写第一个kernel第5章开始引入线程协作后面再慢慢讲性能、原子操作、流和多GPU。这种节奏和现代GPU编程的进阶路径完全一致源码本身就是最好的课程大纲。1.2 源码从哪拿原版目录和GitHub备份这本书原书配套的代码下载链接如今已经不太好用了但GitHub上搜“cuda-by-example”能拿到大量备份仓库。需要注意的是一部分仓库是原封不动的老代码另一部分是好心人改造过、能在新版CUDA上编译通过的版本。我建议两个都下载老代码用来理解原始写法改造版用来快速跑通。原版源码的目录结构一般按章节组织核心文件包括hello_world.cu第一个CUDA程序sum.cuCPU与GPU矢量求和的对比julia.cuJulia集并行计算dot.cu共享内存实现点积ripple.cu共享内存的双缓冲动画histogram.cu原子操作统计直方图thread_fence.cu线程栅栏示例simple_texture.cu纹理内存示例拿到这些源码后不要急着编译先把每个文件打开看一眼主函数和kernel函数的对应关系。我的经验是这几十个例子里真正值得精读的不会超过10个剩下的要么是前面的变体要么是为特定硬件写的花活。2. 环境搭建把书里例子跑起来的完整流程2.1 显卡、驱动和CUDA Toolkit的关系跑CUDA代码你需要搞清楚三件事NVIDIA显卡驱动、CUDA Toolkit、显卡算力Compute Capability。驱动是显卡跑CUDA的基础Toolkit是编译和运行的工具集算力决定了你该用哪个-arch参数。可以简单类比驱动是高速公路Toolkit是车算力是这条路上的限速牌。安装时有个非常关键的原则显卡驱动版本要足够新Toolkit版本和驱动版本匹配。比如你的机器装了最新的NVIDIA驱动那装CUDA 12.x基本没问题但如果驱动很老强行装新Toolkit运行时大概率报cudaErrorInsufficientDriver。老书里的源码对Toolkit版本没那么挑反而是太老的代码配上太新的编译器会出事这点后面细说。安装完Toolkit后打开终端验证一下nvcc --version能看到版本号说明编译器就位。再跑一下nvidia-smi确认你的显卡驱动正常。如果这里已经报错别急着看CUDA先把驱动问题解决。2.2 查看显卡算力并确定编译参数每个例子编译时都会遇到一个绕不开的参数-archsm_XX。这里的sm表示流多处理器架构XX是算力编号比如RTX 3090是sm_86RTX 4090是sm_89A100是sm_80。查看你显卡算力的最简单方式nvidia-smi --query-gpucompute_cap --formatcsv会输出类似8.6这样的结果对应sm_86。编译时就写nvcc -archsm_86 -o hello_world hello_world.cu如果不确定自己的算力或者代码要在多张不同卡上跑建议用这种稳一点的写法nvcc -gencode archcompute_86,codesm_86 -o hello_world hello_world.cu很多老代码自带的Makefile里写的是-archcompute_10这在新版CUDA上已经不被支持必须改成你实际显卡的算力。这一步是最容易踩的第一坑。2.3 老代码依赖的旧库怎么处理书里不少例子在作者写书时用了一个叫cutil的辅助库用来检查CUDA错误、计时、加载纹理等。这个库在现代CUDA Toolkit里没有所以直接编译会报cutil.h: No such file or directory。解决办法有两个。一是用GitHub上已经移除cutil依赖的改造版源码省心二是手把手把cutil调用替换成现代API。我推荐初学者直接选前者等理解代码逻辑后再回头学怎么替换。替换的时候主要做两件事把CUT_CHECK_ERROR(err)这种宏换成cudaGetLastError()判断把cutilDeviceInit(argc, argv)这种初始化函数删掉用cudaSetDevice(0)代替。这两个替换做完90%的旧例子都能在新环境里编译通过。还有一些例子用到了OpenGL和GLUT做可视化比如波纹效果你需要先安装freeglut开发库sudo apt install freeglut3-dev编译时再手动加上-lglut -lGL。这块没有统一答案只能看具体例子。3. 核心例子源码拆解从hello_world到多GPU3.1 第一个kernel搞清楚在干嘛书里的hello_world例子虽然简单但它把CUDA最核心的调用方式演示得明明白白。代码大致长这样__global__ void kernel(void) { printf(Hello, World!\n); } int main(void) { kernel1, 1(); cudaDeviceSynchronize(); return 0; }关键在于1, 1它代表启动kernel时的配置前一个数字是grid里的线程块数量后一个数字是每个线程块里的线程数量。这些方括号里的参数直接决定了GPU上会开多少线程。很多人看完这个例子就过了但我建议你亲手改一改把1, 1改成2, 2看看printf输出几次改成10, 10再观察线程执行顺序。这种“破坏性实验”比背定义管用得多。你会在输出里发现线程顺序是乱序的这正是GPU并行的底层真相你永远不要依赖线程的启动顺序只能依赖显式同步。3.2 矢量求和明白线程和数据的映射关系矢量求和是CUDA入门的“hello world plus”它的核心不是求和的数学而是搞清楚每个线程负责哪个数据。书里的经典做法是int index blockIdx.x * blockDim.x threadIdx.x;这个公式值得你用笔推一遍。blockIdx.x是当前线程块在grid里的编号blockDim.x是块的线程数threadIdx.x是线程在块内的编号。三者组合就能让每个线程拿到自己的全局唯一ID然后去处理数组的第index个元素。一旦这个映射关系想明白后面几乎所有数据并行算法都通了。图像处理、矩阵运算、物理模拟本质都是在做“把一个大数组切成很多份每个线程负责其中一份”这件事。现在OpenCV、PyTorch里很多GPU算子底层也都是这个公式的变体。3.3 点积和共享内存归约操作的原型书里点积dot这个例子我特别喜欢它是所有归约操作的原型。所谓归约就是一大堆数最后归成一个数求和、求最大值、求平均值都是归约。简单做法是每个线程算一部分和再把所有部分和加到一起。但如果这部分和都写回全局内存再读出来效率太低。书里的做法是每个线程块先用共享内存shared缓存部分和__shared__ float partialSum[256];然后通过__syncthreads()等待块内所有线程都完成写入再对共享内存里的值做进一步归约。共享内存是片上资源访问速度远快于全局内存这是GPU性能优化最重要的一课。我在改这个例子时踩过一个小坑__syncthreads()不是随便放的如果某个线程提前return或者分支不一致会导致块内线程等待死锁。书里的例子写法很规范但你自己改造时一定注意别让线程在同步前“跑偏”。3.4 原子操作与多GPU现代大模型推理的底子很多人在学老书时对原子操作和多GPU章节不以为然觉得现在用不上。实际上原子操作在深度学习框架的显存分配、统计归约、量化计算中非常常见。atomicAdd这个函数直到今天依然是CUDA代码里高频出现的东西。书里的直方图例子用原子操作统计像素分布逻辑上很好理解每个线程算出一个结果然后atomicAdd到全局计数器。但要注意原子操作在高并发下会有性能损耗书里的例子可以跑通但实际工程里优化空间很大。多GPU的示例则是通过cudaSetDevice在多个设备间切换现代框架里的“多卡训练”本质上还是这套东西。我把书里的核心例子整理成了一张表方便后面复习的时候对照示例核心知识点对现代开发的参考价值hello_worldkernel启动、线程配置理解最基础的执行模型矢量求和线程与数据映射数据并行的通用套路Julia集网格化并行计算图像/像素级并行点积共享内存、归约深度学习中的Reduce算子波纹双缓冲、同步动画/实时渲染优化直方图原子操作统计、资源分配多GPU示例设备选择、通信多卡训练/推理的雏形4. 实操过程我把每个例子重新编译后的现场记录4.1 官方源码的常见编译错误与修复我使用CUDA 12.3环境显卡是Ampere架构。直接拿某GitHub仓库里的老代码编译第一波就炸了十几个错。最常见的几个问题我记在下面。首先是cutil.h找不到解决方法前面说过找改造版或者替换。其次是nvcc报Unsupported gpu architecture compute_10这个必须去Makefile里把架构参数改成sm_86。还有代码里用了__mul24这类老内置函数在较新版本已被移除改成普通乘法就行。比较隐蔽的一个错误是printf格式问题。老代码里把size_t或者long long直接当成%d打印编译器有时不报错但输出是错的。遇到这种情况用%zu、%lld这类正确格式替换。我修代码的顺序是先找一个能编译的干净例子比如hello_world确认环境没问题再逐个挑战带依赖的例子。千万不要一上来就编译带OpenGL的波纹示例否则你会被链接错误劝退。4.2 运行结果与性能观察所有例子编译通过后我建议把矢量求和那个例子稍微改造一下用来观察GPU并行效果。原书例子里的数组规模比较小GPU的优势根本体现不出来甚至CPU更快。把数组规模调到一千万以上再用clock()计时你才能明显看到差距。我在实测中比较了CPU单线程求和和GPU kernel求和。CPU跑一千万个float大概需要几十毫秒GPU第一次运行要经历JIT编译也差不多这个量级但连续跑第二次就降到几毫秒了。这给了两个启发数据量不够大时GPU并不一定更快数据传输和启动开销很大首次调用kernel会有JIT编译和初始化开销评测性能时不要用第一次的时间。4.3 双显卡和WSL环境的坑我现在开发环境有一部分在WSL2里跑书里例子时遇到过一个典型问题Failed to initialize NVML: GPU access blocked by the operating system。这个问题通常和Windows侧显卡驱动、WSL的GPU挂载权限有关。排查思路分几步。先看看Windows侧nvidia-smi是否正常再看WSL里nvidia-smi是否可用。如果WSL里nvidia-smi正常但CUDA程序起不来检查一下LD_LIBRARY_PATH是否包含/usr/lib/wsl/lib有时候CUDA Toolkit会覆盖这个关键路径。笔记本双显卡同样需要注意。独显和核显同时存在时CUDA默认不一定会选中NVIDIA独显。你可以在程序里显式指定cudaSetDevice(0);或者在运行时用环境变量控制CUDA_VISIBLE_DEVICES0 ./dot这样能减少很多莫名其妙的问题。书里那些老例子没有考虑这种现代笔记本场景你必须自己动手适配。5. 常见问题与排查技巧速查表5.1 编译期问题速查报错信息原因解决办法No such file or directory: cutil.h依赖旧辅助库删除依赖或用改造版源码Unsupported gpu architecture compute_10编译参数太老改为当前显卡算力identifier __mul24 is undefined老内置函数被移除替换为普通乘法cannot find -lglut -lGL缺少可视化库安装freeglut3-deverror: calling a __host__ function from a __global__ function is not allowed在kernel里调用了CPU函数把相关逻辑改为GPU设备函数5.2 运行期问题速查报错信息原因解决办法No kernel image is available for execution on the device编译架构与显卡不匹配用-archsm_XX匹配实际算力cudaErrorInsufficientDriver驱动版本太旧更新NVIDIA驱动Failed to initialize NVMLWSL或驱动环境异常检查WSL和驱动挂载设置LD_LIBRARY_PATHCUDA error: invalid device ordinal显卡编号越界用cudaGetDeviceCount查看设备数量后调整5.3 调试CUDA代码的几点心得编译和运行只是第一步想真正搞懂源码你还需要学会调试和验证。老书年代的cuda-gdb和现在的版本差别很大但核心思路不变先确认数据从CPU到GPU没问题再确认kernel内部计算正确最后确认结果从GPU回传正确。我自己调试的时候特别依赖两个工具cuda-memcheck新版本里是compute-sanitizer用来查越界和未初始化内存Nsight Systems用来分析kernel耗时和内存传输占比。书里老代码没有性能分析的习惯但你跑通之后一定要自己补上这一步。拿dot例子来说如果输入数组只有几千个元素你会看到kernel本身只占很小比例时间大部分时间浪费在cudaMemcpy上这种观察比单纯把代码跑通更有价值。6. 源码不是用来抄的一点实际操作体会书里所有例子源码我都过了一遍最深的体会是代码可以老思想不能老。老代码里的线程映射、共享内存、原子操作这些概念放在今天依然是CUDA编程的基石。但我建议你跑通一个例子后别急着进入下一个而是主动折腾几个地方改线程块大小观察耗时变化把共享内存改成全局内存看性能差多少给kernel里加一个条件分支看看不同线程分支对性能的影响试着把两个例子拼在一起比如用共享内存优化直方图。我在实际使用中发现单纯照着源码抄一遍记住的只是语法只有当你亲手把代码改坏、调试、再修好才能真正理解GPU在想什么。接下来再去看CUDA C Programming Guide很多概念就像打通了任督二脉。这本书的源码说到底就是给你拆着玩的。本文还有配套的精品资源点击获取 CUDA By Example源码实战：从老书代码到现代GPU编程
简介《CUDA BY EXAMPLE》配套源码包是围绕GPU高性能编程经典著作整理的全套示例实现适合刚接触CUDA C/C的开发者、科研人员及高校学生对照学习。压缩包共57个文件包含36个.cu源文件、9个.h头文件、4个.lib库和4个.dll动态库等整体仅443KB轻量便携、解压即用。代码覆盖书中第3章至第11章及附录A的核心示例涉及线程块与网格组织、全局内存与共享内存管理、原子操作、多数据流并发、多GPU协作、OpenGL互操作以及性能分析等关键知识点每个示例与书中章节一一对应可帮助读者在真实可运行代码中理解并行计算的设计与调试方法。包内另附README、许可说明及自我介绍等文档方便快速确认环境与使用方式。目前已有410人学习浏览是深入掌握CUDA编程并提升GPU计算效率的实用参考。 《CUDA BY EXAMPLE》中文版叫《GPU高性能编程CUDA实战》我前前后后刷了三遍书里配套的所有例子源码直到现在还在我硬盘上躺着。这篇东西不是书评而是把我整理源码、编译运行、踩坑排查的全过程记录一遍希望能帮到同样拿这本书入门GPU高性能编程的人。这本书是Jason Sanders和Edward Kandrot写的成书时间很早对应的CUDA版本大概在1.x到2.x但读下来你会发现一个很扎心的事实CUDA的生态已经膨胀到难以置信可核心编程模型几乎没变依然是线程层次、内存模型、kernel启动、共享内存、原子操作这些老底子。所以你完全不用担心“书太老代码过时”。书里所有例子源码只要能跑起来你对GPU编程的理解会比看十篇教程都扎实。1. 这本书和配套源码为什么现在还要折腾它1.1 老书不老CUDA编程模型二十年没变很多人一看到“CUDA By Example”这种老书就下意识觉得没必要看这恰恰是误区。现在很多做深度学习、AI推理的人用PyTorch、TensorFlowGPU的调度被框架封装得严严实实甚至连__global__都不用写。可一旦你遇到显存优化、自定义算子、多流并发、模型推理延迟这类问题不懂底层线程模型就是抓瞎。书里的例子源码走的是“最朴素路线”从最简单的hello_world开始到矢量求和、Julia集、点积、波纹、直方图、多GPU协同每一个例子都只围绕一个知识点展开。没有花哨的库没有工程化包袱反而非常适合用来理解GPU到底是怎么把数据拆开、并行、再合并的。我推荐这本书的另一个原因是它的“例子导向”风格。第3章的hello_world会教你写第一个kernel第5章开始引入线程协作后面再慢慢讲性能、原子操作、流和多GPU。这种节奏和现代GPU编程的进阶路径完全一致源码本身就是最好的课程大纲。1.2 源码从哪拿原版目录和GitHub备份这本书原书配套的代码下载链接如今已经不太好用了但GitHub上搜“cuda-by-example”能拿到大量备份仓库。需要注意的是一部分仓库是原封不动的老代码另一部分是好心人改造过、能在新版CUDA上编译通过的版本。我建议两个都下载老代码用来理解原始写法改造版用来快速跑通。原版源码的目录结构一般按章节组织核心文件包括hello_world.cu第一个CUDA程序sum.cuCPU与GPU矢量求和的对比julia.cuJulia集并行计算dot.cu共享内存实现点积ripple.cu共享内存的双缓冲动画histogram.cu原子操作统计直方图thread_fence.cu线程栅栏示例simple_texture.cu纹理内存示例拿到这些源码后不要急着编译先把每个文件打开看一眼主函数和kernel函数的对应关系。我的经验是这几十个例子里真正值得精读的不会超过10个剩下的要么是前面的变体要么是为特定硬件写的花活。2. 环境搭建把书里例子跑起来的完整流程2.1 显卡、驱动和CUDA Toolkit的关系跑CUDA代码你需要搞清楚三件事NVIDIA显卡驱动、CUDA Toolkit、显卡算力Compute Capability。驱动是显卡跑CUDA的基础Toolkit是编译和运行的工具集算力决定了你该用哪个-arch参数。可以简单类比驱动是高速公路Toolkit是车算力是这条路上的限速牌。安装时有个非常关键的原则显卡驱动版本要足够新Toolkit版本和驱动版本匹配。比如你的机器装了最新的NVIDIA驱动那装CUDA 12.x基本没问题但如果驱动很老强行装新Toolkit运行时大概率报cudaErrorInsufficientDriver。老书里的源码对Toolkit版本没那么挑反而是太老的代码配上太新的编译器会出事这点后面细说。安装完Toolkit后打开终端验证一下nvcc --version能看到版本号说明编译器就位。再跑一下nvidia-smi确认你的显卡驱动正常。如果这里已经报错别急着看CUDA先把驱动问题解决。2.2 查看显卡算力并确定编译参数每个例子编译时都会遇到一个绕不开的参数-archsm_XX。这里的sm表示流多处理器架构XX是算力编号比如RTX 3090是sm_86RTX 4090是sm_89A100是sm_80。查看你显卡算力的最简单方式nvidia-smi --query-gpucompute_cap --formatcsv会输出类似8.6这样的结果对应sm_86。编译时就写nvcc -archsm_86 -o hello_world hello_world.cu如果不确定自己的算力或者代码要在多张不同卡上跑建议用这种稳一点的写法nvcc -gencode archcompute_86,codesm_86 -o hello_world hello_world.cu很多老代码自带的Makefile里写的是-archcompute_10这在新版CUDA上已经不被支持必须改成你实际显卡的算力。这一步是最容易踩的第一坑。2.3 老代码依赖的旧库怎么处理书里不少例子在作者写书时用了一个叫cutil的辅助库用来检查CUDA错误、计时、加载纹理等。这个库在现代CUDA Toolkit里没有所以直接编译会报cutil.h: No such file or directory。解决办法有两个。一是用GitHub上已经移除cutil依赖的改造版源码省心二是手把手把cutil调用替换成现代API。我推荐初学者直接选前者等理解代码逻辑后再回头学怎么替换。替换的时候主要做两件事把CUT_CHECK_ERROR(err)这种宏换成cudaGetLastError()判断把cutilDeviceInit(argc, argv)这种初始化函数删掉用cudaSetDevice(0)代替。这两个替换做完90%的旧例子都能在新环境里编译通过。还有一些例子用到了OpenGL和GLUT做可视化比如波纹效果你需要先安装freeglut开发库sudo apt install freeglut3-dev编译时再手动加上-lglut -lGL。这块没有统一答案只能看具体例子。3. 核心例子源码拆解从hello_world到多GPU3.1 第一个kernel搞清楚在干嘛书里的hello_world例子虽然简单但它把CUDA最核心的调用方式演示得明明白白。代码大致长这样__global__ void kernel(void) { printf(Hello, World!\n); } int main(void) { kernel1, 1(); cudaDeviceSynchronize(); return 0; }关键在于1, 1它代表启动kernel时的配置前一个数字是grid里的线程块数量后一个数字是每个线程块里的线程数量。这些方括号里的参数直接决定了GPU上会开多少线程。很多人看完这个例子就过了但我建议你亲手改一改把1, 1改成2, 2看看printf输出几次改成10, 10再观察线程执行顺序。这种“破坏性实验”比背定义管用得多。你会在输出里发现线程顺序是乱序的这正是GPU并行的底层真相你永远不要依赖线程的启动顺序只能依赖显式同步。3.2 矢量求和明白线程和数据的映射关系矢量求和是CUDA入门的“hello world plus”它的核心不是求和的数学而是搞清楚每个线程负责哪个数据。书里的经典做法是int index blockIdx.x * blockDim.x threadIdx.x;这个公式值得你用笔推一遍。blockIdx.x是当前线程块在grid里的编号blockDim.x是块的线程数threadIdx.x是线程在块内的编号。三者组合就能让每个线程拿到自己的全局唯一ID然后去处理数组的第index个元素。一旦这个映射关系想明白后面几乎所有数据并行算法都通了。图像处理、矩阵运算、物理模拟本质都是在做“把一个大数组切成很多份每个线程负责其中一份”这件事。现在OpenCV、PyTorch里很多GPU算子底层也都是这个公式的变体。3.3 点积和共享内存归约操作的原型书里点积dot这个例子我特别喜欢它是所有归约操作的原型。所谓归约就是一大堆数最后归成一个数求和、求最大值、求平均值都是归约。简单做法是每个线程算一部分和再把所有部分和加到一起。但如果这部分和都写回全局内存再读出来效率太低。书里的做法是每个线程块先用共享内存shared缓存部分和__shared__ float partialSum[256];然后通过__syncthreads()等待块内所有线程都完成写入再对共享内存里的值做进一步归约。共享内存是片上资源访问速度远快于全局内存这是GPU性能优化最重要的一课。我在改这个例子时踩过一个小坑__syncthreads()不是随便放的如果某个线程提前return或者分支不一致会导致块内线程等待死锁。书里的例子写法很规范但你自己改造时一定注意别让线程在同步前“跑偏”。3.4 原子操作与多GPU现代大模型推理的底子很多人在学老书时对原子操作和多GPU章节不以为然觉得现在用不上。实际上原子操作在深度学习框架的显存分配、统计归约、量化计算中非常常见。atomicAdd这个函数直到今天依然是CUDA代码里高频出现的东西。书里的直方图例子用原子操作统计像素分布逻辑上很好理解每个线程算出一个结果然后atomicAdd到全局计数器。但要注意原子操作在高并发下会有性能损耗书里的例子可以跑通但实际工程里优化空间很大。多GPU的示例则是通过cudaSetDevice在多个设备间切换现代框架里的“多卡训练”本质上还是这套东西。我把书里的核心例子整理成了一张表方便后面复习的时候对照示例核心知识点对现代开发的参考价值hello_worldkernel启动、线程配置理解最基础的执行模型矢量求和线程与数据映射数据并行的通用套路Julia集网格化并行计算图像/像素级并行点积共享内存、归约深度学习中的Reduce算子波纹双缓冲、同步动画/实时渲染优化直方图原子操作统计、资源分配多GPU示例设备选择、通信多卡训练/推理的雏形4. 实操过程我把每个例子重新编译后的现场记录4.1 官方源码的常见编译错误与修复我使用CUDA 12.3环境显卡是Ampere架构。直接拿某GitHub仓库里的老代码编译第一波就炸了十几个错。最常见的几个问题我记在下面。首先是cutil.h找不到解决方法前面说过找改造版或者替换。其次是nvcc报Unsupported gpu architecture compute_10这个必须去Makefile里把架构参数改成sm_86。还有代码里用了__mul24这类老内置函数在较新版本已被移除改成普通乘法就行。比较隐蔽的一个错误是printf格式问题。老代码里把size_t或者long long直接当成%d打印编译器有时不报错但输出是错的。遇到这种情况用%zu、%lld这类正确格式替换。我修代码的顺序是先找一个能编译的干净例子比如hello_world确认环境没问题再逐个挑战带依赖的例子。千万不要一上来就编译带OpenGL的波纹示例否则你会被链接错误劝退。4.2 运行结果与性能观察所有例子编译通过后我建议把矢量求和那个例子稍微改造一下用来观察GPU并行效果。原书例子里的数组规模比较小GPU的优势根本体现不出来甚至CPU更快。把数组规模调到一千万以上再用clock()计时你才能明显看到差距。我在实测中比较了CPU单线程求和和GPU kernel求和。CPU跑一千万个float大概需要几十毫秒GPU第一次运行要经历JIT编译也差不多这个量级但连续跑第二次就降到几毫秒了。这给了两个启发数据量不够大时GPU并不一定更快数据传输和启动开销很大首次调用kernel会有JIT编译和初始化开销评测性能时不要用第一次的时间。4.3 双显卡和WSL环境的坑我现在开发环境有一部分在WSL2里跑书里例子时遇到过一个典型问题Failed to initialize NVML: GPU access blocked by the operating system。这个问题通常和Windows侧显卡驱动、WSL的GPU挂载权限有关。排查思路分几步。先看看Windows侧nvidia-smi是否正常再看WSL里nvidia-smi是否可用。如果WSL里nvidia-smi正常但CUDA程序起不来检查一下LD_LIBRARY_PATH是否包含/usr/lib/wsl/lib有时候CUDA Toolkit会覆盖这个关键路径。笔记本双显卡同样需要注意。独显和核显同时存在时CUDA默认不一定会选中NVIDIA独显。你可以在程序里显式指定cudaSetDevice(0);或者在运行时用环境变量控制CUDA_VISIBLE_DEVICES0 ./dot这样能减少很多莫名其妙的问题。书里那些老例子没有考虑这种现代笔记本场景你必须自己动手适配。5. 常见问题与排查技巧速查表5.1 编译期问题速查报错信息原因解决办法No such file or directory: cutil.h依赖旧辅助库删除依赖或用改造版源码Unsupported gpu architecture compute_10编译参数太老改为当前显卡算力identifier __mul24 is undefined老内置函数被移除替换为普通乘法cannot find -lglut -lGL缺少可视化库安装freeglut3-deverror: calling a __host__ function from a __global__ function is not allowed在kernel里调用了CPU函数把相关逻辑改为GPU设备函数5.2 运行期问题速查报错信息原因解决办法No kernel image is available for execution on the device编译架构与显卡不匹配用-archsm_XX匹配实际算力cudaErrorInsufficientDriver驱动版本太旧更新NVIDIA驱动Failed to initialize NVMLWSL或驱动环境异常检查WSL和驱动挂载设置LD_LIBRARY_PATHCUDA error: invalid device ordinal显卡编号越界用cudaGetDeviceCount查看设备数量后调整5.3 调试CUDA代码的几点心得编译和运行只是第一步想真正搞懂源码你还需要学会调试和验证。老书年代的cuda-gdb和现在的版本差别很大但核心思路不变先确认数据从CPU到GPU没问题再确认kernel内部计算正确最后确认结果从GPU回传正确。我自己调试的时候特别依赖两个工具cuda-memcheck新版本里是compute-sanitizer用来查越界和未初始化内存Nsight Systems用来分析kernel耗时和内存传输占比。书里老代码没有性能分析的习惯但你跑通之后一定要自己补上这一步。拿dot例子来说如果输入数组只有几千个元素你会看到kernel本身只占很小比例时间大部分时间浪费在cudaMemcpy上这种观察比单纯把代码跑通更有价值。6. 源码不是用来抄的一点实际操作体会书里所有例子源码我都过了一遍最深的体会是代码可以老思想不能老。老代码里的线程映射、共享内存、原子操作这些概念放在今天依然是CUDA编程的基石。但我建议你跑通一个例子后别急着进入下一个而是主动折腾几个地方改线程块大小观察耗时变化把共享内存改成全局内存看性能差多少给kernel里加一个条件分支看看不同线程分支对性能的影响试着把两个例子拼在一起比如用共享内存优化直方图。我在实际使用中发现单纯照着源码抄一遍记住的只是语法只有当你亲手把代码改坏、调试、再修好才能真正理解GPU在想什么。接下来再去看CUDA C Programming Guide很多概念就像打通了任督二脉。这本书的源码说到底就是给你拆着玩的。本文还有配套的精品资源点击获取