Skip to content

Cut allocations when launching CPU kernels - #766

Closed
giordano wants to merge 3 commits into
JuliaGPU:mainfrom
giordano:mg/launch-allocations
Closed

giordano wants to merge 3 commits into
JuliaGPU:mainfrom
giordano:mg/launch-allocations

Conversation

@giordano

Copy link
Copy Markdown
Collaborator

Cutting down a warm CPU kernel launch from about 50 allocations to 1.

Where the allocations came from (allocation profiler on a warm two-argument launch, 52 allocations, 2208 bytes)

  • Julia does not specialize on args... that a method only splats onward. That applied to the Kernel{POCLBackend} call, call and set_args!, so every kernel argument (CompilerMetadata, device arrays, KernelState) was boxed on its way to clSetKernelArg, and set_args! dispatched dynamically over enumerate(args). About 30 allocations.
  • Per launch: two Vector{Csize_t} for the work sizes, a Ref for each argument passed to clSetKernelArg, an empty SVM pointer vector that no adaptor rule ever fills, a Base.Pairs from splatting keywords through the generated kernel call, and a clGetDeviceInfo query for the device's maximum work dimensions.
  • clfunction, on every launch: a compiler config, a method-instance lookup, a CompilerJob, GPUCompiler's cache query, Ref and Some temporaries, plus a Core.Box for a variable captured by the get! closure, all to return the same HostKernel as last time.

Result

Case allocations main → branch median main → branch
launch, static wg, dyn ndrange 48 → 1 348 → 347 μs
launch, static wg, static ndrange 39 → 2 349 → 347 μs
launch 3-D, dyn ndrange 48 → 1 350 → 348 μs
launch, dyn wg autotune 58 → 15 319 → 349 μs
throughput 1-D linear, 2^23 49 → 1 1.677 → 1.612 ms
throughput 3-D Cartesian 49 → 1 1.682 → 1.697 ms

The remaining allocation is the Ref receiving the event handle from clEnqueueNDRangeKernel. The autotune path keeps 15 because it queries the kernel's work-group info and partitions twice each launch; that can be cached per kernel if wanted. Launch latency itself did not move: the ~300 μs median for a 16-element kernel is POCL's enqueue and event wait, not Julia-side work, so the gain is GC pressure rather than wall time on this backend.

giordano and others added 2 commits September 11, 2026 05:38
A warm launch of a two-argument kernel on the CPU backend allocated about
50 times (2.2 KiB). Most of it came from Julia not specializing on
`args...` that a method only splats onward, which boxed every kernel
argument on the way from `Kernel` to `clSetKernelArg`; the rest from
per-launch vectors for the work sizes and SVM pointers, a keyword splat
through the generated kernel call, and a device property query.

- Annotate the splatted arguments as `Vararg{Any, N}` in the `Kernel`
  call, the generated `AbstractKernel` call, `call` and `set_args!`, and
  unroll `set_args!` by recursion instead of `enumerate`.
- Pass each kernel argument to `clSetKernelArg` as a `Ref{T}` `ccall`
  argument, which copies it to the stack.
- Pass the global and local work sizes as three-element tuples through a
  `ccall` taking `Ref{NTuple{3, Csize_t}}` when there is no offset or
  device RNG state, skipping the `max_work_item_dims` query.
- The generated kernel call takes `global_size` and `local_size` as
  explicit keywords.
- `clconvert` and `call` take `nothing` for the SVM pointers instead of
  allocating an empty vector; no adaptor rule collects such pointers.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01VHciC8x39gm97sABrSvBkt
Every launch went through `clfunction`, which built a compiler config,
looked up the method instance, allocated a `CompilerJob`, queried
GPUCompiler's cache and hashed its way into `_kernel_instances` before
returning the `HostKernel` it had returned last time. Keep the resolved
`HostKernel` per `HostKernel{F, tt}` type together with the world age and
context it was resolved in, and return it directly while no method has
been defined since and the context is unchanged. Reflection (a set
`compile_hook`) and explicit compiler keywords bypass the fast path.

Reading the `compile_hook` scoped value allocates, so it is only read
when a dynamic scope is active. The `Ref` used to find the kernel for the
current context and the `@something` in `compile_or_lookup` are replaced
with plain control flow, which also removes a boxed variable captured by
the `get!` closure.

A warm launch now allocates once, for the event handle.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01VHciC8x39gm97sABrSvBkt
Comment thread src/pocl/compiler/execution.jl Outdated
Comment on lines +207 to +212
# Reading a `ScopedValue` allocates; outside of any dynamic scope it holds its default.
@static if VERSION >= v"1.11"
@inline compile_hook_set() = Core.current_scope() !== nothing && GPUCompiler.compile_hook[] !== nothing
else
@inline compile_hook_set() = GPUCompiler.compile_hook[] !== nothing
end

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Uhm, I think this was resolved in some version of Julia, right? Don't remember which one offhand though, would need to check

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Probably only v1.14: JuliaLang/julia#61053

Reading a `ScopedValue` no longer allocates on Julia 1.14, and checking
`Core.current_scope()` first costs as much as the read itself there, so
the guard is limited to the versions that need it.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01VHciC8x39gm97sABrSvBkt
@vchuravy

Copy link
Copy Markdown
Member

Can we split this up? In particular the cache fast path needs a closer look

@giordano

Copy link
Copy Markdown
Collaborator Author

Can we split this up? In particular the cache fast path needs a closer look

They're separate commits already. Could open separate PRs (although I can't do stacks, as I can't push here)

@codecov

codecov Bot commented Sep 11, 2026

Copy link
Copy Markdown

Codecov Report

❌ Patch coverage is 95.83333% with 2 lines in your changes missing coverage. Please review.
✅ Project coverage is 64.04%. Comparing base (f68a281) to head (0afb29a).
⚠️ Report is 1 commits behind head on main.

Files with missing lines Patch % Lines
src/pocl/nanoOpenCL.jl 88.88% 2 Missing ⚠️
Additional details and impacted files
@@            Coverage Diff             @@
##             main     #766      +/-   ##
==========================================
+ Coverage   63.72%   64.04%   +0.31%     
==========================================
  Files          23       23              
  Lines        1935     1955      +20     
==========================================
+ Hits         1233     1252      +19     
- Misses        702      703       +1     

☔ View full report in Codecov by Harness.
📢 Have feedback on the report? Share it here.

🚀 New features to boost your workflow:
  • ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.

@vchuravy

Copy link
Copy Markdown
Member

as I can't push here)

You should now be able to

@giordano

Copy link
Copy Markdown
Collaborator Author

Closing in favour of #768

@giordano giordano closed this Sep 11, 2026
@giordano
giordano deleted the mg/launch-allocations branch September 11, 2026 22:19
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants