CUDA 线程层次：grid、thread block 与 thread 的三级组织

更新时间：2026-09-25。写第一个 CUDA 程序时，很多人能照着教程敲出 kernel<<<N, M>>>(...)，却说不清尖括号里两个数到底切出了多少线程、线程之间谁认识谁。本文是「集成电路大类 · 芯片设计」进阶层（intermediate） 的一篇，专门拆开 CUDA 编程模型最核心的抽象——CUDA 线程层次（thread hierarchy）：一次 kernel 启动是一个 grid，grid 切成若干 thread block，block 内再排满 thread。这个层次不是语法糖，它直接规定了哪些线程能共享内存、哪些线程能互相等待、以及线程怎样被装进 GPU 的流多处理器。第一个 kernel 的写法见 CUDA 第一个 kernel，硬件侧 SM 的内部结构见 SM 与 CU 的内部结构，warp 调度的整体架构见 GPU 架构与 warp 调度器。

本文要回答的问题

• grid、thread block、thread 三级各是什么，一次 kernel 启动到底创建了多少线程？
• gridDim、blockDim、blockIdx、threadIdx 四个内置变量分别描述哪个层级，为什么都支持 x/y/z 三个分量？
• 处理二维图像时，线程怎样从「block 内坐标」换算成「整张图的像素坐标」？
• warp 和 block 是什么关系，为什么说 block 内连续 32 个线程总被绑在一起执行？
• 为什么 block 内可以用 __syncthreads() 做屏障、共享 shared memory，而 block 之间却几乎不能同步？block 大小为什么常用 256？

一、三级层次：grid、block、thread

CUDA 编程模型用三级结构描述线程。thread（线程） 是最小执行单位，每个 thread 执行一份 kernel 代码、拥有自己的寄存器和局部变量。thread block（线程块，简称 block） 是一批被承诺「可以协同」的线程：它们共享一块片上 shared memory（共享内存），可以用 __syncthreads() 做屏障同步，运行时整个 block 被调度到同一个 SM（streaming multiprocessor，流多处理器） 上。grid（线程网格） 则是一次 kernel 启动（launch）产生的全部 block 的总和——启动语句 kernel<<<gridDim, blockDim>>> 返回后，grid 的形状就固定了。

读图抓住两条线：纵向上 grid 只是 block 的容器，grid 本身不提供任何共享资源；横向上 block 内部的线程再按连续 32 个一组划成 warp。以最常见的配置为例：启动时每个 block 配 256 个线程、grid 配若干 block，那么一个 block 就是 256 个 thread、恰好 8 个 warp；grid 里有多少个 block，总线程数就是 block 数乘以 256。block 的数量往往按数据规模倒推，而不是按 GPU 大小硬编码——这正是 CUDA 可扩展性的关键，第六节会展开。

可以用一个 host 与 device 的对比加深理解：CPU 端写循环时，循环变量 i 在时间上依次取遍 0、1、2……由一个核顺序执行；GPU 端则把「这一整串 i」一次性物化成成千上万个并发 thread，每个 thread 拿到一个固定编号，空间上同时铺开。CPU 程序员关心的是「循环现在转到第几轮」，CUDA 程序员关心的是「我这个线程的编号对应哪份数据」。三级层次就是这套编号体系：grid 回答总共铺开了多少批，block 回答我属于哪一批，thread 回答我在批里排第几。总线程数的算式固定为 gridDim.x × gridDim.y × gridDim.z × blockDim.x × blockDim.y × blockDim.z，启动后不变，排错时可以先拿这个数和数据元素数对一遍。

三级在「作用域、共享资源、同步手段」上的差别，是写 kernel 时最容易踩错的地方：

层级	数量/形状由谁决定	可见的共享存储	同步与通信手段	调度归宿
thread	block 形状切分	私有寄存器、local memory	原子操作（低效）	block 内某条 lane
block	grid 形状切分	本 block 的 shared memory	__syncthreads() 屏障、shared memory	整体驻留同一个 SM
grid	host 端 <<<...>>> 启动参数	全局 global memory（显存）	kernel 之间按启动次序；kernel 内无原生屏障	block 散布到全部 SM

表里还有一个值得点明的不对称：thread 之间虽然各自独立，但处在同一 warp 内时执行节奏被 SIMT 绑死，严格说并不存在「线程间完全无关」这一档；而 grid 层面连「另一个 block 是否已经开始」都无从得知。初学者常把 grid 想象成 CPU 上的线程池、以为 block 像池里的 worker 一样可以互相打招呼，这个类比在共享与同步两栏上正好是错的——grid 更像一张一次性派发出去的任务传单，每张传单（block）独立做完，靠 host 最后回收。

记住一条判据：凡是写在 __shared__ 内存里的数据，出了本 block 没人看得见；凡是 __syncthreads()，只能拦住同一个 block 的线程。 grid 级别的协作只能绕道 global memory，且要拆成多次 kernel 启动，第六节解释原因。

二、四个内置变量：每个线程怎样知道「我是谁」

kernel 代码在成千上万线程上跑同一份，线程怎样分辨自己该处理哪份数据？靠四个编译器内置的向量型变量，类型都是 uint3（含 x、y、z 三个无符号分量）：

• gridDim：grid 的形状，即每个维度上 block 的数量；
• blockDim：每个 block 的形状，即每个维度上 thread 的数量；
• blockIdx：本线程所在 block 在 grid 里的三维坐标；
• threadIdx：本线程在 block 内的三维坐标。

前两个对所有线程读到的值相同，描述「舞台有多大」；后两个每个线程读到的值不同，回答「我站在哪」。host 端启动时用 dim3 描述形状：dim3 是个三维无符号整数结构，未写出的分量默认为 1。kernel<<<dim3(4,2), dim3(16,16)>>>(...) 表示 grid 有 4×2×1 共 8 个 block，每个 block 有 16×16×1 共 256 个线程，整次启动 8×256=2048 个线程。启动参数与内置变量严格对应：尖括号第一个参数就是设备端读到的 gridDim，第二个就是 blockDim。

一个常被忽略的细节：block 内线程数和形状必须在编译后能被 host 端确定地作为启动参数传入，kernel 内部不能改变自己的 blockDim；线程一旦启动，所属的 block 和 block 内编号终生不变，不存在「线程跳槽到另一个 block」的机制。因此 block 是线程身份的一部分，而不只是调度分组。另外，线程私有的并不只有寄存器：寄存器不够时变量会溢落到 local memory——名字叫「local」，物理上却在显存里，访问延迟与 global memory 相当，这也是第六节寄存器压力影响占用率之外的另一条惩罚。

三维不是为了好看，而是让线程坐标天然贴合数据的维度：图像用二维 (x, y)，体数据或视频帧序列用三维 (x, y, z)。匹配后索引换算就是几条乘法，少写一层手写的扁平化。当然三维也有上限：现行硬件上 block 的 z 维最大 64，单 block 线程总数不超过 1024；grid 的 x、y 维可达 2³¹−1。所以 16×16=256、32×8=256 这类二维 block 合法，而 16×16×4=1024 是单 block 的天花板。

三、二维与三维索引映射：以图像 tile 为例

CUDA 里最经典的模式是让 block 对应数据的一个 tile（瓦片/分块），block 内线程对应 tile 内的元素。处理一张宽 W、高 H 的图像，每个 block 负责 16×16 的 tile，线程到像素坐标的换算如下：

dim3 block(16, 16);
dim3 grid((W + 15) / 16, (H + 15) / 16);   // 向上取整，覆盖边缘
kernel<<<grid, block>>>(d_img, W, H);

__global__ void kernel(unsigned char* img, int W, int H) {
  int x = blockIdx.x * blockDim.x + threadIdx.x;  // 列：全局 x
  int y = blockIdx.y * blockDim.y + threadIdx.y;  // 行：全局 y
  if (x < W && y < H)        // 边缘 tile 越界线程直接退出
    img[y * W + x] += 1;
}

公式可以读成「前面所有 block 的线程数，加上我在 block 内的偏移」：blockIdx.x * blockDim.x 是本 block 左边界之前有多少列，threadIdx.x 是 block 内第几列。取整公式 (W+15)/16 保证宽不是 16 整数倍时边缘也有 block 覆盖，多出的线程靠 if (x < W && y < H) 挡掉——这是每个 CUDA 程序员都要背的边界判断。

用具体数字走一遍。1024×1024 的图像、16×16 的 tile：grid 形状为 64×64，共 4096 个 block；每 block 256 线程，整次启动 4096×256=1 048 576 个线程，恰好一个线程对应一个像素。换成 1920×1080：grid 为 120×68（1080/16=67.5，向上取整），共 8160 个 block、2 088 960 个线程，其中最后一行 tile 里有一半线程会被边界判断挡掉。浪费这一点线程换来代码无条件分支之外的简洁，完全值得。

三维场景用体数据（volumetric data）走一遍：512×512×64 的三维 CT 体数据，block 取 8×8×4=256 线程，则 grid 为 64×64×16，共 65 536 个 block、16 777 216 个线程，与体素总数相等。线程的全局坐标写成 (blockIdx*blockDim + threadIdx) 的三个分量后，体素地址就是 z*W*H + y*W + x。可以看到维度越高，手写扁平化越容易出错，而用内置变量分量逐项换算，代码几乎是数据形状的直译。反过来，若数据本身是一维的，就老实用一维 block，硬套三维只会让索引导入无谓的乘法。

三维的扁平化同理。若用 dim3(Tx, Ty, Tz) 描述 block，block 内线性编号是 threadIdx.z * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x。注意这个次序不是随便选的：x 维变化最快，与硬件分配 warp 号的方式一致——第四、五节会看到，连续的线性编号正好对应连续的 warp 成员。

四、warp：block 内连续 32 个线程

block 是程序员看到的协作单位，硬件真正调度的单位却是 warp（线程束）：一个 block 内线性编号连续的 32 个线程组成一个 warp，warp 号为 threadIdx 线性编号 / 32，线程在 warp 内的位置叫 lane（0–31）。256 线程的 block 天然切成 8 个 warp：T0–T31 是 warp0，T32–T63 是 warp1，一直到 T224–T255 是 warp7。warp 的划分完全由 threadIdx 决定，程序员不能指定，也不能拆。

执行方式叫 SIMT（single-instruction, multiple-thread，单指令多线程）：一个 warp 内 32 个线程在 SM 上共享同一个程序计数器、同一条指令，各自作用于自己的寄存器数据。正因为绑在一起，当 warp 内线程走到 if 的不同分支时，硬件只能把两条路径串行跑完，这就是 warp divergence（分支发散）——同 warp 有人真有人假时，实际指令数接近两条路径之和。完整机制与优化见 SIMT 与 warp 分支发散。

为什么偏偏是 32？这个数字由硬件 SIMT 通道数决定，NVIDIA 各代 GPU 一直保持 warp 宽度为 32（别家产品有 16 或 64 的波前），可以理解为调度器一次发一条指令、最多驱动 32 个数据通路。对程序员来说它的实际意义是「block 内资源以 32 为粒度成团」：分支按 warp 发散，访存按 warp 合并，warp 内线程的起始编号永远是 32 的倍数。warp 内的 lane 号可以直接用 threadIdx.x % 32 之类的写法推出，CUDA 也提供 __shfl_* 系列 warp 级原语让 32 个 lane 不经过 shared memory 直接交换寄存器数据——这类原语能用的前提，正是线程清楚自己处在哪一个 32 人小队里。

这也反过来解释了二维 block 的线程排布为什么讲究。16×16 的 block 中，warp0 是 x 坐标 0–15、y=0 的 16 个线程加 y=1 的前 16 个，即 warp 总是横跨两行；因此让相邻 x 坐标的线程访问相邻内存地址，一个 warp 的 32 次访存就能合并成少数几个连续事务（coalescing）。索引公式里让 x 维变化最快，正是在配合这条硬件规则。排错时一个实用检查是手算几个代表性线程的 warp 号与 lane 号：16×16 block 中 threadIdx=(0,0) 与 (15,0) 同属 warp0，(0,1) 也在 warp0 里，(0,2) 才进入 warp1；把这张对应关系在纸上列一次，对安排访存和判断分支粒度很有帮助。

五、block 内协同：shared memory 与 __syncthreads

block 被承诺「整个 block 调度到同一个 SM」，于是同 block 的线程共享一块 SM 片上的高速存储 shared memory，并用屏障 __syncthreads() 协作。典型套路是分块（tiling）：block 内所有线程先合力把一片 global memory 数据搬进 shared memory，调用一次 __syncthreads() 等所有人搬完，再各自在片上高速地复用这份数据。__syncthreads() 的语义是「block 内所有线程都到达这里之后，才允许任何线程越过它」，并且保证屏障前对 shared memory/global memory 的写，对屏障后的所有线程可见——它同时是执行屏障和可见性屏障。

图中要点是「私有」二字：shared memory 按 block 分配、随 block 结束而回收，Block B 即使驻留在同一个 SM 上也看不到 Block A 的内容。屏障同理只在 block 内闭合。两个使用陷阱值得牢记：其一，__syncthreads() 必须被 block 内所有线程无条件到达——把它放进某个线程不一定走的分支里，没到达的线程永远等不到放行，程序直接挂死；其二，它只管 block，想让 block 0 算完再让 block 1 读，屏障帮不上忙。

把 tiling 套路压成最小骨架，下面的代码让 block 给一行数据做就地放大：每个线程负责搬一个元素进 tile，屏障保证复用发生在搬完之后：

__shared__ float tile[256];
tile[threadIdx.x] = input[blockIdx.x * 256 + threadIdx.x];
__syncthreads();                       // 256 线程全到齐才往下
float v = tile[threadIdx.x] + tile[(threadIdx.x + 1) & 255];

如果删掉中间的屏障，线程 1 可能在线程 0 的写入对它可见之前就读 tile[0]，结果随调度时序漂移——这种「有时对有时错」的 bug 在 CUDA 里极难复现，屏障与 shared memory 必须成对出现。

线程在 SM 上的排布到此完整：256 线程的 block 进 SM，先按 32 切成 8 个 warp，由 warp scheduler 在 warp 间切换隐藏访存延迟；block 的 shared memory 和寄存器在整个驻留期间被这 256 个线程独占。SM 内部寄存器堆、CUDA 核心与调度器的物理分工见 SM 与 CU 的内部结构。

六、block 大小与占用率，以及 block 间为什么不能同步

先回答「为什么常用 256」。SM 同时驻留的 warp 越多，scheduler 越容易在某个 warp 等内存时找到就绪 warp，延迟隐藏越好；驻留 warp 数占硬件上限的比例叫 occupancy（占用率）。SM 的并发能力同时受四道资源约束：最大驻留线程数（当代主流数据中心 GPU 为每 SM 2048）、最大驻留 block 数（常见 16 或 32）、每线程寄存器数、每 block 的 shared memory 容量。256 线程的 block 在 2048 上限下恰好可驻留 8 个，凑成 64 个 warp；取 128 则可能撞 block 数量上限，取 512 则每个 SM 只有 4 个 block、调度弹性变差，取 1024 时单个 block 就要吃掉一半线程名额。256 同时是 32 的整数倍、二维上方便拆成 16×16，因此成为事实标准默认值，但不是教条——寄存器用得多时减小 block、shared memory 是瓶颈时调整 block，都要以实际 occupancy 为准。

再回答「block 间为什么不能同步」。根源是 CUDA 对执行次序的承诺极弱：grid 内的 block 之间没有任何先后保证，硬件按 SM 空闲情况随意派发，同一个 block 甚至可能在资源紧张时整体换出。如果语言提供 block 间屏障，就可能出现 block A 占着 SM 等 block B、而 block B 永远排不上 SM 的全局死锁；为避免这种情况，编程模型干脆不保证 block 并发存活，于是 block 间屏障在语义上无法成立。工程上的替代方案有三种：把协作阶段拆成两次 kernel 启动（kernel 之间天然有序，global memory 传递中间结果，最常用）；需要大数据复用的场景使用 cooperative groups 的 grid-wide 同步（要求所有 block 同时驻留，可启动规模受限）；或者干脆按「block 之间完全独立」重设计算法。

把寄存器账也算一遍。SM 的寄存器堆总量固定（各代约 64K 个 32 位寄存器/ SM 量级），256 线程的 block 若每线程用 64 个寄存器，一个 block 就要 16 384 个；8 个 block 全部驻留需 131 072 个，超过容量，硬件就会少放 block，占用率随之跌破满载。编译器可以用 -Xptxas -v 打印每 kernel 的寄存器用量，超了就精简 kernel 内的活变量、或显式用 __launch_bounds__ 提示上限。shared memory 同理：每 block 分到 48KB 时，SM 上百 KB 的片上共享内存只够养两个 block，block 大小再合理也救不回占用率。所以「256 是默认值」的完整说法是：256 只解决线程和 warp 粒度，真正的占用率要拿 profiler 把寄存器、shared memory、block 数四项一起看。

这也解释了第一节那句「block 数量按数据规模倒推」。正因为 block 互不依赖、可任意派发，同一个 kernel 启动配置在高端 GPU 的上百个 SM 和低端芯片的几个 SM 上都能正确运行，只是快慢不同——CUDA 的可扩展性正是建立在「block 是独立执行单元」这一承诺上。代价就是：凡是跨 block 的协作，都得由程序员显式拆阶段。

这一篇在知识体系里的位置

本篇属于芯片设计进阶层的 GPU 编程模型专题，回答「一次 kernel 启动的线程按什么规则组织、协作边界划在哪」。动手前置是 CUDA 第一个 kernel；横向配套，block 驻留 SM 的物理图景见 SM 与 CU 的内部结构，warp 调度与延迟隐藏的整体架构见 GPU 架构与 warp 调度器，warp divergence 的机制与应对见 SIMT 与 warp 分支发散；向后读，shared memory 分块、bank conflict、occupancy 调优与内存合并，都是在本篇三级层次之上的具体优化手段。

总结

一句话总结：CUDA 线程层次把一次 kernel 启动组织成 grid→thread block→thread 三级——grid 是全部线程的容器、不提供共享资源，block 是共享 shared memory、可用 __syncthreads() 屏障且整体驻留同一 SM 的协作单位，thread 靠 gridDim、blockDim、blockIdx、threadIdx 四个内置变量定位；二维/三维坐标让 block 对应数据 tile，全局索引按「前面 block 的线程数加块内偏移」换算，边缘线程用边界判断挡掉；block 内连续 32 个线程组成一个 warp 并以 SIMT 方式绑定执行，256 线程/block 即 8 个 warp，这个大小能在 SM 的线程、block、寄存器、shared memory 四重限制下取得较好占用率，因而成为默认选择；block 间没有执行次序保证、也不保证同时存活，所以无法安全实现 block 间同步，跨 block 协作应拆成多次 kernel 启动或使用 cooperative groups。