From 7a1140be958a23be970699691e74f250f6c41704 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/6] 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 1aca77181310f360e2a3806e5edb5eac0775460d 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 2/6] 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, defaulting to 32.
---
 src/compiler/compilation.jl | 19 ++++++++++++++++---
 src/compiler/execution.jl   |  2 +-
 2 files changed, 17 insertions(+), 4 deletions(-)

diff --git a/src/compiler/compilation.jl b/src/compiler/compilation.jl
index 8f463103..7305e049 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)
 
+    # Set the subgroup size
+    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,16 @@ 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=32, 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 ∉ oneL0.compute_properties(dev).subGroupSizes
+        @error("$sub_group_size is not a valid sub-group size for this device.")
+    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 +324,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 20742c2d..1f499ddf 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]
 
 """

From f9acd84121c2b4abd9d6c1a2b17a414438540c15 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Sat, 22 Aug 2026 15:34:15 -0300
Subject: [PATCH 3/6] Add a KernelInterface back end

Implement KernelInterface next to the KernelAbstractions back end, which
moves to oneAPIKernelsOld.jl until the port to KernelAbstractions 0.10.

The per-dimension launch limits (`max_work_group_dims`, `max_num_groups`)
come from the Level Zero compute properties. Every auto-sized launch
queries them, so they are cached per device.

The KernelInterface test suite runs as the `kernelinterface` test. Its
events test is skipped when every submission synchronizes (the Aurora LTS
workaround, ONEAPI_SYNC_EACH_SUBMISSION), as the ordering it checks is not
observable then.

Co-authored-by: Tim Besard <tim.besard@gmail.com>
---
 Project.toml            |   2 +
 src/oneAPI.jl           |  11 +-
 src/oneAPIKernels.jl    | 248 +++++++++++++++--------------------
 src/oneAPIKernelsOld.jl | 282 ++++++++++++++++++++++++++++++++++++++++
 src/utils.jl            |   2 +-
 test/Project.toml       |   1 +
 test/kernelinterface.jl |  11 ++
 7 files changed, 409 insertions(+), 148 deletions(-)
 create mode 100644 src/oneAPIKernelsOld.jl
 create mode 100644 test/kernelinterface.jl

diff --git a/Project.toml b/Project.toml
index 88b57ba4..e21f1f71 100644
--- a/Project.toml
+++ b/Project.toml
@@ -13,6 +13,7 @@ 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"
@@ -42,6 +43,7 @@ GPUArrays = "11.5.14"
 GPUCompiler = "2.9"
 GPUToolbox = "3.1"
 KernelAbstractions = "0.9.39"
+KernelInterface = "0.2.3"
 LLVM = "6, 7, 8, 9"
 NEO_jll = "=26.18.38308"
 PrecompileTools = "1"
diff --git a/src/oneAPI.jl b/src/oneAPI.jl
index 6b6db42c..66b42a2a 100644
--- a/src/oneAPI.jl
+++ b/src/oneAPI.jl
@@ -13,6 +13,7 @@ using SpecialFunctions
 import Preferences
 
 import KernelAbstractions: KernelAbstractions
+import KernelInterface
 
 using LLVM
 using LLVM.Interop
@@ -77,12 +78,18 @@ include("gpuarrays.jl")
 include("random.jl")
 include("utils.jl")
 
-include("oneAPIKernels.jl")
+# KernelAbstractions
+include("oneAPIKernelsOld.jl")
 import .oneAPIKernels: oneAPIBackend
+export oneAPIBackend
+
+# KernelInterface
+include("oneAPIKernels.jl")
+import .oneAPIInterface
+
 include("accumulate.jl")
 include("sorting.jl")
 include("indexing.jl")
-export oneAPIBackend
 
 # precompilation workload (warms up the SPIR-V compilation pipeline)
 include("compiler/precompile.jl")
diff --git a/src/oneAPIKernels.jl b/src/oneAPIKernels.jl
index 7b90d2ca..f38be558 100644
--- a/src/oneAPIKernels.jl
+++ b/src/oneAPIKernels.jl
@@ -1,55 +1,46 @@
-module oneAPIKernels
+module oneAPIInterface
 
 using ..oneAPI
-using ..oneAPI: @device_override, SPIRVIntrinsics, method_table
+using ..oneAPI: @device_override, SPIRVIntrinsics, method_table, kernel_convert, zefunction
 
-import KernelAbstractions as KA
+import KernelInterface as KI
 
 import StaticArrays
 
-import Adapt
-
-
 ## Back-end Definition
 
 export oneAPIBackend
 
-struct oneAPIBackend <: KA.GPU
+struct oneAPIBackend <: KI.GPU
     prefer_blocks::Bool
     always_inline::Bool
 end
 
+KI.versioninfo(io::IO, ::oneAPIBackend) = oneAPI.versioninfo(io)
+
 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)
+@inline KI.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 KI.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()
+KI.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
+KI.synchronize(::oneAPIBackend) = oneAPI.oneL0.synchronize()
+KI.supports_float64(::oneAPIBackend) = false  # TODO: Check if this is device dependent
+KI.supports_unified(::oneAPIBackend) = true
 
-KA.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)
+KI.functional(::oneAPIBackend) = oneAPI.functional()
 
 # sparse arrays (oneMKL is only available on Linux)
 @static if Sys.islinux()
-    import GPUArrays, SparseArrays
-    KA.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)
+    KI.get_backend(::oneAPI.oneMKL.oneAbstractSparseMatrix) = oneAPIBackend()
 end
 
 
 ## Memory Operations
 
-function KA.copyto!(::oneAPIBackend, A, B)
+function KI.copyto!(::oneAPIBackend, A, B)
     copyto!(A, B)
     # TODO: Address device to host copies in jl being synchronizing
 end
@@ -57,173 +48,135 @@ end
 
 ## Device Operations
 
-function KA.ndevices(::oneAPIBackend)
+function KI.ndevices(::oneAPIBackend)
     return length(oneAPI.devices())
 end
 
-function KA.device(::oneAPIBackend)::Int
+function KI.device(::oneAPIBackend)::Int
     dev = oneAPI.device()
     devs = oneAPI.devices()
     idx = findfirst(==(dev), devs)
     return idx === nothing ? 1 : idx
 end
 
-function KA.device!(backend::oneAPIBackend, id::Int)
+function KI.device!(backend::oneAPIBackend, id::Int)
     return oneAPI.device!(id)
 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
+KI.argconvert(::oneAPIBackend, arg) = kernel_convert(arg)
 
-function KA.launch_config(kernel::KA.Kernel{oneAPIBackend}, ndrange, workgroupsize)
-    if ndrange isa Integer
-        ndrange = (ndrange,)
-    end
-    if workgroupsize isa Integer
-        workgroupsize = (workgroupsize, )
-    end
+function KI.kernel_function(::oneAPIBackend, f::F, tt::TT=Tuple{}; name = nothing, kwargs...) where {F,TT}
+    kern = zefunction(f, tt; name, kwargs...)
+    KI.Kernel{oneAPIBackend, typeof(kern)}(oneAPIBackend(), kern)
+end
 
-    # partition checked that the ndrange's agreed
-    if KA.ndrange(kernel) <: KA.StaticSize
-        ndrange = nothing
-    end
+function (obj::KI.Kernel{oneAPIBackend})(args...; numworkgroups=(), workgroupsize=(), ndrange=(), max_work_group_size=typemax(Int))
+    KI.check_launch_args(numworkgroups, workgroupsize, ndrange)
+    prod(ndrange) == 0 && return nothing
 
-    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
+    numworkgroups, workgroupsize = KI.auto_launch_sizes(obj, numworkgroups, workgroupsize, ndrange, max_work_group_size)
+    items = (workgroupsize..., ntuple(_ -> 1, 3 - length(workgroupsize))...)
+    groups = (numworkgroups..., ntuple(_ -> 1, 3 - length(numworkgroups))...)
 
-    return ndrange, workgroupsize, iterspace, dynamic
+    obj.kern(args...; items, groups)
+    return nothing
 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
+function KI.kernel_max_work_group_size(kernel::KI.Kernel{<:oneAPIBackend}; max_work_items::Int=typemax(Int))::Int
+    group_size = oneAPI.launch_configuration(kernel.kern)
+    Int(min(group_size, max_work_items))
+end
+# querying the device allocates, so cache the limits that every auto-sized launch needs
+const DeviceLimits = @NamedTuple{max_work_group_size::Int, max_work_group_dims::NTuple{3, Int},
+                                 max_num_groups::NTuple{3, Int}}
+function device_limits()
+    dev = device()::oneAPI.oneL0.ZeDevice
+    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)
+        (; 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))
     end
 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)))
+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
+function KI.sub_group_size(::oneAPIBackend)::Int
+    sg_sizes = oneAPI.oneL0.compute_properties(device()).subGroupSizes
+    if 32 in sg_sizes
+        return 32
+    elseif 64 in sg_sizes
+        return 64
+    elseif 16 in sg_sizes
+        return 16
     else
-        # maxthreads = nothing
-    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
-        end
-
-        workgroupsize = threads_to_workgroupsize(items, ndrange)
-        iterspace, dynamic = KA.partition(obj, ndrange, workgroupsize)
-        ctx = KA.mkcontext(obj, ndrange, iterspace)
+        return 1
     end
+end
+function KI.multiprocessor_count(::oneAPIBackend)::Int
+    oneAPI.oneL0.properties(device()).numSlices
+end
 
-    groups = length(KA.blocks(iterspace))
-    items = length(KA.workitems(iterspace))
-
-    if groups == 0
-        return nothing
-    end
+function KI.shfl_down_types(::oneAPIBackend)
+    res = copy(SPIRVIntrinsics.gentypes)
 
-    # Launch kernel
-    kernel(ctx, args...; items, groups)
+    res = setdiff(res, [Float64])
 
-    return nothing
+    return res
 end
 
-
 ## Indexing Functions
-
-@device_override @inline function KA.__index_Local_Linear(ctx)
-    return get_local_id()
+## COV_EXCL_START
+@device_override @inline function KI.get_local_id(::Type{T}) where {T}
+    return (; x = T(get_local_id(1)), y = T(get_local_id(2)), z = T(get_local_id(3)))
 end
 
-@device_override @inline function KA.__index_Group_Linear(ctx)
-    return get_group_id()
+@device_override @inline function KI.get_group_id(::Type{T}) where {T}
+    return (; x = T(get_group_id(1)), y = T(get_group_id(2)), z = T(get_group_id(3)))
 end
 
-@device_override @inline function KA.__index_Global_Linear(ctx)
-    return get_global_id()
+@device_override @inline function KI.get_global_id(::Type{T}) where {T}
+    return (; x = T(get_global_id(1)), y = T(get_global_id(2)), z = T(get_global_id(3)))
 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 = T(get_local_size(1)), y = T(get_local_size(2)), z = T(get_local_size(3)))
 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 = T(get_num_groups(1)), y = T(get_num_groups(2)), z = T(get_num_groups(3)))
 end
 
-@device_override @inline function KA.__index_Global_Cartesian(ctx)
-    return @inbounds KA.expand(KA.__iterspace(ctx), get_group_id(), get_local_id())
+@device_override @inline function KI.get_global_size(::Type{T}) where {T}
+    return (; x = T(get_global_size(1)), y = T(get_global_size(2)), z = T(get_global_size(3)))
 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
+@device_override KI.get_sub_group_size() = get_sub_group_size() % UInt32
+
+@device_override KI.get_max_sub_group_size() = get_max_sub_group_size() % UInt32
+
+@device_override KI.get_num_sub_groups() = get_num_sub_groups() % UInt32
 
+@device_override KI.get_sub_group_id() = get_sub_group_id() % UInt32
+
+@device_override KI.get_sub_group_local_id() = get_sub_group_local_id() % UInt32
 
 ## 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
@@ -233,18 +186,23 @@ end
     barrier(SPIRVIntrinsics.LOCAL_MEM_FENCE | SPIRVIntrinsics.GLOBAL_MEM_FENCE)
 end
 
-@device_override @inline function KA.__print(args...)
-    oneAPI._print(args...)
+@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
 
-## Other
+@device_override @inline function KI._print(args...)
+    oneAPI._print(args...)
+end
 
-Adapt.adapt_storage(to::KA.ConstAdaptor, a::oneDeviceArray) = Base.Experimental.Const(a)
+## COV_EXCL_STOP
 
-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/oneAPIKernelsOld.jl b/src/oneAPIKernelsOld.jl
new file mode 100644
index 00000000..854e8ab4
--- /dev/null
+++ b/src/oneAPIKernelsOld.jl
@@ -0,0 +1,282 @@
+module oneAPIKernels
+
+using ..oneAPI
+using ..oneAPI: @device_override, SPIRVIntrinsics, method_table
+
+import KernelAbstractions as KA
+
+import StaticArrays
+
+import Adapt
+
+
+## Back-end Definition
+
+export oneAPIBackend
+
+struct oneAPIBackend <: KA.GPU
+    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))
+
+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
+
+KA.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()
+    # 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)
+    copyto!(A, B)
+    # TODO: Address device to host copies in jl being synchronizing
+end
+
+
+## Device Operations
+
+function KA.ndevices(::oneAPIBackend)
+    return length(oneAPI.devices())
+end
+
+function KA.device(::oneAPIBackend)::Int
+    dev = oneAPI.device()
+    devs = oneAPI.devices()
+    idx = findfirst(==(dev), devs)
+    return idx === nothing ? 1 : idx
+end
+
+function KA.device!(backend::oneAPIBackend, id::Int)
+    return oneAPI.device!(id)
+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
+
+    # 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
+
+    return ndrange, workgroupsize, iterspace, dynamic
+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
+    end
+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)))
+    else
+        # maxthreads = nothing
+    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
+        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
+    end
+
+    # Launch kernel
+    kernel(ctx, args...; items, groups)
+
+    return nothing
+end
+
+
+## Indexing Functions
+
+@device_override @inline function KA.__index_Local_Linear(ctx)
+    return get_local_id()
+end
+
+@device_override @inline function KA.__index_Group_Linear(ctx)
+    return get_group_id()
+end
+
+@device_override @inline function KA.__index_Global_Linear(ctx)
+    return get_global_id()
+end
+
+@device_override @inline function KA.__index_Local_Cartesian(ctx)
+    @inbounds KA.workitems(KA.__iterspace(ctx))[get_local_id()]
+end
+
+@device_override @inline function KA.__index_Group_Cartesian(ctx)
+    @inbounds KA.blocks(KA.__iterspace(ctx))[get_group_id()]
+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 and Scratch Memory
+
+@device_override @inline function KA.SharedMemory(::Type{T}, ::Val{Dims}, ::Val{Id}) where {T, Dims, Id}
+    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()
+    # 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
+    # treats as ordering *no* memory — so shared-local or global writes are not guaranteed
+    # visible to other work-items after the barrier. `LOCAL_MEM_FENCE | GLOBAL_MEM_FENCE`
+    # ORs in the WorkgroupMemory/CrossWorkgroupMemory fence bits.
+    barrier(SPIRVIntrinsics.LOCAL_MEM_FENCE | SPIRVIntrinsics.GLOBAL_MEM_FENCE)
+end
+
+@device_override @inline function KA.__print(args...)
+    oneAPI._print(args...)
+end
+
+
+## Other
+
+Adapt.adapt_storage(to::KA.ConstAdaptor, a::oneDeviceArray) = Base.Experimental.Const(a)
+
+KA.argconvert(::KA.Kernel{oneAPIBackend}, arg) = kernel_convert(arg)
+
+function KA.priority!(::oneAPIBackend, prio::Symbol)
+    if !(prio in (:high, :normal, :low))
+        error("priority must be one of :high, :normal, :low")
+    end
+
+    priority_enum = if prio == :high
+        oneAPI.oneL0.ZE_COMMAND_QUEUE_PRIORITY_PRIORITY_HIGH
+    elseif prio == :low
+        oneAPI.oneL0.ZE_COMMAND_QUEUE_PRIORITY_PRIORITY_LOW
+    else
+        oneAPI.oneL0.ZE_COMMAND_QUEUE_PRIORITY_NORMAL
+    end
+
+    ctx = oneAPI.context()
+    dev = oneAPI.device()
+
+    # drain the task's current stream before swapping it out, so operations submitted
+    # to the new stream cannot overtake in-flight work on the old one
+    oneAPI.oneL0.synchronize(oneAPI.global_stream(ctx, dev))
+
+    # Replace the stream in task_local_storage. `create_stream` registers the
+    # replacement so `synchronize_all_streams`/`release` can drain it before freeing a
+    # buffer whose in-flight work it references; otherwise all work after `priority!`
+    # runs on an unregistered stream and a freed buffer can be reused while its kernel
+    # is still running (use-after-free → banned context on the LTS NEO stack). The old
+    # stream stays registered until its task dies, like replaced queues before it.
+    new_stream = oneAPI.create_stream(ctx, dev, priority_enum)
+    task_local_storage((:oneStream, ctx, dev), new_stream)
+
+    # the cached SYCL queue wraps the old stream's companion queue; drop it so the next
+    # oneMKL call recreates it against the new stream (the old one was just drained)
+    delete!(task_local_storage(), (:SYCLQueue, ctx, dev))
+
+    return nothing
+end
+
+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/kernelinterface.jl b/test/kernelinterface.jl
new file mode 100644
index 00000000..0854bcfe
--- /dev/null
+++ b/test/kernelinterface.jl
@@ -0,0 +1,11 @@
+import KernelInterface
+using oneAPI.oneAPIInterface
+
+include(joinpath(dirname(pathof(KernelInterface)), "..", "test", "testsuite.jl"))
+
+skip_tests = Set{String}()
+# the events test checks that a waiting task blocks on work queued elsewhere, which is not
+# observable when every submission synchronizes (the Aurora LTS workaround)
+oneAPI.oneL0.sync_each_submission() && push!(skip_tests, "Events")
+
+Testsuite.testsuite(oneAPIInterface.oneAPIBackend, "oneAPI", oneAPI, oneArray, oneAPI.oneDeviceArray; skip_tests)

From 05c27b1d9ca1a647c8d89b879e9f78ed66b83ae0 Mon Sep 17 00:00:00 2001
From: Simeon David Schaub <simeon@schaub.rocks>
Date: Wed, 10 Jun 2026 07:53:29 +0200
Subject: [PATCH 4/6] Add 'Random' to skip_tests in kernelabstractions

Needed to make https://github.com/JuliaGPU/KernelAbstractions.jl/pull/659 pass CI. We should eventually add RNG support to oneAPI by using a similar approach to OpenCL.jl, but lets skip these tests for now
---
 test/kernelabstractions.jl | 1 +
 1 file changed, 1 insertion(+)

diff --git a/test/kernelabstractions.jl b/test/kernelabstractions.jl
index 48bbe1cb..591fd426 100644
--- a/test/kernelabstractions.jl
+++ b/test/kernelabstractions.jl
@@ -4,5 +4,6 @@ include(joinpath(dirname(pathof(KernelAbstractions)), "..", "test", "testsuite.j
 skip_tests=Set([
     "sparse",
     "Convert", # Need to opt out of i128
+    "Random",
 ])
 Testsuite.testsuite(oneAPIBackend, "oneAPI", oneAPI, oneArray, oneDeviceArray; skip_tests)

From b0f3f34cc420884de7d9b597e6769cc1a79fef39 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Thu, 3 Sep 2026 13:11:00 -0300
Subject: [PATCH 5/6] Port to KernelAbstractions 0.10

KernelAbstractions 0.10 builds on KernelInterface, so the KernelInterface
back end becomes `oneAPIBackend` and the old KernelAbstractions back end
is removed. What remains specific to KernelAbstractions (launching
`@kernel` functions, `__validindex`, scratchpads, adapting to the CPU)
moves into a package extension.

Adapting to the back end (`adapt(oneAPIBackend(), x)`) is now defined next
to the back end, as KernelAbstractions 0.10 requires of every back end.

Since KernelAbstractions 0.10, an ndrange may contain ranges, so the
preliminary and the tuned workgroup size, as well as the number of items
used by the prefer_blocks heuristic, are computed from the number of
indices along each dimension rather than from the raw ndrange. An empty
ndrange returns before compiling.

The KernelAbstractions tests that only apply to the CPU back end are
skipped.

Co-authored-by: Tim Besard <tim.besard@gmail.com>
---
 Project.toml                 |   8 +-
 ext/KernelAbstractionsExt.jl | 145 ++++++++++++++++++
 src/oneAPI.jl                |   9 +-
 src/oneAPIKernels.jl         |  10 +-
 src/oneAPIKernelsOld.jl      | 282 -----------------------------------
 test/kernelabstractions.jl   |   2 +
 test/kernelinterface.jl      |   3 +-
 7 files changed, 166 insertions(+), 293 deletions(-)
 create mode 100644 ext/KernelAbstractionsExt.jl
 delete mode 100644 src/oneAPIKernelsOld.jl

diff --git a/Project.toml b/Project.toml
index e21f1f71..41b48c33 100644
--- a/Project.toml
+++ b/Project.toml
@@ -33,6 +33,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"
@@ -42,7 +48,7 @@ ExprTools = "0.1"
 GPUArrays = "11.5.14"
 GPUCompiler = "2.9"
 GPUToolbox = "3.1"
-KernelAbstractions = "0.9.39"
+KernelAbstractions = "0.10"
 KernelInterface = "0.2.3"
 LLVM = "6, 7, 8, 9"
 NEO_jll = "=26.18.38308"
diff --git a/ext/KernelAbstractionsExt.jl b/ext/KernelAbstractionsExt.jl
new file mode 100644
index 00000000..4b901977
--- /dev/null
+++ b/ext/KernelAbstractionsExt.jl
@@ -0,0 +1,145 @@
+module KernelAbstractionsExt
+
+using oneAPI
+using oneAPI: @device_override, SPIRVIntrinsics, method_table
+
+import KernelAbstractions as KA
+
+import StaticArrays
+
+import Adapt
+
+
+Adapt.adapt_storage(::KA.CPU, a::oneArray) = convert(Array, a)
+@static if Sys.islinux()
+    import SparseArrays
+    Adapt.adapt_storage(::KA.CPU, a::oneAPI.oneMKL.oneAbstractSparseMatrix) = SparseArrays.SparseMatrixCSC(a)
+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
+
+# the number of indices along each dimension of `ndrange`, which may contain ranges
+extents(ndrange) = size(CartesianIndices(ndrange))
+
+function KA.launch_config(kernel::KA.Kernel{oneAPIBackend}, ndrange, workgroupsize)
+    if ndrange isa Integer
+        ndrange = (ndrange,)
+    end
+    if workgroupsize isa Integer
+        workgroupsize = (workgroupsize, )
+    end
+
+    # 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.(extents(ndrange), 1))
+    else
+        KA.partition(kernel, ndrange, workgroupsize)
+    end
+
+    return ndrange, workgroupsize, iterspace, dynamic
+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
+    end
+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)
+    # nothing to launch (or compile) for an empty ndrange
+    length(KA.blocks(iterspace)) == 0 && return nothing
+
+    # 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)))
+    else
+        # maxthreads = nothing
+    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(extents(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
+        end
+
+        workgroupsize = threads_to_workgroupsize(items, extents(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))
+
+    # Launch kernel
+    kernel(ctx, args...; items, groups)
+
+    return nothing
+end
+
+
+## Indexing Functions
+
+@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
+
+
+## Scratch Memory
+
+@device_override @inline function KA.Scratchpad(ctx, ::Type{T}, ::Val{Dims}) where {T, Dims}
+    StaticArrays.MArray{KA.__size(Dims), T}(undef)
+end
+
+## Other
+
+Adapt.adapt_storage(to::KA.ConstAdaptor, a::oneDeviceArray) = Base.Experimental.Const(a)
+
+KA.argconvert(::KA.Kernel{oneAPIBackend}, arg) = kernel_convert(arg)
+
+end
diff --git a/src/oneAPI.jl b/src/oneAPI.jl
index 66b42a2a..7c833ff4 100644
--- a/src/oneAPI.jl
+++ b/src/oneAPI.jl
@@ -12,7 +12,6 @@ using SpecialFunctions
 
 import Preferences
 
-import KernelAbstractions: KernelAbstractions
 import KernelInterface
 
 using LLVM
@@ -78,14 +77,10 @@ include("gpuarrays.jl")
 include("random.jl")
 include("utils.jl")
 
-# KernelAbstractions
-include("oneAPIKernelsOld.jl")
-import .oneAPIKernels: oneAPIBackend
-export oneAPIBackend
-
 # KernelInterface
 include("oneAPIKernels.jl")
-import .oneAPIInterface
+import .oneAPIKernels: oneAPIBackend
+export oneAPIBackend
 
 include("accumulate.jl")
 include("sorting.jl")
diff --git a/src/oneAPIKernels.jl b/src/oneAPIKernels.jl
index f38be558..c499b5f9 100644
--- a/src/oneAPIKernels.jl
+++ b/src/oneAPIKernels.jl
@@ -1,4 +1,4 @@
-module oneAPIInterface
+module oneAPIKernels
 
 using ..oneAPI
 using ..oneAPI: @device_override, SPIRVIntrinsics, method_table, kernel_convert, zefunction
@@ -7,6 +7,8 @@ import KernelInterface as KI
 
 import StaticArrays
 
+import Adapt
+
 ## Back-end Definition
 
 export oneAPIBackend
@@ -32,9 +34,15 @@ KI.supports_unified(::oneAPIBackend) = true
 
 KI.functional(::oneAPIBackend) = oneAPI.functional()
 
+Adapt.adapt_storage(::oneAPIBackend, a::AbstractArray) = Adapt.adapt(oneArray, a)
+Adapt.adapt_storage(::oneAPIBackend, a::oneArray) = a
+
 # sparse arrays (oneMKL is only available on Linux)
 @static if Sys.islinux()
+    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
 end
 
 
diff --git a/src/oneAPIKernelsOld.jl b/src/oneAPIKernelsOld.jl
deleted file mode 100644
index 854e8ab4..00000000
--- a/src/oneAPIKernelsOld.jl
+++ /dev/null
@@ -1,282 +0,0 @@
-module oneAPIKernels
-
-using ..oneAPI
-using ..oneAPI: @device_override, SPIRVIntrinsics, method_table
-
-import KernelAbstractions as KA
-
-import StaticArrays
-
-import Adapt
-
-
-## Back-end Definition
-
-export oneAPIBackend
-
-struct oneAPIBackend <: KA.GPU
-    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))
-
-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
-
-KA.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()
-    # 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)
-    copyto!(A, B)
-    # TODO: Address device to host copies in jl being synchronizing
-end
-
-
-## Device Operations
-
-function KA.ndevices(::oneAPIBackend)
-    return length(oneAPI.devices())
-end
-
-function KA.device(::oneAPIBackend)::Int
-    dev = oneAPI.device()
-    devs = oneAPI.devices()
-    idx = findfirst(==(dev), devs)
-    return idx === nothing ? 1 : idx
-end
-
-function KA.device!(backend::oneAPIBackend, id::Int)
-    return oneAPI.device!(id)
-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
-
-    # 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
-
-    return ndrange, workgroupsize, iterspace, dynamic
-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
-    end
-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)))
-    else
-        # maxthreads = nothing
-    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
-        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
-    end
-
-    # Launch kernel
-    kernel(ctx, args...; items, groups)
-
-    return nothing
-end
-
-
-## Indexing Functions
-
-@device_override @inline function KA.__index_Local_Linear(ctx)
-    return get_local_id()
-end
-
-@device_override @inline function KA.__index_Group_Linear(ctx)
-    return get_group_id()
-end
-
-@device_override @inline function KA.__index_Global_Linear(ctx)
-    return get_global_id()
-end
-
-@device_override @inline function KA.__index_Local_Cartesian(ctx)
-    @inbounds KA.workitems(KA.__iterspace(ctx))[get_local_id()]
-end
-
-@device_override @inline function KA.__index_Group_Cartesian(ctx)
-    @inbounds KA.blocks(KA.__iterspace(ctx))[get_group_id()]
-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 and Scratch Memory
-
-@device_override @inline function KA.SharedMemory(::Type{T}, ::Val{Dims}, ::Val{Id}) where {T, Dims, Id}
-    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()
-    # 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
-    # treats as ordering *no* memory — so shared-local or global writes are not guaranteed
-    # visible to other work-items after the barrier. `LOCAL_MEM_FENCE | GLOBAL_MEM_FENCE`
-    # ORs in the WorkgroupMemory/CrossWorkgroupMemory fence bits.
-    barrier(SPIRVIntrinsics.LOCAL_MEM_FENCE | SPIRVIntrinsics.GLOBAL_MEM_FENCE)
-end
-
-@device_override @inline function KA.__print(args...)
-    oneAPI._print(args...)
-end
-
-
-## Other
-
-Adapt.adapt_storage(to::KA.ConstAdaptor, a::oneDeviceArray) = Base.Experimental.Const(a)
-
-KA.argconvert(::KA.Kernel{oneAPIBackend}, arg) = kernel_convert(arg)
-
-function KA.priority!(::oneAPIBackend, prio::Symbol)
-    if !(prio in (:high, :normal, :low))
-        error("priority must be one of :high, :normal, :low")
-    end
-
-    priority_enum = if prio == :high
-        oneAPI.oneL0.ZE_COMMAND_QUEUE_PRIORITY_PRIORITY_HIGH
-    elseif prio == :low
-        oneAPI.oneL0.ZE_COMMAND_QUEUE_PRIORITY_PRIORITY_LOW
-    else
-        oneAPI.oneL0.ZE_COMMAND_QUEUE_PRIORITY_NORMAL
-    end
-
-    ctx = oneAPI.context()
-    dev = oneAPI.device()
-
-    # drain the task's current stream before swapping it out, so operations submitted
-    # to the new stream cannot overtake in-flight work on the old one
-    oneAPI.oneL0.synchronize(oneAPI.global_stream(ctx, dev))
-
-    # Replace the stream in task_local_storage. `create_stream` registers the
-    # replacement so `synchronize_all_streams`/`release` can drain it before freeing a
-    # buffer whose in-flight work it references; otherwise all work after `priority!`
-    # runs on an unregistered stream and a freed buffer can be reused while its kernel
-    # is still running (use-after-free → banned context on the LTS NEO stack). The old
-    # stream stays registered until its task dies, like replaced queues before it.
-    new_stream = oneAPI.create_stream(ctx, dev, priority_enum)
-    task_local_storage((:oneStream, ctx, dev), new_stream)
-
-    # the cached SYCL queue wraps the old stream's companion queue; drop it so the next
-    # oneMKL call recreates it against the new stream (the old one was just drained)
-    delete!(task_local_storage(), (:SYCLQueue, ctx, dev))
-
-    return nothing
-end
-
-end
diff --git a/test/kernelabstractions.jl b/test/kernelabstractions.jl
index 591fd426..b1627afa 100644
--- a/test/kernelabstractions.jl
+++ b/test/kernelabstractions.jl
@@ -5,5 +5,7 @@ skip_tests=Set([
     "sparse",
     "Convert", # Need to opt out of i128
     "Random",
+    "CPU synchronization",
+    "fallback test: callable types"
 ])
 Testsuite.testsuite(oneAPIBackend, "oneAPI", oneAPI, oneArray, oneDeviceArray; skip_tests)
diff --git a/test/kernelinterface.jl b/test/kernelinterface.jl
index 0854bcfe..58c15d6c 100644
--- a/test/kernelinterface.jl
+++ b/test/kernelinterface.jl
@@ -1,5 +1,4 @@
 import KernelInterface
-using oneAPI.oneAPIInterface
 
 include(joinpath(dirname(pathof(KernelInterface)), "..", "test", "testsuite.jl"))
 
@@ -8,4 +7,4 @@ skip_tests = Set{String}()
 # observable when every submission synchronizes (the Aurora LTS workaround)
 oneAPI.oneL0.sync_each_submission() && push!(skip_tests, "Events")
 
-Testsuite.testsuite(oneAPIInterface.oneAPIBackend, "oneAPI", oneAPI, oneArray, oneAPI.oneDeviceArray; skip_tests)
+Testsuite.testsuite(oneAPIBackend, "oneAPI", oneAPI, oneArray, oneAPI.oneDeviceArray; skip_tests)

From 27c95d13f72e7da367de112d29cc6cd798365b06 Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard@gmail.com>
Date: Mon, 28 Sep 2026 07:09:58 +0200
Subject: [PATCH 6/6] TEMP: test against KernelAbstractions and
 AcceleratedKernels main

KernelAbstractions 0.10 is unreleased, and the released AcceleratedKernels does not support it yet.
---
 .buildkite/pipeline.yml | 4 ++++
 1 file changed, 4 insertions(+)

diff --git a/.buildkite/pipeline.yml b/.buildkite/pipeline.yml
index 7cd9f046..22a0fd11 100644
--- a/.buildkite/pipeline.yml
+++ b/.buildkite/pipeline.yml
@@ -25,6 +25,8 @@ steps:
           queue: "oneapi"
         commands: |
           julia --project=deps deps/build_ci.jl
+          # TEMP: test against KernelAbstractions 0.10 and an AcceleratedKernels supporting it
+          julia --project -e 'using Pkg; Pkg.add([PackageSpec(url="https://github.com/JuliaGPU/AcceleratedKernels.jl", rev="main"), PackageSpec(url="https://github.com/JuliaGPU/KernelAbstractions.jl", rev="main")])'
         if: |
           build.message !~ /\[skip [^\]]*(tests|julia)/ &&
           (build.message =~ /\[only [^\]]*(tests|julia)/ || build.message !~ /\[only /)
@@ -75,6 +77,8 @@ steps:
                 - examples
         command: |
           julia --project=deps deps/build_ci.jl
+          # TEMP: test against KernelAbstractions 0.10 and an AcceleratedKernels supporting it
+          julia --project -e 'using Pkg; Pkg.add([PackageSpec(url="https://github.com/JuliaGPU/AcceleratedKernels.jl", rev="main"), PackageSpec(url="https://github.com/JuliaGPU/KernelAbstractions.jl", rev="main")])'
           julia --project -e '
               # use debug JLLs, for asserts + better backtraces
               using oneAPI