4. CUDA编程基础：从零开始理解GPU编程
说实话，我第一次接触CUDA的时候，心里是有点发怵的。那时候我刚从CPU编程转到GPU，满脑子都是"这玩意儿到底怎么并行？"后来踩了不少坑，才慢慢摸清门道。今天我就把这些年积累的经验，跟你好好聊聊。
4.1 CUDA编程模型：主机与设备
CUDA的编程模型，说白了就是"CPU+GPU"协同工作。CPU叫主机（Host），GPU叫设备（Device）。我习惯把CPU想象成项目经理，GPU就是一群埋头苦干的程序员。
典型的CUDA程序流程是这样的：
- CPU把数据准备好，拷贝到GPU显存
- CPU告诉GPU："来，帮我算这个"
- GPU吭哧吭哧算完
- CPU再把结果拷回来
嗯，这里要注意：数据拷贝是开销最大的操作。我在项目中遇到过，有人把每次迭代的数据都来回拷贝，结果性能还不如CPU。你想想看，这多冤啊。
4.2 线程层次结构：Grid、Block、Thread
这是CUDA最核心的概念，也是新手最容易懵的地方。我刚开始学的时候，Grid和Block的关系绕了好久才搞明白。
简单来说：
- Thread（线程）：最基础的计算单元，每个线程执行相同的核函数
- Block（线程块）：一组线程的集合，可以互相协作
- Grid（网格）：一组Block的集合，构成整个计算任务
你可以把GPU想象成一个工厂：
- Grid = 整个工厂
- Block = 车间
- Thread = 工人
每个工人（Thread）干一样的活，但处理不同的数据。这就是所谓的SIMT（单指令多线程）模型。
来看个具体的例子：
// 定义Grid和Block的维度
dim3 gridDim(2, 2); // 2x2的网格，共4个Block
dim3 blockDim(16, 16); // 16x16的Block，共256个线程
// 启动核函数
myKernel<<<gridDim, blockDim>>>(d_data, N);这段代码启动了4个Block，每个Block有256个线程，总共1024个线程并行执行。每个线程可以通过内置变量知道自己是谁：
threadIdx.x：线程在Block内的索引blockIdx.x：Block在Grid内的索引blockDim.x：Block的维度大小gridDim.x：Grid的维度大小
我经常用这个公式计算全局线程ID：
int tid = blockIdx.x * blockDim.x + threadIdx.x;这个公式你最好记牢了，几乎每个CUDA程序都会用到。
4.3 内存层次结构：Global、Shared、Local
GPU的内存体系比CPU复杂得多。我刚开始做项目时，就因为搞混了内存类型，性能差了10倍不止。
先看这张图，帮你理清思路：
三种内存各有特点，我总结一下：
| 内存类型 | 容量 | 速度 | 作用域 | 典型用途 |
|---|---|---|---|---|
| Global Memory | GB级 | 慢（~400 cycles） | 所有线程 | 存储数据、模型参数 |
| Shared Memory | 几十KB | 快（~20 cycles） | 同一Block | 数据复用、线程协作 |
| Local Memory | 寄存器溢出时 | 同Global | 单个线程 | 自动使用，尽量避免 |
4.4 CUDA核函数编写入门
核函数（Kernel Function）就是在GPU上执行的函数。写法上有个标志性的关键字——__global__。
来看一个最简单的例子：向量加法。
// 核函数定义
__global__ void vecAdd(float* A, float* B, float* C, int N) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < N) {
C[i] = A[i] + B[i];
}
}
// 主函数调用
int main() {
int N = 1 << 20; // 约100万个元素
float *d_A, *d_B, *d_C;
// 在GPU上分配内存
cudaMalloc(&d_A, N * sizeof(float));
cudaMalloc(&d_B, N * sizeof(float));
cudaMalloc(&d_C, N * sizeof(float));
// 拷贝数据到GPU
cudaMemcpy(d_A, h_A, N * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, N * sizeof(float), cudaMemcpyHostToDevice);
// 启动核函数
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
// 拷贝结果回CPU
cudaMemcpy(h_C, d_C, N * sizeof(float), cudaMemcpyDeviceToHost);
// 释放GPU内存
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
return 0;
}这段代码有几个关键点：
__global__：告诉编译器这是个核函数，在GPU上执行<<<grid, block>>>：配置启动参数，指定线程组织方式cudaMalloc/cudaMemcpy：显存管理和数据传输if (i < N)：边界检查，防止越界
核函数有几个限制要记住：
- 不能调用C++标准库函数（除了少数数学函数）
- 不能使用静态变量
- 不能使用变长参数
- 返回值必须是void
说白了，核函数就是"轻装上阵"。它只负责计算，不负责管理资源。
最后分享一个调试技巧：我经常用cudaGetLastError()来检查核函数是否启动成功。如果返回cudaSuccess，说明启动没问题。但要注意，这个函数只能检查启动错误，不能检查运行时的逻辑错误。
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("核函数启动失败: %s\n", cudaGetErrorString(err));
}嗯，CUDA编程的基础就这些。说白了就是三件事：理解线程怎么组织、搞清内存怎么用、写好核函数。把这些搞明白了，后面的优化技巧才能派上用场。