make all tests pass, rename backends

lib/cl/kernel.jl

@@ -87,27 +87,6 @@ function set_arg!(k::Kernel, idx::Integer, backend::Type{<:CLBackend}, arg::Abst
     return k
 end
 
-#=
-# OpenCL USM Buffers
-## when passing using `cl.call`
-function set_arg!(k::Kernel, idx::Integer, arg::AbstractBuffer)
-    ext_clSetKernelArgMemPointerINTEL(k, cl_uint(idx - 1), arg.ptr)
-    return k
-end
-## when passing with `clcall`, which has pre-converted the buffer
-function set_arg!(k::Kernel, idx::Integer, arg::CLPtr{T}) where {T}
-    arg = reinterpret(Ptr{Cvoid}, arg)
-    if arg != C_NULL
-        # XXX: this assumes that the receiving argument is pointer-typed, which is not the
-        #      case with Julia's `Ptr` ABI. Instead, one should reinterpret the pointer as a
-        #      `Core.LLVMPtr`, which _is_ pointer-valued. We retain this handling for `Ptr`
-        #      for users passing pointers to OpenCL C, and because `Ptr` is pointer-valued
-        #      starting with Julia 1.12.
-        ext_clSetKernelArgMemPointerINTEL(k, cl_uint(idx - 1), arg)
-    end
-    return k
-end
-=#
 function set_arg!(k::Kernel, idx::Integer, backend::Type{<:CLBackend}, arg::LocalMem)
     clSetKernelArg(k, cl_uint(idx - 1), arg.nbytes, C_NULL)
     return k
@@ -206,7 +185,7 @@ end
 
 function call(
         k::Kernel, args...; global_size = (1,), local_size = nothing,
-        global_work_offset = nothing, wait_on::Vector{Event} = Event[], backend = Ref{Type{<:CLBackend}}(USM),
+        global_work_offset = nothing, wait_on::Vector{Event} = Event[], backend = Ref{Type{<:CLBackend}}(USMBackend),
         pointers::Vector{CLPtr} = CLPtr[]
     )
     set_args!(k, backend[], args...)

lib/cl/memory/backend.jl

@@ -1,27 +1,29 @@
 abstract type CLBackend end
 
-struct SVM <: CLBackend end
+struct SVMBackend <: CLBackend end
 
-struct USM <: CLBackend end
+struct USMBackend <: CLBackend end
+
+struct MemBackend <: CLBackend end
 
 function get_backend_from_buffer(x::Type{<:AbstractBuffer})
     if x == cl.SVMBuffer
-        SVM
+        SVMBackend
     else
-        USM
+        USMBackend
     end
 end
 
 function abstract_kernel_exec_info_ptrs(backend::Type{<:CLBackend})
-    if backend == SVM
+    if backend == SVMBackend
         CL_KERNEL_EXEC_INFO_SVM_PTRS
     else
         CL_KERNEL_EXEC_INFO_USM_PTRS_INTEL
     end
 end
 
 function set_kernel_arg_abstract_pointer(backend::Type{<:CLBackend})
-    if backend == SVM
+    if backend == SVMBackend
         clSetKernelArgSVMPointer
     else
         ext_clSetKernelArgMemPointerINTEL

lib/cl/memory/memory.jl

@@ -52,18 +52,18 @@ end
 # generic memory operations for different buffers
 
 function enqueue_abstract_memcpy(dst::Union{Ptr, CLPtr}, src::Union{Ptr, CLPtr}, nbytes::Integer; queu::CmdQueue=queue(), blocking::Bool=false,
-    wait_for::Vector{Event}=Event[], backend = USM)
-    if backend == USM
+    wait_for::Vector{Event}=Event[], backend = USMBackend)
+    if backend == USMBackend
         enqueue_usm_memcpy(dst, src, nbytes; queu = queu, blocking = blocking, wait_for = wait_for)
-    elseif backend == SVM
+    elseif backend == SVMBackend
         enqueue_usm_memcpy(dst, src, nbytes; queu = queu, blocking = blocking, wait_for = wait_for)
     end
 end
 
 function enqueue_abstract_fill(ptr::Union{Ptr, CLPtr}, pattern::Union{Ptr, CLPtr}, pattern_size::Integer, nbytes::Integer; queu::CmdQueue=queue(), wait_for::Vector{Event}=Event[], backend = USM)
-    if backend == USM
+    if backend == USMBackend
         enqueue_usm_memfill(ptr, pattern, pattern_size, nbytes; queu = queu, wait_for = wait_for)
-    elseif backend == SVM
+    elseif backend == SVMBackend
         enqueue_svm_fill(ptr, pattern, pattern_size, nbytes; queu = queu, wait_for = wait_for)
     end
 end

lib/cl/modified_fns.jl

@@ -55,7 +55,7 @@ end
 function ext_clSetKernelArgMemPointerINTEL(kernel, arg_index, arg_value)
     ocl_intel = ocl_extension("clSetKernelArgMemPointerINTEL")
 
-    return @ccall $ocl_intel(kernel::cl_kernel, arg_index::cl_uint, arg_value::Ptr{Cvoid})::cl_int
+    return @ccall $ocl_intel(kernel::cl_kernel, arg_index::cl_uint, arg_value::PtrOrCLPtr{Cvoid})::cl_int
 end
 
 function ext_clEnqueueMemAdviseINTEL(command_queue, ptr, size, advice, num_events_in_wait_list, event_wait_list, event)

src/array.jl

@@ -354,7 +354,7 @@ for (srcty, dstty) in [(:Array, :CLArray), (:CLArray, :Array), (:CLArray, :CLArr
         function Base.unsafe_copyto!(
                 dst::$dstty{T}, dst_off::Int,
                 src::$srcty{T}, src_off::Int,
-                N::Int; blocking::Bool = true, backend = cl.USM
+                N::Int; blocking::Bool = true, backend = cl.USMBackend
             ) where {T}
             nbytes = N * sizeof(T)
             println(buftype(dst), buftype(src))

src/compiler/execution.jl

@@ -103,7 +103,7 @@ function Adapt.adapt_storage(to::KernelAdaptor, buf::cl.SVMBuffer)
     ptr = pointer(buf)
     T = get_clptr_type(ptr)
     push!(to.pointers, ptr)
-    to.backend[] = cl.SVM
+    to.backend[] = cl.SVMBackend
     return reinterpret(Ptr{T}, ptr)
 end
 
@@ -153,7 +153,7 @@ input object `x` as-is.
 Do not add methods to this function, but instead extend the underlying Adapt.jl package and
 register methods for the the `OpenCL.KernelAdaptor` type.
 """
-kernel_convert(arg, backend = Ref{Type{<:cl.CLBackend}}(cl.USM), pointers::Vector{CLPtr}=CLPtr[]) = adapt(KernelAdaptor(pointers, backend), arg)
+kernel_convert(arg, backend = Ref{Type{<:cl.CLBackend}}(cl.USMBackend), pointers::Vector{CLPtr}=CLPtr[]) = adapt(KernelAdaptor(pointers, backend), arg)
 
 ## abstract kernel functionality
 
@@ -183,7 +183,7 @@ abstract type AbstractKernel{F, TT} end
 
     quote
         pointers = CLPtr[]
-        backend = Ref{Type{<:cl.CLBackend}}(cl.USM)
+        backend = Ref{Type{<:cl.CLBackend}}(cl.USMBackend)
         clcall(kernel.fun, $call_tt, $(call_args...); backend, pointers, call_kwargs...)
     end
 end

src/memory.jl

@@ -39,12 +39,12 @@ function maybe_synchronize(managed::Managed)
     end
 end
 
-function managed_buftype(x::Managed{M}) where M
+function managed_buftype(::Managed{M}) where M
     return M
 end
 
 function get_backend(x::Managed)
-    managed_buftype(x) == cl.SVMBuffer ? cl.SVM : cl.USM
+    cl.get_backend_from_buffer(managed_buftype(x))
 end
 
 function Base.convert(::Type{CLPtr{T}}, managed::Managed{M}) where {T, M}
@@ -84,7 +84,7 @@ end
 
 function Base.unsafe_copyto!(
         ::cl.Context, ::cl.Device, dst::Union{CLPtr{T}, Ptr{T}}, src::Union{CLPtr{T}, Ptr{T}}, N::Integer;
-        queu::cl.CmdQueue = cl.queue(), backend = cl.USM
+        queu::cl.CmdQueue = cl.queue(), backend = cl.USMBackend
     ) where {T}
     cl.enqueue_abstract_memcpy(dst, src, N * sizeof(T); queu = queu, backend = backend)
     cl.finish(queu)
@@ -93,7 +93,7 @@ end
 
 function unsafe_fill!(
         ctx::cl.Context, dev::cl.Device, ptr::Union{Ptr{T}, CLPtr{T}},
-        pattern::Union{Ptr{T}, CLPtr{T}}, N::Integer; queu::cl.CmdQueue = cl.queue(), backend::Type{<:cl.CLBackend} = cl.USM
+        pattern::Union{Ptr{T}, CLPtr{T}}, N::Integer; queu::cl.CmdQueue = cl.queue(), backend::Type{<:cl.CLBackend} = cl.USMBackend
     ) where {T}
     pattern_bytes = N * sizeof(T)
     pattern_bytes == 0 && return

test/buffer.jl

@@ -97,7 +97,7 @@ end
     end
 
     # memory map
-    #=
+
     let buf = cl.svm_alloc(cl.context(), sizeof(Int))
         ptr = pointer(buf)
 
@@ -106,7 +106,7 @@ end
 
         evt = cl.enqueue_svm_map(ptr, sizeof(src), :rw)
         wait(evt)
-        mapped = unsafe_wrap(Array, ptr, 1; own=false)
+        mapped = unsafe_wrap(Array, Ptr{Int}(UInt(ptr)), 1; own=false)
         @test mapped[] == 42
         mapped[] = 100
         cl.enqueue_svm_unmap(ptr) |> cl.wait
@@ -115,7 +115,6 @@ end
         cl.enqueue_svm_memcpy(pointer(dst), ptr, sizeof(dst); blocking=true)
         @test dst == [100]
     end
-    =#
 
     # fill
     let buf = cl.svm_alloc(cl.context(), 3 * sizeof(Int))

test/kernel.jl

@@ -42,7 +42,6 @@
         end
     end
 
-    #=
     @testset "set_arg!/set_args!" begin
         prg = cl.Program(source=test_source) |> cl.build!
         k = cl.Kernel(prg, "sum")
@@ -55,12 +54,13 @@
         A = CLArray(h_ones)
         B = CLArray(h_ones)
         C = CLArray{Float32}(undef, count)
-
+        backend = cl.USMBackend
+
         # we use julia's index by one convention
-        @test cl.set_arg!(k, 1, buffer(A)) != nothing
-        @test cl.set_arg!(k, 2, buffer(B)) != nothing
-        @test cl.set_arg!(k, 3, buffer(C)) != nothing
-        @test cl.set_arg!(k, 4, UInt32(count)) != nothing
+        @test cl.set_arg!(k, 1, backend, A.data[].mem) != nothing
+        @test cl.set_arg!(k, 2, backend, B.data[].mem) != nothing
+        @test cl.set_arg!(k, 3, backend, C.data[].mem) != nothing
+        @test cl.set_arg!(k, 4, backend, UInt32(count)) != nothing
 
         cl.enqueue_kernel(k, count) |> wait
         r = Array(C)
@@ -70,7 +70,7 @@
 
         # test set_args with new kernel
         k2 = cl.Kernel(prg, "sum")
-        cl.set_args!(k2, buffer(A), buffer(B), buffer(C), UInt32(count))
+        cl.set_args!(k2, backend, A.data[].mem, B.data[].mem, C.data[].mem, UInt32(count))
 
         h_twos = fill(2f0, count)
         copyto!(A, h_twos)
@@ -84,7 +84,7 @@
 
         @test all(x -> x == 4.0, Array(C))
     end
-    =#
+
     @testset "enqueue_kernel" begin
         simple_kernel = "
             __kernel void test(__global float *i) {