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 1/3] 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 2/3] 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 3/3] 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)