Guard NEO's scratch allocation with a spill-triggered drain

[michel2323]

NEO allocates a queue's scratch buffer at the first submission of a kernel whose spill exceeds what is already allocated, and that allocation has no error path: prepareScratchAllocation() calls allocateGraphicsMemoryWithProperties() without a null check, and programSurfaceState() aborts the process (UNRECOVERABLE_IF, scratch_space_controller_xehp_and_later.cpp:78, byte-identical from the 25.18 LTS driver through current master) when it failed. Whether it fails depends on how much dead-but-unfinalized driver memory happens to be live at that one instant — a GC lottery. We have logged repeated aborts of exactly this shape on Aurora (PVC, LTS NEO 25.18) under launch storms from real workloads (ExaModels), where the process dies inside zeCommandQueueExecuteCommandLists with no Julia-side recourse.

This PR removes the lottery instead of the (driver-side) bug:

• ZeKernel caches its spillMemSize (seeded by properties(), one Int load on the launch path).
• Each queue tracks a high-water mark of per-thread spill among submitted kernels.
• Before the first submission that crosses the mark, the launch path retires in-flight work, flushes deferred releases (_run_reclaim_callbacks), and runs GC.gc(false) — so the one allocation the driver cannot recover from happens at the cleanest reachable moment.

The hedge fires once per (queue, scratch tier): one synchronize per workload in practice, and a cached-Int compare on the non-spilling fast path. Opt out with ONEAPI_SCRATCH_HEDGE=0 (the high-water mark is maintained regardless).

Tested with a deliberately spilling kernel (256 live accumulators): the hedge fires exactly once, never for no-spill kernels, and the knob-off path still tracks the mark. Suite green on Aurora LTS (PVC) and the kernel spills on both PVC and DG2.

This is deliberately independent of the immediate-command-list rework (#610): it guards a different resource (scratch) against a different garbage source (USM churn), and remains useful on current drivers — the missing null check is in NEO's shared code, not LTS-specific.

[codecov]

Codecov Report

✅ All modified and coverable lines are covered by tests.
✅ Project coverage is 80.97%. Comparing base (cf175bf) to head (2d2a1ca).
⚠️ Report is 2 commits behind head on main.

Additional details and impacted files

@@            Coverage Diff             @@
##             main     #609      +/-   ##
==========================================
+ Coverage   80.48%   80.97%   +0.48%     
==========================================
  Files          50       50              
  Lines        3490     3505      +15     
==========================================
+ Hits         2809     2838      +29     
+ Misses        681      667      -14     

☔ View full report in Codecov by Harness.
📢 Have feedback on the report? Share it here.

🚀 New features to boost your workflow:

• ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.

[github-actions]

Your PR requires formatting changes to meet the project's style guidelines.
Please consider running Runic (git runic main) to apply these changes.

Click here to view the suggested changes.

diff --git a/test/execution.jl b/test/execution.jl
index 8f11ae3..303c8c7 100644
--- a/test/execution.jl
+++ b/test/execution.jl
@@ -635,7 +635,7 @@ function hedge_spill_kernel(out, a, ::Val{N}, ::Val{R}) where {N, R}
     i = get_global_id()
     acc = hedge_rounds(ntuple(j -> a[i] + Float32(j), Val(N)), Val(R))
     s = 0.0f0
-    @inbounds for j = 1:N
+    @inbounds for j in 1:N
         s += acc[j]
     end
     @inbounds out[i] = s
@@ -647,7 +647,7 @@ end
     out = oneAPI.zeros(Float32, 256)
     synchronize()
 
-    k = @oneapi launch=false hedge_spill_kernel(out, a, Val(256), Val(2))
+    k = @oneapi launch = false hedge_spill_kernel(out, a, Val(256), Val(2))
     spill = oneL0.spill_mem_size(k.fun)
 
     if spill == 0
@@ -656,34 +656,38 @@ end
     else
         # queues are task-local, so a fresh task gets a fresh queue with a zero
         # high-water mark; observe there, assert on the test task
-        observed = fetch(@async begin
-            q = global_queue(context(), device())
-            hwm0 = q.scratch_hwm
-            c0 = oneL0.SCRATCH_HEDGE_COUNT[]
-            @oneapi items=64 groups=4 hedge_spill_kernel(out, a, Val(256), Val(2))
-            c1 = oneL0.SCRATCH_HEDGE_COUNT[]
-            # same spill tier: must not fire again
-            @oneapi items=64 groups=4 hedge_spill_kernel(out, a, Val(256), Val(2))
-            c2 = oneL0.SCRATCH_HEDGE_COUNT[]
-            # a storm of no-spill kernels must not fire the hedge either
-            for _ in 1:32
-                @oneapi dummy()
+        observed = fetch(
+            @async begin
+                q = global_queue(context(), device())
+                hwm0 = q.scratch_hwm
+                c0 = oneL0.SCRATCH_HEDGE_COUNT[]
+                @oneapi items = 64 groups = 4 hedge_spill_kernel(out, a, Val(256), Val(2))
+                c1 = oneL0.SCRATCH_HEDGE_COUNT[]
+                # same spill tier: must not fire again
+                @oneapi items = 64 groups = 4 hedge_spill_kernel(out, a, Val(256), Val(2))
+                c2 = oneL0.SCRATCH_HEDGE_COUNT[]
+                # a storm of no-spill kernels must not fire the hedge either
+                for _ in 1:32
+                    @oneapi dummy()
+                end
+                c3 = oneL0.SCRATCH_HEDGE_COUNT[]
+                (hwm0, q.scratch_hwm, c1 - c0, c2 - c1, c3 - c2)
             end
-            c3 = oneL0.SCRATCH_HEDGE_COUNT[]
-            (hwm0, q.scratch_hwm, c1 - c0, c2 - c1, c3 - c2)
-        end)
+        )
         @test observed == (0, spill, 1, 0, 0)
 
         # with the hedge disabled the high-water mark is still maintained
         old = oneL0.SCRATCH_HEDGE[]
         oneL0.SCRATCH_HEDGE[] = false
         try
-            observed = fetch(@async begin
-                q = global_queue(context(), device())
-                c0 = oneL0.SCRATCH_HEDGE_COUNT[]
-                @oneapi items=64 groups=4 hedge_spill_kernel(out, a, Val(256), Val(2))
-                (oneL0.SCRATCH_HEDGE_COUNT[] - c0, q.scratch_hwm)
-            end)
+            observed = fetch(
+                @async begin
+                    q = global_queue(context(), device())
+                    c0 = oneL0.SCRATCH_HEDGE_COUNT[]
+                    @oneapi items = 64 groups = 4 hedge_spill_kernel(out, a, Val(256), Val(2))
+                    (oneL0.SCRATCH_HEDGE_COUNT[] - c0, q.scratch_hwm)
+                end
+            )
             @test observed == (0, spill)
         finally
             oneL0.SCRATCH_HEDGE[] = old