2、CUDA编程模型核心：线程层次结构与内存层次结构

好，咱们直接进入正题。CUDA编程模型，说白了就是你怎么指挥GPU干活。很多初学者一上来就写kernel<<>>，但到底这俩参数怎么配，配错了会怎样，心里没底。今天我就把这块掰开揉碎了讲清楚。

2.1 线程层次结构：Grid / Block / Warp

GPU的线程组织方式，像军队编制。一个Grid是一个战役，Block是一个连队，Warp是一个班。咱们从大到小看。

2.1.1 Grid 和 Block

你启动一个Kernel，就是在定义一个Grid。Grid里包含若干个Block。每个Block里包含若干个线程。

// 定义Grid和Block的维度
dim3 gridDim(2, 2);    // 2x2 = 4个Block
dim3 blockDim(16, 16); // 16x16 = 256个线程/Block

// 启动Kernel
myKernel<<<gridDim, blockDim>>>(d_data);

这里有个关键点：Block内的线程可以协作，通过共享内存和同步操作。但不同Block之间的线程，理论上不能直接通信。为什么？因为Block可能被调度到不同的SM（流多处理器）上执行，它们之间没有硬件同步机制。

2.1.2 Warp：真正的执行单元

嗯，这里要注意。你写代码时感觉线程是独立的，但硬件上不是这么跑的。GPU实际执行的最小单位是Warp——32个线程为一组。

我刚开始做优化时，以为把Block设成100个线程就完事了。结果性能一塌糊涂。后来才意识到，Block大小必须是32的倍数，否则最后一个Warp会空转浪费资源。

Warp还有个特性叫Warp Divergence。同一个Warp里的线程，如果走不同的分支（if-else），那这些分支会串行执行。你想想看，32个线程里16个走if，16个走else，那效率直接砍半。

// 糟糕的写法：同一个Warp内线程ID奇偶不同，走不同分支
if (threadIdx.x % 2 == 0) {
    // 偶数线程
    do_something();
} else {
    // 奇数线程
    do_other();
}

// 更好的写法：尽量让Warp内线程走同一分支
int warpId = threadIdx.x / 32;
if (warpId % 2 == 0) {
    // 偶数Warp
    do_something();
} else {
    // 奇数Warp
    do_other();
}

2.2 内存层次结构

GPU的内存，说白了就是一套金字塔结构。越往上越快，但越小；越往下越慢，但越大。咱们从上往下看。

内存类型	位置	作用域	延迟	大小
寄存器	SM内部	单个线程	~1 cycle	每个SM有限（如65536个）
共享内存	SM内部	单个Block	~5 cycles	每个SM 48KB~164KB
L1缓存	SM内部	单个SM	~10 cycles	与共享内存共用
L2缓存	芯片级	所有SM	~100 cycles	几MB
全局内存	显存	所有线程	~400 cycles	几GB~几十GB
本地内存	显存	单个线程	~400 cycles	寄存器溢出时使用

我曾经在项目里遇到过一个坑：一个Kernel里用了太多局部变量，导致寄存器不够用，数据被自动 spill 到本地内存。性能从200μs直接掉到2ms。查了半天才发现是寄存器压力太大。

2.2.1 共享内存：Block内的"高速缓存"

共享内存是咱们做优化的利器。它比全局内存快两个数量级，而且可以由程序员手动管理。我习惯用它来做数据复用，比如矩阵分块、归约操作。

// 声明共享内存
__shared__ float s_data[256];

// 从全局内存加载到共享内存
int tid = threadIdx.x;
s_data[tid] = d_data[blockIdx.x * blockDim.x + tid];

// 同步，确保所有线程都加载完毕
__syncthreads();

// 在共享内存上做计算
// ...

这里有个细节：__syncthreads() 是Block内所有线程的屏障。它保证所有线程都执行到这一行之后，才能继续往下走。但要注意，它只对同一个Block内的线程有效。不同Block之间不能用它来同步。

2.3 Kernel启动流程

你写一个kernel<<<grid, block>>>()，背后发生了什么？我简单梳理一下。

1. CPU端准备数据：用cudaMalloc分配显存，用cudaMemcpy把数据从主机拷贝到设备。
2. 启动Kernel：CPU把Kernel函数和参数打包成一个"启动命令"，扔到GPU的命令队列里。注意，CPU不会等GPU执行完，它直接返回继续执行后面的代码。这是异步的。
3. GPU调度：GPU的GigaThread引擎把Grid拆成Block，分发给各个SM。每个SM再把Block里的线程分成Warp，调度到CUDA Core上执行。
4. 执行与同步：Warp内的线程以SIMT方式执行。Block内的线程可以通过__syncthreads()同步。整个Kernel执行完后，cudaDeviceSynchronize()会等待所有Block完成。
5. 结果回传：用cudaMemcpy把结果从设备拷贝回主机。

2.4 知识体系总览

下面这张图，我把本章的核心逻辑串起来了。你一看就明白。

这张图里，Grid包含多个Block，每个Block包含多个Warp，每个Warp包含32个线程。内存方面，越靠近线程的越快，但容量越小。这就是你写CUDA代码时，脑子里要时刻绷紧的那根弦。

好了，这一章就到这里。记住：线程层次决定了你怎么组织并行，内存层次决定了你能跑多快。下一章咱们聊Kernel的启动配置和性能调优，到时候我会拿几个实际案例出来讲。