Conversation
Contributor
Benchmark ResultsShow table
Benchmark PlotsA plot of the benchmark results have been uploaded as an artifact to the workflow run for this PR. |
maleadt
removed this pull request from stack #802
September 28, 2026 10:14
maleadt
added this pull request to stack #804
September 28, 2026 10:15
christiangnrd
left a comment
christiangnrd
left a comment
Member
There was a problem hiding this comment.
A few comments.
Also, if it's not too much trouble, could you get claude to split off 7b0276c?
Member
Author
maleadt
removed this pull request from stack #804
September 29, 2026 20:22
maleadt
changed the base branch from
tb/ci-reverse
to
tb/pocl-spirvintrinsics-1.1.3
September 29, 2026 20:22
maleadt
added this pull request to stack #808
September 29, 2026 20:22
`GPU` didn't mean GPU hardware: the CPU backend is PoCL and subtyped it, so `CPU <: GPU`. Backends now subtype `Backend` directly. This is the first breaking change of KernelInterface 0.3.
`testsuite(backend, AT)` replaces `testsuite(Backend, name, mod, AT, DAT)`, whose name, module and device array type were unused, and tests the backend value it is given, e.g. one with non-default options. The events tests now skip backends without events of their own (whose `record_event` returns `nothing`), instead of them having to pass `skip_tests`.
…mgroups `KI.@kernel` prefixes a call and launches it, like `@cuda` or `@metal`, while `KA.@kernel` defines a kernel; sharing the name was a source of confusion. `numgroups` matches `get_num_groups` and `max_num_groups`. `KI.@launch` now also evaluates its backend expression once (it was evaluated once per argument), and passes the keywords it doesn't know to `kernel_function` as compiler options instead of rejecting them.
…dation `kernel_max_work_group_size` returned an occupancy recommendation on CUDA, AMDGPU and oneAPI, but the hard limit on Metal, OpenCL and PoCL: on an RTX 5080, 768 for a trivial kernel that can be launched with 1024 threads. It is replaced by two queries: - `max_work_group_size(kernel)`, the largest work-group the compiled kernel can be launched with; - `launch_configuration(kernel; nitems, max_work_group_size)`, the recommended work-group size for a launch of `nitems` work-items, which `ndrange` launches without a `workgroupsize` use. It falls back to the limit. The problem size is passed separately from the bound, so that heuristics like CUDA's `prefer_blocks` don't have to guess it. `max_work_group_dims` and `max_num_groups` are now required: their `typemax(Int)` fallbacks meant both "unknown" and "unlimited".
…ment KI.launch
Every backend repeated the launch validation and sizing in front of its native
launch, and the copies had drifted apart (return values, a `prod(ndrange) == 0`
check that can overflow). KernelInterface now implements the `Kernel` call and
passes validated 3-D sizes to one required method:
KI.launch(kernel, groups::Dims{3}, items::Dims{3}, args...; kwargs...)
The launch semantics are documented on `Kernel`: `ndrange` is rounded up to
whole work-groups and not masked, a zero in `ndrange` or `numgroups` launches
nothing, work-group sizes are checked against `max_work_group_dims` and
`max_work_group_size(kernel)`, and the number of work-items per dimension has
to fit an `Int`. Invalid launches throw an `ArgumentError` before the backend
sees them. Unknown keywords are passed on to `KI.launch`, so options like CUDA's
`stream` still reach the driver.
The `Kernel` call and PoCL's `launch` declare their arguments as `Vararg{Any, N}`:
Julia doesn't specialize a method on `args...` that it only passes through, which
made every launch dispatch dynamically (on CUDA, 1.2 µs and 1 kB more per
`KI.@launch`). The `launch` docstring recommends the same to backends.
`check_launch_args` and `auto_launch_sizes` are removed. The launch tests now
use different group counts and sizes per dimension, which exposes two tests
whose formulas only worked because they were equal.
The 0.2 docs said that typed queries are "computed in `T`" and undefined when the value doesn't fit, but backends implemented them as `T(x)`, a checked conversion that leaves a `throw_inexacterror` branch in every kernel. The result is now the exact value modulo `T`, as with `x % T`, which is cheap and testable: a `UInt8` query over 384 work-items has to wrap, where `T(x)` throws. Backends implement the four primitive queries (`get_local_id`, `get_group_id`, `get_local_size`, `get_num_groups`); `get_global_id` and `get_global_size` have fallbacks derived from them for backends without a builtin (CUDA, HIP), which backends with one (SPIR-V, Metal) should override.
Sub-groups were underspecified exactly where backends disagree: - Support is a capability: `supports_subgroups(backend)` and `supports_shuffle(backend, T)` replace `shfl_down_types`, whose non-empty result the testsuite used as the flag for sub-group support. - `sub_group_size(backend)` is a guarantee instead of "a reasonable size": kernels from `kernel_function` run with exactly that width, so host code can pick a `Val(N)` for a warp-level reduction. A backend that can't guarantee it reports no sub-group support. PoCL fixes the width when compiling. - Partial sub-groups follow OpenCL and SYCL: `get_sub_group_size()` counts the work-items that are present, and `get_num_sub_groups()` is `cld(items, width)`. - How work-items are assigned to sub-groups is unspecified, but each has a unique `(sub-group id, lane)` pair that doesn't change during the kernel. - A `shfl_down` from a lane that doesn't exist returns an unspecified value, and shuffles are not memory fences. - The sub-group queries take a result type like the index queries, defaulting to `Int`; they returned `UInt32` on most backends but `Int32` on CUDA. The testsuite checks partial sub-groups, lane uniqueness, `shfl_down` on every lane, and that `sub_group_barrier` makes memory writes visible.
- Execution is task-local: each task has an active device per backend and a queue on it, which host queries, compilation, allocations, copies and launches use. `kernel_function` has to store the backend value it was given (oneAPI dropped its compiler options, and OpenCL its platform, by constructing a new default), and a compiled kernel launched after switching devices either works or throws, but never runs on the wrong device. - New: `device(backend, A)`, the device that owns `A`. `device`, `device!`, `ndevices` and `device(backend, A)` are required for backends with more than one device; the single-device fallbacks now throw when `ndevices` reports more instead of answering for the wrong device. - Capability defaults are conservative: `supports_float64` and `supports_atomics` default to `false`, so that a missing method never claims support. `supports_atomics` means Atomix add and CAS on 32-bit integers and floats in global memory. - `copyto!(backend, dst, src)` copies in queue order and returns `dst`, and throws an `ArgumentError` for arrays of different lengths. CUDA's implementation copied `length(dst)` elements from `pointer(src)` without a check. - `localmemory` and `barrier` say what is shared and what becomes visible, and `unsafe_free!` is an optional hint with a legal no-op fallback. The testsuite covers devices, `copyto!`, local memory (two allocations don't alias; contents are visible after a barrier) and global-memory barriers.
- The public API is declared with `public` (Julia 1.11+); nothing is exported, so `KI.` prefixes the interface everywhere. - The docs get a contract table listing what a backend implements, which methods are optional and what their fallbacks do, and a versioning rule: required methods only change in breaking releases, optional methods can be added in any release if their fallback is conservative, and patch releases only add tests for behavior that was already specified. (0.2.3 broke backend CI in a patch release by adding tests for new obligations.) - The testsuite checks that a backend implements the methods without a fallback.
This was referenced Sep 30, 2026
Codecov Report❌ Patch coverage is
Additional details and impacted files@@ Coverage Diff @@
## main #800 +/- ##
==========================================
+ Coverage 67.70% 67.91% +0.21%
==========================================
Files 24 24
Lines 2031 2023 -8
==========================================
- Hits 1375 1374 -1
+ Misses 656 649 -7 ☔ 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
2 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.
KernelInterface (KI) is the layer that backends (CUDA.jl, Metal.jl, oneAPI.jl, AMDGPU.jl, OpenCL.jl, and PoCL in this repository) implement, and that KernelAbstractions builds on. Porting those backends to KI 0.2.3 showed that the interface left too much to the individual backends: every backend validated and sized launches itself, and the copies had drifted apart; the same query meant different things on different backends; and the index and sub-group queries were underspecified exactly where the platforms disagree. A kernel written against KI could therefore behave differently depending on where it ran.
KI 0.3 fixes that before the ports merge. It moves launch validation into KI, and writes down the semantics that backends have to implement, so that the testsuite can check them. The design doesn't change: kernels are plain Julia functions, device functions are overlay stubs, and a backend value comes with a task-local device and queue. This is a breaking release without deprecations, since KI has no users outside the backends yet. Each commit is one of the changes below.
KernelInterface owns the launch
Backends used to implement the whole
Kernelcall, including argument validation and work-group sizing. Now KI implements it, and passes validated 3-D sizes to one method a backend implements:The launch semantics are documented on
KI.Kernel: anndrangeis rounded up to whole work-groups and not masked, a zero anywhere launches nothing, and work-group sizes are checked against the device and the compiled kernel. Invalid launches throw before the driver sees them (here on CUDA):(The last one used to reach the backend as a negative global size.) Keywords that KI doesn't know are passed on to
KI.launch, so backend options like CUDA'sstreamorshmemstill reach the driver.Two renames come with this.
KI.@kernelis nowKI.@launch, since it launches a call like@cudadoes, whileKA.@kerneldefines a kernel.numworkgroupsis nownumgroups, matchingget_num_groupsandmax_num_groups.KI.@launchalso evaluates its backend expression once instead of once per argument, and passes keywords it doesn't know tokernel_functionas compiler options.The limit and the recommendation are separate queries
kernel_max_work_group_sizereturned an occupancy recommendation on CUDA, AMDGPU and oneAPI, and the hard limit on Metal, OpenCL and PoCL. On an RTX 5080 it returns 768 for a trivial kernel that can be launched with 1024 threads. Two queries replace it:max_work_group_size(kernel)is the largest work-group the compiled kernel can be launched with. Launches are validated against it.launch_configuration(kernel; nitems, max_work_group_size)is the recommended work-group size for a launch ofnitemswork-items.ndrangelaunches without aworkgroupsizeuse it, and it falls back to the limit. The problem size is passed separately from the bound, so that heuristics like CUDA'sprefer_blocks(more, smaller blocks) don't have to treat the bound as the problem size.max_work_group_dimsandmax_num_groupsare now required. Theirtypemax(Int)fallbacks meant both "unknown" and "unlimited", and #797 needs them to choose a launch.Typed index queries wrap
The index queries take a result type, e.g.
KI.get_global_id(Int32). The docs said the result is "computed inT" and undefined when it doesn't fit, but the backends implemented it asT(x): a checked conversion, which leaves athrow_inexacterrorbranch in every kernel that uses it. The result is now the exact value moduloT, as withx % T. That costs nothing, and it can be tested: aUInt8query over 384 work-items has to wrap, whereT(x)throws.Backends implement the four primitive queries (
get_local_id,get_group_id,get_local_size,get_num_groups).get_global_idandget_global_sizehave fallbacks derived from them for backends without a builtin (CUDA, HIP), and backends that have one (SPIR-V, Metal) should override them.Sub-groups
supports_subgroups(backend)andsupports_shuffle(backend, T)replaceshfl_down_types. The testsuite used to treat a non-empty type list as the flag for sub-group support.sub_group_size(backend)is a guarantee instead of "a reasonable size": kernels fromkernel_functionrun with exactly that width, so host code can pick aVal(N)for a warp-level reduction. A backend that can't guarantee it reports no sub-group support; PoCL now fixes the width at compile time.get_sub_group_size()counts the work-items that are present, andget_num_sub_groups()iscld(items, width). AMDGPU used÷, reporting 0 sub-groups for a 32-item group on wave64.(sub-group id, lane)pair that doesn't change during the kernel. AMDGPU usedactivelane(), which renumbers lanes under divergence.shfl_downfrom a lane that doesn't exist returns an unspecified value (CUDA returns the caller's own value, OpenCL leaves it undefined), and shuffles are not memory fences.Int. They returnedUInt32on most backends andInt32on CUDA.Execution, devices and memory
kernel_functionhas to store the backend value it was given, so that its options apply (oneAPI dropped its compiler options, and OpenCL its platform, by constructing a new default). A kernel launched after switching devices either works or throws; it never runs on the wrong device.device(backend, A), the device that ownsA. The device functions are required for backends with more than one device, and the single-device fallbacks now throw whenndevices > 1instead of answering for the wrong device.supports_float64andsupports_atomicsnow default tofalse, so that a missing method never claims support.copyto!(backend, dst, src)copies in queue order, returnsdst, and throws anArgumentErrorfor arrays of different lengths. CUDA's implementation copiedlength(dst)elements frompointer(src)without a check.localmemoryandbarrierdocument what memory is shared and which writes become visible.unsafe_free!is an optional hint with a no-op fallback.The contract
The docs now have a "Semantics" section and a contract table that lists, for each area, the required methods and the optional ones with their fallbacks. The public API is declared with
public(Julia 1.11+). There is also a versioning rule: required methods only change in breaking releases, optional methods can be added in any release if their fallback is conservative, and patch releases only add tests for behavior that was already specified. (0.2.3 broke backend CI in a patch release by adding unconditional tests for new obligations.)The testsuite's entry point is now
Testsuite.testsuite(backend, AT), which takes a backend value, so that backends with options can be tested. It checks that the methods without a fallback are implemented, and covers launch validation, non-square launches,copyto!, local memory, barriers, typed-index wrapping, partial sub-groups and shuffles. Two existing launch tests had formulas that only worked because the group count equaled the group size; the tests now use different sizes in every dimension.For KernelAbstractions users
KernelAbstractions.GPUis removed. It didn't mean GPU hardware: KA'sCPUbackend is PoCL and subtyped it, soCPU <: GPU. Backends subtypeKernelAbstractions.Backend, and code that dispatched on::GPUshould dispatch on::Backend, on concrete backend types, or on a capability query. This is in the 0.10 changelog.Backend ports
Each backend's KI PR has commits porting it to 0.3 (they get KI 0.3 from this branch through
[sources]until it's registered):Host launch overhead on CUDA (RTX 5080, 1024-element kernel, per launch) is unchanged, except that launches with an explicit size now also query the kernel's thread limit (about 36 ns):
@cudaKI.@launch, explicit sizesKI.@launch,ndrangeKernel, explicit sizesKernel,ndrangeKA's Buildkite jobs test the backends' KA 0.10 branches, which still require KI 0.2.3, so they can't be installed together with this PR and soft-fail until those branches are rebased onto the ported KI PRs. The ports above were tested against this branch by hand.
Not in this PR
Dynamic local memory, fences and barrier scopes,
max_local_memory, more collectives and sub-group size requests can all be added in a minor release; launch options like dynamic local memory can becomeKernelkeywords without changingKI.launch. #797 and #801, stacked on this PR, use the new contract to launch@kernelkernels on N-d grids with 32-bit indices, and to replace the launch code that every backend copies with one generic implementation in KA.