From 3a61b0f52c268e3011b45e858ae7d7bff41cd8c9 Mon Sep 17 00:00:00 2001 From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com> Date: Mon, 17 Nov 2025 22:02:15 -0400 Subject: [PATCH 1/3] Skip CPU-only tests --- test/kernelabstractions.jl | 2 ++ 1 file changed, 2 insertions(+) diff --git a/test/kernelabstractions.jl b/test/kernelabstractions.jl index fd400bc1..640d79e2 100644 --- a/test/kernelabstractions.jl +++ b/test/kernelabstractions.jl @@ -33,5 +33,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) From 2a946a144f6b1a48b4e7deead273e20fc68bd8a3 Mon Sep 17 00:00:00 2001 From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com> Date: Wed, 19 Aug 2026 15:13:32 -0300 Subject: [PATCH 2/3] KA 0.10 and turn into an extension --- Project.toml | 7 +- .../KernelAbstractionsExt.jl | 94 +------------------ src/OpenCL.jl | 8 +- src/OpenCLKernels.jl | 4 +- test/Project.toml | 1 + test/kernelabstractions.jl | 1 - test/kernelinterface.jl | 4 +- 7 files changed, 12 insertions(+), 107 deletions(-) rename src/OpenCLKernelsOld.jl => ext/KernelAbstractionsExt.jl (53%) diff --git a/Project.toml b/Project.toml index 9315a705..25613a66 100644 --- 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,7 +25,11 @@ 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] +KernelAbstractions = {rev = "main", url = "https://github.com/JuliaGPU/KernelAbstractions.jl"} KernelInterface = {rev = "main", subdir = "lib/KernelInterface", url = "https://github.com/JuliaGPU/KernelAbstractions.jl"} SPIRVIntrinsics = {path = "lib/intrinsics"} @@ -34,7 +37,7 @@ SPIRVIntrinsics = {path = "lib/intrinsics"} 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 b/ext/KernelAbstractionsExt.jl similarity index 53% rename from src/OpenCLKernelsOld.jl rename to ext/KernelAbstractionsExt.jl index 5bca9355..c89133e2 100644 --- a/src/OpenCLKernelsOld.jl +++ b/ext/KernelAbstractionsExt.jl @@ -1,4 +1,4 @@ -module OpenCLKernels +module KernelAbstractionsExt using ..OpenCL using ..OpenCL: @device_override, method_table @@ -9,49 +9,8 @@ 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 index 8f24d85f..5b1ae0e3 100644 --- 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 index 48028beb..d097bff0 100644 --- a/src/OpenCLKernels.jl +++ b/src/OpenCLKernels.jl @@ -1,4 +1,4 @@ -module OpenCLInterface +module OpenCLKernels using ..OpenCL using ..OpenCL: @device_override, method_table, kernel_convert, clfunction @@ -14,7 +14,7 @@ import Adapt ## Back-end Definition -# export OpenCLBackend +export OpenCLBackend struct OpenCLBackend <: KI.GPU end diff --git a/test/Project.toml b/test/Project.toml index 7fe8eb83..d4368d40 100644 --- 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 index 640d79e2..fda63bdb 100644 --- a/test/kernelabstractions.jl +++ b/test/kernelabstractions.jl @@ -32,7 +32,6 @@ end skip_tests=Set([ "sparse", - "Convert", # Need to opt out of i128 "CPU synchronization", "fallback test: callable types" ]) diff --git a/test/kernelinterface.jl b/test/kernelinterface.jl index cc6c2083..bbeb2467 100644 --- a/test/kernelinterface.jl +++ b/test/kernelinterface.jl @@ -1,6 +1,6 @@ import KernelInterface -using OpenCL.OpenCLInterface +using OpenCL include(joinpath(dirname(pathof(KernelInterface)), "..", "test", "testsuite.jl")) -Testsuite.testsuite(OpenCLInterface.OpenCLBackend, "OpenCL", OpenCL, CLArray, OpenCL.CLDeviceArray) +Testsuite.testsuite(OpenCLBackend, "OpenCL", OpenCL, CLArray, OpenCL.CLDeviceArray) From 7188c42a62f71b2ce704a789b25ab5f2c687b414 Mon Sep 17 00:00:00 2001 From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com> Date: Wed, 19 Aug 2026 15:16:22 -0300 Subject: [PATCH 3/3] [temp] source KA 0.10 --- Project.toml | 3 +++ test/Project.toml | 1 + 2 files changed, 4 insertions(+) diff --git a/Project.toml b/Project.toml index 25613a66..80c68b24 100644 --- a/Project.toml +++ b/Project.toml @@ -54,3 +54,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/test/Project.toml b/test/Project.toml index d4368d40..9f10076c 100644 --- a/test/Project.toml +++ b/test/Project.toml @@ -31,6 +31,7 @@ pocl_next_jll = "59abdad9-3cfc-5436-8271-411e8cad6b82" [sources] KernelAbstractions = {rev = "main", url = "https://github.com/JuliaGPU/KernelAbstractions.jl"} +KernelInterface = {rev = "main", subdir = "lib/KernelInterface", url = "https://github.com/JuliaGPU/KernelAbstractions.jl"} OpenCL = {path = ".."} SPIRVIntrinsics = {path = "../lib/intrinsics"}