CUDA 深度八股：grid / block / warp / SIMT 与性能（高频问答）

承接 gpu-cuda.md（概览）。本篇是面试硬核问答版：把「grid、block、thread 三级层次、warp、warp divergence、SIMT vs SIMD、occupancy、合并访问、bank conflict」讲到能扛住连续追问。

一、执行模型三级层次：grid → block → thread

• thread：最小执行单元。
• block（thread block）：一组线程，可协作——块内线程共享 shared memory，可用 __syncthreads() 做 barrier 同步。一个 block 只能驻留在一个 SM 上，绝不跨 SM。
• grid：一次 kernel launch 的所有 block 集合。block 之间不能通信、无法同步（除非拆成多个 kernel 或用 cooperative groups）。

内置索引：threadIdx / blockIdx / blockDim / gridDim（各有 .x/.y/.z），线程如何映射到数据由程序员决定。

为什么设计成三级？（高频追问）

1. 协作粒度：共享内存 + barrier 只在 block 内成立，block 是"能互相通信"的最小单位，所以同步/复用的边界划在 block。
2. 可扩展性：block 可被调度到任意 SM，同一份代码/同一个 grid 能在不同核心数的 GPU 上跑——硬件负责把 block 分派到空闲 SM，程序不感知具体卡。这是 CUDA 能在各代 GPU 通用的关键。
3. 调度灵活：以 block 为调度单位，SM 上能同时驻留多个 block（受资源限制），便于延迟隐藏。

约束：每个 block 线程数上限 1024；一个 SM 能驻留的 block 数、warp 数、寄存器/共享内存都有上限。

二、warp 与 SIMT

warp = 32 个线程，是 GPU 硬件调度与执行的基本单位：SM 以 warp 为单位取指、发射、执行。一个 block 被切成若干 warp（如 256 线程 = 8 个 warp）。

SIMT（Single Instruction, Multiple Thread）：一条指令同时驱动 32 个线程，每个线程有自己的寄存器和数据地址。程序员写的是"标量线程代码"，硬件在底层把 32 个线程打包成 warp 并行执行。

warp 内线程共享同一个 PC（程序计数器），因此同一时刻 warp 内线程执行同一条指令——这正是 warp 内分支发散代价的根源。

SIMT vs SIMD（高频对比题）

维度	SIMD（如 SSE/AVX/NEON）	SIMT（GPU）
抽象层次	指令级显式向量化，一条指令操作向量寄存器	线程级，每条线程有独立 PC/寄存器，硬件打包成 warp
编程方式	程序员/编译器显式用向量类型与 intrinsics	像写普通标量代码，由硬件/编译器隐式并行
数据宽度	固定（如 128/256/512 bit）	warp 大小固定 32，但每线程数据宽度灵活
分支处理	无"分支"概念，靠代码处理	warp 内分支会发散，需串行化
灵活性	向量必须整齐，处理不规则数据难	每线程地址可不同（gather/scatter 友好）
相同点	——	都是"单指令多数据"思想；都是延迟隐藏/吞吐优先

一句话给面试官：SIMD 是"一条指令吃一个向量"，SIMT 是"一条指令吃一整个 warp 的线程，且每个线程有自己的寄存器和地址"——SIMT 更通用（能处理不规则数据），代价是分支发散。

三、warp divergence（分支发散）——最常被深挖

现象：当一个 warp 内的线程走了不同的分支路径（如 if 条件在线程间不一致），硬件无法同时执行两条路径，只能串行执行各分支：先让一部分线程活动（其余用 mask 屏蔽），执行完再让另一部分活动。

代价：原本 1 次执行的代码，可能变成 N 条路径各执行一遍，warp 利用率下降，等价于"并行度被打折"。极端情况（32 条路径）退化到近似串行。

// 若 t 在 warp 内不一致 → 两分支被串行执行
if (threadIdx.x % 2 == 0) { /* 路径 A */ } else { /* 路径 B */ }

如何减少发散：

1. 让分支条件在 warp 内尽量一致（同一 warp 走同一条路）。
2. 用算术/查表替代分支（如用 min/max、select、预计算表）。
3. 重排数据，让做同类操作的线程聚在同一 warp。
4. 循环边界/长度对齐 warp，避免尾部分叉。
5. 实在无法避免时，评估"两段都执行 + 结果选择"是否比发散更划算。

注意：短小的分支可能被编译器用 predication（谓词化） 处理（两条路径都执行、按条件选结果），此时不一定产生真正的串行化，但会浪费执行槽。

四、内存层级与访存优化

层级（从上到下，越快越小）：寄存器 → shared memory/L1 → L2 → 全局内存(DRAM) → 主机内存(PCIe)；另有只读的常量内存(constant)、纹理内存(texture)。

合并访问 coalescing（必答）

• 相邻线程访问相邻地址 → 硬件把整个 warp 的访问合并成少量内存事务，带宽利用率高。
• 若线程跨步/乱序 → 每次访问单独成一个事务 → 有效带宽暴跌。
• 写法：索引 a[threadIdx.x]（连续），矩阵行主序时让 col = threadIdx.x 连续。

shared memory 与 bank conflict

• shared memory 分 32 个 bank，每 bank 4 字节（4B）。
• 同一 warp 内，若多个线程访问同一 bank 的不同地址 → bank conflict，访问被串行化。
• 若多线程访问同一 bank 的同一地址 → 广播，无冲突。
• 缓解：padding（如数组第二维 +1）、调整访问步长、改变数据排布。

shared memory 复用模式

需跨线程复用的数据先拷进 shared（软件管理的 cache），配 __syncthreads()；典型：矩阵分块(tiling)乘法。

五、Occupancy（占用率）

occupancy = 当前 SM 上驻留的 warp 数 / SM 支持的最大 warp 数。

• 受限因素：每线程寄存器数、每 block 共享内存、每 block 线程数、block 数上限。
• 作用：更多驻留 warp → 一个 warp stall（等内存）时可切换其它 warp 执行 → 更好地隐藏延迟。
• 并非越高越好：为提 occupancy 硬压寄存器可能导致 register spill（溢出到 local memory，反而更慢）。实践是在"不溢出"前提下找平衡。
• 用 __launch_bounds__ 或 -maxrregcount 控制寄存器。

六、同步与原子

• __syncthreads()：block 内 barrier，等所有线程到达再继续（常用于 shared 读写前后）。
• 原子操作：atomicAdd/atomicCAS/...，跨线程/跨 block 累加等。
• 没有跨 block 的 barrier；需要全局同步时要拆成两个 kernel（隐式全局 barrier），或用 cooperative groups。

七、性能优化顺序（背给面试官）

1. 先判定瓶颈：用 profiler（Nsight Compute / Nsight Systems / nvprof）看是 memory bound 还是 compute bound，以及 copy / kernel / 通信各占多少。
2. Memory bound（最常见）：合并访问、shared 复用、降精度（fp16/bf16/fp8）、减少冗余访存、向量化（float4）。
3. Compute bound：提高 occupancy、减少寄存器压力、减少发散、利用 FMA、切分指令。
4. 调 grid/block 尺寸使 SM 满载；减少 __syncthreads 与分支。
5. 用 stream 重叠 H2D/D2H 拷贝与计算。

八、高频追问速答

Q：为什么 block 不能跨 SM / 为什么没有 grid 级同步？ A：block 是调度与协作单元，划到某个 SM 上执行；跨 block 同步会破坏可扩展性（不同 block 何时被调度不可控），需要时用多 kernel 或 cooperative groups。

Q：grid/block 尺寸怎么选？ A：block 取 128/256（warp 的整数倍，避免尾部浪费），使 SM 尽量多驻留 block；grid 覆盖数据总量；最终用 profiler 微调。

Q：warp 大小为什么是 32？ A：硬件设计折中——寄存器堆/调度器宽度与利用率平衡的结果；写代码时按 32 对齐访问与分支最有利。

Q：SIMT 为什么能处理不规则数据而 SIMD 难？ A：SIMT 每线程有独立地址/寄存器，可 gather/scatter；SIMD 向量必须整齐。

Q：occupancy 低一定是坏事吗？ A：不一定。若 kernel 不靠访存延迟隐藏（计算密集、寄存器受限），低 occupancy 也能跑满算力；关键是看是否有效隐藏了延迟。

Q：怎么定位 bank conflict？ A：Nsight Compute 的 shared memory 相关指标（如 bank conflict 计数/波前数），结合访问模式分析。

九、自评一句

「我能讲清 grid/block/warp 三级模型与 SIMT 语义，知道 warp divergence 与合并访问是性能关键，懂 occupancy 的权衡和 shared memory bank conflict，能用 profiler 判断 memory/compute bound，并按降访存→提占用→重叠传输的顺序优化。」