3. CUDA程序结构与编译流程

一个CUDA程序从源文件到可执行程序，并不是由单一编译器一次性完成的。它同时包含两种代码：运行在CPU上的主机代码（Host Code）和运行在GPU上的设备代码（Device Code）。这两类代码的指令集、内存模型和执行方式完全不同，因此必须由不同的编译器分别处理，再由链接器合并为一个可执行文件。理解这条编译链，是排查“核函数找不到”“链接失败”“架构不匹配”等问题的前提。

主机与设备代码的分离

在CUDA源文件（通常以 .cu 为扩展名）中，主机代码和设备代码可以写在同一个文件里，但编译器会把它们拆开处理。判断依据是函数前的限定符：

• __host__：仅在主机上执行，由CPU调用，是默认限定符，可省略。
• __global__：核函数，在设备上执行，但必须由主机代码启动，返回类型必须是 void。
• __device__：在设备上执行，只能被核函数或其他 __device__ 函数调用，不能由主机直接调用。

当 nvcc 处理一个 .cu 文件时，它首先把文件中的设备代码抽取出来，交给设备编译器生成GPU的PTX或SASS指令；同时把主机代码中的核函数启动语法（<<<...>>>）替换为对CUDA运行时库的调用，再交给主机编译器（如gcc/clang）编译成普通的目标文件。最后，链接器把主机目标文件与CUDA运行时库链接在一起，形成可执行程序。

nvcc编译驱动流程

nvcc 本身是一个编译驱动程序，它并不直接完成所有编译工作，而是负责拆分代码、调用底层工具链、协调各阶段产物。其典型流程如下：

# 一步完成编译与链接
nvcc -o vector_add vector_add.cu

# 分步观察中间产物
nvcc -c vector_add.cu -o vector_add.o        # 生成目标文件
nvcc -ptx vector_add.cu -o vector_add.ptx    # 只生成PTX
nvcc -cubin vector_add.cu -o vector_add.cubin # 生成cubin
nvcc -o vector_add vector_add.o -lcudart      # 链接CUDA运行时

其中 -arch=sm_80 用于指定目标GPU的计算能力，例如A100对应 sm_80，RTX 4090对应 sm_89。如果编译时指定的架构与运行时的GPU不匹配，可能出现“no kernel image is available for execution on the device”错误。使用 -arch=native 可以让nvcc自动检测当前机器的GPU架构。

核函数定义与启动语法

核函数的定义形式为：

__global__ void vector_add(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

启动核函数使用三尖括号语法：

vector_add<<<gridDim, blockDim>>>(d_a, d_b, d_c, n);

其中 gridDim 表示网格中的线程块数量，blockDim 表示每个线程块中的线程数量。例如 <<<(n+255)/256, 256>>> 表示每个块256个线程，网格大小按元素总数向上取整。核函数启动是异步的：主机发出启动命令后立即继续执行，除非显式同步。因此，在核函数执行完成前读取结果，可能得到未完成的数据。

错误处理机制

CUDA的错误分为两类：启动错误和执行错误。启动错误通常由参数非法、内存不足或架构不匹配引起，可以通过检查 cudaGetLastError() 捕获；执行错误发生在核函数运行期间，例如越界访问，需要通过 cudaDeviceSynchronize() 或后续的CUDA API调用返回的错误码来发现。

vector_add<<<grid, block>>>(d_a, d_b, d_c, n);
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
    printf("Kernel launch failed: %s\n", cudaGetErrorString(err));
}
err = cudaDeviceSynchronize();
if (err != cudaSuccess) {
    printf("Kernel execution failed: %s\n", cudaGetErrorString(err));
}

一个常见的工程习惯是封装 CUDA_CHECK 宏，把错误码检查与文件名、行号绑定，便于定位问题。需要注意的是，cudaGetLastError() 会清除当前错误状态，因此应在每次核函数启动后立即调用，而不是等到程序末尾统一检查。

编译流程中的关键判断

在实际工程中，编译流程是否顺利，往往取决于几个关键判断。第一，目标架构是否覆盖运行环境：如果编译时只指定了 sm_80，而程序运行在 sm_89 的GPU上，可能因缺少对应二进制而失败。第二，主机与设备代码的边界是否清晰：__device__ 函数不能被主机直接调用，__global__ 函数不能有非void返回值，这些限制在编译期就会报错。第三，错误检查是否到位：核函数启动是异步的，如果只检查启动错误而不做同步，执行期的越界访问可能被忽略，导致结果错误却难以定位。

把编译流程、代码分离和错误处理放在一起看，可以发现它们共同指向一个工程原则：CUDA程序的正确性不仅取决于核函数逻辑，还取决于主机与设备之间的协作边界是否被正确建立和维护。编译期负责把两类代码放到正确的位置，运行期负责在边界处检查状态，二者缺一不可。