2、CUDA编程模型：CUDA架构简介、主机与设备、核函数概念、线程层次结构
好，咱们正式开始聊CUDA编程模型。说实话，很多初学者一上来就被各种概念砸晕了——Grid、Block、Thread，还有核函数、主机设备什么的。别急，我当年也是这么过来的。咱们一个一个拆开讲，你会发现其实没那么复杂。
2.1 CUDA架构简介
先说说CUDA到底是什么。CUDA是NVIDIA推出的并行计算平台和编程模型。说白了，它让你能用GPU做通用计算，不只是画图打游戏。
我个人习惯把CUDA架构理解成三层：硬件层、驱动层、编程层。硬件层就是GPU芯片本身，驱动层负责操作系统和硬件之间的通信，编程层就是我们写代码时用的API和工具。
CUDA架构的核心思想是异构计算。什么意思呢？就是CPU和GPU各司其职，CPU处理串行任务和复杂逻辑，GPU处理大规模并行计算。我在项目中遇到过不少新手，总想把所有逻辑都扔给GPU，结果性能反而更差。嗯，这里要注意：不是所有代码都适合上GPU。
关键点：CUDA架构支持细粒度并行和粗粒度并行。细粒度指线程级别的并行，粗粒度指任务级别的并行。实际开发中，两者往往结合使用。
2.2 主机与设备
在CUDA编程中，主机（Host）指的是CPU及其内存系统，设备（Device）指的是GPU及其显存。这两个角色分工明确：
- 主机：负责控制流程、数据准备、调用核函数
- 设备：负责执行大规模并行计算任务
你想想看，主机和设备之间通过PCIe总线通信。数据传输是个大开销，我刚开始做项目时没注意这点，结果数据传输时间比计算时间还长，得不偿失。
避坑指南：我曾经在做一个图像处理项目时，每帧图像都从主机拷贝到设备，处理完再拷回来。后来发现，尽量减少主机与设备之间的数据传输才是性能优化的关键。能一次传完就别分多次传。
典型的CUDA程序流程是这样的：
- 分配主机内存并初始化数据
- 分配设备显存
- 将数据从主机拷贝到设备
- 调用核函数在设备上执行计算
- 将结果从设备拷贝回主机
- 释放内存和显存
代码示例：
// 主机端代码
float *h_data, *d_data;
int size = 1024 * sizeof(float);
// 1. 分配主机内存
h_data = (float*)malloc(size);
// 2. 分配设备显存
cudaMalloc(&d_data, size);
// 3. 数据从主机拷贝到设备
cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);
// 4. 调用核函数（稍后讲解）
myKernel<<<grid, block>>>(d_data);
// 5. 结果从设备拷贝回主机
cudaMemcpy(h_data, d_data, size, cudaMemcpyDeviceToHost);
// 6. 释放内存
free(h_data);
cudaFree(d_data);2.3 核函数概念
核函数（Kernel）是CUDA编程的核心。它是在设备上执行的函数，用__global__修饰符声明。调用核函数时，用三括号语法<<<grid, block>>>指定执行配置。
核函数有几个特点：
- 必须返回
void类型 - 不能使用可变参数
- 不能是类的成员函数（静态成员函数可以）
- 在设备上异步执行，主机继续执行后续代码
举个例子：
// 核函数定义
__global__ void vectorAdd(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 n = 1024;
int threadsPerBlock = 256;
int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, n);我建议你在写核函数时，先确保逻辑正确，再考虑优化。曾经有个同事一上来就搞各种骚操作，结果调试了三天才发现是索引算错了。
小技巧：核函数内部可以通过threadIdx、blockIdx、blockDim等内置变量获取当前线程的索引信息。这些变量是CUDA自动提供的，直接用就行。
2.4 线程层次结构
这是CUDA编程模型中最重要也最容易混淆的部分。线程层次结构分为三级：Grid（网格）、Block（线程块）、Thread（线程）。
| 层次 | 说明 | 最大维度 | 典型大小 |
|---|---|---|---|
| Grid | 由多个Block组成，对应一个核函数调用 | 3维（x, y, z） | 根据数据量决定 |
| Block | 由多个Thread组成，共享共享内存 | 3维（x, y, z） | 128-512个线程 |
| Thread | 最小的执行单元 | 1维 | 1个线程 |
为什么会设计成三层？说白了是为了可扩展性。不同GPU的计算能力不同，有的GPU有更多SM（流多处理器），有的更少。通过Grid-Block-Thread的层次结构，同一个程序可以在不同GPU上运行，只是并行度不同。
我记得刚开始学CUDA时，总搞不清gridDim和blockDim的区别。后来想通了：Grid是宏观视角，Block是微观视角。Grid决定了总共有多少Block，Block决定了每个Block里有多少线程。
线程索引的计算公式：
// 一维情况
int tid = threadIdx.x + blockIdx.x * blockDim.x;
// 二维情况
int tid_x = threadIdx.x + blockIdx.x * blockDim.x;
int tid_y = threadIdx.y + blockIdx.y * blockDim.y;
int tid = tid_y * gridDim.x * blockDim.x + tid_x;
// 三维情况类似，依次类推重要提示：Block内的线程可以通过共享内存（Shared Memory）进行通信和同步。但不同Block之间的线程无法直接通信，只能通过全局内存间接通信。这个设计是为了保证可扩展性——Block可以独立地在任何SM上执行。
实际项目中，我一般这样配置线程层次：
- Block大小选256或512，这是大多数GPU的warp大小（32）的整数倍
- Grid大小根据数据总量和Block大小计算，确保覆盖所有数据
- 如果数据维度是2D或3D，用对应的2D/3D索引会更直观
举个例子，处理一张1024x768的图像：
dim3 blockSize(16, 16); // 每个Block 16x16 = 256个线程
dim3 gridSize(
(1024 + blockSize.x - 1) / blockSize.x,
(768 + blockSize.y - 1) / blockSize.y
);
imageProcess<<<gridSize, blockSize>>>(d_image, width, height);嗯，这里要注意：Block内的线程数不能超过设备限制。不同计算能力的GPU，每个Block最大线程数不同，一般是1024。我曾经在旧设备上踩过这个坑，核函数怎么都不执行，查了半天才发现是Block大小超限了。
常见错误：忘记处理边界情况。当数据总量不是Block大小的整数倍时，最后一个Block的部分线程会访问越界。一定要在核函数里加边界检查，就像前面vectorAdd例子中的if (i < n)那样。
总结一下，CUDA编程模型的核心就是：主机控制、设备计算、核函数执行、线程层次管理。掌握了这些，你就迈出了GPU并行计算的第一步。下一章我们会深入讨论内存管理，那才是真正考验功力的地方。