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_RegVars 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.
__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.
oclocdies with SIGSEGV insidelibigcwhen the GRF register allocator exhausts its iteration budget while fail-safe RA is enabled:coloringRegAlloc()returnsVISA_SUCCESSwithG4_RegVars 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-g11andbmg.Expected a build, or a failure naming the kernel that could not be allocated. Not a crash.
-maxRAIterations 2and 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.