1. CUDA与GPU计算基础

GPU架构演进：从图形卡到通用计算引擎

说起GPU架构，我最早接触的是Fermi架构。那时候还在做分子动力学模拟，说实话，当时的GPU远没有现在这么智能。你想想看，从最初的固定渲染管线，到可编程着色器，再到统一的CUDA架构，这条路走了将近二十年。

我个人习惯把GPU架构演进分成三个阶段：

• 固定功能阶段（2000年前）：GPU只能干图形渲染的活，程序员想用它做计算？门都没有。
• 可编程着色器阶段（2001-2006）：出现了顶点着色器和像素着色器，但编程模型极其别扭。我记得当时为了用GPU做矩阵乘法，得把数据伪装成纹理，那代码写得...嗯，不堪回首。
• 统一架构阶段（2007至今）：NVIDIA发布了CUDA，GPU终于有了统一的编程模型。从Fermi到Turing，再到Ampere和Hopper，每一代都在提升计算能力和内存带宽。

为什么GPU能跑这么快？说白了，就是堆核心。一个高端GPU有上万个CUDA核心，而CPU最多也就几十个核心。但每个CUDA核心都很"瘦"，没有复杂的分支预测、乱序执行这些花哨功能。GPU靠的是人多力量大。

CUDA编程模型：你只需要理解这三个层次

CUDA编程模型，我建议你从三个层次去理解：线程层次、内存层次、执行层次。

线程层次

CUDA把线程组织成三层结构：

• Thread（线程）：最基础的执行单元
• Block（线程块）：一组线程，可以共享内存和同步
• Grid（网格）：一组线程块，构成整个计算任务

我在项目中遇到过一个问题：刚开始写CUDA程序时，总想把所有线程都塞到一个block里。结果呢？一个block最多1024个线程，数据量一大就崩了。后来才明白，Grid才是真正用来描述大规模并行的。

内存层次

GPU的内存层次，说白了就是"越近越快，越远越慢"：

内存类型	位置	访问速度	容量	作用域
寄存器	每个线程私有	~1周期	几十KB	单个线程
共享内存	每个Block私有	~5周期	几十KB	Block内线程
全局内存	所有线程共享	~400周期	几十GB	所有线程
常量内存	只读	~10周期（缓存命中）	64KB	所有线程

执行层次

GPU执行的最小单位是Warp（线程束），通常是32个线程。一个Warp里的线程执行相同的指令，这就是所谓的SIMT（单指令多线程）模型。

你想想看，如果Warp里的32个线程走了不同的分支路径，会发生什么？嗯，所有路径都会串行执行。这就是所谓的"Warp发散"，是性能杀手之一。

GPU与CPU协同工作原理

CPU和GPU的关系，我经常比喻成"将军和士兵"。CPU是将军，负责决策和调度；GPU是士兵，负责执行大量重复的计算任务。

典型的协同工作流程是这样的：

1. 数据准备：CPU把数据从硬盘加载到内存
2. 数据传输：CPU通过PCIe总线把数据拷贝到GPU显存
3. 内核启动：CPU调用CUDA内核函数，GPU开始执行
4. 结果回传：GPU计算完成后，把结果拷贝回CPU内存
5. 后续处理：CPU对结果进行进一步处理或输出

这里有一个简单的CUDA程序示例，展示了CPU和GPU的协同：

// CPU端：分配内存
float *h_a, *h_b, *h_c;
h_a = (float*)malloc(N * sizeof(float));
h_b = (float*)malloc(N * sizeof(float));
h_c = (float*)malloc(N * sizeof(float));

// CPU端：初始化数据
for (int i = 0; i < N; i++) {
    h_a[i] = i * 1.0f;
    h_b[i] = i * 2.0f;
}

// CPU端：分配GPU显存
float *d_a, *d_b, *d_c;
cudaMalloc(&d_a, N * sizeof(float));
cudaMalloc(&d_b, N * sizeof(float));
cudaMalloc(&d_c, N * sizeof(float));

// CPU端：把数据从内存拷贝到显存
cudaMemcpy(d_a, h_a, N * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b, N * sizeof(float), cudaMemcpyHostToDevice);

// CPU端：启动GPU内核
vec_add<<<grid_size, block_size>>>(d_c, d_a, d_b, N);

// CPU端：把结果从显存拷贝回内存
cudaMemcpy(h_c, d_c, N * sizeof(float), cudaMemcpyDeviceToHost);

// CPU端：验证结果
for (int i = 0; i < N; i++) {
    printf("h_c[%d] = %f\n", i, h_c[i]);
}

这个流程看起来简单，但实际项目中要考虑的东西很多。比如：如何选择block_size？如何避免bank conflict？如何利用stream实现并发？这些我们后面章节会详细讲。

本章知识体系

下面这张图展示了本章的核心知识结构：

嗯，以上就是本章的核心内容。GPU架构演进让我们知道"从哪里来"，CUDA编程模型告诉我们"怎么用"，CPU与GPU协同则解决了"怎么配合"的问题。这三块内容，是后续所有章节的基础。