📚 课程目录第3章 CUDA编程模型入门

3. CUDA编程模型入门：CUDA程序结构，主机与设备代码，核函数（kernel）的编写与启动

3.1 CUDA程序的基本结构：主机与设备的协作

CUDA 编程模型的核心思想是异构计算：CPU（主机，Host）负责逻辑控制与串行任务，GPU（设备，Device）负责大规模并行计算。一个完整的 CUDA 程序由两部分代码组成：主机代码（在 CPU 上运行）和设备代码（在 GPU 上运行）。这两部分代码在同一个源文件（.cu）中编写，由 nvcc 编译器区分并分别编译。

主机代码负责：分配设备内存、将输入数据从主机拷贝到设备、启动核函数（Kernel）、将结果从设备拷回主机、释放设备内存。设备代码则特指在 GPU 上执行的核函数，它由大量线程并行执行。理解两者的边界是编写高效 CUDA 程序的第一步。

  关键认知：核函数是异步的。主机在启动核函数后不会等待 GPU 执行完毕，而是立即继续执行后续指令。若需要同步，必须显式调用 cudaDeviceSynchronize() 或使用同步内存拷贝（如 cudaMemcpy 在默认情况下会阻塞主机直到拷贝完成）。

3.2 核函数（Kernel）的编写规范

核函数是 CUDA 程序的灵魂。它使用 __global__ 修饰符声明，返回类型必须为 void。核函数内部通过内置变量 threadIdx、blockIdx、blockDim 和 gridDim 来索引当前线程的全局 ID，从而让每个线程处理不同的数据元素。

一个典型的核函数编写模板如下：

// 核函数定义：向量加法
__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];
    }
}

核函数编写有以下几个硬性约束：

只能访问设备内存（显存），不能直接访问主机内存指针。
不能调用标准 C 库函数（如 printf 除外，但需在计算能力 2.0 以上），不能使用 malloc 动态分配设备内存。
必须声明返回类型为 void，不能向主机返回任何值。
支持递归，但受限于栈空间，实际使用极少。

3.3 核函数的启动配置：网格与线程块

核函数通过特殊的尖括号语法启动：kernel<<<gridDim, blockDim>>>(参数列表)。其中 gridDim 和 blockDim 可以是 int 或 dim3 类型，分别指定网格中线程块的数量和每个线程块中线程的数量。这种两级层次结构是 CUDA 并行模型的核心设计，它使得 GPU 能够高效地调度和管理成千上万个线程。

启动配置决定了线程的组织方式，直接影响并行效率和资源占用。以下是一个完整的向量加法示例：

int main() {
    int n = 1 << 20;  // 1M 个元素
    size_t bytes = n * sizeof(float);

    // 1. 分配主机内存并初始化
    float *h_a, *h_b, *h_c;
    h_a = (float*)malloc(bytes);
    h_b = (float*)malloc(bytes);
    h_c = (float*)malloc(bytes);
    for (int i = 0; i < n; i++) { h_a[i] = i * 0.5f; h_b[i] = i * 1.5f; }

    // 2. 分配设备内存
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);

    // 3. 拷贝数据到设备
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);

    // 4. 配置启动参数并调用核函数
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, n);

    // 5. 同步并拷贝结果回主机
    cudaDeviceSynchronize();
    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);

    // 6. 验证结果并释放内存
    for (int i = 0; i < 10; i++) printf("%.2f ", h_c[i]);
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    free(h_a); free(h_b); free(h_c);
    return 0;
}

线程块大小选择建议：通常选择 128 或 256 个线程/块。线程块大小应为 32 的倍数（一个 warp 的大小），以保证线程调度效率。过小的块（如 32）会导致资源利用率不足，过大的块（如 1024）可能超出硬件限制。

3.4 线程索引计算与内存访问模式

在多维网格中，线程的全局索引计算需要区分 x、y、z 三个维度。对于二维或三维问题（如图像处理、矩阵运算），使用多维索引可以更自然地映射数据。以下是一个二维矩阵加法的核函数示例：

// 二维核函数：矩阵加法
__global__ void matrixAdd(const float *A, const float *B, float *C,
                          int width, int height) {
    // 计算当前线程的二维坐标
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;

    // 边界检查
    if (col < width && row < height) {
        int idx = row * width + col;  // 行主序存储
        C[idx] = A[idx] + B[idx];
    }
}

// 启动配置
dim3 block(16, 16);          // 16x16 = 256 线程/块
dim3 grid((width + 15) / 16, (height + 15) / 16);
matrixAdd<<<grid, block>>>(d_A, d_B, d_C, width, height);

在编写核函数时，内存访问模式对性能影响极大。GPU 以 warp（32 线程）为单位执行指令，当同一 warp 中的线程访问连续内存地址时，硬件会将多次访问合并为一次事务，这称为内存合并访问。反之，若线程访问跳跃地址，则会产生多次事务，严重降低带宽利用率。

访问模式	示例	带宽利用率	建议
合并访问（连续）	`C[idx] = A[idx] + B[idx]`，其中 idx 连续	接近 100%	优先采用
步长访问	`C[idx * 2]`	50% 左右	尽量避免
随机访问	按索引数组取值	极低	考虑使用共享内存或纹理内存

3.5 常见误区与调试复盘

初学者在编写 CUDA 程序时，经常会遇到以下问题，这里逐一进行复盘分析：

误区一：忘记错误检查。CUDA 的大多数 API 函数都会返回错误码，但核函数启动不返回错误。如果核函数内部出错（如越界访问），主机端往往毫无察觉，直到程序崩溃或产生错误结果。建议在每次 CUDA API 调用后检查错误，并在核函数启动后调用 cudaGetLastError() 捕获异步错误。

// 推荐的错误检查宏
#define CHECK_CUDA(call) do { \
    cudaError_t err = (call); \
    if (err != cudaSuccess) { \
        fprintf(stderr, "CUDA error at %s:%d: %s\n", \
                __FILE__, __LINE__, cudaGetErrorString(err)); \
        exit(EXIT_FAILURE); \
    } \
} while(0)

// 使用方式
CHECK_CUDA(cudaMalloc(&d_a, bytes));
vectorAdd<<<grid, block>>>(d_a, d_b, d_c, n);
CHECK_CUDA(cudaGetLastError());  // 捕获核函数启动错误
CHECK_CUDA(cudaDeviceSynchronize());  // 等待核函数执行完毕并检查错误

误区二：主机与设备内存混用。在核函数中直接访问主机指针（如 h_a）会导致程序崩溃，因为 GPU 无法访问主机内存。反之，在主机代码中直接访问设备指针（如 d_a）同样非法。必须通过 cudaMemcpy 显式地在两者之间传输数据。

误区三：忽略边界条件。当数据总量不是线程块大小的整数倍时，最后一个线程块中的部分线程会超出数据范围。如果不做边界检查，这些线程会访问非法内存地址，导致不可预知的结果。务必在核函数中加入 if (idx < n) 之类的判断。

  复盘总结：一个健壮的 CUDA 程序应遵循以下流程——(1) 分配主机和设备内存；(2) 数据拷贝到设备；(3) 配置合理的网格/线程块参数；(4) 启动核函数并立即检查错误；(5) 同步并拷贝结果回主机；(6) 验证结果正确性；(7) 释放所有内存。每一步都不可省略，尤其是错误检查和结果验证，它们是排查问题的关键手段。

通过本节的学习，你已经掌握了 CUDA 程序的基本骨架：主机与设备的职责划分、核函数的编写规范、启动配置的语法以及线程索引的计算方法。这些知识是后续深入学习 GPU 内存层次结构、线程同步和性能优化的基础。在实际编程中，建议从简单的向量运算开始，逐步过渡到矩阵运算和更复杂的并行算法，并在每一步都做好错误检查和性能分析。