📚 课程目录第4章第一个CUDA程序

第四章：第一个CUDA程序——向量加法实战

本章将带领你完成第一个完整的CUDA程序：向量加法。这不仅是CUDA学习的“Hello World”，更是理解GPU编程核心思维——异构计算（Heterogeneous Computing）的最佳切入点。我们将通过一个简单的数学运算，拆解主机（Host，CPU）与设备（Device，GPU）的职责分工，掌握cudaMalloc、cudaMemcpy等关键内存管理API，并首次启动一个并行执行的内核函数（Kernel）。当你完成本章学习后，将具备独立编写、编译和运行基础CUDA程序的能力，并为后续深入并行计算打下坚实的工程基础。

上图清晰地展示了CUDA编程模型的核心：主机与设备拥有各自独立的内存空间，通过PCIe总线进行数据交换。在编写第一个程序前，我们必须建立这样的心智模型：CPU负责逻辑控制与串行任务，GPU负责大规模并行计算。数据必须显式地从主机内存拷贝到设备内存，经过GPU计算后再拷贝回主机。这个过程看似繁琐，但正是这种显式控制，赋予了开发者对性能的极致掌控力。

4.1 核心原理：异构计算与内存模型

CUDA（Compute Unified Device Architecture）的基石是异构计算。一个典型的CUDA程序由两部分组成：

主机代码（Host Code）：运行在CPU上，遵循C/C++标准，负责逻辑控制、数据准备、内核启动和结果回收。主机代码中，我们使用malloc或new分配主机内存。
设备代码（Device Code）：运行在GPU上，使用__global__关键字修饰的函数（称为内核Kernel），负责并行处理数据。设备代码中，我们使用cudaMalloc分配设备全局内存。

理解内存模型是避免程序崩溃的关键。GPU无法直接访问主机内存（h_指针），CPU也无法直接访问设备内存（d_指针）。所有跨域数据传递都必须通过cudaMemcpy函数完成。这种设计虽然增加了编程复杂度，但保证了内存访问的安全性，并允许GPU使用更高效的高速显存。

关键概念：在CUDA中，h_前缀通常表示主机（Host）内存指针，d_前缀表示设备（Device）内存指针。这是一种约定俗成的命名规范，强烈建议在代码中遵循，以提高可读性。

4.2 操作流程：编写第一个向量加法程序

我们将编写一个计算 c[i] = a[i] + b[i] 的程序，其中向量长度为N（例如N=1024）。完整的开发流程分为五个步骤：

包含头文件与错误检查宏：引入<cuda_runtime.h>，并定义一个宏用于检查CUDA API调用是否成功。
编写内核函数：使用__global__修饰，通过threadIdx.x获取当前线程ID。
主机端内存准备：分配并初始化主机内存h_a、h_b、h_c。
设备端内存准备与数据拷贝：使用cudaMalloc分配设备内存，使用cudaMemcpy将输入数据从主机拷贝到设备。
启动内核并回收结果：使用kernel<<<1, N>>>(d_a, d_b, d_c)启动内核，计算完成后将结果拷贝回主机，最后释放所有内存。

下面是完整的代码实现，请仔细阅读每一行注释：

#include <cuda_runtime.h>
#include <iostream>
#include <cstdlib>

// 错误检查宏：封装CUDA API调用，出错则打印文件与行号
#define CHECK_CUDA(call) \
    do { \
        cudaError_t err = (call); \
        if (err != cudaSuccess) { \
            std::cerr << "CUDA Error at " << __FILE__ << ":" << __LINE__ \
                      << " - " << cudaGetErrorString(err) << std::endl; \
            exit(EXIT_FAILURE); \
        } \
    } while(0)

// 内核函数：每个线程计算一个元素
__global__ void vectorAdd(const float *a, const float *b, float *c, int n) {
    int idx = threadIdx.x;  // 获取当前线程的全局ID（此处仅使用block内索引）
    if (idx < n) {          // 边界检查，防止越界
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    const int N = 1024;
    size_t bytes = N * sizeof(float);

    // 1. 分配并初始化主机内存
    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] = static_cast<float>(i);
        h_b[i] = static_cast<float>(i * 2);
    }

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

    // 3. 将输入数据从主机拷贝到设备（H2D）
    CHECK_CUDA(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice));
    CHECK_CUDA(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice));

    // 4. 启动内核：1个block，N个线程
    vectorAdd<<<1, N>>>(d_a, d_b, d_c, N);
    // 检查内核启动是否成功（注意：内核启动是异步的，需要同步检查）
    CHECK_CUDA(cudaGetLastError());
    // 等待内核执行完成
    CHECK_CUDA(cudaDeviceSynchronize());

    // 5. 将结果从设备拷贝回主机（D2H）
    CHECK_CUDA(cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost));

    // 6. 验证结果
    bool correct = true;
    for (int i = 0; i < N; ++i) {
        if (h_c[i] != h_a[i] + h_b[i]) {
            correct = false;
            break;
        }
    }
    std::cout << "Test " << (correct ? "PASSED" : "FAILED") << std::endl;

    // 7. 释放内存
    free(h_a); free(h_b); free(h_c);
    CHECK_CUDA(cudaFree(d_a));
    CHECK_CUDA(cudaFree(d_b));
    CHECK_CUDA(cudaFree(d_c));

    return 0;
}

4.3 案例拆解：从编译到运行的完整闭环

假设我们在一台配备NVIDIA GPU的Linux服务器上工作，使用nvcc编译器。以下是完整的编译与运行过程，以及可能遇到的典型问题。

步骤	命令/操作	预期输出/说明
编译	`nvcc -o vector_add vector_add.cu`	生成可执行文件 `vector_add`。若代码有语法错误，此处会报错。
运行	`./vector_add`	输出 `Test PASSED`。
性能初探	`nvprof ./vector_add` (或 `ncu`)	显示内核执行时间、内存拷贝时间等。你会看到`cudaMemcpy`耗时远大于内核计算耗时，这是PCIe带宽瓶颈的体现。
常见错误1	忘记调用`cudaMemcpy`直接启动内核	程序可能崩溃或输出错误结果。因为设备内存`d_a`未初始化，包含随机值。
常见错误2	在主机代码中直接访问`d_a`	报错 `Segmentation fault`。因为`d_a`是设备内存地址，CPU无法访问。
常见错误3	内核中未做边界检查 `if (idx < n)`	当N不是线程数的整数倍时，会导致越界访问，可能破坏内存或产生`cudaError_t`错误。

在这个案例中，我们使用了cudaGetLastError()来捕获内核启动时的异步错误。这是一个非常重要的习惯：内核启动是异步的，CPU不会等待GPU执行完毕，因此错误可能在下一次同步点才暴露。通过显式调用cudaDeviceSynchronize()，我们强制CPU等待GPU完成，从而确保后续的cudaMemcpy操作读取到的是正确结果。

风险提示：在调试阶段，务必使用CHECK_CUDA宏包裹所有CUDA API调用。不要使用裸的cudaMalloc或cudaMemcpy，否则错误信息将难以定位。此外，cudaDeviceSynchronize()会阻塞CPU，在性能敏感的生产代码中应谨慎使用，但在学习阶段它是保证正确性的必要手段。

4.4 常见误区与总结复盘

初学者在编写第一个CUDA程序时，最容易陷入以下几个误区。我们逐一剖析，并给出规避策略。

误区一：混淆主机与设备内存。这是最致命的问题。记住：malloc分配的是主机内存，cudaMalloc分配的是设备内存。两者不能混用。在代码中，通过h_和d_前缀严格区分。
误区二：忽略错误检查。CUDA API返回cudaError_t类型，必须检查。内核启动错误是异步的，必须通过cudaGetLastError()捕获。建议封装CHECK_CUDA宏，形成肌肉记忆。
误区三：忘记释放内存。设备内存是稀缺资源，忘记cudaFree会导致显存泄漏，最终程序崩溃。务必与cudaMalloc成对使用。
误区四：线程ID使用不当。在只有一个Block的情况下，threadIdx.x的范围是0到N-1。但当我们后续引入多个Block时，需要使用blockIdx.x * blockDim.x + threadIdx.x来计算全局线程ID。本章先不展开，但必须建立这个意识。