From 9c1b44adf4059f6f71146b1dc7ed32ac67697ecd Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Wed, 30 Sep 2026 18:04:22 +0200 Subject: [PATCH 01/10] Pass kernel arguments on as tuples MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Each layer of the launch path took the kernel arguments as varargs and splatted them into the next: the kernel object's call, `launch_with_exception_mailbox`, `cl.call` and `cl.set_args!`. Julia doesn't turn a splat of more than 32 elements into a direct call, and a method with both varargs and keyword arguments splats them into its body, so launching a kernel with many arguments went through `Core._apply_iterate` several times. Pass the arguments along as one tuple instead, and generate the per-argument code. The kernel object keeps its call syntax, with its keyword method defined explicitly, and `OpenCL.launch_tuple` takes the tuple directly. `cl.call` now takes the arguments as a tuple; `cl.set_args!` and `clcall` keep their varargs signatures. On PoCL, launching a kernel with 40 arguments with `@opencl` goes from 10.3 µs and 21.5 kB to 5.2 µs and 976 bytes, which is what a kernel with one argument allocates (it takes 4.3 µs). --- lib/cl/kernel.jl | 17 ++++++++++------- src/compiler/exceptions.jl | 7 +++++-- src/compiler/execution.jl | 21 +++++++++++++++++---- 3 files changed, 32 insertions(+), 13 deletions(-) diff --git a/lib/cl/kernel.jl b/lib/cl/kernel.jl index c914f0e2..04b4e585 100644 --- a/lib/cl/kernel.jl +++ b/lib/cl/kernel.jl @@ -260,14 +260,15 @@ function enqueue_task(k::Kernel; wait_for=nothing) return ret_event[] end +# the arguments are passed as a tuple, see `set_args!` function call( - k::Kernel, args...; global_size = (1,), local_size = nothing, + k::Kernel, args::Tuple; global_size = (1,), local_size = nothing, global_work_offset = nothing, wait_on::Vector{Event} = Event[], indirect_memory::Vector{AbstractMemory} = AbstractMemory[], rng_state=false, ) return Base.@lock k begin - set_args!(k, args...) + _set_args!(k, args) if !isempty(indirect_memory) svm_pointers = CLPtr{Cvoid}[] usm_pointers = CLPtr{Cvoid}[] @@ -360,7 +361,7 @@ clcall(f::F, types::Tuple, args::Vararg{Any,N}; kwargs...) where {N,F} = function clcall(k::Kernel, types::Type{T}, args::Vararg{Any,N}; kwargs...) where {T,N} call_closure = function (converted_args::Vararg{Any,N}) - call(k, converted_args...; kwargs...) + call(k, converted_args; kwargs...) end convert_arguments(call_closure, types, args...) end @@ -415,10 +416,12 @@ unlock_arguments(locks) = foreach(unlock, Iterators.reverse(locks)) return ex end -function set_args!(k::Kernel, args...) - for (i, a) in enumerate(args) - set_arg!(k, i, a) - end +set_args!(k::Kernel, args::Vararg{Any,N}) where {N} = _set_args!(k, args) + +# one call per argument: a splat of more than 32 arguments isn't a direct call +@inline @generated function _set_args!(k::Kernel, args::Tuple) + calls = (:(set_arg!(k, $i, args[$i])) for i in 1:fieldcount(args)) + return :($(calls...); nothing) end function set_arg!(k::Kernel, idx::Integer, arg::T) where {T} diff --git a/src/compiler/exceptions.jl b/src/compiler/exceptions.jl index ba3f4344..8092f668 100644 --- a/src/compiler/exceptions.jl +++ b/src/compiler/exceptions.jl @@ -169,11 +169,14 @@ function with_exception_mailbox(f, mailbox::ExceptionMailbox) end end +# `(x, t...)`, without splatting +@inline @generated prepend(x, t::Tuple) = :((x, $((:(t[$i]) for i in 1:fieldcount(t))...))) + # Launch while holding the mailbox lock, so a concurrent synchronization cannot finish the # queue between enqueue and recording it as pending (or map the mailbox during submission). # Keep this in an ordinary function: the generated `AbstractKernel` call cannot contain a # closure or `do` block on Julia 1.13. -function launch_with_exception_mailbox(kernel::cl.Kernel, args...; +function launch_with_exception_mailbox(kernel::cl.Kernel, args::Tuple; indirect_memory::Vector{cl.AbstractMemory}, rng_state::Bool, kwargs...) ctx, dev, queue = cl.context(), cl.device(), cl.queue() @@ -187,7 +190,7 @@ function launch_with_exception_mailbox(kernel::cl.Kernel, args...; launch_id = (Base.@atomic mailbox.launch_id) + UInt64(1) state = KernelState(rng_state ? Base.rand(UInt32) : UInt32(0), mailbox.address, launch_id) - result = cl.call(kernel, state, args...; indirect_memory, rng_state, kwargs...) + result = cl.call(kernel, prepend(state, args); indirect_memory, rng_state, kwargs...) # only publish the launch once it has been submitted, so that `check_exceptions` # can tell whether kernels have been submitted while it waited without the lock Base.@atomic mailbox.launch_id = launch_id diff --git a/src/compiler/execution.jl b/src/compiler/execution.jl index ac7d79a5..1b6e8255 100644 --- a/src/compiler/execution.jl +++ b/src/compiler/execution.jl @@ -179,11 +179,24 @@ abstract type AbstractKernel{F, TT} end pass_arg(@nospecialize dt) = !(isghosttype(dt) || Core.Compiler.isconstType(dt)) -@inline @generated function (kernel::AbstractKernel{F,TT})(args...; - call_kwargs...) where {F,TT} +# The arguments are passed on as a tuple: Julia doesn't turn a splat of more than 32 +# elements into a direct call, and a method with both varargs and keyword arguments splats +# them into its body. So the keyword method is defined explicitly. +(kernel::AbstractKernel)(args::Vararg{Any,N}) where {N} = launch_tuple(kernel, args) +Core.kwcall(kwargs::NamedTuple, kernel::AbstractKernel, args::Vararg{Any,N}) where {N} = + launch_tuple(kernel, args; kwargs...) + +""" + OpenCL.launch_tuple(kernel, args::Tuple; kwargs...) + +Launch `kernel`, as calling it with the arguments `args...` and the launch keywords +`kwargs` does, without splatting the arguments. +""" +@inline @generated function launch_tuple(kernel::AbstractKernel{F,TT}, args::Tuple; + call_kwargs...) where {F,TT} sig = Tuple{F, TT.parameters...} # Base.signature_type with a function type args = (:(kernel_convert(source, indirect_memory, managed)), - (:(kernel_convert(args[$i], indirect_memory, managed)) for i in 1:length(args))...) + (:(kernel_convert(args[$i], indirect_memory, managed)) for i in 1:fieldcount(args))...) # filter out ghost arguments that shouldn't be passed to_pass = map(pass_arg, sig.parameters) @@ -215,7 +228,7 @@ pass_arg(@nospecialize dt) = !(isghosttype(dt) || Core.Compiler.isconstType(dt)) locked = lock_managed(managed) try foreach(take_ownership!, locked) - launch_with_exception_mailbox(kernel.fun, $(converted...); + launch_with_exception_mailbox(kernel.fun, ($(converted...),); indirect_memory, rng_state=kernel.rng_state, call_kwargs...) finally From 30e87014a47813c8eb5ff8f18705a7a78c66ee78 Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Wed, 30 Sep 2026 18:34:00 +0200 Subject: [PATCH 02/10] Implement KernelInterface, and support KernelAbstractions 0.10 KernelAbstractions 0.10 builds on KernelInterface, which defines what a back end provides: memory and device management, compiling and launching kernels, and the device-side intrinsics. KernelAbstractions then launches `@kernel` kernels itself on any KernelInterface back end, which replaces OpenCL.jl's copy of that launch path: partitioning the ndrange, building the kernel's context, and tuning the workgroup size. `OpenCLBackend` now implements KernelInterface, and OpenCL.jl depends on it instead of on KernelAbstractions. What KernelAbstractions still needs from a back end moves to an extension: the `MArray` behind `@private`, and the Adapt rules for `@Const` and for moving arrays to the CPU. The platform is part of an `OpenCLBackend`'s configuration, but it was only checked, with a warning, before allocating and on every launch. Instead, work for a backend whose platform isn't active now activates the platform's default device, as `KI.device!` would, so that the arrays it creates can be used afterwards, and host queries answer for that device. `get_backend` returns a backend for the array's platform. Kernel queries answer for the device the kernel was compiled for, and launching a kernel on another device throws an error instead of failing in the driver. `KI.launch` passes the kernel arguments on as a tuple, so kernels with many arguments stay cheap to launch. It rejects `global_size` and `local_size`, which would override the launch geometry that KernelInterface validated, but passes OpenCL's other launch options on. The device properties that launches query are cached per device. `KI.kernel_function` compiles the converted callable, and keeps the original: it is converted again at every launch, like the arguments, as `@opencl` already does, so that the arrays a closure captures stay alive. `KI.synchronize` still blocks the thread while waiting for the device, rather than waiting cooperatively as KernelInterface asks. Co-authored-by: Christian Guinard <28689358+christiangnrd@users.noreply.github.com> --- Project.toml | 11 +- ext/KernelAbstractionsExt.jl | 25 ++++ src/OpenCL.jl | 2 +- src/OpenCLKernels.jl | 283 +++++++++++++++++++---------------- src/util.jl | 3 +- test/Project.toml | 1 + test/kernelabstractions.jl | 46 +++++- test/kernelinterface.jl | 102 +++++++++++++ 8 files changed, 338 insertions(+), 135 deletions(-) create mode 100644 ext/KernelAbstractionsExt.jl create mode 100644 test/kernelinterface.jl diff --git a/Project.toml b/Project.toml index 1db064ba..9923464f 100644 --- a/Project.toml +++ b/Project.toml @@ -10,7 +10,7 @@ Adapt = "79e6a3ab-5dfb-504d-930d-738a2a938a0e" 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" LinearAlgebra = "37e2e46d-f89d-539d-b4ee-838fcccc9c8e" OpenCL_jll = "6cb37087-e8b6-5417-8430-1f242f1e46e4" @@ -27,6 +27,12 @@ StaticArrays = "90137ffa-7385-5640-81b9-e52037218182" UnsafeAtomics = "013be700-e6cd-48c3-b4a1-df204f14c38f" spirv2clc_jll = "f0274c0c-8c8a-59f1-85b7-f7d60330c5fb" +[weakdeps] +KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c" + +[extensions] +KernelAbstractionsExt = "KernelAbstractions" + [sources] SPIRVIntrinsics = {path = "lib/intrinsics"} @@ -35,7 +41,8 @@ Adapt = "4" GPUArrays = "12.1" GPUCompiler = "2.13" GPUToolbox = "3.3.2" -KernelAbstractions = "0.9.38" +KernelAbstractions = "0.10" +KernelInterface = "0.4" LLVM = "10" LinearAlgebra = "1" OpenCL_jll = "=2024.10.24" diff --git a/ext/KernelAbstractionsExt.jl b/ext/KernelAbstractionsExt.jl new file mode 100644 index 00000000..fa406c1d --- /dev/null +++ b/ext/KernelAbstractionsExt.jl @@ -0,0 +1,25 @@ +module KernelAbstractionsExt + +using OpenCL +using OpenCL: @device_override, method_table + +import KernelAbstractions as KA + +import StaticArrays + +import Adapt + +Adapt.adapt_storage(::KA.CPU, a::CLArray) = convert(Array, a) + +# `@Const` applies `constify` inside the kernel, where arguments have already been +# converted to device arrays, so the rule has to be registered for `CLDeviceArray` +# rather than for `CLArray`. +Adapt.adapt_storage(::KA.ConstAdaptor, a::CLDeviceArray) = Base.Experimental.Const(a) + +## scratch memory + +@device_override @inline function KA.Scratchpad(ctx, ::Type{T}, ::Val{Dims}) where {T, Dims} + StaticArrays.MArray{Tuple{Dims...}, T}(undef) +end + +end diff --git a/src/OpenCL.jl b/src/OpenCL.jl index e0dae368..c7a49ce5 100644 --- a/src/OpenCL.jl +++ b/src/OpenCL.jl @@ -11,7 +11,7 @@ using GPUArrays using Random using Preferences -import KernelAbstractions: KernelAbstractions +import KernelInterface using Core: LLVMPtr diff --git a/src/OpenCLKernels.jl b/src/OpenCLKernels.jl index 98220c9f..c258d6ea 100644 --- a/src/OpenCLKernels.jl +++ b/src/OpenCLKernels.jl @@ -1,11 +1,9 @@ module OpenCLKernels using ..OpenCL -using ..OpenCL: @device_override, method_table +using ..OpenCL: @device_override, method_table, kernel_convert, clfunction -import KernelAbstractions as KA - -import StaticArrays +import KernelInterface as KI import Adapt @@ -14,17 +12,36 @@ import Adapt export OpenCLBackend -Base.@kwdef struct OpenCLBackend <: KA.GPU +""" + OpenCLBackend(; platform=cl.platform()) + +KernelInterface back end for the OpenCL devices of `platform`. + +A backend works with the task's active device if that is on its platform. Otherwise, work +for the backend (allocations, copies, compilation and launches) first activates the default +device of its platform, as `KernelInterface.device!` would, so that the arrays it creates +can be used afterwards. +""" +Base.@kwdef struct OpenCLBackend <: KI.Backend platform::cl.Platform = cl.platform() end -@noinline function platform_mismatch_warning(expected::cl.Platform, active::cl.Platform) - @warn "OpenCLBackend platform \"$(expected.name)\" is not the active platform \"$(active.name)\"" - return nothing +# the device that `b` works with +function backend_device(b::OpenCLBackend) + cl.platform() == b.platform && return cl.device() + dev = cl.default_device(b.platform) + dev === nothing && throw(ArgumentError("OpenCL platform \"$(b.platform.name)\" has no devices")) + return dev +end + +# make the backend's device the task's active device +@inline function activate(b::OpenCLBackend) + cl.platform() == b.platform || cl.platform!(b.platform) + return end -function KA.allocate(b::OpenCLBackend, ::Type{T}, dims::Tuple; unified::Bool = false) where T - b.platform === cl.platform() || platform_mismatch_warning(b.platform, cl.platform()) +function KI.allocate(b::OpenCLBackend, ::Type{T}, dims::Tuple; unified::Bool = false) where T + activate(b) if unified memory_backend = cl.unified_memory_backend() if memory_backend === cl.USMBackend() @@ -39,193 +56,199 @@ function KA.allocate(b::OpenCLBackend, ::Type{T}, dims::Tuple; unified::Bool = f end end -KA.supports_unified(::OpenCLBackend) = cl.default_memory_backend(cl.device(); unified=true) !== nothing +# OpenCL.jl creates a context per device +context_device(ctx::cl.Context) = ctx == cl.context() ? cl.device() : only(ctx.devices) -KA.get_backend(::CLArray) = OpenCLBackend() -KA.synchronize(::OpenCLBackend) = OpenCL.synchronize() -KA.supports_float64(::OpenCLBackend) = in("cl_khr_fp64", cl.device().extensions) +function KI.get_backend(A::CLArray) + ctx = OpenCL.context(A) + ctx == cl.context() && return OpenCLBackend(cl.platform()) + return OpenCLBackend(context_device(ctx).platform) +end + +function KI.synchronize(b::OpenCLBackend) + activate(b) + OpenCL.synchronize() + return +end -Adapt.adapt_storage(::OpenCLBackend, a::Array) = Adapt.adapt(CLArray, a) +function Adapt.adapt_storage(b::OpenCLBackend, a::Array) + activate(b) + return Adapt.adapt(CLArray, a) +end Adapt.adapt_storage(::OpenCLBackend, a::CLArray) = a -Adapt.adapt_storage(::KA.CPU, a::CLArray) = convert(Array, a) -# `@Const` applies `constify` inside the kernel, where arguments have already been -# converted to device arrays, so the rule has to be registered for `CLDeviceArray` -# rather than for `CLArray`. -Adapt.adapt_storage(::KA.ConstAdaptor, a::CLDeviceArray) = Base.Experimental.Const(a) ## Device Selection # devices are numbered consecutively within the backend's platform, in enumeration order -function KA.ndevices(b::OpenCLBackend) +function KI.ndevices(b::OpenCLBackend) Int(cl.ndevices(b.platform)) end -function KA.device(b::OpenCLBackend) - current = cl.device() - for (i, d) in enumerate(cl.devices(b.platform)) - d == current && return i - end - error("Active OpenCL device $current not found in the OpenCLBackend's platform \"$(b.platform.name)\".") +function device_index(b::OpenCLBackend, dev::cl.Device) + id = findfirst(==(dev), cl.devices(b.platform)) + id === nothing && + throw(ArgumentError("OpenCL device $(dev.name) is not on the backend's platform \"$(b.platform.name)\"")) + return id end -function KA.device!(b::OpenCLBackend, id::Int) - 0 < id <= KA.ndevices(b) || throw(ArgumentError("Device id $id out of bounds.")) +KI.device(b::OpenCLBackend) = device_index(b, backend_device(b)) + +KI.device(b::OpenCLBackend, A::CLArray) = device_index(b, context_device(OpenCL.context(A))) + +function KI.device!(b::OpenCLBackend, id::Int) + 0 < id <= KI.ndevices(b) || throw(ArgumentError("Device id $id out of bounds.")) devs = cl.devices(b.platform) cl.device!(devs[id]) return nothing end + ## Memory Operations -function KA.copyto!(::OpenCLBackend, A, B) +function KI.copyto!(b::OpenCLBackend, A, B) + length(A) == length(B) || + throw(ArgumentError("Arrays must have the same length, got $(length(A)) and $(length(B))")) + activate(b) copyto!(A, B) - # TODO: Address device to host copies in jl being synchronizing + return A end +KI.unsafe_free!(A::CLArray) = OpenCL.unsafe_free!(A) -## Kernel Launch - -function KA.mkcontext(kernel::KA.Kernel{OpenCLBackend}, _ndrange, iterspace) - KA.CompilerMetadata{KA.ndrange(kernel), KA.DynamicCheck}(_ndrange, iterspace) -end -function KA.mkcontext(kernel::KA.Kernel{OpenCLBackend}, I, _ndrange, iterspace, - ::Dynamic) where Dynamic - KA.CompilerMetadata{KA.ndrange(kernel), Dynamic}(I, _ndrange, iterspace) -end -function KA.launch_config(kernel::KA.Kernel{OpenCLBackend}, 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 +## Kernel Launch - 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 +KI.argconvert(::OpenCLBackend, arg) = kernel_convert(arg) - return ndrange, workgroupsize, iterspace, dynamic +# `f` is the host-side callable: the kernel keeps it as its `source`, which is converted +# again at every launch, because the converted callable only holds pointers to the arrays +# it captures +function KI.kernel_function(backend::OpenCLBackend, f::F, tt::TT=Tuple{}; name = nothing, kwargs...) where {F,TT} + activate(backend) + kern = GC.@preserve f clfunction(kernel_convert(f), tt; source=f, name, kwargs...) + KI.Kernel{OpenCLBackend, typeof(kern)}(backend, kern) 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 +# the context that a kernel was compiled for +function kernel_context(kernel::KI.Kernel{OpenCLBackend}) + ctx = Ref{cl.cl_context}() + cl.clGetKernelInfo(kernel.kern.fun, cl.CL_KERNEL_CONTEXT, sizeof(cl.cl_context), ctx, C_NULL) + return ctx[] end -function (obj::KA.Kernel{OpenCLBackend})(args...; ndrange=nothing, workgroupsize=nothing) - obj.backend.platform === cl.platform() || platform_mismatch_warning(obj.backend.platform, cl.platform()) +function kernel_device(kernel::KI.Kernel{OpenCLBackend}) + ctx = kernel_context(kernel) + ctx == cl.context().id && return cl.device() + return context_device(cl.Context(ctx; retain=true)) +end - ndrange, workgroupsize, iterspace, dynamic = - KA.launch_config(obj, ndrange, workgroupsize) +@noinline function throw_device_mismatch(kernel) + throw(ArgumentError("Cannot launch a kernel compiled for $(kernel_device(kernel).name) on $(cl.device().name)")) +end - # this might not be the final context, since we may tune the workgroupsize - ctx = KA.mkcontext(obj, ndrange, iterspace) - kernel = @opencl launch=false obj.f(ctx, args...) +@noinline function throw_geometry_keyword() + throw(ArgumentError("KernelInterface kernels take `numgroups`, `workgroupsize` or `ndrange`, not `global_size` or `local_size`")) +end - # figure out the optimal workgroupsize automatically - if KA.workgroupsize(obj) <: KA.DynamicSize && workgroupsize === nothing - wg_info = cl.work_group_info(kernel.fun, cl.device()) - wg_size_nd = threads_to_workgroupsize(wg_info.size, ndrange) - iterspace, dynamic = KA.partition(obj, ndrange, wg_size_nd) - ctx = KA.mkcontext(obj, ndrange, iterspace) +# passes the arguments on as a tuple, like calling the `HostKernel` does +function KI.launch(kernel::KI.Kernel{OpenCLBackend}, groups::Dims{3}, items::Dims{3}, + args::Tuple; kwargs...) + # KernelInterface has validated the launch geometry + if haskey(kwargs, :global_size) || haskey(kwargs, :local_size) + throw_geometry_keyword() end - - groups = length(KA.blocks(iterspace)) - items = length(KA.workitems(iterspace)) - - if groups == 0 - return nothing + activate(kernel.backend) + kernel_context(kernel) == cl.context().id || throw_device_mismatch(kernel) + OpenCL.launch_tuple(kernel.kern, args; local_size = items, global_size = items .* groups, + kwargs...) + return +end + +function KI.max_work_group_size(kernel::KI.Kernel{OpenCLBackend})::Int + wginfo = cl.work_group_info(kernel.kern.fun, kernel_device(kernel)) + Int(wginfo.size) +end + + +## Device Properties + +# querying the device allocates, so cache what launches and kernels need. the cache is +# keyed on the device, because the task-local device can be switched. +const DeviceProperties = @NamedTuple{ + max_work_group_size::Int, max_work_group_dims::NTuple{3, Int}, compute_units::Int, + float64::Bool, unified::Bool, +} +function device_properties(dev::cl.Device) + cache = get!(task_local_storage(), :CLDeviceProperties) do + Dict{cl.Device, DeviceProperties}() + end::Dict{cl.Device, DeviceProperties} + return get!(cache, dev) do + sizes = dev.max_work_item_size + (; max_work_group_size = Int(dev.max_work_group_size), + max_work_group_dims = ntuple(d -> d <= length(sizes) ? Int(sizes[d]) : 1, 3), + compute_units = Int(dev.max_compute_units), + float64 = "cl_khr_fp64" in dev.extensions, + unified = cl.default_memory_backend(dev; unified=true) !== nothing) end +end - # Launch kernel - global_size = groups * items - local_size = items - kernel(ctx, args...; global_size, local_size) +device_properties(b::OpenCLBackend) = device_properties(backend_device(b)) - return nothing +KI.max_work_group_size(b::OpenCLBackend)::Int = device_properties(b).max_work_group_size +KI.max_work_group_dims(b::OpenCLBackend)::NTuple{3, Int} = device_properties(b).max_work_group_dims +# OpenCL doesn't limit the number of work-groups, only the global size (to `size_t`) +function KI.max_num_groups(b::OpenCLBackend)::NTuple{3, Int} + return typemax(Int) .÷ KI.max_work_group_dims(b) end +KI.multiprocessor_count(b::OpenCLBackend)::Int = device_properties(b).compute_units +KI.supports_float64(b::OpenCLBackend) = device_properties(b).float64 +KI.supports_unified(b::OpenCLBackend) = device_properties(b).unified +# 32-bit integer atomics are core OpenCL; float atomics fall back to compare-and-swap +KI.supports_atomics(::OpenCLBackend) = true -## 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 +## Indexing Functions -@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 +# computed with `% T`, which unlike `T(x)` has no error path. KernelInterface derives the +# global queries from these. -@device_override @inline function KA.__index_Local_Cartesian(ctx) - @inbounds KA.workitems(KA.__iterspace(ctx))[get_local_id(1)] +@device_override @inline function KI.get_local_id(::Type{T}) where {T} + return (; x = get_local_id(1) % T, y = get_local_id(2) % T, z = get_local_id(3) % T) end -@device_override @inline function KA.__index_Group_Cartesian(ctx) - @inbounds KA.blocks(KA.__iterspace(ctx))[get_group_id(1)] +@device_override @inline function KI.get_group_id(::Type{T}) where {T} + return (; x = get_group_id(1) % T, y = get_group_id(2) % T, z = get_group_id(3) % T) 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)) +@device_override @inline function KI.get_local_size(::Type{T}) where {T} + return (; x = get_local_size(1) % T, y = get_local_size(2) % T, z = get_local_size(3) % T) end -@device_override @inline function KA.__validindex(ctx) - if KA.__dynamic_checkbounds(ctx) - I = KA.__index_Global_Cartesian(ctx) - return I in KA.__ndrange(ctx) - else - return true - end +@device_override @inline function KI.get_num_groups(::Type{T}) where {T} + return (; x = get_num_groups(1) % T, y = get_num_groups(2) % T, z = get_num_groups(3) % T) end -## Shared and Scratch Memory +## Shared 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} 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() +@device_override @inline function KI.barrier() work_group_barrier(OpenCL.LOCAL_MEM_FENCE | OpenCL.GLOBAL_MEM_FENCE) end -@device_override @inline function KA.__print(args...) +@device_override @inline function KI._print(args...) OpenCL._print(args...) end - -## Other - -KA.argconvert(::KA.Kernel{OpenCLBackend}, arg) = OpenCL.kernel_convert(arg) - end diff --git a/src/util.jl b/src/util.jl index a1900e4f..57ce840c 100644 --- a/src/util.jl +++ b/src/util.jl @@ -72,7 +72,8 @@ function versioninfo(io::IO=stdout) (pkg[2], get(Base.loaded_modules, id, nothing)) end - for pkg in [:GPUArrays, :GPUCompiler, ("63c18a36-062a-441e-b654-da1e3ab1ce7c", "KernelAbstractions"), + for pkg in [:GPUArrays, :GPUCompiler, :KernelInterface, + ("63c18a36-062a-441e-b654-da1e3ab1ce7c", "KernelAbstractions"), :LLVM, :SPIRVIntrinsics, ("627d6b7a-bbe6-5189-83e7-98cc0a5aeadd", "pocl_jll"), ("59abdad9-3cfc-5436-8271-411e8cad6b82", "pocl_next_jll"), ("6e3bd4d7-b6fc-52f9-a000-aa432bc9a0d7", "intel_opencl_rt_jll")] diff --git a/test/Project.toml b/test/Project.toml index 35daa4a5..61736319 100644 --- a/test/Project.toml +++ b/test/Project.toml @@ -9,6 +9,7 @@ IOCapture = "b5f81e59-6552-4d32-b1f0-c071b021bf89" InteractiveUtils = "b77e0a4c-d291-57a0-90e8-8db25a27a240" JLD2 = "033835bb-8acc-5ee8-8aae-3f567f8a3819" KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c" +KernelInterface = "4ee993da-d684-4d17-a7dd-4e58e78d92bf" LinearAlgebra = "37e2e46d-f89d-539d-b4ee-838fcccc9c8e" OpenCL = "08131aa3-fb12-5dee-8b74-c09406e224a2" ParallelTestRunner = "d3525ed8-44d0-4b2c-a655-542cee43accc" diff --git a/test/kernelabstractions.jl b/test/kernelabstractions.jl index 4f2d3de7..db2594ee 100644 --- a/test/kernelabstractions.jl +++ b/test/kernelabstractions.jl @@ -64,8 +64,52 @@ const KATestSuite = let mod.Testsuite end +# the others run kernels on KernelAbstractions' PoCL-based CPU back 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) + +@kernel function store_global_linear!(A) + I = @index(Global, Linear) + @inbounds A[I] = I +end + +@kernel function store_last_index!(A) + I = @index(Global, Linear) + if I == prod(@ndrange()) + @inbounds A[1] = I + @inbounds A[2] = @index(Global, Cartesian)[2] + end +end + +@testset "launch configuration" begin + backend = OpenCLBackend() + function select(kernel, ndrange, workgroupsize=nothing) + ndrange, workgroupsize, iterspace, _ = KernelAbstractions.launch_config(kernel, ndrange, workgroupsize) + KernelAbstractions.select_launch(kernel, workgroupsize, iterspace) + end + + # kernels are launched on an N-d grid, computing indices in 32 bits + kernel = store_global_linear!(backend) + @test select(kernel, (64, 32, 16)) === KernelAbstractions.NDLaunch{Int32}() + @test select(kernel, (4, 4, 4, 4)) === KernelAbstractions.LinearLaunch{Int32}() + + # which doesn't need divisions to compute the index of a dynamic N-d range + A = OpenCL.zeros(Int, 64, 32, 16) + ir = sprint(io -> @device_code_llvm io=io kernel(A; ndrange=size(A))) + @test !occursin(r"\b[su](div|rem) ", ir) + @test Array(A) == LinearIndices(A) + + # iteration spaces that don't fit 32 bits use 64-bit indices + kernel = store_last_index!(backend) + A = OpenCL.zeros(Int, 2) + for (dims, launch) in (((2^16 + 1, 2^15), KernelAbstractions.NDLaunch{Int}()), + ((2^11 + 1, 2^10, 2^10, 1), KernelAbstractions.LinearLaunch{Int}())) + @test select(kernel, dims) === launch + kernel(A; ndrange=dims) + @test Array(A) == [prod(dims), dims[2]] + end +end diff --git a/test/kernelinterface.jl b/test/kernelinterface.jl new file mode 100644 index 00000000..ee38ec2a --- /dev/null +++ b/test/kernelinterface.jl @@ -0,0 +1,102 @@ +import KernelInterface +import KernelInterface as KI +import Adapt + +include(joinpath(dirname(pathof(KernelInterface)), "..", "test", "testsuite.jl")) + +Testsuite.testsuite(OpenCLBackend(), CLArray) + +function ki_fill_kernel(a, val) + i = KI.get_global_id().x + if i <= length(a) + @inbounds a[i] = val + end + return +end + +@testset "backend platform" begin + backend = OpenCLBackend() + a = KI.zeros(backend, Int32, 4) + @test KI.get_backend(a) == backend + + # a kernel belongs to the device it was compiled for + kernel = KI.@launch backend launch=false ki_fill_kernel(a, Int32(1)) + cl.context!(cl.Context(cl.device())) do + @test_throws ArgumentError kernel(a, Int32(1); ndrange = 4) + end + kernel(a, Int32(1); ndrange = 4) + @test Array(a) == ones(Int32, 4) + + # a backend for another platform activates it + others = filter(!=(cl.platform()), cl.platforms()) + if !isempty(others) + platform, device = cl.platform(), cl.device() + try + other = OpenCLBackend(; platform = first(others)) + @test KI.device(other) == 1 + b = KI.zeros(other, Int32, 4) + @test cl.platform() == first(others) + @test KI.get_backend(b) == other + @test KI.device(other, b) == KI.device(other) + + # so does moving arrays to it + cl.platform!(platform) + c = Adapt.adapt(other, Int32[1, 2]) + @test cl.platform() == first(others) + @test KI.get_backend(c) == other + KI.@launch other ndrange = 4 ki_fill_kernel(b, Int32(2)) + @test Array(b) == fill(Int32(2), 4) + + # arrays keep their platform + cl.device!(device) + @test KI.get_backend(b) == other + @test KI.get_backend(a) == backend + finally + cl.platform!(platform) + cl.device!(device) + end + end +end + +struct KICapturedArray{A} + array::A +end +Adapt.@adapt_structure KICapturedArray +function (f::KICapturedArray)() + @inbounds f.array[1] = 7 + return +end + +@testset "captured arrays" begin + function captured_kernel() + array = CLArray(Int32[0]) + kernel = KI.@launch OpenCLBackend() launch=false KICapturedArray(array)() + return kernel, WeakRef(array) + end + + # the kernel keeps the callable alive, and converts it again at launch, which makes the + # launch's queue the owner of the arrays it captures + kernel, owner = captured_kernel() + GC.gc(true) + @test owner.value !== nothing + queue = cl.CmdQueue() + cl.queue!(queue) do + kernel() + @test owner.value.data[].queue === queue + cl.finish(queue) + end + @test Array(owner.value) == Int32[7] +end + +@testset "launch keywords" begin + a = KI.zeros(OpenCLBackend(), Int32, 4) + kernel = KI.@launch OpenCLBackend() launch=false ki_fill_kernel(a, Int32(1)) + + # OpenCL's launch options are passed on + kernel(a, Int32(3); ndrange = 4, wait_on = cl.Event[]) + @test Array(a) == fill(Int32(3), 4) + + # but not ones that would override the launch geometry + @test_throws ArgumentError kernel(a, Int32(1); ndrange = 4, global_size = 8) + @test_throws ArgumentError kernel(a, Int32(1); ndrange = 4, local_size = 2) +end From 7ad7fb2df3e251012a47ddacf42250fced9a578f Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Wed, 30 Sep 2026 18:38:21 +0200 Subject: [PATCH 03/10] Support KernelInterface's sub-groups Implement the sub-group queries, `sub_group_barrier` and `shfl_down` with OpenCL's sub-group builtins, whose semantics already match KernelInterface's: the last sub-group of a work-group can be partial, and `get_sub_group_size` counts the work-items that are present. KernelInterface requires kernels to execute with the sub-group width that `KI.sub_group_size` reports. `clfunction` requests that width on devices with `cl_intel_required_subgroup_size`, so only those report sub-group support, and shuffles for the types that `cl_khr_subgroup_shuffle` covers on the device. A `sub_group_size` compiler option asking for another width is rejected. How work-items form sub-groups differs between devices: Intel's CPU runtime forms them per row of a multi-dimensional work-group. KernelInterface doesn't promise more than that (JuliaGPU/KernelAbstractions.jl#815). Co-authored-by: Christian Guinard <28689358+christiangnrd@users.noreply.github.com> --- src/OpenCLKernels.jl | 46 ++++++++++++++++++++++++++++++++++++++++- test/kernelinterface.jl | 22 ++++++++++++++++++++ 2 files changed, 67 insertions(+), 1 deletion(-) diff --git a/src/OpenCLKernels.jl b/src/OpenCLKernels.jl index c258d6ea..6aedea14 100644 --- a/src/OpenCLKernels.jl +++ b/src/OpenCLKernels.jl @@ -128,10 +128,20 @@ KI.argconvert(::OpenCLBackend, arg) = kernel_convert(arg) # it captures function KI.kernel_function(backend::OpenCLBackend, f::F, tt::TT=Tuple{}; name = nothing, kwargs...) where {F,TT} activate(backend) + check_sub_group_size(backend, kwargs) kern = GC.@preserve f clfunction(kernel_convert(f), tt; source=f, name, kwargs...) KI.Kernel{OpenCLBackend, typeof(kern)}(backend, kern) end +# kernels have to execute with the sub-group width that `KI.sub_group_size` reports +function check_sub_group_size(backend::OpenCLBackend, kwargs) + haskey(kwargs, :sub_group_size) && KI.supports_subgroups(backend) || return + width = KI.sub_group_size(backend) + kwargs[:sub_group_size] == width || + throw(ArgumentError("KernelInterface kernels execute with sub-group width $width, got `sub_group_size=$(kwargs[:sub_group_size])`")) + return +end + # the context that a kernel was compiled for function kernel_context(kernel::KI.Kernel{OpenCLBackend}) ctx = Ref{cl.cl_context}() @@ -180,6 +190,8 @@ end const DeviceProperties = @NamedTuple{ max_work_group_size::Int, max_work_group_dims::NTuple{3, Int}, compute_units::Int, float64::Bool, unified::Bool, + # 0 if the device doesn't support sub-groups of a fixed width + sub_group_size::Int, shuffle_types::Vector{DataType}, } function device_properties(dev::cl.Device) cache = get!(task_local_storage(), :CLDeviceProperties) do @@ -187,11 +199,17 @@ function device_properties(dev::cl.Device) end::Dict{cl.Device, DeviceProperties} return get!(cache, dev) do sizes = dev.max_work_item_size + # the sub-group width is only fixed for kernels that request it, which `clfunction` + # does for devices with `cl_intel_required_subgroup_size` + fixed_sub_groups = cl.sub_groups_supported(dev) && + "cl_intel_required_subgroup_size" in dev.extensions (; max_work_group_size = Int(dev.max_work_group_size), max_work_group_dims = ntuple(d -> d <= length(sizes) ? Int(sizes[d]) : 1, 3), compute_units = Int(dev.max_compute_units), float64 = "cl_khr_fp64" in dev.extensions, - unified = cl.default_memory_backend(dev; unified=true) !== nothing) + unified = cl.default_memory_backend(dev; unified=true) !== nothing, + sub_group_size = fixed_sub_groups ? cl.sub_group_size(dev) : 0, + shuffle_types = fixed_sub_groups ? cl.sub_group_shuffle_supported_types(dev) : DataType[]) end end @@ -210,6 +228,10 @@ KI.supports_unified(b::OpenCLBackend) = device_properties(b).unified # 32-bit integer atomics are core OpenCL; float atomics fall back to compare-and-swap KI.supports_atomics(::OpenCLBackend) = true +KI.supports_subgroups(b::OpenCLBackend) = device_properties(b).sub_group_size > 0 +KI.sub_group_size(b::OpenCLBackend)::Int = device_properties(b).sub_group_size +KI.supports_shuffle(b::OpenCLBackend, ::Type{T}) where {T} = T in device_properties(b).shuffle_types + ## Indexing Functions @@ -232,6 +254,19 @@ end return (; x = get_num_groups(1) % T, y = get_num_groups(2) % T, z = get_num_groups(3) % T) end +# OpenCL's sub-group queries already have KernelInterface's semantics: the last sub-group +# of a work-group can be partial, and `get_sub_group_size` counts the work-items present + +@device_override KI.get_sub_group_size(::Type{T}) where {T} = get_sub_group_size() % T + +@device_override KI.get_max_sub_group_size(::Type{T}) where {T} = get_max_sub_group_size() % T + +@device_override KI.get_num_sub_groups(::Type{T}) where {T} = get_num_sub_groups() % T + +@device_override KI.get_sub_group_id(::Type{T}) where {T} = get_sub_group_id() % T + +@device_override KI.get_sub_group_local_id(::Type{T}) where {T} = get_sub_group_local_id() % T + ## Shared Memory @@ -247,6 +282,15 @@ end work_group_barrier(OpenCL.LOCAL_MEM_FENCE | OpenCL.GLOBAL_MEM_FENCE) end +@device_override @inline function KI.sub_group_barrier() + sub_group_barrier(OpenCL.LOCAL_MEM_FENCE | OpenCL.GLOBAL_MEM_FENCE) +end + +# out-of-range source lanes give an unspecified value, as KernelInterface allows +@device_override function KI.shfl_down(val::T, offset::Integer) where T + sub_group_shuffle(val, get_sub_group_local_id() % UInt32 + offset % UInt32) +end + @device_override @inline function KI._print(args...) OpenCL._print(args...) end diff --git a/test/kernelinterface.jl b/test/kernelinterface.jl index ee38ec2a..c883943c 100644 --- a/test/kernelinterface.jl +++ b/test/kernelinterface.jl @@ -100,3 +100,25 @@ end @test_throws ArgumentError kernel(a, Int32(1); ndrange = 4, global_size = 8) @test_throws ArgumentError kernel(a, Int32(1); ndrange = 4, local_size = 2) end + +@testset "sub-groups" begin + backend = OpenCLBackend() + dev = cl.device() + # kernels only execute with a fixed sub-group width if they can request one + if cl.sub_groups_supported(dev) && "cl_intel_required_subgroup_size" in dev.extensions + @test KI.supports_subgroups(backend) + @test KI.sub_group_size(backend) == cl.sub_group_size(dev) + @test KI.supports_shuffle(backend, Int32) == + ("cl_khr_subgroup_shuffle" in dev.extensions) + + # which kernels can't opt out of + a = KI.zeros(backend, Int32, 4) + width = KI.sub_group_size(backend) + KI.@launch backend launch=false sub_group_size=width ki_fill_kernel(a, Int32(1)) + @test_throws ArgumentError KI.@launch backend launch=false sub_group_size=nothing ki_fill_kernel(a, Int32(1)) + @test_throws ArgumentError KI.@launch backend launch=false sub_group_size=2width ki_fill_kernel(a, Int32(1)) + else + @test !KI.supports_subgroups(backend) + end + @test !KI.supports_shuffle(backend, Complex{Float32}) +end From a41d1863b6c8eaa45d79d4825eb3e28ab83e09b6 Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Wed, 30 Sep 2026 18:44:12 +0200 Subject: [PATCH 04/10] Order work across tasks with OpenCL events Implement `KI.record_event` with a marker on the task's queue, and `KI.wait_event` with a barrier on the waiting task's queue. `KernelAbstractions.@spawn` uses them to order a new task's work after the work its parent had queued, without synchronizing the parent; the default is a full synchronization. A queue can only wait for events of its own context, and OpenCL.jl creates a context per device, so waiting for an event of another device waits on the host. --- src/OpenCLKernels.jl | 31 +++++++++++++++++++++++++++++++ test/kernelinterface.jl | 15 +++++++++++++++ 2 files changed, 46 insertions(+) diff --git a/src/OpenCLKernels.jl b/src/OpenCLKernels.jl index 6aedea14..acc6d438 100644 --- a/src/OpenCLKernels.jl +++ b/src/OpenCLKernels.jl @@ -71,6 +71,37 @@ function KI.synchronize(b::OpenCLBackend) return end +# queues are task-local, so work is ordered across tasks with a marker event on the +# recording task's queue +function KI.record_event(b::OpenCLBackend) + activate(b) + event = cl.enqueue_marker_with_wait_list(cl.AbstractEvent[]) + # the waiting queue only makes progress if this one is submitted + cl.flush(cl.queue()) + return event +end + +function event_context(event::cl.Event) + ctx = Ref{cl.cl_context}() + cl.clGetEventInfo(event, cl.CL_EVENT_CONTEXT, sizeof(cl.cl_context), ctx, C_NULL) + return ctx[] +end + +function KI.wait_event(b::OpenCLBackend, event::cl.Event) + activate(b) + # the event has to stay alive until the driver has retained it + GC.@preserve event begin + if event_context(event) == cl.context().id + cl.enqueue_barrier_with_wait_list(cl.AbstractEvent[event]) + else + # XXX: queues can only wait for events of their own context, and OpenCL.jl + # creates a context per device, so wait for other devices on the host. + wait(event) + end + end + return +end + function Adapt.adapt_storage(b::OpenCLBackend, a::Array) activate(b) return Adapt.adapt(CLArray, a) diff --git a/test/kernelinterface.jl b/test/kernelinterface.jl index c883943c..e09cbf0a 100644 --- a/test/kernelinterface.jl +++ b/test/kernelinterface.jl @@ -122,3 +122,18 @@ end end @test !KI.supports_shuffle(backend, Complex{Float32}) end + +@testset "events" begin + backend = OpenCLBackend() + a = KI.zeros(backend, Int32, 4) + KI.@launch backend ndrange = 4 ki_fill_kernel(a, Int32(5)) + event = KI.record_event(backend) + @test event isa cl.Event + + # queues of another context can't wait for the event, so the host does + cl.context!(cl.Context(cl.device())) do + @test KI.wait_event(backend, event) === nothing + @test event.status == :complete + end + @test Array(a) == fill(Int32(5), 4) +end From 92b3267c0a5b51caa4b5cf24d23ac670bf808742 Mon Sep 17 00:00:00 2001 From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com> Date: Wed, 30 Sep 2026 18:44:41 +0200 Subject: [PATCH 05/10] Report OpenCL's version info through KernelInterface `KI.versioninfo(OpenCLBackend())` prints `OpenCL.versioninfo()`. --- src/OpenCLKernels.jl | 2 ++ test/kernelinterface.jl | 4 ++++ 2 files changed, 6 insertions(+) diff --git a/src/OpenCLKernels.jl b/src/OpenCLKernels.jl index acc6d438..48e5e987 100644 --- a/src/OpenCLKernels.jl +++ b/src/OpenCLKernels.jl @@ -26,6 +26,8 @@ Base.@kwdef struct OpenCLBackend <: KI.Backend platform::cl.Platform = cl.platform() end +KI.versioninfo(io::IO, ::OpenCLBackend) = OpenCL.versioninfo(io) + # the device that `b` works with function backend_device(b::OpenCLBackend) cl.platform() == b.platform && return cl.device() diff --git a/test/kernelinterface.jl b/test/kernelinterface.jl index e09cbf0a..8242cff2 100644 --- a/test/kernelinterface.jl +++ b/test/kernelinterface.jl @@ -137,3 +137,7 @@ end end @test Array(a) == fill(Int32(5), 4) end + +@testset "versioninfo" begin + @test occursin("OpenCL.jl version", sprint(KI.versioninfo, OpenCLBackend())) +end From 440d07b447fa98da1805f0299abe369afc65c7ce Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Fri, 2 Oct 2026 15:23:38 +0200 Subject: [PATCH 06/10] Include the KernelInterface back-end in the API reference Documenter failed on the OpenCLBackend docstring missing from the manual. --- docs/src/95-reference.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs/src/95-reference.md b/docs/src/95-reference.md index e88779f0..f4da4eb1 100644 --- a/docs/src/95-reference.md +++ b/docs/src/95-reference.md @@ -13,5 +13,5 @@ Pages = ["95-reference.md"] ``` ```@autodocs -Modules = [OpenCL, OpenCL.cl] +Modules = [OpenCL, OpenCL.cl, OpenCL.OpenCLKernels] ``` From dc5354ec589571a95f61b21855fb88c9a42a6341 Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Tue, 6 Oct 2026 11:18:06 +0200 Subject: [PATCH 07/10] Use KernelAbstractions' @private storage KernelAbstractions now backs @private with a stack allocation through an overlay in GPUCompiler's shared method table, so drop our MArray-based Scratchpad override along with the StaticArrays dependency. This relies on OpenCL consulting the shared table, which OpenCL.jl#530 does by declaring our method tables through method_tables. --- Project.toml | 2 -- ext/KernelAbstractionsExt.jl | 9 --------- 2 files changed, 11 deletions(-) diff --git a/Project.toml b/Project.toml index 9923464f..950bbf8a 100644 --- a/Project.toml +++ b/Project.toml @@ -23,7 +23,6 @@ Reexport = "189a3867-3050-52da-a836-e630ba90ab69" SPIRVIntrinsics = "71d1d633-e7e8-4a92-83a1-de8814b09ba8" SPIRV_LLVM_Backend_jll = "4376b9bf-cff8-51b6-bb48-39421dff0d0c" SPIRV_Tools_jll = "6ac6d60f-d740-5983-97d7-a4482c0689f4" -StaticArrays = "90137ffa-7385-5640-81b9-e52037218182" UnsafeAtomics = "013be700-e6cd-48c3-b4a1-df204f14c38f" spirv2clc_jll = "f0274c0c-8c8a-59f1-85b7-f7d60330c5fb" @@ -55,7 +54,6 @@ Reexport = "1" SPIRVIntrinsics = "1.4" SPIRV_LLVM_Backend_jll = "23" SPIRV_Tools_jll = "2025.1" -StaticArrays = "1" UnsafeAtomics = "0.4" julia = "1.10" spirv2clc_jll = "0.2.2" diff --git a/ext/KernelAbstractionsExt.jl b/ext/KernelAbstractionsExt.jl index fa406c1d..14e97632 100644 --- a/ext/KernelAbstractionsExt.jl +++ b/ext/KernelAbstractionsExt.jl @@ -1,12 +1,9 @@ module KernelAbstractionsExt using OpenCL -using OpenCL: @device_override, method_table import KernelAbstractions as KA -import StaticArrays - import Adapt Adapt.adapt_storage(::KA.CPU, a::CLArray) = convert(Array, a) @@ -16,10 +13,4 @@ Adapt.adapt_storage(::KA.CPU, a::CLArray) = convert(Array, a) # rather than for `CLArray`. Adapt.adapt_storage(::KA.ConstAdaptor, a::CLDeviceArray) = Base.Experimental.Const(a) -## scratch memory - -@device_override @inline function KA.Scratchpad(ctx, ::Type{T}, ::Val{Dims}) where {T, Dims} - StaticArrays.MArray{Tuple{Dims...}, T}(undef) -end - end From a2e7d807d3e0cf3b87ab319d42dcfee6f1c4972c Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Sat, 10 Oct 2026 19:06:52 +0200 Subject: [PATCH 08/10] Implement KernelInterface's sub-group shuffles, votes and collectives Override the shuffles for the types `sub_group_shuffle` supports (with `cl_khr_subgroup_shuffle`), and only those, so that other types reach KernelInterface's word and field fallbacks; `shfl_down` and `shfl_up` return the own value past the width. Add the votes, with the ballot from `cl_khr_subgroup_ballot` or, without it, from 32-bit reductions, and the native reductions and scans where they have Julia's semantics. Sub-groups still need `cl_intel_required_subgroup_size` for a fixed width, now also when reported as an OpenCL 3.0 core feature, and of at most 64 lanes for the ballot. The linear layout is reported for PoCL and Intel's CPU and GPU runtimes, which the layout tests confirm; independent sub-groups everywhere but PoCL (whose sub-group operations are work-group barriers) and Intel's CPU runtime (which loses the writes after a `sub_group_barrier` that other sub-groups returned before, intel/llvm#23449). OpenCL C compilation through spirv2clc needs 0.2.5, which translates the sub-group operations. --- Project.toml | 2 +- lib/cl/device.jl | 10 ++- src/OpenCLKernels.jl | 137 ++++++++++++++++++++++++++++++++--- src/compiler/capabilities.jl | 1 + test/kernelinterface.jl | 34 ++++++++- 5 files changed, 168 insertions(+), 16 deletions(-) diff --git a/Project.toml b/Project.toml index 950bbf8a..03d985ba 100644 --- a/Project.toml +++ b/Project.toml @@ -56,4 +56,4 @@ SPIRV_LLVM_Backend_jll = "23" SPIRV_Tools_jll = "2025.1" UnsafeAtomics = "0.4" julia = "1.10" -spirv2clc_jll = "0.2.2" +spirv2clc_jll = "0.2.5" diff --git a/lib/cl/device.jl b/lib/cl/device.jl index 17806f82..a6c30888 100644 --- a/lib/cl/device.jl +++ b/lib/cl/device.jl @@ -369,7 +369,15 @@ function cl_device_type(dtype::Symbol) return cl_dtype end -sub_groups_supported(d::Device) = "cl_khr_subgroups" in d.extensions || "cl_intel_subgroups" in d.extensions +# sub-groups are an extension before OpenCL 3.0, and an optional core feature since, which +# devices without it report with a maximum of 0 sub-groups +function sub_groups_supported(d::Device) + ("cl_khr_subgroups" in d.extensions || "cl_intel_subgroups" in d.extensions) && return true + d.opencl_version >= v"3.0" || return false + num = Ref{cl_uint}(0) + err = unchecked_clGetDeviceInfo(d, CL_DEVICE_MAX_NUM_SUB_GROUPS, sizeof(cl_uint), num, C_NULL) + return err == CL_SUCCESS && num[] > 0 +end function sub_group_size(d::Device) sub_groups_supported(d) || return 0 if "cl_amd_device_attribute_query" in d.extensions diff --git a/src/OpenCLKernels.jl b/src/OpenCLKernels.jl index 48e5e987..e6069e66 100644 --- a/src/OpenCLKernels.jl +++ b/src/OpenCLKernels.jl @@ -1,7 +1,7 @@ module OpenCLKernels using ..OpenCL -using ..OpenCL: @device_override, method_table, kernel_convert, clfunction +using ..OpenCL: @device_override, method_table, kernel_convert, clfunction, has_feature import KernelInterface as KI @@ -21,6 +21,21 @@ A backend works with the task's active device if that is on its platform. Otherw for the backend (allocations, copies, compilation and launches) first activates the default device of its platform, as `KernelInterface.device!` would, so that the arrays it creates can be used afterwards. + +# Sub-groups + +KernelInterface's sub-groups need a width that is the same for every kernel. For now, they +are only supported on devices with `cl_intel_required_subgroup_size`, through which kernels +request the width that `KernelInterface.sub_group_size` reports (at most 64). The shuffles +also need `cl_khr_subgroup_shuffle`. + +OpenCL doesn't specify which work-items form a sub-group, so +`KernelInterface.supports_linear_subgroups` is only `true` on runtimes known to form them +from consecutive work-items: PoCL and Intel's CPU and GPU runtimes. On PoCL, sub-group operations are +work-group barriers, so `KernelInterface.supports_independent_subgroups` is `false`: all +sub-groups of a work-group have to execute the same sub-group operations, in the same order. +It is also `false` on Intel's CPU runtime, which loses the writes after a +`sub_group_barrier` that other sub-groups returned before (intel/llvm#23449). """ Base.@kwdef struct OpenCLBackend <: KI.Backend platform::cl.Platform = cl.platform() @@ -224,7 +239,8 @@ const DeviceProperties = @NamedTuple{ max_work_group_size::Int, max_work_group_dims::NTuple{3, Int}, compute_units::Int, float64::Bool, unified::Bool, # 0 if the device doesn't support sub-groups of a fixed width - sub_group_size::Int, shuffle_types::Vector{DataType}, + sub_group_size::Int, shuffle::Bool, float16::Bool, + linear_sub_groups::Bool, independent_sub_groups::Bool, } function device_properties(dev::cl.Device) cache = get!(task_local_storage(), :CLDeviceProperties) do @@ -232,20 +248,41 @@ function device_properties(dev::cl.Device) end::Dict{cl.Device, DeviceProperties} return get!(cache, dev) do sizes = dev.max_work_item_size - # the sub-group width is only fixed for kernels that request it, which `clfunction` - # does for devices with `cl_intel_required_subgroup_size` - fixed_sub_groups = cl.sub_groups_supported(dev) && - "cl_intel_required_subgroup_size" in dev.extensions + # KernelInterface needs a fixed sub-group width, which for now means one that kernels + # request with `cl_intel_required_subgroup_size` (as `clfunction` does). Wider + # sub-groups than 64 aren't reported, since `KI.sub_group_ballot` has a bit per lane. + sub_group_size = 0 + if cl.sub_groups_supported(dev) && "cl_intel_required_subgroup_size" in dev.extensions + width = cl.sub_group_size(dev) + width <= 64 && (sub_group_size = width) + end + pocl = occursin("pocl", dev.platform.vendor) (; max_work_group_size = Int(dev.max_work_group_size), max_work_group_dims = ntuple(d -> d <= length(sizes) ? Int(sizes[d]) : 1, 3), compute_units = Int(dev.max_compute_units), float64 = "cl_khr_fp64" in dev.extensions, unified = cl.default_memory_backend(dev; unified=true) !== nothing, - sub_group_size = fixed_sub_groups ? cl.sub_group_size(dev) : 0, - shuffle_types = fixed_sub_groups ? cl.sub_group_shuffle_supported_types(dev) : DataType[]) + sub_group_size, + shuffle = sub_group_size > 0 && "cl_khr_subgroup_shuffle" in dev.extensions, + float16 = "cl_khr_fp16" in dev.extensions, + # OpenCL leaves the layout to the implementation: only report it for the + # runtimes known to form sub-groups from consecutive work-items + linear_sub_groups = sub_group_size > 0 && (pocl || linear_sub_group_runtime(dev)), + # PoCL executes sub-group operations as work-group barriers, and Intel's CPU + # runtime loses the writes after a `sub_group_barrier` that other sub-groups + # returned before (intel/llvm#23449) + independent_sub_groups = sub_group_size > 0 && !pocl && !intel_cpu_runtime(dev)) end end +# Intel's CPU and GPU runtimes form sub-groups from consecutive work-items in the x +# dimension, as KernelInterface's layout tests check +function linear_sub_group_runtime(dev::cl.Device) + return occursin("Intel", dev.platform.vendor) && dev.device_type in (:cpu, :gpu) +end + +intel_cpu_runtime(dev::cl.Device) = occursin("Intel", dev.platform.vendor) && dev.device_type === :cpu + device_properties(b::OpenCLBackend) = device_properties(backend_device(b)) KI.max_work_group_size(b::OpenCLBackend)::Int = device_properties(b).max_work_group_size @@ -263,7 +300,22 @@ KI.supports_atomics(::OpenCLBackend) = true KI.supports_subgroups(b::OpenCLBackend) = device_properties(b).sub_group_size > 0 KI.sub_group_size(b::OpenCLBackend)::Int = device_properties(b).sub_group_size -KI.supports_shuffle(b::OpenCLBackend, ::Type{T}) where {T} = T in device_properties(b).shuffle_types +KI.supports_linear_subgroups(b::OpenCLBackend) = device_properties(b).linear_sub_groups +KI.supports_independent_subgroups(b::OpenCLBackend) = device_properties(b).independent_sub_groups + +# The types `sub_group_shuffle` supports; other types are shuffled as words or field by +# field. `Float16` and `Float64` are reported as unsupported on devices without `cl_khr_fp16` +# or `cl_khr_fp64`: a kernel can't hold such values there, even just to move them, since +# LLVM turns a load of an integer that is reinterpreted as a float into a load of the float. +const ShuffleTypes = Union{OpenCL.SPIRVIntrinsics.gentypes...} + +function KI.supports_shuffle(b::OpenCLBackend, ::Type{T}) where {T <: ShuffleTypes} + props = device_properties(b) + props.shuffle || return false + T === Float64 && return props.float64 + T === Float16 && return props.float16 + return true +end ## Indexing Functions @@ -319,9 +371,70 @@ end sub_group_barrier(OpenCL.LOCAL_MEM_FENCE | OpenCL.GLOBAL_MEM_FENCE) end -# out-of-range source lanes give an unspecified value, as KernelInterface allows -@device_override function KI.shfl_down(val::T, offset::Integer) where T - sub_group_shuffle(val, get_sub_group_local_id() % UInt32 + offset % UInt32) +# `cl_khr_subgroup_shuffle` +@device_override KI.shfl(val::T, lane::Integer) where {T <: ShuffleTypes} = + sub_group_shuffle(val, lane) + +# past the sub-group width, `shfl_down` and `shfl_up` return the work-item's own value, which +# `sub_group_shuffle` (like SPIR-V's `OpGroupNonUniformShuffleDown`) leaves undefined +@device_override function KI.shfl_down(val::T, offset::Integer) where {T <: ShuffleTypes} + lane = get_sub_group_local_id() + # compared before adding, so that large offsets don't overflow + inside = offset <= get_max_sub_group_size() - lane + return sub_group_shuffle(val, ifelse(inside, lane + offset, lane)) +end + +@device_override function KI.shfl_up(val::T, offset::Integer) where {T <: ShuffleTypes} + lane = get_sub_group_local_id() + return sub_group_shuffle(val, ifelse(lane > offset, lane - offset, lane)) +end + +@device_override KI.shfl_xor(val::T, mask::Integer) where {T <: ShuffleTypes} = + sub_group_shuffle_xor(val, mask) + +# `cl_khr_subgroups` +@device_override KI.sub_group_any(pred::Bool) = sub_group_any(pred) + +@device_override KI.sub_group_all(pred::Bool) = sub_group_all(pred) + +@device_override function KI.sub_group_ballot(pred::Bool) + if has_feature(:subgroup_ballot) + mask = sub_group_ballot(pred) + return UInt64(mask[1].value) | (UInt64(mask[2].value) << 32) + else + return reduction_ballot(pred) + end +end + +# The ballot without `cl_khr_subgroup_ballot`: add up a bit per lane, with 32-bit reductions, +# which every device with sub-groups supports. The width is at most 64 (see +# `device_properties`). +@inline function reduction_ballot(pred::Bool) + lane = (get_sub_group_local_id() - 1) % UInt32 + lo = sub_group_reduce_add(ifelse(pred & (lane < 32), UInt32(1) << lane, UInt32(0))) + get_max_sub_group_size() <= 32 && return UInt64(lo) + hi = sub_group_reduce_add(ifelse(pred & (lane >= 32), UInt32(1) << (lane - 32), UInt32(0))) + return UInt64(lo) | (UInt64(hi) << 32) +end + +# Native reductions and scans, for `+` on 32- and 64-bit integers and floats, and `min`/`max` +# on integers (OpenCL's `min` and `max` treat NaN and the sign of zero differently from +# Julia's). +const CollectiveIntTypes = Union{Int32, UInt32, Int64, UInt64} +const CollectiveTypes = Union{CollectiveIntTypes, Float16, Float32, Float64} + +@device_override KI.sub_group_reduce(::typeof(+), val::CollectiveTypes) = sub_group_reduce_add(val) +@device_override KI.sub_group_reduce(::typeof(min), val::CollectiveIntTypes) = sub_group_reduce_min(val) +@device_override KI.sub_group_reduce(::typeof(max), val::CollectiveIntTypes) = sub_group_reduce_max(val) + +@device_override KI.sub_group_scan(::typeof(+), val::CollectiveTypes) = sub_group_scan_inclusive_add(val) +@device_override KI.sub_group_scan(::typeof(min), val::CollectiveIntTypes) = sub_group_scan_inclusive_min(val) +@device_override KI.sub_group_scan(::typeof(max), val::CollectiveIntTypes) = sub_group_scan_inclusive_max(val) + +# the exclusive scans of `cl_khr_subgroups` start from the identity, not from `init` +@device_override function KI.sub_group_exclusive_scan(::typeof(+), val::T, init::T) where {T <: CollectiveTypes} + prefix = sub_group_scan_exclusive_add(val) + return ifelse(get_sub_group_local_id() == 1, init, init + prefix) end @device_override @inline function KI._print(args...) diff --git a/src/compiler/capabilities.jl b/src/compiler/capabilities.jl index 88b3470c..ba8c4fdf 100644 --- a/src/compiler/capabilities.jl +++ b/src/compiler/capabilities.jl @@ -105,6 +105,7 @@ const FEATURES = Feature[ Feature(:fp64, dev -> "cl_khr_fp64" in dev.extensions), Feature(:int64_atomics, dev -> "cl_khr_int64_base_atomics" in dev.extensions), Feature(:subgroups, cl.sub_groups_supported), + Feature(:subgroup_ballot, dev -> "cl_khr_subgroup_ballot" in dev.extensions), Feature(:generic_address_space, dev -> has_opencl_c_feature(dev, "__opencl_c_generic_address_space")), # What the device reports for floating-point atomics in both global and local memory. diff --git a/test/kernelinterface.jl b/test/kernelinterface.jl index 8242cff2..8ac62432 100644 --- a/test/kernelinterface.jl +++ b/test/kernelinterface.jl @@ -105,7 +105,9 @@ end backend = OpenCLBackend() dev = cl.device() # kernels only execute with a fixed sub-group width if they can request one - if cl.sub_groups_supported(dev) && "cl_intel_required_subgroup_size" in dev.extensions + # (of at most 64 lanes, for the ballot) + if cl.sub_groups_supported(dev) && "cl_intel_required_subgroup_size" in dev.extensions && + cl.sub_group_size(dev) <= 64 @test KI.supports_subgroups(backend) @test KI.sub_group_size(backend) == cl.sub_group_size(dev) @test KI.supports_shuffle(backend, Int32) == @@ -117,10 +119,38 @@ end KI.@launch backend launch=false sub_group_size=width ki_fill_kernel(a, Int32(1)) @test_throws ArgumentError KI.@launch backend launch=false sub_group_size=nothing ki_fill_kernel(a, Int32(1)) @test_throws ArgumentError KI.@launch backend launch=false sub_group_size=2width ki_fill_kernel(a, Int32(1)) + # PoCL executes sub-group operations as work-group barriers + occursin("pocl", dev.platform.vendor) && @test !KI.supports_independent_subgroups(backend) else @test !KI.supports_subgroups(backend) + @test !KI.supports_linear_subgroups(backend) + @test !KI.supports_independent_subgroups(backend) + end + # structs are shuffled field by field + @test KI.supports_shuffle(backend, Complex{Float32}) == KI.supports_shuffle(backend, Float32) +end + +# the ballot of devices without `cl_khr_subgroup_ballot` +function reduction_ballot_kernel(out, pred) + i = KI.get_local_id().x + @inbounds out[i, 1] = OpenCL.OpenCLKernels.reduction_ballot(pred[i]) + @inbounds out[i, 2] = KI.sub_group_ballot(pred[i]) + return +end + +@testset "reduction ballot" begin + backend = OpenCLBackend() + if KI.supports_subgroups(backend) + n = 2 * KI.sub_group_size(backend) + 5 + pred = CLArray(rand(Bool, n)) + out = CLArray(zeros(UInt64, n, 2)) + kernel = KI.@launch backend launch=false reduction_ballot_kernel(out, pred) + if n <= KI.max_work_group_size(kernel) + kernel(out, pred; workgroupsize = n) + out = Array(out) + @test out[:, 1] == out[:, 2] + end end - @test !KI.supports_shuffle(backend, Complex{Float32}) end @testset "events" begin From b6b34399b7e73948980543a7cb8c47b04f3b0e6c Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Sun, 11 Oct 2026 11:33:55 +0200 Subject: [PATCH 09/10] Report partitionable sub-groups on Intel GPUs Report `KI.supports_partitionable_subgroups` only on Intel's GPU runtime with `cl_khr_subgroup_ballot`, where a sub-group is the SIMD lanes of one EU thread and the width functions' shuffles and ballot only involve the active lanes. OpenCL has no barrier for part of a sub-group, so a new `:partitionable_sub_groups` device feature lets `KI.sub_group_barrier(width)` use a sub-group-scope memory barrier there for segments narrower than the sub-group; elsewhere, including PoCL (whose sub-group functions are work-group barriers) and Intel's CPU runtime, the whole sub-group keeps executing the width functions together. --- src/OpenCLKernels.jl | 50 ++++++++++++++++++++++++++++++------ src/compiler/capabilities.jl | 8 ++++++ 2 files changed, 50 insertions(+), 8 deletions(-) diff --git a/src/OpenCLKernels.jl b/src/OpenCLKernels.jl index e6069e66..88a6b50a 100644 --- a/src/OpenCLKernels.jl +++ b/src/OpenCLKernels.jl @@ -1,7 +1,8 @@ module OpenCLKernels using ..OpenCL -using ..OpenCL: @device_override, method_table, kernel_convert, clfunction, has_feature +using ..OpenCL: @device_override, method_table, kernel_convert, clfunction, has_feature, + SPIRVIntrinsics import KernelInterface as KI @@ -36,6 +37,12 @@ work-group barriers, so `KernelInterface.supports_independent_subgroups` is `fal sub-groups of a work-group have to execute the same sub-group operations, in the same order. It is also `false` on Intel's CPU runtime, which loses the writes after a `sub_group_barrier` that other sub-groups returned before (intel/llvm#23449). + +OpenCL has no barrier for part of a sub-group, so +`KernelInterface.supports_partitionable_subgroups` is only `true` on Intel's GPU runtime, where +a sub-group is the SIMD lanes of one hardware thread (and with `cl_khr_subgroup_ballot`). +Elsewhere, the segments of a sub-group have to execute the sub-group functions with a `width` +together, like the ones without. """ Base.@kwdef struct OpenCLBackend <: KI.Backend platform::cl.Platform = cl.platform() @@ -240,7 +247,7 @@ const DeviceProperties = @NamedTuple{ float64::Bool, unified::Bool, # 0 if the device doesn't support sub-groups of a fixed width sub_group_size::Int, shuffle::Bool, float16::Bool, - linear_sub_groups::Bool, independent_sub_groups::Bool, + linear_sub_groups::Bool, independent_sub_groups::Bool, partitionable_sub_groups::Bool, } function device_properties(dev::cl.Device) cache = get!(task_local_storage(), :CLDeviceProperties) do @@ -271,7 +278,13 @@ function device_properties(dev::cl.Device) # PoCL executes sub-group operations as work-group barriers, and Intel's CPU # runtime loses the writes after a `sub_group_barrier` that other sub-groups # returned before (intel/llvm#23449) - independent_sub_groups = sub_group_size > 0 && !pocl && !intel_cpu_runtime(dev)) + independent_sub_groups = sub_group_size > 0 && !pocl && !intel_cpu_runtime(dev), + # PoCL limitation: as its sub-group functions are work-group barriers, the segments of + # a sub-group can't diverge either. Nor can they on other runtimes, since OpenCL has no + # barrier for part of a sub-group (see `:partitionable_sub_groups`). It implies + # independent sub-groups. + partitionable_sub_groups = sub_group_size > 0 && !pocl && !intel_cpu_runtime(dev) && + OpenCL.feature_supported(dev, :partitionable_sub_groups)) end end @@ -302,11 +315,12 @@ KI.supports_subgroups(b::OpenCLBackend) = device_properties(b).sub_group_size > KI.sub_group_size(b::OpenCLBackend)::Int = device_properties(b).sub_group_size KI.supports_linear_subgroups(b::OpenCLBackend) = device_properties(b).linear_sub_groups KI.supports_independent_subgroups(b::OpenCLBackend) = device_properties(b).independent_sub_groups +KI.supports_partitionable_subgroups(b::OpenCLBackend) = device_properties(b).partitionable_sub_groups -# The types `sub_group_shuffle` supports; other types are shuffled as words or field by -# field. `Float16` and `Float64` are reported as unsupported on devices without `cl_khr_fp16` -# or `cl_khr_fp64`: a kernel can't hold such values there, even just to move them, since -# LLVM turns a load of an integer that is reinterpreted as a float into a load of the float. +# The types `sub_group_shuffle` supports; other types are shuffled as words. `Float16` and +# `Float64` are reported as unsupported on devices without `cl_khr_fp16` or `cl_khr_fp64`: a +# kernel can't hold such values there, even just to move them, since LLVM turns a load of an +# integer that is reinterpreted as a float into a load of the float. const ShuffleTypes = Union{OpenCL.SPIRVIntrinsics.gentypes...} function KI.supports_shuffle(b::OpenCLBackend, ::Type{T}) where {T <: ShuffleTypes} @@ -371,7 +385,26 @@ end sub_group_barrier(OpenCL.LOCAL_MEM_FENCE | OpenCL.GLOBAL_MEM_FENCE) end +# OpenCL limitation: a sub-group barrier has to be executed by the whole sub-group, and there +# is no barrier for part of one. Where a segment's lanes execute together anyway, a narrower +# segment only fences its memory accesses; elsewhere, the whole sub-group executes this. +@device_override @inline function KI.sub_group_barrier(width::Integer) + if has_feature(:partitionable_sub_groups) && width != get_max_sub_group_size() + atomic_work_item_fence(OpenCL.LOCAL_MEM_FENCE | OpenCL.GLOBAL_MEM_FENCE, + SPIRVIntrinsics.memory_order_acq_rel, + SPIRVIntrinsics.memory_scope_sub_group) + else + KI.sub_group_barrier() + end +end + # `cl_khr_subgroup_shuffle` +# +# `KI.shfl_word_tuple` could exchange the words of a struct with the vector shuffles of +# `SPV_INTEL_subgroups`, as KernelAbstractions' POCL backend does: Intel's GPU runtime supports +# them (`cl_intel_spirv_subgroups`), and so does PoCL's CPU device with `cl_intel_subgroups` +# since pocl/pocl#2394, which matters most there. PoCL limitation: it doesn't advertise +# `cl_intel_spirv_subgroups`, so the words are shuffled one by one until that can be detected. @device_override KI.shfl(val::T, lane::Integer) where {T <: ShuffleTypes} = sub_group_shuffle(val, lane) @@ -408,7 +441,8 @@ end # The ballot without `cl_khr_subgroup_ballot`: add up a bit per lane, with 32-bit reductions, # which every device with sub-groups supports. The width is at most 64 (see -# `device_properties`). +# `device_properties`). The whole sub-group has to execute these reductions, so the segments +# of a sub-group are only independent with the native ballot. @inline function reduction_ballot(pred::Bool) lane = (get_sub_group_local_id() - 1) % UInt32 lo = sub_group_reduce_add(ifelse(pred & (lane < 32), UInt32(1) << lane, UInt32(0))) diff --git a/src/compiler/capabilities.jl b/src/compiler/capabilities.jl index ba8c4fdf..c7a4c049 100644 --- a/src/compiler/capabilities.jl +++ b/src/compiler/capabilities.jl @@ -106,6 +106,14 @@ const FEATURES = Feature[ Feature(:int64_atomics, dev -> "cl_khr_int64_base_atomics" in dev.extensions), Feature(:subgroups, cl.sub_groups_supported), Feature(:subgroup_ballot, dev -> "cl_khr_subgroup_ballot" in dev.extensions), + # Whether a segment of consecutive lanes of a sub-group can communicate (shuffle, ballot and + # fence) without the rest of the sub-group. OpenCL limitation: OpenCL has no barrier for + # part of a sub-group, so this relies on the execution model of Intel's GPU runtime, where a + # sub-group is the SIMD lanes of one EU thread, and on the native ballot, which only reports + # the active lanes. + Feature(:partitionable_sub_groups, + dev -> occursin("Intel", dev.platform.vendor) && dev.device_type === :gpu && + "cl_khr_subgroup_ballot" in dev.extensions), Feature(:generic_address_space, dev -> has_opencl_c_feature(dev, "__opencl_c_generic_address_space")), # What the device reports for floating-point atomics in both global and local memory. From 7081ec7202f7ff04b479a4d47176e738d0c097d5 Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Tue, 13 Oct 2026 06:52:16 +0200 Subject: [PATCH 10/10] Skip the random number tests on Intel's CPU runtime without AVX-512 Without AVX-512, the runtime has no masked 32- and 64-wide `pown` and `rootn`, so kernels calling them in a branch fail to link at the sub-group width we request. GitHub's AVX2-only runners (AMD EPYC 7763) hit this in the random number tests. --- test/runtests.jl | 13 +++++++++++++ 1 file changed, 13 insertions(+) diff --git a/test/runtests.jl b/test/runtests.jl index d11976d4..0c06dd55 100644 --- a/test/runtests.jl +++ b/test/runtests.jl @@ -141,6 +141,19 @@ for t in targets end end +# On hosts without AVX-512, Intel's CPU runtime lacks the masked 32- and 64-wide variants of +# `pown` and `rootn` (e.g. `_Z4pownDv32_fDv32_iS_`), so kernels calling them under divergent +# control flow fail to link at the sub-group width we request (intel/llvm#23456). The random +# number tests do. +if !Base.BinaryPlatforms.CPUID.test_cpu_feature(Base.BinaryPlatforms.CPUID.JL_X86_avx512f) + for t in targets + t.platform == "Intel(R) OpenCL" || continue + for name in ("device/random", "gpuarrays/random") + delete!(testsuite, "$(t.label)/$name") + end + end +end + const init_worker_code = quote using OpenCL, $driver_pkg