1. GPU计算基础：GPU架构概述、CUDA编程模型入门、GPU与CPU的协同工作方式
各位同学，欢迎来到第一节课。
说实话，每次讲GPU计算基础，我都会想起自己刚入行时的困惑。那时候我对着NVIDIA的文档看了三天，愣是没搞明白为什么GPU能比CPU快那么多。后来亲手调了几个模型，踩了不少坑，才算真正摸到门道。
今天这节课，我们就从最底层开始，把GPU的底裤扒开看看。
1.1 GPU架构概述：为什么它天生适合做感知？
先问大家一个问题：CPU和GPU到底差在哪？
你想想看，CPU就像一位全能教授——什么都会，但一次只能处理一两件事。GPU呢？它更像是一群小学生——单个看啥都不行，但几千个一起上，算起加减乘除来，教授拍马也赶不上。
自动驾驶感知里，我们处理的是图像、点云、BEV特征。这些数据有个共同特点：数据量大，但计算模式单一。说白了，就是同样的操作要在几百万个像素点上重复执行。这不就是GPU的强项吗？
核心差异一句话总结：
- CPU：延迟优先，擅长复杂控制流，核心少但单核强
- GPU：吞吐优先，擅长数据并行，核心多但单核弱
我给大家拆解一下现代GPU的硬件结构。以NVIDIA的Ampere架构为例（就是A100用的那个），它大概长这样：
| 层级 | 名称 | 数量（A100为例） | 说明 |
|---|---|---|---|
| 1 | GPC（图形处理簇） | 8个 | 最高层级，每个GPC包含多个TPC |
| 2 | TPC（纹理处理簇） | 每个GPC含2个 | 包含SM和纹理单元 |
| 3 | SM（流式多处理器） | 每个TPC含2个 | 真正的计算单元，包含CUDA Core |
| 4 | CUDA Core | 每个SM含64个 | 执行浮点运算的最小单元 |
算一下：8 GPC × 2 TPC × 2 SM × 64 CUDA Core = 2048个CUDA Core。这还只是A100，H100直接干到18432个。
我在项目中遇到过一个问题：为什么我的卷积在GPU上跑得还不如CPU快？后来一查，发现是数据量太小了。GPU启动有开销，数据量没到一定规模，CPU反而更快。这个阈值大概是多少？我个人经验是，单次计算量低于10万次浮点运算，CPU可能更划算。
1.2 CUDA编程模型入门：从Hello World到并行计算
好，硬件看完了，我们聊聊怎么用它。
CUDA编程模型，说白了就是一套规则——告诉GPU：你要干什么、数据在哪、干完活怎么把结果拿回来。
先看一个最简单的例子。假设我们要把两个数组加起来：
// CPU版本
void add_cpu(float *a, float *b, float *c, int n) {
for (int i = 0; i < n; i++) {
c[i] = a[i] + b[i];
}
}
// GPU版本（CUDA）
__global__ void add_gpu(float *a, float *b, float *c, int n) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < n) {
c[idx] = a[idx] + b[idx];
}
}看到区别了吗？CPU版本用for循环一个个加，GPU版本每个线程只算一个元素。这就是数据并行的核心思想。
调用方式也不一样：
int main() {
float *d_a, *d_b, *d_c;
int n = 1000000;
// 1. 在GPU上分配显存
cudaMalloc(&d_a, n * sizeof(float));
cudaMalloc(&d_b, n * sizeof(float));
cudaMalloc(&d_c, n * sizeof(float));
// 2. 把数据从CPU拷贝到GPU
cudaMemcpy(d_a, h_a, n * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b, n * sizeof(float), cudaMemcpyHostToDevice);
// 3. 启动GPU内核
int threadsPerBlock = 256;
int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
add_gpu<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, n);
// 4. 把结果拷回CPU
cudaMemcpy(h_c, d_c, n * sizeof(float), cudaMemcpyDeviceToHost);
// 5. 释放显存
cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
}个人小技巧：我习惯把cudaMalloc和cudaMemcpy封装成宏或函数，加上错误检查。否则一旦显存泄漏，排查起来真的很痛苦。我曾经有一次在训练循环里忘了释放中间变量，结果跑了12小时后显存爆了，整个实验白做。
这里有几个关键概念，大家一定要记住：
- Grid：一个内核启动对应的所有线程集合
- Block：Grid里的线程块，同一个Block内的线程可以共享显存
- Thread：最小的执行单元
- Kernel：在GPU上执行的函数，用
__global__修饰
为什么Block大小通常选256或512？这不是随便选的。每个SM能同时运行的线程数有限，Block太大太小都不好。我一般从256开始调，看occupancy（占用率）指标，哪个高用哪个。
1.3 GPU与CPU的协同工作方式：别让任何一个闲着
好了，现在CPU和GPU都能干活了。但怎么让它们配合好？
说白了，就是流水线。CPU在准备下一批数据的时候，GPU在算当前这一批。两边都不闲着。
我给大家画个时间线：
| 时间 | CPU在干嘛 | GPU在干嘛 |
|---|---|---|
| T1 | 读取传感器数据，预处理 | 空闲 |
| T2 | 把数据拷贝到GPU（异步） | 空闲 |
| T3 | 准备下一帧数据 | 执行推理 |
| T4 | 从GPU取回结果 | 空闲（或开始下一帧） |
这里有个关键点：异步拷贝。用cudaMemcpyAsync配合CUDA Stream，可以让数据传输和计算重叠。我见过太多新手把拷贝写成同步的，结果GPU在那干等数据，利用率直接砍半。
避坑指南：我曾经在一个项目中，把CPU预处理和GPU推理写成了串行。结果帧率只有15 FPS，怎么优化都上不去。后来改成双缓冲+异步拷贝，直接干到30 FPS。记住：不要让GPU等CPU，也不要让CPU等GPU。
具体怎么做？我给大家一个伪代码框架：
// 双缓冲模式
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
float *buffer1, *buffer2;
cudaMalloc(&buffer1, size);
cudaMalloc(&buffer2, size);
while (running) {
// 当前帧用buffer1，下一帧预加载到buffer2
cudaMemcpyAsync(buffer2, cpu_data_next, size,
cudaMemcpyHostToDevice, stream2);
// 同时，GPU用stream1处理buffer1
kernel<<<grid, block, 0, stream1>>>(buffer1);
// 交换指针
swap(buffer1, buffer2);
swap(stream1, stream2);
}嗯，这里要注意：Stream的数量不是越多越好。我试过4个Stream，结果管理开销太大，反而变慢了。一般来说，2个Stream做双缓冲就够用。
小结
今天的内容就到这里。我们讲了三个东西：
- GPU为什么快——因为核心多，适合数据并行
- CUDA怎么用——Grid-Block-Thread三层结构，记住那5步流程
- CPU和GPU怎么配合——异步拷贝+双缓冲，别让任何一方闲着
下一节课，我们会深入CUDA内存模型，讲讲Global Memory、Shared Memory、Register这些玩意儿到底怎么用才能把性能榨干。到时候我会分享一个我调优YOLO检测器时踩过的坑——因为Shared Memory用错了，推理速度反而慢了3倍。
好，今天就到这。有问题随时问。