一、CUDA入门：GPU计算革命、CUDA架构概览、CUDA编程模型、Hello World从CPU到GPU

1.1 GPU计算革命：为什么是GPU？

说实话，我第一次接触GPU计算是在2012年。那时候我在做图像处理，CPU跑一个滤波算法要好几秒。同事说「试试CUDA」，我当时还纳闷：显卡不是打游戏的吗？

结果你猜怎么着？同样的算法，GPU只用了30毫秒。那一刻我就明白了——计算的世界要变天了。

为什么会这样？说白了，CPU和GPU的设计哲学完全不同：

CPU：像几个全能博士，什么都能干，但数量少（4-16核）
GPU：像成千上万的流水线工人，只会简单操作，但人多力量大

我习惯用一个比喻：CPU是法拉利，GPU是卡车队。运一箱货法拉利快，但运一万箱？卡车队完胜。深度学习训练就是「运一万箱」的场景——成千上万的矩阵运算，天然适合GPU。

核心认知：GPU不是用来取代CPU的，它们是搭档。CPU负责调度、控制流，GPU负责大规模并行计算。我在项目中见过有人非要用GPU跑单线程任务，结果比CPU还慢——嗯，工具要用对地方。

1.2 CUDA架构概览：一张图看懂

先看这张图，我花了不少心思画的。它把CUDA的硬件和软件对应关系讲清楚了：

这张图我建议你多看几遍。左边是硬件，右边是软件，中间虚线就是「映射关系」。我当年学CUDA时，就是靠这张图打通了任督二脉。

我的经验：刚开始别纠结每个细节。先记住「Grid→GPU、Block→SM、Thread→CUDA Core」这个对应关系。剩下的，写几个程序自然就懂了。

1.3 CUDA编程模型：三步走

CUDA编程其实就三步，我总结为「搬运-计算-搬回」：

搬运数据到GPU：用 cudaMalloc 分配显存，cudaMemcpy 把数据从CPU拷贝到GPU
在GPU上计算：写一个核函数（Kernel），用 <<<grid, block>>> 启动
把结果搬回CPU：再用 cudaMemcpy 把结果拷回来

你想想看，这和「外包」是不是很像？你把活派给外包团队（GPU），他们干完活再把成果交回来。中间怎么分工、怎么协作，就是你要设计的。

核函数长什么样？

核函数就是跑在GPU上的函数。有两个标志：

__global__：表示这是一个核函数，从CPU调用，在GPU上执行
__device__：表示这是一个设备函数，只能从GPU上调用

避坑指南：我曾经犯过一个低级错误——在核函数里用 printf 调试。嗯，核函数里确实可以用 printf，但输出顺序是乱的，因为成千上万个线程同时打印。后来我改用 cuda-gdb 调试，舒服多了。

1.4 Hello World：从CPU到GPU

光说不练假把式。咱们写个真正的Hello World——让GPU给数组每个元素加1。

CPU版本（先看基线）

#include <stdio.h>

#define N 10

void add_one_cpu(int *a, int n) {
    for (int i = 0; i < n; i++) {
        a[i] += 1;
    }
}

int main() {
    int a[N] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9};
    
    add_one_cpu(a, N);
    
    for (int i = 0; i < N; i++) {
        printf("%d ", a[i]);
    }
    printf("\n");
    return 0;
}

GPU版本（CUDA改写）

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

#define N 10

// 核函数：每个线程处理一个元素
__global__ void add_one_gpu(int *a, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        a[idx] += 1;
    }
}

int main() {
    int a[N] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9};
    int *d_a;
    
    // 第一步：在GPU上分配显存
    cudaMalloc(&d_a, N * sizeof(int));
    
    // 第二步：把数据从CPU拷贝到GPU
    cudaMemcpy(d_a, a, N * sizeof(int), cudaMemcpyHostToDevice);
    
    // 第三步：启动核函数（1个Grid，1个Block，N个线程）
    add_one_gpu<<<1, N>>>(d_a, N);
    
    // 第四步：把结果从GPU拷回CPU
    cudaMemcpy(a, d_a, N * sizeof(int), cudaMemcpyDeviceToHost);
    
    // 打印结果
    for (int i = 0; i < N; i++) {
        printf("%d ", a[i]);
    }
    printf("\n");
    
    // 第五步：释放显存
    cudaFree(d_a);
    
    return 0;
}

运行结果：1 2 3 4 5 6 7 8 9 10

和CPU版本一模一样。但背后的逻辑完全不同——CPU是串行循环，GPU是并行启动N个线程同时干活。

注意：核函数里的 threadIdx.x、blockIdx.x、blockDim.x 是CUDA内置变量。它们告诉每个线程「你是谁、你在哪」。我刚开始总搞混，后来记住一句话：threadIdx.x 是线程在块内的编号，blockIdx.x 是块在网格内的编号。

代码逐行拆解

代码行	作用	我的注释
`__global__ void add_one_gpu(...)`	定义核函数	__global__ 告诉编译器这是GPU函数
`int idx = threadIdx.x + blockIdx.x * blockDim.x;`	计算全局线程ID	这是CUDA最经典的「寻址公式」，背下来
`if (idx < n)`	边界检查	防止线程数超过数组长度，我吃过亏
`cudaMalloc(&d_a, N * sizeof(int))`	分配显存	和malloc很像，但分配的是GPU内存
`cudaMemcpy(..., cudaMemcpyHostToDevice)`	数据从CPU到GPU	方向别搞反了，我见过有人拷了个寂寞
`<<<1, N>>>`	配置执行配置	第一个参数是Grid大小，第二个是Block大小
`cudaFree(d_a)`	释放显存	不释放会内存泄漏，和free一样

调试小技巧：在每个CUDA API调用后加 cudaGetLastError() 检查错误。我习惯写个宏：

#define CUDA_CHECK(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            printf("CUDA Error: %s at %s:%d\n", \
                   cudaGetErrorString(err), __FILE__, __LINE__); \
            exit(1); \
        } \
    } while(0)

这样出错了能立刻定位。别问我怎么知道的——都是血泪教训。

1.5 本章小结

咱们这一章走完了从「为什么用GPU」到「写出第一个CUDA程序」的全过程。核心就三件事：

GPU适合大规模并行计算，尤其是深度学习里的矩阵运算
CUDA架构是分层的：Grid→Block→Thread，对应GPU→SM→CUDA Core
编程模型三步走：搬数据→算数据→搬回来

说实话，我第一次跑通CUDA程序时，看到控制台输出「1 2 3 4 5 6 7 8 9 10」，激动得差点跳起来。那种「让成千上万个线程同时为我工作」的感觉，真的很爽。

下一章咱们会深入线程层次和内存模型——这才是CUDA性能优化的真正战场。到时候你会发现，今天写的这个Hello World，其实只是冰山一角。

想要更多资料，联系可可

📚 课程目录

一、CUDA入门：GPU计算革命、CUDA架构概览、CUDA编程模型、Hello World从CPU到GPU

1.1 GPU计算革命：为什么是GPU？

1.2 CUDA架构概览：一张图看懂

1.3 CUDA编程模型：三步走

核函数长什么样？

1.4 Hello World：从CPU到GPU

CPU版本（先看基线）

GPU版本（CUDA改写）

代码逐行拆解

1.5 本章小结