From dc54dc9c8828316d547da06c7642f852b3aaf319 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Sat, 22 Aug 2026 15:32:34 -0300
Subject: [PATCH 1/9] Only list KernelAbstractions in versioninfo when it is
 loaded

In preparation for KernelAbstractions becoming a weak dependency.
---
 src/utils.jl | 14 +++++++++++---
 1 file changed, 11 insertions(+), 3 deletions(-)

diff --git a/src/utils.jl b/src/utils.jl
index ddf4b1f9..03d4de12 100644
--- a/src/utils.jl
+++ b/src/utils.jl
@@ -27,11 +27,19 @@ function versioninfo(io::IO=stdout)
     println(io, "- LLVM: $(LLVM.version())")
     println(io)
 
+
+    get_module(name::Symbol) = (name, getfield(oneAPI, name))
+    function get_module(pkg::Tuple{String, String})
+        id = Base.PkgId(Base.UUID(pkg[1]), pkg[2])
+        (pkg[2], get(Base.loaded_modules, id, nothing))
+    end
+
     println(io, "Julia packages:")
     println(io, "- oneAPI.jl: $(Base.pkgversion(oneAPI))")
-    for name in [:GPUArrays, :GPUCompiler, :KernelAbstractions, :LLVM, :SPIRVIntrinsics]
-        mod = getfield(oneAPI, name)
-        println(io, "- $(name): $(Base.pkgversion(mod))")
+    for pkg in [:GPUArrays, :GPUCompiler, ("63c18a36-062a-441e-b654-da1e3ab1ce7c", "KernelAbstractions"),
+                 :LLVM, :SPIRVIntrinsics]
+        name, mod = get_module(pkg)
+        isnothing(mod) || println(io, "- $(name): $(Base.pkgversion(mod))")
     end
     println(io)
 

From b58019e3f54b6333db898b6365227bcecce5808b Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard@gmail.com>
Date: Mon, 28 Sep 2026 10:12:28 +0200
Subject: [PATCH 2/9] Cache a kernel's maximum group size

Like the spill size, cache the kernel's `maxGroupSize` in the `ZeKernel`, so
that `launch_configuration` (and the KernelInterface launch validation, which
checks every explicit work-group size against it) don't query the kernel
properties, which allocates, on every launch.
---
 lib/level-zero/module.jl  | 15 ++++++++++++++-
 src/compiler/execution.jl |  8 ++++----
 2 files changed, 18 insertions(+), 5 deletions(-)

diff --git a/lib/level-zero/module.jl b/lib/level-zero/module.jl
index 91cc0dfc..6e0635df 100644
--- a/lib/level-zero/module.jl
+++ b/lib/level-zero/module.jl
@@ -85,13 +85,17 @@ mutable struct ZeKernel
     # Read on every launch by the scratch hedge, so it must not cost an API call.
     spill::Int
 
+    # cached maxGroupSize, seeded by `properties`; -1 while unqueried, and 0 without the
+    # MAX_GROUP_SIZE extension. Read on every KernelInterface launch.
+    max_group_size::Int
+
     function ZeKernel(mod, name)
         GC.@preserve name begin
             desc_ref = Ref(ze_kernel_desc_t(; pKernelName=pointer(name)))
             handle_ref = Ref{ze_kernel_handle_t}()
             zeKernelCreate(mod, desc_ref, handle_ref)
         end
-        obj = new(mod, handle_ref[], ReentrantLock(), -1)
+        obj = new(mod, handle_ref[], ReentrantLock(), -1, -1)
 
         finalizer(obj) do obj
             zeKernelDestroy(obj)
@@ -268,6 +272,8 @@ function properties(kernel::ZeKernel)
 
     props = props_ref[]
     kernel.spill = Int(props.spillMemSize)
+    kernel.max_group_size = max_group_size_props_ref === nothing ? 0 :
+        Int(max_group_size_props_ref[].maxGroupSize)
     return (
         numKernelArgs=Int(props.numKernelArgs),
         requiredGroupSize=ZeDim3(props.requiredGroupSizeX,
@@ -295,6 +301,13 @@ function spill_mem_size(kernel::ZeKernel)
     return s >= 0 ? s : Int(properties(kernel).spillMemSize)
 end
 
+# Cached access to a kernel's maxGroupSize, `missing` without the MAX_GROUP_SIZE extension.
+function max_group_size(kernel::ZeKernel)
+    s = kernel.max_group_size
+    s < 0 && (s = coalesce(properties(kernel).maxGroupSize, 0))
+    return s > 0 ? s : missing
+end
+
 
 ## execution
 
diff --git a/src/compiler/execution.jl b/src/compiler/execution.jl
index 20742c2d..2d4ae759 100644
--- a/src/compiler/execution.jl
+++ b/src/compiler/execution.jl
@@ -241,9 +241,9 @@ function launch_configuration(kernel::HostKernel{F,TT}) where {F,TT}
     # configurations, so roll our own version that behaves like CUDA's
     # occupancy API and assumes the kernel still does bounds checking.
 
-    kernel_props = oneL0.properties(kernel.fun)
-    group_size = if kernel_props.maxGroupSize !== missing
-        kernel_props.maxGroupSize
+    max_group_size = oneL0.max_group_size(kernel.fun)
+    group_size = if max_group_size !== missing
+        max_group_size
     else
         # without the MAX_GROUP_SIZE extension, we need to be conservative
         dev = kernel.fun.mod.device
@@ -261,7 +261,7 @@ function launch_configuration(kernel::HostKernel{F,TT}) where {F,TT}
     # size but does not fold it into `maxGroupSize`, so account for it here. Rounded down to
     # a power of two, both because group sizes want to be anyway and to stay clear of the
     # limit rather than right at it.
-    spill = kernel_props.spillMemSize
+    spill = oneL0.spill_mem_size(kernel.fun)
     if spill > 0 && group_size * spill > MAX_GROUP_SCRATCH
         group_size = max(1, prevpow(2, max(1, MAX_GROUP_SCRATCH ÷ spill)))
     end

From 83f991b6c364fe30827b4a09af8f4902f1c8c31e Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard@gmail.com>
Date: Mon, 28 Sep 2026 10:18:55 +0200
Subject: [PATCH 3/9] Make synchronization cooperative

`synchronize` blocked the calling thread in the driver until the work had
completed, so no other task could run on it in the meantime, and no other
thread could run the GC either. KernelInterface requires `synchronize` to
be cooperative.

Wait like CUDA.jl does: busy-wait on a non-blocking query first, which keeps
the latency of short operations low, and then block in the driver (GC-safe)
on one of a few dedicated threads, while the calling task waits for it
without blocking the scheduler.

This applies to `synchronize()` and `synchronize(::oneStream)`, i.e., to user
code, KernelAbstractions, KernelInterface and the synchronizing copies;
`synchronize(; blocking=true)` restores the old behavior. The command list
and queue methods, which also run from finalizers, keep blocking. With
`ONEAPI_SYNC_EACH_SUBMISSION` set, as on Aurora's LTS stack, the wait after
every launch is cooperative too; otherwise it would block the thread, and
leave `synchronize` nothing to wait for.
---
 lib/level-zero/oneL0.jl           |   1 +
 lib/level-zero/synchronization.jl | 195 ++++++++++++++++++++++++++++++
 lib/utils/APIUtils.jl             |   4 +-
 src/compiler/execution.jl         |   3 +-
 src/context.jl                    |  24 ++--
 src/oneAPIKernels.jl              |   1 -
 test/execution.jl                 |  45 +++++++
 7 files changed, 259 insertions(+), 14 deletions(-)
 create mode 100644 lib/level-zero/synchronization.jl

diff --git a/lib/level-zero/oneL0.jl b/lib/level-zero/oneL0.jl
index 82e33eec..c36b9eee 100644
--- a/lib/level-zero/oneL0.jl
+++ b/lib/level-zero/oneL0.jl
@@ -124,6 +124,7 @@ end
 include("context.jl")
 include("cmdqueue.jl")
 include("cmdlist.jl")
+include("synchronization.jl")
 include("fence.jl")
 include("event.jl")
 include("barrier.jl")
diff --git a/lib/level-zero/synchronization.jl b/lib/level-zero/synchronization.jl
new file mode 100644
index 00000000..809a03b8
--- /dev/null
+++ b/lib/level-zero/synchronization.jl
@@ -0,0 +1,195 @@
+# cooperative synchronization
+#
+# `zeCommandListHostSynchronize` and `zeCommandQueueSynchronize` block the calling thread
+# until the work has completed, so no other task can run on it in the meantime. As CUDA.jl
+# does, first busy-wait on a non-blocking query, which keeps the latency of short
+# operations low, and then block in the driver on a separate thread, while the calling
+# task waits for that thread without blocking the scheduler.
+
+export nonblocking_synchronize
+
+const SyncObject = Union{ZeImmediateCommandList, ZeCommandQueue}
+
+# with a zero timeout, a synchronization is a query
+function check_done(res::ze_result_t)
+    if res == RESULT_NOT_READY
+        return false
+    elseif res == RESULT_SUCCESS
+        return true
+    else
+        throw_api_error(res)
+    end
+end
+Base.isdone(list::ZeImmediateCommandList) =
+    check_done(unchecked_zeCommandListHostSynchronize(list, 0))
+Base.isdone(queue::ZeCommandQueue) =
+    check_done(unchecked_zeCommandQueueSynchronize(queue, 0))
+
+# the blocking synchronization, marked GC-safe so that it doesn't keep the GC from running
+gcsafe_synchronize(list::ZeImmediateCommandList) =
+    @gcsafe_ccall libze_loader.zeCommandListHostSynchronize(
+        list::ze_command_list_handle_t, typemax(UInt64)::UInt64)::ze_result_t
+gcsafe_synchronize(queue::ZeCommandQueue) =
+    @gcsafe_ccall libze_loader.zeCommandQueueSynchronize(
+        queue::ze_command_queue_handle_t, typemax(UInt64)::UInt64)::ze_result_t
+
+
+## bidirectional channel
+
+# custom, unbuffered channel that supports returning a value to the sender
+# without the need for a second channel
+struct BidirectionalChannel{I,O} <: AbstractChannel{I}
+    cond_take::Threads.Condition                 # waiting for data to become available
+    cond_put::Threads.Condition                  # waiting for a writeable slot
+    cond_ret::Threads.Condition                  # waiting for a data to be returned
+
+    function BidirectionalChannel{I,O}() where {I,O}
+        lock = ReentrantLock()
+        cond_put = Threads.Condition(lock)
+        cond_take = Threads.Condition(lock)
+        cond_ret = Threads.Condition(lock)
+        return new(cond_take, cond_put, cond_ret)
+    end
+end
+
+Base.put!(c::BidirectionalChannel{I}, v) where {I} = put!(c, convert(I, v))
+function Base.put!(c::BidirectionalChannel{I,O}, v::I) where {I,O}
+    lock(c)
+    try
+        # wait for a slot to be available
+        while isempty(c.cond_take)
+            Base.wait(c.cond_put)
+        end
+
+        # pass a value to the consumer
+        notify(c.cond_take, v, false, false)
+
+        # wait for a return value to be produced
+        Base.wait(c.cond_ret)::O
+    finally
+        unlock(c)
+    end
+end
+
+function Base.take!(f::Base.Callable, c::BidirectionalChannel{I,O}) where {I,O}
+    lock(c)
+    try
+        # notify the producer that we're ready to accept a value
+        notify(c.cond_put, nothing, false, false)
+
+        # receive a value from the producer
+        v = Base.wait(c.cond_take)::I
+
+        # return a value to the producer
+        ret = f(v)::O
+        notify(c.cond_ret, ret, false, false)
+    finally
+        unlock(c)
+    end
+end
+
+Base.lock(c::BidirectionalChannel) = lock(c.cond_take)
+Base.unlock(c::BidirectionalChannel) = unlock(c.cond_take)
+
+
+## fast path
+
+# before blocking on a separate thread, which has some overhead, busy-wait on a query of
+# the object to synchronize. when this returns true, the object still has to be
+# synchronized, but that won't block anymore.
+function spinning_synchronization(f, obj)
+    # fast path
+    f(obj) && return true
+
+    # minimize latency of short operations by busy-waiting,
+    # initially without even yielding to other tasks
+    spins = 0
+    while spins < 256
+        if spins < 32
+            ccall(:jl_cpu_pause, Cvoid, ())
+            # temporary solution before we have gc transition support in codegen.
+            ccall(:jl_gc_safepoint, Cvoid, ())
+        else
+            yield()
+        end
+        f(obj) && return true
+        spins += 1
+    end
+
+    return false
+end
+
+
+## slow path: synchronize on a separate thread
+
+const MAX_SYNC_THREADS = 4
+const sync_channels = Array{BidirectionalChannel{SyncObject,ze_result_t}}(undef, MAX_SYNC_THREADS)
+const sync_channel_cursor = Threads.Atomic{UInt32}(1)
+const sync_channel_lock = Base.ReentrantLock()
+
+function synchronization_worker(data)
+    i = Int(data)
+    chan = sync_channels[i]
+
+    while true
+        # wait for work
+        take!(gcsafe_synchronize, chan)
+    end
+end
+
+@noinline function create_synchronization_worker(i)
+    lock(sync_channel_lock) do
+        # test and test-and-set
+        if isassigned(sync_channels, i)
+            return
+        end
+
+        # should be safe to assign before threads are running;
+        # any user will just submit work that makes it block
+        sync_channels[i] = BidirectionalChannel{SyncObject,ze_result_t}()
+
+        # we don't know what the size of uv_thread_t is, so reserve enough space
+        tid = Ref{NTuple{32, UInt8}}(ntuple(i -> 0, 32))
+
+        cb = @cfunction(synchronization_worker, Cvoid, (Ptr{Cvoid},))
+        err = @ccall uv_thread_create(tid::Ptr{Cvoid}, cb::Ptr{Cvoid}, Ptr{Cvoid}(i)::Ptr{Cvoid})::Cint
+        err == 0 || Base.uv_error("uv_thread_create", err)
+        err = @ccall uv_thread_detach(tid::Ptr{Cvoid})::Cint
+        err == 0 || Base.uv_error("uv_thread_detach", err)
+    end
+
+    return
+end
+
+"""
+    nonblocking_synchronize(list_or_queue)
+
+Wait for the work on an immediate command list or command queue to complete, like
+[`synchronize`](@ref), but without blocking the Julia scheduler: other tasks keep running
+while this one waits.
+"""
+function nonblocking_synchronize(obj::SyncObject)
+    if spinning_synchronization(Base.isdone, obj)
+        # done, so this doesn't block
+        synchronize(obj)
+        return
+    end
+
+    # pick a worker channel: sticky per task, so repeated synchronizations from the
+    # same task always hit the same, already running worker thread.
+    tls = task_local_storage()
+    i = get!(tls, :ZeSyncChannel) do
+        mod1(Threads.atomic_add!(sync_channel_cursor, UInt32(1)), MAX_SYNC_THREADS)
+    end::Int
+    if !isassigned(sync_channels, i)
+        create_synchronization_worker(i)
+    end
+    chan = @inbounds sync_channels[i]
+
+    # submit the object to synchronize; unlike with regular channels, this `put!` blocks
+    # until the worker has synchronized it and returned the result
+    res = put!(chan, obj)
+    res == RESULT_SUCCESS || throw_api_error(res)
+
+    return
+end
diff --git a/lib/utils/APIUtils.jl b/lib/utils/APIUtils.jl
index d8d30394..61ec0743 100644
--- a/lib/utils/APIUtils.jl
+++ b/lib/utils/APIUtils.jl
@@ -1,8 +1,8 @@
 module APIUtils
 
 # helpers that facilitate working with C APIs
-using GPUToolbox: @checked, @debug_ccall
-export @checked, @debug_ccall
+using GPUToolbox: @checked, @debug_ccall, @gcsafe_ccall
+export @checked, @debug_ccall, @gcsafe_ccall
 include("enum.jl")
 
 end
diff --git a/src/compiler/execution.jl b/src/compiler/execution.jl
index 2d4ae759..c92398e9 100644
--- a/src/compiler/execution.jl
+++ b/src/compiler/execution.jl
@@ -360,7 +360,8 @@ end
     spill > s.scratch_hwm && scratch_hedge!(s, spill)
 
     append_launch!(s.list, kernel, groups)
-    oneL0.sync_each_submission() && oneL0.synchronize(s.list)
+    # wait cooperatively, as `synchronize` does, or the workaround blocks the thread
+    oneL0.sync_each_submission() && oneL0.nonblocking_synchronize(s.list)
     return
 end
 
diff --git a/src/context.jl b/src/context.jl
index 409a9d7a..6c9e188f 100644
--- a/src/context.jl
+++ b/src/context.jl
@@ -378,12 +378,15 @@ function synchronize_all_streams(ctx::ZeContext, dev::Union{ZeDevice, Nothing})
 end
 
 """
-    synchronize()
-    synchronize(stream::oneStream)
+    synchronize(; blocking=false)
+    synchronize(stream::oneStream; blocking=false)
 
-Block the host thread until all operations on the calling task's stream for the current
-context and device have completed: work appended to the immediate command list as well
-as oneMKL work on the companion queue.
+Block the calling task until all operations on its stream for the current context and
+device have completed: work appended to the immediate command list as well as oneMKL work
+on the companion queue.
+
+Unless `blocking` is set, other tasks keep running while waiting: the host thread is only
+blocked in the driver when the work is already done.
 
 This is useful for timing operations or ensuring that GPU work has finished before
 accessing results on the CPU.
@@ -398,18 +401,19 @@ println("GPU work completed")
 
 See also: [`global_stream`](@ref), [`context`](@ref), [`device`](@ref)
 """
-function oneL0.synchronize(s::oneStream)
-    oneL0.synchronize(s.list)
+function oneL0.synchronize(s::oneStream; blocking::Bool=false)
+    sync = blocking ? oneL0.synchronize : oneL0.nonblocking_synchronize
+    sync(s.list)
     q = s.queue
     if q !== nothing
-        oneL0.synchronize(q)
+        sync(q)
         s.mkl_dirty = false
     end
     return
 end
 
-function oneL0.synchronize()
-    oneL0.synchronize(global_stream(context(), device()))
+function oneL0.synchronize(; blocking::Bool=false)
+    oneL0.synchronize(global_stream(context(), device()); blocking)
 end
 
 # Julia → MKL ordering: everything Julia appended to the task's immediate list must be
diff --git a/src/oneAPIKernels.jl b/src/oneAPIKernels.jl
index 7b90d2ca..ae693caa 100644
--- a/src/oneAPIKernels.jl
+++ b/src/oneAPIKernels.jl
@@ -26,7 +26,6 @@ oneAPIBackend(; prefer_blocks = false, always_inline = false) = oneAPIBackend(pr
 @inline KA.ones(::oneAPIBackend, ::Type{T}, dims::Tuple; unified::Bool = false) where {T} = fill!(oneArray{T, length(dims), unified ? oneAPI.oneL0.SharedBuffer : oneAPI.oneL0.DeviceBuffer}(undef, dims), one(T))
 
 KA.get_backend(::oneArray) = oneAPIBackend()
-# TODO should be non-blocking
 KA.synchronize(::oneAPIBackend) = oneAPI.oneL0.synchronize()
 KA.supports_float64(::oneAPIBackend) = false  # TODO: Check if this is device dependent
 KA.supports_unified(::oneAPIBackend) = true
diff --git a/test/execution.jl b/test/execution.jl
index 3efcdbf1..330a1612 100644
--- a/test/execution.jl
+++ b/test/execution.jl
@@ -743,6 +743,51 @@ end
     @test all(results)
 end
 
+# burns `iters` dependent steps per work-item, which the compiler cannot fold away
+function slow_kernel(a, iters)
+    i = get_global_id()
+    acc = i % UInt32
+    for k in UInt32(1):iters
+        acc = acc * 0x0019660d + k
+    end
+    @inbounds a[i] = acc
+    return
+end
+
+@testset "cooperative synchronize" begin
+    a = oneArray{UInt32}(undef, 64)
+    slow(iters) = @oneapi items=64 slow_kernel(a, UInt32(iters))
+    slow(1)
+    synchronize()
+    # warm up the slow path of `synchronize`: compiling it would end the calibration early
+    slow(2^16)
+    synchronize()
+
+    # make the kernel run for a while
+    iters = 2^16
+    while iters < 2^30 && @elapsed((slow(iters); synchronize())) < 0.1
+        iters *= 4
+    end
+
+    # other tasks keep running while one waits
+    stamps = UInt64[]
+    waiting = Ref(true)
+    ticker = @async while waiting[]
+        push!(stamps, time_ns())
+        sleep(0.001)
+    end
+    slow(iters)
+    synchronize()
+    done = time_ns()
+    waiting[] = false
+    wait(ticker)
+    @test count(<(done), stamps) > 10
+
+    # blocking synchronization is still available
+    slow(1)
+    @test synchronize(; blocking=true) === nothing
+end
+
 ############################################################################################
 
 # Keep allocation consumers at top level so kernels do not capture test state.

From 9120af0b240a1a67bf82235b9de6b4009f2f0f3b Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard@gmail.com>
Date: Wed, 30 Sep 2026 18:03:56 +0200
Subject: [PATCH 4/9] Implement KernelInterface, and support KernelAbstractions
 0.10

KernelAbstractions 0.10 builds on KernelInterface, which defines what a back
end provides: memory and device management, compiling and launching kernels,
and the device-side intrinsics. KernelAbstractions then launches `@kernel`
kernels itself on any KernelInterface back end, which replaces oneAPI's copy
of that launch path: partitioning the ndrange, building the kernel's context,
and tuning the work-group size.

`oneAPIBackend` now implements KernelInterface, and oneAPI depends on it
instead of on KernelAbstractions. What KernelAbstractions still needs from a
back end moves to an extension: the `MArray` behind `@private`, the Adapt
rule for `@Const`, and adapting to KernelAbstractions' CPU back end.
`prefer_blocks` now applies in `KI.launch_configuration`, which receives the
number of work-items.

The index queries are computed in the requested type with `% T`, which needs
SPIRVIntrinsics 1.1.3 to truncate the 3-D built-ins without producing illegal
vector types. `KI.launch` rejects `items` and `groups`, which would override
the launch geometry that KernelInterface validated, and refuses to launch a
kernel on another context or device than the one it was compiled for.
`KI.kernel_function` converts the callable to compile it, and the kernel keeps
the original alive, as the converted form of a closure only holds pointers to
the arrays it captures.

Co-authored-by: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
---
 Project.toml                 |  13 +-
 ext/KernelAbstractionsExt.jl |  32 +++++
 src/context.jl               |   2 +-
 src/oneAPI.jl                |   2 +-
 src/oneAPIKernels.jl         | 266 ++++++++++++++++-------------------
 src/utils.jl                 |   2 +-
 test/Project.toml            |   1 +
 test/kernelabstractions.jl   |  76 +++++++++-
 test/kernelinterface.jl      |  27 ++++
 9 files changed, 272 insertions(+), 149 deletions(-)
 create mode 100644 ext/KernelAbstractionsExt.jl
 create mode 100644 test/kernelinterface.jl

diff --git a/Project.toml b/Project.toml
index 88b57ba4..378ea7bb 100644
--- a/Project.toml
+++ b/Project.toml
@@ -12,7 +12,7 @@ ExprTools = "e2ba6199-217a-4e67-a87a-7c52f15ade04"
 GPUArrays = "0c68f7d7-f131-5f86-a1c3-88cf8149b2d7"
 GPUCompiler = "61eb1bfa-7361-4325-ad38-22787b887f55"
 GPUToolbox = "096a3bc2-3ced-46d0-87f4-dd12716f4bfc"
-KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c"
+KernelInterface = "4ee993da-d684-4d17-a7dd-4e58e78d92bf"
 LLVM = "929cbde3-209d-540e-8aea-75f648917ca0"
 Libdl = "8f399da3-3557-5675-b5ff-fb832c97cbdb"
 LinearAlgebra = "37e2e46d-f89d-539d-b4ee-838fcccc9c8e"
@@ -32,6 +32,12 @@ oneAPI_Level_Zero_Headers_jll = "f4bc562b-d309-54f8-9efb-476e56f0410d"
 oneAPI_Level_Zero_Loader_jll = "13eca655-d68d-5b81-8367-6d99d727ab01"
 oneAPI_Support_jll = "b049733a-a71d-5ed3-8eba-7d323ac00b36"
 
+[weakdeps]
+KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c"
+
+[extensions]
+KernelAbstractionsExt = "KernelAbstractions"
+
 [compat]
 AbstractFFTs = "1.5.0"
 AcceleratedKernels = "0.3.1, 0.4"
@@ -41,12 +47,13 @@ ExprTools = "0.1"
 GPUArrays = "11.5.14"
 GPUCompiler = "2.9"
 GPUToolbox = "3.1"
-KernelAbstractions = "0.9.39"
+KernelAbstractions = "0.10"
+KernelInterface = "0.4"
 LLVM = "6, 7, 8, 9"
 NEO_jll = "=26.18.38308"
 PrecompileTools = "1"
 Preferences = "1"
-SPIRVIntrinsics = "1"
+SPIRVIntrinsics = "1.1.3"
 SPIRV_LLVM_Backend_jll = "23"
 SPIRV_LLVM_Translator_jll = "23"
 SPIRV_Tools_jll = "2025.4.0"
diff --git a/ext/KernelAbstractionsExt.jl b/ext/KernelAbstractionsExt.jl
new file mode 100644
index 00000000..73f4a14d
--- /dev/null
+++ b/ext/KernelAbstractionsExt.jl
@@ -0,0 +1,32 @@
+module KernelAbstractionsExt
+
+using oneAPI
+using oneAPI: @device_override, method_table
+
+import KernelAbstractions as KA
+
+import StaticArrays
+
+import Adapt
+
+Adapt.adapt_storage(::KA.CPU, a::oneArray) = convert(Array, a)
+
+# sparse arrays (oneMKL is only available on Linux)
+@static if Sys.islinux()
+    import SparseArrays
+    Adapt.adapt_storage(::KA.CPU, a::oneAPI.oneMKL.oneAbstractSparseMatrix) = SparseArrays.SparseMatrixCSC(a)
+end
+
+
+## scratch memory
+
+@device_override @inline function KA.Scratchpad(ctx, ::Type{T}, ::Val{Dims}) where {T, Dims}
+    StaticArrays.MArray{Tuple{Dims...}, T}(undef)
+end
+
+
+## other
+
+Adapt.adapt_storage(to::KA.ConstAdaptor, a::oneDeviceArray) = Base.Experimental.Const(a)
+
+end
diff --git a/src/context.jl b/src/context.jl
index 6c9e188f..521adb3e 100644
--- a/src/context.jl
+++ b/src/context.jl
@@ -284,7 +284,7 @@ global_queue(ctx::ZeContext, dev::ZeDevice) = stream_queue(global_stream(ctx, de
 # Register `stream` as a stream targeting (ctx, dev) so `synchronize_all_streams`/
 # `release` can find and drain it before freeing buffers whose in-flight work it may
 # still reference. EVERY stream that becomes a task's active stream must go through here
-# — not just the one `global_stream` creates but also the replacement `KA.priority!`
+# — not just the one `global_stream` creates but also the replacement `KI.priority!`
 # installs — or the unregistered stream's in-flight work can outlive a freed buffer (a
 # use-after-free that faults and bans the context on the LTS NEO stack). Only the LTS
 # stack maintains the registry; on the rolling stack this is a no-op. Returns `stream`.
diff --git a/src/oneAPI.jl b/src/oneAPI.jl
index 6b6db42c..8ac4261e 100644
--- a/src/oneAPI.jl
+++ b/src/oneAPI.jl
@@ -12,7 +12,7 @@ using SpecialFunctions
 
 import Preferences
 
-import KernelAbstractions: KernelAbstractions
+import KernelInterface
 
 using LLVM
 using LLVM.Interop
diff --git a/src/oneAPIKernels.jl b/src/oneAPIKernels.jl
index ae693caa..a4089e6d 100644
--- a/src/oneAPIKernels.jl
+++ b/src/oneAPIKernels.jl
@@ -1,11 +1,9 @@
 module oneAPIKernels
 
 using ..oneAPI
-using ..oneAPI: @device_override, SPIRVIntrinsics, method_table
+using ..oneAPI: @device_override, SPIRVIntrinsics, method_table, kernel_convert, zefunction
 
-import KernelAbstractions as KA
-
-import StaticArrays
+import KernelInterface as KI
 
 import Adapt
 
@@ -14,215 +12,201 @@ import Adapt
 
 export oneAPIBackend
 
-struct oneAPIBackend <: KA.GPU
+"""
+    oneAPIBackend(; prefer_blocks=false, always_inline=false)
+
+KernelInterface back end for oneAPI, which KernelAbstractions uses to launch kernels on
+Intel GPUs. `prefer_blocks=true` spreads small launches over more, smaller work-groups.
+"""
+struct oneAPIBackend <: KI.Backend
     prefer_blocks::Bool
     always_inline::Bool
 end
 
 oneAPIBackend(; prefer_blocks = false, always_inline = false) = oneAPIBackend(prefer_blocks, always_inline)
 
-@inline KA.allocate(::oneAPIBackend, ::Type{T}, dims::Tuple; unified::Bool = false) where {T} = oneArray{T, length(dims), unified ? oneAPI.oneL0.SharedBuffer : oneAPI.oneL0.DeviceBuffer}(undef, dims)
-@inline KA.zeros(::oneAPIBackend, ::Type{T}, dims::Tuple; unified::Bool = false) where {T} = fill!(oneArray{T, length(dims), unified ? oneAPI.oneL0.SharedBuffer : oneAPI.oneL0.DeviceBuffer}(undef, dims), zero(T))
-@inline KA.ones(::oneAPIBackend, ::Type{T}, dims::Tuple; unified::Bool = false) where {T} = fill!(oneArray{T, length(dims), unified ? oneAPI.oneL0.SharedBuffer : oneAPI.oneL0.DeviceBuffer}(undef, dims), one(T))
+@inline KI.allocate(::oneAPIBackend, ::Type{T}, dims::Tuple; unified::Bool = false) where {T} = oneArray{T, length(dims), unified ? oneAPI.oneL0.SharedBuffer : oneAPI.oneL0.DeviceBuffer}(undef, dims)
 
-KA.get_backend(::oneArray) = oneAPIBackend()
-KA.synchronize(::oneAPIBackend) = oneAPI.oneL0.synchronize()
-KA.supports_float64(::oneAPIBackend) = false  # TODO: Check if this is device dependent
-KA.supports_unified(::oneAPIBackend) = true
+KI.get_backend(::oneArray) = oneAPIBackend()
+KI.synchronize(::oneAPIBackend) = oneAPI.oneL0.synchronize()
+KI.supports_float64(::oneAPIBackend) = device_limits().supports_float64
+KI.supports_unified(::oneAPIBackend) = true
+KI.supports_atomics(::oneAPIBackend) = true
 
-KA.functional(::oneAPIBackend) = oneAPI.functional()
+KI.functional(::oneAPIBackend) = oneAPI.functional()
 
 Adapt.adapt_storage(::oneAPIBackend, a::AbstractArray) = Adapt.adapt(oneArray, a)
 Adapt.adapt_storage(::oneAPIBackend, a::oneArray) = a
-Adapt.adapt_storage(::KA.CPU, a::oneArray) = convert(Array, a)
 
 # sparse arrays (oneMKL is only available on Linux)
 @static if Sys.islinux()
-    import GPUArrays, SparseArrays
-    KA.get_backend(::oneAPI.oneMKL.oneAbstractSparseMatrix) = oneAPIBackend()
+    import GPUArrays
+    KI.get_backend(::oneAPI.oneMKL.oneAbstractSparseMatrix) = oneAPIBackend()
     # without this, `adapt_storage(::oneAPIBackend, ::AbstractArray)` would densify sparse arrays
     Adapt.adapt_storage(::oneAPIBackend, a::GPUArrays.AbstractGPUSparseArray) = a
-    Adapt.adapt_storage(::KA.CPU, a::oneAPI.oneMKL.oneAbstractSparseMatrix) = SparseArrays.SparseMatrixCSC(a)
 end
 
 
 ## Memory Operations
 
-function KA.copyto!(::oneAPIBackend, A, B)
+function KI.copyto!(::oneAPIBackend, A, B)
+    length(A) == length(B) ||
+        throw(ArgumentError("Arrays must have the same length, got $(length(A)) and $(length(B))"))
     copyto!(A, B)
     # TODO: Address device to host copies in jl being synchronizing
+    return A
 end
 
+KI.unsafe_free!(A::oneArray) = oneAPI.unsafe_free!(A)
+
 
 ## Device Operations
 
-function KA.ndevices(::oneAPIBackend)
+function KI.ndevices(::oneAPIBackend)
     return length(oneAPI.devices())
 end
 
-function KA.device(::oneAPIBackend)::Int
-    dev = oneAPI.device()
+function device_index(dev)::Int
     devs = oneAPI.devices()
     idx = findfirst(==(dev), devs)
     return idx === nothing ? 1 : idx
 end
+KI.device(::oneAPIBackend)::Int = device_index(oneAPI.device())
+KI.device(::oneAPIBackend, A::oneArray)::Int = device_index(oneAPI.device(A))
 
-function KA.device!(backend::oneAPIBackend, id::Int)
-    return oneAPI.device!(id)
+function KI.device!(backend::oneAPIBackend, id::Int)
+    oneAPI.device!(id)
+    return
 end
 
 
 ## Kernel Launch
 
-function KA.mkcontext(kernel::KA.Kernel{oneAPIBackend}, _ndrange, iterspace)
-    KA.CompilerMetadata{KA.ndrange(kernel), KA.DynamicCheck}(_ndrange, iterspace)
-end
-function KA.mkcontext(kernel::KA.Kernel{oneAPIBackend}, I, _ndrange, iterspace,
-                      ::Dynamic) where Dynamic
-    KA.CompilerMetadata{KA.ndrange(kernel), Dynamic}(I, _ndrange, iterspace)
-end
-
-function KA.launch_config(kernel::KA.Kernel{oneAPIBackend}, ndrange, workgroupsize)
-    if ndrange isa Integer
-        ndrange = (ndrange,)
-    end
-    if workgroupsize isa Integer
-        workgroupsize = (workgroupsize, )
-    end
+KI.argconvert(::oneAPIBackend, arg) = kernel_convert(arg)
 
-    # partition checked that the ndrange's agreed
-    if KA.ndrange(kernel) <: KA.StaticSize
-        ndrange = nothing
-    end
-
-    iterspace, dynamic = if KA.workgroupsize(kernel) <: KA.DynamicSize &&
-        workgroupsize === nothing
-        # use ndrange as preliminary workgroupsize for autotuning
-        # (clamped to 1, since an empty ndrange cannot serve as a workgroup size)
-        KA.partition(kernel, ndrange, max.(ndrange, 1))
-    else
-        KA.partition(kernel, ndrange, workgroupsize)
-    end
+# a compiled kernel, and the callable it was compiled from: the converted callable only holds
+# pointers to the arrays a closure captures, so the kernel has to keep the original alive
+struct oneAPIKernel{F, H <: oneAPI.HostKernel}
+    f::F
+    host::H
+end
 
-    return ndrange, workgroupsize, iterspace, dynamic
+function KI.kernel_function(backend::oneAPIBackend, f::F, tt::TT=Tuple{}; name = nothing, kwargs...) where {F,TT}
+    host = zefunction(kernel_convert(f), tt; name, backend.always_inline, kwargs...)
+    kern = oneAPIKernel(f, host)
+    KI.Kernel{oneAPIBackend, typeof(kern)}(backend, kern)
 end
 
-function threads_to_workgroupsize(threads, ndrange)
-    total = 1
-    return map(ndrange) do n
-        x = max(1, min(div(threads, total), n))
-        total *= x
-        return x
+# XXX: calling a `HostKernel` takes the arguments as varargs, so this splats them
+function KI.launch(obj::KI.Kernel{oneAPIBackend}, groups::Dims{3}, items::Dims{3},
+                   args::Tuple; kwargs...)
+    # KernelInterface has validated the launch geometry
+    if haskey(kwargs, :items) || haskey(kwargs, :groups)
+        throw(ArgumentError("KernelInterface kernels take `numgroups`, `workgroupsize` or `ndrange`, not `items` or `groups`"))
     end
+    # kernels are compiled for a context and device, and by default launched on the task's
+    # stream for the active ones
+    (; f, host) = obj.kern
+    check_target(host.fun.mod, get(kwargs, :queue, nothing))
+    GC.@preserve f host(args...; items, groups, kwargs...)
+    return
 end
 
-function (obj::KA.Kernel{oneAPIBackend})(args...; ndrange=nothing, workgroupsize=nothing)
-    backend = KA.backend(obj)
-
-    ndrange, workgroupsize, iterspace, dynamic = KA.launch_config(obj, ndrange, workgroupsize)
-    # this might not be the final context, since we may tune the workgroupsize
-    ctx = KA.mkcontext(obj, ndrange, iterspace)
-
-    # If the kernel is statically sized we can tell the compiler about that
-    if KA.workgroupsize(obj) <: KA.StaticSize
-        # TODO: maxthreads
-        # maxthreads = prod(KA.get(KA.workgroupsize(obj)))
+function check_target(mod::oneAPI.oneL0.ZeModule, queue)
+    ctx, dev = if queue === nothing
+        context(), device()
+    elseif queue isa oneAPI.oneStream
+        queue.ctx, queue.dev
     else
-        # maxthreads = nothing
+        queue.context, queue.device
     end
+    mod.context == ctx && mod.device == dev ||
+        throw(ArgumentError("Cannot launch a kernel on another context or device than the one it was compiled for"))
+    return
+end
 
-    kernel = @oneapi launch = false always_inline = backend.always_inline obj.f(ctx, args...)
-
-    # figure out the optimal workgroupsize automatically
-    if KA.workgroupsize(obj) <: KA.DynamicSize && workgroupsize === nothing
-        items = oneAPI.launch_configuration(kernel)
-
-        if backend.prefer_blocks
-            # Prefer blocks over threads:
-            # Reducing the workgroup size (items) increases the number of workgroups (blocks).
-            # We use a simple heuristic here since we lack full occupancy info (max_blocks) from launch_configuration.
-
-            # If the total range is large enough, full workgroups are fine.
-            # If the range is small, we might want to reduce 'items' to create more blocks to fill the GPU.
-            # (Simplified logic compared to CUDA.jl which uses explicit occupancy calculators)
-            total_items = prod(ndrange)
-            if total_items < items * 16 # Heuristic factor
-                # Force at least a few blocks if possible by reducing items per block
-                target_blocks = 16 # Target at least 16 blocks
-                items = max(1, min(items, cld(total_items, target_blocks)))
-            end
+function KI.max_work_group_size(kernel::KI.Kernel{oneAPIBackend})::Int
+    fun = kernel.kern.host.fun
+    max_group_size = oneAPI.oneL0.max_group_size(fun)
+    # without the MAX_GROUP_SIZE extension, the device limit is all we know
+    return coalesce(max_group_size, device_limits(fun.mod.device).max_work_group_size)
+end
+function KI.launch_configuration(
+        kernel::KI.Kernel{oneAPIBackend}; nitems::Union{Integer, Nothing} = nothing,
+        max_work_group_size::Integer = typemax(Int)
+    )
+    items = min(oneAPI.launch_configuration(kernel.kern.host), max_work_group_size)
+    if kernel.backend.prefer_blocks && nitems !== nothing
+        # prefer more, smaller work-groups: without an occupancy API that tells how many
+        # work-groups fill the device, make sure that a small launch uses at least a few
+        target_groups = 16
+        if nitems < items * target_groups
+            items = max(1, min(items, cld(nitems, target_groups)))
         end
-
-        workgroupsize = threads_to_workgroupsize(items, ndrange)
-        iterspace, dynamic = KA.partition(obj, ndrange, workgroupsize)
-        ctx = KA.mkcontext(obj, ndrange, iterspace)
     end
-
-    groups = length(KA.blocks(iterspace))
-    items = length(KA.workitems(iterspace))
-
-    if groups == 0
-        return nothing
+    return (; workgroupsize = Int(items))
+end
+# querying the device allocates, so cache what every launch needs
+const DeviceLimits = @NamedTuple{
+    max_work_group_size::Int, max_work_group_dims::NTuple{3, Int}, max_num_groups::NTuple{3, Int},
+    supports_float64::Bool,
+}
+function device_limits(dev::oneAPI.oneL0.ZeDevice = device())
+    limits = get!(task_local_storage(), :oneAPIDeviceLimits) do
+        Dict{oneAPI.oneL0.ZeDevice, DeviceLimits}()
+    end::Dict{oneAPI.oneL0.ZeDevice, DeviceLimits}
+    get!(limits, dev) do
+        props = oneAPI.oneL0.compute_properties(dev)
+        module_props = oneAPI.oneL0.module_properties(dev)
+        (; max_work_group_size = props.maxTotalGroupSize,
+           max_work_group_dims = (props.maxGroupSizeX, props.maxGroupSizeY, props.maxGroupSizeZ),
+           max_num_groups = (props.maxGroupCountX, props.maxGroupCountY, props.maxGroupCountZ),
+           supports_float64 = module_props.flags & oneAPI.oneL0.ZE_DEVICE_MODULE_FLAG_FP64 != 0)
     end
-
-    # Launch kernel
-    kernel(ctx, args...; items, groups)
-
-    return nothing
+end
+KI.max_work_group_size(::oneAPIBackend)::Int = device_limits().max_work_group_size
+KI.max_work_group_dims(::oneAPIBackend)::NTuple{3, Int} = device_limits().max_work_group_dims
+KI.max_num_groups(::oneAPIBackend)::NTuple{3, Int} = device_limits().max_num_groups
+# Level Zero calls Xe cores sub-slices
+function KI.multiprocessor_count(::oneAPIBackend)::Int
+    props = oneAPI.oneL0.properties(device())
+    return props.numSlices * props.numSubslicesPerSlice
 end
 
 
 ## Indexing Functions
+## COV_EXCL_START
 
-@device_override @inline function KA.__index_Local_Linear(ctx)
-    return get_local_id()
-end
+# computed with `% T`, which unlike `T(x)` has no error path
 
-@device_override @inline function KA.__index_Group_Linear(ctx)
-    return get_group_id()
+@device_override @inline function KI.get_local_id(::Type{T}) where {T}
+    return (; x = get_local_id(1) % T, y = get_local_id(2) % T, z = get_local_id(3) % T)
 end
 
-@device_override @inline function KA.__index_Global_Linear(ctx)
-    return get_global_id()
+@device_override @inline function KI.get_group_id(::Type{T}) where {T}
+    return (; x = get_group_id(1) % T, y = get_group_id(2) % T, z = get_group_id(3) % T)
 end
 
-@device_override @inline function KA.__index_Local_Cartesian(ctx)
-    @inbounds KA.workitems(KA.__iterspace(ctx))[get_local_id()]
+@device_override @inline function KI.get_local_size(::Type{T}) where {T}
+    return (; x = get_local_size(1) % T, y = get_local_size(2) % T, z = get_local_size(3) % T)
 end
 
-@device_override @inline function KA.__index_Group_Cartesian(ctx)
-    @inbounds KA.blocks(KA.__iterspace(ctx))[get_group_id()]
+@device_override @inline function KI.get_num_groups(::Type{T}) where {T}
+    return (; x = get_num_groups(1) % T, y = get_num_groups(2) % T, z = get_num_groups(3) % T)
 end
 
-@device_override @inline function KA.__index_Global_Cartesian(ctx)
-    return @inbounds KA.expand(KA.__iterspace(ctx), get_group_id(), get_local_id())
-end
-
-@device_override @inline function KA.__validindex(ctx)
-    if KA.__dynamic_checkbounds(ctx)
-        I = @inbounds KA.expand(KA.__iterspace(ctx), get_group_id(), get_local_id())
-        return I in KA.__ndrange(ctx)
-    else
-        return true
-    end
-end
 
+## Shared Memory
 
-## Shared and Scratch Memory
-
-@device_override @inline function KA.SharedMemory(::Type{T}, ::Val{Dims}, ::Val{Id}) where {T, Dims, Id}
+@device_override @inline function KI.localmemory(::Type{T}, ::Val{Dims}) where {T, Dims}
     ptr = oneAPI.emit_localmemory(T, Val(prod(Dims)))
     oneDeviceArray(Dims, ptr)
 end
 
-@device_override @inline function KA.Scratchpad(ctx, ::Type{T}, ::Val{Dims}) where {T, Dims}
-    StaticArrays.MArray{KA.__size(Dims), T}(undef)
-end
-
 
 ## Synchronization and Printing
 
-@device_override @inline function KA.__synchronize()
+@device_override @inline function KI.barrier()
     # Fence both local and global memory across the workgroup barrier, matching CUDA
     # `__syncthreads` semantics. `barrier(0)` lowers to `OpControlBarrier` with
     # `SequentiallyConsistent` but WITHOUT any storage-class bit, which the SPIR-V spec
@@ -232,18 +216,16 @@ end
     barrier(SPIRVIntrinsics.LOCAL_MEM_FENCE | SPIRVIntrinsics.GLOBAL_MEM_FENCE)
 end
 
-@device_override @inline function KA.__print(args...)
+@device_override @inline function KI._print(args...)
     oneAPI._print(args...)
 end
 
+## COV_EXCL_STOP
 
-## Other
-
-Adapt.adapt_storage(to::KA.ConstAdaptor, a::oneDeviceArray) = Base.Experimental.Const(a)
 
-KA.argconvert(::KA.Kernel{oneAPIBackend}, arg) = kernel_convert(arg)
+## Other
 
-function KA.priority!(::oneAPIBackend, prio::Symbol)
+function KI.priority!(::oneAPIBackend, prio::Symbol)
     if !(prio in (:high, :normal, :low))
         error("priority must be one of :high, :normal, :low")
     end
diff --git a/src/utils.jl b/src/utils.jl
index 03d4de12..05bceb8a 100644
--- a/src/utils.jl
+++ b/src/utils.jl
@@ -37,7 +37,7 @@ function versioninfo(io::IO=stdout)
     println(io, "Julia packages:")
     println(io, "- oneAPI.jl: $(Base.pkgversion(oneAPI))")
     for pkg in [:GPUArrays, :GPUCompiler, ("63c18a36-062a-441e-b654-da1e3ab1ce7c", "KernelAbstractions"),
-                 :LLVM, :SPIRVIntrinsics]
+                 :KernelInterface, :LLVM, :SPIRVIntrinsics]
         name, mod = get_module(pkg)
         isnothing(mod) || println(io, "- $(name): $(Base.pkgversion(mod))")
     end
diff --git a/test/Project.toml b/test/Project.toml
index 190beeab..fe34ca49 100644
--- a/test/Project.toml
+++ b/test/Project.toml
@@ -9,6 +9,7 @@ GPUArrays = "0c68f7d7-f131-5f86-a1c3-88cf8149b2d7"
 InteractiveUtils = "b77e0a4c-d291-57a0-90e8-8db25a27a240"
 JLD2 = "033835bb-8acc-5ee8-8aae-3f567f8a3819"
 KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c"
+KernelInterface = "4ee993da-d684-4d17-a7dd-4e58e78d92bf"
 LinearAlgebra = "37e2e46d-f89d-539d-b4ee-838fcccc9c8e"
 NEO_jll = "700fe977-ac61-5f37-bbc8-c6c4b2b6a9fd"
 ParallelTestRunner = "d3525ed8-44d0-4b2c-a655-542cee43accc"
diff --git a/test/kernelabstractions.jl b/test/kernelabstractions.jl
index 48bbe1cb..e3f3b56c 100644
--- a/test/kernelabstractions.jl
+++ b/test/kernelabstractions.jl
@@ -1,8 +1,82 @@
 import KernelAbstractions
+import KernelAbstractions as KA
+import KernelInterface as KI
 include(joinpath(dirname(pathof(KernelAbstractions)), "..", "test", "testsuite.jl"))
 
 skip_tests=Set([
     "sparse",
     "Convert", # Need to opt out of i128
+    "Random", # oneAPI doesn't support Random's default RNG in kernels yet
+    # these run kernels on KernelAbstractions' POCL-based CPU back-end
+    "CPU synchronization",
+    "fallback test: callable types",
 ])
-Testsuite.testsuite(oneAPIBackend, "oneAPI", oneAPI, oneArray, oneDeviceArray; skip_tests)
+Testsuite.testsuite(()->oneAPIBackend(), "oneAPI", oneAPI, oneArray, oneDeviceArray; skip_tests)
+
+KA.@kernel function store_global_linear!(A)
+    I = KA.@index(Global, Linear)
+    @inbounds A[I] = I
+end
+
+KA.@kernel function store_last_index!(A)
+    I = KA.@index(Global, Linear)
+    if I == prod(KA.@ndrange())
+        @inbounds A[1] = I
+        @inbounds A[2] = KA.@index(Global, Cartesian)[2]
+    end
+end
+
+@testset "launch configuration" begin
+    backend = oneAPIBackend()
+    function select(kernel, ndrange, workgroupsize=nothing)
+        ndrange, workgroupsize, iterspace, _ = KA.launch_config(kernel, ndrange, workgroupsize)
+        KA.select_launch(kernel, workgroupsize, iterspace)
+    end
+
+    # kernels are launched on an N-d grid, computing indices in 32 bits
+    kernel = store_global_linear!(backend)
+    @test select(kernel, (64, 32, 16)) === KA.NDLaunch{Int32}()
+    @test select(kernel, (4, 4, 4, 4)) === KA.LinearLaunch{Int32}()
+
+    # which doesn't need divisions to compute the index of a dynamic N-d range
+    A = oneAPI.zeros(Int, 64, 32, 16)
+    ir = sprint(io -> oneAPI.@device_code_llvm io=io kernel(A; ndrange=size(A)))
+    @test !occursin(r"\b[us](div|rem) ", ir)
+    @test Array(A) == LinearIndices(A)
+
+    # tuning for more work-groups (the testsuite above uses the default)
+    Testsuite.launch_testsuite(()->oneAPIBackend(; prefer_blocks=true), oneArray)
+
+    # iteration spaces that don't fit 32 bits use 64-bit indices
+    kernel = store_last_index!(backend)
+    A = oneAPI.zeros(Int, 2)
+    for (dims, launch) in (((2^16 + 1, 2^15), KA.NDLaunch{Int}()),
+                           ((2^11 + 1, 2^10, 2^10, 1), KA.LinearLaunch{Int}()))
+        @test select(kernel, dims) === launch
+        kernel(A; ndrange=dims)
+        @test Array(A) == [prod(dims), dims[2]]
+    end
+end
+
+function ki_store_index!(A)
+    i = KI.get_global_id().x
+    if i <= length(A)
+        @inbounds A[i] = i
+    end
+    return
+end
+
+@testset "tuning" begin
+    # tuning receives the number of work-items; prefer_blocks launches more, smaller groups
+    A = oneAPI.zeros(Int, 1024)
+    kernel = KI.@launch oneAPIBackend() launch=false ki_store_index!(A)
+    items = KI.launch_configuration(kernel; nitems=1024).workgroupsize
+    @test items <= KI.max_work_group_size(kernel)
+    kernel = KI.@launch oneAPIBackend(; prefer_blocks=true) launch=false ki_store_index!(A)
+    fewer = KI.launch_configuration(kernel; nitems=1024).workgroupsize
+    @test fewer < items
+    # ... but not for a bound on the work-group size alone
+    @test KI.launch_configuration(kernel; max_work_group_size=1024).workgroupsize == items
+    KI.@launch oneAPIBackend(; prefer_blocks=true) ndrange=length(A) ki_store_index!(A)
+    @test Array(A) == 1:1024
+end
diff --git a/test/kernelinterface.jl b/test/kernelinterface.jl
new file mode 100644
index 00000000..30f00df2
--- /dev/null
+++ b/test/kernelinterface.jl
@@ -0,0 +1,27 @@
+import KernelInterface
+import KernelInterface as KI
+
+include(joinpath(dirname(pathof(KernelInterface)), "..", "test", "testsuite.jl"))
+
+Testsuite.testsuite(oneAPIBackend(), oneArray)
+
+function ki_fill!(A)
+    i = KI.get_global_id().x
+    if i <= length(A)
+        @inbounds A[i] = i
+    end
+    return
+end
+
+@testset "launch keywords" begin
+    A = oneAPI.zeros(Int, 4)
+    kernel = KI.@launch oneAPIBackend() launch=false ki_fill!(A)
+
+    # oneAPI's launch options are passed on
+    kernel(A; ndrange=4, queue=oneAPI.global_stream(oneAPI.context(), oneAPI.device()))
+    @test Array(A) == 1:4
+
+    # but not ones that would override the launch geometry
+    @test_throws ArgumentError kernel(A; ndrange=4, items=8)
+    @test_throws ArgumentError kernel(A; ndrange=4, groups=2)
+end

From bcc69567d6a7b31ccf0ca7c965e4ef34adc92062 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Tue, 6 Jan 2026 15:40:48 -0400
Subject: [PATCH 5/9] Allow setting the sub-group size of a kernel

The `sub_group_size` compiler keyword sets the `intel_reqd_sub_group_size`
metadata of the kernel. By default, the compiler still chooses the sub-group
size, and a size the device doesn't support is an error.
---
 src/compiler/compilation.jl | 21 ++++++++++++++++++---
 src/compiler/execution.jl   |  5 ++++-
 test/execution.jl           | 16 ++++++++++++++++
 3 files changed, 38 insertions(+), 4 deletions(-)

diff --git a/src/compiler/compilation.jl b/src/compiler/compilation.jl
index 8f463103..a6a78b08 100644
--- a/src/compiler/compilation.jl
+++ b/src/compiler/compilation.jl
@@ -1,6 +1,9 @@
 ## gpucompiler interface implementation
 
-struct oneAPICompilerParams <: AbstractCompilerParams end
+Base.@kwdef struct oneAPICompilerParams <: AbstractCompilerParams
+    sub_group_size::Union{Nothing,Int} = nothing
+end
+
 const oneAPICompilerConfig = CompilerConfig{SPIRVCompilerTarget, oneAPICompilerParams}
 const oneAPICompilerJob = CompilerJob{SPIRVCompilerTarget,oneAPICompilerParams}
 
@@ -58,6 +61,11 @@ function GPUCompiler.finish_module!(job::oneAPICompilerJob, mod::LLVM.Module,
                    Tuple{CompilerJob{SPIRVCompilerTarget}, typeof(mod), typeof(entry)},
                    job, mod, entry)
 
+    # require a sub-group size, instead of letting the compiler choose one
+    if job.config.params.sub_group_size !== nothing
+        metadata(entry)["intel_reqd_sub_group_size"] = MDNode([ConstantInt(Int32(job.config.params.sub_group_size))])
+    end
+
     # OpenCL 2.0
     push!(metadata(mod)["opencl.ocl.version"],
           MDNode([ConstantInt(Int32(2)),
@@ -273,11 +281,18 @@ function _driver_supports_bfloat16_spirv(dev=device())
     end
 end
 
-@noinline function _compiler_config(dev; kernel=true, name=nothing, always_inline=false, kwargs...)
+@noinline function _compiler_config(dev; kernel=true, name=nothing, always_inline=false,
+                                    sub_group_size=nothing, kwargs...)
     properties = oneL0.module_properties(dev)
     supports_fp16 = properties.fp16flags & oneL0.ZE_DEVICE_MODULE_FLAG_FP16 == oneL0.ZE_DEVICE_MODULE_FLAG_FP16
     supports_fp64 = properties.fp64flags & oneL0.ZE_DEVICE_MODULE_FLAG_FP64 == oneL0.ZE_DEVICE_MODULE_FLAG_FP64
 
+    if sub_group_size !== nothing
+        sizes = oneL0.compute_properties(dev).subGroupSizes
+        sub_group_size in sizes ||
+            throw(ArgumentError("Sub-group size $sub_group_size is not supported by this device, which supports $(join(sizes, ", "))"))
+    end
+
     # SPIR-V codegen path. The Aurora LTS NEO/IGC runtime only accepts SPIR-V from the
     # Khronos translator; the rolling stack uses the LLVM SPIR-V back-end. GPUCompiler picks
     # the tool from the target's `backend` field and loads the JLL lazily, so both can be
@@ -311,7 +326,7 @@ end
     # create GPUCompiler objects
     target = SPIRVCompilerTarget(; backend, extensions = extensions_str, supports_fp16, supports_fp64, supports_bfloat16,
                                    driver = :intel, kwargs...)
-    params = oneAPICompilerParams()
+    params = oneAPICompilerParams(; sub_group_size)
     CompilerConfig(target, params; kernel, name, always_inline)
 end
 
diff --git a/src/compiler/execution.jl b/src/compiler/execution.jl
index c92398e9..4adcdb73 100644
--- a/src/compiler/execution.jl
+++ b/src/compiler/execution.jl
@@ -4,7 +4,7 @@ export @oneapi, zefunction, kernel_convert
 ## high-level @oneapi interface
 
 const MACRO_KWARGS = [:launch]
-const COMPILER_KWARGS = [:kernel, :name, :always_inline]
+const COMPILER_KWARGS = [:kernel, :name, :always_inline, :sub_group_size]
 const LAUNCH_KWARGS = [:groups, :items, :queue]
 
 """
@@ -25,6 +25,9 @@ launches the kernel on the GPU.
 - `kernel::Bool=false`: Whether to compile as a kernel (true) or device function (false)
 - `name::Union{String,Nothing}=nothing`: Explicit name for the kernel
 - `always_inline::Bool=false`: Whether to always inline device functions
+- `sub_group_size::Union{Int,Nothing}=nothing`: The sub-group size the kernel has to be
+  compiled for, one of the device's `oneL0.compute_properties(dev).subGroupSizes`. By
+  default, the compiler chooses one.
 
 ## Launch Keywords (runtime)
 - `groups`: Number of workgroups (required). Can be an integer or tuple.
diff --git a/test/execution.jl b/test/execution.jl
index 330a1612..2f943d76 100644
--- a/test/execution.jl
+++ b/test/execution.jl
@@ -41,6 +41,22 @@ end
 end
 
 
+function store_max_sub_group_size(a)
+    @inbounds a[1] = get_max_sub_group_size()
+    return
+end
+
+@testset "sub-group size" begin
+    a = oneArray{UInt32}(undef, 1)
+    sizes = oneL0.compute_properties(device()).subGroupSizes
+    for sub_group_size in sizes
+        @oneapi items=64 sub_group_size store_max_sub_group_size(a)
+        @test Array(a)[1] == sub_group_size
+    end
+    @test_throws ArgumentError @oneapi sub_group_size=maximum(sizes)+1 dummy()
+end
+
+
 @testset "inference" begin
     foo() = @oneapi dummy()
     @inferred foo()

From b12c6ba6c2d22ebce7bdd2a845d533fdf75707ab Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard@gmail.com>
Date: Wed, 30 Sep 2026 18:15:39 +0200
Subject: [PATCH 6/9] Support KernelInterface's sub-groups

Implement the sub-group queries, `sub_group_barrier` and `shfl_down`, and
report their support to KernelInterface, which requires kernels to execute
with a fixed sub-group width: `kernel_function` compiles for 32 work-items,
like a CUDA warp, if the device supports that width, and for its widest one
otherwise. Shuffles of `Float16` and `Float64` values depend on the device
supporting those types.

KernelInterface leaves unspecified how work-items are grouped into
sub-groups; Intel GPUs form them from consecutive linear work-item indices,
so the last sub-group of a work-group can be partial, which a test checks.

Co-authored-by: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
---
 src/oneAPIKernels.jl    | 45 +++++++++++++++++++++++++++++++++++++++--
 test/kernelinterface.jl | 36 +++++++++++++++++++++++++++++++++
 2 files changed, 79 insertions(+), 2 deletions(-)

diff --git a/src/oneAPIKernels.jl b/src/oneAPIKernels.jl
index a4089e6d..b1507d10 100644
--- a/src/oneAPIKernels.jl
+++ b/src/oneAPIKernels.jl
@@ -32,6 +32,13 @@ KI.synchronize(::oneAPIBackend) = oneAPI.oneL0.synchronize()
 KI.supports_float64(::oneAPIBackend) = device_limits().supports_float64
 KI.supports_unified(::oneAPIBackend) = true
 KI.supports_atomics(::oneAPIBackend) = true
+KI.supports_subgroups(::oneAPIBackend) = device_limits().sub_group_size > 0
+function KI.supports_shuffle(::oneAPIBackend, ::Type{T}) where {T}
+    T in SPIRVIntrinsics.gentypes || return false
+    T === Float64 && return device_limits().supports_float64
+    T === Float16 && return device_limits().supports_float16
+    return true
+end
 
 KI.functional(::oneAPIBackend) = oneAPI.functional()
 
@@ -92,7 +99,16 @@ struct oneAPIKernel{F, H <: oneAPI.HostKernel}
 end
 
 function KI.kernel_function(backend::oneAPIBackend, f::F, tt::TT=Tuple{}; name = nothing, kwargs...) where {F,TT}
-    host = zefunction(kernel_convert(f), tt; name, backend.always_inline, kwargs...)
+    # compile for the sub-group width that `KI.sub_group_size` promises
+    sub_group_size = KI.sub_group_size(backend)
+    if get(kwargs, :sub_group_size, sub_group_size) != sub_group_size
+        throw(ArgumentError("KernelInterface kernels are compiled for a sub-group size of $sub_group_size, got $(kwargs[:sub_group_size])"))
+    end
+    host = if sub_group_size > 0
+        zefunction(kernel_convert(f), tt; name, backend.always_inline, kwargs..., sub_group_size)
+    else
+        zefunction(kernel_convert(f), tt; name, backend.always_inline, kwargs...)
+    end
     kern = oneAPIKernel(f, host)
     KI.Kernel{oneAPIBackend, typeof(kern)}(backend, kern)
 end
@@ -149,7 +165,7 @@ end
 # querying the device allocates, so cache what every launch needs
 const DeviceLimits = @NamedTuple{
     max_work_group_size::Int, max_work_group_dims::NTuple{3, Int}, max_num_groups::NTuple{3, Int},
-    supports_float64::Bool,
+    sub_group_size::Int, supports_float16::Bool, supports_float64::Bool,
 }
 function device_limits(dev::oneAPI.oneL0.ZeDevice = device())
     limits = get!(task_local_storage(), :oneAPIDeviceLimits) do
@@ -158,15 +174,22 @@ function device_limits(dev::oneAPI.oneL0.ZeDevice = device())
     get!(limits, dev) do
         props = oneAPI.oneL0.compute_properties(dev)
         module_props = oneAPI.oneL0.module_properties(dev)
+        # the sub-group width that `kernel_function` compiles for: 32, like a CUDA warp, if
+        # the device supports it, and 0 if the device has no sub-groups
+        sg_sizes = props.subGroupSizes
+        sub_group_size = 32 in sg_sizes ? 32 : maximum(sg_sizes; init = 0)
         (; max_work_group_size = props.maxTotalGroupSize,
            max_work_group_dims = (props.maxGroupSizeX, props.maxGroupSizeY, props.maxGroupSizeZ),
            max_num_groups = (props.maxGroupCountX, props.maxGroupCountY, props.maxGroupCountZ),
+           sub_group_size,
+           supports_float16 = module_props.flags & oneAPI.oneL0.ZE_DEVICE_MODULE_FLAG_FP16 != 0,
            supports_float64 = module_props.flags & oneAPI.oneL0.ZE_DEVICE_MODULE_FLAG_FP64 != 0)
     end
 end
 KI.max_work_group_size(::oneAPIBackend)::Int = device_limits().max_work_group_size
 KI.max_work_group_dims(::oneAPIBackend)::NTuple{3, Int} = device_limits().max_work_group_dims
 KI.max_num_groups(::oneAPIBackend)::NTuple{3, Int} = device_limits().max_num_groups
+KI.sub_group_size(::oneAPIBackend)::Int = device_limits().sub_group_size
 # Level Zero calls Xe cores sub-slices
 function KI.multiprocessor_count(::oneAPIBackend)::Int
     props = oneAPI.oneL0.properties(device())
@@ -195,6 +218,16 @@ end
     return (; x = get_num_groups(1) % T, y = get_num_groups(2) % T, z = get_num_groups(3) % T)
 end
 
+@device_override KI.get_sub_group_size(::Type{T}) where {T} = get_sub_group_size() % T
+
+@device_override KI.get_max_sub_group_size(::Type{T}) where {T} = get_max_sub_group_size() % T
+
+@device_override KI.get_num_sub_groups(::Type{T}) where {T} = get_num_sub_groups() % T
+
+@device_override KI.get_sub_group_id(::Type{T}) where {T} = get_sub_group_id() % T
+
+@device_override KI.get_sub_group_local_id(::Type{T}) where {T} = get_sub_group_local_id() % T
+
 
 ## Shared Memory
 
@@ -216,6 +249,14 @@ end
     barrier(SPIRVIntrinsics.LOCAL_MEM_FENCE | SPIRVIntrinsics.GLOBAL_MEM_FENCE)
 end
 
+@device_override @inline function KI.sub_group_barrier()
+    sub_group_barrier(SPIRVIntrinsics.LOCAL_MEM_FENCE | SPIRVIntrinsics.GLOBAL_MEM_FENCE)
+end
+
+@device_override function KI.shfl_down(val::T, offset::Integer) where T
+    sub_group_shuffle(val, get_sub_group_local_id() + offset)
+end
+
 @device_override @inline function KI._print(args...)
     oneAPI._print(args...)
 end
diff --git a/test/kernelinterface.jl b/test/kernelinterface.jl
index 30f00df2..df0287b5 100644
--- a/test/kernelinterface.jl
+++ b/test/kernelinterface.jl
@@ -5,6 +5,42 @@ include(joinpath(dirname(pathof(KernelInterface)), "..", "test", "testsuite.jl")
 
 Testsuite.testsuite(oneAPIBackend(), oneArray)
 
+function ki_subgroup_kernel(num, sizes, max, id, lane)
+    l = KI.get_local_id()
+    s = KI.get_local_size()
+    i = l.x + (l.y - 1) * s.x
+    @inbounds begin
+        num[i] = KI.get_num_sub_groups()
+        sizes[i] = KI.get_sub_group_size()
+        max[i] = KI.get_max_sub_group_size()
+        id[i] = KI.get_sub_group_id()
+        lane[i] = KI.get_sub_group_local_id()
+    end
+    return
+end
+
+# KernelInterface leaves the formation of sub-groups unspecified; Intel GPUs form them from
+# consecutive linear work-item indices
+@testset "partial sub-groups" begin
+    backend = oneAPIBackend()
+    @test KI.supports_subgroups(backend)
+    width = KI.sub_group_size(backend)
+
+    # a (width+1)x2 work-group is made up of 3 sub-groups, the last one only partially filled
+    workgroupsize = (width + 1, 2)
+    n = prod(workgroupsize)
+    num, sizes, max, id, lane = (oneArray{UInt32}(undef, n) for _ in 1:5)
+    KI.@launch backend workgroupsize=workgroupsize ki_subgroup_kernel(num, sizes, max, id, lane)
+    @test all(==(3), Array(num))
+    @test all(==(width), Array(max))
+    @test Array(sizes) == [i < 2width ? width : 2 for i in 0:n-1]
+    @test Array(id) == [div(i, width) + 1 for i in 0:n-1]
+    @test Array(lane) == [rem(i, width) + 1 for i in 0:n-1]
+
+    # the width is fixed
+    @test_throws ArgumentError KI.@launch backend launch=false sub_group_size=(width == 16 ? 8 : 16) ki_subgroup_kernel(num, sizes, max, id, lane)
+end
+
 function ki_fill!(A)
     i = KI.get_global_id().x
     if i <= length(A)

From 353a168131381a0f956e1eeaf5dbbc17e113364b Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard@gmail.com>
Date: Wed, 30 Sep 2026 18:18:05 +0200
Subject: [PATCH 7/9] Order work across tasks with Level Zero events

Implement `KI.record_event` with an event signaled on the task's stream,
and `KI.wait_event` by waiting for it on the host, cooperatively, like
`synchronize` does. `KernelAbstractions.@spawn` uses them to order a new
task's work after the work its parent had queued, without synchronizing the
parent; the default is a full synchronization.

The wait happens on the host, not by making the waiting task's stream wait
on the device, because a Level Zero event must not be destroyed while a
command list still references it, and the stream would reference it until it
executes the wait. For the same reason, recorded events are kept alive
until they have been signaled, independently of the recording task, which
may end before that.

KernelInterface's testsuite checks them once `record_event` returns an event.

Co-authored-by: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
---
 lib/level-zero/event.jl           |  4 +--
 lib/level-zero/oneL0.jl           |  2 +-
 lib/level-zero/synchronization.jl | 23 +++++++++-----
 src/oneAPIKernels.jl              | 52 +++++++++++++++++++++++++++++++
 4 files changed, 71 insertions(+), 10 deletions(-)

diff --git a/lib/level-zero/event.jl b/lib/level-zero/event.jl
index 92460f2b..d805cf12 100644
--- a/lib/level-zero/event.jl
+++ b/lib/level-zero/event.jl
@@ -36,8 +36,8 @@ mutable struct ZeEvent
     handle::ze_event_handle_t
     pool::ZeEventPool
 
-    function ZeEvent(pool, index::Integer)
-        desc_ref = Ref(ze_event_desc_t(; index=index-1))
+    function ZeEvent(pool, index::Integer; signal=0, wait=0)
+        desc_ref = Ref(ze_event_desc_t(; index=index-1, signal, wait))
         handle_ref = Ref{ze_event_handle_t}()
         zeEventCreate(pool, desc_ref, handle_ref)
         obj = new(handle_ref[], pool)
diff --git a/lib/level-zero/oneL0.jl b/lib/level-zero/oneL0.jl
index c36b9eee..f8e32e8a 100644
--- a/lib/level-zero/oneL0.jl
+++ b/lib/level-zero/oneL0.jl
@@ -124,9 +124,9 @@ end
 include("context.jl")
 include("cmdqueue.jl")
 include("cmdlist.jl")
-include("synchronization.jl")
 include("fence.jl")
 include("event.jl")
+include("synchronization.jl")
 include("barrier.jl")
 include("module.jl")
 include("memory.jl")
diff --git a/lib/level-zero/synchronization.jl b/lib/level-zero/synchronization.jl
index 809a03b8..d5ab74b7 100644
--- a/lib/level-zero/synchronization.jl
+++ b/lib/level-zero/synchronization.jl
@@ -1,14 +1,14 @@
 # cooperative synchronization
 #
-# `zeCommandListHostSynchronize` and `zeCommandQueueSynchronize` block the calling thread
-# until the work has completed, so no other task can run on it in the meantime. As CUDA.jl
+# `zeCommandListHostSynchronize`, `zeCommandQueueSynchronize` and `zeEventHostSynchronize`
+# block the calling thread until the work has completed, so no other task can run on it in the meantime. As CUDA.jl
 # does, first busy-wait on a non-blocking query, which keeps the latency of short
 # operations low, and then block in the driver on a separate thread, while the calling
 # task waits for that thread without blocking the scheduler.
 
 export nonblocking_synchronize
 
-const SyncObject = Union{ZeImmediateCommandList, ZeCommandQueue}
+const SyncObject = Union{ZeImmediateCommandList, ZeCommandQueue, ZeEvent}
 
 # with a zero timeout, a synchronization is a query
 function check_done(res::ze_result_t)
@@ -32,6 +32,14 @@ gcsafe_synchronize(list::ZeImmediateCommandList) =
 gcsafe_synchronize(queue::ZeCommandQueue) =
     @gcsafe_ccall libze_loader.zeCommandQueueSynchronize(
         queue::ze_command_queue_handle_t, typemax(UInt64)::UInt64)::ze_result_t
+gcsafe_synchronize(event::ZeEvent) =
+    @gcsafe_ccall libze_loader.zeEventHostSynchronize(
+        event::ze_event_handle_t, typemax(UInt64)::UInt64)::ze_result_t
+
+# once the work has completed, synchronizing a list or queue doesn't block anymore; a
+# signaled event needs nothing more
+finish_synchronization(obj::Union{ZeImmediateCommandList, ZeCommandQueue}) = synchronize(obj)
+finish_synchronization(::ZeEvent) = nothing
 
 
 ## bidirectional channel
@@ -163,15 +171,16 @@ end
 
 """
     nonblocking_synchronize(list_or_queue)
+    nonblocking_synchronize(event)
 
-Wait for the work on an immediate command list or command queue to complete, like
-[`synchronize`](@ref), but without blocking the Julia scheduler: other tasks keep running
-while this one waits.
+Wait for the work on an immediate command list or command queue to complete, or for an
+event to be signaled, like [`synchronize`](@ref) or `wait`, but without blocking the Julia
+scheduler: other tasks keep running while this one waits.
 """
 function nonblocking_synchronize(obj::SyncObject)
     if spinning_synchronization(Base.isdone, obj)
         # done, so this doesn't block
-        synchronize(obj)
+        finish_synchronization(obj)
         return
     end
 
diff --git a/src/oneAPIKernels.jl b/src/oneAPIKernels.jl
index b1507d10..6cc2c6d3 100644
--- a/src/oneAPIKernels.jl
+++ b/src/oneAPIKernels.jl
@@ -264,6 +264,58 @@ end
 ## COV_EXCL_STOP
 
 
+## Events
+
+# Level Zero has no reusable timeline events, so every `record_event` creates a fresh
+# one-shot event in its own host-visible pool. Its signal has `HOST` scope, the broadest:
+# it flushes caches far enough for the host *and* other devices.
+#
+# `wait_event` waits for the event on the host, cooperatively: making the waiting task's
+# stream wait for it on the device would require keeping the event alive until that stream
+# has executed the wait. Recorded events are kept alive until they have been signaled, as
+# destroying an event that the device still has to signal is undefined; not by the
+# recording task, which may end before that.
+const pending_events = oneAPI.oneL0.ZeEvent[]
+const pending_events_lock = ReentrantLock()
+
+function KI.record_event(::oneAPIBackend)
+    ctx = oneAPI.context()
+    dev = oneAPI.device()
+
+    pool = oneAPI.oneL0.ZeEventPool(ctx, 1;
+                                    flags = oneAPI.oneL0.ZE_EVENT_POOL_FLAG_HOST_VISIBLE)
+    ev = oneAPI.oneL0.ZeEvent(pool, 1;
+                              signal = oneAPI.oneL0.ZE_EVENT_SCOPE_FLAG_HOST,
+                              wait = oneAPI.oneL0.ZE_EVENT_SCOPE_FLAG_HOST)
+
+    @lock pending_events_lock begin
+        filter!(!Base.isdone, pending_events)
+        push!(pending_events, ev)
+    end
+
+    # the task's stream is an in-order immediate command list, so the signal fires once
+    # everything appended before it has completed. `execute!` first drains any pending
+    # oneMKL work on the companion queue (`mkl_wait!`), so that is captured as well.
+    try
+        oneAPI.oneL0.execute!(oneAPI.global_stream(ctx, dev)) do list
+            oneAPI.oneL0.append_signal!(list, ev)
+        end
+    catch
+        # the signal wasn't submitted, so nothing references the event
+        @lock pending_events_lock filter!(!=(ev), pending_events)
+        rethrow()
+    end
+
+    return ev
+end
+
+# the event can come from any device, also one of another driver
+function KI.wait_event(::oneAPIBackend, ev::oneAPI.oneL0.ZeEvent)
+    oneAPI.oneL0.nonblocking_synchronize(ev)
+    return
+end
+
+
 ## Other
 
 function KI.priority!(::oneAPIBackend, prio::Symbol)

From 42a56379e7dcca24e397c016ad3798fc2231709d Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Wed, 30 Sep 2026 18:18:12 +0200
Subject: [PATCH 8/9] Report oneAPI's version info through KernelInterface

`KI.versioninfo(oneAPIBackend())` prints `oneAPI.versioninfo()`.
---
 src/oneAPIKernels.jl    | 2 ++
 test/kernelinterface.jl | 4 ++++
 2 files changed, 6 insertions(+)

diff --git a/src/oneAPIKernels.jl b/src/oneAPIKernels.jl
index 6cc2c6d3..8bc28dbd 100644
--- a/src/oneAPIKernels.jl
+++ b/src/oneAPIKernels.jl
@@ -318,6 +318,8 @@ end
 
 ## Other
 
+KI.versioninfo(io::IO, ::oneAPIBackend) = oneAPI.versioninfo(io)
+
 function KI.priority!(::oneAPIBackend, prio::Symbol)
     if !(prio in (:high, :normal, :low))
         error("priority must be one of :high, :normal, :low")
diff --git a/test/kernelinterface.jl b/test/kernelinterface.jl
index df0287b5..27c01948 100644
--- a/test/kernelinterface.jl
+++ b/test/kernelinterface.jl
@@ -61,3 +61,7 @@ end
     @test_throws ArgumentError kernel(A; ndrange=4, items=8)
     @test_throws ArgumentError kernel(A; ndrange=4, groups=2)
 end
+
+@testset "versioninfo" begin
+    @test occursin("oneAPI.jl", sprint(KI.versioninfo, oneAPIBackend()))
+end

From 914b7629e4c0a9d4d818dc4f8afea68e052d93e5 Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard@gmail.com>
Date: Wed, 30 Sep 2026 18:10:14 +0200
Subject: [PATCH 9/9] [TEMP] Get KernelAbstractions 0.10, KernelInterface 0.4
 and AcceleratedKernels from their main branches

None of them is registered: KernelAbstractions 0.10 and KernelInterface 0.4
come from their development branch, and AcceleratedKernels from its main
branch, whose registered release excludes KernelAbstractions 0.10.
AcceleratedKernels depends on KernelAbstractions, so resolving oneAPI needs
KernelAbstractions from [sources] too, which Pkg only accepts for a regular
dependency: it is one until this commit is dropped. The test project takes
oneAPI from the checkout, so that it doesn't resolve the registered release.

Julia 1.10 ignores [sources], and Julia 1.11 doesn't pick them up for the
test environment, so CI develops all three explicitly there, and the
documentation is built with Julia 1.12.

Drop this commit once they are registered.
---
 .buildkite/pipeline.yml    | 10 ++++++++++
 .github/workflows/docs.yml |  4 +++-
 .gitlab-ci.yml             | 14 ++++++++++++++
 Project.toml               |  7 +++++--
 test/Project.toml          |  5 +++++
 5 files changed, 37 insertions(+), 3 deletions(-)

diff --git a/.buildkite/pipeline.yml b/.buildkite/pipeline.yml
index 7cd9f046..9b275537 100644
--- a/.buildkite/pipeline.yml
+++ b/.buildkite/pipeline.yml
@@ -24,6 +24,16 @@ steps:
         agents:
           queue: "oneapi"
         commands: |
+          # XXX: KernelAbstractions 0.10 and KernelInterface 0.4 aren't registered yet, and
+          #      neither is an AcceleratedKernels that supports them. Julia 1.10 ignores
+          #      [sources], and Julia 1.11 doesn't pick them up for the test environment,
+          #      so check them out and develop them together, as the registered
+          #      AcceleratedKernels excludes KernelAbstractions 0.10.
+          if [[ "{{matrix.julia}}" == "1.10" || "{{matrix.julia}}" == "1.11" ]]; then
+            git clone --depth 1 --branch main https://github.com/JuliaGPU/KernelAbstractions.jl ka
+            git clone --depth 1 https://github.com/JuliaGPU/AcceleratedKernels.jl ak
+            julia --project -e 'using Pkg; Pkg.develop([PackageSpec(; path) for path in ("ka", "ka/lib/KernelInterface", "ak")])'
+          fi
           julia --project=deps deps/build_ci.jl
         if: |
           build.message !~ /\[skip [^\]]*(tests|julia)/ &&
diff --git a/.github/workflows/docs.yml b/.github/workflows/docs.yml
index c853602f..338af859 100644
--- a/.github/workflows/docs.yml
+++ b/.github/workflows/docs.yml
@@ -29,7 +29,9 @@ jobs:
           ssh-strict: false
       - uses: julia-actions/setup-julia@latest
         with:
-          version: 'lts'
+          # XXX: needs Julia 1.11+ to pick up the unregistered KernelAbstractions 0.10 and
+          #      KernelInterface 0.4 from [sources]
+          version: '1.12'
       - uses: julia-actions/cache@v3
       - uses: julia-actions/julia-buildpkg@latest
       - run: julia --project=docs/ docs/make.jl
diff --git a/.gitlab-ci.yml b/.gitlab-ci.yml
index cf54dd92..a9985987 100644
--- a/.gitlab-ci.yml
+++ b/.gitlab-ci.yml
@@ -90,6 +90,18 @@ stages: [build, test, coverage]
   script:
     - !reference [.aurora-env, script]
     - julia --color=yes --project=deps deps/build_local.jl
+    # XXX: KernelAbstractions 0.10 and KernelInterface 0.4 aren't registered yet, nor is an
+    #      AcceleratedKernels that supports them, and Julia 1.10 ignores [sources]: check
+    #      them out and develop them (the checkouts travel to the test job as artifacts,
+    #      like test/Manifest.toml)
+    - |
+      if [ "$JULIA_VERSION" = "1.10" ]; then
+        rm -rf ka ak
+        git clone --depth 1 --branch main https://github.com/JuliaGPU/KernelAbstractions.jl ka
+        git clone --depth 1 https://github.com/JuliaGPU/AcceleratedKernels.jl ak
+        julia --color=yes --project=. -e 'using Pkg; Pkg.develop([PackageSpec(; path) for path in ("ka", "ka/lib/KernelInterface", "ak")])'
+        julia --color=yes --project=test -e 'using Pkg; Pkg.develop([PackageSpec(; path) for path in (".", "ka", "ka/lib/KernelInterface", "ak")])'
+      fi
     # Instantiate (and thereby precompile) both environments here so the 1 h batch job
     # spends its walltime on tests, not on Pkg. Manifests are gitignored, so the test
     # env must be resolved here with oneAPI dev'ed at the checkout (path "..") — a plain
@@ -113,6 +125,8 @@ stages: [build, test, coverage]
       - LocalPreferences.toml
       - test/LocalPreferences.toml
       - test/Manifest.toml
+      - ka
+      - ak
     expire_in: 1 week
 
 .test:lts:
diff --git a/Project.toml b/Project.toml
index 378ea7bb..7420467c 100644
--- a/Project.toml
+++ b/Project.toml
@@ -12,6 +12,7 @@ ExprTools = "e2ba6199-217a-4e67-a87a-7c52f15ade04"
 GPUArrays = "0c68f7d7-f131-5f86-a1c3-88cf8149b2d7"
 GPUCompiler = "61eb1bfa-7361-4325-ad38-22787b887f55"
 GPUToolbox = "096a3bc2-3ced-46d0-87f4-dd12716f4bfc"
+KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c"
 KernelInterface = "4ee993da-d684-4d17-a7dd-4e58e78d92bf"
 LLVM = "929cbde3-209d-540e-8aea-75f648917ca0"
 Libdl = "8f399da3-3557-5675-b5ff-fb832c97cbdb"
@@ -32,8 +33,10 @@ oneAPI_Level_Zero_Headers_jll = "f4bc562b-d309-54f8-9efb-476e56f0410d"
 oneAPI_Level_Zero_Loader_jll = "13eca655-d68d-5b81-8367-6d99d727ab01"
 oneAPI_Support_jll = "b049733a-a71d-5ed3-8eba-7d323ac00b36"
 
-[weakdeps]
-KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c"
+[sources]
+AcceleratedKernels = {url = "https://github.com/JuliaGPU/AcceleratedKernels.jl", rev = "main"}
+KernelAbstractions = {url = "https://github.com/JuliaGPU/KernelAbstractions.jl", rev = "main"}
+KernelInterface = {url = "https://github.com/JuliaGPU/KernelAbstractions.jl", rev = "main", subdir = "lib/KernelInterface"}
 
 [extensions]
 KernelAbstractionsExt = "KernelAbstractions"
diff --git a/test/Project.toml b/test/Project.toml
index fe34ca49..9d8d4da8 100644
--- a/test/Project.toml
+++ b/test/Project.toml
@@ -26,5 +26,10 @@ libigc_jll = "94295238-5935-5bd7-bb0f-b00942e9bdd5"
 oneAPI = "8f75cd03-7ff8-4ecb-9b8f-daf728133b1b"
 oneAPI_Support_jll = "b049733a-a71d-5ed3-8eba-7d323ac00b36"
 
+[sources]
+KernelAbstractions = {url = "https://github.com/JuliaGPU/KernelAbstractions.jl", rev = "main"}
+KernelInterface = {url = "https://github.com/JuliaGPU/KernelAbstractions.jl", rev = "main", subdir = "lib/KernelInterface"}
+oneAPI = {path = ".."}
+
 [compat]
 ParallelTestRunner = "2.8"