第一章 CUDA编程基础回顾：编程模型与层次结构
各位同学，欢迎来到《GPU集群核函数部署与调优实战》的第一章。今天咱们不聊虚的，直接回顾CUDA编程最核心的几个概念。我做了这么多年GPU优化，发现很多性能问题追根溯源，都是对基础理解不透彻导致的。所以这一章，咱们把地基打牢。
1.1 CUDA编程模型：从CPU到GPU的思维转变
CUDA的编程模型，说白了就是「CPU负责指挥，GPU负责干活」。CPU我们叫它Host，GPU叫Device。你想想看，CPU擅长处理复杂的逻辑判断，GPU擅长做大量重复的数学计算。所以典型的CUDA程序流程是：
- 在CPU上准备好数据
- 把数据拷贝到GPU显存
- 在GPU上启动成千上万个线程并行计算
- 把结果拷回CPU
嗯，这里有个关键点：数据传输是瓶颈。我早期做项目时，曾经花了一周调优核函数，结果发现90%的时间都花在PCIe传输上。所以后来我养成了一个习惯——尽量减少CPU和GPU之间的数据搬运，能一次传完就别分十次。
核心原则：CUDA编程的本质是「异构计算」——CPU做控制，GPU做计算。两者各司其职，别让CPU去干GPU的活，也别让GPU去处理复杂的if-else分支。
1.2 线程层次结构：Grid、Block、Thread
CUDA的线程组织方式，是一个三层结构：Grid → Block → Thread。我习惯把它想象成一个军队：
- Grid：整个战役（一个核函数启动）
- Block：一个作战小队（一组线程）
- Thread：一个士兵（单个执行单元）
每个线程都有自己唯一的ID，通过 threadIdx、blockIdx、blockDim 和 gridDim 来定位。举个例子：
// 一个典型的CUDA核函数
__global__ void vecAdd(float *A, float *B, float *C, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {
C[idx] = A[idx] + B[idx];
}
}
// 启动配置
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);这里有个坑，我踩过好几次：Block内的线程数不要超过1024（对于现代GPU）。为什么？因为一个Block内的线程共享同一个SM（流多处理器）的资源，线程太多会导致寄存器不够用，性能反而下降。
我的经验：Block大小通常选128或256。太小了SM利用率低，太大了资源竞争激烈。256是个比较稳妥的起点，然后根据具体算法微调。
1.3 内存层次结构：谁快谁慢，心里要有数
GPU的内存层次，说白了就是「距离计算单元越近，速度越快，容量越小」。我画了一张图，帮你直观理解：
这张图我建议你保存下来。每次写核函数前看一眼，问问自己：数据放对地方了吗？
1.3.1 寄存器（Register）
寄存器是GPU里最快的内存，每个线程私有。我见过不少新手，在核函数里声明了一大堆局部变量，结果寄存器不够用，数据被「溢出」到本地内存，性能直接腰斩。怎么避免？控制每个线程的寄存器使用量。可以用 --maxrregcount=32 编译选项限制，或者用 __launch_bounds__ 提示编译器。
1.3.2 共享内存（Shared Memory）
共享内存是同一个Block内所有线程都能访问的「小快灵」存储。我做过一个矩阵乘法的优化，用共享内存做数据分块，性能提升了5倍。但要注意：共享内存只有几十KB（具体看GPU型号），用的时候要精打细算。
// 共享内存使用示例：矩阵分块
__global__ void matMulShared(float *A, float *B, float *C, int N) {
__shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
__shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
float sum = 0.0f;
for (int tile = 0; tile < N / BLOCK_SIZE; tile++) {
As[threadIdx.y][threadIdx.x] = A[row * N + tile * BLOCK_SIZE + threadIdx.x];
Bs[threadIdx.y][threadIdx.x] = B[(tile * BLOCK_SIZE + threadIdx.y) * N + col];
__syncthreads();
for (int k = 0; k < BLOCK_SIZE; k++) {
sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
}
__syncthreads();
}
C[row * N + col] = sum;
}注意：使用共享内存时，别忘了加 __syncthreads() 同步。我曾经因为漏掉这个同步，调试了整整两天，结果发现是线程执行顺序不一致导致的数据竞争。
1.3.3 全局内存（Global Memory）
全局内存是GPU里最大的存储空间，但也是最慢的。访问全局内存时，要尽量做到合并访问（coalesced access）。什么意思？就是让连续的线程访问连续的内存地址。比如：
// 好的访问模式：合并访问
int idx = threadIdx.x + blockIdx.x * blockDim.x;
float val = data[idx]; // 线程0访问data[0]，线程1访问data[1]...
// 坏的访问模式：非合并访问
int idx = threadIdx.x + blockIdx.x * blockDim.x;
float val = data[idx * 4]; // 线程0访问data[0]，线程1访问data[4]...非合并访问会导致内存带宽利用率极低。我见过一个项目，改了一行代码（把步长从4改成1），性能提升了3倍。
1.3.4 本地内存（Local Memory）
本地内存是个「陷阱」。名字叫「本地」，实际上它存储在全局内存里，速度很慢。什么时候会用？当寄存器不够用的时候，编译器会自动把一些局部变量放到本地内存。怎么发现？用 nvcc --ptxas-options=-v 编译，看输出里的 lmem 使用量。如果超过0，说明有寄存器溢出，需要优化。
1.4 总结：一张表看清所有内存
| 内存类型 | 作用域 | 生命周期 | 速度 | 容量 |
|---|---|---|---|---|
| 寄存器 | 单个线程 | 核函数执行期间 | 最快 (~1 cycle) | 每个线程最多255个（取决于GPU） |
| 共享内存 | 单个Block | 核函数执行期间 | 快 (~5-30 cycles) | 每个SM 48KB-164KB |
| 全局内存 | 所有线程+Host | 程序运行期间 | 慢 (~400-800 cycles) | 最大 (数GB到数十GB) |
| 本地内存 | 单个线程 | 核函数执行期间 | 慢 (同全局内存) | 取决于寄存器溢出量 |
我的建议：写核函数时，优先用寄存器，其次共享内存，最后才用全局内存。如果非要用全局内存，一定要保证合并访问。这个习惯能帮你避免90%的性能问题。
好了，这一章的内容就到这里。CUDA的编程模型和内存层次结构是后续所有优化的基础。你把这些概念吃透了，后面讲核函数调优、多卡通信、性能分析工具的时候，才能跟得上节奏。