Make host USM resident before kernel use

[michel2323]

Summary

Host USM allocations (HostBuffer) and the host pattern buffer used by fill! were never made resident on the device, unlike DeviceBuffer and SharedBuffer allocations (which already call make_resident). A GPU kernel that reads a non-resident host buffer can take a NotPresent pagefault under GC/allocation churn, even though host USM is nominally device-accessible.

This makes host USM resident at allocation time (src/pool.jl) and for fill!'s pattern buffer (src/array.jl), removing the asymmetry with device/shared USM.

Notes

Extracted as a standalone, stack-independent fix from the Aurora LTS work; the residency omission is a real asymmetry on main regardless of runtime.

Reproducer

A kernel reads a HostBuffer-backed oneVector under GC churn, then the array is dropped and freed. On the Aurora LTS NEO stack (where this was isolated) the read intermittently takes a NotPresent read pagefault that bans the kernel context — every later submission then fails with ZE_RESULT_ERROR_UNKNOWN. With this PR all modes pass. The noread mode (the GPU never touches the buffer) is the control and passes either way, isolating the fault to the read of non-resident host USM rather than its allocation or free.

using oneAPI
using oneAPI: oneL0

readk(A, B) = (i = get_global_id(); @inbounds B[i] = A[1]; return)

const MODE = get(ENV, "ONEAPI_MODE", "read")   # "read" or "noread"
const DST  = oneAPI.zeros(Int, 1); oneAPI.synchronize()
const N    = parse(Int, get(ENV, "N", "4000"))

for k in 1:N
    c = oneVector{Int,oneL0.HostBuffer}([k])
    MODE == "read" && @oneapi items=1 readk(c, DST)   # GPU reads the host buffer
    c = nothing                                        # drop; finalizer frees host USM
    k % 16  == 0 && GC.gc(true)
    if k % 200 == 0
        try
            oneAPI.synchronize()
        catch err
            println("FAILED at $k: $(typeof(err))"); exit(1)
        end
    end
end
oneAPI.synchronize(); println("PASSED")

ONEAPI_MODE=read   N=4000 julia --project repro.jl   # faults unpatched, passes with this PR
ONEAPI_MODE=noread N=4000 julia --project repro.jl   # control: passes either way

[codecov]

Codecov Report

✅ All modified and coverable lines are covered by tests.
✅ Project coverage is 80.90%. Comparing base (e995a63) to head (aec725f).
⚠️ Report is 1 commits behind head on main.

Additional details and impacted files

@@            Coverage Diff             @@
##             main     #585      +/-   ##
==========================================
- Coverage   80.92%   80.90%   -0.02%     
==========================================
  Files          48       48              
  Lines        3234     3237       +3     
==========================================
+ Hits         2617     2619       +2     
- Misses        617      618       +1     

☔ 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/src/array.jl b/src/array.jl
index a6bb9b3..4660c99 100644
--- a/src/array.jl
+++ b/src/array.jl
@@ -526,12 +526,12 @@ function Base.fill!(A::oneDenseArray{T}, val) where T
   val = convert(T, val)
   sizeof(T) == 0 && return A
 
-  # execute! is async, so we need to allocate the pattern in USM memory and keep it alive
-  # until the operation completes. The fill reads this host buffer on the GPU, so it must
-  # be made resident on the device like any other USM a kernel reads (see
-  # `allocate(::Type{oneL0.HostBuffer}, ...)`).
+    # execute! is async, so we need to allocate the pattern in USM memory and keep it alive
+    # until the operation completes. The fill reads this host buffer on the GPU, so it must
+    # be made resident on the device like any other USM a kernel reads (see
+    # `allocate(::Type{oneL0.HostBuffer}, ...)`).
   buf = oneL0.host_alloc(context(A), sizeof(T), Base.datatype_alignment(T))
-  oneL0.make_resident(context(A), device(), buf)
+    oneL0.make_resident(context(A), device(), buf)
   unsafe_store!(convert(Ptr{T}, buf), val)
   unsafe_fill!(context(A), device(), pointer(A), convert(ZePtr{T}, buf), length(A))
   synchronize(global_queue(context(A), device()))