`LLVM ERROR: failed to parse IR: multiple definition of local value named '.pn…'` when a Triton kernel is unrolled with a large `tl.static_range` (Arc Pro B70 / Battlemage)

[TSUMUGI-XE]

A Triton kernel that keeps several tables in registers and updates them inside a
tl.static_range loop stops compiling once the unroll count gets large enough. The program is
not rejected — the compiler aborts reading back IR it produced itself:

LLVM ERROR: failed to parse IR: multiple definition of local value named
'.pn687.pn.pn.pn … .pn'          (1,026 characters)

The abort is not catchable, so the host process dies rather than falling back.

The trigger is a name length

.pn is LLVM's phi-node suffix, appended each time a transformation rewrites the value. For
unroll count K_MAX the colliding name is exactly ".pn" + str(2*K_MAX+5) + ".pn"*(K_MAX-1),
which holds for every failing run. So the boundary tracks the length, not the count:

K_MAX   name length   result
 320        963       compiles
 340      1,023       compiles      <- longest name that compiles
 341      1,026       fails
 352      1,059       fails
 512      1,540       fails

LLVM's default -non-global-value-max-name-size is 1,024, which sits in that gap. Not that the
limit is wrong — it presumably exists to stop this kind of growth — only that the boundary
straddles it; whether it is what produces the collision we cannot see from outside.

The unroll count alone is not the trigger: a stripped-down kernel with one register-resident
table and one reduction per iteration compiles fine at K_MAX=512. What matters is how many
times the value gets rewritten, since that is what lengthens the name.

This is not unique to IGC. JuliaGPU/Metal.jl#655 reports the same failure class on the Metal
backend — a different accumulating suffix (.sroa.0 instead of .pn), names running to hundreds
of characters, duplicate local names, and the same parse failure. It was closed as an upstream
LLVM issue. I am reporting it here anyway because the part that is reachable from IGC is not the
name growth but the round-trip: whatever step re-reads the printed IR is the thing that turns a
printed-name collision into a hard abort, and that step is IGC's.

Reproducer

The reproducer is inlined at the end of this issue — torch and triton only, no vLLM, ~40 lines: an LRU cache
update with three register-resident tables (512 / 256 / 256), two reductions and masked updates
per iteration, under a data-dependent if. K_MAX is the only variable; each value ran as a
separate process, since the abort ends the process.

Environment

All versions are inside the container where the crash occurs; torch 2.13.0+xpu,
triton 3.7.2, Intel Arc Pro B70 (Battlemage G31, 8086:e223), kernel 7.2.0 (xe).

Three userspace stacks, same boundary and same colliding name:

IGC 2.40.13 + compute runtime 26.31.39395.13
IGC 2.41.5  + compute runtime 26.31.39395.13
IGC 2.41.5  + compute runtime 26.35.39758.10 + gmmlib 22.10.0
   ^ the matched set named by the 26.35.39758.10 release notes

The runtime dlopen()s IGC and the same soname ships at two prefixes, so the reproducer prints
which IGC the failing process mapped (LD_DEBUG=libs + /proc/self/maps); all three resolved
/usr/local/lib/libigc.so.2.<version>. Tags v2.41.6..v2.41.9 have no packages; their diff is
four commits (LLVM 22 default, MemOpt alignment) that do not touch naming, so I have not built
them.

Where the depth comes from: K_MAX is 16x the per-call token budget of a VRAM expert cache in
vLLM, so 8 concurrent sequences compile (K_MAX=128, 387 characters) and 32 abort during
CUDA-graph capture (K_MAX=512, 1,540).

What would help

The invariant rather than the limit: IR the compiler emitted should be readable by the compiler.
Bounding the .pn accumulation, or not depending on printed local names being unique where the
IR is read back, would each do it — which is right we cannot judge from here. A diagnostic
naming the kernel and the unroll depth, instead of LLVM ERROR plus abort, would also have made
this much cheaper to find.

igc_static_range_pn.py (reproducer)

#!/usr/bin/env python3
"""Minimal reproducer: IGC fails to parse its own IR when a Triton kernel with
register-resident tables is unrolled with a large `tl.static_range`.

  LLVM ERROR: failed to parse IR: multiple definition of local value
              named '.pn1029.pn.pn.pn.pn.pn.pn...'   lineno: 137807
  PLEASE submit a bug report to https://software.intel.com/en-us/support/priority-support

`.pn` is LLVM's phi-node name suffix; each transformation appends another one.
Past a certain unroll depth two distinct values end up with the same printed
name, and IGC's round-trip through textual IR then fails to re-parse.

The kernel below is an LRU cache update: three tables are held in registers,
and each unrolled iteration does two reductions plus masked updates on all
three.  Only UNROLL (`K_MAX`) changes between the passing and failing runs --
every other constexpr is held fixed.

Depends on torch + triton only.
Usage:  python3 igc_static_range_pn.py [unroll ...]     (default: 128 256 512)
"""
import sys
import torch
import triton
import triton.language as tl


@triton.jit
def _lru_update(
    ids_ptr, map_ptr, slot_expert_ptr, lastuse_ptr, step_ptr,
    out_slots_ptr, fill_src_ptr, fill_dst_ptr,
    n_ids,
    E: tl.constexpr, E_P2: tl.constexpr,
    C: tl.constexpr, C_P2: tl.constexpr, K_MAX: tl.constexpr,
):
    eoffs = tl.arange(0, E_P2)
    coffs = tl.arange(0, C_P2)
    emask = eoffs < E
    cmask = coffs < C
    map_v = tl.load(map_ptr + eoffs, mask=emask, other=-1).to(tl.int32)
    se_v = tl.load(slot_expert_ptr + coffs, mask=cmask, other=-1).to(tl.int32)
    lu_v = tl.load(lastuse_ptr + coffs, mask=cmask, other=2147483647).to(tl.int32)
    step = tl.load(step_ptr).to(tl.int32)

    for i in tl.static_range(K_MAX):
        if i < n_ids:
            e = tl.load(ids_ptr + i).to(tl.int32)
            is_e = eoffs == e
            slot = tl.sum(tl.where(is_e, map_v, 0), axis=0).to(tl.int32)
            miss = slot < 0
            victim = tl.argmin(lu_v, axis=0).to(tl.int32)
            old = tl.sum(tl.where(coffs == victim, se_v, 0), axis=0).to(tl.int32)
            chosen = tl.where(miss, victim, slot).to(tl.int32)
            map_v = tl.where((eoffs == old) & miss & (old >= 0), -1, map_v)
            map_v = tl.where(is_e, chosen, map_v)
            se_v = tl.where((coffs == chosen) & miss, e, se_v)
            lu_v = tl.where(coffs == chosen, step, lu_v)
            tl.store(out_slots_ptr + i, chosen)
            tl.store(fill_src_ptr + i, tl.where(miss, e, -1))
            tl.store(fill_dst_ptr + i, tl.where(miss, chosen, -1))
        else:
            tl.store(fill_dst_ptr + i, -1)

    tl.store(map_ptr + eoffs, map_v, mask=emask)
    tl.store(slot_expert_ptr + coffs, se_v, mask=cmask)
    tl.store(lastuse_ptr + coffs, lu_v, mask=cmask)
    tl.store(step_ptr, step + 1)


# The shapes the real workload uses: 512 experts, a 160-slot cache.
E, C = 512, 160
E_P2 = triton.next_power_of_2(E)      # 512
C_P2 = triton.next_power_of_2(C)      # 256


def try_unroll(k_max: int) -> str:
    dev = "xpu"
    ids = torch.zeros(k_max, dtype=torch.int32, device=dev)
    mp = torch.full((E_P2,), -1, dtype=torch.int32, device=dev)
    se = torch.full((C_P2,), -1, dtype=torch.int32, device=dev)
    lu = torch.zeros(C_P2, dtype=torch.int32, device=dev)
    st = torch.zeros(1, dtype=torch.int32, device=dev)
    o1 = torch.zeros(k_max, dtype=torch.int32, device=dev)
    o2 = torch.zeros(k_max, dtype=torch.int32, device=dev)
    o3 = torch.zeros(k_max, dtype=torch.int32, device=dev)
    try:
        _lru_update[(1,)](ids, mp, se, lu, st, o1, o2, o3, 8,
                          E=E, E_P2=E_P2, C=C, C_P2=C_P2, K_MAX=k_max)
        torch.xpu.synchronize()
        return "OK"
    except Exception as e:                       # noqa: BLE001
        return f"FAIL {type(e).__name__}: {str(e)[:200]}"


def loaded_igc() -> list[str]:
    """Which libigc/libigdfcl this very process has mapped.

    The compute runtime dlopen()s IGC, so `ldd` does not show it, and the same
    soname is shipped by two packages at two prefixes -- Intel's deb installs
    under /usr/local/lib and does not replace the distro one under
    /usr/lib/x86_64-linux-gnu.  Read it from the process instead of assuming.
    """
    out = []
    try:
        for line in open("/proc/self/maps"):
            if "libigc" in line or "libigdfcl" in line:
                path = line.split()[-1]
                if path not in out:
                    out.append(path)
    except OSError:
        pass
    return out


if __name__ == "__main__":
    print("torch  ", torch.__version__, flush=True)
    print("triton ", triton.__version__, flush=True)
    print("device ", torch.xpu.get_device_name(0), flush=True)
    # Warm the compiler up on a trivial kernel so IGC is mapped, then say which
    # IGC this process is actually using -- printed BEFORE the failing compile,
    # because a hard abort would take the evidence with it.
    try_unroll(8)
    for _p in loaded_igc():
        print("IGC in use:", _p, flush=True)
    print(f"shapes  E={E} E_P2={E_P2} C={C} C_P2={C_P2}", flush=True)
    for u in ([int(a) for a in sys.argv[1:]] or [128, 256, 512]):
        # A hard IGC crash aborts the process, so announce before compiling:
        # the last "..." line with no result is the one that died.
        print(f"--- K_MAX={u} ...", flush=True)
        print(f"--- K_MAX={u} -> {try_unroll(u)}", flush=True)

[obrotowy]

It is a known issue in LLVM https://github.com/llvm/llvm-project/blob/ca7933e47d3a3451d81e72ac174dcb5aa28b59d1/llvm/lib/AsmParser/LLParser.cpp#L3766:L3768

However, either reducing name lengths and increasing -non-global-value-max-name-size could be done on Triton side, not IGC one, as IGC handles SPIR-Vs with such long names correctly. In the reported case, the error gets emitted before the SPIR-V for IGC. I've prepared a patch for triton, which allows to manipulate -non-global-value-max-name-size= by setting appropriate env flag, and confirmed that with this flag set to -1 (infinity), your reproducer worked fine. Feel free to submit an issue for Triton.

triton.patch