CUDA编程模型概述
各位同学好，我是你们这节课程的讲师。今天咱们聊聊CUDA编程模型——说白了，就是NVIDIA为了让GPU干活而设计的一套规则。我刚开始接触CUDA时，觉得这东西挺玄乎的，后来做多了才发现，理解清楚编程模型，后面写算子就跟搭积木一样顺手。
CUDA编程模型的核心思想其实很简单：把CPU当作指挥官，GPU当作打工仔。CPU负责逻辑控制、串行任务，GPU负责大规模并行计算。这种设计我特别喜欢，因为它把复杂问题拆成了两块——你不需要操心GPU内部怎么调度，只需要告诉它“干什么”和“怎么干”。
核心要点：CUDA不是让GPU独立工作，而是CPU+GPU协同作战。CPU叫Host，GPU叫Device。
Host与Device
先说说Host和Device。Host就是CPU及其内存系统，Device就是GPU及其显存。我在项目中遇到过不少新手，上来就把所有数据都扔到GPU上，结果发现有些逻辑判断根本不适合并行——这就是没搞清楚Host和Device的分工。
典型的CUDA程序流程是这样的：
- 在Host上准备好数据
- 把数据从Host拷贝到Device
- 在Device上执行Kernel函数
- 把结果从Device拷贝回Host
嗯，这里要注意：数据传输是开销最大的环节。我见过有人为了省事，每次循环都来回拷贝数据，结果性能还不如纯CPU版本。你想想看，PCIe带宽再快，也赶不上显存带宽啊。
避坑指南：我曾经在一个图像处理项目里，每帧都把整张图从Host拷到Device，再从Device拷回Host。后来改成只拷贝变化区域，性能提升了3倍。记住：尽量减少Host-Device之间的数据传输。
Kernel函数概念
Kernel函数，就是跑在GPU上的函数。它的定义方式很特别——前面加个__global__修饰符。我习惯把Kernel函数想象成“给每个线程执行的同一段代码”。
看个最简单的例子：
// 定义一个Kernel函数
__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];
}
}
// 调用Kernel函数
int main() {
// ... 数据准备和拷贝 ...
vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);
// ... 结果拷贝回Host ...
}注意那个<<<gridSize, blockSize>>>语法，这是CUDA特有的调用方式。它告诉GPU：我要启动多少个线程块（Block），每个块里有多少个线程（Thread）。
个人经验：我刚开始写Kernel时，总喜欢把逻辑写得特别复杂。后来发现，Kernel函数越简单越好——最好只做一件事，而且数据访问模式要规整。复杂的逻辑交给Host去处理。
线程层次结构（Grid/Block/Thread）
这是CUDA编程模型里最核心的概念。我画个图帮你理解：
这张图展示了CUDA的线程层次结构：
- Grid（网格）：一个Kernel启动时创建的所有线程的集合
- Block（线程块）：Grid被分成若干Block，每个Block内的线程可以协作（共享内存、同步）
- Thread（线程）：最小的执行单元，每个线程执行相同的Kernel代码
每个线程都有唯一的ID，通过threadIdx、blockIdx、blockDim、gridDim这些内置变量来定位。我习惯用一维索引，但实际项目中二维、三维索引也很常见——比如处理图像时用二维索引就特别自然。
关键理解：Block内的线程可以共享数据，但Block之间是独立的。这意味着你不能假设两个不同Block的线程执行顺序——它们可能同时跑，也可能先后跑。
CUDA程序编译与执行流程
最后聊聊编译和执行。CUDA程序的编译不是一步到位的，它分两个阶段：
| 阶段 | 做什么 | 产出 |
|---|---|---|
| 离线编译 | nvcc编译器把.cu文件中的Host代码和Device代码分开处理 | Host代码编译成普通CPU可执行文件，Device代码编译成PTX或cubin |
| 运行时编译 | GPU驱动把PTX代码编译成具体GPU架构的机器码 | 可执行的GPU指令 |
为什么搞这么复杂？说白了就是为了兼容性。你写的CUDA代码可以在不同代的GPU上跑，驱动会根据具体硬件做优化。我记得有一次在旧卡上调试好的代码，换到新卡上性能反而下降了——后来发现是PTX编译时没指定架构参数。
实用建议：编译时加上-arch=sm_xx参数，指定目标GPU架构。比如-arch=sm_80对应A100。如果不确定，用-arch=all生成所有架构的代码，但编译时间会长一些。
执行流程就更直观了：
- Host代码运行，准备数据
- 调用
cudaMemcpy把数据拷到Device - 启动Kernel（
<<<grid, block>>>） - GPU执行Kernel，所有线程并行跑
- 调用
cudaMemcpy把结果拷回Host - 释放资源
这里有个坑：Kernel启动是异步的。也就是说，你调用vecAdd<<<...>>>(...)之后，CPU不会等GPU算完就继续往下执行。我曾经因为这个bug调试了一整天——在Kernel还没执行完就去读结果，读到的全是垃圾数据。解决方案是用cudaDeviceSynchronize()同步，或者用cudaMemcpy隐式同步。
避坑指南：我曾经在循环里反复启动Kernel，每次启动后都忘了同步。结果GPU队列里积压了几百个Kernel，显存爆了。记住：异步是好事，但别忘了同步点。
好了，这一章的内容就到这里。CUDA编程模型是后面所有算子的基础——理解清楚Host/Device分工、Kernel函数写法、线程层次结构、编译执行流程，后面写各种算子才能游刃有余。
想要更多资料，联系可可