From a05b1d3fbcbd46c29a37139e9df17d4088568945 Mon Sep 17 00:00:00 2001
From: Christian <28689358+christiangnrd@users.noreply.github.com>
Date: Sat, 22 Aug 2026 12:43:53 -0300
Subject: [PATCH 01/10] Optionally display KernelAbstractions in versioninfo

---
 CUDATools/src/utilities.jl | 10 ++++++++--
 1 file changed, 8 insertions(+), 2 deletions(-)

diff --git a/CUDATools/src/utilities.jl b/CUDATools/src/utilities.jl
index 4ef938f2ba..01c9d5ee6e 100644
--- a/CUDATools/src/utilities.jl
+++ b/CUDATools/src/utilities.jl
@@ -56,10 +56,16 @@ function versioninfo(io::IO=stdout)
     println(io, "- NVML: ", has_nvml() ? NVML.version() : "missing")
     println(io)
 
+    get_module(name::Symbol) = (name, getfield(Metal, 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, "- CUDACore: $(Base.pkgversion(CUDACore))")
-    for name in [:GPUArrays, :GPUCompiler, :KernelAbstractions, :CUDA_Driver_jll,
-                 :CUDA_Compiler_jll, :CUDA_Runtime_jll, :CUDA_Runtime_Discovery,
+    for name in [:GPUArrays, :GPUCompiler, ("63c18a36-062a-441e-b654-da1e3ab1ce7c", "KernelAbstractions"),
+                 :CUDA_Driver_jll, :CUDA_Compiler_jll, :CUDA_Runtime_jll, :CUDA_Runtime_Discovery,
                  :NVPTX_LLVM_Backend_jll]
         isdefined(CUDACore, name) || continue
         mod = getfield(CUDACore, name)

From 5eff07bb5103ba0c8ab1a13542f4b912e7ce5900 Mon Sep 17 00:00:00 2001
From: Christian <28689358+christiangnrd@users.noreply.github.com>
Date: Sat, 22 Aug 2026 13:22:39 -0300
Subject: [PATCH 02/10] Optional printing of some packages in versioninfo

---
 CUDATools/src/utilities.jl | 13 ++++++-------
 1 file changed, 6 insertions(+), 7 deletions(-)

diff --git a/CUDATools/src/utilities.jl b/CUDATools/src/utilities.jl
index 01c9d5ee6e..ee86d8aafd 100644
--- a/CUDATools/src/utilities.jl
+++ b/CUDATools/src/utilities.jl
@@ -56,7 +56,7 @@ function versioninfo(io::IO=stdout)
     println(io, "- NVML: ", has_nvml() ? NVML.version() : "missing")
     println(io)
 
-    get_module(name::Symbol) = (name, getfield(Metal, name))
+    get_module(name::Symbol) = (name, getfield(CUDACore, 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))
@@ -64,12 +64,11 @@ function versioninfo(io::IO=stdout)
 
     println(io, "Julia packages: ")
     println(io, "- CUDACore: $(Base.pkgversion(CUDACore))")
-    for name in [:GPUArrays, :GPUCompiler, ("63c18a36-062a-441e-b654-da1e3ab1ce7c", "KernelAbstractions"),
-                 :CUDA_Driver_jll, :CUDA_Compiler_jll, :CUDA_Runtime_jll, :CUDA_Runtime_Discovery,
-                 :NVPTX_LLVM_Backend_jll]
-        isdefined(CUDACore, name) || continue
-        mod = getfield(CUDACore, name)
-        println(io, "- $(name): $(Base.pkgversion(mod))")
+    for pkg in [:GPUArrays, :GPUCompiler, ("63c18a36-062a-441e-b654-da1e3ab1ce7c", "KernelAbstractions"),
+                 :CUDA_Driver_jll, :CUDA_Compiler_jll, ("76a88914-d11a-5bdc-97e0-2f5a05c973a2", "CUDA_Runtime_jll"),
+                 ("1af6417a-86b4-443c-805f-a4643ffb695f", "CUDA_Runtime_Discovery"), :NVPTX_LLVM_Backend_jll]
+        name, mod = get_module(pkg)
+        isnothing(mod) || println(io, "- $(name): $(Base.pkgversion(mod))")
     end
     println(io)
 

From d670a43056d2f97181970ac60362aa4872336f98 Mon Sep 17 00:00:00 2001
From: Christian <28689358+christiangnrd@users.noreply.github.com>
Date: Sat, 22 Aug 2026 13:32:19 -0300
Subject: [PATCH 03/10] Add support for KernelInterface

---
 CUDACore/Project.toml          |   2 +
 CUDACore/src/CUDACore.jl       |   7 +-
 CUDACore/src/CUDAKernels.jl    | 244 ++++++++++++++++---------------
 CUDACore/src/CUDAKernelsOld.jl | 257 +++++++++++++++++++++++++++++++++
 CUDATools/src/utilities.jl     |   2 +-
 test/Project.toml              |   1 +
 test/core/kernelinterface.jl   |   6 +
 7 files changed, 394 insertions(+), 125 deletions(-)
 create mode 100644 CUDACore/src/CUDAKernelsOld.jl
 create mode 100644 test/core/kernelinterface.jl

diff --git a/CUDACore/Project.toml b/CUDACore/Project.toml
index 7e52054edf..29f2f3a52d 100644
--- a/CUDACore/Project.toml
+++ b/CUDACore/Project.toml
@@ -15,6 +15,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"
 LLVMLoopInfo = "8b046642-f1f6-4319-8d3c-209ddc03c586"
 LazyArtifacts = "4af54fe1-eca0-43a8-85a7-787d91b784e3"
@@ -57,6 +58,7 @@ GPUArrays = "11.5.4"
 GPUCompiler = "2.4"
 GPUToolbox = "3"
 KernelAbstractions = "0.9.38"
+KernelInterface = "0.1.0"
 LLVM = "9.6"
 LLVMLoopInfo = "1"
 LazyArtifacts = "1"
diff --git a/CUDACore/src/CUDACore.jl b/CUDACore/src/CUDACore.jl
index 8919d33242..5a577f1394 100644
--- a/CUDACore/src/CUDACore.jl
+++ b/CUDACore/src/CUDACore.jl
@@ -18,6 +18,7 @@ using LLVM.Interop
 using Core: LLVMPtr
 
 import KernelAbstractions
+import KernelInterface
 
 using Adapt: Adapt, adapt, WrappedArray
 
@@ -124,10 +125,14 @@ include("complex.jl")
 include("library_types.jl")
 
 # KernelAbstractions
-include("CUDAKernels.jl")
+include("CUDAKernelsOld.jl")
 import .CUDAKernels: CUDABackend
 export CUDABackend
 
+# KernelInterface - NOT PUBLIC. Use KernelInterface.get_backend on an CLArray to get the backend
+include("CUDAKernels.jl")
+import .CUDAInterface
+
 # StaticArrays is still a direct dependency, so directly include the extension
 include("../ext/StaticArraysExt.jl")
 # NOTE: StaticArrays is a direct dep, so extension is directly included
diff --git a/CUDACore/src/CUDAKernels.jl b/CUDACore/src/CUDAKernels.jl
index b8bc1f16cb..a171231c50 100644
--- a/CUDACore/src/CUDAKernels.jl
+++ b/CUDACore/src/CUDAKernels.jl
@@ -1,9 +1,9 @@
-module CUDAKernels
+module CUDAInterface
 
 using ..CUDACore
 using ..CUDACore: @device_override, default_memory, UnifiedMemory, GPUArrays
 
-import KernelAbstractions as KA
+import KernelInterface as KI
 
 import StaticArrays
 
@@ -13,31 +13,30 @@ import Adapt
 
 export CUDABackend
 
-struct CUDABackend <: KA.GPU
+struct CUDABackend <: KI.GPU
     prefer_blocks::Bool
     always_inline::Bool
 end
 
 CUDABackend(; prefer_blocks=false, always_inline=false) = CUDABackend(prefer_blocks, always_inline)
 
-@inline KA.allocate(::CUDABackend, ::Type{T}, dims::Tuple; unified::Bool = false) where T = CuArray{T, length(dims), unified ? UnifiedMemory : default_memory}(undef, dims)
-@inline KA.zeros(::CUDABackend, ::Type{T}, dims::Tuple; unified::Bool = false) where T = fill!(CuArray{T, length(dims), unified ? UnifiedMemory : default_memory}(undef, dims), zero(T))
-@inline KA.ones(::CUDABackend, ::Type{T}, dims::Tuple; unified::Bool = false) where T = fill!(CuArray{T, length(dims), unified ? UnifiedMemory : default_memory}(undef, dims), one(T))
+@inline KI.allocate(::CUDABackend, ::Type{T}, dims::Tuple; unified::Bool = false) where T = CuArray{T, length(dims), unified ? UnifiedMemory : default_memory}(undef, dims)
+@inline KI.zeros(::CUDABackend, ::Type{T}, dims::Tuple; unified::Bool = false) where T = fill!(CuArray{T, length(dims), unified ? UnifiedMemory : default_memory}(undef, dims), zero(T))
+@inline KI.ones(::CUDABackend, ::Type{T}, dims::Tuple; unified::Bool = false) where T = fill!(CuArray{T, length(dims), unified ? UnifiedMemory : default_memory}(undef, dims), one(T))
 
-KA.get_backend(::CuArray) = CUDABackend()
-KA.synchronize(::CUDABackend) = synchronize()
+KI.get_backend(::CuArray) = CUDABackend()
+KI.synchronize(::CUDABackend) = synchronize()
 
-KA.functional(::CUDABackend) = CUDACore.functional()
+KI.functional(::CUDABackend) = CUDACore.functional()
 
-KA.supports_unified(::CUDABackend) = true
+KI.supports_unified(::CUDABackend) = true
 
 Adapt.adapt_storage(::CUDABackend, a::AbstractArray) = Adapt.adapt(CuArray, a)
 Adapt.adapt_storage(::CUDABackend, a::Union{CuArray,GPUArrays.AbstractGPUSparseArray}) = a
-Adapt.adapt_storage(::KA.CPU, a::Union{CuArray,GPUArrays.AbstractGPUSparseArray}) = Adapt.adapt(Array, a)
 
 ## memory operations
 
-function KA.copyto!(::CUDABackend, A, B)
+function KI.copyto!(::CUDABackend, A, B)
     GC.@preserve A B begin
         destptr = pointer(A)
         srcptr  = pointer(B)
@@ -47,172 +46,175 @@ function KA.copyto!(::CUDABackend, A, B)
     return A
 end
 
-function KA.pagelock!(::CUDABackend, A::Array)
+function KI.pagelock!(::CUDABackend, A::Array)
     CUDACore.pin(A)
     return nothing
 end
 
 ## device operations
 
-function KA.ndevices(::CUDABackend)
+function KI.ndevices(::CUDABackend)
     return Int(ndevices())
 end
 
-function KA.device(::CUDABackend)::Int
+function KI.device(::CUDABackend)::Int
     deviceid(CUDACore.active_state().device) + 1
 end
 
-function KA.device!(backend::CUDABackend, id::Int)
-    if !(0 < id <= KA.ndevices(backend))
+function KI.device!(backend::CUDABackend, id::Int)
+    if !(0 < id <= KI.ndevices(backend))
         throw(ArgumentError("Device id $id out of bounds."))
     end
     device!(id - 1)
 end
 
-## kernel launch
+# function (obj::KA.Kernel{CUDABackend})(args...; ndrange=nothing, workgroupsize=nothing)
+#     backend = KA.backend(obj)
 
-function KA.mkcontext(kernel::KA.Kernel{CUDABackend}, _ndrange, iterspace)
-    KA.CompilerMetadata{KA.ndrange(kernel), KA.DynamicCheck}(_ndrange, iterspace)
-end
+#     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)
 
-function KA.launch_config(kernel::KA.Kernel{CUDABackend}, ndrange, workgroupsize)
-    if ndrange isa Integer
-        ndrange = (ndrange,)
-    end
-    if workgroupsize isa Integer
-        workgroupsize = (workgroupsize, )
-    end
+#     # If the kernel is statically sized we can tell the compiler about that
+#     if KA.workgroupsize(obj) <: KA.StaticSize
+#         maxthreads = prod(KA.get(KA.workgroupsize(obj)))
+#     else
+#         maxthreads = nothing
+#     end
 
-    # partition checked that the ndrange's agreed
-    if KA.ndrange(kernel) <: KA.StaticSize
-        ndrange = nothing
-    end
+#     call = CUDACore.kernel_call(obj.f, (ctx, args...))
+#     kernel = CUDACore.kernel_compile(call; always_inline=backend.always_inline, maxthreads)
 
-    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
+#     # figure out the optimal workgroupsize automatically
+#     if KA.workgroupsize(obj) <: KA.DynamicSize && workgroupsize === nothing
+#         config = CUDACore.launch_configuration(kernel.fun; max_threads=prod(ndrange))
+#         if backend.prefer_blocks
+#             # Prefer blocks over threads
+#             threads = min(prod(ndrange), config.threads)
+#             # XXX: Some kernels performs much better with all blocks active
+#             cu_blocks = max(cld(prod(ndrange), threads), config.blocks)
+#             threads = cld(prod(ndrange), cu_blocks)
+#         else
+#             threads = config.threads
+#         end
 
-    return ndrange, workgroupsize, iterspace, dynamic
-end
+#         workgroupsize = threads_to_workgroupsize(threads, ndrange)
+#         iterspace, dynamic = KA.partition(obj, ndrange, workgroupsize)
+#         ctx = KA.mkcontext(obj, ndrange, iterspace)
+#         call = CUDACore.rebind(call, ctx, 1)
+#     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
+#     blocks = length(KA.blocks(iterspace))
+#     threads = length(KA.workitems(iterspace))
 
-function (obj::KA.Kernel{CUDABackend})(args...; ndrange=nothing, workgroupsize=nothing)
-    backend = KA.backend(obj)
+#     if blocks == 0
+#         return nothing
+#     end
 
-    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)
+#     # Launch kernel
+#     CUDACore.kernel_launch(kernel, call; threads, blocks)
 
-    # If the kernel is statically sized we can tell the compiler about that
-    if KA.workgroupsize(obj) <: KA.StaticSize
-        maxthreads = prod(KA.get(KA.workgroupsize(obj)))
-    else
-        maxthreads = nothing
-    end
+#     return nothing
+# end
 
-    call = CUDACore.kernel_call(obj.f, (ctx, args...))
-    kernel = CUDACore.kernel_compile(call; always_inline=backend.always_inline, maxthreads)
-
-    # figure out the optimal workgroupsize automatically
-    if KA.workgroupsize(obj) <: KA.DynamicSize && workgroupsize === nothing
-        config = CUDACore.launch_configuration(kernel.fun; max_threads=prod(ndrange))
-        if backend.prefer_blocks
-            # Prefer blocks over threads
-            threads = min(prod(ndrange), config.threads)
-            # XXX: Some kernels performs much better with all blocks active
-            cu_blocks = max(cld(prod(ndrange), threads), config.blocks)
-            threads = cld(prod(ndrange), cu_blocks)
-        else
-            threads = config.threads
-        end
-
-        workgroupsize = threads_to_workgroupsize(threads, ndrange)
-        iterspace, dynamic = KA.partition(obj, ndrange, workgroupsize)
-        ctx = KA.mkcontext(obj, ndrange, iterspace)
-        call = CUDACore.rebind(call, ctx, 1)
-    end
+KI.argconvert(::CUDABackend, arg) = cudaconvert(arg)
 
-    blocks = length(KA.blocks(iterspace))
-    threads = length(KA.workitems(iterspace))
+function KI.kernel_function(::CUDABackend, f::F, tt::TT=Tuple{}; name=nothing, kwargs...) where {F,TT}
+    kern = cufunction(f, tt; name, kwargs...)
+    KI.Kernel{CUDABackend, typeof(kern)}(CUDABackend(), kern)
+end
 
-    if blocks == 0
-        return nothing
-    end
+function (obj::KI.Kernel{CUDABackend})(args...; numworkgroups=(), workgroupsize=(), ndrange=(), max_work_group_size=typemax(Int))
+    KI.check_launch_args(numworkgroups, workgroupsize, ndrange)
+    prod(ndrange) == 0 && return nothing
 
-    # Launch kernel
-    CUDACore.kernel_launch(kernel, call; threads, blocks)
+    numworkgroups, workgroupsize = KI.auto_launch_sizes(obj, numworkgroups, workgroupsize, ndrange, max_work_group_size)
 
+    obj.kern(args...; threads=workgroupsize, blocks=numworkgroups)
     return nothing
 end
 
-## indexing
 
-## COV_EXCL_START
-@device_override @inline function KA.__index_Local_Linear(ctx)
-    return threadIdx().x
+function KI.kernel_max_work_group_size(kernel::KI.Kernel{<:CUDABackend}; max_work_items::Int=typemax(Int))::Int
+    kernel_config = launch_configuration(kernel.kern.fun)
+
+    Int(min(kernel_config.threads, max_work_items))
+end
+function KI.max_work_group_size(::CUDABackend)::Int
+    Int(attribute(device(), CUDACore.DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK))
+end
+function KI.sub_group_size(::CUDABackend)::Int
+    warpsize(device())
+end
+function KI.multiprocessor_count(::CUDABackend)::Int
+    Int(attribute(device(), CUDACore.DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT))
 end
 
+KI.shfl_down_types(::CUDABackend) = DataType[Bool,
+                                             UInt8, UInt16, UInt32, UInt64, UInt128,
+                                             Int8, Int16, Int32, Int64, Int128,
+                                             Float16, Float32, Float64,
+                                             ComplexF16, ComplexF32, ComplexF64]
 
-@device_override @inline function KA.__index_Group_Linear(ctx)
-    return blockIdx().x
+## indexing
+
+## COV_EXCL_START
+@device_override @inline function KI.get_local_id()
+    return (; x = Int(threadIdx().x), y = Int(threadIdx().y), z = Int(threadIdx().z))
 end
 
-@device_override @inline function KA.__index_Global_Linear(ctx)
-    I =  @inbounds KA.expand(KA.__iterspace(ctx), blockIdx().x, threadIdx().x)
-    # TODO: This is unfortunate, can we get the linear index cheaper
-    @inbounds LinearIndices(KA.__ndrange(ctx))[I]
+@device_override @inline function KI.get_group_id()
+    return (; x = Int(blockIdx().x), y = Int(blockIdx().y), z = Int(blockIdx().z))
 end
 
-@device_override @inline function KA.__index_Local_Cartesian(ctx)
-    @inbounds KA.workitems(KA.__iterspace(ctx))[threadIdx().x]
+@device_override @inline function KI.get_global_id()
+    return (; x = Int((blockIdx().x-1)*blockDim().x + threadIdx().x), y = Int((blockIdx().y-1)*blockDim().y + threadIdx().y), z = Int((blockIdx().z-1)*blockDim().z + threadIdx().z))
 end
 
-@device_override @inline function KA.__index_Group_Cartesian(ctx)
-    @inbounds KA.blocks(KA.__iterspace(ctx))[blockIdx().x]
+@device_override @inline function KI.get_local_size()
+    return (; x = Int(blockDim().x), y = Int(blockDim().y), z = Int(blockDim().z))
 end
 
-@device_override @inline function KA.__index_Global_Cartesian(ctx)
-    return @inbounds KA.expand(KA.__iterspace(ctx), blockIdx().x, threadIdx().x)
+@device_override @inline function KI.get_num_groups()
+    return (; x = Int(gridDim().x), y = Int(gridDim().y), z = Int(gridDim().z))
 end
 
-@device_override @inline function KA.__validindex(ctx)
-    if KA.__dynamic_checkbounds(ctx)
-        I = @inbounds KA.expand(KA.__iterspace(ctx), blockIdx().x, threadIdx().x)
-        return I in KA.__ndrange(ctx)
-    else
-        return true
-    end
+@device_override @inline function KI.get_global_size()
+    return (; x = Int(blockDim().x * gridDim().x), y = Int(blockDim().y * gridDim().y), z = Int(blockDim().z * gridDim().z))
 end
 
+@device_override KI.get_sub_group_size() = UInt32(warpsize())
+
+@device_override KI.get_max_sub_group_size() = UInt32(warpsize())
+
+@device_override KI.get_num_sub_groups() = UInt32(prod(blockDim()) ÷ warpsize())
+
+@device_override KI.get_sub_group_id() = UInt32(((threadIdx().x - 1) + blockDim().x * (threadIdx().y - 1) + blockDim().x * blockDim().y * (threadIdx().z - 1)) ÷ warpsize()) + 0x1
+
+@device_override KI.get_sub_group_local_id() = UInt32(laneid())
+
+
 ## 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}
     CuStaticSharedArray(T, Dims)
 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()
     sync_threads()
 end
 
-@device_override @inline function KA.__print(args...)
+@device_override @inline function KI.sub_group_barrier()
+    sync_warp()
+end
+
+@device_override function KI.shfl_down(val::T, offset::Integer) where T
+    shfl_down_sync(0xffffffff, val, offset)
+end
+
+@device_override @inline function KI._print(args...)
     CUDACore._cuprint(args...)
 end
 
@@ -220,11 +222,7 @@ end
 
 ## other
 
-Adapt.adapt_storage(to::KA.ConstAdaptor, a::CuDeviceArray) = Base.Experimental.Const(a)
-
-KA.argconvert(k::KA.Kernel{CUDABackend}, arg) = cudaconvert(arg)
-
-function KA.priority!(::CUDABackend, prio::Symbol)
+function KI.priority!(::CUDABackend, prio::Symbol)
     if !(prio in (:high, :normal, :low))
         error("priority must be one of :high, :normal, :low")
     end
diff --git a/CUDACore/src/CUDAKernelsOld.jl b/CUDACore/src/CUDAKernelsOld.jl
new file mode 100644
index 0000000000..b8bc1f16cb
--- /dev/null
+++ b/CUDACore/src/CUDAKernelsOld.jl
@@ -0,0 +1,257 @@
+module CUDAKernels
+
+using ..CUDACore
+using ..CUDACore: @device_override, default_memory, UnifiedMemory, GPUArrays
+
+import KernelAbstractions as KA
+
+import StaticArrays
+
+import Adapt
+
+## back-end
+
+export CUDABackend
+
+struct CUDABackend <: KA.GPU
+    prefer_blocks::Bool
+    always_inline::Bool
+end
+
+CUDABackend(; prefer_blocks=false, always_inline=false) = CUDABackend(prefer_blocks, always_inline)
+
+@inline KA.allocate(::CUDABackend, ::Type{T}, dims::Tuple; unified::Bool = false) where T = CuArray{T, length(dims), unified ? UnifiedMemory : default_memory}(undef, dims)
+@inline KA.zeros(::CUDABackend, ::Type{T}, dims::Tuple; unified::Bool = false) where T = fill!(CuArray{T, length(dims), unified ? UnifiedMemory : default_memory}(undef, dims), zero(T))
+@inline KA.ones(::CUDABackend, ::Type{T}, dims::Tuple; unified::Bool = false) where T = fill!(CuArray{T, length(dims), unified ? UnifiedMemory : default_memory}(undef, dims), one(T))
+
+KA.get_backend(::CuArray) = CUDABackend()
+KA.synchronize(::CUDABackend) = synchronize()
+
+KA.functional(::CUDABackend) = CUDACore.functional()
+
+KA.supports_unified(::CUDABackend) = true
+
+Adapt.adapt_storage(::CUDABackend, a::AbstractArray) = Adapt.adapt(CuArray, a)
+Adapt.adapt_storage(::CUDABackend, a::Union{CuArray,GPUArrays.AbstractGPUSparseArray}) = a
+Adapt.adapt_storage(::KA.CPU, a::Union{CuArray,GPUArrays.AbstractGPUSparseArray}) = Adapt.adapt(Array, a)
+
+## memory operations
+
+function KA.copyto!(::CUDABackend, A, B)
+    GC.@preserve A B begin
+        destptr = pointer(A)
+        srcptr  = pointer(B)
+        N       = length(A)
+        unsafe_copyto!(destptr, srcptr, N, async=true)
+    end
+    return A
+end
+
+function KA.pagelock!(::CUDABackend, A::Array)
+    CUDACore.pin(A)
+    return nothing
+end
+
+## device operations
+
+function KA.ndevices(::CUDABackend)
+    return Int(ndevices())
+end
+
+function KA.device(::CUDABackend)::Int
+    deviceid(CUDACore.active_state().device) + 1
+end
+
+function KA.device!(backend::CUDABackend, id::Int)
+    if !(0 < id <= KA.ndevices(backend))
+        throw(ArgumentError("Device id $id out of bounds."))
+    end
+    device!(id - 1)
+end
+
+## kernel launch
+
+function KA.mkcontext(kernel::KA.Kernel{CUDABackend}, _ndrange, iterspace)
+    KA.CompilerMetadata{KA.ndrange(kernel), KA.DynamicCheck}(_ndrange, iterspace)
+end
+
+function KA.launch_config(kernel::KA.Kernel{CUDABackend}, 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{CUDABackend})(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
+        maxthreads = prod(KA.get(KA.workgroupsize(obj)))
+    else
+        maxthreads = nothing
+    end
+
+    call = CUDACore.kernel_call(obj.f, (ctx, args...))
+    kernel = CUDACore.kernel_compile(call; always_inline=backend.always_inline, maxthreads)
+
+    # figure out the optimal workgroupsize automatically
+    if KA.workgroupsize(obj) <: KA.DynamicSize && workgroupsize === nothing
+        config = CUDACore.launch_configuration(kernel.fun; max_threads=prod(ndrange))
+        if backend.prefer_blocks
+            # Prefer blocks over threads
+            threads = min(prod(ndrange), config.threads)
+            # XXX: Some kernels performs much better with all blocks active
+            cu_blocks = max(cld(prod(ndrange), threads), config.blocks)
+            threads = cld(prod(ndrange), cu_blocks)
+        else
+            threads = config.threads
+        end
+
+        workgroupsize = threads_to_workgroupsize(threads, ndrange)
+        iterspace, dynamic = KA.partition(obj, ndrange, workgroupsize)
+        ctx = KA.mkcontext(obj, ndrange, iterspace)
+        call = CUDACore.rebind(call, ctx, 1)
+    end
+
+    blocks = length(KA.blocks(iterspace))
+    threads = length(KA.workitems(iterspace))
+
+    if blocks == 0
+        return nothing
+    end
+
+    # Launch kernel
+    CUDACore.kernel_launch(kernel, call; threads, blocks)
+
+    return nothing
+end
+
+## indexing
+
+## COV_EXCL_START
+@device_override @inline function KA.__index_Local_Linear(ctx)
+    return threadIdx().x
+end
+
+
+@device_override @inline function KA.__index_Group_Linear(ctx)
+    return blockIdx().x
+end
+
+@device_override @inline function KA.__index_Global_Linear(ctx)
+    I =  @inbounds KA.expand(KA.__iterspace(ctx), blockIdx().x, threadIdx().x)
+    # TODO: This is unfortunate, can we get the linear index cheaper
+    @inbounds LinearIndices(KA.__ndrange(ctx))[I]
+end
+
+@device_override @inline function KA.__index_Local_Cartesian(ctx)
+    @inbounds KA.workitems(KA.__iterspace(ctx))[threadIdx().x]
+end
+
+@device_override @inline function KA.__index_Group_Cartesian(ctx)
+    @inbounds KA.blocks(KA.__iterspace(ctx))[blockIdx().x]
+end
+
+@device_override @inline function KA.__index_Global_Cartesian(ctx)
+    return @inbounds KA.expand(KA.__iterspace(ctx), blockIdx().x, threadIdx().x)
+end
+
+@device_override @inline function KA.__validindex(ctx)
+    if KA.__dynamic_checkbounds(ctx)
+        I = @inbounds KA.expand(KA.__iterspace(ctx), blockIdx().x, threadIdx().x)
+        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}
+    CuStaticSharedArray(T, Dims)
+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()
+    sync_threads()
+end
+
+@device_override @inline function KA.__print(args...)
+    CUDACore._cuprint(args...)
+end
+
+## COV_EXCL_STOP
+
+## other
+
+Adapt.adapt_storage(to::KA.ConstAdaptor, a::CuDeviceArray) = Base.Experimental.Const(a)
+
+KA.argconvert(k::KA.Kernel{CUDABackend}, arg) = cudaconvert(arg)
+
+function KA.priority!(::CUDABackend, prio::Symbol)
+    if !(prio in (:high, :normal, :low))
+        error("priority must be one of :high, :normal, :low")
+    end
+
+    range = priority_range()
+    # 0:-1:-5
+    # lower number is higher priority, default is 0
+    # there is no "low"
+    if prio === :high
+        priority = last(range)
+    elseif prio === :normal || prio === :low
+        priority = first(range)
+    end
+
+    old_stream = stream()
+    r_flags = Ref{Cuint}()
+    CUDACore.cuStreamGetFlags(old_stream, r_flags)
+    flags = CUDACore.CUstream_flags_enum(r_flags[])
+
+    event = CuEvent(CUDACore.EVENT_DISABLE_TIMING)
+    record(event, old_stream)
+
+    @debug "Switching default stream" flags priority
+    new_stream = CuStream(; flags, priority)
+    CUDACore.wait(event, new_stream)
+    stream!(new_stream)
+    return nothing
+end
+
+end
diff --git a/CUDATools/src/utilities.jl b/CUDATools/src/utilities.jl
index ee86d8aafd..b042432ced 100644
--- a/CUDATools/src/utilities.jl
+++ b/CUDATools/src/utilities.jl
@@ -64,7 +64,7 @@ function versioninfo(io::IO=stdout)
 
     println(io, "Julia packages: ")
     println(io, "- CUDACore: $(Base.pkgversion(CUDACore))")
-    for pkg in [:GPUArrays, :GPUCompiler, ("63c18a36-062a-441e-b654-da1e3ab1ce7c", "KernelAbstractions"),
+    for pkg in [:GPUArrays, :GPUCompiler, ("63c18a36-062a-441e-b654-da1e3ab1ce7c", "KernelAbstractions"), :KernelInterface,
                  :CUDA_Driver_jll, :CUDA_Compiler_jll, ("76a88914-d11a-5bdc-97e0-2f5a05c973a2", "CUDA_Runtime_jll"),
                  ("1af6417a-86b4-443c-805f-a4643ffb695f", "CUDA_Runtime_Discovery"), :NVPTX_LLVM_Backend_jll]
         name, mod = get_module(pkg)
diff --git a/test/Project.toml b/test/Project.toml
index 2b0579e832..1cff754b82 100644
--- a/test/Project.toml
+++ b/test/Project.toml
@@ -18,6 +18,7 @@ InteractiveUtils = "b77e0a4c-d291-57a0-90e8-8db25a27a240"
 Interpolations = "a98d9a8b-a2ab-59e6-89dd-64a1c18fca59"
 JLD2 = "033835bb-8acc-5ee8-8aae-3f567f8a3819"
 KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c"
+KernelInterface = "4ee993da-d684-4d17-a7dd-4e58e78d92bf"
 LinearAlgebra = "37e2e46d-f89d-539d-b4ee-838fcccc9c8e"
 NVML = "611af6d1-644e-4c5d-bd58-854d7d1254b9"
 NVTX = "5da4648a-3479-48b8-97b9-01cb529c0a1f"
diff --git a/test/core/kernelinterface.jl b/test/core/kernelinterface.jl
new file mode 100644
index 0000000000..d3004bae17
--- /dev/null
+++ b/test/core/kernelinterface.jl
@@ -0,0 +1,6 @@
+import KernelInterface
+using CUDACore.CUDAInterface
+
+include(joinpath(dirname(pathof(KernelInterface)), "..", "test", "testsuite.jl"))
+
+Testsuite.testsuite(CUDAInterface.CUDABackend, "CUDACore", CUDACore, CuArray, CUDACore.CuDeviceArray)

From 610c1f79c23ac5f5215f943292c7ecb43f3cc925 Mon Sep 17 00:00:00 2001
From: Christian <28689358+christiangnrd@users.noreply.github.com>
Date: Sat, 22 Aug 2026 17:02:03 -0300
Subject: [PATCH 04/10] KernelAbstractions 0.10 as extension

---
 CUDACore/Project.toml                 |   5 +-
 CUDACore/ext/KernelAbstractionsExt.jl | 130 +++++++++++++
 CUDACore/src/CUDACore.jl              |   9 +-
 CUDACore/src/CUDAKernels.jl           |   2 +-
 CUDACore/src/CUDAKernelsOld.jl        | 257 --------------------------
 lib/cusparse/Project.toml             |   2 +-
 test/core/kernelinterface.jl          |   4 +-
 7 files changed, 139 insertions(+), 270 deletions(-)
 create mode 100644 CUDACore/ext/KernelAbstractionsExt.jl
 delete mode 100644 CUDACore/src/CUDAKernelsOld.jl

diff --git a/CUDACore/Project.toml b/CUDACore/Project.toml
index 29f2f3a52d..7c91c438eb 100644
--- a/CUDACore/Project.toml
+++ b/CUDACore/Project.toml
@@ -14,7 +14,6 @@ ExprTools = "e2ba6199-217a-4e67-a87a-7c52f15ade04"
 GPUArrays = "0c68f7d7-f131-5f86-a1c3-88cf8149b2d7"
 GPUCompiler = "61eb1bfa-7361-4325-ad38-22787b887f55"
 GPUToolbox = "096a3bc2-3ced-46d0-87f4-dd12716f4bfc"
-KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c"
 KernelInterface = "4ee993da-d684-4d17-a7dd-4e58e78d92bf"
 LLVM = "929cbde3-209d-540e-8aea-75f648917ca0"
 LLVMLoopInfo = "8b046642-f1f6-4319-8d3c-209ddc03c586"
@@ -35,11 +34,13 @@ StaticArrays = "90137ffa-7385-5640-81b9-e52037218182"
 CUDA = "052768ef-5323-5732-b1bb-66c8b64840ba"
 ChainRulesCore = "d360d2e6-b24c-11e9-a2a3-2a2ae2dbcce4"
 EnzymeCore = "f151be2c-9106-41f4-ab19-57ee4f262869"
+KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c"
 SpecialFunctions = "276daf66-3868-5448-9aa4-cd146d93841b"
 
 [extensions]
 ChainRulesCoreExt = "ChainRulesCore"
 EnzymeCoreExt = "EnzymeCore"
+KernelAbstractionsExt = "KernelAbstractions"
 SpecialFunctionsExt = "SpecialFunctions"
 
 [compat]
@@ -57,7 +58,7 @@ ExprTools = "0.1"
 GPUArrays = "11.5.4"
 GPUCompiler = "2.4"
 GPUToolbox = "3"
-KernelAbstractions = "0.9.38"
+KernelAbstractions = "0.10"
 KernelInterface = "0.1.0"
 LLVM = "9.6"
 LLVMLoopInfo = "1"
diff --git a/CUDACore/ext/KernelAbstractionsExt.jl b/CUDACore/ext/KernelAbstractionsExt.jl
new file mode 100644
index 0000000000..b3efa00b92
--- /dev/null
+++ b/CUDACore/ext/KernelAbstractionsExt.jl
@@ -0,0 +1,130 @@
+module KernelAbstractionsExt
+
+using CUDACore
+using CUDACore: @device_override, default_memory, UnifiedMemory, GPUArrays
+
+import KernelAbstractions as KA
+
+
+import StaticArrays
+
+import Adapt
+
+# TODO: Move AbstractGPUSparseArray stuff out
+Adapt.adapt_storage(::KA.CPU, a::Union{CuArray,GPUArrays.AbstractGPUSparseArray}) = Adapt.adapt(Array, a)
+
+## kernel launch
+
+function KA.mkcontext(kernel::KA.Kernel{CUDABackend}, _ndrange, iterspace)
+    KA.CompilerMetadata{KA.ndrange(kernel), KA.DynamicCheck}(_ndrange, iterspace)
+end
+
+function KA.launch_config(kernel::KA.Kernel{CUDABackend}, 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 = Ref(1)
+    return map(ndrange) do n
+        x = min(div(threads, total[]), n)
+        total[] *= x
+        return x
+    end
+end
+
+function (obj::KA.Kernel{CUDABackend})(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
+        maxthreads = prod(KA.get(KA.workgroupsize(obj)))
+    else
+        maxthreads = nothing
+    end
+
+    call = CUDACore.kernel_call(obj.f, (ctx, args...))
+    kernel = CUDACore.kernel_compile(call; always_inline=backend.always_inline, maxthreads)
+
+    # figure out the optimal workgroupsize automatically
+    if KA.workgroupsize(obj) <: KA.DynamicSize && workgroupsize === nothing
+        config = CUDACore.launch_configuration(kernel.fun; max_threads=prod(ndrange))
+        if backend.prefer_blocks
+            # Prefer blocks over threads
+            threads = min(prod(ndrange), config.threads)
+            # XXX: Some kernels performs much better with all blocks active
+            cu_blocks = max(cld(prod(ndrange), threads), config.blocks)
+            threads = cld(prod(ndrange), cu_blocks)
+        else
+            threads = config.threads
+        end
+
+        workgroupsize = threads_to_workgroupsize(threads, ndrange)
+        iterspace, dynamic = KA.partition(obj, ndrange, workgroupsize)
+        ctx = KA.mkcontext(obj, ndrange, iterspace)
+        call = CUDACore.rebind(call, ctx, 1)
+    end
+
+    blocks = length(KA.blocks(iterspace))
+    threads = length(KA.workitems(iterspace))
+
+    if blocks == 0
+        return nothing
+    end
+
+    # Launch kernel
+    CUDACore.kernel_launch(kernel, call; threads, blocks)
+
+    return nothing
+end
+
+## indexing
+
+## COV_EXCL_START
+@device_override @inline function KA.__validindex(ctx)
+    if KA.__dynamic_checkbounds(ctx)
+        I = @inbounds KA.expand(KA.__iterspace(ctx), blockIdx().x, threadIdx().x)
+        return I in KA.__ndrange(ctx)
+    else
+        return true
+    end
+end
+
+## shared and scratch memory
+
+@device_override @inline function KA.Scratchpad(ctx, ::Type{T}, ::Val{Dims}) where {T, Dims}
+    StaticArrays.MArray{KA.__size(Dims), T}(undef)
+end
+
+## COV_EXCL_STOP
+
+## other
+
+Adapt.adapt_storage(to::KA.ConstAdaptor, a::CuDeviceArray) = Base.Experimental.Const(a)
+
+KA.argconvert(k::KA.Kernel{CUDABackend}, arg) = cudaconvert(arg)
+
+end
diff --git a/CUDACore/src/CUDACore.jl b/CUDACore/src/CUDACore.jl
index 5a577f1394..462c01a1b8 100644
--- a/CUDACore/src/CUDACore.jl
+++ b/CUDACore/src/CUDACore.jl
@@ -17,7 +17,6 @@ using LLVM
 using LLVM.Interop
 using Core: LLVMPtr
 
-import KernelAbstractions
 import KernelInterface
 
 using Adapt: Adapt, adapt, WrappedArray
@@ -124,15 +123,11 @@ include("sorting.jl")
 include("complex.jl")
 include("library_types.jl")
 
-# KernelAbstractions
-include("CUDAKernelsOld.jl")
+# KernelInterface
+include("CUDAKernels.jl")
 import .CUDAKernels: CUDABackend
 export CUDABackend
 
-# KernelInterface - NOT PUBLIC. Use KernelInterface.get_backend on an CLArray to get the backend
-include("CUDAKernels.jl")
-import .CUDAInterface
-
 # StaticArrays is still a direct dependency, so directly include the extension
 include("../ext/StaticArraysExt.jl")
 # NOTE: StaticArrays is a direct dep, so extension is directly included
diff --git a/CUDACore/src/CUDAKernels.jl b/CUDACore/src/CUDAKernels.jl
index a171231c50..07d99e0975 100644
--- a/CUDACore/src/CUDAKernels.jl
+++ b/CUDACore/src/CUDAKernels.jl
@@ -1,4 +1,4 @@
-module CUDAInterface
+module CUDAKernels
 
 using ..CUDACore
 using ..CUDACore: @device_override, default_memory, UnifiedMemory, GPUArrays
diff --git a/CUDACore/src/CUDAKernelsOld.jl b/CUDACore/src/CUDAKernelsOld.jl
deleted file mode 100644
index b8bc1f16cb..0000000000
--- a/CUDACore/src/CUDAKernelsOld.jl
+++ /dev/null
@@ -1,257 +0,0 @@
-module CUDAKernels
-
-using ..CUDACore
-using ..CUDACore: @device_override, default_memory, UnifiedMemory, GPUArrays
-
-import KernelAbstractions as KA
-
-import StaticArrays
-
-import Adapt
-
-## back-end
-
-export CUDABackend
-
-struct CUDABackend <: KA.GPU
-    prefer_blocks::Bool
-    always_inline::Bool
-end
-
-CUDABackend(; prefer_blocks=false, always_inline=false) = CUDABackend(prefer_blocks, always_inline)
-
-@inline KA.allocate(::CUDABackend, ::Type{T}, dims::Tuple; unified::Bool = false) where T = CuArray{T, length(dims), unified ? UnifiedMemory : default_memory}(undef, dims)
-@inline KA.zeros(::CUDABackend, ::Type{T}, dims::Tuple; unified::Bool = false) where T = fill!(CuArray{T, length(dims), unified ? UnifiedMemory : default_memory}(undef, dims), zero(T))
-@inline KA.ones(::CUDABackend, ::Type{T}, dims::Tuple; unified::Bool = false) where T = fill!(CuArray{T, length(dims), unified ? UnifiedMemory : default_memory}(undef, dims), one(T))
-
-KA.get_backend(::CuArray) = CUDABackend()
-KA.synchronize(::CUDABackend) = synchronize()
-
-KA.functional(::CUDABackend) = CUDACore.functional()
-
-KA.supports_unified(::CUDABackend) = true
-
-Adapt.adapt_storage(::CUDABackend, a::AbstractArray) = Adapt.adapt(CuArray, a)
-Adapt.adapt_storage(::CUDABackend, a::Union{CuArray,GPUArrays.AbstractGPUSparseArray}) = a
-Adapt.adapt_storage(::KA.CPU, a::Union{CuArray,GPUArrays.AbstractGPUSparseArray}) = Adapt.adapt(Array, a)
-
-## memory operations
-
-function KA.copyto!(::CUDABackend, A, B)
-    GC.@preserve A B begin
-        destptr = pointer(A)
-        srcptr  = pointer(B)
-        N       = length(A)
-        unsafe_copyto!(destptr, srcptr, N, async=true)
-    end
-    return A
-end
-
-function KA.pagelock!(::CUDABackend, A::Array)
-    CUDACore.pin(A)
-    return nothing
-end
-
-## device operations
-
-function KA.ndevices(::CUDABackend)
-    return Int(ndevices())
-end
-
-function KA.device(::CUDABackend)::Int
-    deviceid(CUDACore.active_state().device) + 1
-end
-
-function KA.device!(backend::CUDABackend, id::Int)
-    if !(0 < id <= KA.ndevices(backend))
-        throw(ArgumentError("Device id $id out of bounds."))
-    end
-    device!(id - 1)
-end
-
-## kernel launch
-
-function KA.mkcontext(kernel::KA.Kernel{CUDABackend}, _ndrange, iterspace)
-    KA.CompilerMetadata{KA.ndrange(kernel), KA.DynamicCheck}(_ndrange, iterspace)
-end
-
-function KA.launch_config(kernel::KA.Kernel{CUDABackend}, 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{CUDABackend})(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
-        maxthreads = prod(KA.get(KA.workgroupsize(obj)))
-    else
-        maxthreads = nothing
-    end
-
-    call = CUDACore.kernel_call(obj.f, (ctx, args...))
-    kernel = CUDACore.kernel_compile(call; always_inline=backend.always_inline, maxthreads)
-
-    # figure out the optimal workgroupsize automatically
-    if KA.workgroupsize(obj) <: KA.DynamicSize && workgroupsize === nothing
-        config = CUDACore.launch_configuration(kernel.fun; max_threads=prod(ndrange))
-        if backend.prefer_blocks
-            # Prefer blocks over threads
-            threads = min(prod(ndrange), config.threads)
-            # XXX: Some kernels performs much better with all blocks active
-            cu_blocks = max(cld(prod(ndrange), threads), config.blocks)
-            threads = cld(prod(ndrange), cu_blocks)
-        else
-            threads = config.threads
-        end
-
-        workgroupsize = threads_to_workgroupsize(threads, ndrange)
-        iterspace, dynamic = KA.partition(obj, ndrange, workgroupsize)
-        ctx = KA.mkcontext(obj, ndrange, iterspace)
-        call = CUDACore.rebind(call, ctx, 1)
-    end
-
-    blocks = length(KA.blocks(iterspace))
-    threads = length(KA.workitems(iterspace))
-
-    if blocks == 0
-        return nothing
-    end
-
-    # Launch kernel
-    CUDACore.kernel_launch(kernel, call; threads, blocks)
-
-    return nothing
-end
-
-## indexing
-
-## COV_EXCL_START
-@device_override @inline function KA.__index_Local_Linear(ctx)
-    return threadIdx().x
-end
-
-
-@device_override @inline function KA.__index_Group_Linear(ctx)
-    return blockIdx().x
-end
-
-@device_override @inline function KA.__index_Global_Linear(ctx)
-    I =  @inbounds KA.expand(KA.__iterspace(ctx), blockIdx().x, threadIdx().x)
-    # TODO: This is unfortunate, can we get the linear index cheaper
-    @inbounds LinearIndices(KA.__ndrange(ctx))[I]
-end
-
-@device_override @inline function KA.__index_Local_Cartesian(ctx)
-    @inbounds KA.workitems(KA.__iterspace(ctx))[threadIdx().x]
-end
-
-@device_override @inline function KA.__index_Group_Cartesian(ctx)
-    @inbounds KA.blocks(KA.__iterspace(ctx))[blockIdx().x]
-end
-
-@device_override @inline function KA.__index_Global_Cartesian(ctx)
-    return @inbounds KA.expand(KA.__iterspace(ctx), blockIdx().x, threadIdx().x)
-end
-
-@device_override @inline function KA.__validindex(ctx)
-    if KA.__dynamic_checkbounds(ctx)
-        I = @inbounds KA.expand(KA.__iterspace(ctx), blockIdx().x, threadIdx().x)
-        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}
-    CuStaticSharedArray(T, Dims)
-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()
-    sync_threads()
-end
-
-@device_override @inline function KA.__print(args...)
-    CUDACore._cuprint(args...)
-end
-
-## COV_EXCL_STOP
-
-## other
-
-Adapt.adapt_storage(to::KA.ConstAdaptor, a::CuDeviceArray) = Base.Experimental.Const(a)
-
-KA.argconvert(k::KA.Kernel{CUDABackend}, arg) = cudaconvert(arg)
-
-function KA.priority!(::CUDABackend, prio::Symbol)
-    if !(prio in (:high, :normal, :low))
-        error("priority must be one of :high, :normal, :low")
-    end
-
-    range = priority_range()
-    # 0:-1:-5
-    # lower number is higher priority, default is 0
-    # there is no "low"
-    if prio === :high
-        priority = last(range)
-    elseif prio === :normal || prio === :low
-        priority = first(range)
-    end
-
-    old_stream = stream()
-    r_flags = Ref{Cuint}()
-    CUDACore.cuStreamGetFlags(old_stream, r_flags)
-    flags = CUDACore.CUstream_flags_enum(r_flags[])
-
-    event = CuEvent(CUDACore.EVENT_DISABLE_TIMING)
-    record(event, old_stream)
-
-    @debug "Switching default stream" flags priority
-    new_stream = CuStream(; flags, priority)
-    CUDACore.wait(event, new_stream)
-    stream!(new_stream)
-    return nothing
-end
-
-end
diff --git a/lib/cusparse/Project.toml b/lib/cusparse/Project.toml
index 5911310cb5..db445cb91d 100644
--- a/lib/cusparse/Project.toml
+++ b/lib/cusparse/Project.toml
@@ -32,7 +32,7 @@ CUDA_Runtime_Discovery = "2"
 CUDA_Runtime_jll = "0.24"
 GPUArrays = "11.4.1"
 GPUToolbox = "3"
-KernelAbstractions = "0.9.38"
+KernelAbstractions = "0.10"
 LinearAlgebra = "1"
 SparseArrays = "1"
 SparseMatricesCSR = "0.6.9"
diff --git a/test/core/kernelinterface.jl b/test/core/kernelinterface.jl
index d3004bae17..109ab6e5a9 100644
--- a/test/core/kernelinterface.jl
+++ b/test/core/kernelinterface.jl
@@ -1,6 +1,6 @@
 import KernelInterface
-using CUDACore.CUDAInterface
+using CUDACore
 
 include(joinpath(dirname(pathof(KernelInterface)), "..", "test", "testsuite.jl"))
 
-Testsuite.testsuite(CUDAInterface.CUDABackend, "CUDACore", CUDACore, CuArray, CUDACore.CuDeviceArray)
+Testsuite.testsuite(CUDABackend, "CUDACore", CUDACore, CuArray, CUDACore.CuDeviceArray)

From ea8b69ac1476e70ead6a7f2e73db5ca6b601a58c Mon Sep 17 00:00:00 2001
From: Christian <28689358+christiangnrd@users.noreply.github.com>
Date: Sat, 22 Aug 2026 17:02:11 -0300
Subject: [PATCH 05/10] Temp source KA

---
 test/Project.toml | 1 +
 1 file changed, 1 insertion(+)

diff --git a/test/Project.toml b/test/Project.toml
index 1cff754b82..f3715b2953 100644
--- a/test/Project.toml
+++ b/test/Project.toml
@@ -48,6 +48,7 @@ CUDA = {path = ".."}
 CUDACore = {path = "../CUDACore"}
 CUDATools = {path = "../CUDATools"}
 CUPTI = {path = "../lib/cupti"}
+KernelAbstractions = {rev = "main", url = "https://github.com/JuliaGPU/KernelAbstractions.jl"}
 NVML = {path = "../lib/nvml"}
 cuBLAS = {path = "../lib/cublas"}
 cuDNN = {path = "../lib/cudnn"}

From 23ea6712bee4fb44442b723c9420cb2b17c856d7 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Mon, 17 Nov 2025 21:26:56 -0400
Subject: [PATCH 06/10] Dogfood

---
 CUDACore/src/CUDACore.jl      |  2 +-
 CUDACore/src/accumulate.jl    | 26 ++++++++++++------------
 CUDACore/src/device/random.jl |  6 +++---
 CUDACore/src/indexing.jl      |  2 +-
 CUDACore/src/mapreduce.jl     | 37 ++++++++++++++++++++++-------------
 5 files changed, 41 insertions(+), 32 deletions(-)

diff --git a/CUDACore/src/CUDACore.jl b/CUDACore/src/CUDACore.jl
index 462c01a1b8..81d9bba6f1 100644
--- a/CUDACore/src/CUDACore.jl
+++ b/CUDACore/src/CUDACore.jl
@@ -17,7 +17,7 @@ using LLVM
 using LLVM.Interop
 using Core: LLVMPtr
 
-import KernelInterface
+import KernelInterface as KI
 
 using Adapt: Adapt, adapt, WrappedArray
 
diff --git a/CUDACore/src/accumulate.jl b/CUDACore/src/accumulate.jl
index a20d76e828..3b3646752b 100644
--- a/CUDACore/src/accumulate.jl
+++ b/CUDACore/src/accumulate.jl
@@ -15,16 +15,16 @@
 function partial_scan(op::Function, output::AbstractArray{T}, input::AbstractArray,
                       Rdim, Rpre, Rpost, Rother, neutral, init,
                       ::Val{inclusive}=Val(true)) where {T, inclusive}
-    threads = blockDim().x
-    thread = threadIdx().x
-    block = blockIdx().x
+    threads = KI.get_local_size().x
+    thread = KI.get_local_id().x
+    block = KI.get_group_id().x
 
     temp = CuDynamicSharedArray(T, (2*threads,))
 
     # iterate the main dimension using threads and the first block dimension
-    i = (blockIdx().x-1i32) * blockDim().x + threadIdx().x
+    i = (KI.get_group_id().x-1i32) * KI.get_local_size().x + KI.get_local_id().x
     # iterate the other dimensions using the remaining block dimensions
-    j = (blockIdx().z-1i32) * gridDim().y + blockIdx().y
+    j = (KI.get_group_id().z-1i32) * KI.get_num_groups().y + KI.get_group_id().y
 
     if j > length(Rother)
         return
@@ -47,7 +47,7 @@ function partial_scan(op::Function, output::AbstractArray{T}, input::AbstractArr
     offset = 1
     d = threads>>1
     while d > 0
-        sync_threads()
+        KI.barrier()
         @inbounds if thread <= d
             ai = offset * (2*thread-1)
             bi = offset * (2*thread)
@@ -66,7 +66,7 @@ function partial_scan(op::Function, output::AbstractArray{T}, input::AbstractArr
     d = 1
     while d < threads
         offset >>= 1
-        sync_threads()
+        KI.barrier()
         @inbounds if thread <= d
             ai = offset * (2*thread-1)
             bi = offset * (2*thread)
@@ -78,7 +78,7 @@ function partial_scan(op::Function, output::AbstractArray{T}, input::AbstractArr
         d *= 2
     end
 
-    sync_threads()
+    KI.barrier()
 
     # write results to device memory
     @inbounds if i <= length(Rdim)
@@ -100,14 +100,14 @@ end
 function aggregate_partial_scan(op::Function, output::AbstractArray,
                                 aggregates::AbstractArray, Rdim, Rpre, Rpost, Rother,
                                 init)
-    threads = blockDim().x
-    thread = threadIdx().x
-    block = blockIdx().x
+    threads = KI.get_local_size().x
+    thread = KI.get_local_id().x
+    block = KI.get_group_id().x
 
     # iterate the main dimension using threads and the first block dimension
-    i = (blockIdx().x-1i32) * blockDim().x + threadIdx().x
+    i = (KI.get_group_id().x-1i32) * KI.get_local_size().x + KI.get_local_id().x
     # iterate the other dimensions using the remaining block dimensions
-    j = (blockIdx().z-1i32) * gridDim().y + blockIdx().y
+    j = (KI.get_group_id().z-1i32) * KI.get_num_groups().y + KI.get_group_id().y
 
     @inbounds if i <= length(Rdim) && j <= length(Rother)
         I = Rother[j]
diff --git a/CUDACore/src/device/random.jl b/CUDACore/src/device/random.jl
index e3267694c4..0c2fcf483e 100644
--- a/CUDACore/src/device/random.jl
+++ b/CUDACore/src/device/random.jl
@@ -72,9 +72,9 @@ end
     elseif field === :ctr1
         @inbounds global_random_counters()[warpId]
     elseif field === :ctr2
-        blockId = blockIdx().x + (blockIdx().y - 1i32) * gridDim().x +
-                                 (blockIdx().z - 1i32) * gridDim().x * gridDim().y
-        globalId = threadId + (blockId - 1i32) * (blockDim().x * blockDim().y * blockDim().z)
+        globalId = KI.get_global_id().x +
+                   (KI.get_global_id().y - 1i32) * KI.get_global_size().x +
+                   (KI.get_global_id().z - 1i32) * KI.get_global_size().x * KI.get_global_size().y
         globalId%UInt32
     end::UInt32
 end
diff --git a/CUDACore/src/indexing.jl b/CUDACore/src/indexing.jl
index 46f5593ebf..9b9a5353ba 100644
--- a/CUDACore/src/indexing.jl
+++ b/CUDACore/src/indexing.jl
@@ -33,7 +33,7 @@ function Base.findall(bools::AnyCuArray{Bool})
     if n > 0
         ## COV_EXCL_START
         function kernel(ys::CuDeviceArray, bools, indices)
-            i = threadIdx().x + (blockIdx().x - 1i32) * blockDim().x
+            i = KI.get_local_id().x + (KI.get_group_id().x - 1i32) * KI.get_local_size().x
 
             @inbounds if i <= length(bools) && bools[i]
                 i′ = CartesianIndices(bools)[i]
diff --git a/CUDACore/src/mapreduce.jl b/CUDACore/src/mapreduce.jl
index 7f30984db5..3bd5efdb52 100644
--- a/CUDACore/src/mapreduce.jl
+++ b/CUDACore/src/mapreduce.jl
@@ -19,9 +19,9 @@ end
 @inline function reduce_block(op, val::T, neutral, shuffle::Val{true}) where T
     # shared mem for partial sums
     assume(warpsize() == 32)
-    shared = CuStaticSharedArray(T, 32)
+    shared = KI.localmemory(T, 32)
 
-    wid, lane = fldmod1(threadIdx().x, warpsize())
+    wid, lane = fldmod1(KI.get_local_id().x, warpsize())
 
     # each warp performs partial reduction
     val = reduce_warp(op, val)
@@ -32,10 +32,10 @@ end
     end
 
     # wait for all partial reductions
-    sync_threads()
+    KI.barrier()
 
     # read from shared memory only if that warp existed
-    val = if threadIdx().x <= fld1(blockDim().x, warpsize())
+    val = if KI.get_local_id().x <= fld1(KI.get_local_size().x, warpsize())
          @inbounds shared[lane]
     else
         neutral
@@ -49,8 +49,8 @@ end
     return val
 end
 @inline function reduce_block(op, val::T, neutral, shuffle::Val{false}) where T
-    threads = blockDim().x
-    thread = threadIdx().x
+    threads = KI.get_local_size().x
+    thread = KI.get_local_id().x
 
     # shared mem for a complete reduction
     shared = CuDynamicSharedArray(T, (threads,))
@@ -59,7 +59,7 @@ end
     # perform a reduction
     d = 1
     while d < threads
-        sync_threads()
+        KI.barrier()
         index = 2 * d * (thread-1) + 1
         @inbounds if index <= threads
             other_val = if index + d <= threads
@@ -92,10 +92,10 @@ function partial_mapreduce_grid(f, op, neutral, Rreduce, Rother, shuffle, R::Abs
 
     # decompose the 1D hardware indices into separate ones for reduction (across threads
     # and possibly blocks if it doesn't fit) and other elements (remaining blocks)
-    threadIdx_reduce = threadIdx().x
-    blockDim_reduce = blockDim().x
-    blockIdx_reduce, blockIdx_other = fldmod1(blockIdx().x, length(Rother))
-    gridDim_reduce = gridDim().x ÷ length(Rother)
+    threadIdx_reduce = KI.get_local_id().x
+    blockDim_reduce = KI.get_local_size().x
+    blockIdx_reduce, blockIdx_other = fldmod1(KI.get_group_id().x, length(Rother))
+    gridDim_reduce = KI.get_num_groups().x ÷ length(Rother)
 
     # block-based indexing into the values outside of the reduction dimension
     # (that means we can safely synchronize threads within this block)
@@ -134,7 +134,7 @@ function partial_mapreduce_grid(f, op, neutral, Rreduce, Rother, shuffle, R::Abs
 end
 
 function serial_mapreduce_kernel(f, op, neutral, Rreduce, Rother, R, As)
-    grid_idx = threadIdx().x + (blockIdx().x - 1i32) * blockDim().x
+    grid_idx = KI.get_local_id().x + (KI.get_group_id().x - 1i32) * KI.get_local_size().x
     @inbounds if grid_idx <= length(Rother)
         Iother = Rother[grid_idx]
 
@@ -160,14 +160,14 @@ end
 
 # factored out for use in tests
 function serial_mapreduce_threshold(dev)
-    max_concurrency = attribute(dev, DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK) *
-                      attribute(dev, DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT)
+    max_concurrency = KI.max_work_group_size(CUDABackend()) * KI.multiprocessor_count(CUDABackend())
     return max_concurrency
 end
 
 function GPUArrays.mapreducedim!(f::F, op::OP, R::AnyCuArray{T},
                                  A::Union{AbstractArray,Broadcast.Broadcasted};
                                  init=nothing) where {F, OP, T}
+    backend = CUDABackend()
     if !isa(A, Broadcast.Broadcasted)
         # XXX: Base.axes isn't defined anymore for Broadcasted, breaking this check
         Base.check_reducedims(R, A)
@@ -202,11 +202,14 @@ function GPUArrays.mapreducedim!(f::F, op::OP, R::AnyCuArray{T},
     # If `Rother` is large enough, then a naive loop is more efficient than partial reductions.
     if length(Rother) >= serial_mapreduce_threshold(dev)
         args = (f, op, init, Rreduce, Rother, R, A)
+        # kernel = KI.KIKernel(backend, serial_mapreduce_kernel, args...)
+        # kernel_config = launch_configuration(kernel.kern.fun)
         call = KernelCall(serial_mapreduce_kernel, args...)
         kernel = kernel_compile(call)
         kernel_config = launch_configuration(kernel.fun)
         threads = kernel_config.threads
         blocks = cld(length(Rother), threads)
+        # kernel(args...; workgroupsize=threads, numworkgroups=blocks)
         kernel_launch(kernel, call; threads, blocks)
         return R_old
     end
@@ -230,7 +233,9 @@ function GPUArrays.mapreducedim!(f::F, op::OP, R::AnyCuArray{T},
     # we might not be able to launch all those threads to reduce each slice in one go.
     # that's why each threads also loops across their inputs, processing multiple values
     # so that we can span the entire reduction dimension using a single thread block.
+    # kernel = KI.KIKernel(backend, partial_mapreduce_grid, f, op, init, Rreduce, Rother, Val(shuffle), R, A)
     compute_shmem(threads) = shuffle ? 0 : threads*sizeof(T)
+    # kernel_config = launch_configuration(kernel.kern.fun; shmem=compute_shmem∘compute_threads)
     call = KernelCall(partial_mapreduce_grid, f, op, init, Rreduce, Rother, Val(shuffle), R, A)
     kernel = kernel_compile(call)
     kernel_config = launch_configuration(kernel.fun; shmem=compute_shmem∘compute_threads)
@@ -258,6 +263,7 @@ function GPUArrays.mapreducedim!(f::F, op::OP, R::AnyCuArray{T},
     # perform the actual reduction
     if reduce_blocks == 1
         # we can cover the dimensions to reduce using a single block
+        # kernel(f, op, init, Rreduce, Rother, Val(shuffle), R, A; workgroupsize=threads, numworkgroups=blocks, shmem)
         kernel_launch(kernel, call; threads, blocks, shmem)
     else
         # TODO: provide a version that atomically reduces from different blocks
@@ -268,9 +274,11 @@ function GPUArrays.mapreducedim!(f::F, op::OP, R::AnyCuArray{T},
         # NOTE: we can't use the previously-compiled kernel, or its launch configuration,
         #       since the type of `partial` might not match the original output container
         #       (e.g. if that was a view).
+        # partial_kernel = KI.KIKernel(backend, partial_mapreduce_grid, f, op, init, Rreduce, Rother, Val(shuffle), partial, A)
         partial_call = KernelCall(partial_mapreduce_grid, f, op, init, Rreduce, Rother,
                                   Val(shuffle), partial, A)
         partial_kernel = kernel_compile(partial_call)
+        # partial_kernel_config = launch_configuration(partial_kernel.kern.fun; shmem=compute_shmem∘compute_threads)
         partial_kernel_config = launch_configuration(partial_kernel.fun; shmem=compute_shmem∘compute_threads)
         partial_reduce_threads = compute_threads(partial_kernel_config.threads)
         partial_reduce_shmem = compute_shmem(partial_reduce_threads)
@@ -293,6 +301,7 @@ function GPUArrays.mapreducedim!(f::F, op::OP, R::AnyCuArray{T},
         partial_call = rebind(partial_call, partial, 7)
         kernel_launch(partial_kernel, partial_call;
                threads=partial_threads, blocks=partial_blocks, shmem=partial_shmem)
+               # workgroupsize=partial_threads, numworkgroups=partial_blocks, shmem=partial_shmem)
 
         GPUArrays.mapreducedim!(identity, op, R, partial; init)
     end

From 2482f74e7b052496aaf6f964856b51d0aa0046bd Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Mon, 17 Nov 2025 21:25:09 -0400
Subject: [PATCH 07/10] Temp CI

---
 CUDACore/Project.toml | 3 +++
 test/runtests.jl      | 5 +++++
 2 files changed, 8 insertions(+)

diff --git a/CUDACore/Project.toml b/CUDACore/Project.toml
index 7c91c438eb..4727182d9e 100644
--- a/CUDACore/Project.toml
+++ b/CUDACore/Project.toml
@@ -77,6 +77,9 @@ SpecialFunctions = "1.3, 2"
 StaticArrays = "1.9.8"
 julia = "1.10"
 
+[sources]
+KernelAbstractions = {rev = "main", url = "https://github.com/JuliaGPU/KernelAbstractions.jl"}
+
 [extras]
 ChainRulesCore = "d360d2e6-b24c-11e9-a2a3-2a2ae2dbcce4"
 EnzymeCore = "f151be2c-9106-41f4-ab19-57ee4f262869"
diff --git a/test/runtests.jl b/test/runtests.jl
index e0ef61fb76..128b266fce 100644
--- a/test/runtests.jl
+++ b/test/runtests.jl
@@ -1,3 +1,8 @@
+@static if VERSION < v"1.11" && get(ENV, "BUILDKITE_PIPELINE_NAME", "CUDA.jl") == "CUDA.jl"
+    using Pkg
+    Pkg.add(url="https://github.com/JuliaGPU/KernelAbstractions.jl", rev="main")
+end
+
 using CUDA
 using CUDACore
 using cuBLAS, cuSPARSE, cuSOLVER, cuFFT, cuRAND

From c52be1d7a88e4109383378814f8362fd32792731 Mon Sep 17 00:00:00 2001
From: Christian <28689358+christiangnrd@users.noreply.github.com>
Date: Sat, 22 Aug 2026 17:11:19 -0300
Subject: [PATCH 08/10] Fixup KA tests

---
 test/core/kernelabstractions.jl | 5 ++---
 1 file changed, 2 insertions(+), 3 deletions(-)

diff --git a/test/core/kernelabstractions.jl b/test/core/kernelabstractions.jl
index 4b0693e0e3..7ed73bc0f0 100644
--- a/test/core/kernelabstractions.jl
+++ b/test/core/kernelabstractions.jl
@@ -25,9 +25,8 @@ end
 
 include(joinpath(dirname(pathof(KernelAbstractions)), "..", "test", "testsuite.jl"))
 
-ka_skip_tests = Set{String}(["sparse"])
-Testsuite.testsuite(()->CUDABackend(false, false), "CUDA", CUDA, CuArray, CuDeviceArray;
-                    skip_tests=ka_skip_tests)
+ka_skip_tests = Set{String}(["CPU synchronization", "fallback test: callable types"])
+Testsuite.testsuite(()->CUDABackend(false, false), "CUDA", CUDA, CuArray, CuDeviceArray; skip_tests=ka_skip_tests)
 for (PreferBlocks, AlwaysInline) in Iterators.product((true, false), (true, false))
     Testsuite.unittest_testsuite(()->CUDABackend(PreferBlocks, AlwaysInline), "CUDA", CUDA, CuDeviceArray;
                                  skip_tests=ka_skip_tests)

From 5f5b5f63281dda05ac4234f85c6312b4f624f269 Mon Sep 17 00:00:00 2001
From: Christian <28689358+christiangnrd@users.noreply.github.com>
Date: Sat, 22 Aug 2026 18:33:41 -0300
Subject: [PATCH 09/10] Fix

---
 CUDACore/Project.toml | 5 +----
 1 file changed, 1 insertion(+), 4 deletions(-)

diff --git a/CUDACore/Project.toml b/CUDACore/Project.toml
index 4727182d9e..57720b6fb3 100644
--- a/CUDACore/Project.toml
+++ b/CUDACore/Project.toml
@@ -58,7 +58,7 @@ ExprTools = "0.1"
 GPUArrays = "11.5.4"
 GPUCompiler = "2.4"
 GPUToolbox = "3"
-KernelAbstractions = "0.10"
+KernelAbstractions = "0.9.38, 0.10"
 KernelInterface = "0.1.0"
 LLVM = "9.6"
 LLVMLoopInfo = "1"
@@ -77,9 +77,6 @@ SpecialFunctions = "1.3, 2"
 StaticArrays = "1.9.8"
 julia = "1.10"
 
-[sources]
-KernelAbstractions = {rev = "main", url = "https://github.com/JuliaGPU/KernelAbstractions.jl"}
-
 [extras]
 ChainRulesCore = "d360d2e6-b24c-11e9-a2a3-2a2ae2dbcce4"
 EnzymeCore = "f151be2c-9106-41f4-ab19-57ee4f262869"

From c6602995d0e310f3c2fd68ef0c5dac5bede71c60 Mon Sep 17 00:00:00 2001
From: Christian <28689358+christiangnrd@users.noreply.github.com>
Date: Thu, 27 Aug 2026 15:34:33 -0300
Subject: [PATCH 10/10] [temp]

---
 lib/cusparse/Project.toml | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/lib/cusparse/Project.toml b/lib/cusparse/Project.toml
index db445cb91d..6046a8244d 100644
--- a/lib/cusparse/Project.toml
+++ b/lib/cusparse/Project.toml
@@ -32,7 +32,7 @@ CUDA_Runtime_Discovery = "2"
 CUDA_Runtime_jll = "0.24"
 GPUArrays = "11.4.1"
 GPUToolbox = "3"
-KernelAbstractions = "0.10"
+KernelAbstractions = "0.9.38, 0.10"
 LinearAlgebra = "1"
 SparseArrays = "1"
 SparseMatricesCSR = "0.6.9"