vISA: report GRF RA failure when the iteration budget leaves operands unallocated

IGC/ocloc_tests/ra-iteration-limit-complete-failsafe-result.cl

@@ -0,0 +1,71 @@
+/*========================== begin_copyright_notice ============================
+
+Copyright (C) 2026 Intel Corporation
+
+SPDX-License-Identifier: MIT
+
+============================= end_copyright_notice ===========================*/
+
+// REQUIRES: regkeys, dg2-supported
+// RUN: ocloc compile -file %s -options "-cl-opt-disable -igc_opts 'VISAOptions=-maxRAIterations 2 -asmToConsole'" -device dg2 | FileCheck %s
+
+// Companion to ra-iteration-limit-graceful-failure.cl, which covers
+// -maxRAIterations 1. Both invocations end the GRF RA loop the same way:
+// GlobalRA::setupFailSafeIfNeeded() enables fail-safe RA for iteration
+// maxRAIterations - 1, that iteration still has to insert spill code, and
+// iterationNo reaches maxRAIterations with no iteration left. The two cases
+// differ in what that leaves behind.
+//
+// With -maxRAIterations 1 the fail-safe iteration is iteration 0 and it leaves
+// operands whose G4_RegVar has a null physical register, so the post-RA passes
+// dereference null. That must be reported as an RA failure.
+//
+// With -maxRAIterations 2 the fail-safe iteration is iteration 1 and it
+// allocates everything: every operand has a physical register and the kernel
+// is complete and correct. Fail-safe RA is designed to produce exactly that -
+// it assigns the spill/fill temporaries it creates out of the reserved GRFs on
+// the spot, which is why it needs no follow-up iteration. Running the
+// iteration counter out is not by itself evidence that anything is wrong, so
+// this result must be kept rather than thrown away.
+//
+// If the out-of-iterations report is raised unconditionally here,
+// coloringRegAlloc() returns VISA_SPILL, Optimizer::optimization() bails out,
+// and CEncoder::Compile() records SIMD_SKIP_SPILL and produces no kernel.
+// intel_reqd_sub_group_size(32) rules out SIMD16 and SIMD8, so every attempt
+// fails the same way and IGC emits a program containing no kernel at all
+// (a zebin with no .text.k, no .symtab and no .ze_info) while ocloc still
+// prints "Build succeeded". -asmToConsole makes that visible: a compile that
+// emits the kernel dumps its assembly, a compile that discarded it dumps
+// nothing.
+//
+// Test setup is identical to ra-iteration-limit-graceful-failure.cl:
+// 8 simultaneously live predicates is the point at which the allocator first
+// has to spill on this target, -cl-opt-disable keeps them from being optimized
+// away, and intel_reqd_sub_group_size(32) pins SIMD32 so the pressure is not
+// relieved by a narrower dispatch.
+
+// CHECK: //.kernel k
+// CHECK: Build succeeded
+
+__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;
+  if (p7) acc += 8; else acc -= 8;
+  out[gid] = acc;
+}

IGC/ocloc_tests/ra-iteration-limit-graceful-failure.cl

@@ -0,0 +1,65 @@
+/*========================== begin_copyright_notice ============================
+
+Copyright (C) 2026 Intel Corporation
+
+SPDX-License-Identifier: MIT
+
+============================= end_copyright_notice ===========================*/
+
+// REQUIRES: regkeys, dg2-supported
+// RUN: ocloc compile -file %s -options "-cl-opt-disable -igc_opts 'VISAOptions=-maxRAIterations 1'" -device dg2 | FileCheck %s
+
+// This test verifies that the GRF register allocator reports failure, rather
+// than returning success with unallocated variables, when it runs out of RA
+// iterations while fail-safe RA is enabled.
+//
+// GlobalRA::setupFailSafeIfNeeded() turns fail-safe RA on when
+// getIterNo() == maxRAIterations - 1. With -maxRAIterations 1 that is
+// iteration 0, so fail-safe RA is enabled on the very first iteration. If that
+// iteration still has to insert spill code it creates new spill/fill
+// temporaries and bumps iterationNo to maxRAIterations, and there is no further
+// iteration in which to allocate them. The post-loop failure check in
+// coloringRegAlloc() used to be guarded by !reserveSpillReg, so it skipped the
+// report, returned VISA_SUCCESS, and left G4_RegVars with a null physical
+// register. Post-RA passes then dereferenced it: SWSB's
+// G4_BB_SB::getFootprintForOperand() and the local scheduler's
+// DDD::getBucketsForOperand() both do
+//   phyReg = base->asRegVar()->getPhyReg(); switch (phyReg->getKind())
+// with no null check, so the compiler segfaulted.
+//
+// With the failure correctly reported, RA returns VISA_SPILL, Optimizer
+// bails out before those passes, and IGC's retry logic recompiles the kernel
+// successfully. So the expected result is a clean successful build; before the
+// fix this invocation died with SIGSEGV.
+//
+// Test setup:
+// - 8 simultaneously live predicates is the point at which flag/GRF pressure
+//   first forces the allocator to spill on this target; 7 does not trigger it.
+// - -cl-opt-disable keeps the predicates from being optimized away.
+// - intel_reqd_sub_group_size(32) pins SIMD32 so pressure is not reduced by
+//   selecting a narrower dispatch.
+
+// CHECK: Build succeeded
+
+__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;
+  if (p7) acc += 8; else acc -= 8;
+  out[gid] = acc;
+}

visa/GraphColor.cpp

@@ -12240,6 +12240,22 @@ bool GlobalRA::kernelUsesDpas() const {
   return false;
 }
 
+// True if any operand still refers to a G4_RegVar with no physical register,
+// i.e. register allocation is incomplete. Post-RA passes dereference
+// G4_RegVar::getPhyReg() without a null check (G4_BB_SB::getFootprintForOperand,
+// DDD::getBucketsForOperand), so such a result must not be reported as success.
+static bool hasUnallocatedOperand(G4_Kernel &kernel) {
+  for (G4_BB *bb : kernel.fg)
+    for (G4_INST *inst : *bb)
+      for (int i = 0; i < Opnd_total_num; ++i) {
+        G4_Operand *opnd = inst->getOperand(Gen4_Operand_Number(i));
+        G4_VarBase *base = opnd ? opnd->getBase() : nullptr;
+        if (base && base->isRegVar() && !base->asRegVar()->getPhyReg())
+          return true;
+      }
+  return false;
+}
+
 //
 // graph coloring entry point.  returns nonzero if RA fails
 //
@@ -12690,9 +12706,17 @@ int GlobalRA::coloringRegAlloc() {
   // Report failure to allocate due to excessive register pressure.
   //
   // Failed to spill, or there is no stack call and the loop iterates to
-  // maxRAIterations
-  if (!reserveSpillReg &&
-      (failedToSpill || (!hasStackCall && iterationNo == maxRAIterations))) {
+  // maxRAIterations. The latter is normally benign under fail-safe RA
+  // (reserveSpillReg), which assigns the spill/fill temporaries it creates
+  // out of the reserved GRFs - but entered on the very first iteration
+  // (-maxRAIterations 1) it can itself leave operands unallocated, and
+  // returning success would make the post-RA passes dereference a null
+  // physical register. So report the exhaustion under fail-safe RA when, and
+  // only when, the allocation really is incomplete; a complete fail-safe
+  // result is a usable kernel and must not be discarded.
+  if ((!reserveSpillReg && failedToSpill) ||
+      (!hasStackCall && iterationNo == maxRAIterations &&
+       (!reserveSpillReg || hasUnallocatedOperand(kernel)))) {
     std::stringstream spilledVars;
     for (auto dcl : kernel.Declares) {
       if (dcl->isSpilled() && dcl->getRegFile() == G4_GRF) {