2、CUDA编程入门：向量加法、核函数、线程与块配置

好，咱们正式开始动手写GPU代码。

很多人觉得CUDA门槛高，说白了，就是把一段可以并行的代码，告诉GPU怎么拆成小块、怎么调度。你想想看，CPU是几个强壮的大厨，GPU是一群切菜工。你要让切菜工干活，就得把任务切碎，分给大家。这就是CUDA的核心思维。

为什么需要理解GPU编程？

我最早接触GPU加速时，心里也犯嘀咕：Python调调库不香吗？直到有一次项目，基因组注释模型跑一轮训练要72小时，光等结果就让人崩溃。

后来我把关键算子用CUDA重写，训练时间压缩到8小时以内。嗯，从那以后，我再也不敢小看这层“底层功夫”了。

对于基因组注释这种计算密集型任务，GPU优化不是锦上添花，是雪中送炭。你迟早要面对它。

CUDA编程模型：主机与设备

先明确两个概念：

• 主机（Host）：CPU + 系统内存，负责控制逻辑和串行任务。
• 设备（Device）：GPU + 显存，负责大规模并行计算。

数据必须从主机复制到设备，计算完再复制回来。这个搬运过程，往往是性能瓶颈。我在项目中就吃过这个亏——数据传输没管好，GPU算得再快也白搭。

第一个CUDA程序：向量加法

咱们用向量加法来切入。这是CUDA界的"Hello, World"。

CPU版本大家都会：

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

GPU版本怎么写？我直接给代码，你感受一下：

__global__ void vector_add_gpu(float *a, float *b, float *c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

看到了吗？__global__ 这个修饰符，告诉编译器这是核函数，跑在设备端。
核心变化是：不再写循环，而是通过线程索引来定位每个元素。

核函数：GPU上的并行单元

核函数（Kernel）就是你在GPU上执行的函数。有几个要点：

• 用 __global__ 声明，返回类型必须是 void。
• 不能调用C++标准库（比如 printf 除外）。
• 所有线程执行同一段代码，但操作不同的数据。

我个人习惯把核函数写得尽量简短。一个核函数只做一件事，容易调试，也容易优化。

调用核函数的语法也特别：

vector_add_gpu<<>>(d_a, d_b, d_c, n);

尖括号里的两个参数，就是我们要重点说的——线程与块配置。

线程与块配置：三层的世界

CUDA的线程组织分三层：

层级	说明	常见范围
线程（Thread）	最小的执行单元	单条指令
线程块（Block）	一组线程，可以共享内存、同步	128~1024个线程
网格（Grid）	一组线程块	根据数据量决定

你想想看，这个结构其实很直观：网格是工厂，线程块是车间，线程是工人。

内置变量帮你定位：

• threadIdx.x：线程在线程块内的索引。
• blockIdx.x：线程块在网格内的索引。
• blockDim.x：每个线程块的线程数量。

那个经典的索引公式：

int i = blockIdx.x * blockDim.x + threadIdx.x;

个人觉得，理解这个公式就算入门了一半。

怎么配置？别瞎填

我曾经在配置线程块大小时吃过亏。一开始图省事，直接写：

int block_size = 256;
int grid_size = (n + block_size - 1) / block_size;
vector_add_gpu<<>>(d_a, d_b, d_c, n);

当时没注意 n 可能不是 block_size 整数倍，结果数组末尾越界。所以核函数里一定要加边界检查：

if (i < n) { ... }

完整代码与编译

我把完整的向量加法代码贴出来：

#include <stdio.h>

__global__ void vector_add(float *a, float *b, float *c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

int main() {
    int n = 1 << 20;
    size_t size = n * sizeof(float);

    float *h_a, *h_b, *h_c;
    float *d_a, *d_b, *d_c;

    // 主机分配
    h_a = (float*)malloc(size);
    h_b = (float*)malloc(size);
    h_c = (float*)malloc(size);

    // 设备分配
    cudaMalloc(&d_a, size);
    cudaMalloc(&d_b, size);
    cudaMalloc(&d_c, size);

    // 初始化并拷贝到设备
    for (int i = 0; i < n; i++) {
        h_a[i] = 1.0f; h_b[i] = 2.0f;
    }
    cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);

    // 配置并调用核函数
    int block_size = 256;
    int grid_size = (n + block_size - 1) / block_size;
    vector_add<<>>(d_a, d_b, d_c, n);

    // 拷贝结果回主机
    cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);

    // 验证
    for (int i = 0; i < n; i++) {
        if (h_c[i] != 3.0f) {
            printf("Error at %d\n", i);
            break;
        }
    }
    printf("Done\n");

    // 清理
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    free(h_a); free(h_b); free(h_c);
    return 0;
}

编译命令：

nvcc -o vector_add vector_add.cu

本章知识结构

下面这张图总结了CUDA编程入门阶段你需要掌握的核心概念和它们之间的关系：

避坑指南

最后，分享几个我踩过的坑：

• 忘记cudaMalloc的返回值检查——我曾经在显存不足时没检查，程序直接崩了，排查半天。用cudaError_t判断一下，省心很多。
• 核函数里的printf只在debug模式下可靠——正式跑大数据时别依赖它，用cuda-memcheck更靠谱。
• 同步问题——核函数是异步的，如果后面紧跟CPU操作，要加cudaDeviceSynchronize()。我一开始没加，结果拿到的数据是错的。