4. 第一个CUDA程序：向量加法

好，咱们直接上手写代码。

学GPU编程，第一个程序基本都是向量加法。为什么？因为它简单、直观，而且能清楚展示CPU和GPU的协作流程。说白了，就是两个数组对应元素相加。CPU也能做，但GPU做起来更快——尤其是数据量大的时候。

4.1 先看完整代码

我先把代码贴出来，然后一行行拆开讲。你最好跟着敲一遍，光看是学不会的。

#include <stdio.h>
#include <cuda_runtime.h>

__global__ void vecAdd(float *A, float *B, float *C, int n)
{
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < n) {
        C[i] = A[i] + B[i];
    }
}

int main()
{
    int n = 1000000;
    size_t size = n * sizeof(float);

    // 分配主机内存
    float *h_A = (float*)malloc(size);
    float *h_B = (float*)malloc(size);
    float *h_C = (float*)malloc(size);

    // 初始化数据
    for (int i = 0; i < n; i++) {
        h_A[i] = 1.0f;
        h_B[i] = 2.0f;
    }

    // 分配设备内存
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, size);
    cudaMalloc(&d_B, size);
    cudaMalloc(&d_C, size);

    // 拷贝数据到GPU
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

    // 调用核函数
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, n);

    // 检查错误
    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess) {
        printf("CUDA error: %s\n", cudaGetErrorString(err));
    }

    // 拷贝结果回CPU
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

    // 验证结果
    for (int i = 0; i < 10; i++) {
        printf("C[%d] = %f\n", i, h_C[i]);
    }

    // 释放内存
    free(h_A); free(h_B); free(h_C);
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);

    return 0;
}

4.2 内存分配：cudaMalloc

GPU不能直接用CPU的内存。你得先在GPU上划一块地。

cudaMalloc 就是干这个的。用法和 malloc 几乎一样，只是它分配的是设备端（GPU）的内存。

float *d_A;
cudaMalloc((void**)&d_A, size);

注意第一个参数是二级指针。为什么？因为 cudaMalloc 要修改指针本身的值，让它指向GPU内存的地址。我刚开始写的时候老忘这个 &，结果程序跑起来直接崩。

4.3 数据拷贝：cudaMemcpy

数据在CPU和GPU之间传输，靠 cudaMemcpy。

cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

四个参数：目标地址、源地址、字节数、方向。

方向有四种：

• cudaMemcpyHostToDevice — CPU到GPU
• cudaMemcpyDeviceToHost — GPU到CPU
• cudaMemcpyDeviceToDevice — GPU到GPU
• cudaMemcpyHostToHost — CPU到CPU（基本不用）

我曾经犯过一个低级错误：把方向写反了。结果数据没拷过去，核函数算出来全是0。排查了半天才发现是 HostToDevice 写成了 DeviceToHost。嗯，这种错误你迟早也会遇到。

4.4 核函数调用语法

核函数就是跑在GPU上的函数。定义时前面加 __global__。

__global__ void vecAdd(float *A, float *B, float *C, int n)
{
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < n) {
        C[i] = A[i] + B[i];
    }
}

调用时用三括号语法：

vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, n);

三括号里第一个参数是网格中的线程块数量，第二个是每个线程块中的线程数。

为什么是 threadIdx.x + blockIdx.x * blockDim.x？

你想想看，每个线程都需要知道自己处理哪个元素。线程在自己的块内有编号 threadIdx.x，块在整个网格内有编号 blockIdx.x，每个块有 blockDim.x 个线程。所以全局ID就是 blockIdx.x * blockDim.x + threadIdx.x。

4.5 错误检查：cudaGetLastError

GPU编程最烦人的就是调试。核函数调用失败，CPU端不会直接报错。你得主动检查。

cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
    printf("CUDA error: %s\n", cudaGetErrorString(err));
}

我建议你在每个CUDA API调用后都检查一下错误。尤其是 cudaMalloc 和 cudaMemcpy。我曾经在项目里漏了检查，结果核函数参数传错了，跑出来的结果全是错的。查了两天才发现是设备内存没分配成功。

#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_CHECK(cudaMalloc(&d_A, size));
CUDA_CHECK(cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice));

4.6 完整流程总结

一个完整的CUDA程序，流程其实就五步：

1. 分配主机内存 — malloc 或 new
2. 分配设备内存 — cudaMalloc
3. 拷贝数据到设备 — cudaMemcpy HostToDevice
4. 调用核函数 — <<<grid, block>>>
5. 拷贝结果回主机 — cudaMemcpy DeviceToHost

最后别忘了释放内存。CPU端用 free，GPU端用 cudaFree。

这个流程你以后写任何CUDA程序都会用到。说白了，就是「搬数据上去 → 算 → 搬数据下来」。理解了这一点，你就掌握了GPU编程的核心。