第1章：CUDA编程模型
1.1 CUDA架构简介
CUDA，全称是Compute Unified Device Architecture。说白了，就是NVIDIA给GPU开的一扇门——让咱们能用C语言去指挥GPU干活。
我刚开始接触CUDA时，有个很深的感触：CPU和GPU就像两个性格迥异的搭档。CPU是那种什么都会一点的通才，但一次只能处理几件事；GPU则是专才，擅长同时处理成千上万件简单任务。
CUDA架构的核心思想，就是让CPU和GPU各司其职。CPU负责控制流、逻辑判断这些复杂活；GPU则专注于数据并行计算——也就是把同样的操作，应用到大量数据上。
核心要点：CUDA不是让GPU取代CPU，而是让两者协同工作。CPU是大脑，GPU是肌肉。
我记得第一次用CUDA做矩阵乘法时，CPU版本跑了2秒，GPU版本只用了0.05秒。当时我就知道，这条路走对了。
1.2 主机与设备
在CUDA的世界里，有两个重要的概念：主机（Host）和设备（Device）。
- 主机：就是CPU及其内存系统。负责程序控制、逻辑判断、I/O操作。
- 设备：就是GPU及其显存。负责大规模并行计算。
主机和设备有各自独立的内存空间。你不能直接在CPU代码里访问GPU显存，反之亦然。数据需要显式地在两者之间传输。
我见过不少新手踩过这个坑——在GPU核函数里直接访问CPU内存，结果程序崩溃得一塌糊涂。嗯，这里要注意：主机和设备的内存是隔离的。
典型的数据处理流程是这样的：
- 在主机端分配内存，准备好输入数据
- 在设备端分配显存
- 将数据从主机拷贝到设备
- 调用核函数在GPU上执行计算
- 将结果从设备拷贝回主机
- 释放主机和设备的内存
个人经验：我习惯在代码开头用 cudaSetDevice(0) 指定使用的GPU设备。如果机器上有多个GPU，这一步能避免很多莫名其妙的问题。
1.3 核函数（Kernel）概念
核函数，就是运行在GPU上的函数。用 __global__ 修饰符声明，调用时用 <<<...>>> 语法指定执行配置。
来看一个最简单的例子：
// 核函数定义
__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() {
// ... 数据准备 ...
vecAdd<<<numBlocks, blockSize>>>(d_A, d_B, d_C, N);
// ... 结果回收 ...
}你想想看，这个核函数会被GPU上成千上万个线程同时执行。每个线程处理一个元素，这就是所谓的「单指令多线程（SIMT）」模型。
核函数有几个特点：
- 返回类型必须是
void - 不能使用可变参数
- 不能是类的成员函数（静态成员函数可以）
- 在设备上执行，从主机端调用
避坑指南：我曾经在核函数里用了 printf 来调试，结果发现输出顺序完全随机。因为GPU线程的执行顺序是不确定的。所以，不要在核函数里依赖任何顺序相关的操作。
1.4 线程层次结构
CUDA的线程组织方式，是我觉得最巧妙的设计之一。它把线程分成三个层次：
| 层次 | 说明 | 最大数量 |
|---|---|---|
| 线程（Thread） | 最小的执行单元 | 每个块最多1024个线程 |
| 线程块（Block） | 一组线程，可以协作 | 每个网格最多2^31-1个块 |
| 网格（Grid） | 一组线程块，构成一次核函数调用 | 取决于硬件 |
每个线程都有唯一的ID，通过内置变量访问：
threadIdx.x：线程在块内的索引blockIdx.x：块在网格内的索引blockDim.x：每个块的线程数gridDim.x：网格内的块数
全局线程ID的计算公式：int tid = threadIdx.x + blockIdx.x * blockDim.x;
为什么要有这种层次结构？说白了，是为了适应GPU的硬件架构。同一个块内的线程可以共享数据（通过共享内存），也可以同步（通过 __syncthreads()）。不同块之间的线程则完全独立。
我的建议：块大小一般设为32的倍数（比如128、256），因为GPU的warp大小就是32。这样能充分利用硬件资源。我通常用256，这个值在大多数场景下表现都不错。
线程层次结构可以用一张图来理解：
这张图展示了CUDA的线程层次：一个网格包含多个块，每个块包含多个线程。实际编程中，我们通常用一维、二维或三维的索引来组织线程，方便处理不同维度的数据。
举个例子，处理一张1024×1024的图像，你可以用 dim3 block(32, 32) 和 dim3 grid(32, 32) 来组织线程。每个线程处理一个像素，清晰又直观。
关键理解：线程层次结构不是凭空设计的，它直接映射到GPU的硬件架构——块对应SM（流多处理器），线程对应CUDA核心。理解这一点，你就能写出更高效的代码。
好了，这一章的内容就到这里。CUDA编程模型是后续所有内容的基础，建议你亲手写几个小例子跑一跑。代码跑通了，理解自然就深了。