1、CUDA入门：GPU架构概览、CUDA编程模型、第一个CUDA程序（向量加法）

大家好，欢迎来到这门实战课的第一章。

说实话，我见过太多同学一上来就写 cudaMalloc，结果跑出来的结果全是错的。为什么？因为根本不懂 GPU 在背后怎么干活。所以这一章，咱们先把地基打牢。

1.1 GPU 架构概览：它凭什么比 CPU 快？

先问个问题：CPU 和 GPU 到底差在哪？

CPU 是「通才」。它要处理操作系统、网络请求、分支预测、乱序执行……说白了，它什么活都能干，但一次只能干几件。GPU 是「专才」。它脑子简单，但数量多。一个 GPU 里有几千个计算核心，每个核心只会做加减乘除，但一起上，算力就爆了。

我当年第一次接触 GPU 编程时，总觉得它跟 CPU 差不多。后来调一个矩阵乘法，CPU 版本跑了 3 秒，GPU 版本跑了 30 秒——我差点把电脑砸了。后来才发现，是我把数据拷贝来回搞了 100 次。嗯，这就是不懂架构的代价。

核心差异一句话： CPU 是「少而强」，GPU 是「多而弱」。但「多」到一定程度，就能碾压「强」。

具体到 NVIDIA 的 GPU 架构，有几个关键概念你必须记住：

  • SM（Streaming Multiprocessor）：GPU 里的计算单元组。一个 SM 包含多个 CUDA Core。
  • CUDA Core：最基础的执行单元，做浮点运算的。
  • Global Memory：全局显存，所有线程都能访问，但慢。
  • Shared Memory：共享内存，同一个 Block 内的线程共享，快得多。
  • Warp：32 个线程为一组，一起执行指令。这是 GPU 调度的基本单位。

你想想看，如果一个 Warp 里的 32 个线程走了不同的分支（if-else），那 GPU 就得把两条路都跑一遍，浪费一半算力。这就是所谓的「warp divergence」。我在项目里见过有人写了个 switch-case，32 个线程走了 8 个分支，性能直接掉了 70%。

1.2 CUDA 编程模型：你该怎么指挥这些核心？

CUDA 的编程模型，说白了就是「三句话」：

  1. 把数据从 CPU 搬到 GPU。
  2. 在 GPU 上启动成千上万个线程干活。
  3. 把结果从 GPU 搬回 CPU。

这里有个关键概念叫 Grid-Block-Thread 层次结构：

  • 一个 Grid 包含多个 Block。
  • 一个 Block 包含多个 Thread。
  • 线程之间通过 threadIdx、blockIdx、blockDim 来定位自己。

举个例子：你要处理一个 1024 个元素的数组。你可以开 4 个 Block，每个 Block 里 256 个线程。那么：

  • 线程 (blockIdx=0, threadIdx=0) 处理元素 0
  • 线程 (blockIdx=0, threadIdx=1) 处理元素 1
  • ……
  • 线程 (blockIdx=1, threadIdx=0) 处理元素 256

计算公式很简单：int idx = blockIdx.x * blockDim.x + threadIdx.x;

我的习惯： 写 kernel 之前，先在纸上画一下线程索引怎么映射到数据。画清楚了再写代码，能省一半调试时间。

1.3 第一个 CUDA 程序：向量加法

好了，理论说完了，咱们直接上手写代码。向量加法是最简单的 kernel，没有之一。

先看 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 版本：

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

注意那个 if (idx < n)。为什么要加？因为数组长度不一定能被线程数整除。多出来的线程什么都不干，直接返回。这叫「边界检查」，写 kernel 时必须养成习惯。

接下来是主函数：

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

    // 1. 分配主机内存
    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] = i * 1.0f;
        h_b[i] = i * 2.0f;
    }

    // 2. 分配设备内存
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);

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

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

    // 5. 拷贝结果回 CPU
    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);

    // 6. 验证结果
    for (int i = 0; i < 10; i++) {
        printf("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;
}
我曾经踩过的坑： 第一次写 CUDA 程序时，我忘了调用 cudaMemcpy 把结果拷回来，直接打印了 h_c。结果全是 0。我当时还以为是 kernel 没跑起来，查了整整一下午。后来才发现，GPU 算完的结果还在显存里，根本没回来。

1.4 编译与运行

编译命令很简单：

nvcc -o vec_add vec_add.cu

然后运行：

./vec_add

如果一切正常，你会看到前 10 个元素的结果。比如 c[0] = 0.0，c[1] = 3.0，c[2] = 6.0……

调试小技巧： 在 kernel 后面加一句 cudaDeviceSynchronize();，可以捕获异步错误。我习惯在每个 kernel 调用后都加一句错误检查： 
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
    printf("CUDA Error: %s\n", cudaGetErrorString(err));
}
这能帮你快速定位是哪个 kernel 出了问题。

1.5 本章小结

这一章我们干了三件事：

  • 搞懂了 GPU 为什么快——靠的是堆核心数。
  • 学会了 CUDA 的线程模型——Grid、Block、Thread 怎么映射数据。
  • 亲手写了一个向量加法的 kernel，并且跑通了。

说实话，向量加法虽然简单，但它包含了 CUDA 编程的所有核心步骤：内存分配、数据拷贝、kernel 启动、结果回传、资源释放。后面所有复杂的算子，都是在这个框架上叠加的。

下一章，我们会深入讲 内存管理。你会学到为什么 cudaMalloc 这么慢，以及怎么用 cudaMallocManaged 偷懒。嗯，到时候你会感谢我的。

先写到这。去把代码跑通，有问题随时翻这一章。

← 上一章 📖 返回目录 下一章 →