3. 第一个CUDA程序：向量加法实现、核函数编写、内存分配与数据传输

好，咱们直接动手写第一个CUDA程序。我见过太多人一上来就啃理论，结果连个向量加法都跑不通。说实话，CUDA入门最忌讳的就是「光看不练」。今天我们就拿向量加法开刀，把整个流程走一遍。

3.1 向量加法：CUDA版的"Hello World"

向量加法在CPU上太简单了：

void vec_add_cpu(float *a, float *b, float *c, int n) {
    for (int i = 0; i < n; i++) {
        c[i] = a[i] + b[i];
    }
}

但到了GPU上，事情就变得有意思了。你想想看，GPU有几千个核心，每个核心只算一个加法，那速度得多快？

不过这里有个关键点：每个线程必须知道自己该算哪个元素。这就是CUDA内置变量 threadIdx.x、blockIdx.x、blockDim.x 的用武之地。

3.2 核函数怎么写？

核函数就是跑在GPU上的函数。写法上有个标志性的东西——__global__ 修饰符。

__global__ void vec_add_kernel(float *d_a, float *d_b, float *d_c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 边界检查——这个绝对不能省！
    if (idx < n) {
        d_c[idx] = d_a[idx] + d_b[idx];
    }
}

嗯，这里要注意几点：

• 返回类型必须是void。核函数不能有返回值，这是硬性规定。
• 参数都是指针。因为GPU不能直接访问CPU内存，你得传设备内存的地址。
• 边界检查。我曾经犯过这个错——线程总数比数据量大，结果写到了别人的内存里，查了三天才找到原因。

3.3 内存分配：CPU vs GPU

CUDA编程里，内存是分两家的：

内存类型	分配函数	释放函数	所在位置
主机内存（CPU）	malloc / new	free / delete	系统RAM
设备内存（GPU）	cudaMalloc	cudaFree	显存

我个人习惯在分配设备内存时，变量名前面加个 d_ 前缀。比如 d_a、d_b、d_c。这样一眼就能看出哪些数据在GPU上，避免搞混。

3.4 数据传输：CPU ↔ GPU

数据不会自己飞过去。你得用 cudaMemcpy 手动搬运：

// 主机到设备
cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);

// 设备到主机
cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);

这里有个小细节：cudaMemcpy 是同步的。也就是说，它执行完之前，CPU会一直等着。但核函数是异步的——你调用完核函数，CPU马上就继续往下跑了。

为什么会这样？因为GPU是个独立的处理器，CPU把任务丢给它，自己就可以干别的事。但这也意味着，如果你想在核函数执行完后立刻读取结果，必须加个同步：

vec_add_kernel<<<grid, block>>>(d_a, d_b, d_c, n);
cudaDeviceSynchronize();  // 等GPU干完活
// 现在可以安全地拷贝结果了

3.5 完整代码示例

来，我们把所有东西拼起来：

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

__global__ void vec_add_kernel(float *d_a, float *d_b, float *d_c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        d_c[idx] = d_a[idx] + d_b[idx];
    }
}

int main() {
    int n = 1024;
    size_t size = n * sizeof(float);
    
    // 1. 分配主机内存
    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] = i * 1.0f;
        h_b[i] = i * 2.0f;
    }
    
    // 2. 分配设备内存
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, size);
    cudaMalloc(&d_b, size);
    cudaMalloc(&d_c, size);
    
    // 3. 数据传输：CPU → GPU
    cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);
    
    // 4. 启动核函数
    int threads_per_block = 256;
    int blocks_per_grid = (n + threads_per_block - 1) / threads_per_block;
    
    vec_add_kernel<<<blocks_per_grid, threads_per_block>>>(d_a, d_b, d_c, n);
    
    // 5. 同步并拷贝结果回CPU
    cudaDeviceSynchronize();
    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.6 常见错误与调试

写第一个CUDA程序，你大概率会遇到这些问题：

• 核函数没执行：检查一下是不是忘了加 <<<grid, block>>> 启动配置。
• 结果全是0：八成是 cudaMemcpy 方向搞反了。我犯过这个错，把 cudaMemcpyHostToDevice 写成了 cudaMemcpyDeviceToHost。
• 程序崩溃：先检查边界条件。用 cuda-memcheck 工具跑一下，它能帮你定位越界访问。

说实话，CUDA调试比CPU调试麻烦不少。你不能直接 printf（虽然新版本支持了，但性能很差）。我的建议是：先在CPU上把逻辑跑通，再移植到GPU。这样至少能保证算法本身没问题。

好了，第一个CUDA程序就讲到这里。你把这个向量加法跑通了，CUDA的大门就算正式打开了。下一章我们聊聊线程、块和网格的配置策略——这可是性能调优的重头戏。