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 0236d89d..d88d9faa 100644
--- a/Project.toml
+++ b/Project.toml
@@ -10,7 +10,7 @@ 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"
 OpenCL_jll = "6cb37087-e8b6-5417-8430-1f242f1e46e4"
@@ -26,15 +26,22 @@ 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.9, 0.10"
+KernelInterface = "0.2.3"
 LLVM = "9.6"
 LinearAlgebra = "1"
 OpenCL_jll = "=2024.10.24"
@@ -50,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 da994572..25a1764f 100644
--- a/src/OpenCL.jl
+++ b/src/OpenCL.jl
@@ -10,7 +10,7 @@ using GPUArrays
 using Random
 using Preferences
 
-import KernelAbstractions: KernelAbstractions
+import KernelInterface
 
 using Core: LLVMPtr
 
@@ -52,4 +52,5 @@ include("random.jl")
 include("OpenCLKernels.jl")
 import .OpenCLKernels: OpenCLBackend
 export OpenCLBackend
+
 end
diff --git a/src/OpenCLKernels.jl b/src/OpenCLKernels.jl
index b1d4eaa7..c13904b2 100644
--- a/src/OpenCLKernels.jl
+++ b/src/OpenCLKernels.jl
@@ -1,9 +1,11 @@
 module OpenCLKernels
 
 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
 
@@ -14,7 +16,8 @@ import Adapt
 
 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/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..f0efeaf8 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"
@@ -34,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
new file mode 100644
index 00000000..bbeb2467
--- /dev/null
+++ b/test/kernelinterface.jl
@@ -0,0 +1,6 @@
+import KernelInterface
+using OpenCL
+
+include(joinpath(dirname(pathof(KernelInterface)), "..", "test", "testsuite.jl"))
+
+Testsuite.testsuite(OpenCLBackend, "OpenCL", OpenCL, CLArray, OpenCL.CLDeviceArray)