第一章：CUDA入门——GPU架构概览、CUDA编程模型、Hello World从CPU到GPU

大家好，我是你们的讲师。今天咱们正式开讲《CUDA核心与算子优化实战》的第一章。

说实话，每次讲这一章，我都挺感慨的。记得我第一次接触CUDA，是在一个图像处理的项目里。那时候CPU算不过来，老板丢给我一块显卡说：“你试试这个。”我对着文档折腾了三天，才让第一个“Hello World”跑起来。嗯，那种感觉，就像打开了一扇新世界的大门。

好，废话不多说。咱们今天的目标很明确：搞懂GPU凭什么比CPU快，CUDA到底怎么编程，以及亲手把第一个程序从CPU搬到GPU上。

1.1 GPU架构概览：为什么它这么快？

先问大家一个问题：CPU和GPU，本质区别在哪？

你可能会说：“CPU核少，GPU核多。”对，但不全对。

我习惯用一个比喻：CPU就像几个博士生，什么难题都能解，但数量少。GPU就像几千个小学生，只会简单的加减法，但人多力量大。

看这张图，你就明白了：

说白了，CPU是“大核少核”，GPU是“小核多核”。CPU擅长处理复杂的逻辑分支，GPU擅长做大量重复的数学运算。

GPU里最核心的单元叫SM（Streaming Multiprocessor，流式多处理器）。每个SM里有一堆CUDA核心，还有共享内存、寄存器、调度器这些东西。你写的CUDA程序，最终就是被分配到这些SM上执行的。

关键概念：

SM（流式多处理器）：GPU的基本计算单元，相当于一个“小CPU集群”
CUDA Core：SM里的基本运算单元，执行简单的算术逻辑
Global Memory：显存，所有线程都能访问，但速度慢
Shared Memory：同一个Block内的线程共享，速度快，但容量小

我的经验：刚开始学的时候，别纠结SM内部的具体电路。你只需要记住：GPU靠“人多”取胜。一个任务如果能拆成成千上万个独立的小任务，那就适合GPU干。

1.2 CUDA编程模型：主机与设备

好，架构看完了。接下来咱们聊聊怎么编程。

CUDA的编程模型，说白了就是“CPU+GPU”协同工作。CPU叫主机（Host），GPU叫设备（Device）。

流程大概是这样的：

CPU把数据从内存拷贝到显存
CPU告诉GPU：“干活！”（启动Kernel）
GPU的成千上万个线程同时执行
GPU干完活，CPU把结果从显存拷回内存

你想想看，这像什么？像不像一个项目经理（CPU）把任务分给一群工人（GPU线程），工人干完活再汇报结果？

这里有个重要的概念叫线程层次结构：

层次	说明	类比
Thread（线程）	最小的执行单元	一个工人
Block（线程块）	一组线程，可共享内存	一个班组
Grid（网格）	一组Block，构成整个任务	整个工地

每个线程都有一个唯一的ID，通过 threadIdx 和 blockIdx 来区分。这样每个线程就知道自己该处理哪部分数据了。

注意：我曾经见过新手把Block和Grid的维度搞反，结果数据全乱套了。记住：blockIdx 是Block的索引，threadIdx 是线程在Block内的索引。

1.3 Hello World：从CPU到GPU

理论说完了，咱们动手写代码。

先看一个最简单的例子：给数组每个元素加1。

CPU版本：

#include <stdio.h>

#define N 10

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

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

这个很简单，一个for循环搞定。但数据量大了怎么办？比如1000万个元素？CPU就得一个一个算，慢得很。

GPU版本：

#include <stdio.h>

#define N 10

// GPU上执行的函数，叫Kernel
__global__ void add_one_gpu(int *a) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        a[idx] = a[idx] + 1;
    }
}

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

看到了吗？关键区别在于：

__global__ 关键字：告诉编译器这个函数在GPU上运行
<<<1, N>>>：启动配置，第一个参数是Block数，第二个参数是每个Block的线程数
threadIdx.x：当前线程在Block内的索引
blockIdx.x：当前Block在Grid内的索引
blockDim.x：每个Block的线程数

每个线程都执行同样的代码，但通过 threadIdx.x + blockIdx.x * blockDim.x 计算出自己负责的元素下标。这样，N个线程同时干活，速度自然就快了。

避坑指南：我曾经犯过一个低级错误——忘了加 cudaMemcpy 把数据拷回CPU，结果打印出来的全是0。记住：GPU和CPU的内存是隔离的，数据不会自动同步！

1.4 编译与运行

写好了代码，怎么跑？

用 nvcc 编译器：

nvcc -o hello_gpu hello_gpu.cu
./hello_gpu

注意文件后缀是 .cu，不是 .c 或 .cpp。nvcc 会自动识别哪些代码跑在GPU上，哪些跑在CPU上。

如果你看到输出 1 2 3 4 5 6 7 8 9 10，恭喜你，你的第一个CUDA程序跑通了！

总结一下本章核心：

GPU靠大量简单核心并行计算，适合数据并行任务
CUDA编程模型是Host（CPU）+ Device（GPU）协同
线程层次：Grid → Block → Thread
Kernel用 <<<grid, block>>> 启动
数据需要显式地在Host和Device之间拷贝

嗯，第一章就到这里。代码不多，但背后的思想很重要。你想想看，从单核CPU到多核CPU，再到GPU的成千上万个核心，这其实就是并行计算的发展史。

下一章咱们会深入聊聊线程、Block和Grid的配置，以及怎么写出更高效的Kernel。到时候见。

想要更多资料，联系可可

📚 课程目录

第一章：CUDA入门——GPU架构概览、CUDA编程模型、Hello World从CPU到GPU

1.1 GPU架构概览：为什么它这么快？

1.2 CUDA编程模型：主机与设备

1.3 Hello World：从CPU到GPU

1.4 编译与运行