From 2124760be9da593c28db0fda4a04935d4f9e11b4 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Mos=C3=A8=20Giordano?= Date: Fri, 16 Oct 2026 15:09:59 -0500 Subject: [PATCH 1/4] POCL: hash the compiler configuration's keyword arguments as a NamedTuple `compiler_config` is looked up on every launch, and keyed by `hash(kwargs)`. The generic `hash(::AbstractDict)` iterates over the `Pairs`, whose values have different types (e.g. `name::Nothing` and `sub_group_size::Int`), so every element was boxed: 16 bytes per launch. Hash `values(kwargs)`, a `NamedTuple`, instead, like the other back-ends. `compiler_config` and `clfunction` of a compiled kernel no longer allocate, which the new tests check where Julia allows: on 1.10 both allocate, and on 1.11 GPUCompiler's lookup of the compiled code does. Bump the version to 0.10.1. Co-Authored-By: Claude Opus 5.5 Claude-Session: https://claude.ai/code/session_01D6fxP52AyNqfcKSSQXFxfB --- Project.toml | 2 +- src/pocl/compiler/compilation.jl | 3 ++- test/runtests.jl | 23 +++++++++++++++++++++++ 3 files changed, 26 insertions(+), 2 deletions(-) diff --git a/Project.toml b/Project.toml index ae8d6eef7..40c6d03fe 100644 --- a/Project.toml +++ b/Project.toml @@ -1,6 +1,6 @@ name = "KernelAbstractions" uuid = "63c18a36-062a-441e-b654-da1e3ab1ce7c" -version = "0.10.0" +version = "0.10.1" authors = ["Valentin Churavy and contributors"] [workspace] diff --git a/src/pocl/compiler/compilation.jl b/src/pocl/compiler/compilation.jl index 60844c17d..eedf937de 100644 --- a/src/pocl/compiler/compilation.jl +++ b/src/pocl/compiler/compilation.jl @@ -190,7 +190,8 @@ const default_extensions = "+SPV_KHR_no_integer_wrap_decoration" const _toolchain = Ref{Any}() const _compiler_configs = Dict{UInt, OpenCLCompilerConfig}() function compiler_config(dev::cl.Device; kwargs...) - h = hash(dev, hash(kwargs)) + # (as a `NamedTuple`: hashing the `Pairs` boxes each value, on every launch) + h = hash(dev, hash(values(kwargs))) # launches already hold this (reentrant) lock, but reflection doesn't return @lock clfunction_lock begin config = get(_compiler_configs, h, nothing) diff --git a/test/runtests.jl b/test/runtests.jl index 088b010a2..4af6e462c 100644 --- a/test/runtests.jl +++ b/test/runtests.jl @@ -423,6 +423,29 @@ end end end +# launching a compiled kernel doesn't allocate +@testset "POCL launch allocations" begin + mod = @eval module $(gensym()) + import KernelAbstractions.POCL + kernel(A) = (A[POCL.get_global_id()] = 1; return) + compiler_config(dev, sub_group_size) = POCL.compiler_config(dev; name = nothing, sub_group_size) + # `tt` is only specialized on as a type parameter + clfunction(::Type{TT}, sub_group_size) where {TT} = + POCL.clfunction(kernel, TT; name = nothing, sub_group_size) + end + # with `args...`, Julia 1.11 and older box the result + allocated(f::F, args::Vararg{Any, N}) where {F, N} = (f(args...); @allocated f(args...)) + + # these allocate on Julia 1.10, and those that look up compiled code in GPUCompiler on 1.11 + skip = VERSION < v"1.11" + skip_lookup = VERSION < v"1.12" + tt = Tuple{typeof(POCL.clconvert(zeros(Int, 4)))} + sub_group_size = POCL.device_limits().sub_group_size + sub_group_size = sub_group_size > 0 ? sub_group_size : nothing + @test allocated(mod.compiler_config, POCL.device(), sub_group_size) == 0 skip = skip + @test allocated(mod.clfunction, tt, sub_group_size) == 0 skip = skip_lookup +end + @testset "POCL exceptions" begin @kernel function exception_kernel!(A) I = @index(Global, Linear) From 173d323d610d62ec8164b8d939443a6ac21ac20f Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Sat, 17 Oct 2026 11:10:27 +0200 Subject: [PATCH 2/4] POCL: determine the default sub-group width with the compiler configuration MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit `KI.kernel_function` passed the device's sub-group width to `clfunction`, as `nothing` on devices without sub-groups. The keyword arguments' type was thus only known at run time, so the keyword call to `clfunction` was dynamic: building the `NamedTuple` and boxing the returned kernel allocated 64 bytes on every launch. Instead, default the width to the device's in `_compiler_config`, which only runs when creating a configuration, and don't pass it at all. That also changes the default of `@opencl` from 32 to that width, which is the same on devices that support 32, and works on those that don't. `KI.kernel_function` of a compiled kernel no longer allocates, which a new test checks. Co-authored-by: Mosè Giordano --- src/pocl/backend.jl | 9 ++++----- src/pocl/compiler/compilation.jl | 5 +++-- src/pocl/pocl.jl | 16 +++++++++++----- test/runtests.jl | 16 ++++++++-------- 4 files changed, 26 insertions(+), 20 deletions(-) diff --git a/src/pocl/backend.jl b/src/pocl/backend.jl index 72ce21602..32137671d 100644 --- a/src/pocl/backend.jl +++ b/src/pocl/backend.jl @@ -165,11 +165,10 @@ struct POCLKernel{K, F} end function KI.kernel_function(backend::POCLBackend, f::F, tt::TT = Tuple{}; name = nothing, kwargs...) where {F, TT} - # fix the sub-group width, as `KI.sub_group_size` promises. pass it even if the device - # has no sub-groups, so that `clfunction` is only compiled for one set of keywords. - sub_group_size = device_limits().sub_group_size - sub_group_size = sub_group_size > 0 ? sub_group_size : nothing - kernel = clfunction(clconvert(f), tt; name, sub_group_size, kwargs...) + # the sub-group width defaults to the one `KI.sub_group_size` reports. it is determined + # when creating the compiler configuration, since passing it here would make the keyword + # arguments' type depend on the device, and the call to `clfunction` dynamic. + kernel = clfunction(clconvert(f), tt; name, kwargs...) kern = POCLKernel(kernel, f) return KI.Kernel{POCLBackend, typeof(kern)}(backend, kern) end diff --git a/src/pocl/compiler/compilation.jl b/src/pocl/compiler/compilation.jl index eedf937de..7ac9db287 100644 --- a/src/pocl/compiler/compilation.jl +++ b/src/pocl/compiler/compilation.jl @@ -165,7 +165,8 @@ end The GPUCompiler configuration for compiling kernels for `dev`, cached per device and keyword arguments. Besides those of `CompilerConfig` (`kernel`, `name`, `always_inline`, -`debug_level`) and `sub_group_size`, it takes: +`debug_level`) and `sub_group_size` (the sub-group width to require, or `nothing` for none, +defaulting to the width `KI.sub_group_size` reports), it takes: - `atomics`: override the atomic capabilities GPUCompiler may select directly. This replaces the whole device-derived `SPIRVAtomics` (see `spirv_atomics`), and defaults to @@ -205,7 +206,7 @@ end @noinline function _compiler_config( dev; kernel = true, name = nothing, always_inline = false, debug_level = Base.JLOptions().debug_level, - sub_group_size::Union{Nothing, Int} = 32, + sub_group_size::Union{Nothing, Int} = default_sub_group_size(dev), atomics::SPIRVAtomics = spirv_atomics(dev), extensions::Union{Nothing, String} = nothing, kwargs... ) diff --git a/src/pocl/pocl.jl b/src/pocl/pocl.jl index a46d0ecbf..25e6b72fe 100644 --- a/src/pocl/pocl.jl +++ b/src/pocl/pocl.jl @@ -23,6 +23,15 @@ struct Session } end +# the sub-group width kernels are compiled for, or `nothing` if the device has no sub-groups. +# POCL can technically support any sub-group size; prefer the common GPU ones, and only +# widths of up to 64, so that `KI.sub_group_ballot` (a `UInt64` mask) covers every lane. +function default_sub_group_size(device::cl.Device) + sg_sizes = filter(<=(64), device.sub_group_sizes) + common = filter(in(sg_sizes), [32, 64, 16, sg_sizes...]) + return isempty(common) ? nothing : first(common) +end + function Session() cl.pocl_standalone_jll.is_available() || error("The CPU back-end needs PoCL, which is not available for $(Sys.MACHINE)") @@ -39,15 +48,12 @@ function Session() context = cl.Context(device) sizes = device.max_work_item_size - # POCL can technically support any sub-group size; prefer the common GPU ones, and only - # widths of up to 64, so that `KI.sub_group_ballot` (a `UInt64` mask) covers every lane - sg_sizes = filter(<=(64), device.sub_group_sizes) - common = filter(in(sg_sizes), [32, 64, 16, sg_sizes...]) + sub_group_size = default_sub_group_size(device) limits = (; max_work_group_size = Int(device.max_work_group_size), max_work_group_dims = ntuple(d -> d <= length(sizes) ? sizes[d] : 1, 3), # 0 if the device has no sub-groups - sub_group_size = isempty(common) ? 0 : first(common), + sub_group_size = something(sub_group_size, 0), # the threads PoCL runs work-groups on compute_units = Int(device.max_compute_units), ) diff --git a/test/runtests.jl b/test/runtests.jl index 4af6e462c..1f9c30951 100644 --- a/test/runtests.jl +++ b/test/runtests.jl @@ -426,12 +426,13 @@ end # launching a compiled kernel doesn't allocate @testset "POCL launch allocations" begin mod = @eval module $(gensym()) - import KernelAbstractions.POCL + import KernelAbstractions, KernelAbstractions.POCL kernel(A) = (A[POCL.get_global_id()] = 1; return) - compiler_config(dev, sub_group_size) = POCL.compiler_config(dev; name = nothing, sub_group_size) + compiler_config(dev) = POCL.compiler_config(dev; name = nothing) # `tt` is only specialized on as a type parameter - clfunction(::Type{TT}, sub_group_size) where {TT} = - POCL.clfunction(kernel, TT; name = nothing, sub_group_size) + clfunction(::Type{TT}) where {TT} = POCL.clfunction(kernel, TT; name = nothing) + kernel_function(::Type{TT}) where {TT} = + KernelAbstractions.KernelInterface.kernel_function(POCL.POCLBackend(), kernel, TT) end # with `args...`, Julia 1.11 and older box the result allocated(f::F, args::Vararg{Any, N}) where {F, N} = (f(args...); @allocated f(args...)) @@ -440,10 +441,9 @@ end skip = VERSION < v"1.11" skip_lookup = VERSION < v"1.12" tt = Tuple{typeof(POCL.clconvert(zeros(Int, 4)))} - sub_group_size = POCL.device_limits().sub_group_size - sub_group_size = sub_group_size > 0 ? sub_group_size : nothing - @test allocated(mod.compiler_config, POCL.device(), sub_group_size) == 0 skip = skip - @test allocated(mod.clfunction, tt, sub_group_size) == 0 skip = skip_lookup + @test allocated(mod.compiler_config, POCL.device()) == 0 skip = skip + @test allocated(mod.clfunction, tt) == 0 skip = skip_lookup + @test allocated(mod.kernel_function, tt) == 0 skip = skip_lookup end @testset "POCL exceptions" begin From 895c47d1d64806eafc1cee716f6fcb6889e39f13 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Mos=C3=A8=20Giordano?= Date: Fri, 16 Oct 2026 15:19:48 -0500 Subject: [PATCH 3/4] POCL: inline `check` Every `@checked` OpenCL wrapper runs the API call in a closure passed to `check`. When `check` isn't inlined, a `Ref` the call writes its result to is captured by the closure and escapes, so it is heap-allocated even in callers that inline the wrapper itself. Inline `check`, so that such a `Ref` can live on the stack. On its own this doesn't remove an allocation, since the wrappers aren't inlined where it matters; it allows inlining the wrapper for `clEnqueueNDRangeKernel`, whose `Ref` for the event is allocated on every launch. Co-Authored-By: Claude Opus 5.5 Claude-Session: https://claude.ai/code/session_01D6fxP52AyNqfcKSSQXFxfB --- src/pocl/nanoOpenCL.jl | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/src/pocl/nanoOpenCL.jl b/src/pocl/nanoOpenCL.jl index b7bfcfb28..d1b5e2fc0 100644 --- a/src/pocl/nanoOpenCL.jl +++ b/src/pocl/nanoOpenCL.jl @@ -426,7 +426,8 @@ end throw(CLError(res)) end -function check(f) +# inlined, so that `Ref`s passed to API calls in `f` can be allocated on the stack +@inline function check(f) res = f() if res != CL_SUCCESS From da38f0d60381ff6014bafc5cbae1364594d27cc4 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Mos=C3=A8=20Giordano?= Date: Fri, 16 Oct 2026 15:24:04 -0500 Subject: [PATCH 4/4] POCL: allocate the event of a kernel launch on the stack `enqueue_kernel` passes a `Ref` for the launch's event to `clEnqueueNDRangeKernel`, which wasn't inlined, so the `Ref` escaped and every launch allocated 16 bytes. Inline the wrapper at this call site (which, with `check` inlined, lets the `Ref` live on the stack). With this, launching a compiled kernel through KernelAbstractions no longer allocates (down from 96 bytes on 0.10.0). Waiting for it doesn't either since GPUToolbox 3.4.1, also when the kernel takes long enough for the wait to involve a completion callback, so require that. Test that a launch through `launch_tuple` and `@kernel` launches of a short and a long kernel don't allocate. Co-authored-by: Tim Besard Co-Authored-By: Claude Opus 5.5 Claude-Session: https://claude.ai/code/session_01D6fxP52AyNqfcKSSQXFxfB --- Project.toml | 2 +- src/pocl/nanoOpenCL.jl | 3 ++- test/runtests.jl | 35 +++++++++++++++++++++++++++++++++++ 3 files changed, 38 insertions(+), 2 deletions(-) diff --git a/Project.toml b/Project.toml index 40c6d03fe..cdddadd42 100644 --- a/Project.toml +++ b/Project.toml @@ -43,7 +43,7 @@ StaticArraysExt = "StaticArrays" Adapt = "0.4, 1.0, 2.0, 3.0, 4" Atomix = "1.5" GPUCompiler = "2.13" -GPUToolbox = "3.3.3" +GPUToolbox = "3.4.1" KernelInterface = "0.4" LLVM = "10" LinearAlgebra = "1.6" diff --git a/src/pocl/nanoOpenCL.jl b/src/pocl/nanoOpenCL.jl index d1b5e2fc0..6b74897e0 100644 --- a/src/pocl/nanoOpenCL.jl +++ b/src/pocl/nanoOpenCL.jl @@ -1356,8 +1356,9 @@ function enqueue_kernel( if length(local_work_size) != work_dim throw(ArgumentError("global_work_size and local_work_size have differing dims")) end + # inlined, so that `ret_event` can be allocated on the stack ret_event = Ref{cl_event}() - clEnqueueNDRangeKernel( + @inline clEnqueueNDRangeKernel( queue(), k, cl_uint(work_dim), work_size(global_work_size), work_size(local_work_size), ret_event ) diff --git a/test/runtests.jl b/test/runtests.jl index 1f9c30951..2d90fe9fc 100644 --- a/test/runtests.jl +++ b/test/runtests.jl @@ -427,12 +427,27 @@ end @testset "POCL launch allocations" begin mod = @eval module $(gensym()) import KernelAbstractions, KernelAbstractions.POCL + using KernelAbstractions: @kernel, @index kernel(A) = (A[POCL.get_global_id()] = 1; return) compiler_config(dev) = POCL.compiler_config(dev; name = nothing) # `tt` is only specialized on as a type parameter clfunction(::Type{TT}) where {TT} = POCL.clfunction(kernel, TT; name = nothing) kernel_function(::Type{TT}) where {TT} = KernelAbstractions.KernelInterface.kernel_function(POCL.POCLBackend(), kernel, TT) + function launch_tuple(kernel, A, info) + event = POCL.launch_tuple(kernel, (A,), info; global_size = (4,), local_size = (4,)) + wait(event) + POCL.cl.clReleaseEvent(event) + end + @kernel function ka_kernel!(A, n) + I = @index(Global, Linear) + x = A[I] + for _ in 1:n + x = sin(x) + end + @inbounds A[I] = x + end + launch(kernel, A, n) = kernel(A, n; ndrange = length(A)) end # with `args...`, Julia 1.11 and older box the result allocated(f::F, args::Vararg{Any, N}) where {F, N} = (f(args...); @allocated f(args...)) @@ -444,6 +459,26 @@ end @test allocated(mod.compiler_config, POCL.device()) == 0 skip = skip @test allocated(mod.clfunction, tt) == 0 skip = skip_lookup @test allocated(mod.kernel_function, tt) == 0 skip = skip_lookup + + # neither does waiting for the kernel, also when it takes long enough for the wait to + # involve a completion callback (except on 1.11, and to shield it from task cancellation + # on 1.14+) + skip_wait = VERSION < v"1.12" || isdefined(Base, :CANCEL_TOKEN) + A = zeros(Int, 4) + info = POCL.exception_info() + GC.@preserve info begin + info_ptr = Base.unsafe_convert(Ptr{POCL.ExceptionInfo_st}, info) + kernel = mod.clfunction(tt) + mod.launch_tuple(kernel, A, info_ptr) + @test all(==(1), A) + @test allocated(mod.launch_tuple, kernel, A, info_ptr) == 0 skip = skip_wait + end + B = ones(4) + kernel = mod.ka_kernel!(CPU(), 4) + for n in (1, 100_000) + mod.launch(kernel, B, n) + @test allocated(mod.launch, kernel, B, n) == 0 skip = skip_wait + end end @testset "POCL exceptions" begin