Full LTO default on AMDGCN causes register spills and 7.6x slowdown in MIOpen's hipconv solver (gfx950)

[amd-alavin]

Summary

Since the AMDGCN default LTO mode became full, MIOpen's ConvHipConv convolution kernels regress by up to 7.6x on gfx950. The emitted computation is unchanged — identical MFMA and LDS instruction counts — but the register allocator hits the 256-VGPR ceiling and spills, including inside a software-pipelined inner loop. -fno-offload-lto recovers all of it.

Cause appears to be llvm#201457, "[Clang] Set default LTO mode for AMDGCN/SPIR-V targets to full", carried in this fork as fb37e56.

Reproducer (compile only — no GPU, no runtime, no link)

Clone ROCm/rocm-libraries at branch users/amd-alavin/hipconv-v0.3.0 (830389b1065). The config indices and the ConvParams type name below match that ref; develop carries an older hipconv snapshot where both differ. hipconv lives at projects/miopen/src/hipconv. Create repro.cpp, one template instantiation:

#include "direct/direct_wgrad/config_table.h"
#include "direct/direct_wgrad/kernel.h"
#include <hip/hip_runtime.h>
namespace hipconv::cdna4::direct_wgrad {
template void launch_impl<configs[45]>(const LaunchParams&, const ConvParams&,
                                       const void*, const void*, void*, void*, hipStream_t);
}

configs[45] is Config{3, 3, 2, 2, 1, 2, 4, ...}, the same kernel the runtime measurement below exercises.

Compile it twice, -fno-offload-lto the only difference:

HC=<clone>/projects/miopen/src/hipconv
INC="-I$HC/src/arch/cdna4 -I$HC/src -I$HC/include -I$HC/.."

# 1. current default (full LTO)
clang++ -x hip --offload-arch=gfx950 -O3 -std=gnu++20 $INC \
  --save-temps=obj -c repro.cpp -o out.o

# 2. with offload LTO disabled
clang++ -x hip --offload-arch=gfx950 -O3 -std=gnu++20 $INC \
  -fno-offload-lto --save-temps=obj -c repro.cpp -o out.o

The two runs leave differently named device assembly, so they do not overwrite
each other:

# 1 -> repro.cpp-hip-amdgcn-amd-amdhsa.hipfb.amdgcn.gfx950.img.lto.s
# 2 -> repro-hip-amdgcn-amd-amdhsa-gfx950.s
grep -c scratch_ *gfx950*.s

Then count spills in the generated device assembly:

	scratch ops	v_mfma	s_waitcnt vmcnt(0)
default (full LTO)	260	360	122
-fno-offload-lto	4	360	10

Same MFMA count both ways — the computation is identical. 65x the spills and 12x the full pipeline drains.

Not every config in that table is affected. Of the fourteen in this shard, configs[37], [45] and [85] regress and the other eleven are unchanged, so the trigger depends on the tile shape.

Impact (needs a gfx950)

export MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipConv

MIOpenDriver convbfp16 -n 8 -c 512 -H 384 -W 384 -k 512 -y 3 -x 3 \
  -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -g 1 -F 4 \
  -I NHWC -O NHWC -f NHWC -t 1 -V 1

toolchain	solver	time
ROCm 7.16.26362, clang 24.0.0git d6f6cb6	220/ConvHipConv	24.17 ms
ROCm 7.13.0, clang 23.0.0git 43215c7	220/ConvHipConv	3.18 ms

Both report Verifies OK — results are correct in each case, only performance differs. MI355X (gfx950), stock ROCm, nothing rebuilt.

Which toolchains are affected

The ROCm version string does not identify this — two ROCm 7.15.0 installs here differ. The clang commit does:

	clang
affected	24.0.0git d6f6cb6 (ROCm 7.16.26362)
affected	24.0.0git 5a93c2a
affected	23.0.0git f3f3b49
unaffected	23.0.0git aa451e1
unaffected	23.0.0git 46fcb33
unaffected	23.0.0git 43215c7

Detection:

echo '__global__ void k(float*p){*p=1;}' > /tmp/t.hip
clang++ -x hip --offload-arch=gfx950 -O3 -c /tmp/t.hip -o /dev/null -### 2>&1 \
  | grep -q -- -flto=full && echo AFFECTED || echo "not affected"

What changes in the generated code

The same kernel, as it appears in the full library build:

	-fno-offload-lto	default (full LTO)
total instructions	1800	2012
v_mfma	180	180
ds_read	106	106
ds_write	36	36
s_waitcnt	100	160
scratch access	0	126
VGPRs	247	256
ScratchSize	0 B	268 B
Occupancy	2	2

Every added instruction is spill machinery. All spills are marked 4-byte Folded Spill and hold address arithmetic, not accumulators:

    s_add_i32 s95, s55, 0xc000
    scratch_store_dword off, v18, off offset:236 ; 4-byte Folded Spill
    v_bitop3_b32 v18, v19, v1, 16 bitop3:0x36

Four spill/reload pairs land inside the kernel's software-pipelined inner loop, together with four s_waitcnt vmcnt(0). That loop keeps a prefetch in flight and waits only on a partial vmcnt; the full drains defeat it.

conv2d_direct_l1_cdna4 behaves the same way: 768 MFMA either way, +142 scratch accesses under LTO.

A second, separate effect appears in a different kernel family: loop structure changes (v_mfma 576 -> 1152) and ld.lld emits 150 loop not unrolled: the optimizer was unable to perform the requested transformation warnings, all at one #pragma unroll site. That family is the least affected of the three, so it does not explain the main regression.

Scope

Only build-time-compiled HIP C++ device code is affected. MIOpen kernels compiled at runtime through comgr are untouched — ConvDirectNaiveConv is HIP C++ and measures flat. Across a 300-layer sweep, restricted to calls of at least 100 us, hipconv kernels regress to 0.54x while assembly (0.999), Winograd (1.002), naive (1.000) and CK (1.006) solvers are unchanged.

Workaround

-fno-offload-lto restores 7.13.0-level performance to within 1% on every layer measured, passes the full functional suite, and is accepted without warning by every toolchain tested (7.3.0, 7.13.0, 7.15.0, 7.16.26362). It does not leave the new offloading model — --offload-new-driver and clang-linker-wrapper are still used; only the device -cc1 action returns to -emit-obj.

No minimal test case yet

The reproducer above is one template instantiation, but it still pulls in hipconv's kernel headers, so it is small in effort rather than minimal. We tried to shrink it to a standalone kernel and failed: a small kernel with #pragma unroll loops over a fixed trip count, holding a private array indexed by the loop variable, compiles to identical code with and without -flto=full. Whatever triggers the spilling needs something the real kernel has that our small version does not.

[amd-alavin]

Attached the IR fed to codegen:
hipconv_direct_wgrad_gfx950.precodegen.bc.gz

gunzip hipconv_direct_wgrad_gfx950.precodegen.bc.gz
clang++ -x ir hipconv_direct_wgrad_gfx950.precodegen.bc \
  --target=amdgcn-amd-amdhsa -mcpu=gfx950 -O3 -S -o out.s
grep -c scratch_ out.s      # 126, with 180 v_mfma

Preprocessed source won't reproduce it -- under -flto=full the device cc1 emits bitcode and codegen happens in clang-linker-wrapper.

[jhuber6]

This appears to be a scheduler cliff exposed by full LTO. The pre-link pipeline inlines helpers before link-time internalization, producing a wider CFG and losing some downstream SimplifyCFG/CSE opportunities. This then uses all 256 VGPRs and we spill into the stack, unable to spill SGPRs to a VGPR + v_readlane.

Unfortunately, I'm not sure if there is a simple solution that solves this, as it seems to be primarily related to already thin thresholds. I noticed that changing the scheduler approach at least recovered the stack spill, but I was unable to run it. @amd-alavin would it be possible for you to test with -Xoffload-linker --mllvm=-amdgpu-sched-strategy=max-ilp and compare it to -fno-offload-lto? I would assume that we would still get slightly wider CFGs but it might recover the scheduler enough to minimize the impact.

[petar-jovanovic]

This appears to be a scheduler cliff exposed by full LTO. The pre-link pipeline inlines helpers before link-time internalization, producing a wider CFG and losing some downstream SimplifyCFG/CSE opportunities. This then uses all 256 VGPRs and we spill into the stack, unable to spill SGPRs to a VGPR + v_readlane.

Unfortunately, I'm not sure if there is a simple solution that solves this, as it seems to be primarily related to already thin thresholds. I noticed that changing the scheduler approach at least recovered the stack spill, but I was unable to run it. @amd-alavin would it be possible for you to test with -Xoffload-linker --mllvm=-amdgpu-sched-strategy=max-ilp and compare it to -fno-offload-lto? I would assume that we would still get slightly wider CFGs but it might recover the scheduler enough to minimize the impact.

@jhuber6 my findings are a little bit different.

SlotIndexes never removes the index list entry of an instruction that has been erased. SlotIndexes::removeMachineInstrFromMaps only clears the instruction pointer:

  } else {
    // FIXME: Eventually we want to actually delete these indexes.
    MIEntry.setInstr(nullptr);
  }

LiveInterval::getSize() is measured in slot indexes, and DefaultPriorityAdvisor::getPriority ranks global live ranges by that size, longest first. Because the leftover entries are not spread evenly over the function, each range's size is inflated by a different factor, so the ranking
no longer reflects the ranges' real extents. On a register-starved function that is enough to turn a clean allocation into heavy spilling.

[petar-jovanovic]

@jhuber6 Here are the changes https://github.com/llvm/llvm-project/tree/users/petar-jovanovic/slotindexes-compaction that add SlotIndexes::compactIndexes(), which drops those entries before register allocation so that sizes measure real code distance again. It is off by default behind -greedy-compact-slot-indexes.

[jhuber6]

Seems promising, my observations were very surface level about where the LTO pass pipeline diverges. I checked out your commit and observed the spills went away with -Xoffload-linker --mllvm=-greedy-compact-slot-indexes.