`make_llir` ~8-minute O(n²) stall after LLVM bump #9746 (BuiltinFuncToLLVM / simplifyRegions) on AMD backend
bug
### Describe the bug
## Summary
Compiling a fairly small reduction kernel for the AMD backend (`gfx1151`) regressed
from **~1.4 s to ~500 s** in the `make_llir` (ttgir → llir) stage. The compile
spins on a **single CPU core at ~99%** the entire time (no memory growth, no spill),
which points at an algorithmic (O(n²)) blow-up rather than I/O or thrashing.
I bisected this to the LLVM bump in #9746 (`cmake/llvm-hash.txt`
`278dba37d0ac…` → `7f77ca0dbda4…`).
## Reproducer
The kernel is an Inductor-generated fused reduction (two sequential RMSNorm-style
variance reductions guarded by `if pid < n0 / elif pid < n1`, i.e. **two top-level
`scf.if` regions** in one function). Feeding the cached `.ttgir` straight into
`triton.compile()` starts the pipeline at the ttgir stage and isolates `make_llir`:
```python
import os, time, triton
from triton.backends.compiler import GPUTarget
import triton.backends.amd.compiler as amd
TTGIR = "triton_red_fused_1.ttgir" # attached
# time make_llir specifically
_orig = amd.HIPBackend.make_llir
def _timed(src, metadata, options):
t0 = time.time()
try:
return _orig(src, metadata, options)
finally:
print(f"make_llir: {time.time()-t0:.1f}s", flush=True)
amd.HIPBackend.make_llir = staticmethod(_timed)
target = GPUTarget("hip", "gfx1151", 32)
options = {"num_warps": 1, "num_stages": 1, "num_ctas": 1, "waves_per_eu": 0,
"kpack": 1, "matrix_instr_nonkdim": 0, "enable_fp_fusion": True}
t0 = time.time()
triton.compile(TTGIR, target=target, options=options)
print(f"TOTAL compile from ttgir: {time.time()-t0:.1f}s")
```
## Observed behavior
| Triton commit | LLVM hash | `make_llir` |
|---|---|---|
| `b4977f1d0f` (#9773, parent of the bump) | `278dba37d0ac` | **1.4 s** |
| `ba295c0218` (#9746, the bump) | `7f77ca0dbda4` | **> 120 s (measured ~500 s)** |
| current `main` | `87717bf9…` | ~498 s |
- 99% of one core for the whole duration; no measurable memory growth.
- The same kernel/options that take ~500 s here took **< 2 s** before the bump.
## Suspected root cause
The cost shows up in `make_llir`, in the AMD `BuiltinFuncToLLVM` pass, which drives
the greedy pattern rewriter with region simplification set to
`GreedySimplifyRegionLevel::Aggressive`
(`third_party/amd/lib/TritonAMDGPUToLLVM/BuiltinFuncToLLVM.cpp`). Backtrace sits in
`simplifyRegions` -> `OperandStorage::eraseOperands`.
`triton_red_fused_1.ttgir`:
```
#blocked = #ttg.blocked<{sizePerThread = [1, 8], threadsPerWarp = [2, 16], warpsPerCTA = [1, 1], order = [1, 0]}>
#blocked1 = #ttg.blocked<{sizePerThread = [1, 1], threadsPerWarp = [32, 1], warpsPerCTA = [1, 1], order = [1, 0]}>
#blocked2 = #ttg.blocked<{sizePerThread = [1, 1], threadsPerWarp = [32, 1], warpsPerCTA = [1, 1], order = [0, 1]}>
module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 1 : i32, ttg.target = "hip:gfx1151", "ttg.threads-per-warp" = 32 : i32} {
tt.func public @triton_red_fused_1(%in_ptr0: !tt.ptr<f16> {tt.divisibility = 16 : i32}, %out_ptr0: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %out_ptr1: !tt.ptr<f32> {tt.divisibility = 16 : i32}, %xnumel_0: i32 {tt.divisibility = 16 : i32}, %xnumel_1: i32) attributes {noinline = false} {
%cst = arith.constant dense<0.000000e+00> : tensor<32x128xf32, #blocked>
%c31_i32 = arith.constant 31 : i32
%cst_0 = arith.constant dense<0.000000e+00> : tensor<32x128xf16, #blocked>
%c32_i32 = arith.constant 32 : i32
%cst_1 = arith.constant dense<32> : tensor<32x1xi32, #blocked>
%cst_2 = arith.constant dense<4096> : tensor<1x128xi32, #blocked>
%cst_3 = arith.constant dense<8> : tensor<32x1xi32, #blocked>
%cst_4 = arith.constant dense<6144> : tensor<32x1xi32, #blocked>
%cst_5 = arith.constant dense<128> : tensor<32x1xi32, #blocked>
%pid = tt.get_program_id x : i32
%num_xblocks_0 = arith.addi %xnumel_0, %c31_i32 : i32
%num_xblocks_0_6 = arith.divsi %num_xblocks_0, %c32_i32 : i32
%num_xblocks_1 = arith.addi %xnumel_1, %c31_i32 : i32
%num_xblocks_1_7 = arith.divsi %num_xblocks_1, %c32_i32 : i32
%num_xblocks_1_8 = arith.addi %num_xblocks_0_6, %num_xblocks_1_7 : i32
%0 = arith.cmpi slt, %pid, %num_xblocks_0_6 : i32
scf.if %0 {
%xoffset = arith.muli %pid, %c32_i32 : i32
%xindex = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked}>>
%xindex_9 = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked1}>>
%xindex_10 = tt.expand_dims %xindex {axis = 1 : i32} : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<32x1xi32, #blocked>
%xindex_11 = tt.expand_dims %xindex_9 {axis = 1 : i32} : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked1}>> -> tensor<32x1xi32, #blocked1>
%xindex_12 = tt.splat %xoffset : i32 -> tensor<32x1xi32, #blocked>
%xindex_13 = tt.splat %xoffset : i32 -> tensor<32x1xi32, #blocked1>
%xindex_14 = arith.addi %xindex_12, %xindex_10 : tensor<32x1xi32, #blocked>
%xindex_15 = arith.addi %xindex_13, %xindex_11 : tensor<32x1xi32, #blocked1>
%xmask = tt.splat %xnumel_0 : i32 -> tensor<32x1xi32, #blocked>
%xmask_16 = tt.splat %xnumel_0 : i32 -> tensor<32x1xi32, #blocked1>
%xmask_17 = arith.cmpi slt, %xindex_14, %xmask : tensor<32x1xi32, #blocked>
%xmask_18 = arith.cmpi slt, %xindex_15, %xmask_16 : tensor<32x1xi32, #blocked1>
%r0_base = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32, #ttg.slice<{dim = 0, parent = #blocked}>>
%r0_base_19 = tt.expand_dims %r0_base {axis = 0 : i32} : tensor<128xi32, #ttg.slice<{dim = 0, parent = #blocked}>> -> tensor<1x128xi32, #blocked>
%x0 = arith.remsi %xindex_14, %cst_1 : tensor<32x1xi32, #blocked>
%x1 = arith.divsi %xindex_14, %cst_1 : tensor<32x1xi32, #blocked>
%tmp0 = arith.muli %x0, %cst_5 : tensor<32x1xi32, #blocked>
%tmp0_20 = tt.broadcast %r0_base_19 : tensor<1x128xi32, #blocked> -> tensor<32x128xi32, #blocked>
%tmp0_21 = tt.broadcast %tmp0 : tensor<32x1xi32, #blocked> -> tensor<32x128xi32, #blocked>
%tmp0_22 = arith.addi %tmp0_20, %tmp0_21 : tensor<32x128xi32, #blocked>
%tmp0_23 = arith.muli %x1, %cst_4 : tensor<32x1xi32, #blocked>
%tmp0_24 = tt.broadcast %tmp0_23 : tensor<32x1xi32, #blocked> -> tensor<32x128xi32, #blocked>
%tmp0_25 = arith.addi %tmp0_22, %tmp0_24 : tensor<32x128xi32, #blocked>
%tmp0_26 = tt.splat %in_ptr0 : !tt.ptr<f16> -> tensor<32x128x!tt.ptr<f16>, #blocked>
%tmp0_27 = tt.addptr %tmp0_26, %tmp0_25 : tensor<32x128x!tt.ptr<f16>, #blocked>, tensor<32x128xi32, #blocked>
%tmp0_28 = tt.broadcast %xmask_17 : tensor<32x1xi1, #blocked> -> tensor<32x128xi1, #blocked>
%tmp0_29 = tt.load %tmp0_27, %tmp0_28, %cst_0 evictionPolicy = evict_first : tensor<32x128x!tt.ptr<f16>, #blocked>
%tmp0_30 = arith.extf %tmp0_29 : tensor<32x128xf16, #blocked> to tensor<32x128xf32, #blocked>
%tmp2 = arith.mulf %tmp0_30, %tmp0_30 : tensor<32x128xf32, #blocked>
%tmp5 = arith.addf %tmp2, %cst : tensor<32x128xf32, #blocked>
%_tmp4 = arith.select %tmp0_28, %tmp5, %cst : tensor<32x128xi1, #blocked>, tensor<32x128xf32, #blocked>
%tmp4 = "tt.reduce"(%_tmp4) <{axis = 1 : i32}> ({
^bb0(%tmp4_33: f32, %tmp4_34: f32):
%tmp4_35 = arith.addf %tmp4_33, %tmp4_34 : f32
tt.reduce.return %tmp4_35 : f32
}) : (tensor<32x128xf32, #blocked>) -> tensor<32xf32, #ttg.slice<{dim = 1, parent = #blocked}>>
%tmp4_31 = ttg.convert_layout %tmp4 : tensor<32xf32, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<32xf32, #ttg.slice<{dim = 1, parent = #blocked2}>>
%tmp4_32 = tt.expand_dims %tmp4_31 {axis = 1 : i32} : tensor<32xf32, #ttg.slice<{dim = 1, parent = #blocked2}>> -> tensor<32x1xf32, #blocked2>
%1 = tt.splat %out_ptr0 : !tt.ptr<f32> -> tensor<32x1x!tt.ptr<f32>, #blocked1>
%2 = tt.addptr %1, %xindex_15 : tensor<32x1x!tt.ptr<f32>, #blocked1>, tensor<32x1xi32, #blocked1>
%3 = ttg.convert_layout %2 : tensor<32x1x!tt.ptr<f32>, #blocked1> -> tensor<32x1x!tt.ptr<f32>, #blocked2>
%4 = ttg.convert_layout %xmask_18 : tensor<32x1xi1, #blocked1> -> tensor<32x1xi1, #blocked2>
tt.store %3, %tmp4_32, %4 : tensor<32x1x!tt.ptr<f32>, #blocked2>
} else {
%1 = arith.cmpi slt, %pid, %num_xblocks_1_8 : i32
scf.if %1 {
%pid_offset = arith.subi %pid, %num_xblocks_0_6 : i32
%xoffset = arith.muli %pid_offset, %c32_i32 : i32
%xindex = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked}>>
%xindex_9 = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked1}>>
%xindex_10 = tt.expand_dims %xindex {axis = 1 : i32} : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<32x1xi32, #blocked>
%xindex_11 = tt.expand_dims %xindex_9 {axis = 1 : i32} : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked1}>> -> tensor<32x1xi32, #blocked1>
%xindex_12 = tt.splat %xoffset : i32 -> tensor<32x1xi32, #blocked>
%xindex_13 = tt.splat %xoffset : i32 -> tensor<32x1xi32, #blocked1>
%xindex_14 = arith.addi %xindex_12, %xindex_10 : tensor<32x1xi32, #blocked>
%xindex_15 = arith.addi %xindex_13, %xindex_11 : tensor<32x1xi32, #blocked1>
%xmask = tt.splat %xnumel_1 : i32 -> tensor<32x1xi32, #blocked>
%xmask_16 = tt.splat %xnumel_1 : i32 -> tensor<32x1xi32, #blocked1>
%xmask_17 = arith.cmpi slt, %xindex_14, %xmask : tensor<32x1xi32, #blocked>
%xmask_18 = arith.cmpi slt, %xindex_15, %xmask_16 : tensor<32x1xi32, #blocked1>
%r0_base = tt.make_range {end = 128 : i32, start = 0 : i32} : tensor<128xi32, #ttg.slice<{dim = 0, parent = #blocked}>>
%r0_base_19 = tt.expand_dims %r0_base {axis = 0 : i32} : tensor<128xi32, #ttg.slice<{dim = 0, parent = #blocked}>> -> tensor<1x128xi32, #blocked>
%x4 = arith.remsi %xindex_14, %cst_3 : tensor<32x1xi32, #blocked>
%x5 = arith.divsi %xindex_14, %cst_3 : tensor<32x1xi32, #blocked>
%tmp6 = arith.addi %r0_base_19, %cst_2 : tensor<1x128xi32, #blocked>
%tmp6_20 = arith.muli %x4, %cst_5 : tensor<32x1xi32, #blocked>
%tmp6_21 = tt.broadcast %tmp6 : tensor<1x128xi32, #blocked> -> tensor<32x128xi32, #blocked>
%tmp6_22 = tt.broadcast %tmp6_20 : tensor<32x1xi32, #blocked> -> tensor<32x128xi32, #blocked>
%tmp6_23 = arith.addi %tmp6_21, %tmp6_22 : tensor<32x128xi32, #blocked>
%tmp6_24 = arith.muli %x5, %cst_4 : tensor<32x1xi32, #blocked>
%tmp6_25 = tt.broadcast %tmp6_24 : tensor<32x1xi32, #blocked> -> tensor<32x128xi32, #blocked>
%tmp6_26 = arith.addi %tmp6_23, %tmp6_25 : tensor<32x128xi32, #blocked>
%tmp6_27 = tt.splat %in_ptr0 : !tt.ptr<f16> -> tensor<32x128x!tt.ptr<f16>, #blocked>
%tmp6_28 = tt.addptr %tmp6_27, %tmp6_26 : tensor<32x128x!tt.ptr<f16>, #blocked>, tensor<32x128xi32, #blocked>
%tmp6_29 = tt.broadcast %xmask_17 : tensor<32x1xi1, #blocked> -> tensor<32x128xi1, #blocked>
%tmp6_30 = tt.load %tmp6_28, %tmp6_29, %cst_0 evictionPolicy = evict_first : tensor<32x128x!tt.ptr<f16>, #blocked>
%tmp6_31 = arith.extf %tmp6_30 : tensor<32x128xf16, #blocked> to tensor<32x128xf32, #blocked>
%tmp8 = arith.mulf %tmp6_31, %tmp6_31 : tensor<32x128xf32, #blocked>
%tmp11 = arith.addf %tmp8, %cst : tensor<32x128xf32, #blocked>
%_tmp10 = arith.select %tmp6_29, %tmp11, %cst : tensor<32x128xi1, #blocked>, tensor<32x128xf32, #blocked>
%tmp10 = "tt.reduce"(%_tmp10) <{axis = 1 : i32}> ({
^bb0(%tmp10_34: f32, %tmp10_35: f32):
%tmp10_36 = arith.addf %tmp10_34, %tmp10_35 : f32
tt.reduce.return %tmp10_36 : f32
}) : (tensor<32x128xf32, #blocked>) -> tensor<32xf32, #ttg.slice<{dim = 1, parent = #blocked}>>
%tmp10_32 = ttg.convert_layout %tmp10 : tensor<32xf32, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<32xf32, #ttg.slice<{dim = 1, parent = #blocked2}>>
%tmp10_33 = tt.expand_dims %tmp10_32 {axis = 1 : i32} : tensor<32xf32, #ttg.slice<{dim = 1, parent = #blocked2}>> -> tensor<32x1xf32, #blocked2>
%2 = tt.splat %out_ptr1 : !tt.ptr<f32> -> tensor<32x1x!tt.ptr<f32>, #blocked1>
%3 = tt.addptr %2, %xindex_15 : tensor<32x1x!tt.ptr<f32>, #blocked1>, tensor<32x1xi32, #blocked1>
%4 = ttg.convert_layout %3 : tensor<32x1x!tt.ptr<f32>, #blocked1> -> tensor<32x1x!tt.ptr<f32>, #blocked2>
%5 = ttg.convert_layout %xmask_18 : tensor<32x1xi1, #blocked1> -> tensor<32x1xi1, #blocked2>
tt.store %4, %tmp10_33, %5 : tensor<32x1x!tt.ptr<f32>, #blocked2>
}
}
tt.return
}
}
```
### Environment details
- Triton: built from `main` at 3b5446d4a8
- Backend / target: `hip`, `gfx1151` (RDNA 3.5), `threads-per-warp = 32`
- Host: Linux x86_64
0 条评论