[AMDGPU] Clang 23: 11 min / 2.4 GB to compile one CK FMHA HIP instance at -O1 (MachineScheduler + Greedy RA)

## Summary

Clang 23's AMDGPU backend takes **~11 minutes and 2.4 GB RSS** to compile **one** Composable Kernel FMHA backward explicit instantiation at `-O1`, for a single offload arch, with device LTO and `-amdgpu-early-inline-all` both **off**.

Official [xformers](https://github.com/facebookresearch/xformers) ROCm wheels (`v0.0.35`, `.github/workflows/wheels.yml` + `wheels_build.yml`) build the **same sources** under **ROCm 7.1**, whose `amdclang++` is **LLVM 20.0.0**, on a 16-core / 64 GB runner, `MAX_JOBS=2`, `HIP_ARCHITECTURES='gfx90a gfx942'`, 6 hour timeout. That job would be impossible if each of the ~624 instance TUs took 11 minutes (even at `-j2`, and they compile **two** archs).

This is **not** https://github.com/llvm/llvm-project/issues/86332 (`-amdgpu-early-inline-all=true` 8× memory). That flag is omitted here.

`hipcc` is not involved in the cost: it is a ~5 MB process asleep on a pipe. The time is `clang -cc1 -triple amdgcn-amd-amdhsa`.

## Environment

```
OpenMandriva 23.1.0~rc3-2 clang version 23.1.0
(llvm-project-llvmorg-23.1.0-rc3/clang cf7382be50b5d4f7c32d8cd69388f47ef1c06642)
HIP 7.14.0-therock-7.14
Target: gfx90a (one arch)
Host: 32-thread Ryzen, 64 GB RAM
```

## Test case

A 15-line explicit instantiation — the file xformers already splits out as one ninja TU:

```hip
#include
#include "ck_tiled_fmha_batched_backward_hip.h"

template void run_batched_backward_mask_bias_dropout_dispatch<
ck_tile::bf16_t,
true, true, true, true,
64>(BatchedBackwardParams& param, hipStream_t stream);
```

Headers:

- xformers 0.0.35 `xformers/csrc/attention/hip_fmha/`
- composable_kernel `50fad035248b154cdfa4505cf5de7465ce146149` (xformers v0.0.35 pin)

No PyTorch headers.

```sh
hipcc -c -x hip instance.hip -o instance.o \
-I$XFORMERS/xformers/csrc/attention/hip_fmha \
-I$XFORMERS/third_party/composable_kernel_tiled/include \
-std=c++20 -O1 --offload-arch=gfx90a \
-fno-offload-lto -fno-gpu-rdc -fPIC \
-DCK_TILE_FMHA_FWD_FAST_EXP2=1 \
-fgpu-flush-denormals-to-zero \
-ftime-report
```

A packaged copy of the test file and this script is in the first comment / gist (see below).

## Clang 23.1.0-rc3 measurements (`/usr/bin/time -v` + `-ftime-report`)

| | |
|---|---|
| Wall | **10:49** |
| User CPU | 638.8 s |
| Max RSS | **2 425 284 KB (~2.4 GB)** |
| Exit | 0 |

Device `-cc1` (616.8 s wall):

| Phase | Wall | Share |
|---|---:|---:|
| **Machine code generation** | **454.8 s** | **73.7 %** |
| Optimizer | 119.0 s | 19.3 % |
| Front end | 36.6 s | 5.9 % |
| LLVM IR generation | 6.3 s | 1.0 % |

Top passes (device, 554 s pass timer):

| Pass | Wall | Share |
|---|---:|---:|
| **Machine Instruction Scheduler** | **182.6 s** | **32.9 %** |
| **Greedy Register Allocator #3** | **123.0 s** | **22.2 %** |
| Structurize control flow | 40.5 s | 7.3 % |
| LICMPass | 36.3 s | 6.5 % |
| InstCombinePass | 17.2 s | 3.1 % |
| AMDGPU DAG→DAG ISel | 17.1 s | 3.1 % |
| Loop Strength Reduction | 10.3 s | 1.9 % |
| InlinerPass | 8.6 s | 1.5 % |
| Live Interval Analysis | 8.2 s | 1.5 % |

Host x86_64 `-cc1` of the same TU is 32 s (almost all frontend). The cost is AMDGPU backend.

During the compile, samples of `clang-23 -cc1` were in `PreRARematStage::initGCNSchedStage` and `LiveRangeCalc::findReachingDefs`, which matches the scheduler / RA dominance.

## Why this sinks a real package build

xformers HIP/CK flash-attention is 643 ninja TUs. 624 of them are instances like this one (one template specialization each). The remaining ~12 are *dispatchers* that include the same CK FMHA headers and `BOOL_SWITCH` over many specializations; those sit at ninja `[6/643]` for much longer (~5 GB RSS, samples in `ModuleInlinerWrapperPass` / `LICMPass`) because ninja only increments after the whole `hipcc --offload-arch=A --offload-arch=B` command exits.

At 11 minutes × 624 instances × 2 archs / `MAX_JOBS=2` this is tens of hours, plus the dispatcher TUs. Official ROCm 7.1 / LLVM 20 wheels finish in the 6 hour CI budget.

The same host previously OOM-killed `-O3` + device LTO + `-amdgpu-early-inline-all` builds of the dispatcher TUs at 14–19 GB RSS. Those flags make it worse; they are not required to see the 11 minute / 2.4 GB instance compile.

## Related

- https://github.com/llvm/llvm-project/issues/86332 — `-amdgpu-early-inline-all` memory (CK batchnorm, Clang 17→18). **Different flag, still open.**
- https://github.com/llvm/llvm-project/issues/168849 — AMDGPU InstCombine compile-time.
- https://github.com/llvm/llvm-project/issues/166665 — `amdgpu-uniform-intrinsic-combine`.

## Notes on a Clang 20 A/B on this machine

Upstream `clang++` 20.1.8 cannot compile this TU against ROCm 7.14 HIP headers (`__builtin_elementwise_exp10` is unknown; device bitcode is LLVM 23). A clean local 20-vs-23 A/B needs ROCm 7.1 `amdclang++` (or HIP headers from that era). The official wheel job is that A/B.

Reproducer files + full `-ftime-report` log:
https://gist.github.com/berolinux/38656c50d32c69fe1c94f58ffa01ec1b

`-O0 --offload-device-only -emit-llvm` of this TU is 139 MB (1.2 MB xz). Happy to put that somewhere if a `.ll` + `llc`/`opt` split would help.

Start with the packaged instance.hip and script in the linked gist, reproducing the command with -ftime-report and profiling the device clang -cc1 process. Read the Machine Instruction Scheduler and Greedy Register Allocator timings, along with samples in PreRARematStage::initGCNSchedStage and LiveRangeCalc::findReachingDefs. Done means demonstrating a measured improvement in compile time or memory for this AMDGPU reproducer.

Written by the indexing model from the issue text.