4. CUDA线程模型与索引：线程层次结构详解、多维网格与块索引计算、线程束（Warp）与分支发散、Occupancy计算与优化
好，咱们今天聊点硬核的。CUDA的线程模型，说白了就是GPU怎么组织那成千上万个干活的小工。我刚开始接触CUDA时，总觉得这东西不就是开一堆线程嘛，有啥好讲的？结果第一次写核函数，数据全算错了——索引搞反了，数据读串了。嗯，从那以后，我再也不敢小看线程模型了。
4.1 线程层次结构：Grid、Block、Thread
CUDA把线程分成三层：Grid（网格）、Block（线程块）、Thread（线程）。你想想看，一个Grid就是一个大任务，里面包含多个Block，每个Block又包含一堆Thread。
我个人习惯把Grid想象成一个工厂车间，Block是车间里的工作台，Thread就是工作台上的工人。工人之间可以共享工作台上的工具（共享内存），但不同工作台的工人不能直接交流。
关键点：线程层次结构决定了数据的组织方式和访问模式。选错了层次，性能直接腰斩。
代码里怎么声明呢？很简单：
// 定义网格和块的大小
dim3 gridDim(2, 2); // 2x2的网格，共4个块
dim3 blockDim(16, 16); // 16x16的块，共256个线程
// 启动核函数
myKernel<<<gridDim, blockDim>>>(d_data);这里dim3是个三维向量，但第三维默认为1。实际项目中，我很少用到第三维，除非处理3D体数据。
4.2 多维网格与块索引计算
索引计算是新手最容易翻车的地方。每个线程都有自己唯一的ID，怎么算？
先看内置变量：
threadIdx.x、threadIdx.y、threadIdx.z：线程在块内的索引blockIdx.x、blockIdx.y、blockIdx.z：块在网格内的索引blockDim.x、blockDim.y、blockDim.z：每个块的维度大小gridDim.x、gridDim.y、gridDim.z：网格的维度大小
计算全局线程ID的公式：
// 一维情况
int tid = blockIdx.x * blockDim.x + threadIdx.x;
// 二维情况
int tid_x = blockIdx.x * blockDim.x + threadIdx.x;
int tid_y = blockIdx.y * blockDim.y + threadIdx.y;
int global_id = tid_y * (gridDim.x * blockDim.x) + tid_x;我的经验：处理二维矩阵时，我习惯让x对应列，y对应行。这样内存访问是连续的，性能更好。曾经有个项目，同事把xy搞反了，结果带宽利用率从80%掉到20%，排查了两天才找到原因。
三维索引类似，只是多了一个z维度。不过说实话，三维用得少，因为大多数问题可以映射到一维或二维。
4.3 线程束（Warp）与分支发散
这是GPU架构的核心概念。线程束（Warp）是GPU调度和执行的基本单位，NVIDIA的GPU里，一个Warp包含32个线程。
什么意思呢？就是这32个线程在硬件上是“绑在一起”的，它们执行同一条指令，只是处理不同的数据。这就是SIMT（单指令多线程）模型。
注意：如果Warp内的线程走了不同的分支（比如if-else），就会发生分支发散。GPU会串行执行所有分支路径，性能直接打折扣。
举个例子：
__global__ void divergent_kernel(float* data, int N) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid < N/2) {
data[tid] = data[tid] * 2.0f; // 路径A
} else {
data[tid] = data[tid] * 0.5f; // 路径B
}
}这段代码里，如果Warp内的线程有的走路径A，有的走路径B，那这个Warp就要执行两次：先执行路径A，再执行路径B，最后合并结果。性能直接减半。
怎么避免？我常用的方法：
- 让分支条件基于Warp对齐：比如用
tid / 32而不是tid做判断 - 重新组织数据：把走同一分支的数据放在一起
- 使用三元运算符：有时候编译器能优化成无分支代码
避坑指南：我曾经在一个图像处理项目里，用线程ID的奇偶性做不同处理，结果性能惨不忍睹。后来改成按块分配，每个块处理连续区域，分支发散问题就解决了。
4.4 Occupancy计算与优化
Occupancy，中文叫“占用率”，指的是SM（流多处理器）上活跃的Warp数量与最大支持Warp数量的比值。说白了，就是SM有多“忙”。
占用率越高越好吗？不一定。我见过有人为了追求100%占用率，把共享内存用得极少，结果每个线程能用的寄存器也少了，反而导致寄存器溢出到全局内存，性能更差。
占用率受三个因素限制：
| 资源 | 说明 | 影响 |
|---|---|---|
| 寄存器 | 每个线程使用的寄存器数量 | 寄存器越多，能同时运行的Block越少 |
| 共享内存 | 每个Block使用的共享内存大小 | 共享内存越大，能同时运行的Block越少 |
| 线程数 | 每个Block的线程数量 | 线程数太少，SM无法填满 |
怎么计算占用率？NVIDIA提供了CUDA Occupancy Calculator，也可以用cudaOccupancyMaxPotentialBlockSize函数：
int minGridSize, blockSize;
cudaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize,
myKernel, 0, 0);
// blockSize就是推荐值我的建议：不要盲目追求高占用率。我一般先算一下理论占用率，然后实际跑一下看看吞吐量。有时候70%的占用率配合更好的内存访问模式，比90%占用率但内存乱访问要快得多。
优化占用率的常用手段：
- 调整Block大小：通常是128或256的倍数，且是32的倍数
- 减少寄存器使用：用
__launch_bounds__限制每个线程的寄存器数 - 减少共享内存：能复用就复用，别每个Block都开一大块
- 使用动态共享内存：按需分配，避免浪费
嗯，这里要注意：__launch_bounds__这个属性可以告诉编译器你期望的Block大小，编译器会尽量优化寄存器分配。但别设得太小，否则寄存器不够用，编译器会 spill 到本地内存，反而更慢。
4.5 知识体系总览
下面这张图是我自己画的，把本章的核心逻辑串起来了：
这张图把Grid、Block、Thread的层次关系画清楚了。Warp和Occupancy是影响性能的两个关键因素，一个管执行效率，一个管资源利用率。
好了，这一章的内容就到这儿。记住：线程模型是CUDA编程的基石，索引算不对，后面全白费。多写几个测试核函数，把索引打印出来看看，比死记公式管用得多。