JuliaGPU · christiangnrd · Nov 18, 2025 · Aug 19, 2026 · Aug 22, 2026
diff --git a/Project.toml b/Project.toml
@@ -9,7 +9,6 @@ projects = ["lib/intrinsics", "test", "docs", "res"]
 Adapt = "79e6a3ab-5dfb-504d-930d-738a2a938a0e"
 GPUArrays = "0c68f7d7-f131-5f86-a1c3-88cf8149b2d7"
 GPUCompiler = "61eb1bfa-7361-4325-ad38-22787b887f55"
-KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c"
 KernelInterface = "4ee993da-d684-4d17-a7dd-4e58e78d92bf"
 LLVM = "929cbde3-209d-540e-8aea-75f648917ca0"
 LinearAlgebra = "37e2e46d-f89d-539d-b4ee-838fcccc9c8e"
@@ -26,14 +25,20 @@ 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"
-KernelAbstractions = "0.9.38"
+KernelAbstractions = "0.10"
 KernelInterface = "0.1"
 LLVM = "9.6"
 LinearAlgebra = "1"

diff --git a/src/OpenCLKernelsOld.jl → ext/KernelAbstractionsExt.jl b/src/OpenCLKernelsOld.jl → ext/KernelAbstractionsExt.jl
@@ -1,57 +1,16 @@
-module OpenCLKernels
+module KernelAbstractionsExt
 
-using ..OpenCL
-using ..OpenCL: @device_override, method_table
+using OpenCL
+using OpenCL: @device_override, method_table
 
 import KernelAbstractions as KA
 
 import StaticArrays
 
 import Adapt
 
-
-## Back-end Definition
-
-export OpenCLBackend
-
-struct OpenCLBackend <: KA.GPU
-end
-
-function KA.allocate(::OpenCLBackend, ::Type{T}, dims::Tuple; unified::Bool = false) where T
-    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)
 
-
-## 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)
@@ -126,35 +85,6 @@ function (obj::KA.Kernel{OpenCLBackend})(args...; ndrange=nothing, workgroupsize
     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)
@@ -164,32 +94,10 @@ end
     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/OpenCL.jl b/src/OpenCL.jl
@@ -9,8 +9,6 @@ using GPUArrays
 using Random
 using Preferences
 
-import KernelAbstractions: KernelAbstractions
-
 import KernelInterface
 
 using Core: LLVMPtr
@@ -49,12 +47,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
@@ -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
 
 struct OpenCLBackend <: KI.GPU
 end

diff --git a/test/Project.toml b/test/Project.toml
@@ -30,6 +30,7 @@ pocl_jll = "627d6b7a-bbe6-5189-83e7-98cc0a5aeadd"
 pocl_next_jll = "59abdad9-3cfc-5436-8271-411e8cad6b82"
 
 [sources]
+KernelAbstractions = {rev = "main", url = "https://github.com/JuliaGPU/KernelAbstractions.jl"}
 OpenCL = {path = ".."}
 SPIRVIntrinsics = {path = "../lib/intrinsics"}
 

diff --git a/test/kernelabstractions.jl b/test/kernelabstractions.jl
@@ -32,6 +32,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
@@ -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)