1. CUDA编程模型概述：CPU与GPU架构差异、异构计算概念、CUDA编程基本流程
各位同学，咱们今天开篇。说实话，CUDA这东西，我刚开始接触的时候也是一头雾水。CPU玩得好好的，干嘛非要折腾GPU？后来真香了。这一章，咱们先把地基打牢。
1.1 CPU与GPU：天生不同
CPU和GPU，说白了就是两种完全不同的“打工人”。
CPU，像是一位全能教授。什么都能干，逻辑判断、分支预测、大缓存、乱序执行……样样精通。但问题是，教授就一个，一次只能处理一个复杂任务。它的核心设计目标，是降低延迟——让单个任务跑得飞快。
GPU，则像是一支千人工厂流水线。每个工人只会简单的加减乘除，但人多啊！成千上万个工人一起干，吞吐量惊人。它的核心设计目标，是提高吞吐量——让海量简单任务并行完成。
我举个例子你就明白了。假设你要算10000个数的加法：
- CPU：一个教授，算完一个再算下一个，每个都算得飞快。
- GPU：一万个工人，每人算一个，一秒搞定。
这就是差异的本质。CPU擅长串行、复杂、依赖性强的任务。GPU擅长并行、简单、相互独立的任务。
核心差异总结：
| 特性 | CPU | GPU |
|---|---|---|
| 核心数 | 少（4-16核） | 极多（数千核） |
| 缓存 | 大（MB级别） | 小（KB级别） |
| 控制单元 | 复杂（分支预测、乱序执行） | 简单（顺序执行） |
| 设计目标 | 降低延迟 | 提高吞吐量 |
| 适合任务 | 串行、复杂逻辑 | 并行、数据密集型 |
1.2 异构计算：让对的人做对的事
既然CPU和GPU各有千秋，那为什么不把它们结合起来？这就是异构计算的核心思想。
异构计算，说白了就是“让对的人做对的事”。
- CPU：负责控制逻辑、串行部分、I/O操作、任务调度。
- GPU：负责大规模并行计算、数据密集型运算。
我做过一个图像处理项目，一开始全用CPU跑，一张4K图片要处理3秒。后来改成CPU负责读取图片和调度，GPU负责像素级并行处理，直接降到30毫秒。你想想看，这差距有多大？
我的经验：异构计算不是简单的“把任务扔给GPU”。你需要仔细分析：哪些部分必须串行？哪些可以并行？串行部分留在CPU，并行部分交给GPU。这个“任务划分”的功夫，决定了最终性能。
1.3 CUDA编程基本流程
好了，理论讲完，咱们来点实际的。CUDA编程到底怎么玩？我总结了一个“三步走”流程：
- 数据准备（CPU端）：在主机内存中分配空间，初始化数据。
- 数据搬运（CPU→GPU）：把数据从主机内存拷贝到设备显存。
- GPU执行（核函数）：在GPU上启动核函数，并行计算。
- 结果回传（GPU→CPU）：把计算结果从显存拷贝回主机内存。
- 资源清理：释放分配的内存和显存。
嗯，这里要注意：数据搬运是开销最大的环节。我曾经有个项目，计算本身只用了1毫秒，但数据拷贝花了10毫秒。你说亏不亏？所以，尽量减少CPU和GPU之间的数据传输次数。
来看一个最简单的CUDA程序，向量加法：
// 核函数：在GPU上执行
__global__ void vecAdd(float *A, float *B, float *C, int N) {
int i = threadIdx.x + blockIdx.x * blockDim.x;
if (i < N) {
C[i] = A[i] + B[i];
}
}
int main() {
int N = 1024;
size_t size = N * sizeof(float);
// 1. 主机端分配内存
float *h_A = (float*)malloc(size);
float *h_B = (float*)malloc(size);
float *h_C = (float*)malloc(size);
// 初始化数据
for (int i = 0; i < N; i++) {
h_A[i] = i;
h_B[i] = i * 2;
}
// 2. 设备端分配显存
float *d_A, *d_B, *d_C;
cudaMalloc(&d_A, size);
cudaMalloc(&d_B, size);
cudaMalloc(&d_C, size);
// 3. 数据从主机拷贝到设备
cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
// 4. 启动核函数
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
// 5. 结果从设备拷贝回主机
cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
// 6. 释放资源
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
free(h_A);
free(h_B);
free(h_C);
return 0;
}避坑指南：我曾经犯过一个低级错误——忘记检查cudaMalloc的返回值。结果在显存不足时，程序静默失败，数据全乱了。记住：每次CUDA API调用后，都要检查错误码。可以用cudaGetLastError()或者定义宏来统一处理。
1.4 线程、块、网格：CUDA的并行层次
CUDA的并行模型，其实就三个层次：
- 线程（Thread）：最小的执行单元。每个线程执行核函数中的代码。
- 线程块（Block）：一组线程的集合。块内的线程可以共享内存、同步协作。
- 网格（Grid）：一组线程块的集合。一个核函数对应一个网格。
你想想看，这就像一家公司：
- 线程 = 普通员工，各干各的活。
- 线程块 = 项目组，组内可以开会讨论（共享内存）。
- 网格 = 整个公司，多个项目组并行推进。
在代码中，我们通过<<<grid, block>>>来配置：
// gridDim: 网格的维度（有多少个block）
// blockDim: 块的维度（每个block有多少个thread）
kernel<<<gridDim, blockDim>>>(...);每个线程可以通过内置变量知道自己是谁：
threadIdx.x：在线程块内的索引blockIdx.x：在网格内的块索引blockDim.x：每个块的线程数gridDim.x：网格的块数
全局线程ID的计算公式：int tid = blockIdx.x * blockDim.x + threadIdx.x;
我的习惯：刚开始学CUDA时，我总搞不清这些索引。后来我画了一张图——把网格想象成一个二维表格，block是单元格，thread是单元格里的小格子。每次写代码前，先画图，再写索引公式。这招很管用，你可以试试。
1.5 本章小结
这一章，咱们把CUDA的“世界观”建立起来了：
- CPU和GPU，一个精于串行，一个强于并行。
- 异构计算，就是让两者各司其职。
- CUDA编程，核心是“数据搬运 + 核函数执行”。
- 线程、块、网格，构成了CUDA的并行层次。
下一章，咱们会深入线程块和网格的配置细节。你会学到：到底该设多少个线程？多少个块？为什么设错了性能会暴跌？嗯，这些都是实战中血泪换来的经验。
先消化今天的内容。有问题随时问。