intel / intel/intel-graphics-compiler

SIGSEGV in vISA when the GRF RA iteration budget is exhausted under fail-safe RA (dg2, bmg)

Open
#421 0 comments 0 reactions 0 assignees View on GitHub
Dominant language
C++
Stars
718
Forks
191
PR merge metrics
No merged PRs in 30d

Description

`ocloc` dies with SIGSEGV inside `libigc` when the GRF register allocator exhausts its iteration budget while fail-safe RA is enabled: `coloringRegAlloc()` returns `VISA_SUCCESS` with `G4_RegVar`s whose physical register is still null, and the first post-RA pass to touch one dereferences it.

Stock IGC, no local patches. Reproduces on IGC 2.36.3 (ocloc 26.22.38646.4, Arc A380) and 2.38.2 (ocloc 26.27.39122.11, Arc B570), `-device dg2`, `dg2-g11` and `bmg`.

```c
__attribute__((intel_reqd_sub_group_size(32)))
kernel void k(global const int *in, global int *out, global const uint *off) {
int gid = get_global_id(0);
bool p0 = in[gid + off[0]] > 0; bool p1 = in[gid + off[1]] > 1;
bool p2 = in[gid + off[2]] > 2; bool p3 = in[gid + off[3]] > 3;
bool p4 = in[gid + off[4]] > 4; bool p5 = in[gid + off[5]] > 5;
bool p6 = in[gid + off[6]] > 6; bool p7 = in[gid + off[7]] > 7;
int acc = 0;
if (p0) acc += 1; else acc -= 1; if (p1) acc += 2; else acc -= 2;
if (p2) acc += 3; else acc -= 3; if (p3) acc += 4; else acc -= 4;
if (p4) acc += 5; else acc -= 5; if (p5) acc += 6; else acc -= 6;
if (p6) acc += 7; else acc -= 7;
out[gid] = acc;
}
```

```
$ ocloc compile -file repro.cl -options "-cl-opt-disable -igc_opts 'VISAOptions=-maxRAIterations 1'" -device dg2
Segmentation fault
$ echo $?
139
```

Expected a build, or a failure naming the kernel that could not be allocated. Not a crash.

`-maxRAIterations 2` and above pass, so the exhausted budget is the trigger. Not specific to `-cl-opt-disable`: the same kernel extended to 32 predicates crashes identically with optimization enabled.

I have a fix and can open a PR if useful.

Contributor guide

Open the contributing guide

Research direction

Start by reproducing the crash with the provided repro.cl and ocloc command using -maxRAIterations 1. Trace coloringRegAlloc() and the first post-RA pass that handles unallocated G4_RegVars; done means the compiler builds successfully or reports a failure naming the affected kernel instead of dereferencing null state.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
compilers
Issue type
Bug
Difficulty
4/5
Estimated time
3-5 days
Activity status
Quiet
Clarity
Mostly clear
Newbie friendliness
55/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.