From 9b94708706dd397bb808a6347591e7636263bb67 Mon Sep 17 00:00:00 2001 From: Richard-Qin-X Date: Thu, 10 Sep 2026 00:39:43 +0800 Subject: Add register allocation note series --- .../register-allocation/gpu-occupancy/index.md | 384 +++++++++++++++++++++ 1 file changed, 384 insertions(+) create mode 100644 content/notes/register-allocation/gpu-occupancy/index.md (limited to 'content/notes/register-allocation/gpu-occupancy/index.md') diff --git a/content/notes/register-allocation/gpu-occupancy/index.md b/content/notes/register-allocation/gpu-occupancy/index.md new file mode 100644 index 0000000..ea000b7 --- /dev/null +++ b/content/notes/register-allocation/gpu-occupancy/index.md @@ -0,0 +1,384 @@ +--- +title: "第十部分: GPU 寄存器分配与 Occupancy" +slug: "gpu-occupancy" +lang: zh-Hans +series: register-allocation +weight: 110 +created: 2026-08-18 +updated: 2026-09-09 +license: CC-BY-SA-4.0 +--- + +CPU 和 GPU 都需要把数量近乎无限的 virtual registers 映射到有限的 physical registers, 因而 liveness, interference, splitting, spilling 和 coalescing 这些基本问题并没有因为执行设备变成 GPU 而消失. 真正发生变化的是寄存器资源参与性能的方式. 由于 GCC 在 GPU 上已经落后很多了, 所以这一部分我们主要以 LLVM 来分析. + +CPU allocator 主要考虑单个 instruction stream 中的寄存器复用成本. GPU 同时维持大量 threads, 并把它们组织成 warp 或 wavefront 执行. 每个 thread 使用的寄存器最终都要占用 SM 或 CU 上有限的 register file. 当单个 kernel 的寄存器需求增加时, 一个执行单元能够同时驻留的 warps 或 waves 可能减少, 进而降低 GPU 隐藏 memory latency 和 execution latency 的能力. NVIDIA 的 CUDA 执行模型明确把 registers 和 shared memory 都视为决定 SM resident blocks 和 warps 数量的资源. AMD GPU 上的 SGPR 和 VGPR 同样从有限 register pools 中分配给正在执行的 wavefronts. + +所以 GPU register allocator 面对一个 CPU 上弱得多的反馈关系: + +```text +register allocation + | + v +register usage per thread / wave + | + v +resident warps / waves + | + v +occupancy + | + v +latency hiding and throughput +``` + +这使 GPU RA 的优化目标发生了变化. 减少一次 spill 当然仍然有价值, 但如果为了避免 spill 而让整个 kernel 跨过一个 register-allocation threshold, 使 resident waves 从 4 个下降到 3 个, 最终性能可能反而下降. 反过来, 强行减少寄存器数量也可能引入大量 memory traffic 或额外 instructions. GPU RA 因此长期处在 register pressure, spill cost, instruction-level parallelism 和 occupancy 之间的权衡中. NVIDIA 的 `--maxrregcount` 文档直接把 per-thread register usage 和 thread parallelism 描述为一项 trade-off, AMD 的 GCN scheduler 也专门根据 register pressure 是否会降低 wave occupancy 调整调度策略. + +## 1 从 Thread Register 到整个 Register File + +先把 GPU 的执行层次和寄存器资源对应起来. CUDA 中的 threads 被组织成 thread blocks, block 进入 SM 后再被划分成 warps. 一个 warp 包含 32 个 threads, warp scheduler 从 ready warps 中选择下一条 instruction 发射. 每个 SM 拥有一组 32-bit registers, 这些寄存器资源会在驻留于该 SM 的 warps 之间分配. + +从程序员视角看, register 是 thread-private 的: + +```text +thread 0: r0 r1 r2 ... +thread 1: r0 r1 r2 ... +thread 2: r0 r1 r2 ... +... +``` + +但物理资源来自同一个有限 register file. 如果一个 warp 中每个 lane 都需要一个 32-bit value, 那么硬件需要为整个 warp 保存这一组 lane values. 所以 "kernel 使用 64 个 registers" 这类指标通常不是说整个 kernel 只占 64 个 hardware words, 而是在描述每个 thread 或每个 wave 的 architectural register requirement, 最终硬件会把它乘到同时驻留的执行上下文上. NVIDIA 也因此把 register usage 作为 per-thread kernel resource 报告, 并用它限制同时可以驻留的 threads 和 blocks. + +AMD 的 register organization 更能体现 GPU RA 与普通 CPU RA 的差别. AMD GCN/CDNA/RDNA 系列存在 SGPR 和 VGPR 两类主要 general-purpose register resources. VGPR 保存每个 lane 可以不同的 vector value, SGPR 保存编译器能够证明在整个 wave 中一致的 scalar value. 因而一个 wave-uniform 的值如果进入 SGPR, 整个 wave 只需要保存一份; 一个 divergent value 进入 VGPR 后, 每个 lane 拥有自己的 component. + +例如: + +```text +wave: + +lane 0: x = 3 +lane 1: x = 7 +lane 2: x = 5 +lane 3: x = 9 +... +``` + +这种 `x` 是 divergent value, 需要 vector storage. 另一个值: + +```text +wave: + +lane 0: n = 128 +lane 1: n = 128 +lane 2: n = 128 +lane 3: n = 128 +... +``` + +如果 compiler 能证明 `n` 在整个 wave 中 uniform, AMDGPU 可以让它走 scalar register bank. Uniformity analysis 因而和 register pressure 直接连接. 一个 value 从 VGPR 需求转化成 SGPR 需求时, 减少的不只是一个普通颜色冲突, 还改变了它消耗哪一类硬件资源. LLVM 的 AMDGPU register-bank logic 也明确区分 SGPR, VGPR 和 AGPR, 并要求只有已知 uniform 的 value 才能合法从 VGPR 语义转入 SGPR. + +部分 AMD GPU 还拥有 AGPR, 即 accumulator register resource, 用于 matrix accumulation 等操作. LLVM AMDGPU 的 pressure model 区分 SGPR, VGPR, AGPR 和 unified/mixed vector register categories, 并额外追踪 tuple pressure. 这意味着一个 AMDGPU kernel 的 register pressure 从一开始就不是一个单独整数. + +## 2 Register Pressure 为什么会改变 Occupancy + +Occupancy 表示一个 execution unit 上实际能够同时驻留多少 waves/warps 相对于硬件允许上限的程度. 限制 occupancy 的资源不只有 registers, 还包括 shared memory 或 LDS, thread-block/workgroup size, architectural wave slots 等. 可以把这种关系抽象写成 $W_{\text{resident}}=\min(W_{\text{arch}},W_{\text{reg}},W_{\text{smem}},W_{\text{block}},\ldots)$. + +其中 register-limited occupancy 可以粗略理解成 register file 容量除以每个 resident execution context 需要的 register storage. 假设某种抽象 GPU 每个 execution unit 有 $R$ 个可用 register units, 每个 wave 分配 $r$ 个 units, 那么单从寄存器容量得到的上限近似为 $\lfloor R/r\rfloor$. 真实硬件还会按照一定 granularity 分配 registers, 因而 occupancy 往往表现成阶梯函数, 而不是随着 register count 连续变化. LLVM AMDGPU 中专门存在 VGPR/SGPR allocation granule 以及根据 register 数量计算 waves-per-EU 的 target functions. + +用一个纯粹的抽象例子说明这个现象. 假设 register file 容量为 256 个 units, 分配粒度为 16. 如果一个 wave 需要 64 个 units, 最多可以容纳 4 个 waves. 如果最大 pressure 从 64 增长到 65, 硬件实际需要向上取整到 80, register-limited occupancy 就可能直接从 4 个 waves 下降到 3 个: + +```text +logical pressure allocated amount resident waves + + 64 64 4 + 65 80 3 + 79 80 3 + 80 80 3 +``` + +于是一个只增加了 1 个 live register 的代码变换, 有可能产生远大于 "多占一个寄存器" 的资源后果. 这类位置常被称为 occupancy cliff 或 occupancy threshold. Register allocator 和 scheduler真正关心的经常不是把 pressure 从 53 降到 52, 而是能否跨过一个会增加 resident-wave count 的 threshold. + +AMD 的公开架构资料能够直接看到这种效果. MI100 的 VGPR 和 SGPR 会动态分配给 wavefronts, VGPR 使用量直接限制并发 wavefront 数. AMD 给出的一个具体例子中, wave 使用 119 个 VGPR 时只能达到 2 个 active wavefronts. + +但 occupancy 不能单独当成 performance metric. 更多 resident waves 可以提供更多独立工作, 用一个 wave 的执行去覆盖另一个 wave 的 memory 或 pipeline latency. 当 kernel 本来已经有足够 latency hiding, 继续提高 occupancy 未必产生明显收益. 如果为了获得更高 occupancy 而减少 registers, compiler 还可能被迫增加 reload, rematerialization 或 dependency chains. NVIDIA 的 launch-bounds 文档甚至允许 compiler 在 occupancy requirement 已经满足后增加 register usage, 用更多 registers 换取更少 instructions 或更好的 single-thread latency hiding. + +所以 GPU allocator的目标不能写成简单的 "最小化 register count". 更接近实际问题的表达是: 在不引入过高 spill 和 instruction cost 的前提下, 控制 register usage 落在合适的 occupancy range 内. + +## 3 Liveness, Scheduling 与 Divergence + +GPU register pressure 仍然从 liveness 产生. 如果某个程序位置上有 80 个 VGPR-width values 同时 live, allocator 无法通过重新给这些 values 换 register number 来把 pressure 变成 40. 真正能够改变 maximum pressure 的办法仍然是缩短 live range, splitting, rematerialization, spilling, 改变 instruction schedule, 或者让某些 values 转移到其他 register bank. + +GPU instruction scheduling 因此和 RA 结合得非常紧. 假设有几个高 latency loads: + +```text +load a +load b +load c + +compute a +compute b +compute c +``` + +把 loads 提前可以制造更多 memory-level parallelism, 让 GPU 同时等待多个 memory operations. 代价是 `a`, `b`, `c` 从 load 完成到真正消费之间都必须保持 live: + +```text +a: |----------------| +b: |-------------| +c: |----------| +``` + +如果改成: + +```text +load a +compute a + +load b +compute b + +load c +compute c +``` + +live ranges 会缩短, register pressure 下降, 但 memory-level parallelism 和 instruction-level parallelism 也可能下降. + +AMDGPU 的 machine scheduler 直接把这个矛盾放进 scheduling policy. `GCNSchedStrategy` 会对一个 kernel进行多阶段 scheduling, 原因就是某些 scheduling regions 的 register pressure 会决定整个 kernel 的 occupancy. 当某个 region 接近会降低 occupancy 的 pressure threshold 时, scheduler 更倾向于压低 SGPR/VGPR pressure; 在不会影响 occupancy 的区域, 约束可以放松, 让 scheduler继续追求 ILP. + +这说明 GPU 上 "schedule first, register allocate later" 并不意味着两者彼此独立. Scheduler 虽然在 RA 前工作, 却必须预测 RA 后可能达到的 register usage. RA 随后又在 scheduler塑造出的 live ranges 上工作. + +Divergence 则进一步改变 register bank pressure. 一个基于 lane ID 的 value 通常会沿不同 lanes 产生不同结果, 后续计算也继续保持 divergent. 在 AMDGPU 上, 这种数据通常进入 VGPR. Workgroup ID, 某些 kernel arguments, constant addresses 或其他 wave-uniform data 则有机会进入 SGPR. AMD 的编程和 profiling 文档明确把 VGPR 描述为保存 wave 内不同 work-items 的数据, SGPR 保存编译期已知 wave-uniform 的数据. + +因此下面两种 IR 虽然从普通 SSA liveness 看都只有一个 value: + +```text +v = thread_id.x + 1 +``` + +和: + +```text +v = workgroup_id.x + 1 +``` + +对 AMD GPU 的 register allocation 可能完全不同. 前者随 lane 变化, 通常形成 vector value. 后者对一个 wave 中的 lanes 可以保持 uniform, 有机会使用 scalar resources. Uniformity analysis 做错或过于保守时, 本来只需要一个 scalar value 的数据可能长期占据 VGPR, 增加 vector pressure. + +SGPR 也不是无限资源. LLVM AMDGPU 的 occupancy model能够分别计算 SGPR 和 VGPR 对 occupancy 的影响, 某些架构代际上 SGPR allocation本身也会形成 occupancy limit. `GCNRegPressure` 因此比较 register-pressure states 时先比较 occupancy, 再比较 spilling, tuple pressure 和 raw register pressure, 而不是简单把所有 registers 加在一起. + +控制流 divergence 还会让 live range 形状更加复杂. 一个 value 可能只在部分 active lanes 上真正参与计算, 但 physical vector register仍然必须为整个 wave 保留对应 storage. AMD 的 `EXEC` mask 决定哪些 lanes 在当前 vector instruction 中 active, register allocation却不能把同一 VGPR 的不同 lanes随意分给无关 virtual values, 除非 backend 有专门的 lane/subregister packing 机制. 所以前面 CPU 上 "两个值在不同 CFG paths 上互斥, 可以共享寄存器" 的思想到了 SIMT 环境还需要同时理解 control-flow mask 和 register-lane semantics. + +## 4 LLVM AMDGPU 如何做 Register Allocation + +AMDGPU 的 physical register allocation发生在 LLVM machine-code backend 中. 上层可以经过 Clang, OpenMP/HIP lowering, MLIR GPU/ROCDL 等多种路径, 但进入 `lib/Target/AMDGPU` 后, allocator处理的已经是目标相关 Machine IR, 最终生成 AMD GPU ISA. + +AMDGPU 没有重新实现一套完全独立的 allocator. 它复用了 LLVM target-independent Greedy RA infrastructure, 同时把 GPU register banks 拆成多个 allocation stages. 优化编译中, SGPR 和 VGPR 分别运行 Greedy allocator, whole-wave/whole-quad mode 所需 registers 也有独立 allocation stage. AGPR 相关 values 在 RA 前经过专门 preparation, allocation 完成后还有 AGPR copy rewrite 等 target passes. + +主干可以简化成: + +```text +Machine IR + | + v +AMDGPU pre-RA optimizations +and occupancy-aware scheduling + | + v +prepare AGPR allocation + | + v +SGPR register allocation + | + v +lower SGPR spills + | + v +WWM / special register allocation + | + v +VGPR register allocation + | + v +rewrite virtual registers + | + v +post-RA AMDGPU passes +``` + +SGPR 和 VGPR 分开 allocation 很符合 AMD 硬件的资源结构. Scalar values 和 vector values 不只是两个不同 register classes, 它们拥有不同 storage semantics, 不同 spill possibilities, 对 occupancy 的作用也不同. Backend 因而可以在 SGPR allocation结束后先 lower 一部分 scalar spills, 再继续处理 VGPR allocation. + +Greedy RA 本身仍然使用第九部分讨论过的 `LiveInterval`, `LiveRegMatrix`, eviction 和 splitting. GPU-specific behavior主要通过 register classes, allocation order, target-specific spiller, scheduler 和 pressure model进入这套 framework. 也就是说, AMDGPU 并没有把 LLVM Greedy 换成一个名字叫 "GPU allocator" 的完全不同算法, 它改变的是 Greedy 所面对的 physical-resource topology 和 cost structure. + +`GCNRegPressure` 很适合观察这种 target-specific cost model. 它分别维护 SGPR, VGPR, AGPR 等 register kinds, 同时维护 tuple pressure. 比较两个 pressure states 时, 优先考虑哪个状态能够获得更高 occupancy, 随后才比较 spill risk, tuple pressure 和 raw register count. + +这里可以看到 CPU 和 GPU RA 的一个结构差异. CPU 上 allocator通常希望某个热点 region 使用较少 registers, 但一个函数最大 pressure 增加 1 不一定对全函数产生离散的执行资源变化. GPU 上一个 region 的 maximum VGPR pressure 可能决定整个 kernel 的 allocated register count, 最终改变所有 waves 的 residency. 所以 AMDGPU scheduler 和 allocator需要把局部 live-range decision 和 kernel-wide resource consequence联系起来. + +AGPR 又增加了一层 register-bank problem. 某些 matrix instructions 可以使用 accumulator registers, backend需要判断一些 values 保持在 VGPR 还是转换到 AGPR 更合适, 并考虑为这种转换插入 copies. AMDGPU 的 scheduling 和 RA machinery因此还会处理 VGPR/AGPR register-bank choice 和相关 copy cost. + +## 5 AMDGPU Spill 为什么比普通 Stack Spill 更复杂 + +CPU spilling 常被简化成 "store 到 stack slot, use 前 load 回来". AMDGPU 同样有 scratch/private memory 作为寄存器溢出的 backing storage, 但不同 register banks 的 spill path 并不相同. AMDGPU backend为 scratch access准备相应的 private-segment state, VGPR spill 可以通过 scratch memory保存每个 lane 的 private value. + +VGPR spill 的直观形态可以画成: + +```text +VGPR live range + +|-------------------------------| + + high pressure + | + v + + store lane value -> scratch + ... + load lane value <- scratch +``` + +Scratch 对每个 lane提供 private address semantics, 所以一个 spilled VGPR实际上对应 wave 中各 lanes 各自的 spilled value. Memory backing 和 register file 的组织方式不同, backend需要根据 wave size 和 scratch addressing convention产生正确地址. + +SGPR spill 更有 GPU 特征. 因为 SGPR 中保存的是整个 wave 共用的一份 scalar value, 直接把它按每个 lane复制到 scratch 会浪费 memory bandwidth. LLVM AMDGPU 可以优先把 SGPR spill 到 VGPR 的某个 lane, 使用 lane read/write operations 保存和恢复 scalar value. 如果最终必须进入 memory, backend可以先把 SGPR value写入 temporary VGPR lane, 再把 VGPR 保存到 scratch; restore 时反向执行. + +于是 SGPR spilling 可能形成这样的资源转换: + +```text +SGPR pressure + | + | spill scalar value + v +VGPR lane +``` + +如果 VGPR 也已经紧张, 才进一步变成: + +```text +SGPR + | + v +temporary VGPR lane + | + v +scratch memory +``` + +这种设计减少 memory spill 的机会, 但它把 scalar-register pressure转移成 vector-register pressure. 假设 kernel 正好位于 VGPR occupancy threshold 附近, 多占一个 VGPR 可能比一次普通 SGPR spill 更昂贵. 所以 GPU spilling 的成本不能只按照 "load/store 次数" 建模, 还需要考虑 spill medium 本身占用了哪一类 register resource. + +AMDGPU 还存在 VGPR 与 AGPR 之间利用另一类 register storage 避免 memory spill 的机会. Backend 中有为 VGPR-to-AGPR 或 AGPR-to-VGPR spill保留寄存器的机制. 这使 register-bank pressure能够在一定条件下互相转移. + +Scratch 的使用本身也可能要求 kernel准备额外的 scratch-related SGPR state. AMDGPU backend 会根据 kernel 是否需要 private/scratch access建立相应的 scratch buffer 或 flat-scratch state. 因而 spilling 还可能反过来增加少量固定 register resource 和 prologue setup. + +这就是 GPU allocator里一个经常出现的反馈环: + +```text +high VGPR pressure + | + v +spill + | + v +scratch instructions + | + v +more address / temporary state + | + v +new register pressure +``` + +好的 spiller需要避免让 spill rewrite 本身制造新的大规模 pressure problem. Rematerialization, splitting 和 spill-to-other-register-bank 都是在尝试切断这种反馈. + +## 6 NVIDIA: LLVM 到 PTX, 再到真正的 Physical Register Allocation + +NVIDIA 的 LLVM 路径与 AMDGPU 有一个架构上的根本差异. LLVM `NVPTX` backend 的目标是 PTX, 而 PTX 本身是 virtual ISA. LLVM NVPTX 把 LLVM IR 转换成 PTX assembly, PTX 中仍然可以声明大量 virtual registers. NVIDIA 的 PTX ISA甚至专门提供参数化 `.reg` 声明, 方便 compiler生成大量 `%r0`, `%r1`, `%r2` 这样的 virtual register names. + +例如 PTX 可以出现: + +```text +.reg .b32 %r<100>; +``` + +这里声明的是 100 个 PTX virtual registers, 不是说 GPU hardware 已经为该 thread分配了 100 个 physical registers. PTX 之后还要经过 NVIDIA backend compiler. + +LLVM/NVIDIA pipeline 可以表示成: + +```text +LLVM IR + | + v +LLVM NVPTX backend + | + v +PTX virtual ISA + | + +----------------------+ + | | + v v + ptxas CUDA Driver JIT + | | + +----------+-----------+ + | + v + cubin + | + v + native GPU machine code +``` + +CUDA Driver 可以把 PTX JIT compile 成 native GPU machine code, `nvcc` 文档把 `ptxas` 定义为 PTX optimizing assembler. `ptxas --resource-usage` 的结果会报告最终使用多少 registers, 以及多少 spill loads/stores因为 variables 无法放进 physical registers 而产生. 因而 NVIDIA 路径中的最终 physical register assignment发生在 PTX optimizing backend这一侧, 不是 LLVM NVPTX 在生成 PTX 时完成. + +对 AMDGPU, LLVM Machine IR 中可以观察 Greedy allocator如何把 virtual VGPR/SGPR 映射成最终 architectural registers. 对 NVIDIA, LLVM NVPTX 输出仍然处于 virtual-register层, 最终 SASS register allocation由 NVIDIA backend完成. 修改 LLVM target-independent Greedy allocator不会直接替换 NVIDIA SASS allocator. + +NVIDIA 仍然把 register pressure控制接口暴露在 PTX 和 CUDA compilation model 中. PTX 的 `.maxnreg` 指令可以限制 kernel 每个 thread最多分配多少 registers, `.minnctapersm` 可以向 backend表达希望每个 SM 至少容纳多少 CTAs. 这些 directives 会让 optimizing backend在 per-thread register count 和 SM utilization之间做取舍. + +CUDA 层还有 `--maxrregcount` 和 `__launch_bounds__`. `--maxrregcount=N` 设置 GPU function 的 register upper bound. `__launch_bounds__(maxThreadsPerBlock,minBlocksPerMultiprocessor)` 则让 compiler根据希望同时驻留的 blocks数量推导 register budget. 如果原始 allocation 超过这个 budget, compiler会降低 register usage, 代价通常表现为更多 local-memory traffic 或更多 instructions. + +这提供了一个很直接的 RA experiment. 对同一个 kernel 编译多个版本: + +```text +max registers/thread: + +128 +96 +80 +64 +48 +``` + +然后观察: + +```text +register count +spill stores +spill loads +local memory usage +occupancy +kernel execution time +``` + +结果往往不会呈现 "register count 越低越快" 的单调关系. 在某个点之前, 减少 registers可能提高 occupancy而没有严重 spills; 继续压缩后, spill traffic开始主导性能. CUDA 文档也把 `maxrregcount` 明确描述为 individual-thread performance 和 thread parallelism 之间的 trade-off. + +NVIDIA register spill进入 local memory. CUDA 中的 local memory 是 thread-local address space, 但它的物理 backing 位于 device memory. Compiler在寄存器不足时可以把 value spill 到 local memory, 所以这种 spill拥有远高于 register access 的 latency, 同时经过 GPU memory hierarchy. + +PTX 还有 `.pred` 等不同 register types以及各种 special registers. 从 allocator角度看, 它们再次说明 GPU machine state 并不只是一个统一的整数 register file. 最终 NVIDIA backend还需要处理 predicate state, general registers, special architectural state以及不同 machine instructions 的 operand requirements. PTX virtual ISA把这些细节留给后端 compiler继续降低. + +## 7 CPU RA 与 GPU RA 的目标函数 + +经过 AMDGPU 和 NVIDIA 两种实现以后, 可以把 GPU register allocation 和前面的 CPU register allocation放在同一个框架下比较. + +|问题|CPU|GPU| +|---|---|---| +|Register pressure 的主要后果|spill, copy, instruction constraints|spill之外还可能降低 occupancy| +|一个 register assignment 的影响范围|主要影响当前 instruction stream|per-thread/per-wave usage 会乘到大量 resident execution contexts| +|Spill backing|通常是 stack memory|private/local/scratch memory, 某些 GPU 还可利用其他 register bank| +|Scheduler 与 RA|ILP, latency 和 pressure 权衡|ILP, latency, pressure 和 occupancy threshold 权衡| +|Register classes|GPR, SIMD, special regs 等|scalar/vector/accumulator/predicate 等多类资源| +|Control flow|CFG liveness|CFG liveness再叠加 SIMT divergence 和 execution mask| +|评价指标|runtime, spills, copies, code size|runtime, spills, register count, occupancy, wave residency, code size| + +CPU allocator中一个常见目标是尽量避免 hot-path spill, 同时减少 copies 和 callee-saved cost. GPU allocator仍然追求这些目标, 但 register count本身开始成为 kernel-level resource metric. 一次看起来很便宜的 live-range extension可能让 maximum pressure跨过 occupancy threshold, 影响整个 kernel 的 resident-wave capacity. + +反过来, GPU allocator也不能把 occupancy最大化作为唯一目标. 一个 compute-heavy kernel如果拥有大量独立 arithmetic instructions, 更多 per-thread registers可以保存更多 intermediate values, 增加 ILP并避免 scratch traffic. 一个 memory-latency-bound kernel可能更愿意牺牲一些 single-thread efficiency, 换取更多 resident waves. AMD 的 occupancy-aware scheduler和 NVIDIA 的 launch-bounds机制都反映了这种取舍. + +所以评价 GPU register allocator时至少要把 register count和程序性能一起测量. 单独报告 "VGPR 从 72 降到 64" 这句话就没有包含足够信息. 如果 72 和 64 位于同一个 occupancy bucket, 这次 reduction可能没有任何 residency收益; 如果 64 刚好跨过一个 threshold, 效果可能很大. 同样, occupancy 从 50% 提高到 75% 也不能自动说明优化成功, 因为实现可能用大量 scratch spill换来了这个数字. + +GPU RA 的很多高级优化最终都围绕这种非线性成本展开. Scheduler可以重新排列 instructions 缩短 pressure peak, allocator可以围绕 high-pressure region splitting, rematerialization可以用 computation换 register residency, uniformity analysis可以把 VGPR pressure转移到 SGPR, register-bank choice可以在 VGPR 和 AGPR 之间重新分布资源. 有些研究甚至会有意识地接受少量 spill, 只为了把 maximum register count压过一个 occupancy threshold. + +到这里, 前十部分已经把 register allocation 从最简单的 virtual-to-physical mapping一路推进到了 CPU 和 GPU 的工业实现. 下一部分可以在这些基础上讨论更高层的协同优化问题: register-pressure-aware scheduling, profile-guided allocation, loop-aware splitting, rematerialization, instruction selection与 RA 的相互作用, occupancy-aware GPU allocation, allocator benchmark方法. -- cgit v1.2.3