[all-commits] [llvm/llvm-project] 73a366: [AMDGPU] Add SIPromoteGlobalLoadSAddr: promote loo...

[Matthias Gehre via All-commits]

  Branch: refs/heads/users/mgehre-amd/si-promote-global-load-saddr
  Home:   https://github.com/llvm/llvm-project
  Commit: 73a366fa244d8ad9e75e19e4a5c09a8d9869921d
      https://github.com/llvm/llvm-project/commit/73a366fa244d8ad9e75e19e4a5c09a8d9869921d
  Author: Matthias Gehre <matthias.gehre at amd.com>
  Date:   2026-06-24 (Wed, 24 Jun 2026)

  Changed paths:
    M llvm/lib/Target/AMDGPU/AMDGPU.h
    M llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
    M llvm/lib/Target/AMDGPU/CMakeLists.txt
    A llvm/lib/Target/AMDGPU/SIPromoteGlobalLoadToSAddr.cpp
    A llvm/test/CodeGen/AMDGPU/si-promote-global-load-saddr.ll

  Log Message:
  -----------
  [AMDGPU] Add SIPromoteGlobalLoadSAddr: promote loop-carried ptr phi to SADDR

Adds a pre-RA peephole pass that promotes global loads whose address is a
loop-carried vreg_64 phi of the form (sgpr_base + divergent_tid_offset) to the
SADDR addressing mode, replacing the VGPR address advance with scalar
S_ADD_U32/S_ADDC_U32 instructions.

Problem
-------
When a uniform global pointer is advanced by a uniform stride inside a loop
and indexed by threadIdx.x, the compiler fails to select the
`global_load ... SADDR+VADDR` form. Instead it keeps the full 64-bit address
as a per-lane VGPR induction variable and emits a pure-VGPR
`global_load ... off`, advancing the address with two V_ADD_CO_U32 per
iteration. This is the textbook SADDR+VADDR shape, but the compiler's
SelectionDAG divergence analysis conservatively treats the loop-carried phi
as divergent even though all incoming values are uniform.

Root cause
----------
The miss is in the LLVM mid-end: LICM cannot hoist the load (base is
loop-carried, not invariant), IndVarSimplify does not normalize pointer IVs
of the form {base, +, stride} into an integer IV, and ISel does not recover
the split from a loop-carried phi node. The fix is a post-ISel, pre-RA
peephole that operates on Machine IR where SGPR/VGPR assignment is concrete.

Pattern matched (after SIFoldOperands expands V_ADD_U64_PSEUDO)
---------------------------------------------------------------
Preheader:
  %lo, %carry = V_ADD_CO_U32_e64  sgpr_base.sub0, vgpr_tid_bytes
  %hi, dead _ = V_ADDC_U32_e64    sgpr_base.sub1, 0, carry
  %addr_init:vreg_64 = REG_SEQUENCE %lo, sub0, %hi, sub1

Loop:
  %addr:vreg_64 = PHI [%addr_init, ph], [%addr_next, loop]
  GLOBAL_LOAD_*  vdst, %addr:vreg_64, off
  %lo2, %c2 = V_ADD_CO_U32_e64  %addr.sub0, stride_imm
  %hi2, _   = V_ADDC_U32_e64    %addr.sub1, 0, c2
  %addr_next:vreg_64 = REG_SEQUENCE %lo2, sub0, %hi2, sub1

Transformation
--------------
Loop (after):
  %sbase:sreg_64 = PHI [sgpr_base, ph], [%snext, loop]
  GLOBAL_LOAD_*_SADDR  vdst, %sbase:sreg_64, vgpr_tid_bytes, off
  %snext_lo = S_ADD_U32  %sbase.sub0, stride_imm
  %snext_hi = S_ADDC_U32 %sbase.sub1, 0
  %snext:sreg_64 = REG_SEQUENCE %snext_lo, sub0, %snext_hi, sub1

The dead vreg_64 phi cycle is explicitly removed to avoid leaving dead
V_ADD_CO_U32 instructions that DCE cannot break due to the self-referential
phi structure.

Covers all common GLOBAL_LOAD widths (d16/b16, ubyte, ushort, dword, dwordx2,
dwordx3, dwordx4) and their _t16 variants. Tested on gfx1010 and gfx1100;
end-to-end correctness verified on gfx1151 (exact bitwise match).

Changes:
- SIPromoteGlobalLoadToSAddr.cpp: new pre-RA MachineFunctionPass
- AMDGPU.h: declare pass creation function and initialize symbol
- AMDGPUTargetMachine.cpp: register pass and insert after SIFoldOperands
  in GCNPassConfig::addMachineSSAOptimization
- CMakeLists.txt: add new source file
- si-promote-global-load-saddr.ll: lit test covering the promoted loop body
  with no v_add_co_u32 remaining (gfx1010, gfx1100)



To unsubscribe from these emails, change your notification settings at https://github.com/llvm/llvm-project/settings/notifications