本文为 GPGPU 编译器工程师面试通关指南 系列第二期（Day 2）。
建议在阅读“面试官参考解析”前，先独立思考每道题目的考察意图与答题脉络（每题思考 3~5 分钟）。

💡 核心考察维度：标量/向量双套硬件管线协同（UniformityAnalysis）、片上共享内存（Shared Memory / LDS）Bank 冲突与 Swizzling 消除、软流水（Software Pipelining）与异步拷贝（Async Copy）硬件演进、物理寄存器体冲突（Register Bank Conflict）与操作数收集器（Operand Collector）、非确定延迟指令的记分牌（Scoreboard）与 Waitcnt 编译下发。

📝 今日高频面试题清单

【Q1 · LLVM 核心分析与双执行单元】Divergence Analysis（发散分析）与 Uniformity（一致性）优化

现代 GPGPU 架构（如 AMD CDNA/RDNA 的 SALU/VALU + SGPR/VGPR，以及 NVIDIA Turing 之后的 Uniform Datapath + UR 寄存器）都引入了专门的标量/Uniform 硬件单元。

1. 请详细说明 LLVM 中的 UniformityAnalysis（或旧版的 DivergenceAnalysis）是如何判断一个变量是 Uniform 还是 Divergent 的？
2. 当存在发散的控制流（Divergent Branch）时，为什么一个由常量赋值的变量也可能变成 Divergent？
3. 将变量识别为 Uniform 对后端指令选择和寄存器分配有什么巨大收益？

【Q2 · GPGPU 体系结构与访存优化】Shared Memory (LDS) Bank Conflict 的产生机理与编译期消除

Shared Memory / LDS 是实现高性能算子（如 GEMM、Reduction）的核心，但如果访问模式不当触发 Bank Conflict，带宽会成倍下降。

1. 请从硬件底层解释什么是 $N$-way Bank Conflict？
2. 哪些情况下同一 Warp 内多个线程访问同一个 Bank 却不会发生冲突？
3. 在编译器（如 LLVM 后端或 MLIR/Triton）中，有哪些常见的自动优化手段（如 Padding、Swizzling）可以用来静态消除 Bank Conflict？

【Q3 · 编译优化与延迟隐藏】GPU 上的 Software Pipelining（软流水）与异步拷贝（Async Copy）

在算力密度极高的现代 GPU（含 Tensor Core / Matrix Core）上，仅靠硬件 Warp 切换（TLP）往往不足以完全掩盖显存（Global Memory）到寄存器/Shared Memory 的长延迟，编译器必须引入软件流水（Software Pipelining / Modulo Scheduling）。

1. 请解释双缓冲（Double Buffering）或多级流水在循环编译中是如何展开与重排读写依赖的？
2. 这会给寄存器生命周期（Live Range）带来什么负面影响？
3. NVIDIA Ampere+ 引入的 cp.async（或 Hopper TMA）从硬件和编译器角度是如何缓解这一痛点的？

【Q4 · LLVM 后端与底层硬件】Register Bank Conflict（寄存器体冲突）与 Operand Collector

除了 Shared Memory 有 Bank，GPU SM 内部庞大的物理寄存器文件（Register File，每个 SM 高达 256KB）为了节省面积和功耗，通常并不是全相联的多端口 SRAM，而是采用分 Bank 的单端口/少端口 SRAM 设计。

1. 什么是 Register Bank Conflict？硬件中的 Operand Collector（操作数收集器）是如何工作的？
2. LLVM 后端在 Register Allocation（寄存器分配）和 Post-RA Scheduling（分配后调度）阶段是如何规避寄存器 Bank 冲突的？

【Q5 · 计算机底层原理与指令发射】Scoreboarding（记分牌）、Waitcnt 插入与 Control Stalls

GPU 采用顺序发射（In-Order Issue）、乱序完成（Out-of-Order Completion）的管线设计。对比 NVIDIA 架构（PTX 虚拟 ISA vs SASS 物理 ISA 中的 Control Code / Stall Counts）和 AMD GCN/CDNA 架构（显式的 s_waitcnt 指令）：

1. 编译器在最后生成机器码（MC 层或 Post-RA 阶段）时，是如何处理非确定性长延迟指令（如 Global Load, Shared Load）的数据依赖的？
2. 如果编译器漏插了等待计数或算错了 Stall 周期，硬件会发生什么？

🎯 面试官深度解析与追问点

mindmap
  root((Day 2 核心考点图谱))
    标量与发散分析
      发散源与数据依赖传播
      控制依赖与 IPDOM PHI 发散
      SGPR与UR节省
      标量指令合并与广播
    Shared Memory 访存
      32-Bank 交错寻址
      Broadcast 与 Multicast
      Padding 步长错位
      XOR Swizzling 仿射重映射
    软流水与异步拷贝
      双缓冲跨迭代重叠
      在途数据 VGPR 膨胀
      cp.async 旁路直写 LDS
      TMA 硬件地址计算卸载
    寄存器 Bank 冲突
      单端口 SRAM 端口争用
      Operand Collector 轮询收集
      RA 阶数与着色启发式
      Post-RA 操作数交换与复用
    管线依赖与指令发射
      顺序发射与乱序回包
      AMD SIInsertWaitcnts
      NVIDIA SASS Control Codes
      RAW 危害与静默数据损坏

题 1 深度解析：Uniformity Analysis（一致性分析）与标量通道

1. 发散源（Sources of Divergence）的界定

LLVM 的 UniformityAnalysis 采用基于数据流图的前向标记与不动点迭代。在内核入口处：

• 天然发散源（Divergent Sources）：线程私有 ID（如 threadIdx.x、Lane ID）、依赖线程执行结果的原子操作返回值（如 atomicAdd）、特定硬件内建函数等。
• 天然一致源（Uniform Sources）：内核参数（Kernel Arguments）、网格/线程块维度（blockIdx.x、gridDim.x）、字面常量（Constant Literals）。

2. 数据依赖 vs 控制依赖（Temporal / Sync Divergence）

• 数据依赖（Data Dependency）：若指令 $I$ 的操作数中包含 Divergent 变量，则 $I$ 的计算结果也是 Divergent（除非属于特殊的 Warp 级规约原语，如 __all_sync 将结果统一广播）。
• 控制依赖（Sync Divergence，高频考点）：
  看下面这段极其经典的代码：

1
2
3
4
5
6
7

int x = 0;
if (threadIdx.x < 16) {
    x = 1;
}
// 汇聚点：x = phi(1, 0)

• 变量 x 在两个分支被赋的值都是常量（1 和 0 本身皆为 Uniform）。
• 但因为判断条件 threadIdx.x < 16 是发散的，导致 Warp 内部前 16 个线程走了 Then 块，后 16 个线程走了默认路径。
• 在紧邻后支配节点（IPDOM）处的 $\phi$ 节点（phi [1, %then], [0, %entry]）汇总时，不同线程在不同时间/空间获得了不同的值。
• 编译器处理逻辑：UniformityAnalysis 结合 Dominator Tree 与 Post-Dominator Tree 计算发散分支的影响区间（Influence Region）。凡是在该区间内被赋值并在区间外汇聚点使用的变量（即穿过汇聚边界的 $\phi$ 节点），均必须被强制标记为 Divergent。

flowchart TD
    Entry["Entry: x_0 = 0"] --> Cond{"threadIdx.x < 16 (Divergent)"}
    Cond -->|"True (Lane 0..15)"| ThenBlock["Then: x_1 = 1 (Const Uniform)"]
    Cond -->|"False (Lane 16..31)"| Join["IPDOM: x_2 = phi(x_1, x_0)"]
    ThenBlock --> Join
    Join --> DivergentNode["x_2 最终为 Divergent<br/>(跨发散控制流汇总)"]

    style DivergentNode fill:#ffeef0,stroke:#ff7875,stroke-width:2px;

3. 识别为 Uniform 的巨大编译收益

1. 寄存器分流（节省 VGPR）：
   • AMDGPU：Uniform 变量分配至 SGPR（Scalar GPR）。一个 Wavefront（32 或 64 线程）只需占用 1 个 32-bit SGPR 物理单元，极大减少 VGPR 消耗，避免触发 Occupancy 阶梯暴跌。
   • NVIDIA（Turing+）：分配至专属的 UR（Uniform Register），减轻主通用寄存器文件压力。
2. 利用标量执行单元（SALU / Uniform Datapath）：
   • 循环迭代变量累加、基地址偏移计算全部旁路到 SALU 执行，完全不挤占向量算术单元（VALU / Tensor Core）的发射槽。
3. 标量访存与单次广播（Scalar Load / Broadcast）：
   • 当指针被证明为 Uniform 时，直接生成标量加载（如 AMD s_buffer_load_dword 或 NV LDC 指令）。
   • 硬件只需发起 1 次 32 字节或 64 字节的标量 Cache 访问并向所有线程硬件广播（Broadcast），彻底绕过了复杂的向量地址生成和访存合并（Coalescing）单元。

题 2 深度解析：Shared Memory Bank Conflict 与 Swizzling

1. 硬件机理：$N$-way Bank Conflict

• 物理划分：以主流架构为例，Shared Memory 被均匀切分为 32 个 Banks，每个 Bank 单周期提供 4 Bytes（32-bit）的访问吞吐。
• 地址映射公式：
  $$\text{Bank ID} = \left(\frac{\text{Byte Address}}{4}\right) \pmod{32}$$
• 冲突产生：当同一 Warp 的 32 个线程在同一周期并发访存时，若有 $N$ 个线程访问了同一个 Bank 内的不同字地址（Row），硬件片上交叉开关（Crossbar）无法在一个周期内并发出数，只能强制将该请求拆解为 $N$ 个周期串行执行。此时发生 $N$-way Bank Conflict，有效带宽跌至原先的 $1/N$。

2. 无冲突特例：Broadcast 与 Multicast

• 若同一 Warp 内的多个线程访问的是同一个 Bank 内的完全相同的数据地址，硬件仲裁器不会串行化，而是激活片上广播（Broadcast）或多播（Multicast）网络，在单周期内将同一份数据分发给所有请求线程，耗时依然为 1 周期。

3. 编译期消除手段（Padding vs XOR Swizzling）

flowchart LR
    subgraph Padding方案
    P1["tile[32][32] 列步长=32"] -->|扩容| P2["tile[32][33] 列步长=33"]
    P2 -->|效果| P3["Bank 依次偏移，消除冲突<br/>(但破坏了 16 字节对齐)"]
    end

    subgraph Swizzling方案
    S1["保持 tile[32][32] 大小不变"] -->|逻辑映射| S2["Col' = Col ^ (Row % 8)"]
    S2 -->|效果| S3["打散列访问至各 Bank<br/>(且完美保持 128-bit 对齐)"]
    end

💡 Triton / MLIR 实践：在 Triton 的 TritonGPUToLLVM 或 MLIR 的 VectorToGPU 转换管线中，通常会在将数据从 Global Load 存入 Shared Memory 时直接下发 triton_gpu.swizzle 属性，通过一条 xor 指令完成地址混淆，使得后续从 Shared Memory 进行 Dot（矩阵乘）操作读取时天然零冲突。

题 3 深度解析：Software Pipelining（软流水）与异步拷贝

1. 软件流水与双缓冲（Double Buffering）循环重组

以典型的 GEMM 主循环为例，朴素执行流为：
$$\text{Load Global} \longrightarrow \text{Store Shared} \longrightarrow \text{Sync} \longrightarrow \text{Load Shared} \longrightarrow \text{Compute}$$
为了掩盖从 Global 显存到片上 SRAM 的数百个周期的高延迟，编译器采用软件流水（Software Pipelining / Modulo Scheduling）将跨迭代操作重叠：

sequenceDiagram
    participant Prologue as 序言 (Prologue)
    participant Loop as 稳态主循环 (Steady State)
    participant Epilogue as 尾声 (Epilogue)

    Note over Prologue: 发射 Load Global(0)<br/>预取第 0 块数据
    Note over Loop: 迭代 i: 边算边取<br/>- 发射 Load Global(i+1)<br/>- 基于第 i 块数据执行 Compute(i)
    Note over Epilogue: 处理最后一块数据 Compute(N-1)<br/>收尾不发射多余 Load

2. 负面影响：寄存器生命周期（Live Range）急剧膨胀

• 在途数据（In-flight Data）的占用：
  在传统硬件路径下，数据通路为：$\text{Global Memory} \xrightarrow{\text{LDG}} \mathbf{VGPR} \xrightarrow{\text{STS}} \text{Shared Memory}$。
• 当编译器将预取提前 1 级、2 级甚至 4 级时，这些尚未消费的中间数据必须全天候驻留在通用寄存器（VGPR）中。
• 连锁反应：Live Range 大幅横跨整个循环体，导致单线程寄存器配额剧增，极易触发 Occupancy 断崖式下跌 或直接引发恶性的 Register Spilling。

3. 硬件与编译演进：cp.async 与 TMA 的技术革命

1. NVIDIA Ampere (cp.async) / AMD Direct-to-LDS：
   • 硬件旁路（Bypass VGPR）：硬件在 L1/Crossbar 处开辟通路，数据从显存直通写入 Shared Memory，完全不占用任何 VGPR！
   • 非阻塞流水控制：编译器生成 cp.async.commit_group 与 cp.async.wait_group N，使得编译器能够轻松实现 3~5 级的极深软流水（Multi-stage Pipeline）而无需担心寄存器爆炸。
2. **NVIDIA Hopper (TMA - Tensor Memory Accelerator)**：
   • 地址计算卸载：传统 cp.async 仍需要线程执行指令来计算每次加载的多维张量地址；Hopper TMA 将多维 Tensor 切片计算直接交由专用硬件单元完成，进一步释放了大量标量与通用寄存器。

题 4 深度解析：Register Bank Conflict 与 Operand Collector

1. 产生根因：单端口 SRAM 端口争用

• GPU 的每个 SM 拥有高达 256KB 的寄存器文件。若做成多端口（Multi-port）SRAM，芯片面积和静态漏电将无法承受。
• 因此工业界通常采用 Banked 单端口/双端口 SRAM（如 4 个 Bank，按寄存器编号 $R_{\text{id}} \pmod 4$ 分布）。
• 考虑一条三操作数指令：FMA R0, R1, R2, R3：
  • 若分配器碰巧给源操作数分配了物理寄存器 R1, R5, R9（全在 Bank 1），单端口 Bank 在同一周期只能读出 1 个数据，硬件必须强制停顿（Stall）3 个周期才能把 3 个源操作数读齐。

2. 硬件机制：Operand Collector（操作数收集器）

• 位于发射单元（Instruction Dispatch）与执行单元（ALU）之间。
• 收集池（Collector Units）：接收指令后，将操作数读取请求分发到各个 Bank 的读队列。收集器跨多个周期分别从不同 Bank 读取数据，待所有源操作数就绪后才允许指令进入 ALU 执行管道。

flowchart TD
    Issue["指令发射 (例如: FMA R0, R4, R8)"] --> Collector["Operand Collector 缓冲池"]
    subgraph RegisterFile["分体寄存器文件 (Register File)"]
        B0["Bank 0 (R0, R4, R8...)"]
        B1["Bank 1 (R1, R5, R9...)"]
        B2["Bank 2 (R2, R6, R10...)"]
        B3["Bank 3 (R3, R7, R11...)"]
    end
    Collector -->|"周期 1: 读取 R0"| B0
    Collector -->|"周期 2: 读取 R4 (冲突等待)"| B0
    Collector -->|"周期 3: 读取 R8 (冲突等待)"| B0
    Collector -->|"操作数全部集齐"| ALU["进入 ALU 执行"]

    style B0 fill:#ffeef0,stroke:#ff7875,stroke-width:2px;

3. 编译器的规避策略

1. RegAlloc 阶段的 Bank 启发式着色：
   • 在 LLVM 的 RAGreedy 分配器中，目标后端通过定制 AllocationOrder：在为当前虚拟寄存器选择物理寄存器时，探测同条指令中已分配的兄弟操作数所在 Bank，惩罚同 Bank 物理寄存器的权重，优先选择落在不同 Bank 的物理编号。
2. Post-RA 指令级微调与操作数交换（Commuting）：
   • 利用算术对称律交换操作数位置（如将 ADD R0, R1, R2 交换为 ADD R0, R2, R1）。
   • 复用缓存（Operand Reuse Cache）利用：硬件通常维护一个极小的 Reuse Cache。若相邻两条指令在同一槽位复用同一个寄存器，硬件可直接命中缓存而无需访问 Bank。编译器通过重排指令或交换操作数槽位，最大化该缓存命中率。

题 5 深度解析：Scoreboarding、Waitcnt 与指令发射

1. 为什么 GPU 必须采用轻量化依赖追踪？

• CPU 采用庞大的重排序缓冲区（ROB）和寄存器重命名（Tomasulo 算法），硬件开销极大；
• GPU 核心数极多，硬件管线设计为 顺序发射（In-Order Issue）、乱序完成（Out-of-Order Completion）。
• 固定延迟指令（如 VALU 算术运算通常 4~8 周期）与非确定长延迟指令（如 Global Load 命中 30 周期、Miss 达 400 周期）混杂，必须在编译器与极简硬件记分牌之间达成协同。

2. AMD 架构：LLVM 编译期显式插入 s_waitcnt

• 硬件维护计数器：
  • vmcnt：在途向量显存读取（Vector Memory）的请求计数。
  • lgkmcnt：在途 LDS、GDS、常量（K）以及标量（M）访存计数。
• LLVM SIInsertWaitcnts Pass：
  • 该 Pass 位于 CodeGen 的尾声。它在基本块内维护数据流跟踪，记录每个寄存器的写入指令及其在途状态。
  • 当一条指令即将使用（Use）某个由 global_load 写入的寄存器时，编译器计算从该 Load 到当前指令之间还积压了多少条未返回的访存，精准插入等待指令：

1
2
3
4
5
6
7
8
9

global_load_dword v1, v0, s[0:1] ; 发起访存 1 (vmcnt 增加)
global_load_dword v2, v0, s[0:1] ; 发起访存 2 (vmcnt 增加)
; ... 执行其他不相关的算术指令 ...
s_waitcnt vmcnt(1)               ; 仅需等待访存 1 回包，允许访存 2 继续在途！
v_add_f32 v3, v1, 1.0            ; 安全消费 v1
s_waitcnt vmcnt(0)               ; 等待所有访存全部回包
v_mul_f32 v4, v2, 2.0            ; 安全消费 v2

3. NVIDIA 架构：SASS 物理指令控制码（Control Codes）与记分牌

• PTX 是虚拟指令集，不暴露硬件细节。在底层汇编（ptxas 生成的 SASS）中，每条 128 位指令捆绑了专门的 Control Codes（控制字）：
  • Stall Counts：静态写入当前指令发射后必须等待的固定时钟周期数（0~15）。
  • **Hardware Scoreboards (sb0 ~ sb5)**：
    • 当发射一条异步长延迟指令（如 LDG）时，编译器指定其置位某个记分牌的 Write Barrier（例如设置 sb0）。
    • 在后续第一条消费该结果的指令控制字中，设置对应的 Wait Barrier Mask（要求流水线在 sb0 清除前暂停发射该指令）。

4. 漏插或算错的灾难后果

• 在 CPU 上依赖硬件自动保证正确性，算错只会影响流水线效率；
• 在 GPU 底层物理 ISA 上，如果编译器算错了 Stall Count，或漏插了 s_waitcnt / Barrier Mask，硬件不会报异常，而是会直接拿寄存器中的旧数据（旧值/脏数据）进行后续计算！
• 这会引发极其致命且极难定位的 RAW（Read-After-Write）数据竞争与静默数据错误（Silent Data Corruption）。

💡 面试实战总结（复盘金句）

1. 谈发散分析时：*”常量赋值之所以会发散，本质是控制依赖在 IPDOM 处的空间/时间投影；LLVM 的 UniformityAnalysis 通过识别发散分支的影响区间，精准捕捉跨边界 PHI 节点，这是驱动 SGPR 和标量加载广播优化的核心基石。”*
2. 谈 Bank 冲突时：*”Memory Padding 是以空间换时间，但往往破坏了 128-bit 向量化对齐；而现代编译器首选的 XOR Swizzling 则是利用线性空间仿射变换，在保证极致对齐的前提下彻底打散地址投影。”*
3. 谈软流水时：*”传统多级流水的最大杀手是中转在途数据导致的 VGPR 爆炸；现代架构通过 cp.async / TMA 实现显存到片上共享内存的硬件直通，让编译器能够在零寄存器压力惩罚下构建极深的微流水。”*
4. 谈指令发射与依赖时：*”GPU 顺序发射的极简硬件哲学，把非确定延迟的依赖解析全权交给了编译器；从 AMD 的 SIInsertWaitcnts 到 NVIDIA 的 SASS 记分牌控制字，漏算一个周期就是静默数据损坏的灾难。”*

👈 上一期回顾：**Day 1：体系结构、控制流、LLVM 优化与内存模型**
👉 下一期推荐：**Day 3：算子融合、One-Shot Bufferize、寄存器溢出与集群同步**
👉 专栏返回：查看更多模块内容与更新计划，请访问 **GPGPU 编译器面试通关指南目录**。

本文为 GPGPU 编译器工程师面试通关指南 系列第一期（Day 1）。
建议在阅读“面试官参考解析”前，先独立思考每道题目的考察意图与答题脉络（每题思考 3~5 分钟）。

💡 核心考察维度：硬件资源约束（Occupancy & Spilling）、SIMT 执行模型与控制流结构化、LLVM 异构地址空间体系、中后端访存合并与向量化安全分析、细粒度 GPU 内存模型及调度约束。

📝 今日高频面试题清单

【Q1 · GPGPU 体系结构与编译优化】寄存器压力与 Occupancy（占用率）的权衡

在 GPGPU 架构中，SM（Streaming Multiprocessor）的寄存器文件总量是固定的。编译器在进行指令调度（Instruction Scheduling）和循环展开（Loop Unrolling）时，往往会增加单线程的寄存器使用量。请解释：

1. 单线程寄存器增加如何影响 Occupancy？
2. 编译器在 Backend 阶段通常如何权衡 Instruction-Level Parallelism (ILP) 与 Thread-Level Parallelism (TLP)？
3. 如果发生 Register Spilling，在 GPU 上它的底层代价具体是什么？

【Q2 · SIMT 控制流】Branch Divergence（分支发散）与重汇聚（Reconvergence）机制

GPGPU 采用 SIMT（单指令多线程）执行模型。当一个 Warp/Wavefront 内的线程遇到条件分支（如 if-else 或复杂的循环嵌套）发生分歧时：

1. 编译器是如何配合硬件生成谓词指令（Predication）或控制流指令来保证正确性的？
2. 请结合 LLVM IR 到目标 ISA 的转换，谈谈 Structurization（控制流结构化）的作用。
3. Volta 架构引入 Independent Thread Scheduling 后，对编译器分析与优化产生了什么深远影响？

【Q3 · LLVM 核心机制】Address Space（地址空间）在 GPGPU Backend 中的作用

不同于 CPU 的平坦内存模型（Flat Memory Model），GPGPU 拥有复杂的存储层次（Global, Shared/LDS, Local, Constant, Private）。

1. 在 LLVM IR 中是如何表征这些不同存储空间的？
2. 一个 Generic Pointer（通用指针）是如何通过 LLVM Pass 优化并降级（Lower）为特定地址空间指令的？
3. 如果编译器无法静态推导出地址空间，会产生什么底层性能损耗？

【Q4 · LLVM 向量化与访存优化】Memory Coalescing（访存合并）与 Vectorization

在 GPU 中，将标量访存合并为宽向量访存（例如将 4 个 ld.global.f32 合并为 ld.global.v4.f32）能显著提升带宽利用率。

1. LLVM 中的 LoadStoreVectorizer Pass 是如何工作的？
2. 它依赖哪些前置分析（Analysis Passes）来证明合并的安全性（如 Alias Analysis、Alignment Analysis 等）？

【Q5 · 计算机底层原理与并发】GPU Memory Model（内存模型）与 Fence 指令编译

现代 GPGPU 支持细粒度的内存一致性模型（如 PTX 中的 .cta, .gpu, .sys 作用域以及 acquire/release 语义）。

1. 从编译器角度看，一个带有 memory_order_acquire 的原子操作在 LLVM IR 中如何表示？
2. 编译器在指令选择（Instruction Selection）和指令调度（Instruction Scheduling）阶段，必须遵守哪些约束以防止破坏内存序？

🎯 面试官深度解析与追问点

mindmap
  root((GPGPU Compiler 面试考察点))
    资源与调度
      Occupancy 阶梯效应
      ILP vs TLP 权衡
      Local Memory 缓存抖动
    控制流
      SIMT Stack 与 IPDOM
      IfConversion 谓词化
      StructurizeCFG 规约
      Volta 独立线程调度影响
    内存体系与优化
      addrspace 属性与消除
      InferAddressSpaces 数据流推导
      LoadStoreVectorizer 向量化
      SCEV 与 AA 依赖检查
    并发与内存模型
      LLVM Atomic & syncscope
      SelectionDAG 链式依赖
      Cache Invalidation 指令下发

题 1 深度解析：寄存器压力、Occupancy 与 Spilling

1. Occupancy 影响与“阶梯效应”（Cliff Effect）

• 硬件约束原理：每个 SM 的物理寄存器总数是固定的（例如 NVIDIA SM 通常为 64K 个 32-bit 寄存器，AMD Compute Unit 则有 64KB/128KB VGPR 文件）。
• 量化关系：
  $$\text{Max Active Warps} = \min\left(\text{Max Warps per SM}, \left\lfloor \frac{\text{SM Total Registers}}{\text{Registers per Thread} \times \text{Threads per Warp}} \right\rfloor\right)$$
• 阶梯效应（Cliff Effect）：硬件在分配寄存器时存在颗粒度限制（例如以 256/512 个寄存器或一定数量的 Register Block 为步进）。单线程仅仅多使用 1 个寄存器，就可能触发分配块跨级，导致 SM 能同时容纳的 Active Warps 数量瞬间断崖式下跌（例如从 32 个 Warp 骤降到 16 个），造成 Occupancy 减半。

2. ILP vs TLP 的编译器权衡策略

• ILP 激进调度的代价：Pre-RA Scheduler（寄存器分配前调度）为了掩盖长延迟算术或访存指令，往往采用激进的列表调度（List Scheduling）将独立指令拉开。但这会大幅拉长变量的 Live Range（生命周期），导致并发活跃的虚拟寄存器数量激增。
• 编译器的应对机制：
  • 压力感知调度（Pressure-Aware Scheduling）：如 LLVM 的 GCNIterativeScheduler 或 NVPTX 相关策略，实时跟踪当前的 Register Pressure（RP）。当 RP 接近设定的阈值或 __launch_bounds__ 上限时，调度算法从“最大化 ILP”退化为“最小化生命周期的 DFS/Topological 优先调度”。
  • 重计算（Rematerialization）：如果某个变量的生成代价极低（如常数载入、基地址偏移计算、小跨度位移），编译器会在其被使用的地方重新生成一份指令，而不是为其保留一个物理寄存器。以微小的算力为代价，换取关键路径上的寄存器释放。

3. Register Spilling 的真正代价

• 本质真相：在 CPU 上，Spill 是写入栈帧，由于 CPU 拥有超大的私有 L1/L2 Cache，多数溢出仅命中片上缓存。而在 GPU 上，Local Memory 在物理上位于 Global Memory（外部 DRAM/HBM）中。
• 灾难性后果：
  1. 吞吐断崖与高延迟：尽管 Local Memory 会经过 L1/L2 Cache 过滤，但海量并发线程同时发生 Spill 会瞬间冲垮 L1/L2，引发严重的 Cache Thrashing（缓存抖动）。
  2. 数百周期延迟暴露：一旦未命中片上缓存，单次 Spill 涉及 200~800 个时钟周期的访存等待，完全摧毁吞吐。
     (注：部分现代架构支持将未用尽的片上 Shared Memory/LDS 作为 Spill 溢出区，以缓解访存延迟灾难。)

题 2 深度解析：SIMT 分支发散与控制流结构化

1. SIMT 硬件执行模型与重汇聚（Reconvergence）

• Active Mask / EXEC 掩码：同一 Warp 内的 32/64 个线程共享同一个 PC。当执行条件分支时，硬件通过掩码标记当前处于活动状态的线程。
• SIMT Stack 机制：硬件维护分支执行栈。遇到 Divergence 时：
  1. 压入 False 分支的目标 PC、掩码以及重汇聚点（Reconvergence Point）。
  2. 掩码切换为 True 线程集合，继续向下执行。
  3. 执行完 True 分支后，弹出栈顶恢复 False 路径。
  4. 到达重汇聚点时，恢复分支前的完整掩码。
• IPDOM 原则：最优重汇聚点在理论上通常是控制流图中的 Immediate Post-Dominator（紧邻后支配节点）。如果重汇聚点选错或下移，会导致线程无谓地串行化执行更长时间。

flowchart TD
    Cond{"Condition (Divergent)"}
    ThenBlock["Then Block (Exec Mask: Active True)"]
    ElseBlock["Else Block (Exec Mask: Active False)"]
    IPDOM["Immediate Post-Dominator (Reconverge)"]

    Cond -->|True Paths| ThenBlock
    Cond -->|False Paths| ElseBlock
    ThenBlock --> IPDOM
    ElseBlock --> IPDOM

2. Predication vs Jump（IfConversion）

• 短分支谓词化：对于代码行数极少、无副作用的分支（如 y = (x > 0) ? a : b），编译器通过 IfConversion Pass 将控制依赖转化为数据依赖，生成带谓词的指令（如 PTX @p0 mov.b32，AMDGPU v_cndmask_b32），完全消除硬件分支跳转和压栈开销。

3. LLVM StructurizeCFG 的核心使命

• 硬件要求：GPU 硬件在处理底层掩码保存/恢复（如 AMDGPU 的 s_and_saveexec_b64 / s_or_b64 exec）时，严格依赖结构化（Reducible）的循环与分支拓扑。
• 处理不可约流图（Irreducible CFG）：若源码中包含复杂的 goto 或跨循环层级的跳转，CFG 会呈现不可约形态。LLVM 的 StructurizeCFG Pass 会识别这些结构，显式插入 Flow 块、引入状态变量（State Variable）与循环重整，将复杂 CFG 规整为单一出口/入口的标准循环与分支，保证后端能够正确插入 EXEC 掩码操作指令。

4. Volta Independent Thread Scheduling (ITS) 的冲击

• 架构变革：Volta 及后续架构为每个线程配备了独立的 PC 和 Call Stack，允许 Warp 内线程以细粒度交错执行，甚至支持 Warp 内线程之间的锁与交互。
• 编译器影响：
  • Warp-Synchronous 隐式同步失效：以前开发者认为同一 Warp 在同一个基本块内天然锁步执行，编译器也可以放心地跨指令做公共子表达式消除（CSE）。在 ITS 下，隐式锁步假设彻底打破，必须依赖显式屏障指令（如 PTX bar.warp.sync 或 __syncwarp()）。
  • 内存序与死锁分析：编译器在移动含有同步/自旋依赖的代码时，必须严格遵守屏障范围，严禁将代码外提（Hoist）越过同步点。

题 3 深度解析：Address Space 与指针优化

1. LLVM IR 中的地址空间表示

在 LLVM 中，不同硬件存储体由指针类型的 addrspace(N) 属性区分（即便是 LLVM 15+ 的 Opaque Pointers 时代，类型依然写作 ptr addrspace(N)）。
常见映射关系对照表：

2. 为什么必须消除 Generic Pointer？

• 指令差异显著：
  • 访问明确的 Shared 内存：指令为 ld.shared（NVPTX）或 ds_read（AMDGPU）。地址只需 32 位，硬件直接走片内互联总线，消耗寄存器少、发射延迟极低。
  • 访问 Generic 内存：指令为通用的 ld 或 flat_load。由于地址可能指向 Shared 也可能指向 Global，硬件必须使用 64 位宽地址，且硬件执行流水线必须在运行时提取高位 Address Bits 进行动态分支判断与路由，不仅多占一倍地址寄存器，还会引入额外的译码延迟与能耗。

3. InferAddressSpaces Pass 原理

• 推导机制：
  1. 定位 Source：寻找具有确定地址空间的源头指针（例如全局变量 @shared_buf 天然携带 addrspace(3)，内核入参指针通过属性标记为 addrspace(1)）。
  2. 数据流图传播：沿着 Def-Use 链向后传播（遍历 getelementptr, bitcast, phi, select 等操作）。
  3. Rewrite & Eliminate：将中间不必要的 addrspacecast (ptr addrspace(3) to ptr) 逆向消除，并将终端的 load ptr 重写为 load ptr addrspace(3)。
• 推导失败的惩罚：如果因为复杂跨模块函数调用、动态指针别名导致推导中断，编译器只能在后端生成保底的 Flat/Generic Load，直接带来 15%~30% 的寻址性能开销。

题 4 深度解析：访存合并（Memory Coalescing）与向量化

1. LoadStoreVectorizer 的工作原理

• 目标：将连续的标量 Load/Store（如 4 个连续的 32-bit float 读写）提升为宽向量访存（如 128-bit 的 float4 对应 ld.global.v4.f32 或 global_load_dwordx4）。
• 优势：
  • 减少发往 MIO/LSU（加载存储单元）的指令条数，降低指令发射队列拥塞。
  • 充分发挥 GPU 内存总线单次请求的突发带宽（Burst Transfer）。

2. 核心前置依赖分析（Safety & Legality）

flowchart LR
    BB["扫描 BasicBlock 指令"] --> SCEV{"SCEV 分析<br/>连续常量偏移?"}
    SCEV -->|Yes| AA{"Alias Analysis<br/>中间有冲突 Store?"}
    AA -->|No Alias| Align{"Alignment 分析<br/>满足 16 字节对齐?"}
    Align -->|Safe| Vectorize["合并为 128-bit 向量访存"]
    SCEV -->|No| Abort["放弃合并"]
    AA -->|May/Must Alias| Abort
    Align -->|Misaligned| Abort

1. **ScalarEvolution (SCEV)**：
   • 判定两个指针表达式的差值：$\text{Ptr}_2 - \text{Ptr}_1 = \text{Constant Offset}$。
   • 必须确保偏移量刚好等于上一个访存类型的字节大小（如 $+4, +8, +12$）。
2. AliasAnalysis (AA) 与 MemorySSA 依赖检验：
   • 假设有序列：L1 = Load ptr; Store ptr_x; L2 = Load ptr+4;。
   • 要将 L1 与 L2 合并，必须将 L2 提升移动到 Store 之前（或者将 L1 下沉）。
   • 编译器必须通过 AA->isNoAlias(ptr+4, ptr_x) 证明该移动绝对不会越过任何存在别名依赖的写操作，否则会导致读写乱序，违背 RAW（Read-After-Write）规则。
3. Alignment（对齐分析）：
   • 硬件向量指令（如 128-bit 向量访存）通常强制要求自然对齐（16-byte alignment）。
   • 编译器必须通过 getOrEnforceKnownAlignment、KnownBits 或用户假设（__builtin_assume_aligned）确信基地址在运行时满足 16 字节对齐。如果基址对齐不足，强行合并会导致运行时直接抛出硬件非法访存异常（Hardware Misaligned Exception）。

题 5 深度解析：GPU 内存模型与 Fence 编译

1. LLVM IR 内存模型规范

LLVM 原生支持多层次并发语义模型：

• Memory Ordering：monotonic, acquire, release, acq_rel, seq_cst。
• **Synchronization Scope (syncscope)**：
  • "workgroup"（OpenCL）/ "cta"（CUDA）：作用域仅限于同一个线程块。
  • "agent" / "gpu"：作用域涵盖整颗 GPU 设备内部所有核心。
  • "system"：作用域跨越 PCIe/CXL 总线，与 Host CPU 保持统一视图。

示例 LLVM IR 片段：

1
2
3
4
5

; 线程块内 Release 语义的原子写
store atomic i32 %val, ptr addrspace(1) %ptr release, syncscope("workgroup"), align 4

; GPU 全局 Acquire 语义的原子读
%res = load atomic i32, ptr addrspace(1) %ptr acquire, syncscope("gpu"), align 4

2. 指令选择（ISel）与 Cache 控制指令下发

GPU 的各级 Cache 架构（如 L1/Texture Cache）默认通常是弱一致性甚至是只读/非全相联硬件失效的。不同于 CPU 自动的 MESI 硬件缓存一致性嗅探总线，GPU 强依赖编译器下发显式指令维护缓存可见性：

• Acquire 编译：
  • 目标是将全局最新数据拉入本地。
  • 编译器在下发读取指令时，不仅生成带 .acquire 修饰符的硬件操作，对于 AMDGPU 还需设置指令属性位（如 glc 全局连贯位、slc 系统级连贯位），甚至在必要时插入 buffer_wbinvl1_vol（无效化本地 L1 Cache，强制穿透到 L2/全局拉取最新值）并插入 s_waitcnt vmcnt(0) 等待访存回包。
• Release 编译：
  • 目标是保证之前的写操作对其他线程全部可见。
  • 编译器会先插入写屏障指令（Drain Pending Writes），确保所有数据由本地 Store Buffer 刷入 L2 或全局内存，再发射释放锁的 Store。

3. 指令调度器（Instruction Scheduler）的约束

• SelectionDAG / MachineScheduler 规则：
  • 带有 Acquire 属性的节点，在 DAG 中会被赋予类似 MayLoad 以及特殊的 Order Dependency / Chain 依赖边。
  • 调度器在拓扑排序和计算就绪列表（Ready Queue）时，严格禁止把依赖图位于 Acquire 之后的所有指令提升到 Acquire 之前（单向屏障 One-way Barrier）。
  • 同样，对于 Release 指令，严格禁止把其前面的读写指令下沉重排到其后，从而在编译期确保了语言规范级内存序不被乱序优化破坏。

💡 面试实战总结（复盘金句）

1. 谈 Occupancy 时：*”不要一味追求 100% 的理论占用率；在具备足够 ILP 能够隐藏访存延迟的前提下，适度降低 Occupancy 换取更大的单线程寄存器配额，往往能获得更好的真实吞吐。”*
2. 谈控制流时：*”GPU 控制流优化的终极目标是减少 Active Mask 切换与串行化开销；短分支用 IfConversion 数据依赖替代控制依赖，复杂分支必须经由 StructurizeCFG 规整以保证掩码栈操作的正确性。”*
3. 谈指针优化时：*”Generic Pointer 是 GPGPU 性能隐形杀手；中端必须穷尽 InferAddressSpaces 的数据流推导，消除动态地址路由。”*
4. 谈向量化时：*”LoadStoreVectorizer 并不是简单的指令打包，它本质是 SCEV 跨步推导、AA 别名无依赖证明与硬件 Strict Alignment 的三位一体。”*

👉 下一期推荐：**Day 2：标量优化、Bank Conflict、软流水与指令发射**
👉 专栏返回：查看更多模块内容与更新计划，请访问 **GPGPU 编译器面试通关指南目录**。