CUDA 的线程层次是一套四级组织：Grid（网格）→ Block（线程块）→ Warp（线程束）→ Thread（线程），其中 Warp 是硬件真正调度的最小单位，固定 32 个线程一组。
一句话讲清调度逻辑
GPU 的调度器不按单个线程发指令，而是以 32 个线程为一束，同时丢给它们同一条指令——这套机制叫 SIMT（Single Instruction, Multiple Threads，单指令多线程）。就像一位工头对着并排站的 32 个工人喊口令，同一时刻只能喊一条。
打个比方：32 个工人共用一位工头
让他们一起搬砖，口令只有一条，32 个人同时动手，效率拉满。但如果你写的是「戴安全帽的往前一步，没戴的往后退一步」，工头没法同时喊两条，只能先喊「戴帽的走」，让没戴的站着发呆；再喊「没戴的退」，让戴帽的站着发呆。两条路串起来跑，耗时接近两倍——这就是 warp 分歧（warp divergence）。
关键点在于：分歧的代价不取决于分支数量，而取决于同一个 warp 里有没有线程走了不同的路。如果 32 个线程全走同一条分支，一点浪费都没有；只要有一个线程掉队，整个 warp 就得陪着把两条路都走一遍。
四个层次分别由谁决定
| 层次 | 规模 | 谁说了算 | 主要职责 |
|---|---|---|---|
| Thread | 1 个执行单元 | 硬件按 warp 打包调度 | 处理一个元素 |
| Warp | 32 个线程 | 硬件固定切分 | 真正被发射指令的单位 |
| Block | 若干 warp | 程序员指定 | 共享内存、__syncthreads() 同步 |
| Grid | 若干 Block | 程序员指定 | 一次内核启动的全部线程 |
由此可以推出几条常见规律：Block 内线程能通过共享内存（shared memory）交换数据并同步，Warp 内线程靠 __shfl 这类指令互换；Block 之间基本只能走全局内存（global memory），所以通信密集的计算要尽量塞进同一个 Block。块大小取 32 的倍数（128、256 这类）也是这个道理，别切出半截 warp 白占名额。
对从业者的实际意义
同一段内核代码性能差十倍，往往不是算法问题，而是 warp 里那几个线程走岔了。可操作的写法有三种：
1. 消灭分支：把 if/else 改成算术表达式、查表或谓词赋值，让 32 个线程执行同一条指令流。
2. 对齐分支：分支躲不掉时，按条件对数据分桶或分区，让判断结果相同的线程落进同一个 warp，这样每个 warp 只走一条路。
3. 验证假设：用性能分析工具确认分歧是不是真瓶颈，别凭感觉优化。工具的具体名称与指标以官方文档为准。
对不写内核的职场人，这套模型也有借用价值：并行任务里最贵的往往不是计算本身，而是大家步调不一致。分派工作时，先想清楚哪些人必须共用同一节奏，比单纯堆人力更管用。