KernelInterface: per-dimension launch limits, inlined index queries - #798
Merged
Merged
Conversation
KernelAbstractions' index functions call the zero-argument forms, which weren't inlined on CUDA: every work-item made two calls, returning the ids through local memory, making e.g. broadcasts up to 2x slower.
Add `max_work_group_dims` and `max_num_groups`, and respect the former when distributing work-items over a workgroup, so that e.g. an ndrange of (1, 1, 5000) no longer gets a (1, 1, 1024) workgroup that CUDA can't launch. Also specify that the typed index queries compute in `T`. POCL reports its per-dimension workgroup limits, and the testsuite checks that back ends can launch at their reported limits and that automatically sized launches respect them.
Only CUDA.jl ran it so far. Skip the events tests, which need asynchronous launches.
On Julia 1.10, which ignores [sources], building the package already resolves the environment, which fails if KernelAbstractions requires a KernelInterface version that isn't registered yet.
maleadt
added this pull request to stack #799
September 27, 2026 20:17
vchuravy
approved these changes
Sep 27, 2026
Contributor
Benchmark ResultsShow table
Benchmark PlotsA plot of the benchmark results have been uploaded as an artifact to the workflow run for this PR. |
Member
|
I say we merge and tag and then fix the kernelinterface PRs |
christiangnrd
approved these changes
Sep 28, 2026
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
3 participants
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.
Two KernelInterface changes, split off from #797 so that they can be released as KernelInterface 0.2.3 first.
Per-dimension launch limits. KernelInterface only knows a back end's total workgroup size, so
auto_launch_sizescan pick workgroups the hardware can't launch: for an ndrange of(1, 1, 5000)it chooses(1, 1, 1024), while CUDA allows at most 64 threads along z. Back ends can now report their per-dimension workgroup and grid limits withmax_work_group_dims(backend)andmax_num_groups(backend)(unlimited by default), andthreads_to_workgroupsizerespects the former. KernelAbstractions needs these limits to decide whether it can launch a kernel on an N-d grid (#797). POCL reports its limits.Inlined index queries. The zero-argument
KI.get_local_id(),get_group_id()etc. weren't inlined on CUDA, so every work-item made real calls that returned the ids through local memory. With KAmain, that makes a 1-D broadcast in CUDA.jl's KA 0.10 port (JuliaGPU/CUDA.jl#3302) 1.65× slower than with KA 0.9, and some kernels up to 2.4× slower. They're now@inline.The typed queries (
get_group_id(Int32)etc.) are now documented to compute in the requested type.Tests. The KernelInterface testsuite, which back ends run, now checks that launches at the reported limits work and that automatically sized launches respect them. Back ends with a small limit along some dimension (CUDA: 64 along z) need to implement
max_work_group_dimsto pass. KA's CI now also runs the testsuite on POCL (so far only CUDA.jl ran it), except for the events tests, which need asynchronous launches.CI: Julia 1.10 ignores
[sources], so the workflow now developslib/KernelInterfacebefore building KernelAbstractions, which requires the unregistered 0.2.3.