From 24aa1f25c1ad4b6213dd194927d18ed11da8b772 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Wed, 19 Aug 2026 11:28:59 -0300
Subject: [PATCH 1/3] Add a KernelInterface back end

Implement the KernelInterface (KI) API for OpenCL.jl in an internal
`OpenCLInterface` module, reachable through `KI.get_backend` on a CLArray.
The existing KernelAbstractions 0.9 back end moves unchanged to
OpenCLKernelsOld.jl. Run the KI testsuite as part of the test suite.

The device's work-group limits, including the per-dimension limits that
KI 0.2.3 uses to bound automatically chosen workgroup sizes (NVIDIA, for
example, only allows 64 work-items along the third dimension), are
cached per device because querying them allocates on every launch.

Co-authored-by: Tim Besard <tim.besard@gmail.com>
---
 Project.toml            |   2 +
 src/OpenCL.jl           |   9 +-
 src/OpenCLKernels.jl    | 218 ++++++++++++++++++-------------------
 src/OpenCLKernelsOld.jl | 232 ++++++++++++++++++++++++++++++++++++++++
 src/util.jl             |   2 +-
 test/Project.toml       |   1 +
 test/kernelinterface.jl |   6 ++
 7 files changed, 352 insertions(+), 118 deletions(-)
 create mode 100644 src/OpenCLKernelsOld.jl
 create mode 100644 test/kernelinterface.jl

diff --git a/Project.toml b/Project.toml
index 0236d89d..d74ddf83 100644
--- a/Project.toml
+++ b/Project.toml
@@ -11,6 +11,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"
 LinearAlgebra = "37e2e46d-f89d-539d-b4ee-838fcccc9c8e"
 OpenCL_jll = "6cb37087-e8b6-5417-8430-1f242f1e46e4"
@@ -35,6 +36,7 @@ GPUArrays = "11.2.1"
 GPUCompiler = "2.9"
 GPUToolbox = "3.1"
 KernelAbstractions = "0.9.38"
+KernelInterface = "0.2.3"
 LLVM = "9.6"
 LinearAlgebra = "1"
 OpenCL_jll = "=2024.10.24"
diff --git a/src/OpenCL.jl b/src/OpenCL.jl
index da994572..ca02eb1b 100644
--- a/src/OpenCL.jl
+++ b/src/OpenCL.jl
@@ -12,6 +12,8 @@ using Preferences
 
 import KernelAbstractions: KernelAbstractions
 
+import KernelInterface
+
 using Core: LLVMPtr
 
 # library wrappers
@@ -49,7 +51,12 @@ include("mapreduce.jl")
 include("gpuarrays.jl")
 include("random.jl")
 
-include("OpenCLKernels.jl")
+include("OpenCLKernelsOld.jl")
 import .OpenCLKernels: OpenCLBackend
 export OpenCLBackend
+
+# KernelInterface - NOT PUBLIC. Use KernelInterface.get_backend on an CLArray to get the backend
+include("OpenCLKernels.jl")
+import .OpenCLInterface
+
 end
diff --git a/src/OpenCLKernels.jl b/src/OpenCLKernels.jl
index b1d4eaa7..cdd18a3f 100644
--- a/src/OpenCLKernels.jl
+++ b/src/OpenCLKernels.jl
@@ -1,9 +1,11 @@
-module OpenCLKernels
+module OpenCLInterface
 
 using ..OpenCL
-using ..OpenCL: @device_override, method_table
+using ..OpenCL: @device_override, method_table, kernel_convert, clfunction
 
-import KernelAbstractions as KA
+import KernelInterface as KI
+
+import SPIRVIntrinsics
 
 import StaticArrays
 
@@ -12,9 +14,10 @@ import Adapt
 
 ## Back-end Definition
 
-export OpenCLBackend
+# export OpenCLBackend
+
 
-Base.@kwdef struct OpenCLBackend <: KA.GPU
+Base.@kwdef struct OpenCLBackend <: KI.GPU
     platform::cl.Platform = cl.platform()
 end
 
@@ -23,7 +26,9 @@ end
     return nothing
 end
 
-function KA.allocate(b::OpenCLBackend, ::Type{T}, dims::Tuple; unified::Bool = false) where T
+KI.versioninfo(io::IO, ::OpenCLBackend) = OpenCL.versioninfo(io)
+
+function KI.allocate(b::OpenCLBackend, ::Type{T}, dims::Tuple; unified::Bool = false) where T
     b.platform === cl.platform() || platform_mismatch_warning(b.platform, cl.platform())
     if unified
         memory_backend = cl.unified_memory_backend()
@@ -39,31 +44,22 @@ function KA.allocate(b::OpenCLBackend, ::Type{T}, dims::Tuple; unified::Bool = f
     end
 end
 
-KA.supports_unified(::OpenCLBackend) = cl.default_memory_backend(cl.device(); unified=true) !== nothing
+KI.supports_unified(::OpenCLBackend) = cl.default_memory_backend(cl.device(); unified=true) !== nothing
 
-KA.get_backend(::CLArray) = OpenCLBackend()
+KI.get_backend(::CLArray) = OpenCLBackend()
 # TODO should be non-blocking
-KA.synchronize(::OpenCLBackend) = cl.finish(cl.queue())
-KA.supports_float64(::OpenCLBackend) = in("cl_khr_fp64", cl.device().extensions)
-
-Adapt.adapt_storage(::OpenCLBackend, a::Array) = Adapt.adapt(CLArray, a)
-Adapt.adapt_storage(::OpenCLBackend, a::CLArray) = a
-Adapt.adapt_storage(::KA.CPU, a::CLArray) = convert(Array, a)
-
-# `@Const` applies `constify` inside the kernel, where arguments have already been
-# converted to device arrays, so the rule has to be registered for `CLDeviceArray`
-# rather than for `CLArray`.
-Adapt.adapt_storage(::KA.ConstAdaptor, a::CLDeviceArray) = Base.Experimental.Const(a)
+KI.synchronize(::OpenCLBackend) = cl.finish(cl.queue())
+KI.supports_float64(::OpenCLBackend) = in("cl_khr_fp64", cl.device().extensions)
 
 ## Device Selection
 
 # devices are numbered consecutively within the backend's platform, in enumeration order
 
-function KA.ndevices(b::OpenCLBackend)
+function KI.ndevices(b::OpenCLBackend)
     Int(cl.ndevices(b.platform))
 end
 
-function KA.device(b::OpenCLBackend)
+function KI.device(b::OpenCLBackend)
     current = cl.device()
     for (i, d) in enumerate(cl.devices(b.platform))
         d == current && return i
@@ -71,8 +67,8 @@ function KA.device(b::OpenCLBackend)
     error("Active OpenCL device $current not found in the OpenCLBackend's platform \"$(b.platform.name)\".")
 end
 
-function KA.device!(b::OpenCLBackend, id::Int)
-    0 < id <= KA.ndevices(b) || throw(ArgumentError("Device id $id out of bounds."))
+function KI.device!(b::OpenCLBackend, id::Int)
+    0 < id <= KI.ndevices(b) || throw(ArgumentError("Device id $id out of bounds."))
     devs = cl.devices(b.platform)
 
     cl.device!(devs[id])
@@ -81,7 +77,7 @@ end
 
 ## Memory Operations
 
-function KA.copyto!(::OpenCLBackend, A, B)
+function KI.copyto!(::OpenCLBackend, A, B)
     copyto!(A, B)
     # TODO: Address device to host copies in jl being synchronizing
 end
@@ -89,144 +85,134 @@ end
 
 ## Kernel Launch
 
-function KA.mkcontext(kernel::KA.Kernel{OpenCLBackend}, _ndrange, iterspace)
-    KA.CompilerMetadata{KA.ndrange(kernel), KA.DynamicCheck}(_ndrange, iterspace)
-end
-function KA.mkcontext(kernel::KA.Kernel{OpenCLBackend}, I, _ndrange, iterspace,
-                      ::Dynamic) where Dynamic
-    KA.CompilerMetadata{KA.ndrange(kernel), Dynamic}(I, _ndrange, iterspace)
+
+KI.argconvert(::OpenCLBackend, arg) = kernel_convert(arg)
+
+function KI.kernel_function(::OpenCLBackend, f::F, tt::TT=Tuple{}; name = nothing, kwargs...) where {F,TT}
+    kern = clfunction(f, tt; name, kwargs...)
+    KI.Kernel{OpenCLBackend, typeof(kern)}(OpenCLBackend(), kern)
 end
 
-function KA.launch_config(kernel::KA.Kernel{OpenCLBackend}, ndrange, workgroupsize)
-    if ndrange isa Integer
-        ndrange = (ndrange,)
-    end
-    if workgroupsize isa Integer
-        workgroupsize = (workgroupsize, )
-    end
+function (obj::KI.Kernel{OpenCLBackend})(args...; numworkgroups=(), workgroupsize=(), ndrange=(), max_work_group_size=typemax(Int))
+    obj.backend.platform === cl.platform() || platform_mismatch_warning(obj.backend.platform, cl.platform())
+    KI.check_launch_args(numworkgroups, workgroupsize, ndrange)
+    prod(ndrange) == 0 && return nothing
 
-    # partition checked that the ndrange's agreed
-    if KA.ndrange(kernel) <: KA.StaticSize
-        ndrange = nothing
-    end
+    numworkgroups, workgroupsize = KI.auto_launch_sizes(obj, numworkgroups, workgroupsize, ndrange, max_work_group_size)
+    local_size = (workgroupsize..., ntuple(_ -> 1, 3 - length(workgroupsize))...)
+    numworkgroups = (numworkgroups..., ntuple(_ -> 1, 3 - length(numworkgroups))...)
+    global_size = local_size .* numworkgroups
+
+    obj.kern(args...; local_size, global_size)
+    return nothing
+end
 
-    iterspace, dynamic = if KA.workgroupsize(kernel) <: KA.DynamicSize &&
-        workgroupsize === nothing
-        # use ndrange as preliminary workgroupsize for autotuning
-        KA.partition(kernel, ndrange, ndrange)
-    else
-        KA.partition(kernel, ndrange, workgroupsize)
-    end
 
-    return ndrange, workgroupsize, iterspace, dynamic
+function KI.kernel_max_work_group_size(kernel::KI.Kernel{<:OpenCLBackend}; max_work_items::Int=typemax(Int))::Int
+    wginfo = cl.work_group_info(kernel.kern.fun, cl.device())
+    Int(min(wginfo.size, max_work_items))
 end
 
-function threads_to_workgroupsize(threads, ndrange)
-    total = 1
-    return map(ndrange) do n
-        x = min(div(threads, total), n)
-        total *= x
-        return x
+# querying the device allocates, so cache the limits that every launch needs. the cache is
+# keyed on the device, because the task-local device can be switched.
+const DeviceLimits = @NamedTuple{max_work_group_size::Int, max_work_group_dims::NTuple{3, Int}}
+function device_limits()
+    dev = cl.device()
+    cache = get!(task_local_storage(), :CLDeviceLimits) do
+        Dict{cl.Device, DeviceLimits}()
+    end::Dict{cl.Device, DeviceLimits}
+    return get!(cache, dev) do
+        sizes = dev.max_work_item_size
+        (; max_work_group_size = Int(dev.max_work_group_size),
+           max_work_group_dims = ntuple(d -> d <= length(sizes) ? sizes[d] : 1, 3))
     end
 end
+KI.max_work_group_size(::OpenCLBackend)::Int = device_limits().max_work_group_size
+KI.max_work_group_dims(::OpenCLBackend)::NTuple{3, Int} = device_limits().max_work_group_dims
+function KI.sub_group_size(::OpenCLBackend)::Int
+    cl.sub_group_size(cl.device())
+end
+function KI.multiprocessor_count(::OpenCLBackend)::Int
+    Int(cl.device().max_compute_units)
+end
 
-function (obj::KA.Kernel{OpenCLBackend})(args...; ndrange=nothing, workgroupsize=nothing)
-    obj.backend.platform === cl.platform() || platform_mismatch_warning(obj.backend.platform, cl.platform())
-
-    ndrange, workgroupsize, iterspace, dynamic =
-        KA.launch_config(obj, ndrange, workgroupsize)
+function KI.shfl_down_types(::OpenCLBackend)
+    backend_extensions = cl.device().extensions
+    "cl_khr_subgroup_shuffle" in backend_extensions || return DataType[]
 
-    # this might not be the final context, since we may tune the workgroupsize
-    ctx = KA.mkcontext(obj, ndrange, iterspace)
-    kernel = @opencl launch=false obj.f(ctx, args...)
+    res = copy(SPIRVIntrinsics.gentypes)
 
-    # figure out the optimal workgroupsize automatically
-    if KA.workgroupsize(obj) <: KA.DynamicSize && workgroupsize === nothing
-        wg_info = cl.work_group_info(kernel.fun, cl.device())
-        wg_size_nd = threads_to_workgroupsize(wg_info.size, ndrange)
-        iterspace, dynamic = KA.partition(obj, ndrange, wg_size_nd)
-        ctx = KA.mkcontext(obj, ndrange, iterspace)
+    if "cl_khr_fp64" ∉ backend_extensions
+        res = setdiff(res, [Float64])
     end
-
-    groups = length(KA.blocks(iterspace))
-    items = length(KA.workitems(iterspace))
-
-    if groups == 0
-        return nothing
+    if "cl_khr_fp16" ∉ backend_extensions
+        res = setdiff(res, [Float16])
     end
 
-    # Launch kernel
-    global_size = groups * items
-    local_size = items
-    kernel(ctx, args...; global_size, local_size)
-
-    return nothing
+    return res
 end
 
-
 ## Indexing Functions
+## COV_EXCL_START
 
-@device_override @inline function KA.__index_Local_Linear(ctx)
-    return get_local_id(1)
+@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(1)
+@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(1)    # JuliaGPU/OpenCL.jl#346
-    I = KA.__index_Global_Cartesian(ctx)
-    @inbounds LinearIndices(KA.__ndrange(ctx))[I]
+@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(1)]
+@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(1)]
+@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(1), get_local_id(1))
+@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 = KA.__index_Global_Cartesian(ctx)
-        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 = OpenCL.emit_localmemory(T, Val(prod(Dims)))
     CLDeviceArray(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()
     work_group_barrier(OpenCL.LOCAL_MEM_FENCE | OpenCL.GLOBAL_MEM_FENCE)
 end
 
-@device_override @inline function KA.__print(args...)
-    OpenCL._print(args...)
+@device_override @inline function KI.sub_group_barrier()
+    sub_group_barrier(OpenCL.LOCAL_MEM_FENCE | OpenCL.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
-
-KA.argconvert(::KA.Kernel{OpenCLBackend}, arg) = OpenCL.kernel_convert(arg)
+@device_override @inline function KI._print(args...)
+    OpenCL._print(args...)
+end
+## COV_EXCL_STOP
 
 end
diff --git a/src/OpenCLKernelsOld.jl b/src/OpenCLKernelsOld.jl
new file mode 100644
index 00000000..b1d4eaa7
--- /dev/null
+++ b/src/OpenCLKernelsOld.jl
@@ -0,0 +1,232 @@
+module OpenCLKernels
+
+using ..OpenCL
+using ..OpenCL: @device_override, method_table
+
+import KernelAbstractions as KA
+
+import StaticArrays
+
+import Adapt
+
+
+## Back-end Definition
+
+export OpenCLBackend
+
+Base.@kwdef struct OpenCLBackend <: KA.GPU
+    platform::cl.Platform = cl.platform()
+end
+
+@noinline function platform_mismatch_warning(expected::cl.Platform, active::cl.Platform)
+    @warn "OpenCLBackend platform \"$(expected.name)\" is not the active platform \"$(active.name)\""
+    return nothing
+end
+
+function KA.allocate(b::OpenCLBackend, ::Type{T}, dims::Tuple; unified::Bool = false) where T
+    b.platform === cl.platform() || platform_mismatch_warning(b.platform, cl.platform())
+    if unified
+        memory_backend = cl.unified_memory_backend()
+        if memory_backend === cl.USMBackend()
+            return CLArray{T, length(dims), cl.UnifiedSharedMemory}(undef, dims)
+        elseif memory_backend === cl.SVMBackend()
+            return CLArray{T, length(dims), cl.SharedVirtualMemory}(undef, dims)
+        else
+            throw(ArgumentError("Unified memory not supported"))
+        end
+    else
+        return CLArray{T}(undef, dims)
+    end
+end
+
+KA.supports_unified(::OpenCLBackend) = cl.default_memory_backend(cl.device(); unified=true) !== nothing
+
+KA.get_backend(::CLArray) = OpenCLBackend()
+# TODO should be non-blocking
+KA.synchronize(::OpenCLBackend) = cl.finish(cl.queue())
+KA.supports_float64(::OpenCLBackend) = in("cl_khr_fp64", cl.device().extensions)
+
+Adapt.adapt_storage(::OpenCLBackend, a::Array) = Adapt.adapt(CLArray, a)
+Adapt.adapt_storage(::OpenCLBackend, a::CLArray) = a
+Adapt.adapt_storage(::KA.CPU, a::CLArray) = convert(Array, a)
+
+# `@Const` applies `constify` inside the kernel, where arguments have already been
+# converted to device arrays, so the rule has to be registered for `CLDeviceArray`
+# rather than for `CLArray`.
+Adapt.adapt_storage(::KA.ConstAdaptor, a::CLDeviceArray) = Base.Experimental.Const(a)
+
+## Device Selection
+
+# devices are numbered consecutively within the backend's platform, in enumeration order
+
+function KA.ndevices(b::OpenCLBackend)
+    Int(cl.ndevices(b.platform))
+end
+
+function KA.device(b::OpenCLBackend)
+    current = cl.device()
+    for (i, d) in enumerate(cl.devices(b.platform))
+        d == current && return i
+    end
+    error("Active OpenCL device $current not found in the OpenCLBackend's platform \"$(b.platform.name)\".")
+end
+
+function KA.device!(b::OpenCLBackend, id::Int)
+    0 < id <= KA.ndevices(b) || throw(ArgumentError("Device id $id out of bounds."))
+    devs = cl.devices(b.platform)
+
+    cl.device!(devs[id])
+    return nothing
+end
+
+## Memory Operations
+
+function KA.copyto!(::OpenCLBackend, A, B)
+    copyto!(A, B)
+    # TODO: Address device to host copies in jl being synchronizing
+end
+
+
+## Kernel Launch
+
+function KA.mkcontext(kernel::KA.Kernel{OpenCLBackend}, _ndrange, iterspace)
+    KA.CompilerMetadata{KA.ndrange(kernel), KA.DynamicCheck}(_ndrange, iterspace)
+end
+function KA.mkcontext(kernel::KA.Kernel{OpenCLBackend}, I, _ndrange, iterspace,
+                      ::Dynamic) where Dynamic
+    KA.CompilerMetadata{KA.ndrange(kernel), Dynamic}(I, _ndrange, iterspace)
+end
+
+function KA.launch_config(kernel::KA.Kernel{OpenCLBackend}, 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
+        KA.partition(kernel, ndrange, ndrange)
+    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 = min(div(threads, total), n)
+        total *= x
+        return x
+    end
+end
+
+function (obj::KA.Kernel{OpenCLBackend})(args...; ndrange=nothing, workgroupsize=nothing)
+    obj.backend.platform === cl.platform() || platform_mismatch_warning(obj.backend.platform, cl.platform())
+
+    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)
+    kernel = @opencl launch=false obj.f(ctx, args...)
+
+    # figure out the optimal workgroupsize automatically
+    if KA.workgroupsize(obj) <: KA.DynamicSize && workgroupsize === nothing
+        wg_info = cl.work_group_info(kernel.fun, cl.device())
+        wg_size_nd = threads_to_workgroupsize(wg_info.size, ndrange)
+        iterspace, dynamic = KA.partition(obj, ndrange, wg_size_nd)
+        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
+    global_size = groups * items
+    local_size = items
+    kernel(ctx, args...; global_size, local_size)
+
+    return nothing
+end
+
+
+## Indexing Functions
+
+@device_override @inline function KA.__index_Local_Linear(ctx)
+    return get_local_id(1)
+end
+
+@device_override @inline function KA.__index_Group_Linear(ctx)
+    return get_group_id(1)
+end
+
+@device_override @inline function KA.__index_Global_Linear(ctx)
+    #return get_global_id(1)    # JuliaGPU/OpenCL.jl#346
+    I = KA.__index_Global_Cartesian(ctx)
+    @inbounds LinearIndices(KA.__ndrange(ctx))[I]
+end
+
+@device_override @inline function KA.__index_Local_Cartesian(ctx)
+    @inbounds KA.workitems(KA.__iterspace(ctx))[get_local_id(1)]
+end
+
+@device_override @inline function KA.__index_Group_Cartesian(ctx)
+    @inbounds KA.blocks(KA.__iterspace(ctx))[get_group_id(1)]
+end
+
+@device_override @inline function KA.__index_Global_Cartesian(ctx)
+    return @inbounds KA.expand(KA.__iterspace(ctx), get_group_id(1), get_local_id(1))
+end
+
+@device_override @inline function KA.__validindex(ctx)
+    if KA.__dynamic_checkbounds(ctx)
+        I = KA.__index_Global_Cartesian(ctx)
+        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 = OpenCL.emit_localmemory(T, Val(prod(Dims)))
+    CLDeviceArray(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()
+    work_group_barrier(OpenCL.LOCAL_MEM_FENCE | OpenCL.GLOBAL_MEM_FENCE)
+end
+
+@device_override @inline function KA.__print(args...)
+    OpenCL._print(args...)
+end
+
+
+## Other
+
+KA.argconvert(::KA.Kernel{OpenCLBackend}, arg) = OpenCL.kernel_convert(arg)
+
+end
diff --git a/src/util.jl b/src/util.jl
index 725fc1f4..8fbe58d8 100644
--- a/src/util.jl
+++ b/src/util.jl
@@ -73,7 +73,7 @@ function versioninfo(io::IO=stdout)
     end
 
     for pkg in [:GPUArrays, :GPUCompiler, ("63c18a36-062a-441e-b654-da1e3ab1ce7c", "KernelAbstractions"),
-                 :LLVM, :SPIRVIntrinsics, ("627d6b7a-bbe6-5189-83e7-98cc0a5aeadd", "pocl_jll"),
+                 :KernelInterface, :LLVM, :SPIRVIntrinsics, ("627d6b7a-bbe6-5189-83e7-98cc0a5aeadd", "pocl_jll"),
                  ("59abdad9-3cfc-5436-8271-411e8cad6b82", "pocl_next_jll")]
         name, mod = get_module(pkg)
         isnothing(mod) && continue
diff --git a/test/Project.toml b/test/Project.toml
index 14eef6fd..c4e439a9 100644
--- a/test/Project.toml
+++ b/test/Project.toml
@@ -9,6 +9,7 @@ IOCapture = "b5f81e59-6552-4d32-b1f0-c071b021bf89"
 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"
 OpenCL = "08131aa3-fb12-5dee-8b74-c09406e224a2"
 ParallelTestRunner = "d3525ed8-44d0-4b2c-a655-542cee43accc"
diff --git a/test/kernelinterface.jl b/test/kernelinterface.jl
new file mode 100644
index 00000000..cc6c2083
--- /dev/null
+++ b/test/kernelinterface.jl
@@ -0,0 +1,6 @@
+import KernelInterface
+using OpenCL.OpenCLInterface
+
+include(joinpath(dirname(pathof(KernelInterface)), "..", "test", "testsuite.jl"))
+
+Testsuite.testsuite(OpenCLInterface.OpenCLBackend, "OpenCL", OpenCL, CLArray, OpenCL.CLDeviceArray)

From 1d9efa9c2437cca7f80eafd2310d441fa0cda11a Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Wed, 19 Aug 2026 14:50:05 -0300
Subject: [PATCH 2/3] Support KernelAbstractions 0.10 through a package
 extension

KernelAbstractions 0.10 builds on KernelInterface, so the KernelInterface
back end becomes the public `OpenCLBackend` (in `OpenCLKernels`), and the
KernelAbstractions-specific parts move to a `KernelAbstractionsExt`
extension. The KernelAbstractions 0.9 back end is removed.

Kernel launches with an automatically tuned workgroup size compute it
from the number of indices along each dimension, as an ndrange may now
contain ranges, and return before compiling when the ndrange is empty.

Skip the KernelAbstractions tests that only apply to the CPU back end.

Co-authored-by: Tim Besard <tim.besard@gmail.com>
---
 Project.toml                 |  12 +-
 ext/KernelAbstractionsExt.jl | 113 +++++++++++++++++
 src/OpenCL.jl                |   8 +-
 src/OpenCLKernels.jl         |   4 +-
 src/OpenCLKernelsOld.jl      | 232 -----------------------------------
 test/Project.toml            |   2 +-
 test/kernelabstractions.jl   |   3 +-
 test/kernelinterface.jl      |   4 +-
 8 files changed, 131 insertions(+), 247 deletions(-)
 create mode 100644 ext/KernelAbstractionsExt.jl
 delete mode 100644 src/OpenCLKernelsOld.jl

diff --git a/Project.toml b/Project.toml
index d74ddf83..bbe0e8ce 100644
--- a/Project.toml
+++ b/Project.toml
@@ -10,7 +10,6 @@ Adapt = "79e6a3ab-5dfb-504d-930d-738a2a938a0e"
 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"
 LinearAlgebra = "37e2e46d-f89d-539d-b4ee-838fcccc9c8e"
@@ -27,15 +26,21 @@ SPIRV_Tools_jll = "6ac6d60f-d740-5983-97d7-a4482c0689f4"
 StaticArrays = "90137ffa-7385-5640-81b9-e52037218182"
 spirv2clc_jll = "f0274c0c-8c8a-59f1-85b7-f7d60330c5fb"
 
+[weakdeps]
+KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c"
+
 [sources]
 SPIRVIntrinsics = {path = "lib/intrinsics"}
 
+[extensions]
+KernelAbstractionsExt = "KernelAbstractions"
+
 [compat]
 Adapt = "4"
 GPUArrays = "11.2.1"
 GPUCompiler = "2.9"
 GPUToolbox = "3.1"
-KernelAbstractions = "0.9.38"
+KernelAbstractions = "0.10"
 KernelInterface = "0.2.3"
 LLVM = "9.6"
 LinearAlgebra = "1"
@@ -52,3 +57,6 @@ SPIRV_Tools_jll = "2025.1"
 StaticArrays = "1"
 julia = "1.10"
 spirv2clc_jll = "0.2.0"
+
+[extras]
+KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c"
diff --git a/ext/KernelAbstractionsExt.jl b/ext/KernelAbstractionsExt.jl
new file mode 100644
index 00000000..295ef755
--- /dev/null
+++ b/ext/KernelAbstractionsExt.jl
@@ -0,0 +1,113 @@
+module KernelAbstractionsExt
+
+using OpenCL
+using OpenCL: @device_override, method_table
+
+import KernelAbstractions as KA
+
+import StaticArrays
+
+import Adapt
+
+Adapt.adapt_storage(::OpenCLBackend, a::Array) = Adapt.adapt(CLArray, a)
+Adapt.adapt_storage(::OpenCLBackend, a::CLArray) = a
+Adapt.adapt_storage(::KA.CPU, a::CLArray) = convert(Array, a)
+
+# `@Const` applies `constify` inside the kernel, where arguments have already been
+# converted to device arrays, so the rule has to be registered for `CLDeviceArray`
+# rather than for `CLArray`.
+Adapt.adapt_storage(::KA.ConstAdaptor, a::CLDeviceArray) = Base.Experimental.Const(a)
+
+## Kernel Launch
+
+function KA.mkcontext(kernel::KA.Kernel{OpenCLBackend}, _ndrange, iterspace)
+    KA.CompilerMetadata{KA.ndrange(kernel), KA.DynamicCheck}(_ndrange, iterspace)
+end
+function KA.mkcontext(kernel::KA.Kernel{OpenCLBackend}, I, _ndrange, iterspace,
+                      ::Dynamic) where Dynamic
+    KA.CompilerMetadata{KA.ndrange(kernel), Dynamic}(I, _ndrange, iterspace)
+end
+
+function KA.launch_config(kernel::KA.Kernel{OpenCLBackend}, 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
+        KA.partition(kernel, ndrange, ndrange)
+    else
+        KA.partition(kernel, ndrange, workgroupsize)
+    end
+
+    return ndrange, workgroupsize, iterspace, dynamic
+end
+
+# the number of indices along each dimension of `ndrange`, which may contain ranges
+extents(ndrange) = size(CartesianIndices(ndrange))
+
+function threads_to_workgroupsize(threads, ndrange)
+    total = 1
+    return map(ndrange) do n
+        x = max(min(div(threads, total), n), 1)
+        total *= x
+        return x
+    end
+end
+
+function (obj::KA.Kernel{OpenCLBackend})(args...; ndrange=nothing, workgroupsize=nothing)
+    obj.backend.platform === cl.platform() || OpenCL.OpenCLKernels.platform_mismatch_warning(obj.backend.platform, cl.platform())
+
+    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)
+    kernel = @opencl launch=false obj.f(ctx, args...)
+
+    # figure out the optimal workgroupsize automatically
+    if KA.workgroupsize(obj) <: KA.DynamicSize && workgroupsize === nothing
+        wg_info = cl.work_group_info(kernel.fun, cl.device())
+        wg_size_nd = threads_to_workgroupsize(wg_info.size, extents(ndrange))
+        iterspace, dynamic = KA.partition(obj, ndrange, wg_size_nd)
+        ctx = KA.mkcontext(obj, ndrange, iterspace)
+    end
+
+    groups = length(KA.blocks(iterspace))
+    items = length(KA.workitems(iterspace))
+
+    # Launch kernel
+    global_size = groups * items
+    local_size = items
+    kernel(ctx, args...; global_size, local_size)
+
+    return nothing
+end
+
+@device_override @inline function KA.__validindex(ctx)
+    if KA.__dynamic_checkbounds(ctx)
+        I = KA.__index_Global_Cartesian(ctx)
+        return I in KA.__ndrange(ctx)
+    else
+        return true
+    end
+end
+
+@device_override @inline function KA.Scratchpad(ctx, ::Type{T}, ::Val{Dims}) where {T, Dims}
+    StaticArrays.MArray{KA.__size(Dims), T}(undef)
+end
+
+KA.argconvert(::KA.Kernel{OpenCLBackend}, arg) = OpenCL.kernel_convert(arg)
+
+end
diff --git a/src/OpenCL.jl b/src/OpenCL.jl
index ca02eb1b..25a1764f 100644
--- a/src/OpenCL.jl
+++ b/src/OpenCL.jl
@@ -10,8 +10,6 @@ using GPUArrays
 using Random
 using Preferences
 
-import KernelAbstractions: KernelAbstractions
-
 import KernelInterface
 
 using Core: LLVMPtr
@@ -51,12 +49,8 @@ include("mapreduce.jl")
 include("gpuarrays.jl")
 include("random.jl")
 
-include("OpenCLKernelsOld.jl")
+include("OpenCLKernels.jl")
 import .OpenCLKernels: OpenCLBackend
 export OpenCLBackend
 
-# KernelInterface - NOT PUBLIC. Use KernelInterface.get_backend on an CLArray to get the backend
-include("OpenCLKernels.jl")
-import .OpenCLInterface
-
 end
diff --git a/src/OpenCLKernels.jl b/src/OpenCLKernels.jl
index cdd18a3f..c13904b2 100644
--- a/src/OpenCLKernels.jl
+++ b/src/OpenCLKernels.jl
@@ -1,4 +1,4 @@
-module OpenCLInterface
+module OpenCLKernels
 
 using ..OpenCL
 using ..OpenCL: @device_override, method_table, kernel_convert, clfunction
@@ -14,7 +14,7 @@ import Adapt
 
 ## Back-end Definition
 
-# export OpenCLBackend
+export OpenCLBackend
 
 
 Base.@kwdef struct OpenCLBackend <: KI.GPU
diff --git a/src/OpenCLKernelsOld.jl b/src/OpenCLKernelsOld.jl
deleted file mode 100644
index b1d4eaa7..00000000
--- a/src/OpenCLKernelsOld.jl
+++ /dev/null
@@ -1,232 +0,0 @@
-module OpenCLKernels
-
-using ..OpenCL
-using ..OpenCL: @device_override, method_table
-
-import KernelAbstractions as KA
-
-import StaticArrays
-
-import Adapt
-
-
-## Back-end Definition
-
-export OpenCLBackend
-
-Base.@kwdef struct OpenCLBackend <: KA.GPU
-    platform::cl.Platform = cl.platform()
-end
-
-@noinline function platform_mismatch_warning(expected::cl.Platform, active::cl.Platform)
-    @warn "OpenCLBackend platform \"$(expected.name)\" is not the active platform \"$(active.name)\""
-    return nothing
-end
-
-function KA.allocate(b::OpenCLBackend, ::Type{T}, dims::Tuple; unified::Bool = false) where T
-    b.platform === cl.platform() || platform_mismatch_warning(b.platform, cl.platform())
-    if unified
-        memory_backend = cl.unified_memory_backend()
-        if memory_backend === cl.USMBackend()
-            return CLArray{T, length(dims), cl.UnifiedSharedMemory}(undef, dims)
-        elseif memory_backend === cl.SVMBackend()
-            return CLArray{T, length(dims), cl.SharedVirtualMemory}(undef, dims)
-        else
-            throw(ArgumentError("Unified memory not supported"))
-        end
-    else
-        return CLArray{T}(undef, dims)
-    end
-end
-
-KA.supports_unified(::OpenCLBackend) = cl.default_memory_backend(cl.device(); unified=true) !== nothing
-
-KA.get_backend(::CLArray) = OpenCLBackend()
-# TODO should be non-blocking
-KA.synchronize(::OpenCLBackend) = cl.finish(cl.queue())
-KA.supports_float64(::OpenCLBackend) = in("cl_khr_fp64", cl.device().extensions)
-
-Adapt.adapt_storage(::OpenCLBackend, a::Array) = Adapt.adapt(CLArray, a)
-Adapt.adapt_storage(::OpenCLBackend, a::CLArray) = a
-Adapt.adapt_storage(::KA.CPU, a::CLArray) = convert(Array, a)
-
-# `@Const` applies `constify` inside the kernel, where arguments have already been
-# converted to device arrays, so the rule has to be registered for `CLDeviceArray`
-# rather than for `CLArray`.
-Adapt.adapt_storage(::KA.ConstAdaptor, a::CLDeviceArray) = Base.Experimental.Const(a)
-
-## Device Selection
-
-# devices are numbered consecutively within the backend's platform, in enumeration order
-
-function KA.ndevices(b::OpenCLBackend)
-    Int(cl.ndevices(b.platform))
-end
-
-function KA.device(b::OpenCLBackend)
-    current = cl.device()
-    for (i, d) in enumerate(cl.devices(b.platform))
-        d == current && return i
-    end
-    error("Active OpenCL device $current not found in the OpenCLBackend's platform \"$(b.platform.name)\".")
-end
-
-function KA.device!(b::OpenCLBackend, id::Int)
-    0 < id <= KA.ndevices(b) || throw(ArgumentError("Device id $id out of bounds."))
-    devs = cl.devices(b.platform)
-
-    cl.device!(devs[id])
-    return nothing
-end
-
-## Memory Operations
-
-function KA.copyto!(::OpenCLBackend, A, B)
-    copyto!(A, B)
-    # TODO: Address device to host copies in jl being synchronizing
-end
-
-
-## Kernel Launch
-
-function KA.mkcontext(kernel::KA.Kernel{OpenCLBackend}, _ndrange, iterspace)
-    KA.CompilerMetadata{KA.ndrange(kernel), KA.DynamicCheck}(_ndrange, iterspace)
-end
-function KA.mkcontext(kernel::KA.Kernel{OpenCLBackend}, I, _ndrange, iterspace,
-                      ::Dynamic) where Dynamic
-    KA.CompilerMetadata{KA.ndrange(kernel), Dynamic}(I, _ndrange, iterspace)
-end
-
-function KA.launch_config(kernel::KA.Kernel{OpenCLBackend}, 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
-        KA.partition(kernel, ndrange, ndrange)
-    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 = min(div(threads, total), n)
-        total *= x
-        return x
-    end
-end
-
-function (obj::KA.Kernel{OpenCLBackend})(args...; ndrange=nothing, workgroupsize=nothing)
-    obj.backend.platform === cl.platform() || platform_mismatch_warning(obj.backend.platform, cl.platform())
-
-    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)
-    kernel = @opencl launch=false obj.f(ctx, args...)
-
-    # figure out the optimal workgroupsize automatically
-    if KA.workgroupsize(obj) <: KA.DynamicSize && workgroupsize === nothing
-        wg_info = cl.work_group_info(kernel.fun, cl.device())
-        wg_size_nd = threads_to_workgroupsize(wg_info.size, ndrange)
-        iterspace, dynamic = KA.partition(obj, ndrange, wg_size_nd)
-        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
-    global_size = groups * items
-    local_size = items
-    kernel(ctx, args...; global_size, local_size)
-
-    return nothing
-end
-
-
-## Indexing Functions
-
-@device_override @inline function KA.__index_Local_Linear(ctx)
-    return get_local_id(1)
-end
-
-@device_override @inline function KA.__index_Group_Linear(ctx)
-    return get_group_id(1)
-end
-
-@device_override @inline function KA.__index_Global_Linear(ctx)
-    #return get_global_id(1)    # JuliaGPU/OpenCL.jl#346
-    I = KA.__index_Global_Cartesian(ctx)
-    @inbounds LinearIndices(KA.__ndrange(ctx))[I]
-end
-
-@device_override @inline function KA.__index_Local_Cartesian(ctx)
-    @inbounds KA.workitems(KA.__iterspace(ctx))[get_local_id(1)]
-end
-
-@device_override @inline function KA.__index_Group_Cartesian(ctx)
-    @inbounds KA.blocks(KA.__iterspace(ctx))[get_group_id(1)]
-end
-
-@device_override @inline function KA.__index_Global_Cartesian(ctx)
-    return @inbounds KA.expand(KA.__iterspace(ctx), get_group_id(1), get_local_id(1))
-end
-
-@device_override @inline function KA.__validindex(ctx)
-    if KA.__dynamic_checkbounds(ctx)
-        I = KA.__index_Global_Cartesian(ctx)
-        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 = OpenCL.emit_localmemory(T, Val(prod(Dims)))
-    CLDeviceArray(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()
-    work_group_barrier(OpenCL.LOCAL_MEM_FENCE | OpenCL.GLOBAL_MEM_FENCE)
-end
-
-@device_override @inline function KA.__print(args...)
-    OpenCL._print(args...)
-end
-
-
-## Other
-
-KA.argconvert(::KA.Kernel{OpenCLBackend}, arg) = OpenCL.kernel_convert(arg)
-
-end
diff --git a/test/Project.toml b/test/Project.toml
index c4e439a9..f0efeaf8 100644
--- a/test/Project.toml
+++ b/test/Project.toml
@@ -35,7 +35,7 @@ SPIRVIntrinsics = {path = "../lib/intrinsics"}
 
 [compat]
 IOCapture = "1"
+ParallelTestRunner = "2.2"
 Statistics = "1"
 pocl_jll = "7.2"
 pocl_next_jll = "7.2"
-ParallelTestRunner = "2.2"
diff --git a/test/kernelabstractions.jl b/test/kernelabstractions.jl
index 4f2d3de7..ccea9748 100644
--- a/test/kernelabstractions.jl
+++ b/test/kernelabstractions.jl
@@ -66,6 +66,7 @@ end
 
 skip_tests=Set([
     "sparse",
-    "Convert", # Need to opt out of i128
+    "CPU synchronization",
+    "fallback test: callable types"
 ])
 KATestSuite.testsuite(OpenCLBackend, "OpenCL", OpenCL, CLArray, CLDeviceArray; skip_tests)
diff --git a/test/kernelinterface.jl b/test/kernelinterface.jl
index cc6c2083..bbeb2467 100644
--- a/test/kernelinterface.jl
+++ b/test/kernelinterface.jl
@@ -1,6 +1,6 @@
 import KernelInterface
-using OpenCL.OpenCLInterface
+using OpenCL
 
 include(joinpath(dirname(pathof(KernelInterface)), "..", "test", "testsuite.jl"))
 
-Testsuite.testsuite(OpenCLInterface.OpenCLBackend, "OpenCL", OpenCL, CLArray, OpenCL.CLDeviceArray)
+Testsuite.testsuite(OpenCLBackend, "OpenCL", OpenCL, CLArray, OpenCL.CLDeviceArray)

From ec09404f615cc2057f4407c56531c5f20def86a6 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Sat, 22 Aug 2026 12:31:53 -0300
Subject: [PATCH 3/3] [temp] CI: develop KernelAbstractions until 0.10 is
 registered

The compat also admits 0.9, because Pkg orders 0.10.0-dev before 0.10.0.
Pkg < 1.12 drops a developed weak dependency from the project, so restore
the project after developing KernelAbstractions on those versions.

Co-authored-by: Tim Besard <tim.besard@gmail.com>
---
 .buildkite/pipeline.yml    | 2 ++
 .github/workflows/Test.yml | 6 +++++-
 Project.toml               | 2 +-
 3 files changed, 8 insertions(+), 2 deletions(-)

diff --git a/.buildkite/pipeline.yml b/.buildkite/pipeline.yml
index 6d5d5c17..f1bc6631 100644
--- a/.buildkite/pipeline.yml
+++ b/.buildkite/pipeline.yml
@@ -16,6 +16,7 @@ steps:
         # Julia 1.10 does not support [sources], so dev the in-tree
         # SPIRVIntrinsics; Pkg.test then carries it into the test sandbox.
         Pkg.develop(path="lib/intrinsics")
+        Pkg.develop("KernelAbstractions")
 
         println("+++ :julia: Running tests")
         Pkg.test(; coverage=true, test_args=`--platform=cuda`)'
@@ -48,6 +49,7 @@ steps:
             # against the registry, which has no SPIRVIntrinsics 1. Pkg.test
             # then carries the in-tree copy into the test sandbox.
             Pkg.develop(path="lib/intrinsics")
+            Pkg.develop("KernelAbstractions")
             Pkg.add("{{matrix.pocl}}_jll")
             Pkg.add("InteractiveUtils")
 
diff --git a/.github/workflows/Test.yml b/.github/workflows/Test.yml
index ed5a15bc..09675c84 100644
--- a/.github/workflows/Test.yml
+++ b/.github/workflows/Test.yml
@@ -142,7 +142,11 @@ jobs:
             using Pkg
             # Julia 1.10 does not support [sources], so dev the in-tree
             # SPIRVIntrinsics; Pkg.test then carries it into the test sandbox.
-            Pkg.develop(path="lib/intrinsics")'
+            Pkg.develop(path="lib/intrinsics")
+            Pkg.develop("KernelAbstractions")
+            # Pkg < 1.12 drops a developed weak dependency from the project, breaking the
+            # extension; restore the project, keeping the developed package in the manifest.
+            VERSION < v"1.12" && run(`git checkout Project.toml`)'
 
       - name: Test OpenCL.jl
         uses: julia-actions/julia-runtest@v1
diff --git a/Project.toml b/Project.toml
index bbe0e8ce..d88d9faa 100644
--- a/Project.toml
+++ b/Project.toml
@@ -40,7 +40,7 @@ Adapt = "4"
 GPUArrays = "11.2.1"
 GPUCompiler = "2.9"
 GPUToolbox = "3.1"
-KernelAbstractions = "0.10"
+KernelAbstractions = "0.9, 0.10"
 KernelInterface = "0.2.3"
 LLVM = "9.6"
 LinearAlgebra = "1"