diff --git a/IGC/ocloc_tests/ra-iteration-limit-complete-failsafe-result.cl b/IGC/ocloc_tests/ra-iteration-limit-complete-failsafe-result.cl new file mode 100644 index 000000000000..a5d5556e048b --- /dev/null +++ b/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; +} diff --git a/IGC/ocloc_tests/ra-iteration-limit-graceful-failure.cl b/IGC/ocloc_tests/ra-iteration-limit-graceful-failure.cl new file mode 100644 index 000000000000..d8550c4d7bb5 --- /dev/null +++ b/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; +} diff --git a/visa/GraphColor.cpp b/visa/GraphColor.cpp index 7c4b24e588d8..f7782896bbe5 100644 --- a/visa/GraphColor.cpp +++ b/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) {