向量加法实战：从CPU到GPU的第一次亲密接触

说实话，向量加法这个例子，在CUDA学习里就跟编程界的"Hello World"一样经典。我当年刚接触GPU编程时，第一个跑通的程序就是它。别看它简单，这里面藏着的门道可不少——从CPU串行到GPU并行，从带宽计算到性能分析，每一步都能让你对GPU架构有更深的理解。

先来个CPU版本，找找感觉

我们先写个最朴素的CPU版本。两个长度为N的浮点数组相加，这活儿CPU干起来其实挺轻松的。

// CPU版本：向量加法
void vecAddCPU(const float* a, const float* b, float* c, int n) {
    for (int i = 0; i < n; i++) {
        c[i] = a[i] + b[i];
    }
}

这个循环，说白了就是一条一条地加。N=100万的时候，在我的笔记本上大概跑个2-3毫秒。嗯，看起来还行？别急，等会儿看看GPU的表现，你就知道什么叫"降维打击"了。

GPU版本：让成千上万个线程一起干活

GPU的思路完全不同。它不搞"一个一个来"，而是"大家一起上"。每个线程负责一个元素的计算，成千上万个线程同时开工。

// GPU版本：向量加法kernel
__global__ void vecAddGPU(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

这里有个小细节我特别想强调：if (i < n)这个边界检查。为什么要有它？因为数据量不一定能被线程块大小整除。比如你有1000个元素，每个块256个线程，那最后一个块只有232个线程干活，剩下的24个线程得闲着。不加这个判断，它们就会访问越界内存——我曾经因为这个bug排查了一整个下午，血的教训啊。

主函数调用：把数据搬到GPU上去

GPU不能直接访问CPU的内存，得先把数据拷过去。这个"搬运"过程，往往是性能瓶颈所在。

int main() {
    int n = 1 << 20;  // 1048576个元素
    size_t bytes = n * sizeof(float);

    // 分配主机内存
    float *h_a, *h_b, *h_c;
    h_a = (float*)malloc(bytes);
    h_b = (float*)malloc(bytes);
    h_c = (float*)malloc(bytes);

    // 初始化数据
    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, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);

    // 拷贝数据到GPU
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);

    // 启动kernel
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    vecAddGPU<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, n);

    // 拷回结果
    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);

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

    return 0;
}

小提示：线程块大小选256是个经验值。我试过128、256、512，在我的卡上256表现最好。为什么？因为每个SM（流多处理器）能同时调度的线程数有限，256刚好能填满它的胃口。你可以在自己的卡上试试不同值，看看效果。

性能对比：GPU到底快了多少？

我用NVIDIA Nsight做了个简单的计时。N=100万，结果如下：

版本	执行时间	加速比
CPU (单核)	2.3 ms	1x
GPU (不含数据传输)	0.035 ms	~65x
GPU (含数据传输)	0.52 ms	~4.4x

看到没？光算的话，GPU快了65倍。但加上数据搬运，加速比就掉到4.4倍了。这就是我常说的：数据传输是GPU计算的阿喀琉斯之踵。

核心结论：向量加法这种计算密集型任务，GPU的计算能力远强于CPU。但数据搬运的开销不容忽视。如果你的计算量太小，可能连搬运的时间都赚不回来——这就是所谓的"计算密度"问题。

带宽计算：看看你的GPU有多"宽"

带宽，说白了就是GPU读写数据的速度。我们来算算这个例子里的实际带宽：

每次向量加法，需要读取两个数组（a和b），写入一个数组（c）。每个元素4字节（float），总共3次访问。N=100万时：

数据量 = 3 * 1,048,576 * 4 = 12,582,912 字节 ≈ 12 MB
带宽 = 数据量 / 时间 = 12 MB / 0.035 ms ≈ 343 GB/s

我用的这张卡理论带宽是400 GB/s，实际跑出来343 GB/s，利用率大概85%。这个数字还算不错。如果你跑出来利用率低于50%，那就要检查一下内存访问模式了——是不是有大量的非对齐访问？是不是每个线程访问的数据太分散？

注意：带宽计算时，一定要用kernel的执行时间，而不是包含数据传输的总时间。否则算出来的带宽会偏低，误导你对kernel性能的判断。我曾经犯过这个错，还以为自己写的kernel很慢，结果发现是数据传输占了大部分时间。

画个图，理清整个流程

下面这张图展示了向量加法从CPU到GPU的完整数据流。我习惯用这种图来跟团队沟通，一目了然。

这张图把整个流程拆得很清楚：数据从CPU内存搬到GPU显存，GPU算完再搬回来。中间的kernel执行时间虽然短，但两头的搬运时间加起来，才是你真正需要关心的总时间。

几个实战中踩过的坑

最后分享几个我实际项目中遇到的问题，希望能帮你少走弯路：

内存分配失败：cudaMalloc返回错误？先检查显存够不够。我遇到过有人一次分配2GB，结果卡只有1.5GB显存，直接崩了。
数据没同步：kernel启动是异步的。如果你在kernel返回后立刻拷数据，可能拿到的是旧值。记得加cudaDeviceSynchronize()。
索引算错：blockIdx.x * blockDim.x + threadIdx.x这个公式，写错一个符号就是灾难。我见过有人写成blockIdx.x + blockDim.x * threadIdx.x，结果所有线程访问同一个地址。
边界条件：数据量不是线程块大小的整数倍时，一定要做边界检查。否则最后一个块会访问越界，轻则数据错乱，重则程序崩溃。

向量加法虽然简单，但它是理解GPU编程模型的最佳起点。把这里面的概念吃透了，后面学矩阵乘法、卷积这些复杂算子，你会觉得顺滑很多。

想要更多资料，联系可可

📚 课程目录

向量加法实战：从CPU到GPU的第一次亲密接触

先来个CPU版本，找找感觉

GPU版本：让成千上万个线程一起干活

主函数调用：把数据搬到GPU上去

性能对比：GPU到底快了多少？

带宽计算：看看你的GPU有多"宽"

画个图，理清整个流程

几个实战中踩过的坑