Skip to content

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

Description

@pvelesko

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.

Activity

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Metadata

Metadata

Assignees

No one assigned

    Labels

    No labels
    No labels

    Type

    No type

    Projects

    No projects

      Milestone

      No milestone

      Relationships

      None yet

      Development

      No branches or pull requests

      Issue actions