3、第一个CUDA程序：编写向量加法Kernel，理解<<<>>>语法，编译与运行。

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

很多初学者学CUDA，第一个坎儿就是看不懂那个尖括号语法。我当年第一次看到 <<< >>> 的时候，心里直犯嘀咕：这玩意儿到底是啥？别急，咱们一步步来。

3.1 向量加法：CPU版本回顾

先看一个最朴素的C语言向量加法。假设有两个数组 A 和 B，长度都是 N，我们要把对应元素相加，结果存到 C 里。

// CPU版本
void vecAdd_cpu(float *A, float *B, float *C, int N) {
    for (int i = 0; i < N; i++) {
        C[i] = A[i] + B[i];
    }
}

这个循环，说白了就是挨个算。N 等于 100 万的时候，CPU 得老老实实跑 100 万次。GPU 怎么干？它把 100 万个线程同时扔上去，每个线程只算一个元素。

3.2 第一个CUDA Kernel

Kernel 就是跑在 GPU 上的函数。写法上有个标志：前面加个 __global__。

// GPU版本：向量加法Kernel
__global__ void vecAdd_gpu(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];
    }
}

这里有个关键点：threadIdx.x、blockIdx.x、blockDim.x 这三个变量。它们不是我们手动赋值的，而是CUDA运行时自动提供的。每个线程拿到的值都不一样，从而实现了「并行」。

3.3 尖括号语法：<<< >>> 到底在传什么？

好，重点来了。调用 Kernel 的时候，写法是这样的：

vecAdd_gpu<<<gridDim, blockDim>>>(d_A, d_B, d_C, N);

尖括号里两个参数，分别指定了「网格维度」和「块维度」。我习惯这么理解：

• blockDim：一个块里有多少个线程。比如 256 个。
• gridDim：一个网格里有多少个块。比如 N / 256 + 1 个。

那么总线程数 = gridDim * blockDim。每个线程通过 threadIdx.x + blockIdx.x * blockDim.x 算出自己的全局索引。

3.4 完整的主机端代码

光有 Kernel 不行，还得有主机端代码来分配显存、拷贝数据、启动 Kernel。看下面这个完整示例：

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

#define N 1000000

int main() {
    float *h_A, *h_B, *h_C;      // 主机端指针
    float *d_A, *d_B, *d_C;      // 设备端指针
    size_t size = N * sizeof(float);

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

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

    // 2. 分配设备显存
    cudaMalloc(&d_A, size);
    cudaMalloc(&d_B, size);
    cudaMalloc(&d_C, size);

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

    // 4. 启动Kernel
    int blockSize = 256;
    int gridSize = (N + blockSize - 1) / blockSize;
    vecAdd_gpu<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);

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

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

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

    return 0;
}

3.5 编译与运行

CUDA 程序用 nvcc 编译，不是 gcc。命令很简单：

nvcc -o vecAdd vecAdd.cu
./vecAdd

如果你看到类似这样的输出：

h_C[0] = 1000000.000000
h_C[1] = 1000000.000000
...

恭喜你，第一个 CUDA 程序跑通了！

3.6 常见错误与避坑指南

嗯，这里我总结几个新手最容易踩的坑：

错误类型	现象	原因
段错误	程序崩溃	指针没分配显存就用了
结果全0	输出全是0	忘了 cudaMemcpy 拷回主机
索引越界	部分结果正确，部分错误	gridSize 算小了，有些元素没被处理
编译报错	找不到 cuda_runtime.h	环境变量没配好，或者用了 gcc 而不是 nvcc

3.7 小结

这一章我们干了三件事：

• 写了一个最简单的向量加法 Kernel
• 理解了 <<<gridDim, blockDim>>> 的含义
• 走通了编译和运行的完整流程

说白了，CUDA 编程就是把一个 for 循环拆成无数个线程并行执行。你想想看，是不是这个理儿？下一章我们会深入聊聊线程、块、网格的更多细节，以及怎么用 cudaError_t 来捕获错误。到时候你会发现，调试 CUDA 程序其实没那么可怕。