企业官网建设流程全解析

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和device

CUDA程序里有两个角色：主机（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在设备上计算 myKernel<<<gridDim, 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函数为什么必须返回void

CUDA的kernel函数有几个硬性规定：必须用__global__修饰，必须返回void，不能是类的成员函数（除非是静态的）。这些限制背后是有原因的。

__global__表示这个函数在设备上执行、从主机调用。它返回void是因为kernel是异步启动的——主机发出启动命令后就继续往下走了，根本不等kernel执行完。既然主机不等待返回值，返回值就没有意义。如果你需要从kernel获取结果，只能通过设备指针写回显存，再拷贝回主机。

这里有个很重要的概念：kernel启动是异步的。下面这段代码：

myKernel<<<1, 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提供了几个内置变量：

对于一维的情况，全局线程索引就是：

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; vectorAdd<<<gridSize, 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慢得离谱

全局内存的访问有一个重要特性：如果同一个warp（32个线程）访问连续的内存地址，硬件可以把这些访问合并成一次大事务。这叫合并访问（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 -arch=sm_75 --ptxas-options=-v mykernel.cu

输出会显示每个kernel用了多少个寄存器、多少共享内存、多少本地内存。如果寄存器数量接近255（上限），说明寄存器压力很大，可能需要优化代码结构。

我的一般经验是：对于计算密集型的kernel，寄存器多用一点没关系，因为计算本身就能隐藏延迟；对于内存密集型的kernel，occupancy更重要，需要控制寄存器使用量。

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()来检查：

myKernel<<<grid, 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); myKernel<<<grid, 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 myKernel<<<grid, block, 0, stream1>>>(d_data1); // 在stream2上启动另一个kernel myKernel<<<grid, block, 0, stream2>>>(d_data2); // 等待两个流都完成 cudaStreamSynchronize(stream1); cudaStreamSynchronize(stream2);

<<<>>>的第三个参数是共享内存大小，第四个参数就是流。如果不指定流，就用默认流（stream 0）。

多流并发的效果取决于kernel的资源占用。如果每个kernel都把GPU占满了，并发就没有意义。一般来说，小kernel或者资源占用不高的kernel适合用多流来重叠执行。

4.4 常见错误码的含义与排查方向

CUDA的错误码有几十个，但新手最常遇到的就那么几个：

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); myKernel<<<grid, block>>>(d_data); cudaFree(d_data); }

正确的做法是在循环外分配一次，循环内重复使用：

// 正确做法 float *d_data; cudaMalloc(&d_data, bytes); for (int i = 0; i < 1000; i++) { myKernel<<<grid, 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大小来避免寄存器溢出。对于共享内存使用量大的kernel，block大小又会影响共享内存的分配。

最可靠的方法还是实测。写一个简单的benchmark，遍历几个不同的block大小，看哪个性能最好。这个测试花不了多少时间，但收益可能很大。

5.4 坑四：忘记检查kernel启动是否成功

kernel启动失败不会导致程序崩溃，只会设置一个错误状态。如果不检查，后续的操作可能会基于错误的状态执行，导致更难排查的问题。

myKernel<<<grid, 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查询硬件参数，根据实际情况调整配置，而不是硬编码某个特定的值。