想象你是一家超大规模“并行工厂”的车间主任。GPU 就是这座工厂，里面没有上百人的生产线，而是上万个工位同时开工。作为主任，你手里有一份超大订单，比如生产 10 万件 T 恤。你不会把订单直接砸给一个工人，而是把它分成大项目（Grid）、中项目（Block）、再分到一个个工人（Thread）手里。工人也不是单独行动的——32 个工人必须同乘一辆“班车”（Warp）进出工位，班车停，所有人同步开工；中途遇到岔路（if/else），班车只能先送一批人去剪裁、再送另一批人去熨烫，谁先谁后指令说了算。小组内部有一张公共工作台（共享内存），组员们把半成品放上去互相传递；组长吹一声哨（Barrier），所有人先放下手里的活，等组员都准备好了再一起进入下一道工序。工厂的“上座率”（Occupancy）则决定了工位是否被充分利用——班车明明能坐 32 人，却因为工具箱（寄存器）塞得太满只能坐 16 人，那另一半工位就白白空着。这篇内容，就是带着你把这个工厂的每个层级、每束调度、每个资源约束都走一遍。

岗位关联

掌握 GPU 线程模型，是写出高性能 CUDA 内核的必修课。无论是深度学习算子优化、图像处理还是科学计算，它决定你能否把 GPU 的上万核心真正“喂饱”。

目录
  1. 可视化演示
  2. 线程层级结构
  3. SIMT执行模型
  4. Thread执行单元
  5. Block线程块
  6. Grid网格
  7. Warp调度机制
  8. 占用率因素
  9. 硬件参数组织线程
  10. AI 应用场景
  11. Python 演示
  12. 常见误区
  13. 面试题

可视化演示

展示CUDA三大核心抽象：Thread-Block-Grid三级层次结构、32线程Warp的SIMT调度机制、以及不同Block大小下共享内存与寄存器的占用率变化。交互切换不同视角。

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

Thread-Block-Grid三级层次 Grid包含多个Block，Block包含多个Thread 中心Block展开显示4×4=16个Thread total = gridDim × blockDim GPU通过三级层次组织并行计算资源

线程层级结构

回到那座并行工厂。车间主任接了一个生产 10 万件 T 恤的大单，他不会把订单直接丢给某个工人，而是先把单子拆成“大项目”（Grid），每个大项目派给若干个“生产小组”（Block），小组里再分到一个个“工人”（Thread）。这一层套一层的嵌套，就是 CUDA 线程模型最核心的组织结构：一个 Grid 由多个 Block 组成，一个 Block 由多个 Thread 组成。

用 CUDA 的话说，Grid 和 Block 都可以是一维、二维或三维的索引空间。当我们启动内核时，写出 gridDim.x 表示网格在 x 方向上有几个 Block，blockDim.x 表示每个 Block 在 x 方向上有几个 Thread。这种嵌套关系决定了任务如何被逐级分解：

这个嵌套结构就像俄罗斯套娃——打开大项目是小组，打开小组是工人。理解了这层嵌套，后面的共享内存、同步、调度才有安放的位置。

Nthreads=gridDim.x×blockDim.x\displaystyle N_{\text{threads}} = \text{gridDim.x} \times \text{blockDim.x}
一个网格中所有线程的总数，等于网格里块的数量乘以每个块里线程的数量。
  1. gridDim.x网格在 x 维上包含的块数。就像订单里分出了几组生产小组，这个数告诉 GPU 总共有多少个 Block 等着被调度。
    gridDim.x=网格中的块数\displaystyle \text{gridDim.x} = \text{网格中的块数}
  2. blockDim.x每个块在 x 维上包含的线程数。一个生产小组里分配了几名工人，决定了每个 Block 能并行处理多少份数据。
    blockDim.x=每个块中的线程数\displaystyle \text{blockDim.x} = \text{每个块中的线程数}
  3. 两者相乘把“组数 × 每组人数”得到所有能同时参与生产的工人总数，也就是整个内核启动后一共会产生多少个线程。
    Nthreads=gridDim.x×blockDim.x\displaystyle N_{\text{threads}} = \text{gridDim.x} \times \text{blockDim.x}

SIMT执行模型

工厂里有一条规定：同乘一辆班车的 32 个工人，必须同步开工、同步收工、做一模一样的动作。这就是 GPU 的 SIMT（单指令多线程）执行模型。CPU 像一位博学的教授，一次只专心讲一章；GPU 更像教官，一声口令“向左转！”，几百个工位上的工人齐刷刷同时转身——一条指令驱动成百上千个线程同步执行。

不过，工人会遇到岔路。假设指令是“如果布料是红色的就裁剪，否则熨烫”，一个 warp 里有 16 个红色布料工人、16 个非红色布料工人，教官怎么办？他只能先喊“裁剪的去做”，16 人动；再喊“熨烫的去做”，另外 16 人动。两条路径被串行执行，时间翻倍。这就是 SIMT 里著名的“分支发散”（divergence）问题：

理解了 SIMT，你就明白了 GPU 高性能的秘诀和陷阱：它强在“同一动作复制万份”，弱在“各做各的”。写内核时，让线程尽量走同一条路，班车就能一趟跑完。

Twarp=tif+telse\displaystyle T_{\text{warp}} = t_{\text{if}} + t_{\text{else}}
当 Warp 内线程走向不同的分支时，总耗时等于各分支串行执行时间之和，而不是并行。
  1. t_if走 if 分支的那部分线程执行所需的时间。教官先带 16 个工人完成“裁剪”动作，耗时 t_if。
    tif=if 分支的串行执行时间\displaystyle t_{\text{if}} = \text{if 分支的串行执行时间}
  2. t_else走 else 分支的那部分线程执行所需的时间。完成后教官再带另外 16 个工人做“熨烫”，耗时 t_else。
    telse=else 分支的串行执行时间\displaystyle t_{\text{else}} = \text{else 分支的串行执行时间}
  3. 相加由于同一时刻 warp 只能执行一条指令，两个分支无法并行，只能先做完一支再做另一支，总耗时是二者之和。这就是为什么写 CUDA 代码时要尽量避免 warp 内的线程分道扬镳。
    Twarp=tif+telse\displaystyle T_{\text{warp}} = t_{\text{if}} + t_{\text{else}}

Thread执行单元

每个工人（Thread）是工厂里最小的执行单元。他有一个只属于自己的私人工具箱（寄存器），里面放着他正在处理的数据；他的工位上有独立的操作通道，可以自己取布料、自己放成品。但注意：工人的私人工具箱非常小——一个线程通常只有几十个 32 位寄存器可用，放不了多少东西。如果你让一个线程背上几千个私人物品，GPU 就不得不把多余的“行李”寄存到远处的仓库（本地内存），速度瞬间掉下来。

在 CUDA 中，每个线程有唯一的 threadIdx（在块内的编号）和 blockIdx（块在网格中的编号）。工人干活时，通过这两个编号知道自己是谁、该处理哪块布料。线程的私有属性还包括它消耗的寄存器数量，这个数字直接影响后面要讲的“上座率”：

线程是“最小但独立”的工人：有自己的工具、自己的编号、自己的活。但一个人干不成大项目，工人必须加入小组才能协作——这就引出 Block。

Rblock=Rthread×blockDim.x\displaystyle R_{\text{block}} = R_{\text{thread}} \times \text{blockDim.x}
一个线程块总共需要的寄存器个数，等于每个线程的寄存器用量乘以块内线程数。
  1. R_thread每个线程运行内核时申请的寄存器个数，由编译器根据代码复杂度决定。工人开工前要领取自己的工具数量。
    Rthread=单个线程的寄存器配额\displaystyle R_{\text{thread}} = \text{单个线程的寄存器配额}
  2. blockDim.x块内线程数。一个小组里有几名工人，每个人都要领一套工具。
    blockDim.x=线程块中的线程总数\displaystyle \text{blockDim.x} = \text{线程块中的线程总数}
  3. 相乘整个块的寄存器需求。GPU 每个 SM（流式多处理器）的寄存器总量是固定的——相当于车间每个工位区的工具柜总数固定，小组领走的工具多了，能同时开工的小组就少了。
    Rblock=Rthread×blockDim.x\displaystyle R_{\text{block}} = R_{\text{thread}} \times \text{blockDim.x}

Block线程块

工人不是孤军奋战的，他们以“生产小组”（Block）为单位坐在同一个工位区。小组内部有一张公共工作台（共享内存），所有组员都能把半成品放上去，也能拿别人放下的料。这张工作台比私人工具箱大得多，而且就在工位旁边，大家拿取极快——这正是 Block 内线程通信高效的根本原因。

但公共工作台有个规矩：组员不能各拿各的乱操作。假设 32 个工人都在往工作台上放零件，有人还没放完，有人就去拿了，那拿到的可能是残缺品。所以组长要吹哨（__syncthreads() 同步栅栏）：所有人必须把手里的零件都放上工作台、举手示意，组长确认全员到齐后吹哨，大家才能开始取料。这个机制保证了共享数据的一致性：

Block 是“协作的单元”：共享内存让组员高效配合，同步栅栏让大家配合得不出错。但一个小组最多 1024 人，一个订单需要上千个小组共同完成——这些小组由 Grid 统一编排。

Sblock=Sthread×blockDim.x\displaystyle S_{\text{block}} = S_{\text{thread}} \times \text{blockDim.x}
一个线程块需要的共享内存总量，等于每个线程申请的共享内存大小乘以线程数。
  1. S_thread每个线程在工作台上占用的空间，比如它往共享内存数组里写一个 float 就占 4 字节。工人放零件占的工作台面积。
    Sthread=每个线程申请的共享内存字节数\displaystyle S_{\text{thread}} = \text{每个线程申请的共享内存字节数}
  2. blockDim.x块内线程数。组里多少人，每人占一块地方，总面积自然按人数累加。
    blockDim.x=块内线程数量\displaystyle \text{blockDim.x} = \text{块内线程数量}
  3. 相乘得到整个块的共享内存申请量。SM 的工作台总面积有限，一个组占多了，能同时驻留的组就少了。
    Sblock=Sthread×blockDim.x\displaystyle S_{\text{block}} = S_{\text{thread}} \times \text{blockDim.x}

Grid网格

Grid（网格）是整个工厂的“总订单”。它把所有 Block 组织在一起，形成一个完整的任务空间。车间主任（硬件调度器）拿到 Grid 后，会把里面的 Block 一个个分配给不同的工位区（SM）执行。关键点来了：Block 与 Block 之间是彼此独立的——不同小组不需要知道对方在干什么，也不共享工作台，它们可以按任意顺序开工、完工，互不干扰。

这种“块间独立”是 GPU 能大规模并行的前提：如果所有小组都要互相通信，那整个 GPU 都会卡在等待上。正因独立，GPU 才能像流水线一样，哪个工位区空了就往里塞一个 Block。那么一个线程在全局唯一的编号如何计算？答案是双重编号：

Grid 是“全局的调度蓝图”：它定义了任务总量，而块间独立保证了 GPU 可以自由地动态调度。现在问题来了：调度的最小单位是什么？答案是 Warp。

tid=blockIdx.x×blockDim.x+threadIdx.x\displaystyle \text{tid} = \text{blockIdx.x} \times \text{blockDim.x} + \text{threadIdx.x}
线程的全局编号 = 块编号 × 每块线程数 + 块内线程编号，由此定位它负责哪份数据。
  1. blockIdx.x当前线程所在 Block 在整个 Grid 中的编号。相当于第几组工人，从 0 数起。
    blockIdx.x=块在网格中的序号\displaystyle \text{blockIdx.x} = \text{块在网格中的序号}
  2. blockDim.x每块线程数。每个小组的人数相同，前面几组一共占了多少个工人，用组号乘人数得到。
    blockDim.x=线程块的大小\displaystyle \text{blockDim.x} = \text{线程块的大小}
  3. threadIdx.x线程在块内的编号。自己在这组里的第几个位置。
    threadIdx.x=线程在块内的序号\displaystyle \text{threadIdx.x} = \text{线程在块内的序号}
  4. 相加组的偏移加上组内偏移，得到的 tid 就是这位工人在整个工厂里的工号。访问数组时，常用 tid 直接定位 data[tid]。
    tid=blockIdx.x×blockDim.x+threadIdx.x\displaystyle \text{tid} = \text{blockIdx.x} \times \text{blockDim.x} + \text{threadIdx.x}

Warp调度机制

接下来看工厂最精妙的环节——班车调度。GPU 在硬件层面并不是一个一个地调度线程，而是以 32 个线程为一束打包调度，这一束就是 Warp。32 名工人必须同乘一辆班车进出工位区，班车停靠时，32 人同时开始执行同一条指令；下一时刻，整车人同时执行下一条指令。硬件调度器就是那个排班员，它手里有很多班车，当一辆班车因等待数据而停下时，它立刻调度另一辆班车上工，用“车海战术”掩盖延迟——这就是 GPU 高吞吐的秘诀。

一个线程在 Warp 里的身份由它在线程块内的编号决定：前 32 个线程属于 Warp 0，第 33 到 64 个属于 Warp 1，以此类推。具体位置这样计算：

Warp 是 GPU 的真正执行单位：同一辆班车内的 32 人永远步调一致。这也意味着，如果你的 Block 线程数是 66，那么最后一个 Warp 只坐了 2 个人，剩下 30 个座位空着——资源被白白浪费。所以写代码时，Block 尺寸尽量凑成 32 的整数倍。

warp_id=threadIdx.x÷32,lane=threadIdx.x mod 32\displaystyle \text{warp\_id} = \text{threadIdx.x} \div 32, \qquad \text{lane} = \text{threadIdx.x} \bmod 32
线程所在的 Warp 编号 = 块内线程号整除 32，线程在 Warp 内的车道号 = 块内线程号对 32 取余。
  1. threadIdx.x ÷ 32整除运算。块内线程号每 32 个归入同一个 Warp，就像每 32 名工人坐同一辆班车，整除结果就是班车编号。
    warp_id=⌊threadIdx.x/32⌋\displaystyle \text{warp\_id} = \lfloor \text{threadIdx.x} / 32 \rfloor
  2. threadIdx.x mod 32取余运算。确定线程在这辆班车上坐的是第几个座位（0 到 31 号），这个座位号称为 lane。
    lane=threadIdx.x mod 32\displaystyle \text{lane} = \text{threadIdx.x} \bmod 32

占用率因素

工厂老板最关心的指标是“上座率”——Occupancy（占用率）。它衡量的是：每个工位区（SM）最多能同时容纳多少个 Warp，而当前你的内核实际用上了多少个。上座率 100% 意味着班车全部满座出发，GPU 的每个角落都在干活；上座率只有 50% 意味着有一半的班车停在车库里，工位空置。公式如下：

影响上座率的有三样“硬资源”：寄存器配额太大，每人领的工具太多，工具柜装不下那么多人；共享内存申请太多，工作台被占满；线程数超过 SM 上限，工位不够。想提高上座率，就要学会“精打细算”——少用寄存器、少占共享内存，让更多 Warp 挤上车。

Occupancy=Nactive_warpNmax_warp\displaystyle \text{Occupancy} = \frac{N_{\text{active\_warp}}}{N_{\text{max\_warp}}}
占用率 = 活跃 Warp 数 ÷ 每个 SM 能容纳的最大 Warp 数。
  1. N_active_warp你的内核实际能在 SM 上同时存活的 Warp 数。受寄存器、共享内存等资源约束，不是想开几辆班车就开几辆。
    Nactive_warp=实际可驻留的Warp数\displaystyle N_{\text{active\_warp}} = \text{实际可驻留的Warp数}
  2. N_max_warp硬件允许的 Warp 上限。不同计算能力的 GPU 有不同的上限，比如常见的 64。
    Nmax_warp=每SM最大Warp数\displaystyle N_{\text{max\_warp}} = \text{每SM最大Warp数}
  3. 比值两者相除得到占用率。它不是性能的绝对保证，但占用率太低时，GPU 一定“吃不饱”，延迟无法被掩盖。
    Occupancy=Nactive_warpNmax_warp\displaystyle \text{Occupancy} = \frac{N_{\text{active\_warp}}}{N_{\text{max\_warp}}}

硬件参数组织线程

既然上座率这么重要，那具体怎么配置 Grid 和 Block 的大小？答案藏在硬件的规格书里。GPU 的计算能力（Compute Capability）参数表写明了各种上限：每个 Block 最多 1024 个线程、每个 SM 最多 2048 个线程、Grid 每维最大 2³¹-1、共享内存每 SM 通常 48KB 到 228KB。这些数字就像工厂的建筑图纸：车间能放多少台机器、每台机器能站几个人，都是设计死的。

三维 Block 的维度也有限制，也就是说一个 Block 的 x、y、z 三个方向的线程数相乘不能超过总上限：

有了这些数字，配置策略就清晰了：先根据数据规模定 Grid 总线程数；再选 Block 大小为 128 或 256（32 的倍数，资源利用率高）；最后用 cudaOccupancyMaxPotentialBlockSize 这类 API 帮你算出理论最优值。工厂的每一寸工位都被充满，GPU 才真正物尽其用。

blockDim.x×blockDim.y×blockDim.z≤1024\displaystyle \text{blockDim.x} \times \text{blockDim.y} \times \text{blockDim.z} \le 1024
线程块三个维度的乘积不得超过 1024，这是硬件对单块规模的硬约束。
  1. blockDim.x / y / z块在三个方向上的线程数。就像小组在车间里的长、宽、高三维排布，三个方向人数相乘等于小组总人数。
    blockDim.x×blockDim.y×blockDim.z=块内总线程数\displaystyle \text{blockDim.x} \times \text{blockDim.y} \times \text{blockDim.z} = \text{块内总线程数}
  2. ≤ 1024硬件规定的单块人数上限。超过这个数内核直接启动失败。
    约束条件:≤1024\displaystyle \text{约束条件}: \le 1024

AI 应用场景

深度学习矩阵乘法的 Kernel 设计

阶段：模型训练与推理的算子优化层

解决的问题：矩阵乘法是神经网络最核心的算子。传统写法让每个线程独立从全局内存取数，访存延迟高、重复读多。利用线程模型重新编排后，把一个大矩阵切成多个子矩阵 Tile，每个 Block 负责一个 Tile，Block 内线程通过共享内存协作复用数据，性能可提升数倍甚至数十倍。

具体案例：在优化 LLM 训练中的 GEMM（通用矩阵乘法）时，工程师将 4096×4096 的矩阵划分为 16×16 的 Block 网格，每个 Block 内 256 个线程（8×8 排列），每个线程处理 4×4 的微块。计算时先把 Tile 载入共享内存，再用 __syncthreads() 同步后重复利用，访存次数从上千次降到几十次。

图像卷积的共享内存加速

阶段：CV 预处理与推理部署

解决的问题：卷积操作需要反复读取相邻像素，直接访问全局内存带宽不够。利用 Block 内共享内存，让每个 Block 负责输出图像的一个小块，把该块及其边界像素一次性载入共享内存，然后所有线程从共享内存取数做卷积。

具体案例：在 YOLO 推理的图像预处理中，把 1920×1080 的图像分成 15×8 个 Block，每块 32×32 线程，共享内存里放 34×34 的带 halo 数据块。这样每个像素只从全局内存读一次，卷积计算全程走共享内存，整体推理延迟降低约 40%。

大模型推理中 FlashAttention 的 Warp 协作

阶段：LLM 推理的注意力算子优化

解决的问题：Attention 计算中 QK^T 矩阵的每一行与一个 Warp 对应，Warp 内的 32 个线程分别负责不同的头维和序列分块，通过 Warp Shuffle 指令（同一 Warp 内线程交换数据）避免把中间矩阵写回全局内存，从而省下整个注意力矩阵的显存读写。

具体案例：在部署 70B 参数模型时，推理引擎用 FlashAttention-2 把 Q 矩阵按照 Warp 划分：一个 Warp 负责 64 行的输出，32 个线程通过 shuffle 在寄存器间传递部分和，最终只需一次全局写回。显存占用降低 75%，长序列生成速度提升 2 倍以上。

Python 演示

python56 行
import math class CUDAThread:    """模拟单个CUDA线程，持有全局编号与寄存器占用"""    def __init__(self, global_id, regs):        self.id = global_id        self.registers = regs  class BlockSim:    """模拟一个线程块：包含多个Thread，拥有共享内存"""    def __init__(self, block_id, block_dim, shared_bytes):        self.block_id = block_id        self.block_dim = block_dim        self.shared_bytes = shared_bytes        self.thread_list = []        for i in range(block_dim):            gid = block_id * block_dim + i            self.thread_list.append(CUDAThread(gid, regs=8))  class GridSim:    """模拟整个Grid：按Warp粒度调度线程执行"""    def __init__(self, grid_dim, block_dim, sm_max_warps=64):        self.grid_dim = grid_dim        self.block_dim = block_dim        self.block_list = []        for b in range(grid_dim):            self.block_list.append(BlockSim(b, block_dim, shared_bytes=1024))        total_threads = grid_dim * block_dim        self.warp_count = math.ceil(total_threads / 32)        self.max_warps = sm_max_warps     def run(self):        """按Warp推进：每次让32个线程同步执行一条指令"""        executed = 0        ticks = 0        while executed < self.grid_dim * self.block_dim:            executed += 32            ticks += 1        return ticks     def occupancy(self):        """计算占用率 = 活跃Warp数 / SM最大Warp数"""        active = min(self.warp_count, self.max_warps)        return active / self.max_warps  if __name__ == "__main__":    gpu = GridSim(grid_dim=16, block_dim=256)    total = gpu.grid_dim * gpu.block_dim    ticks = gpu.run()    print("总线程数:", total)    print("总Warp数:", gpu.warp_count)    print("模拟指令轮次:", ticks)    print("占用率: %.1f%%" % (gpu.occupancy() * 100))
目的：用 Python 模拟 CUDA 的 Grid-Block-Thread 三级结构，并以 Warp 为最小单位推进执行，直观展示线程总数、Warp 数量和占用率的计算方式。
  1. gid = block_id * block_dim + i每个线程的全局编号：块编号乘以块大小再加上块内偏移，对应正文里的 tid 映射公式。
  2. self.thread_list.append(...)把线程装进 Block 的列表中，模拟 Block 内包含 Thread 的层级结构。
  3. math.ceil(total_threads / 32)线程总数除以 32 向上取整得到 Warp 数：不足 32 人的尾巴也得算一辆班车，这就是“虚位 Warp”的来源。
  4. while executed < totalSIMT 主循环：每轮只推进 32 个线程，模拟一个 Warp 整束同步执行一条指令。
  5. active = min(self.warp_count, self.max_warps)占用率计算：活跃 Warp 数不能超过 SM 硬件上限，这个比值就是上座率。

常见误区

误区很多人以为 GPU 有上万核心，那就把 Grid 和 Block 开到最大，线程总数越多 GPU 越忙，性能自然越高。

正解线程多只是“名义上的并行度”，硬件真正执行的最小单位是 Warp。如果每个线程寄存器占用过高、共享内存申请过大，SM 能同时驻留的 Warp 反而变少，占用率下降。就像工厂订单写得再大，工位区只放得下两组人，多出来的订单只能排队。性能瓶颈往往不是线程不够，而是资源装不下。

为何常见：因为 CPU 编程习惯里“任务越多越忙”的直觉惯性，忽略了 GPU 的物理资源墙。

误区初学者常把 GPU 线程想象成 32 个独立的 CPU 核心，以为它们可以一个跑加法、一个跑乘法，互不干涉。

正解SIMT 模型下，一个 Warp 在同一时钟周期只能执行同一条指令。如果代码里有 if/else 分支，GPU 会先让满足条件的线程执行 if 分支、屏蔽其余线程，再反过来执行 else 分支——相当于班车分两趟跑，总耗时翻倍。写代码时应尽量让 Warp 内线程走同一路径，或通过重新组织数据避免分歧。

为何常见：因为“多线程”这个词容易让人联想到 CPU 的多核并行，忽略了 GPU 是“单指令广播”的架构。

误区有人以为共享内存和全局内存一样，写完立刻别人就能读到最新值，所以直接写共享内存数组再让别的线程读，结果拿到脏数据。

正解共享内存没有自动的一致性保证。一个线程写、另一个线程读，必须用 __syncthreads() 在中间做栅栏同步——所有线程到达这个点后才能继续。就像工作台上放零件，必须等所有人都放完、组长吹哨，大家才能取料。遗漏同步会导致数据竞争，而且这类 bug 时隐时现，极难排查。

为何常见：因为常规并发编程里“可见性”是语言运行时保证的，而 CUDA 把同步的职责完全交给了程序员。

误区既然 Block 内通信方便，那干脆把 Block 设成 1024 线程，能多协作就多协作。

正解Block 过大有两个问题：一是每个 SM 能容纳的 Block 数量有限，大块会让调度不够灵活，资源碎片化；二是共享内存和寄存器按块分配，块越大每个线程分到的份额越小，甚至导致内核启动失败。业界常用 128 或 256 的 Block，既能跑满 Warp，又给资源留有余地。工厂里不是组越大越好，而是要让每个工位区都能塞进合适数量的小组。

为何常见：因为“通信方便”听起来诱人，忽略了硬件对 Block 驻留数量的限制。

面试题

Q1. 在 CUDA 线程层次结构中，GPU 调度的最小执行单元是？

单选题 · easy

  • A.Thread 线程是执行的最小单元← 正确答案
  • B.Block 块是执行的最小单元
  • C.Grid 网格是执行的最小单元
  • D.Kernel 内核是执行的最小单元

答案：A。Thread 是真正执行计算的最小单元，Block 和 Grid 只是逻辑组织容器，内核则是整个程序。就像工厂里真正干活的是工人，小组和订单只是把人组织起来的方式。

Q2. 一个 Warp 中包含多少个线程？

单选题 · easy

  • A.8 个线程为一束共同调度
  • B.16 个线程为一束共同调度
  • C.32 个线程为一束共同调度← 正确答案
  • D.64 个线程为一束共同调度

答案：C。Warp 是硬件定义的调度单位，固定为 32 个线程。32 人同乘一辆班车，班车停大家一起停。其他数字都不是 CUDA 的 Warp 规格。

Q3. Block 内线程之间最常用的高效通信方式是？

单选题 · easy

  • A.通过全局内存传递中间数据
  • B.通过共享内存协作交换数据← 正确答案
  • C.通过网络接口远程交换数据
  • D.通过文件系统落盘后再读取

答案：B。共享内存在 SM 内部，延迟远低于全局内存，是 Block 内协作的首选。相当于小组的工作台在工位边上，比跨车间传递快得多。全局内存虽可用但慢，网络和文件系统根本不在 GPU 编程的考虑范围。

Q4. 全局线程编号计算式 tid = blockIdx.x * blockDim.x + threadIdx.x 中，blockDim.x 的含义是？

单选题 · mid

  • A.网格中块在 x 方向的个数
  • B.块中线程在 x 方向的数量← 正确答案
  • C.设备上流多处理器的总数
  • D.线程在 x 方向上的步长间隙

答案：B。blockDim 描述的是“块的大小”，即每个块里有多少线程。blockIdx 是第几块，threadIdx 是块内第几个。三者相乘相加得到全局编号，就像“组号 × 每组人数 + 组内序号”算出总工号。

Q5. SIMT 模型中，若一个 Warp 内一半线程走 if 分支、一半走 else 分支，执行结果是？

单选题 · mid

  • A.两条分支同时并行执行互不干扰
  • B.两条分支串行执行总耗时翻倍← 正确答案
  • C.只执行人数多的那个分支路径
  • D.整个内核抛出异常终止运行

答案：B。SIMT 规定同一 Warp 同时刻只能执行一条指令，所以 GPU 会先执行 if 路径、屏蔽另一半线程，再执行 else 路径。班车只能先送一批人再去接另一批，无法同时去两个地方。

Q6. 下列哪些因素会影响 GPU 的 Occupancy（占用率）？

多选题 · mid

  • A.每个线程使用的寄存器数量← 正确答案
  • B.每个块申请的共享内存大小← 正确答案
  • C.SM 最大可容纳的线程总数← 正确答案
  • D.主机端硬盘的连续读写速度

答案：ABC。占用率由 SM 内三大资源决定：寄存器总量、共享内存总量、线程驻留上限。硬盘读写与 GPU 内核执行无关，属于主机端 I/O，不参与 SM 资源配额。上座率能坐多少人，取决于工人工具箱、工作台和工位三个硬限制。

Q7. 现代 CUDA GPU 中，一个线程块最多允许有多少个线程？

单选题 · mid

  • A.最多允许 256 个线程
  • B.最多允许 512 个线程
  • C.最多允许 1024 个线程← 正确答案
  • D.最多允许 2048 个线程

答案：C。硬件规定 blockDim 三个维度乘积不超过 1024。这是工厂设计图纸上的工位区上限，超过则内核启动失败。256 和 512 是常用配置但不是上限，2048 超出了硬件约束。

Q8. 以下哪些情况会导致 Warp 执行效率明显下降？

多选题 · hard

  • A.Warp 内出现 if/else 分支并且两组都执行← 正确答案
  • B.内存访问时相邻线程读取相邻地址
  • C.线程等待共享内存数据产生束内停滞← 正确答案
  • D.Block 内线程数不是 32 的整数倍← 正确答案

答案：ACD。分支发散让 Warp 串行执行多条路径；等待共享内存会停住整辆班车；线程数不是 32 的倍数会造成最后一个 Warp 虚位。而相邻线程读相邻地址正是合并访问，是最高效的内存模式，属于加速而非降效。

Q9. 每个线程寄存器用量增大时，对 Occupancy 的影响是？

单选题 · hard

  • A.寄存器多能提升运行速度，占用率变高
  • B.每线程寄存器越多同时驻留的线程越少← 正确答案
  • C.寄存器数量与占用率之间完全无关
  • D.共享内存变大就能抵消寄存器的影响

答案：B。SM 的寄存器文件总量固定（如 65536 个），每个线程分到的寄存器越多，能同时容纳的线程总数就越少，占用率随之下降。工具箱占地方太多，工位区自然站不下那么多人。共享内存和寄存器是两套独立资源，无法互相抵消。

Q10. 关于 Grid 与 Block 的关系，下列说法正确的是？

多选题 · extreme

  • A.Block 间共享内存可直接同步协作
  • B.Block 间执行顺序不确定且相互独立← 正确答案
  • C.Grid 的维度决定了任务的总规模← 正确答案
  • D.Block 内协作依赖共享内存加同步← 正确答案

答案：BCD。共享内存只存在于 Block 内部，Block 之间无法直接用它通信，A 错。GPU 不保证 Block 的执行顺序，它们相互独立，这让调度器能动态填充空闲 SM，B 对。Grid 维度 × Block 维度 = 总线程数，C 对。Block 内线程用共享内存交换数据时必须搭配 __syncthreads() 同步，D 对。