1. 从一张显卡的算力清单说起为什么我决定系统啃CUDA手里有一块支持CUDA的显卡装完驱动跑个nvidia-smi看到算力版本、显存大小、CUDA版本号都正常很多人到这一步就停了。真正开始写代码才发现torch.cuda.is_available()返回True只是万里长征第一步——想自己写个自定义算子、想优化一个卷积实现、想把一段循环从CPU搬到GPU上跑出几十倍加速中间隔着一整套并行编程的思维转换。我最初接触CUDA的动机很朴素一个图像处理的小项目CPU上跑一帧要几百毫秒实时性完全不够。当时想当然地觉得“把for循环改成GPU版本就行”结果第一版代码写出来比CPU还慢。后来才明白GPU编程不是把代码翻译一遍而是要把问题重新拆解成成千上万个线程能同时干的活。这个认知转变花了我不少时间也踩了不少坑。这篇笔记面向的是有C/C基础、了解计算机体系结构基本概念、但还没系统写过CUDA程序的开发者。我会从最核心的执行模型讲起把线程组织、内存层次、同步机制这些容易混淆的概念用实际代码串起来再补充几个新手最容易翻车的地方。读完之后你应该能独立写出一个结构正确、性能合理的CUDA程序并且知道往哪个方向继续深入。需要提前说明的是CUDA的版本迭代很快不同版本之间API有细微差异。我下面提到的内容以CUDA 11.x和12.x为主如果你用的是更早的版本个别函数名可能不同但核心概念是通用的。2. 主机与设备的分工逻辑CUDA执行模型拆解2.1 从“谁在干活”理解host和deviceCUDA程序里有两个角色主机host和设备device。主机就是CPU和它的内存设备就是GPU和它的显存。这两个角色是物理隔离的——CPU不能直接读写显存GPU也不能直接访问主存。所有的数据交换都必须显式地通过PCIe总线拷贝。这个设计初看很麻烦但理解之后就会发现它其实很合理。GPU的显存带宽远高于主存带宽如果允许CPU随意访问显存反而会成为性能瓶颈。强制分离让开发者必须认真考虑数据在哪里、什么时候需要搬。一个典型的CUDA程序结构是这样的// 1. 在主机上分配内存并初始化数据 float *h_data (float *)malloc(N * sizeof(float)); // ... 填充数据 ... // 2. 在设备上分配显存 float *d_data; cudaMalloc(d_data, N * sizeof(float)); // 3. 把数据从主机拷贝到设备 cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice); // 4. 启动kernel在设备上计算 myKernelgridDim, blockDim(d_data, N); // 5. 把结果从设备拷贝回主机 cudaMemcpy(h_data, d_data, N * sizeof(float), cudaMemcpyDeviceToHost); // 6. 释放显存和主机内存 cudaFree(d_data); free(h_data);这六步是CUDA程序的骨架。新手最容易犯的错误是忘记第3步或第5步直接在kernel里访问主机指针结果要么程序崩溃要么得到一堆垃圾数据。更隐蔽的错误是拷贝方向写反了cudaMemcpyHostToDevice和cudaMemcpyDeviceToHost搞混编译器不会报错但结果完全不对。2.2 kernel函数为什么必须返回voidCUDA的kernel函数有几个硬性规定必须用__global__修饰必须返回void不能是类的成员函数除非是静态的。这些限制背后是有原因的。__global__表示这个函数在设备上执行、从主机调用。它返回void是因为kernel是异步启动的——主机发出启动命令后就继续往下走了根本不等kernel执行完。既然主机不等待返回值返回值就没有意义。如果你需要从kernel获取结果只能通过设备指针写回显存再拷贝回主机。这里有个很重要的概念kernel启动是异步的。下面这段代码myKernel1, 1(d_data); printf(kernel finished\n);printf几乎肯定会在kernel执行完之前就打印出来。如果你需要确保kernel执行完毕必须调用cudaDeviceSynchronize()。这个函数会阻塞主机线程直到设备上所有之前的任务完成。我见过有人在kernel启动后立刻cudaMemcpy回主机然后疑惑为什么数据不对。原因就是拷贝操作可能在kernel执行完之前就开始了。正确的做法是在拷贝前加同步或者利用CUDA的流机制来管理依赖关系。2.3 线程层次grid、block、thread的三层结构CUDA的线程组织是三层结构grid包含多个block每个block包含多个thread。这个结构和硬件是直接对应的——一个block会被调度到一个SM流多处理器上执行block内的线程可以通过共享内存通信不同block之间则完全独立。启动kernel时的gridDim, blockDim语法就是指定这两个维度。比如4, 256表示启动4个block每个block有256个线程总共1024个线程。线程的索引计算是新手最容易搞混的地方。每个线程需要知道“我是谁”才能决定处理哪部分数据。CUDA提供了几个内置变量变量名含义维度threadIdx线程在block内的索引三维blockIdxblock在grid内的索引三维blockDimblock的维度大小三维gridDimgrid的维度大小三维对于一维的情况全局线程索引就是int idx blockIdx.x * blockDim.x threadIdx.x;这个公式必须刻在脑子里。我刚开始写的时候经常写成blockIdx.x threadIdx.x * blockDim.x结果就是线程访问了错误的内存位置程序要么崩溃要么结果错乱。对于二维的情况比如处理图像索引计算变成int x blockIdx.x * blockDim.x threadIdx.x; int y blockIdx.y * blockDim.y threadIdx.y; int offset y * width x;这里width是图像的宽度。注意blockDim.x和blockDim.y可能不同不要想当然地认为它们相等。2.4 一个完整的向量加法示例光说概念太抽象来看一个完整的向量加法实现#include stdio.h #include cuda_runtime.h __global__ void vectorAdd(const float *a, const float *b, float *c, int n) { int idx blockIdx.x * blockDim.x threadIdx.x; if (idx n) { c[idx] a[idx] b[idx]; } } int main() { int n 100000; size_t bytes n * sizeof(float); // 主机内存分配 float *h_a (float *)malloc(bytes); float *h_b (float *)malloc(bytes); float *h_c (float *)malloc(bytes); // 初始化数据 for (int i 0; i n; i) { h_a[i] i * 1.0f; h_b[i] i * 2.0f; } // 设备内存分配 float *d_a, *d_b, *d_c; cudaMalloc(d_a, bytes); cudaMalloc(d_b, bytes); cudaMalloc(d_c, bytes); // 数据拷贝到设备 cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice); cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice); // 启动kernel int blockSize 256; int gridSize (n blockSize - 1) / blockSize; vectorAddgridSize, blockSize(d_a, d_b, d_c, n); // 结果拷贝回主机 cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost); // 验证结果 for (int i 0; i 10; i) { printf(c[%d] %f\n, i, h_c[i]); } // 释放内存 cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); free(h_a); free(h_b); free(h_c); return 0; }这段代码里有几个值得注意的细节。gridSize的计算用了(n blockSize - 1) / blockSize这是向上取整的标准写法。如果直接写n / blockSize当n不是blockSize的整数倍时最后一部分数据就不会被处理。kernel里的if (idx n)判断也是必须的。因为gridSize是向上取整的实际启动的线程数可能略多于n多出来的线程如果不加判断就会访问越界内存。3. 内存层次与访问效率显存不是一块平坦的草地3.1 从全局内存到寄存器的金字塔CUDA的内存模型是一个层次结构从下到上依次是全局内存Global Memory容量最大几个GB延迟最高几百个时钟周期所有线程都能访问。常量内存Constant Memory容量小64KB有缓存适合存放所有线程都读取的只读数据。共享内存Shared Memory每个block独有容量小几十KB延迟低几十个时钟周期是block内线程通信的主要手段。寄存器Register每个线程独有速度最快但数量有限。这个层次结构决定了CUDA程序的性能优化方向尽量让数据在离计算单元近的地方。全局内存访问是性能杀手能避免就避免不能避免就想办法合并访问。3.2 合并访问为什么你的kernel慢得离谱全局内存的访问有一个重要特性如果同一个warp32个线程访问连续的内存地址硬件可以把这些访问合并成一次大事务。这叫合并访问coalesced access。如果访问是分散的硬件就需要发起多次事务效率成倍下降。举个例子假设有一个数组data[1024]warp里的32个线程分别访问data[threadIdx.x]这是完美的合并访问——32个连续地址一次事务搞定。但如果访问的是data[threadIdx.x * 2]地址就变成了0, 2, 4, 6...硬件需要两次事务才能覆盖。如果访问的是data[threadIdx.x * 32]地址跨度太大可能需要32次事务。我在优化一个矩阵转置kernel时深刻体会到了这一点。最初的实现是直接读一行写一列读的时候是合并的写的时候完全分散性能惨不忍睹。后来用共享内存做中转先把数据合并读到共享内存再从共享内存合并写到全局内存性能提升了将近10倍。3.3 共享内存的正确打开方式共享内存的使用需要显式声明和同步。来看一个用共享内存优化矩阵乘法的例子#define TILE_SIZE 16 __global__ void matMulShared(const float *A, const float *B, float *C, int N) { __shared__ float tileA[TILE_SIZE][TILE_SIZE]; __shared__ float tileB[TILE_SIZE][TILE_SIZE]; int row blockIdx.y * TILE_SIZE threadIdx.y; int col blockIdx.x * TILE_SIZE threadIdx.x; float sum 0.0f; for (int t 0; t N / TILE_SIZE; t) { // 协作加载tile到共享内存 tileA[threadIdx.y][threadIdx.x] A[row * N t * TILE_SIZE threadIdx.x]; tileB[threadIdx.y][threadIdx.x] B[(t * TILE_SIZE threadIdx.y) * N col]; // 等待所有线程完成加载 __syncthreads(); // 计算tile内的部分和 for (int k 0; k TILE_SIZE; k) { sum tileA[threadIdx.y][k] * tileB[k][threadIdx.x]; } // 等待所有线程完成计算才能加载下一个tile __syncthreads(); } C[row * N col] sum; }这段代码里有几个关键点。__shared__声明的数组在每个block内是独立的不同block的同名数组互不干扰。__syncthreads()是block内的同步屏障所有线程必须都到达这个点才能继续。如果漏掉第二个__syncthreads()可能会出现某些线程还在读tileA另一些线程已经开始写下一轮数据了结果就是数据竞争。共享内存的容量是有限的。在大多数GPU上每个block最多使用48KB的共享内存可配置到更高但需要特殊设置。如果声明了超过限制的共享内存kernel启动会直接失败。我遇到过有人声明了一个__shared__ float buffer[1024][1024]编译能过运行时报错查了半天才发现是共享内存超了。3.4 寄存器压力与occupancy的权衡寄存器是每个线程私有的数量有限。如果kernel里用了太多局部变量寄存器不够用编译器会把一些变量放到本地内存local memory而本地内存实际上是全局内存的一部分访问速度极慢。寄存器的使用量会影响occupancy也就是一个SM上能同时驻留多少个warp。occupancy越高延迟隐藏效果越好但每个线程能用的寄存器就越少。这是一个需要权衡的问题。你可以用nvcc的--ptxas-options-v选项查看kernel的寄存器使用情况nvcc -archsm_75 --ptxas-options-v mykernel.cu输出会显示每个kernel用了多少个寄存器、多少共享内存、多少本地内存。如果寄存器数量接近255上限说明寄存器压力很大可能需要优化代码结构。我的一般经验是对于计算密集型的kernel寄存器多用一点没关系因为计算本身就能隐藏延迟对于内存密集型的kerneloccupancy更重要需要控制寄存器使用量。4. 同步、错误处理与性能测量那些文档里不会细说的细节4.1 错误处理不是可选项CUDA的API调用几乎都会返回一个cudaError_t类型的错误码。很多人写CUDA代码时完全不检查返回值出了问题就一脸茫然。我强烈建议在开发的每个阶段都加上错误检查。最基础的检查方式是cudaError_t err cudaMalloc(d_data, bytes); if (err ! cudaSuccess) { printf(cudaMalloc failed: %s\n, cudaGetErrorString(err)); return -1; }但这样每个调用都写一遍太啰嗦。可以定义一个宏#define CUDA_CHECK(call) \ do { \ cudaError_t err call; \ if (err ! cudaSuccess) { \ printf(CUDA error at %s:%d: %s\n, __FILE__, __LINE__, \ cudaGetErrorString(err)); \ exit(1); \ } \ } while (0)然后所有CUDA调用都包在CUDA_CHECK()里。这样一旦出错能立刻知道是哪个文件哪一行出的问题。kernel启动的错误比较特殊因为它不返回错误码。需要用cudaGetLastError()来检查myKernelgrid, block(args); CUDA_CHECK(cudaGetLastError());注意cudaGetLastError()会清除错误状态所以不要连续调用两次。4.2 用事件来精确计时CPU上的clock()或std::chrono不能用来测量kernel的执行时间因为kernel是异步的。正确的做法是使用CUDA事件cudaEvent_t start, stop; cudaEventCreate(start); cudaEventCreate(stop); cudaEventRecord(start); myKernelgrid, block(args); cudaEventRecord(stop); cudaEventSynchronize(stop); float milliseconds 0; cudaEventElapsedTime(milliseconds, start, stop); printf(kernel time: %.3f ms\n, milliseconds); cudaEventDestroy(start); cudaEventDestroy(stop);cudaEventRecord在流中插入一个标记cudaEventElapsedTime计算两个标记之间的时间。这个时间包含了kernel的实际执行时间但不包含主机端的启动开销。如果你想测量包含数据传输在内的总时间把cudaMemcpy也放在两个事件之间就行。但要注意第一次运行kernel时可能会有一些初始化开销最好先跑几遍热身再取平均值。4.3 流与并发让多个kernel同时跑默认情况下所有kernel都在同一个流stream里顺序执行。但GPU其实支持多个流并发执行只要资源允许。这对于需要同时处理多个独立任务的场景很有用。cudaStream_t stream1, stream2; cudaStreamCreate(stream1); cudaStreamCreate(stream2); // 在stream1上启动kernel myKernelgrid, block, 0, stream1(d_data1); // 在stream2上启动另一个kernel myKernelgrid, block, 0, stream2(d_data2); // 等待两个流都完成 cudaStreamSynchronize(stream1); cudaStreamSynchronize(stream2);的第三个参数是共享内存大小第四个参数就是流。如果不指定流就用默认流stream 0。多流并发的效果取决于kernel的资源占用。如果每个kernel都把GPU占满了并发就没有意义。一般来说小kernel或者资源占用不高的kernel适合用多流来重叠执行。4.4 常见错误码的含义与排查方向CUDA的错误码有几十个但新手最常遇到的就那么几个错误码含义常见原因cudaErrorMemoryAllocation内存分配失败显存不够或者请求的大小为负数cudaErrorInvalidValue参数无效grid或block维度为0或者超过了硬件限制cudaErrorLaunchFailurekernel启动失败访问了非法内存或者kernel执行超时cudaErrorIllegalAddress非法地址访问数组越界或者访问了未分配的内存cudaErrorSynchronizationError同步错误之前的异步操作出错了cudaErrorIllegalAddress是最常见的错误之一。它通常意味着kernel里访问了越界的内存地址。排查方法是检查所有数组索引的计算特别是blockIdx和threadIdx的组合。可以用compute-sanitizer工具来定位具体的越界位置compute-sanitizer ./my_program这个工具会报告每一次非法访问的详细信息包括是读还是写、地址是多少、在哪个kernel里发生的。虽然它会拖慢程序运行速度但排查内存问题时非常有用。5. 从能跑到跑得快新手最该避开的五个坑5.1 坑一在循环里反复分配和释放显存cudaMalloc和cudaFree是相对昂贵的操作涉及与驱动程序的交互。如果在循环里反复分配释放开销会非常大。// 错误做法 for (int i 0; i 1000; i) { float *d_data; cudaMalloc(d_data, bytes); myKernelgrid, block(d_data); cudaFree(d_data); }正确的做法是在循环外分配一次循环内重复使用// 正确做法 float *d_data; cudaMalloc(d_data, bytes); for (int i 0; i 1000; i) { myKernelgrid, block(d_data); } cudaFree(d_data);如果确实需要动态管理显存可以考虑实现一个简单的内存池预先分配一大块显存然后自己管理分配和释放。5.2 坑二忽略数据传输的开销GPU计算很快但PCIe传输很慢。如果计算量不够大数据传输的时间可能比计算时间还长整体反而比CPU慢。我做过一个测试对一个100万元素的数组做简单加法CPU版本耗时约2毫秒GPU版本的计算只用了0.1毫秒但数据传输花了3毫秒总共3.1毫秒反而更慢。判断是否值得用GPU的一个粗略标准是计算量是否足够大能掩盖数据传输的开销。对于简单的逐元素操作通常需要至少几百万元素才能体现出优势。对于计算密集型的操作如矩阵乘法、卷积门槛会低很多。另外如果数据本来就在GPU上比如是上一个kernel的输出就不需要传输这时候GPU的优势就很明显了。5.3 坑三block大小设置得太随意block大小blockDim的选择对性能有显著影响。太小会导致SM利用率不足太大则可能限制occupancy。一般来说block大小应该是32的倍数warp的大小常用的值是128、256、512。我个人的经验是256是一个比较安全的默认值适用于大多数场景。但这不是绝对的。对于寄存器使用量大的kernel可能需要减小block大小来避免寄存器溢出。对于共享内存使用量大的kernelblock大小又会影响共享内存的分配。最可靠的方法还是实测。写一个简单的benchmark遍历几个不同的block大小看哪个性能最好。这个测试花不了多少时间但收益可能很大。5.4 坑四忘记检查kernel启动是否成功kernel启动失败不会导致程序崩溃只会设置一个错误状态。如果不检查后续的操作可能会基于错误的状态执行导致更难排查的问题。myKernelgrid, block(args); // 如果不检查这里可能已经出错了 cudaMemcpy(h_data, d_data, bytes, cudaMemcpyDeviceToHost); // 拷贝回来的可能是垃圾数据养成习惯每次kernel启动后都加一句CUDA_CHECK(cudaGetLastError())。这个检查几乎不花时间但能帮你省下大量调试时间。5.5 坑五在kernel里用printf调试CUDA确实支持在kernel里调用printf但输出顺序是混乱的而且大量printf会严重拖慢kernel执行。更麻烦的是printf的输出缓冲区大小有限如果输出太多可能会丢失部分内容。我的建议是调试阶段可以用printf但只打印关键信息并且加上线程索引作为前缀方便定位。生产代码里绝对不要留printf。更好的调试方式是先把中间结果写回全局内存拷贝回主机后再检查。虽然麻烦一点但更可靠。6. 下一步往哪走从能写kernel到写好kernel把上面这些内容消化完你应该能写出一个功能正确的CUDA程序了。但“能跑”和“跑得快”之间还有很大的距离。接下来可以深入的方向有几个。一是性能分析工具nsight compute和nsight systems能告诉你kernel的瓶颈在哪里是内存带宽受限还是计算受限是occupancy不够还是warp分歧严重。二是高级内存优化比如用__ldg做只读缓存、用向量化加载float4提高带宽利用率、用异步拷贝cudaMemcpyAsync重叠传输和计算。三是多GPU编程当单卡显存不够或者算力不足时如何把任务拆分到多张卡上。我个人的学习路径是先写一个能跑的版本然后用nsight compute看瓶颈针对瓶颈做优化再看效果。这个循环重复几次对CUDA的理解就会深入很多。最怕的是一上来就追求极致性能结果连基本的正确性都没保证调到最后都不知道是哪里出了问题。另外CUDA的官方文档和样例代码是最好的学习资源。cuda-samples仓库里有大量精心编写的示例覆盖了从基础到高级的各种技术。遇到不确定的API用法直接查文档比在网上搜答案靠谱得多。最后分享一个我踩过的坑不要假设不同GPU的行为是一致的。我在一张卡上跑得好好的kernel换到另一张卡上可能因为寄存器数量、共享内存大小、warp调度策略的差异而表现完全不同。写代码时尽量用cudaGetDeviceProperties查询硬件参数根据实际情况调整配置而不是硬编码某个特定的值。