想象一下，你是一家大型快递分拣中心的调度主管。整个中心（Grid）接到一批包裹，需要分成好几个班组（Block）同时干活；每个班组里又有几十名分拣员（Thread）各司其职，有人扫码、有人搬运、有人装车。为了让协作顺畅，同一班组的工人可以围在公共操作台（shared memory）旁快速传递包裹，而每个工人自己手里的小抄本（register）则记录专属信息。在这座GPU工厂里，硬件不会一个个单独指挥工人，而是让组长（warp）带着固定32人的小队同步行动——这就是CUDA线程层次的核心直觉。

岗位关联

理解Thread/Block/Grid的职责与限制，是写出高效CUDA算子的第一道门槛。掌握它，你就能像GPU调度主管一样，把成千上万个线程安排得明明白白，为AI大模型、图像处理等高性能场景提供计算引擎。

目录
  1. 可视化演示
  2. 层次总览
  3. Thread职责
  4. Block协作
  5. Grid调度
  6. Warp执行
  7. 维度限制
  8. 资源约束
  9. 实战配置
  10. AI 应用场景
  11. Python 演示
  12. 常见误区
  13. 面试题

可视化演示

展示Thread→Block→Grid三级线程层次嵌套、Warp调度机制（32线程）、共享内存与寄存器对Block大小的资源约束。切换模式深入理解GPU并行计算的组织结构。

本节配有交互动效（进入视口后自动加载，约 3KB）。

Thread→Block→Grid三级线程层次 Grid包含多个Block，Block包含多个Thread，三级嵌套构成CUDA并行计算核心组织结构 Block内线程共享内存和同步，Grid内Block独立调度，每层有明确职责和资源约束 Grid → Block → Thread, blockDim × gridDim = total_threads 三级层次是编写高效CUDA程序的基础抽象

层次总览

回到那家快递分拣中心。整个中心接收来自全国各地的包裹，这是一个完整的任务（Grid）；中心里分成多个班组（Block），每个班组负责一条流水线；流水线上每个工人（Thread）只完成“扫码”“搬运”“装车”中的一件小事。工人被管理在班组里，班组被管理在中心里，三级嵌套，缺一不可。CUDA的线程层次正是这样：Grid包含若干Block，Block包含若干Thread。GPU运行一个核函数时，会同时启动一个Grid，Grid里所有Block被分配到各个流式多处理器（SM）上，每个Block里的Thread则被硬件调度执行。

这种层级不是人为设定的仪式，而是为了同时解决“规模”和“协作”两个问题：Grid决定了任务的总规模，Block提供了可通信的小团队，Thread就是实际的执行者。硬件上，一个Block会被放到一个SM上，同一Block的线程能共享快速存储并同步；不同Block之间则相互独立，只能通过全局内存间接沟通。

一个Grid总共能调度多少个并行线程？用乘法就能算出来：

了解总览之后再看具体的每一层：先看最底层干活的人——Thread。

total_threads=gridDim.x×gridDim.y×gridDim.z×blockDim.x×blockDim.y×blockDim.z\text{total\_threads} = \text{gridDim.x}\times\text{gridDim.y}\times\text{gridDim.z}\times\text{blockDim.x}\times\text{blockDim.y}\times\text{blockDim.z}
公式的意义：GPU的并行规模 = 网格的块数 × 每块的线程数，二者相乘就是一次核函数启动能产生的线程总数。
  1. gridDim.x/y/zGrid在三维方向各有多少个Block，例如 gridDim.x=4 表示 x 方向有 4 个 Block。
    gridDim.x×gridDim.y×gridDim.z\text{gridDim.x}\times\text{gridDim.y}\times\text{gridDim.z}
  2. blockDim.x/y/z每个Block在三维方向各有多少个Thread，例如 blockDim.x=256 表示一块有 256 个线程。
    blockDim.x×blockDim.y×blockDim.z\text{blockDim.x}\times\text{blockDim.y}\times\text{blockDim.z}
  3. 两部分相乘Block总数与每Block线程数相乘，得到一次启动的全部线程总数。
    total_threads=blocks_per_grid×threads_per_block\text{total\_threads} = \text{blocks\_per\_grid} \times \text{threads\_per\_block}

Thread职责

工人是整个分拣中心最基础的力量。每个工人手里有一张小工牌，上面写着自己在班组里的编号：threadIdx.x、threadIdx.y、threadIdx.z。工人们也想知道自己在整个中心里的唯一编号——比如“你是第二车间、第5个工人”——这样才能知道自己该处理哪件包裹。CUDA里，每个Thread用一个三元组 threadIdx 定位自己在Block中的位置，再用 blockIdx 知道自己属于哪个Block。

Thread是CUDA的最小执行单元。你写的那段“核函数（kernel）”代码，会被GPU上成千上万个Thread同时执行。每个Thread通过 threadIdx 与 blockIdx 计算出自己负责的数据位置，从而避免重复劳动。最常见的1D写法是：

工人不能永远单干。同一班组的工人需要传递半成品、互相核对数据，这就引出Block的核心价值——协作。

global_id=blockIdx.x×blockDim.x+threadIdx.x\text{global\_id} = \text{blockIdx.x} \times \text{blockDim.x} + \text{threadIdx.x}
只要知道“我属于第几个块、每块多少人、我是块内第几个”，就能算出全Grid的唯一ID。
  1. blockIdx.x当前线程所在Block在Grid中的序号，相当于“第几个车间”。
    blockIdx.x\text{blockIdx.x}
  2. blockDim.x每个Block在x方向上有多少个线程，相当于“一个车间有多少人”。
    blockDim.x\text{blockDim.x}
  3. 乘积排在当前Block之前的所有Block的线程总数。
    blockIdx.x×blockDim.x\text{blockIdx.x} \times \text{blockDim.x}
  4. 加偏移加上自己在Block内的位置，得到全局唯一的线程ID。
    global_id=blockIdx.x×blockDim.x+threadIdx.x\text{global\_id} = \text{blockIdx.x} \times \text{blockDim.x} + \text{threadIdx.x}

Block协作

一个班组之所以是“组”，是因为工人需要共用一张操作台。这张台子就是CUDA里的共享内存（shared memory）。同一Block内的所有线程都能快速读写共享内存，速度远快于全局内存，就像工人从身边的台子拿料而不是跑几百米去仓库。GPU把一个Block分配给一个SM后，Block内的线程就能通过共享内存协作：有人负责从全局内存加载数据，有人负责计算，有人负责写回。

为了让协作不出错，CUDA提供了同步屏障 __syncthreads()。它像班组长吹哨子：“所有人都把手上的活干完，再一起进入下一步！” 只有同一Block内的所有线程到达这个点，程序才会继续。不同的Block之间不能互相干预，也不能同步，就像不同车间有不同的班组长，互不指挥。

这也解释了为什么Block不能太大也不能太小：太小则共享内存和同步带来的收益被摊薄；太大则硬件资源放不下。通常Block维度会设计成能对应数据的一个“切片”或“瓦片（tile）”。比如处理一张图像时，让一个Block负责一块 16×16 的像素区域，Block内线程用共享内存互相配合。

Grid调度

Grid是整个任务的最高指挥层。当你启动一个核函数时，实际上是在告诉GPU：“我要建一个分拣中心，里面放多少个车间。” 这些车间（Block）会被分配到GPU的各个SM上并行执行。一个SM可能同时驻留多个Block，但一个Block绝不会跨SM运行——这是硬件规则。

每个Block有自己的编号 blockIdx，也是三维的。Block数量由 gridDim.x/gridDim.y/gridDim.z 决定。在1D场景下，Block总数可以用“数据总量 / 每块线程数”向上取整得到：

Grid把任务分给Block后，硬件层面还需要一个“调度员”来真正指挥线程——它就是Warp。

blocks_per_grid=⌈NblockDim.x⌉\text{blocks\_per\_grid} = \left\lceil \frac{N}{\text{blockDim.x}} \right\rceil
数据有 N 份，每块处理 blockDim.x 份，需要的块数就是 N 除以块线程数后向上取整。
  1. N待处理的数据总量，比如 1000 个向量元素。
    N=1000N = 1000
  2. blockDim.x每个Block的线程数，比如 256。
    blockDim.x=256\text{blockDim.x} = 256
  3. 相除向上取整整数除法必须向上取整，否则最后一块数据没人处理。
    ⌈NblockDim.x⌉\left\lceil \frac{N}{\text{blockDim.x}} \right\rceil

Warp执行

工人们虽然按班组组织，但真正执行时，GPU并不一个一个指挥工人，而是让“组长”一次带32个工人同步行动。这个32人的小队就是Warp。它是GPU硬件调度的最小单位：一个线程束（warp）里的32个线程，在大多数情况下会同时执行同一条指令。这被称为SIMT（单指令多线程）执行模式——组长的口令只有一句，32个工人同时听令。

一个Block由若干个warp组成。块内线程数不一定是32的倍数，但非整数倍会造成最后一个warp部分线程空闲，浪费计算能力。warp数量计算公式为：

既然warp是硬件调度的真实单位，我们在设计Block时就应该考虑32的倍数，少留空位。

warps_per_block=⌈threads_per_block32⌉\text{warps\_per\_block} = \left\lceil \frac{\text{threads\_per\_block}}{32} \right\rceil
比如 64 线程的Block = 2 个warp；1000 线程的Block = 32 个warp（最后一个warp不满员）。
  1. threads_per_block当前Block的线程总数，即三种线程维度乘积。
    threads_per_block=blockDim.x×blockDim.y×blockDim.z\text{threads\_per\_block} = \text{blockDim.x}\times\text{blockDim.y}\times\text{blockDim.z}
  2. 除以32每32个线程组成一个warp。
    threads_per_block32\frac{\text{threads\_per\_block}}{32}
  3. 向上取整如果线程数不是32的整数倍，仍然需要额外一个warp承载剩余线程。
    ⌈threads_per_block32⌉\left\lceil \frac{\text{threads\_per\_block}}{32} \right\rceil

维度限制

分拣中心不是无限大的。CUDA对Grid和Block的三维规模有硬性上限，就像消防规定一个车间最多容纳多少人、一块场地最多放多少个车间。最常见也最重要的约束是：

  • 每个Block最多 1024 个线程，且三维乘积不能超过1024；
  • Block的每个维度也有限制：x方向最大1024，y方向最大1024，z方向最大64；
  • Grid的x方向最大 2^31-1（约21亿），y和z方向最大 65535。

为什么会有这些限制？因为硬件要给每个Block预留共享内存、寄存器等资源，还要给线程编号留足索引空间。设计配置时，不要试图挑战上限，而应“够用就好、留有余量”。

维度选择上，通常遵循直觉：数据是一维数组就用1D Grid，数据是图像矩阵就用2D Block，数据是体素（比如医疗CT）才用3D。不要为了炫技强行用高维，多维索引的换算也会带来额外开销：

维度限制是“硬约束”，而真正决定一个Block能否成功启动，还要看它吃的“资源”够不够——这就是资源约束。

blockDim.x×blockDim.y×blockDim.z≤1024\text{blockDim.x}\times\text{blockDim.y}\times\text{blockDim.z} \le 1024
这是Block线程数的“总闸门”，无论每个维多大，三者乘积必须控制在1024以内。
  1. 三个维度相乘Block的线程总数是三个方向线程数的乘积。
    blockDim.x×blockDim.y×blockDim.z\text{blockDim.x}\times\text{blockDim.y}\times\text{blockDim.z}
  2. 上限约束乘积不能超过硬件支持的1024。
    ≤1024\le 1024
  3. z维特殊限制即使前两个维度很小，z方向也最多64。
    blockDim.z≤64\text{blockDim.z} \le 64

资源约束

GPU的每个SM就像一块有限的工地：共享内存、寄存器、warp槽位都是稀缺资源。一个Block要占用一块共享内存和一批寄存器，能同时驻留多少个Block，取决于这些资源能被分成几份。

  • 共享内存限制：每个Block最多能用一定大小的共享内存（不同架构约48KB到227KB）。如果Block太大、每个线程都要大量共享内存，就会导致能驻留的Block数减少。
  • 寄存器限制：每个线程最多255个寄存器，SM上的寄存器总量固定。线程用寄存器越多，SM能同时容纳的线程就越少。
  • Occupancy（占用率）：SM上活跃warp数占最大warp数的比例。高占用率并不代表高性能，它只是“有足够多的工人排队”，让内存延迟被更多并行任务掩盖。

占用率公式：

资源约束决定了理论可行性，实战中真正重要的是把数据维度与线程维度对齐——最后一节看实操。

occupancy=active_warpsmax_warps_per_SM×100%\text{occupancy} = \frac{\text{active\_warps}}{\text{max\_warps\_per\_SM}} \times 100\%
占用率衡量SM有多“热闹”。100%意味着闲置很少，但过度追求占用率可能牺牲每个线程的计算效率。
  1. active_warps当前SM上正在运行（未阻塞）的warp数。
    active_warps\text{active\_warps}
  2. max_warps_per_SM该SM架构支持的最大warp数。
    max_warps_per_SM\text{max\_warps\_per\_SM}
  3. 比值乘100%得出活跃warp的占比，即占用率。
    active_warpsmax_warps_per_SM×100%\frac{\text{active\_warps}}{\text{max\_warps\_per\_SM}} \times 100\%

实战配置

假设你是一个AI算子工程师，要给一个长度 N=1000 的向量做逐元素加1。你不需要真的开1000个Block，那样太浪费，更合理的做法是：

1. 选一个常用Block大小，如 blockDim.x=256；

2. 计算需要多少Block：gridDim.x=ceil(1000/256)=4；

3. 在核函数里计算全局ID，并用 if (id<N) 保证越界的线程不干活。

这种“数组越界保护”是CUDA开发的标配，因为1000不能被256整除，最后一个Block只有232个线程真正参与计算。对应的映射公式是：

配置好启动参数后，要用NVIDIA Nsight等工具观察占用率和实际吞吐。好的配置不是公式硬套出来的，而是在资源约束与数据形状之间反复权衡的结果。就像分拣中心要根据包裹量动态调整班组人数，CUDA的线程配置也要跟着数据走。

row=flat_idwidth,col=flat_id mod width\text{row} = \frac{\text{flat\_id}}{\text{width}}, \qquad \text{col} = \text{flat\_id} \bmod \text{width}
处理二维数据时，把逻辑上的行列索引换算成“展平后的全局ID”，再反解出行列值。
  1. flat_id线程的全局一维编号，由 blockIdx 和 threadIdx 算出。
    flat_id=blockIdx.x×blockDim.x+threadIdx.x\text{flat\_id} = \text{blockIdx.x} \times \text{blockDim.x} + \text{threadIdx.x}
  2. 整除取行用全局ID除以宽度得到所在行。
    row=flat_idwidth\text{row} = \frac{\text{flat\_id}}{\text{width}}
  3. 取模得列用全局ID对宽度取余得到所在列。
    col=flat_id mod width\text{col} = \text{flat\_id} \bmod \text{width}

AI 应用场景

场景一：AI算子库的逐元素运算优化

阶段：模型训练与推理中的算子开发阶段

解决的问题：PyTorch/OneFlow等框架里的 Add、LayerNorm 等逐元素算子，如果只用Python循环处理百万级张量，速度慢到无法接受。CUDA用Thread/Block/Grid层次把张量切分成多块并行计算，让吞吐提升数百倍。

具体案例：为长度100万的FP16向量做 y = x + 1，按 blockDim=256、gridDim=ceil(1e6/256) 启动核函数。每个线程计算全局ID，在边界处做保护，避免越界。该模式是几乎所有AI算子优化的基础模板。

场景二：图像卷积的Tile分块

阶段：图像预处理或CNN卷积算子的高性能实现

解决的问题：传统图像卷积需要频繁读取相邻像素，若每个线程直接访问全局内存，带宽会成为瓶颈。Block+共享内存可以把图像切成Tile，每个Block加载一个Tile到共享内存，再边界填充后计算，大幅减少全局内存访问。

具体案例：处理 1024×1024 图像，用 16×16 的Block，每个Block负责一个16×16像素块。Block内256个线程先把Tile数据搬进共享内存，__syncthreads() 后每个线程计算一个输出像素。该模式也是Transformer中 attention 算子“分块softmax”的底层思想。

场景三：推荐系统Embedding批量查询

阶段：推荐模型推理的Embedding特征查询阶段

解决的问题：一次推理要查几十万个ID对应的Embedding向量，串行查询延迟极高。GPU可以把查询请求分配给不同Block，每个Block内部用共享内存合并访问，按数据维度灵活配置1D/2D Grid。

具体案例：某个性化推荐服务，batch=4096，每个样本查询7个Embedding。用 blockDim=256，gridDim=16，每个线程处理一个查询单元，配合CUDA Graph捕获减少启动开销，GPU利用率提升到85%以上。

Python 演示

python51 行
class CudaThreadSim:    def __init__(self, grid_dim=(1, 1, 1), block_dim=(1, 1, 1)):        self.grid_dim = grid_dim        self.block_dim = block_dim     @property    def threads_per_block(self):        bx, by, bz = self.block_dim        return bx * by * bz     def validate(self):        bx, by, bz = self.block_dim        gx, gy, gz = self.grid_dim        assert self.threads_per_block <= 1024, "block 线程数不能超过 1024"        assert bx <= 1024 and by <= 1024 and bz <= 64, "block 单维度超限"        assert gx < 2**31 - 1, "gridDim.x 超限"        assert gy <= 65535 and gz <= 65535, "gridDim.y/z 超限"        return True     def global_id(self, block_idx=(0, 0, 0), thread_idx=(0, 0, 0)):        gx, gy, gz = self.grid_dim        bx, by, bz = self.block_dim        tx, ty, tz = thread_idx        block_linear = (block_idx[2] * gy + block_idx[1]) * gx + block_idx[0]        thread_linear = (tz * by + ty) * bx + tx        return block_linear * self.threads_per_block + thread_linear  def launch_1d(n, block_dim=256):    grid_dim = (n + block_dim - 1) // block_dim    sim = CudaThreadSim(grid_dim=(grid_dim, 1, 1), block_dim=(block_dim, 1, 1))    sim.validate()    ids = []    for bx in range(grid_dim):        for tx in range(block_dim):            gid = sim.global_id((bx, 0, 0), (tx, 0, 0))            if gid < n:                ids.append(gid)    return grid_dim, ids  if __name__ == "__main__":    n = 1000    grid, ids = launch_1d(n, block_dim=256)    print("gridDim.x =", grid)    print("有效线程数 =", len(ids))    print("第一个线程ID =", ids[0], ", 最后一个 =", ids[-1])    assert ids == list(range(n)), "全局ID必须覆盖全部数据"    sim2 = CudaThreadSim(grid_dim=(2, 2, 1), block_dim=(32, 32, 1))    sim2.validate()    print("2D Grid 总线程数 =", 2 * 2 * 32 * 32)
目的：用Python模拟CUDA的一维和二维线程索引，验证 gridDim.x = ceil(N/blockDim) 的正确性。
  1. class CudaThreadSim用面向对象的方式模拟一次CUDA launch，grid_dim 和 block_dim 分别对应三个维度的块数和线程数。
  2. self.threads_per_block计算Block内线程总数，对应硬件要求 <=1024 的检查。
  3. validate()检查维度上限，相当于CUDA运行时的合法性校验。
  4. global_id()把三维block和thread索引展平成全局唯一ID，对应正文核心公式。
  5. block_linear把三元组 blockIdx 展平成Block的全局编号。
  6. thread_linear把三元组 threadIdx 展平成Block内编号。
  7. launch_1d(n, block_dim=256)典型的1D配置模板：数据除以Block大小求Grid。
  8. (n + block_dim - 1) // block_dim向上取整，避免遗漏最后一批数据。
  9. if gid < n边界保护，模拟核函数里的越界检查。
  10. assert ids == list(range(n))验证所有数据都恰好被一个线程处理。

常见误区

误区我以为 `Thread` 和CPU线程一样，每个线程独占一个核心，所以线程数越多GPU核心越多。

正解CUDA Thread是逻辑执行单元，数量可以远超物理核心。GPU的SM通过warp调度把所有线程分时/交错地塞进有限的执行单元。就像工厂里工人（逻辑线程）比工位（物理核心）多得多，大家轮班干活，你的“岗位”不代表你独享一台机器。

为何常见：因为CPU线程常被理解成“一个核跑一个线程”，把这种思维直接平移到了GPU。

误区反正GPU并行能力强，Block直接开满1024线程，性能一定最高。

正解Block过大会占用大量共享内存和寄存器，导致SM能驻留的Block数变少，占用率反而下降，甚至因资源不足无法启动。合理选择Block大小（64、128、256、512常见）能平衡负载与资源。

为何常见：只看到“线程多=并行度高”，忽略了硬件资源的物理上限和调度灵活性。

误区经常听说 `gridDim` 上限是65535，于是所有方向的Block数都不敢超过这个数。

正解只有 y 和 z 方向上限是65535，x 方向上限是 2^31-1（约21亿）。现代CUDA中一维大任务完全可以用一个很长的Grid，不需要拆成多个kernel。

为何常见：文档只看了一半，把y/z方向的限制当成了所有方向共用。

误区既然warp是32线程，那 `blockDim` 必须是32的倍数，否则程序会报错。

正解程序不会报错，硬件会把Block向上取整为若干个warp，但最后一个warp存在空转线程，造成约 (32 - 余数)/32 的浪费。所以“建议”而非“必须”是32的倍数。

为何常见：把性能建议误当成硬性语法约束。

误区把Grid和Block都设到最大，让所有SM塞满线程，跑出来的程序应该最快。

正解性能还受内存带宽、访存局部性、分支发散等因素影响。线程过多可能带来资源争抢和调度开销。正确的做法是根据数据规模选择“够用”的线程数，并把访存模式优化好。

为何常见：初学者容易把“并行度”和“计算量”混为一谈，忽略了GPU是吞吐型处理器，延迟需要靠足够并行来掩盖，但也不是越多越好。

面试题

Q1. CUDA线程层次从大到小正确排序是？

单选题 · easy

  • A.Thread → Block → Grid
  • B.Grid → Thread → Block
  • C.Grid → Block → Thread← 正确答案
  • D.Block → Grid → Thread

答案：C。Grid是最大层级，包含多个Block；每个Block又包含多个Thread。记住“工厂→车间→工人”的顺序即可。

Q2. 某Block的 blockDim.x=128，blockIdx.x=2，threadIdx.x=5，该线程的全局ID是？

单选题 · easy

  • A.261← 正确答案
  • B.135
  • C.256
  • D.250

答案：A。全局ID = 2×128+5=261。先算前面所有区块的线程总数，再加上块内偏移。

Q3. 一个Block最多可以包含多少线程？

单选题 · easy

  • A.32
  • B.256
  • C.1024← 正确答案
  • D.65535

答案：C。CUDA规定块内线程数上限为1024，且三维乘积不得超过该值。warp大小32，Block容量是它的32倍。

Q4. N=1000，blockDim.x=256，用1D Grid处理，需要多少个Block？

单选题 · mid

  • A.3
  • B.4← 正确答案
  • C.5
  • D.2

答案：B。ceil(1000/256)=4。前三个块分别处理256个数据，最后一块处理剩余232个，不足256但必须单独开一块。

Q5. 一个Block内有64个线程，它至少会占用几个warp？

单选题 · mid

  • A.1
  • B.2← 正确答案
  • C.3
  • D.4

答案：B。每32个线程组成一个warp，64线程恰好是2个满warp。如果线程数不是32的倍数，则向上取整，但会有部分线程空转。

Q6. gridDim.x 的最大合法值是？

单选题 · mid

  • A.65535
  • B.1024
  • C.2^31 - 1← 正确答案
  • D.无限制

答案：C。只有 y/z 方向上限是65535，x方向最大为 2^31-1。这是容易记混的考点。

Q7. 某核函数每个线程用16个寄存器，SM有65536个寄存器，每个Block有512线程，仅考虑寄存器限制，SM最多驻留几个Block？

单选题 · hard

  • A.4
  • B.8← 正确答案
  • C.16
  • D.32

答案：B。每Block寄存器用量=16×512=8192，65536÷8192=8。但实际还会受线程数等其他限制，这里只考察寄存器这一个维度。

Q8. 下列说法正确的有哪些？

多选题 · mid

  • A.同一Block的线程可用共享内存通信← 正确答案
  • B.不同Block可以直接同步
  • C.Warp是不可控的逻辑概念
  • D.共享内存比全局内存快← 正确答案

答案：AD。共享内存位于SM内部，速度远快于全局内存；不同Block不能直接同步，只能通过kernel边界或原子操作间接协作；warp是硬件调度单位，不由程序员直接控制，但B错误、C表述也不准确。

Q9. 下列哪些因素会影响SM上活跃Block的数量？

多选题 · mid

  • A.共享内存占用← 正确答案
  • B.寄存器使用量← 正确答案
  • C.全局内存大小
  • D.每Block线程数← 正确答案

答案：ABD。共享内存、寄存器、线程数都是决定Block驻留数量的资源约束；全局内存容量一般很大，不是SM内部的限制因素。

Q10. 关于Occupancy，下列理解正确的有哪些？

多选题 · hard

  • A.占用率越高，性能一定越好
  • B.占用率受共享内存与寄存器限制← 正确答案
  • C.适度降低占用率可以减少资源竞争← 正确答案
  • D.占用率只由Block大小决定

答案：BC。高占用率能掩盖延迟，但可能带来缓存/共享内存争用，所以不一定最好；占用率由多种资源共同决定，不只是Block大小。