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

💡 核心考察维度：MLIR 渐进式 Lowering 与 GPU 目标硬件属性保留（Address Space / CallingConv / Intrinsics）、One-Shot Bufferization 别名分析与静态/动态 Shared Memory 降级、SCF/Affine 循环到 GPU Grid/Block 硬件拓扑映射与 Trip Count 裁决、GreedyPatternRewriter 匹配收敛机制与 Tensor Core 原语 Fallback 降级、现代 GPU 编译器中 MLIR 与传统 LLVM 后端的职责边界划分与交接架构。

📝 今日高频面试题清单

【Q1 · 多级中间表示与属性透传】MLIR Dialect Lowering 与 GPU 目标属性保留

在将高层 MLIR（如 gpu、nvvm 或 rocdl Dialect）逐步降级（Lowering）至 llvm Dialect 的过程中，编译器如何确保目标硬件相关的关键属性不丢失？请具体说明：

1. 地址空间（Address Spaces）：memref 的 Memory Space 属性如何准确映射到 LLVM IR 的 Pointer Address Space（如 CUDA 的 Shared/Global Memory 或 ROCm 的 Local/Global Memory）？
2. Kernel 调用约定与属性：gpu.func 在转换到 llvm.func 时，其作为入口函数的标记（gpu.kernel）、Block/Thread 尺寸约束以及硬件专属属性（如 CUDA 的 maxntid，ROCm 的 amdgpu_flat_work_group_size）是如何透传与附加的？
3. 线程内置函数（Builtins）：gpu.thread_id 或 gpu.block_id 等抽象内置操作如何转译为 nvvm / rocdl 的专有 Intrinsic，并在最终降级至 llvm Dialect 时保留正确的硬件寄存器映射？

【Q2 · 内存抽象与物理存储映射】Tensor 到 Memref 的 Bufferization 与 GPU 共享内存映射

在基于 MLIR 的 GPU 编译器（如 IREE 或 MLIR-TensorFlow）中，程序通常先在 Tensor 语义下进行高层优化，随后通过 Bufferization 转化为 Memref 语义。请回答：

1. In-place Bufferization 原理：MLIR 的 One-Shot Bufferize 如何利用读写分析（One-Shot Analysis）与生命周期分析（Liveness Analysis）来实现内存就地复用（In-place Reuse）？如何规避 GPU 上的 Read-After-Write (RAW) 冲突？
2. Dynamic/Static Shared Memory 分配：
   • 静态 Shared Memory 如何在 Bufferization 阶段映射为 memref.alloc() 并最终绑定到 GPU 的 Local Data Share (LDS) / Workgroup 内存？
   • 动态 Shared Memory（尺寸在运行时决定）如何通过 gpu.dynamic_shared_memory Dialect 操作建模？在 Lowering 到 LLVM IR 时，它是如何转化为外部动态共享内存声明指针（如 PTX 中的 extern __shared__）的？

【Q3 · 并行拓扑与循环映射】SCF / Affine 循环到 GPU Grid/Block 的映射

MLIR 中通常使用 scf.for 或 affine.for 表示多层嵌套循环。GPU 编译器的核心 Pass 之一是将这些结构化循环映射到 GPU 的 Hardware Thread Hierarchy（Grid 与 Block）。

1. Mapping Pass 的算法机制：以 affine.parallel 或 scf.parallel 到 gpu.launch 的映射为例，编译器如何决定哪些循环层级映射到 blockIdx（x, y, z），哪些映射到 threadIdx（x, y, z）？如何处理 Loop Trip Count 与 Thread Hierarchy 维度不匹配（如 Loop 迭代次数大于 Thread Block 尺寸）的情况？
2. 对下游循环展开（Unroll & Loop Transformations）的影响：
   • 循环被映射到 GPU 线程拓扑后，原本连续的循环控制流发生了什么变化？
   • 映射后的代码如何影响后端（如 LLVM NVPTX 后端）的 Vectorizer 和 SIMT 谓词化（Predication）优化？

【Q4 · 图模式匹配与指令特化】MLIR Pattern Rewriting 与硬件原语对接及 Fallback 机制

MLIR 依赖 PatternRewriter 体系实现 Dialect 之间的转换与优化（例如将高层的矩阵乘法操作匹配并转换为 NVIDIA PTX 的 wmma / mma.sync，或 AMD 的 mfma 指令）。

1. 匹配与收敛机制：GreedyPatternRewriteDriver 如何维护 Worklist？当多个 Pattern 发生匹配冲突或重叠时，编译器如何确定优先级？如何确保 Rewrite 过程不会陷入无限死循环？
2. 硬件原语匹配与 Fallback 机制：
   • 以 GEMM 为例，当高层 linalg.matmul 或 vector.contract 的 Shape/数据类型（如 i4 或不规则 Shape 13x17x5）与目标硬件矩阵原语（如 CUDA mma.sync.16x8x16）不符合时，Pattern Rewrite 如何触发降级/回退（Fallback）路径？
   • 详细描述从硬件专有原语匹配失败，到生成通用标量/向量循环（FMA）的完整降级逻辑。

【Q5 · 编译器系统架构设计】混合编译架构下 MLIR 与传统 LLVM 后端的职责分工

现代 GPGPU 编译器（如 TensorFlow XLA、Triton、Mojo 等）普遍采用了”MLIR 高层 IR + 传统 LLVM（NVPTX/AMDGPU）后端”的混合编译架构。

1. 底层优化的不可替代性：为什么诸如寄存器分配（Register Allocation）、指令调度（Instruction Scheduling）、软件流水线（Software Pipelining）以及硬件 Hazard 消除等优化不适合在 MLIR 高层 Dialect 中完成？
2. 理想交接点（Handoff Point）：
   • MLIR 与传统 LLVM 后端之间最合理的交接点在哪里？
   • 在进入 LLVM 后端之前，MLIR 必须完成哪些工作？LLVM 后端又接管了哪些工作？

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

mindmap
  root((Day 8 核心考点图谱))
    MLIR 渐进降级与属性透传
      memref 内存空间映射 addrspace
      gpu.kernel 与专属 CallingConv
      nvvm / rocdl 内置 Intrinsic 转译
      DataLayout 与 Target 属性绑定
    Bufferization 与共享内存
      One-Shot Analysis 别名与 RaW 检测
      In-place 原地复用与零拷贝
      静态 Shared Memory 映射全局变量
      gpu.dynamic_shared_memory 外部指针
    循环映射与并行拓扑
      SCF / Affine 映射 Grid / Block
      Grid-Stride Loop 与边界裁决
      高维嵌套循环剥离为微内核
      SIMT 谓词化与向量展开基础
    Pattern Rewriting 与硬件对接
      GreedyPatternRewriteDriver 队列收敛
      Benefit 优先级与死锁预防
      Tensor Core MMA 模式匹配
      Outerproduct 到标量 FMA 的 Fallback
    混合架构职责分工
      MLIR 擅长领域级图变换与内存布局
      LLVM 后端专精硬件 ISA 资源分配
      llvm Dialect 理想交接边界
      两级编译器协同演进的工业实践

题 1 深度解析：MLIR Dialect Lowering 与 GPU 目标属性保留

1. 地址空间映射机制

在 MLIR 中，memref 类型包含一个可选的 memorySpace 属性（如 Attribute 或 IntegerAttr）。在向 llvm Dialect 转换时，LLVMTypeConverter 机制会调用专用的 Address Space Mapping 接口：

• 类型转换：MemRefType 转换为 LLVM::LLVMPointerType。如果 memref 的 memory space 是一个整数或代表特定硬件存储区域的 Dialect 属性（例如 #gpu.address_space<workgroup>），类型转换器会通过硬件 target 提供的 DataLayout 或 TargetAttr 将其映射为具体的数值。
• 硬件映射示例：
  • NVPTX：global $\rightarrow$ addrspace(1)，shared $\rightarrow$ addrspace(3)，local $\rightarrow$ addrspace(5)。
  • AMDGPU：global $\rightarrow$ addrspace(1)，region/lds $\rightarrow$ addrspace(3)，private $\rightarrow$ addrspace(5)。
• 指令重写：读写该 memref 的 memref.load/store 会转换为 llvm.load/store，指针类型严格保留 addrspace(N) 标记，从而强制 LLVM PTX/AMDGPU 后端在 Codegen 时生成对应的硬件访存指令（如 PTX 的 ld.shared 或 ld.global）。

2. Kernel 调用约定与属性透传

gpu.func 到 llvm.func 的转换由 ConvertGPUToLLVMPatterns 驱动：

• Calling Convention 转换：如果 gpu.func 带有 gpu.kernel 单元属性（UnitAttr），转换器在生成 llvm.func 时，会为其赋予特定的调用约定。例如在 NVPTX 后端中设为 LLVM::cconv::C 并添加 nvvm.kernel 属性；在 AMDGPU 中则赋予 LLVM::cconv::AMDGPU_KERNEL (CallingConv 85)。
• 属性翻译：
  • CUDA/NVPTX：gpu.known_block_size 属性会被重写为 nvvm.maxntidX / Y / Z 属性节点。进入 LLVM IR 后生成 !nvvm.annotations 命名元数据。
  • ROCm/AMDGPU：转换为 rocdl.max_flat_work_group_size 节点，最终转化为 LLVM IR 中的 attributes #0 = { "amdgpu-flat-work-group-size"="1,1024" }。

3. 线程内置函数转译

高层 gpu.thread_id x / gpu.block_dim x 等操作在 Lowering 时遵循两阶段映射：

• 阶段一（Dialect 到 Hardware Dialect）：gpu.thread_id {x} 被 Pattern Matching 重写为 nvvm.read.ptx.sreg.tid.x（NVPTX）或 rocdl.workitem.id.x（ROCm）。
• 阶段二（Hardware Dialect 到 LLVM IR Dialect）：这些 Dialect 中的 Op 实现了 LLVM::PointerBitcastOp 或 Intrinsics 接口，被翻译为 llvm.call @llvm.nvvm.read.ptx.sreg.tid.x() 或 llvm.call @llvm.amdgcn.workitem.id.x()。这些是 LLVM 内置的 Intrinsic，能够在后端寄存器分配阶段直接映射到硬件物理寄存器（如 PTX 的 %tid.x 或 AMDGCN 的 v0 寄存器）。

题 2 深度解析：Tensor 到 Memref 的 Bufferization 与 GPU 共享内存映射

1. One-Shot Bufferization 与就地复用

• SSA 破坏与分析：Tensor 是不可变（Immutable）的 SSA 值，而 Memref 是可变（Mutable）的内存缓冲区。Bufferization 的核心任务是将 Tensor 操作转换为对 Memref 的读写。
• One-Shot Analysis（Alias & RaW 分析）：
  • 编译器自底向上/自顶向下遍历 AST，为每个 Tensor SSA 值构建 Write-To-Memory 图。
  • 冲突检测：如果两个 Tensor 操作绑定到同一内存区域，且存在一个读操作在写操作之后发生（甚至存在跨 Loop 的 RaW 依赖），One-Shot 分析会将该节点标记为 “Must Allocate”。
  • 就地复用（In-Place）：若无 RaW 冲突（即旧 Tensor 值在写操作后不再被任何后续节点读取），编译器直接让输出 Tensor 复用输入 Tensor 的 Memref 内存块，避免分配新的显存。
• 规避 GPU 上的 RAW 冲突：在 GPU 多线程环境下，若跨线程/跨 Block 复用同一个 Shared Memory 区域，除了依赖 SSA 分析外，Bufferization 阶段还会引入 gpu.barrier（或 memref.copy 后插 barrier）。若检测到跨 Thread 读写交织，无法原地安全消去，则强制插入 memref.alloc() 申请额外 Buffer。

2. Shared Memory 分配与 LLVM IR 映射

静态 Shared Memory 映射

• 当高层代码声明一个固定 shape 的 Tensor 并指定存放在 Shared Memory 时（如带有 #gpu.address_space<workgroup> 标记），Bufferization 将其转换为：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96

%0 = memref.alloc() : memref<128x64xf32, #gpu.address_space<workgroup>>
```  
* 降级至 `llvm` Dialect 时，该 `memref.alloc` 不会被转换为 `malloc` 或 `cudaMalloc` 调用，而是直接转换为 **Global Variable**，其 Storage Class 设置为 `InternalLinkage`，Address Space 设为 3（Shared/Workgroup）。在 PTX 中生成 `.shared .align 4 .b8 alloc_0[32768]`。

#### **动态 Shared Memory 映射**

* 当 Shared Memory 的大小依赖运行时参数（如 Block 维度）时，无法在静态全局作用域分配。  
* MLIR 使用 `gpu.dynamic_shared_memory` 操作：

```mlir
%shmem = gpu.dynamic_shared_memory : memref<?xi8, #gpu.address_space<workgroup>>
```  
* **Lowering 到 LLVM IR 的转换流程**：  
  1. `gpu.dynamic_shared_memory` 被重写为一个指向外部符号的 `llvm.mlir.addressof` 操作。  
  2. 编译器在 LLVM Module 头部隐式插入一个外部动态共享内存全局变量声明：

```llvm
@dynamic_shmem = external addrspace(3) global [0 x i8], align 16
```  
  3. 通过 `llvm.getelementptr` 获取该全局变量的起始地址，并将其 `bitcast` 转换为实际所需的数据类型指针（如 `float*`）。  
  4. 最终在 PTX 中输出为 `extern .shared .align 16 .b8 dynamic_shmem[];`，其真实物理大小在 Kernel 发射（`cudaLaunchKernel` 的第 4 个参数 `dynamicSharedMemSize`）时由 Runtime 动态传入。

---

### 题 3 深度解析：SCF / Affine 循环到 GPU Grid/Block 的映射

#### 1. 循环映射机制与 Trip Count 匹配算法

#### **映射判定机制（Mapping Strategy）**

* MLIR 的 `ConvertAffineToGPU` 或 `SCFToGPU` Pass 采用属性标记（Attribute-driven）或启发式分析来进行映射。  
* 程序员或自动调度器（如 Polyhedral Scheduler）会给 `scf.parallel` 绑定映射属性，例如 `{mapping = [#gpu.loop_mapping<processor = block_x>, #gpu.loop_mapping<processor = thread_x>]}`。  
* 降级过程：  
  1. 识别顶层的 `parallel` 循环，将其归纳变量（Induction Variables）重新计算为全局/局部线程索引。  
  2. 创建 `gpu.launch blocks(...), threads(...)` 包装块。  
  3. 循环变量 $i$ 被重写为：$i = \\text{lower_bound} + (\\text{blockIdx} \\times \\text{blockDim} + \\text{threadIdx}) \\times \\text{step}$。

#### **Trip Count 与 HW 维度不匹配的处理**

当循环迭代次数 $N$ 大于硬件线程配置（如 $N=1024$，但 `blockDim.x=256`）时，编译器采用两种模式：

| 模式 | 机制 | 生成的代码结构 |
| :---- | :---- | :---- |
| **Grid-Stride Loop (网格步长循环)** | 将超出硬件维度的循环保留为 Kernel 内部的 `scf.for`，跨度为 $GridDim \\times BlockDim$。 | `for (int i = tid; i < N; i += stride)` |
| **Thread Tile + Bound Guard (边界裁决)** | 将 Block 尺寸设为固定值，发射足够多的 Blocks，并在 Kernel 内部插入谓词分支。 | `if (global_id < N) { ... }` |

#### 2. 对下游 Loop 展开与后端 Codegen 的影响

#### **控制流变化**

* 映射前：显式的高层嵌套循环（`scf.for` / `affine.for`），具有完备的依赖关系图与归纳变量。  
* 映射后：原始的循环结构被**拆解/剥离（Strip-mined）**。部分循环轴完全消失，转化为隐式的硬件多线程并行；内核内部仅留下余数的跨段循环（Grid-Stride Loop）或纯粹的单线程顺序执行块（Basic Blocks）。

#### **对下游 LLVM 后端优化的影响**

1. **打破连续循环结构，依赖 SIMT 谓词化（Predication）**：  
   * 原本平滑的循环被插入了 `if (threadIdx.x < boundary)` 的条件分支。  
   * 下游 NVPTX/AMDGPU 后端必须利用 **SIMT Predication** 将这些分支重构为硬件 Mask（如 PTX 中的 `@p0 bra` 或 Execution Mask `EXEC`），消除 Warp Divergence。  
2. **构建向量化与 Tensor Core 化的基础**：  
   * 映射消除了高层的大循环后，留存在 Thread 内部的微小循环（如 $4 \\times 4$ 的寄存器 Tile 循环）变得极易被完全展开（Unroll All）。  
   * 展开后的连续 `memref.load` 操作可以被匹配为向量加载（`vector.load` $\\rightarrow$ PTX `ld.global.v4`），为后端的 Register Allocation 和 Instruction Scheduling 创造了高度定量的指令序列。

---

### 题 4 深度解析：MLIR Pattern Rewriting 与硬件原语对接

#### 1. Pattern Rewriting 匹配收敛机制

* **Worklist 驱动架构**：`GreedyPatternRewriteDriver` 内部维护一个 FIFO/LIFO 的 Worklist，初始状态下将指定区域（Region）内的所有 Operation 入队。  
* **匹配与应用（Match and Rewrite）**：  
  * 驱动器不断从 Worklist 弹出 Op，遍历注册到该 Op 的 `RewritePattern` 集合。  
  * **优先级判定（Benefit）**：每个 Pattern 声明了一个 `PatternBenefit`（无符号整数）。驱动器优先匹配 `Benefit` 值更高的 Pattern。如果 `Benefit` 相同，则按 Pattern 注册的先后顺序尝试。  
  * **状态更新**：一旦某个 Pattern 匹配成功并调用了 `rewrite()`，被修改/新生成的 Op，以及其受影响的邻居 Op（Parent、Users、Operands），会被重新推入 Worklist。  
* **收敛与死锁预防**：  
  * **Iteration Limit**：驱动器设定了最大迭代次数上限（如 `max_iterations = 10`）。  
  * **DAG 严格单调递减/改动检测**：PatternRewriter 强制要求通过 `rewriter.replaceOp` 或 `rewriter.eraseOp` 修改 IR。如果一个 Rewrite 动作没有改变 IR 结构，或者 Pattern 只是不断创建等价的 Op，驱动器通过比较匹配历史或达到 Iteration Limit 终止，抛出编译失败。

#### 2. 硬件原语匹配失败（Fallback）路径机制

针对矩阵乘法（GEMM）下发至 GPU 硬件 Tensor Core 原语的匹配，MLIR 采用分层的降级策略：

```text
               [linalg.matmul / vector.contract]
                               |
                   (Shape & Type Check Pass)
                               |
             +-----------------+-----------------+
             | Match Success                     | Match Fail
             v                                   v
[vector.mma.sync / vector.mfma]     [vector.outerproduct / Unroll]
             |                                   |
             v                                   v
[LLVM NVVM/ROCDL Intrinsics]        [scf.for + Elementwise FMA]
             |                                   |
             v                                   v
(PTX: mma.sync.aligned...)          (PTX: fma.rn.f32...)

1. 匹配条件检查（Pattern Matching）

VectorToGPU 或 LinalgToNVVM Pattern 会检查 vector.contract 的维度与类型：

• Shape 是否匹配：如 $M=16, N=8, K=16$。
• Type 是否支持：如输入为 f16，累加器为 f32。

2. 补齐与切块（Padding & Tiling - 预防 Fallback）

如果 Shape 是 13x17x5，在进入最终 Pattern 之前，高层 Pass（如 LinalgTilingAndPadddingPass）会先尝试进行 Padding（填充 0 至 16x32x16）并切片。若 Padding 失败，才真正进入 Fallback。

3. 硬件原语匹配失败（Fallback 降级）

若 Shape 或类型无法满足硬件 Tensor Core 要求（例如 target 不支持 i4，或者 Tile 尺寸无法对齐）：

• **Step A (Vector Dialect 拆解)**：匹配 vector.contract 到 nvvm.mma 的 High-Benefit Pattern 宣布失败 (failure())。驱动器回退至低 Benefit 的通用 Pattern，将 vector.contract 拆解为外积操作（vector.outerproduct）。
• **Step B (Vector 到 Loop 降级)**：vector.outerproduct 被进一步 Lower 到 scf.for 循环，并在最内层包含向量化/标量的 math.fma 或 arith.mulf + arith.addf。
• **Step C (LLVM Target 展开)**：scf.for 生成标准的控制流，fma 被 Lower 为 llvm.fmuladd 或标量 IR 指令。
• 后端结果：最终在 NVPTX 后端输出标准的 CUDA Core 标量/向量 FMA 指令（如 fma.rn.f32 %f3, %f1, %f2, %f4;），保证了即使缺少硬件 Accelerator 原语支持，程序依然具备语义正确性的底层兜底实现。

题 5 深度解析：混合编译架构下 MLIR 与传统 LLVM 后端的职责分工

1. 为什么底层优化不适合在 MLIR 高层完成？

虽然 MLIR 具备极强的抽象与表达能力，但将某些微架构相关的优化放在 MLIR 中实现是不高效且极其危险的：

• 缺少物理硬件状态建模：
  • 寄存器分配：依赖于具体的物理寄存器文件（Register File）数量、寄存器重命名机制、溢出（Spill）代价以及 Call Convention。MLIR 的 SSA 形式拥有无限的虚拟寄存器/SSA 值，在 MLIR 层无法精确评估 Register Pressure。若在 MLIR 层强行分配，后续 LLVM IR 到 Target Machine 的 Lowering 可能会因为引入新的临时变量而破坏分配结果。
  • 指令调度与 HW Hazard：需要精确的硬件 Execute Pipeline 延迟表（Latency Table）、Issue Width、乱序/顺序执行特性以及 Register Read-After-Write (RAW) Stall 周期（例如 AMD GCN 的 s_waitcnt 或 NVIDIA GPU 的 Instruction Dependency Barriers）。这些属于芯片内部微架构细节，频繁变更且极具私有性，不适合抽象为通用的 MLIR Pass。

2. MLIR 与 LLVM 后端的职责划分与交接点

双方理想的交接点位于 llvm Dialect（转换为标准 LLVM IR 文本或 Bitcode 的临界点）：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23

+-------------------------------------------------------------------+
|                        MLIR Framework                             |
|                                                                   |
|   [High-Level Dialects: Tosa / Linalg / Vector / GPU / SCF]      |
|                               |                                   |
|   - Tiling, Fusing, Vectorization                                 |
|   - Memory Layout Optimization & Bufferization                    |
|   - Explicit GPU Thread/Block Mapping                             |
|                               v                                   |
|                        [LLVM Dialect]                             |
+-------------------------------+-----------------------------------+
                                | (TranslateToLLVMIR)
================================v===================================
                                | (LLVM IR / Bitcode)
+-------------------------------+-----------------------------------+
|                     LLVM Backend (NVPTX / AMDGPU)                 |
|                                                                   |
|   - Target-Specific LLVM IR Passes                                |
|   - Instruction Selection (SelectionDAG / GlobalISel)            |
|   - Register Allocation & Spilling                                |
|   - Machine Instruction Scheduling & Hazard Recognition           |
|   - Assembly / Binary Assembly Emission (PTX / SASS / HSACO)      |
+-------------------------------------------------------------------+

MLIR 必须完成的工作（高层与中层）

1. 结构化控制流平铺（Domain-Specific Transformations）：完成 GEMM/Conv 的 Tile、Fuse（融合）、Unroll 以及 Auto-Tuning。
2. 显式内存层级管理：显式完成从 Global Memory 到 Shared Memory/LDS 的 Copy 操作插入与 Double Buffering（异步拷贝 async.copy）调度。
3. 并行拓扑映射：将抽象逻辑并行映射为明确的 GPU Grid/Block/Thread 维度以及 Thread Group (Warp) 级别的离散操作。
4. 降级为单层表达：将所有的抽象类型（tensor、memref、vector）完全展开为扁平化的指针、数值标量及 llvm.struct，输出符合 C/LLVM ABI 的 llvm Dialect。

LLVM 后端接管的工作（底层与硬件 Target）

1. LLVM IR 标量优化：运行 LLVM 标准的 InstCombine、LICM、GVN 以及 SimplifyCFG 等，消除 MLIR 降级过程中产生的冗余 Load/Store 或 Pointer 计算。
2. 指令选择（Instruction Selection）：将 llvm.call @llvm.nvvm... 或标准 LLVM Ops 转换为 Target Machine Instruction (MachineInstr/MC)。
3. 寄存器分配与溢出优化（Register Allocation & Spill）：在真实物理寄存器约束下分配 Register，并尽可能通过 Liveness Analysis 减少显存 Spill。
4. 机器指令调度与汇编发射（Machine Scheduling & Emission）：
   • 针对特定 GPU 架构进行 Pipeline 填塞（如 PTX 级别或 SASS 级别的指令重排，隐藏访存延迟）。
   • 最终输出目标平台专有代码（NVIDIA PTX/SASS 或 AMDGPU HSACO 镜像）。

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

1. 谈 MLIR 渐进式降级时：*”从高层张量到硬件二进制，MLIR 的多层 Dialect 在每一步都必须保留硬件元数据；通过类型转换器与属性透传，使高层表达的并行语义在底层精准沉降为硬件地址空间与调用约定。”*
2. 谈 Bufferization 共享内存映射时：*”One-Shot Bufferization 打破了函数边界的内存分配阻碍；静态共享内存沉降为内部全局变量，动态共享内存映射为外部符号，统一在编译期锁定了片上 SRAM 的物理拓扑。”*
3. 谈循环变换与线程映射时：*”SCF 循环映射是将软件嵌套迭代投影至物理线程网格的枢纽；剥离外层并行轴、内层留存微小循环，为下游 LLVM 后端的向量化与寄存器完全展开奠定了坚实基础。”*
4. 谈模式重写与 Fallback 时：*”工业级编译器不仅要有最先进硬件原语的高收益匹配，更要有严谨完备的降级兜底；当 Tensor Core 匹配受阻时，优雅回退至通用标量/向量 FMA 循环是保证语义正确性的基石。”*
5. 谈混合编译架构分工时：*”MLIR 专精于宏观结构优化、张量切块与内存层次划分，LLVM 后端则深耕于微架构寄存器分配、流水线调度与物理指令发射；明确两者在 llvm Dialect 处的交接边界，是构建现代 AI 与高性能计算编译基础设施的最佳实践。”*

👈 上一期回顾：**Day 7：参数传递、发射端口饥饿、动态调用与内存布局**
👉 专栏返回：查看更多模块内容与更新计划，请访问 **GPGPU 编译器面试通关指南目录**。

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

💡 核心考察维度：循环融合依赖合法性判定（Distance/Direction Vector）、MLIR One-Shot Bufferize 别名与就地复用分析、GPU 寄存器溢出惩罚与 Rematerialization 重计算博弈、Hopper/CDNA3 集群同步（Cluster Barrier）调度约束、GPU 弱内存模型与 Release-Acquire 缓存一致性下发。

📝 今日高频面试题清单

【Q1 · 图优化与循环变换】算子融合与 Loop Fusion 依赖合法性分析

在深度学习框架与 AI 编译器（如 TVM, XLA, MLIR）中，算子融合（Operator Fusion）是消除 HBM（High Bandwidth Memory）访问开销、提升访存密集型算子性能的关键技术。算子融合在底层通常体现为循环融合（Loop Fusion）。

1. 合法性分析：请说明编译器如何借助 Distance Vector（距离向量）和 Direction Vector（方向向量）严谨证明两个嵌套循环融合（Loop Fusion）的合法性？如果融合颠倒了依赖方向，会导致什么错误？
2. 代价权衡：当算子融合引入多余计算（如冗余重算/Recomputation）或导致局部缓存/寄存器压力剧增时，编译器应建立怎样的 Cost Model 来做出 Fusion 决策？

【Q2 · MLIR 体系与内存规划】MLIR Dialect 体系与 One-Shot Bufferization 机制

MLIR（Multi-Level Intermediate Representation）通过多级 Dialect 降低了端到端编译 AI 模型的复杂度。

1. Lowering 路径：请详细说明 MLIR 中计算密集型算子从高层 linalg / tensor Dialect 降级（Lowering）到硬件相关 gpu / nvvm / amdgpu Dialect 的典型流水线过程。
2. In-place Bufferization：详细阐述 MLIR 的 One-Shot Bufferize 机制如何通过 Alias Analysis（别名分析）与 Memory Effect Analysis（内存效应分析）实现 Tensor 到 MemRef 的内存复用（In-place Write），以最大限度减少内存分配开销。

【Q3 · 寄存器分配与代价模型】Register Pressure 驱动的 Spill 与 Rematerialization

GPU 架构（如 NVIDIA H100, AMD MI300）依靠庞大的 Register File 来隐藏内存访问延迟。然而，高级编译优化（如 Loop Unrolling、Instruction Scheduling、Software Pipelining）极易大幅增加寄存器压力（Register Pressure），触发寄存器溢出（Register Spilling）。

1. 重计算 vs. 溢出：编译器在处理严重寄存器压力时，如何权衡 Spill 到 Local Memory 与 Rematerialization（重计算）？
2. GPU Cost Model 特殊性：相比于 CPU 编译器，GPU 编译器在建立 Rematerialization 与 Spill 代价模型时有哪些特殊的硬件考量？

【Q4 · GPGPU 协作模型与并发控制】Cooperative Groups 与跨块同步（Cluster Barrier）

随着 NVIDIA Hopper 架构引入 Thread Block Cluster 以及 AMD CDNA3 架构引入 Workgroup Coalition，GPU 编程模型已经从单纯的 Block 内同步扩展到了跨 Block 集群同步（Cluster-level Synchronization）。

1. LLVM IR 中文表征：编译器如何在 LLVM IR 层级表征 Cooperative Groups 及 Cluster 级别的同步原语？
2. 指令调度约束：Cluster Barrier 对编译器的 ScheduleDAG（指令调度有向无环图）会产生哪些强约束？如何防止访存指令跨屏障非法重排？

【Q5 · 弱内存模型与缓存一致性】GPU 弱内存模型与 Release Consistency 编译

现代化 GPU 架构采用高度宽松的弱内存模型（Weak Memory Model）。

1. 内存序映射：编译器如何将高级语言中的 C++11 Memory Order（如 memory_order_acquire, memory_order_release）与 SyncScope 映射到 PTX / AMDGPU 硬件指令修饰符？

2. 缓存一致性履行：对于硬件级的 L1/L2 Cache 一致性，编译器需要在机器码生成阶段插入哪些特定的缓存失效（Invalidate）与写回（Writeback/Flush）指令，以履行 Release-Acquire 语义？

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

mindmap
  root((Day 3 核心考点图谱))
    循环融合与依赖
      距离向量与方向向量
      字典序非负判定
      HBM 带宽收益与重算代价
      寄存器溢出风险评估
    MLIR One-Shot Bufferize
      渐进式 Lowering 路径
      别名分析与读写冲突
      In-place 原地复用
      Memref 物理分配降级
    寄存器溢出与重计算
      Local Memory 物理开销
      Warp 放大效应与 Occupancy
      Rematerialization 低代价重算
      循环深度指数加权
    集群同步 Cluster Barrier
      Hopper / CDNA3 硬件原语
      syncscope cluster 语义
      ScheduleDAG Chain 依赖边
      异步事务令牌同步
    弱内存模型与一致性
      Acquire / Release 硬件映射
      PTX / AMDGPU 修饰符与等待
      L1 失效与写回 Flush
      RAW 竞争与静默数据损坏

题 1 深度解析：算子融合与 Loop Fusion 依赖合法性分析

1. 循环融合的依赖合法性分析

循环融合是将两个具有相同迭代域的循环合并为一个循环的过程。设原程序存在两个相邻循环 Loop $L_1$ 和 Loop $L_2$，融合后的循环为 $L_{12}$。

• 依赖向量表征：
  • 距离向量（Distance Vector, $\vec{d}$）：定义为依赖源迭代与依赖目标迭代之间的索引差值，$\vec{d} = \vec{i}{target} - \vec{i}{source}$。
  • 方向向量（Direction Vector, $\vec{\sigma}$）：定义为距离向量元素的符号，$\vec{\sigma} = (\text{sign}(d_1), \text{sign}(d_2), \dots, \text{sign}(d_n))$，其中包含 <, =, >。
• 依赖合法性条件：
  • 融合前，若存在跨循环的数据依赖（如 $L_1$ 写入数据 $A[i]$， $L_2$ 读取 $A[i]$），此时 $L_1$ 的计算严格先于 $L_2$ 执行，该依赖属于 Loop-Independent Dependency 或正常的跨循环 RAW 依赖（距离向量在融合维度上为正或零）。
  • 融合后，两个循环体在同一迭代步内交叉执行。如果融合导致某个数据依赖的方向向量在最外层融合循环上变为 负数（<），这意味着依赖目标（Target）在依赖源（Source）之前被执行，颠倒了因果关系（写后读变成了读后写，或使用了未初始化的数据）。
  • 法则：融合合法当且仅当融合后产生的循环依赖向量字典序非负（Lexicographically Non-negative），即最左侧的非零分量必须为正（>）。

非法融合示例：

原代码： 融合后（若逻辑错误）：

for (i = 0; i < N; i++) for (i = 0; i < N; i++) {

A[i] = B[i] + 1; (L1) A[i] = B[i] + 1;

for (i = 0; i < N; i++) C[i] = A[i+1] * 2; // 尝试读取未更新的 A[i+1]

C[i] = A[i+1] * 2; (L2) }

2. Cost Model 权衡与决策机制

算子融合并非百利而无一害，编译器需通过多维度代价模型评估以下冲突：

编译器的决策算法框架：

$$\Delta \text{Cost} = \text{Cost}{\text{fused}} - (\text{Cost}{L1} + \text{Cost}_{L2})$$

1. HBM 带宽收益评估：
   $$\text{Benefit}{\text{HBM}} = \text{Volume}(\text{Intermediate_Tensor}) \times \text{Latency}{\text{HBM_Access}}$$
2. 计算冗余开销：
   $$\text{Cost}{\text{Compute}} = \Delta \text{FLOPs} \times \text{Latency}{\text{ALU}}$$
3. 寄存器溢出惩罚：
   $$\text{Cost}{\text{Spill}} = \text{Predicted_Spill_Slots} \times \text{Latency}{\text{Local_Memory}}$$

当 $\Delta \text{Cost} < 0$ 且预测的 Register Count $\le \text{Threshold}_{\text{MaxReg}}$ 时，才允许执行融合。对于包含元素级重算与高寄存器压力的场景，编译器可选用 Partial Fusion（部分融合） 或 Tile-level Fusion（通过 Shared Memory 缓冲而非纯 Register 展开）。

题 2 深度解析：MLIR Dialect 体系与 Bufferization 机制

1. MLIR Dialect 下降（Lowering）典型流水线

MLIR 采用渐进式下降（Progressive Lowering）设计，从抽象的高层图语义逐步转换为贴近硬件特性的机器码生成方言：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15

+------------------------------------+
|   linalg / tensor Dialect          | (高层：表达 Tensor 形状、表达无侧副作用的数据流)
+------------------------------------+
                  |
                  |  1. One-Shot Bufferization (Tensor -> MemRef)
                  v
+------------------------------------+
|   linalg / memref / scf Dialect    | (中层：结构化控制流 scf.for/parallel，显示内存访问)
+------------------------------------+
                  |
                  |  2. Tiling & Thread/Block Mapping (分块并映射到 GPU 网格)
                  v
+------------------------------------+
|   gpu / nvvm / rocdl / llvm        | (底层：GPU 线程索引, NVVM IR / ROCDL IR)
+------------------------------------+

• Step 1: Tensor 到 MemRef 的转换（Bufferization）：将 Value-based 的抽象 Tensor 操作转化为 Pointer-based 的显示内存读写操作（memref）。
• Step 2: Loop Tiling & Distribution：使用 linalg.tile_to_forall_op 或 scf.tile 将高维矩阵运算划分为 Block 级与 Thread 级块，将循环轴映射为 gpu.block_id 和 gpu.thread_id。
• Step 3: Vectorization & Local Memory Promotion：将 Thread 级别的计算向量化为 vector Dialect 指令，并将共享数据提升（Promote）到 gpu.address_space<workgroup>（即 Shared Memory）。
• Step 4: Lowering to GPU/LLVM Target：scf 转换为 cf（Control Flow），vector 和 gpu 翻译为针对特定硬件的 nvvm（Nvidia PTX 抽象）或 rocdl（AMD HSAIL 抽象），最终通过 LLVM Backend 导出 PTX 或 HSA Code Object。

2. One-Shot Bufferization 与 In-place 内存复用

在高层 tensor 语义中，数据是不可变的（Value Semantics），任何修改都会产生新的 Tensor。Bufferization 的核心目标是将其映射到物理内存 memref，并极力避免额外的内存分配（memref.alloc）。

MLIR 的 One-Shot Bufferization 采用基于 SSA 的依赖图分析，决定一个 Tensor 操作能否在原位置写（In-place）：

1. 读后写检查（Read-Before-Write Conflict Analysis）：
   • 分析 Tensor 的生命周期。如果张量 $T$ 的某块内存在步 $K$ 被写入，但在步骤 $K$ 之后，仍有其他指令需要读取 $T$ 在步 $K$ 之前的值，则发生冲突。
2. **别名分析与关系图 (Alias Analysis & Origin Tracking)**：
   • 建立 Tensor 之间的追溯树（Find Root Tensor）。追踪两个 Tensor 操作数是否指向同一块底层物理内存缓冲区（Equivalent / Aliasing）。
3. In-place 决定逻辑：
   • 如果通过冲突检查证明：在写入 Tensor 时，原有数据不再有任何后续读取操作（No Read-After-Write Hazard），则允许 In-place Bufferize。
   • 此时直接复用输入 Tensor 的 memref 指针，将其转换成目标读写操作，避免分配新内存。
   • 若冲突成立，则必须 Out-of-place，插入 memref.alloc()，并生成拷贝指令（memref.copy）。

题 3 深度解析：Register Pressure 驱动的 Spill 与 Rematerialization

1. Spill 与 Rematerialization 权衡策略

在寄存器分配（Register Allocation, RA）阶段，当某点活活跃变量（Live Ranges）数量超过 GPU 硬件限制（如每个 Warp 最大 255 个寄存器）时，编译器必须采取措施削减活跃变量。

1
2
3
4
5
6
7
8
9
10

          寄存器压力过高 (Live Range > Hardware Limits)
                               |
              +----------------+----------------+
              |                                 |
     [Rematerialization]                  [Register Spill]
(代价低/无内存访问操作)               (代价高/涉及 Local Memory)
              |                                 |
* 重新计算基地址偏移                    * 将寄存器内容写回
* 常量 / 简单算术逻辑                  * L1/L2 缓存及 Local Memory
* 重新生成 Thread ID 逻辑              * 产生显存流量与延迟

• Rematerialization（重计算）原理：
  • 若某个活跃变量的值可以通过低成本计算轻松重新得到（例如：常量加载、简单的指针偏移计算、threadIdx.x * stride 等），编译器会在该变量被使用的地方丢弃（Kill） 该寄存器，并在需要该值的位置重新插入计算指令。
  • 优势：消除了对该变量的全局活跃期维持，完全避免了内存读写。
• Register Spill（寄存器溢出）原理：
  • 若活跃变量是通过高成本计算得到（例如矩阵乘中间累加值、复杂 Transcendental 函数），无法低成本重算，编译器会为其在 Local Memory（实际上是 HBM/DRAM 划出的区域，依靠 L1/L2 缓存）分配空间，插入 Store 指令写回，并在后续使用点插入 Load 指令。

2. GPU 架构下 Cost Model 的特殊考量

GPU 编译器的 Cost Model 极其严苛地惩罚 Spill，理由如下：

1. Local Memory 的真实物理属性：
   • GPU 没有单独的”物理 Local Memory 硬件”，Local Memory 实际上位于高延迟的 Global Memory（HBM）中。一旦 L1 Cache 缺失，Spill 将带来数百个 Clock Cycles 的停顿。
2. Warp Divergence 与 Thread-Private 放大效应：
   • 一个 Warp 有 32 个线程。一个单寄存器溢出意味着 $32 \times 4 \text{ Bytes} = 128 \text{ Bytes}$ 的数据读写。如果存在 Bank Conflict 或未对齐访问，性能将指数级下降。
3. Occupancy（线程块占用率）跌落：
   • GPU 依靠多线程并发隐藏延迟。如果一个 Kernel 要求的寄存器数量超出临界点（如从 32 个跳升至 64 个），将直接导致 SM 上能同时调度的 Active Warps 数量减半，大幅降低 GPU 的整体吞吐能力。

决策公式：

$$\text{Cost}{\text{Remat}} = \text{Latency}{\text{ALU_Insts}}$$

$$\text{Cost}{\text{Spill}} = \text{Latency}{\text{Local_Store}} + \text{Latency}{\text{Local_Load}} + \text{Penalty}{\text{Occupancy_Drop}}$$

在 GPU 编译器中，几乎优先选择一切可能的 Rematerialization，甚至宁愿在 Loop Preheader 中进行复杂的地址计算重算，也极力避免在循环体内部发生 Spill。

题 4 深度解析：Cooperative Groups 与跨块同步（Cluster Barrier）

1. LLVM IR 中的 Cluster Barrier 表征

在 LLVM IR 中，硬件屏障被表征为目标特定的 Intrinsic 指令，并结合内存屏障（Memory Fence）及同步作用域（SyncScope）属性。

对于 NVIDIA Hopper 架构的 Cluster Barrier，LLVM IR 通常表征为：m
; 显式 Cluster 级别 barrier 指令

1
2
3
4
5
6

call void @llvm.nvvm.barrier.cluster.wait()

; 带 SyncScope 的硬件屏障与内存屏障组合  
fence syncscope("cluster") release  
call void @llvm.nvvm.barrier.cluster.arrive()  
fence syncscope("cluster") acquire

• **syncscope("cluster")**：告诉编译器分析 pass（如 Alias Analysis, LICM），该内存屏障的作用域限定在同一个 Cluster 内部的所有 Thread Blocks 之间，超越了传统的 syncscope("block")（即 __syncthreads()），但小于 syncscope("agent") 或 syncscope("system")。

2. 对 ScheduleDAG 与指令重排的约束

在 LLVM 后端编译阶段，指令调度器（Instruction Scheduler）将基本块中的指令构建为 ScheduleDAG（Node 表示指令，Edge 表示依赖关系）。

1
2
3
4
5
6
7
8
9

    前序访存指令 (Global Store A)
               |
         [ Chain Edge ] (内存/控制依赖)
               v
Cluster Barrier (屏障指令 Node)
               |
         [ Chain Edge ] (内存/控制依赖)
               v
    后续访存指令 (Global Load B)

1. 构建 Chain 依赖边（Chain Edges）：

   • ScheduleDAG 使用特殊的 Chain 边（非数据流依赖，而是序依赖）来约束有侧副作用（Side-effect）的指令。

   • Cluster Barrier 被标记为具有 hasSideEffects 和 isBarrier 属性。调度器在构建 DAG 时，必须在 Barrier 节点与之前所有的 Memory Load/Store 节点之间建立一条 Chain 边，同时在 Barrier 节点与之后所有的 Memory Load/Store 节点之间建立另一条 Chain 边。

2. 禁止跨屏障重排：

   • 在 DAG 进行拓扑排序（Topological Sort）生成物理指令序列时，Chain 边形成了严格的依赖前驱关系。Barrier 之后的指令拓扑排序值必然大于 Barrier，阻断了指令越过 Barrier 的任何可能性。

3. 防止 Vector / Tensor Core 异步指令漏同步：

   • 对于异步拷贝指令（如 TMA - Tensor Memory Accelerator 或 cp.async），Barrier 还必须绑定异步事务令牌（Transaction Barrier）。编译器必须在 ScheduleDAG 中将 async_commit 和 async_wait 绑定到 Barrier 节点，确保流水线在屏障处彻底清空（Flush）。

题 5 深度解析：GPU 弱内存模型与 Release Consistency 编译

1. Memory Order 与 SyncScope 的硬件映射

GPU 内存模型是分层级的，包含了 Memory Model（如 Acquire/Release, Sequentially Consistent）与 Scope（Workgroup, Device, System）两个维度。

在 NVPTX 与 AMDGPU 后端中，LLVM 将 C++ 语义转换为带修饰符的硬件指令：

| C++11 内存序 + 作用域 | PTX 指令 (NVIDIA) | Inst. Modifier (AMD GFX9/GFX10) |

| :— | :— | :— |

| Acquire (CTA Scope) | ld.acquire.cta.b32 | ds_read_b32 (LDS 隐式保证) |

| Release (GPU Scope) | st.release.gpu.b32 | buffer_store_dword ... sc0 sc1 |

| Acquire (System Scope) | ld.acquire.sys.b32 | global_load_dword ... sc0 + s_waitcnt |

• NVIDIA PTX 映射：PTX 引入了显式的 .acquire 和 .release 语义修饰符，以及 .cta (Block)、.gpu (Device)、.sys (System) 作用域限定符。

• AMD ROCm 映射：AMD 使用指令位标志（如 sc0 - System Coherency, sc1 - Warp Coherency）结合显式的等待计数器指令（s_waitcnt vmcnt(0)）来实现。

2. 缓存一致性履行的底层机制

GPU 的 L1 Cache 通常是 Non-coherent（不保证硬件一致性） 或采用 Write-Through / Write-back 混合策略。因此，Release-Acquire 语义不能仅靠普通的 Store/Load 指令完成，编译器必须插入缓存控制指令。

Thread A (Release Side)                      Thread B (Acquire Side)
───────────────────────                      ───────────────────────
1. Write Data to Memory                      1. Loop until Flag == 1  
2. Store-Release (Flag = 1)                  2. Load-Acquire (Data)  
   ├── Write-back L1 Cache (Flush)              ├── Invalidate L1 Cache  
   └── Memory Barrier                           └── Memory Barrier
```                           └── Memory Barrier

1. **Release 语义（发送端）**：

   * **目标**：保证在 Release 操作之前的所有内存写操作，对其他 Thread 可见。

   * **编译器动作**：

     * 插入 Store 指令，并附加 Release 标志（如 `st.release`）。

     * 在硬件层面，这会强制将 Thread 所在 SM 的 L1 Cache 中脏数据（Dirty Lines）**写回（Flush / Writeback）** 到 L2 Cache 甚至 Global Memory。

     * 防止编译器将 Store-Release 之前的写指令重排到其之后。

2. **Acquire 语义（接收端）**：

   * **目标**：保证 Acquire 操作之后的所有内存读写，都能读取到最新写入的数据，而非旧的缓存副本。

   * **编译器动作**：

     * 插入 Load 指令，并附加 Acquire 标志（如 `ld.acquire`）。

     * 在硬件层面，这会触发 Thread 所在 SM 的 **L1 Cache 失效（Invalidate L1）**。

     * 强制后续的 Read 操作绕过当前 L1 脏缓存，直接从 L2 Cache 或 HBM 中重新拉取最新数据。

     * 在 AMD 架构下，编译器还必须紧跟插入 `s_waitcnt vmcnt(0)` 和 `buffer_wbinvl1_vol` 指令，以显式清空 L1 缓存管道。

---

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

> 1. **谈算子融合时**：*"循环融合绝非简单的代码合并，核心在于方向向量的字典序非负性检验；Cost Model 必须在 HBM 访存消除收益与单线程活范围膨胀引发的 Spill 之间寻求平衡。"*
> 2. **谈 Bufferization 时**：*"One-Shot Bufferization 改变了传统内存分配范式；通过生命周期与别名图分析消除 RaW 冲突，实现 Tensor 值语义到 Memref 指针就地复用的零拷贝跳跃。"*
> 3. **谈寄存器溢出时**：*"GPU 上的 Local Memory 本质是戴着缓存面具的高延迟 Global Memory；面对寄存器瓶颈，宁可在循环预头块甚至循环体内进行复杂的 Remat 重计算，也绝不能轻易放任 Spill。"*
> 4. **谈集群屏障时**：*"Thread Block Cluster 打破了单 CTA 的物理壁垒；编译器在 ScheduleDAG 中必须建立双向 Chain 内存依赖并绑定 Transaction Barrier，才能筑牢跨 SM 异步数据交互的安全防线。"*
> 5. **谈内存模型时**：*"GPU 弱内存模型依赖显式软硬件协同；从 syncscope 范围限定到 L1 Cache 失效与写回，漏插一条 barrier 或 waitcnt 就是静默数据错误的灾难。"*

---

> 👈 **上一期回顾**：**[Day 2：标量优化、Bank Conflict、软流水与指令发射](/gpu-compiler/2026-10-17-gpgpu-compiler-interview-day-2/)**
> 👉 **下一期推荐**：**[Day 4：GlobalISel、Fragment 布局、发散展开与异步流水](/gpu-compiler/2026-10-19-gpgpu-compiler-interview-day-4/)**
> 👉 **专栏返回**：查看更多模块内容与更新计划，请访问 **[GPGPU 编译器面试通关指南目录](/gpgpu-interview/)**。