Repository navigation
POCL: reduce launch memory allocations - #901
Merged
Merged
Conversation
giordano
added
performance
Contributor
Benchmark ResultsShow table
Benchmark PlotsA plot of the benchmark results have been uploaded as an artifact to the workflow run for this PR. |
…uple `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 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01D6fxP52AyNqfcKSSQXFxfB
…ration `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 <mose@gnu.org>
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 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01D6fxP52AyNqfcKSSQXFxfB
maleadt
force-pushed
the
mg/pocl-rm-launch-allocs
branch
from
October 17, 2026 09:50
6e895b4 to
e06876f
Compare
`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 <tim.besard@gmail.com> Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01D6fxP52AyNqfcKSSQXFxfB
maleadt
force-pushed
the
mg/pocl-rm-launch-allocs
branch
from
October 17, 2026 10:26
e06876f to
da38f0d
Compare
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## main #901 +/- ##
=======================================
Coverage 82.53% 82.53%
=======================================
Files 28 28
Lines 2456 2457 +1
=======================================
+ Hits 2027 2028 +1
Misses 429 429 ☔ View full report in Codecov by Harness. 🚀 New features to boost your workflow:
|
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Launching a compiled kernel on the POCL back-end allocated 96 bytes, and waiting for it another 112 bytes when the wait involved a completion callback, which depends on how long the kernel takes and on host load. This made allocation tests of CPU kernels, as in Oceananigans (CliMA/Oceananigans.jl#5799, for the KA 0.10 compatibility), fail on busy machines. With this PR and GPUToolbox 3.4.1 (JuliaGPU/GPUToolbox.jl#32, fixing JuliaGPU/GPUToolbox.jl#31), launches allocate nothing on Julia 1.12 and 1.13:
compiler_confighashed its keyword arguments asPairs, which boxes every value (16 bytes). Hashvalues(kwargs), like the other back-ends.KI.kernel_functionpassed the device's sub-group width toclfunction, asnothingon devices without sub-groups, so the keyword call was dynamic (64 bytes). The default width is now determined with the compiler configuration, as in OpenCL.jl, and not passed at all. This also makes@opencldefault to the device's width instead of 32 (which is the same on devices that support 32).checkis inlined, so thatRefs passed to API calls can live on the stack, and so is the call toclEnqueueNDRangeKernel, whoseReffor the event was allocated on every launch (16 bytes).cooperative_waitno longer allocates on the completion-callback path.Allocations per launch of the reproducer from JuliaGPU/GPUToolbox.jl#31:
The new tests check that
compiler_config,clfunction,KI.kernel_functionand launches of a short and a long kernel don't allocate. They're skipped where Julia still allocates: on 1.10, in GPUCompiler's lookup of compiled code and in waits on 1.11, and on 1.14, where GPUToolbox shields waits from task cancellation with a scoped value (and GPUCompiler allocates when looking up compiled code, see JuliaGPU/GPUCompiler.jl#1025).Bumps the version to 0.10.1.