Skip to content

Implement KernelInterface, and support KernelAbstractions 0.10 - #653

Open
maleadt wants to merge 9 commits into
mainfrom
ka-0.10
Open

maleadt wants to merge 9 commits into
mainfrom
ka-0.10

Conversation

@maleadt

@maleadt maleadt commented Sep 30, 2026 •

Copy link
Copy Markdown
Member

KernelAbstractions 0.10 moves its back-end API into a separate package, KernelInterface. A back end implements KernelInterface: allocating and copying memory, selecting devices, compiling and launching kernels, and the device-side intrinsics. KernelAbstractions then launches @kernel kernels the same way on every back end (JuliaGPU/KernelAbstractions.jl#801). This PR ports oneAPI.jl to that model in one step, the way JuliaGPU/CUDA.jl#3314 does for CUDA. It supersedes #624, which added KernelInterface 0.3 next to the KernelAbstractions 0.9 back end, and the intrinsics branch, which ported an earlier KernelInterface to KernelAbstractions 0.10.

oneAPIBackend now implements KernelInterface, and oneAPI.jl depends on KernelInterface instead of KernelAbstractions. That makes the KernelInterface layer usable on its own, without KernelAbstractions' macros:

using oneAPI
import KernelInterface as KI

function vadd(c, a, b)
    i = KI.get_global_id().x
    if i <= length(c)
        @inbounds c[i] = a[i] + b[i]
    end
    return
end

a = oneArray(rand(Float32, 1000)); b = oneArray(rand(Float32, 1000)); c = similar(a)
KI.@launch oneAPIBackend() ndrange=length(c) vadd(c, a, b)

KernelAbstractions becomes a weak dependency. Its extension only provides what KernelAbstractions still needs from a back end: the MArray behind @private, the Adapt rule behind @Const, and adapting arrays to KernelAbstractions' CPU back end. oneAPI's own copy of the KernelAbstractions launch path goes away: partitioning the ndrange, building the kernel's context, tuning the work-group size, and calling the kernel. oneAPIBackend(; prefer_blocks=true) still spreads small launches over more work-groups; that now happens in KI.launch_configuration, which receives the number of work-items to cover. There is no equivalent of CUDA's maxthreads hint for kernels with a static work-group size, since oneAPI's compiler has no such option.

Kernels behave the same, but they are launched differently. KernelAbstractions now launches them on a 3-D grid, and computes @index in 32 bits when the iteration space fits (JuliaGPU/KernelAbstractions.jl#797), so a kernel over a 3-D ndrange doesn't need divisions to compute its index; the test suite checks that in the generated code. With a KernelAbstractions kernel that copies a Float32 array, on the Iris Xe of a Tiger Lake laptop (median of 7-8 runs, each the minimum of 200 host-timed launches; this GPU is memory-bound and its timings vary by about 5% between runs):

main (KA 0.9) this PR
256×256×256, @index(Global, Cartesian) 6.99 ms 7.01 ms
256×256×256, @index(Global, Linear) 6.97 ms 7.00 ms
255×257×129, @index(Global, Cartesian) 4.05 ms 4.29 ms
255×257×129, @index(Global, Linear) 4.05 ms 4.29 ms
launch of an empty kernel 18.4 µs, 2720 bytes 11.0 µs, 1600 bytes
launch of a kernel with 40 arguments 42.2 µs, 46336 bytes 36.2 µs, 25392 bytes

Unlike on CUDA, the index computation hardly matters on this GPU: the kernel times are within the noise of main. For the odd-sized range, both main and this PR are about 20% slower than the same copy launched over ndrange=length(A), independent of the work-group shape; JuliaGPU/KernelAbstractions.jl#814 tracks that. The launches themselves got cheaper.

KI.launch rejects items and groups, which would override the launch geometry KernelInterface has validated, but passes oneAPI's other launch options on, and refuses to launch a kernel on another context or device than the one it was compiled for:

kernel = KI.@launch oneAPIBackend() launch=false vadd(c, a, b)
kernel(c, a, b; ndrange=length(c), queue=global_stream(context(), device()))   # fine
kernel(c, a, b; ndrange=length(c), items=256)                                  # ArgumentError

KernelInterface 0.4 passes the kernel arguments to the back end as a tuple (JuliaGPU/KernelAbstractions.jl#811). oneAPI has no launch path that takes a tuple, like CUDA got in JuliaGPU/CUDA.jl#3309, so KI.launch splats them into the HostKernel call for now; the launch overhead above includes that.

The first commits prepare for the port: versioninfo only lists KernelAbstractions when it is loaded, a kernel's maximum group size is cached so that launches don't query it, and synchronize becomes cooperative, as KernelInterface requires: it busy-waits briefly and then blocks in the driver on a separate thread, so other tasks keep running (synchronize(; blocking=true) restores the old behavior). Then comes the port itself. The commits after it implement parts of KernelInterface that are optional:

  • @oneapi sub_group_size=N compiles a kernel for a given sub-group size. KernelInterface kernels are compiled for a fixed one, 32 if the device supports it, which enables KernelInterface's sub-group queries, barrier and shfl_down;
  • KI.record_event and KI.wait_event, which KernelAbstractions.@spawn uses to order a new task's work after the work its parent had queued, without synchronizing the parent. The new task waits for the event on the host (cooperatively), not on the device, because a Level Zero event must not be destroyed while a command list still references it;
  • KI.versioninfo, which prints oneAPI.versioninfo().

Some limitations remain. KernelAbstractions converts the arguments twice, once to determine the argument types to compile for and once when launching; kernel_convert is supposed to be pure, so this only shows with a conversion that has side effects. KernelAbstractions' new Random tests are skipped, as oneAPI doesn't support Random's default RNG in kernels yet. With ONEAPI_SYNC_EACH_SUBMISSION (the Aurora LTS workaround), the wait after each kernel launch is cooperative, but the one after other submissions, such as copies, still blocks the thread.

This needs a breaking release, since it requires KernelAbstractions 0.10 and KernelInterface 0.4, which aren't registered yet, and an AcceleratedKernels release that supports them. Until then, the last commit takes all three from their development branches through [sources]. Because Pkg doesn't accept [sources] for a weak dependency, and AcceleratedKernels pulls KernelAbstractions in anyway, that commit also makes KernelAbstractions a regular dependency again. Julia 1.10 and 1.11 don't pick up [sources] for the tests, so CI develops the packages explicitly there, and the documentation is built with Julia 1.12. That commit is dropped before merging.

christiangnrd and others added 8 commits September 30, 2026 17:59
In preparation for KernelAbstractions becoming a weak dependency.
Like the spill size, cache the kernel's `maxGroupSize` in the `ZeKernel`, so
that `launch_configuration` (and the KernelInterface launch validation, which
checks every explicit work-group size against it) don't query the kernel
properties, which allocates, on every launch.
`synchronize` blocked the calling thread in the driver until the work had
completed, so no other task could run on it in the meantime, and no other
thread could run the GC either. KernelInterface requires `synchronize` to
be cooperative.

Wait like CUDA.jl does: busy-wait on a non-blocking query first, which keeps
the latency of short operations low, and then block in the driver (GC-safe)
on one of a few dedicated threads, while the calling task waits for it
without blocking the scheduler.

This applies to `synchronize()` and `synchronize(::oneStream)`, i.e., to user
code, KernelAbstractions, KernelInterface and the synchronizing copies;
`synchronize(; blocking=true)` restores the old behavior. The command list
and queue methods, which also run from finalizers, keep blocking. With
`ONEAPI_SYNC_EACH_SUBMISSION` set, as on Aurora's LTS stack, the wait after
every launch is cooperative too; otherwise it would block the thread, and
leave `synchronize` nothing to wait for.
KernelAbstractions 0.10 builds on KernelInterface, which defines what a back
end provides: memory and device management, compiling and launching kernels,
and the device-side intrinsics. KernelAbstractions then launches `@kernel`
kernels itself on any KernelInterface back end, which replaces oneAPI's copy
of that launch path: partitioning the ndrange, building the kernel's context,
and tuning the work-group size.

`oneAPIBackend` now implements KernelInterface, and oneAPI depends on it
instead of on KernelAbstractions. What KernelAbstractions still needs from a
back end moves to an extension: the `MArray` behind `@private`, the Adapt
rule for `@Const`, and adapting to KernelAbstractions' CPU back end.
`prefer_blocks` now applies in `KI.launch_configuration`, which receives the
number of work-items.

The index queries are computed in the requested type with `% T`, which needs
SPIRVIntrinsics 1.1.3 to truncate the 3-D built-ins without producing illegal
vector types. `KI.launch` rejects `items` and `groups`, which would override
the launch geometry that KernelInterface validated, and refuses to launch a
kernel on another context or device than the one it was compiled for.
`KI.kernel_function` converts the callable to compile it, and the kernel keeps
the original alive, as the converted form of a closure only holds pointers to
the arrays it captures.

Co-authored-by: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
The `sub_group_size` compiler keyword sets the `intel_reqd_sub_group_size`
metadata of the kernel. By default, the compiler still chooses the sub-group
size, and a size the device doesn't support is an error.
Implement the sub-group queries, `sub_group_barrier` and `shfl_down`, and
report their support to KernelInterface, which requires kernels to execute
with a fixed sub-group width: `kernel_function` compiles for 32 work-items,
like a CUDA warp, if the device supports that width, and for its widest one
otherwise. Shuffles of `Float16` and `Float64` values depend on the device
supporting those types.

KernelInterface leaves unspecified how work-items are grouped into
sub-groups; Intel GPUs form them from consecutive linear work-item indices,
so the last sub-group of a work-group can be partial, which a test checks.

Co-authored-by: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Implement `KI.record_event` with an event signaled on the task's stream,
and `KI.wait_event` by waiting for it on the host, cooperatively, like
`synchronize` does. `KernelAbstractions.@spawn` uses them to order a new
task's work after the work its parent had queued, without synchronizing the
parent; the default is a full synchronization.

The wait happens on the host, not by making the waiting task's stream wait
on the device, because a Level Zero event must not be destroyed while a
command list still references it, and the stream would reference it until it
executes the wait. For the same reason, recorded events are kept alive
until they have been signaled, independently of the recording task, which
may end before that.

KernelInterface's testsuite checks them once `record_event` returns an event.

Co-authored-by: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
`KI.versioninfo(oneAPIBackend())` prints `oneAPI.versioninfo()`.
@github-actions

Copy link
Copy Markdown
Contributor

Your PR requires formatting changes to meet the project's style guidelines.
Please consider running Runic (git runic main) to apply these changes.

Click here to view the suggested changes.
diff --git a/lib/level-zero/event.jl b/lib/level-zero/event.jl
index d805cf1..be21064 100644
--- a/lib/level-zero/event.jl
+++ b/lib/level-zero/event.jl
@@ -36,8 +36,8 @@ mutable struct ZeEvent
     handle::ze_event_handle_t
     pool::ZeEventPool
 
-    function ZeEvent(pool, index::Integer; signal=0, wait=0)
-        desc_ref = Ref(ze_event_desc_t(; index=index-1, signal, wait))
+    function ZeEvent(pool, index::Integer; signal = 0, wait = 0)
+        desc_ref = Ref(ze_event_desc_t(; index = index - 1, signal, wait))
         handle_ref = Ref{ze_event_handle_t}()
         zeEventCreate(pool, desc_ref, handle_ref)
         obj = new(handle_ref[], pool)
diff --git a/lib/level-zero/synchronization.jl b/lib/level-zero/synchronization.jl
index d5ab74b..e75e6a0 100644
--- a/lib/level-zero/synchronization.jl
+++ b/lib/level-zero/synchronization.jl
@@ -28,13 +28,16 @@ Base.isdone(queue::ZeCommandQueue) =
 # the blocking synchronization, marked GC-safe so that it doesn't keep the GC from running
 gcsafe_synchronize(list::ZeImmediateCommandList) =
     @gcsafe_ccall libze_loader.zeCommandListHostSynchronize(
-        list::ze_command_list_handle_t, typemax(UInt64)::UInt64)::ze_result_t
+    list::ze_command_list_handle_t, typemax(UInt64)::UInt64
+)::ze_result_t
 gcsafe_synchronize(queue::ZeCommandQueue) =
     @gcsafe_ccall libze_loader.zeCommandQueueSynchronize(
-        queue::ze_command_queue_handle_t, typemax(UInt64)::UInt64)::ze_result_t
+    queue::ze_command_queue_handle_t, typemax(UInt64)::UInt64
+)::ze_result_t
 gcsafe_synchronize(event::ZeEvent) =
     @gcsafe_ccall libze_loader.zeEventHostSynchronize(
-        event::ze_event_handle_t, typemax(UInt64)::UInt64)::ze_result_t
+    event::ze_event_handle_t, typemax(UInt64)::UInt64
+)::ze_result_t
 
 # once the work has completed, synchronizing a list or queue doesn't block anymore; a
 # signaled event needs nothing more
@@ -46,12 +49,12 @@ finish_synchronization(::ZeEvent) = nothing
 
 # custom, unbuffered channel that supports returning a value to the sender
 # without the need for a second channel
-struct BidirectionalChannel{I,O} <: AbstractChannel{I}
+struct BidirectionalChannel{I, O} <: AbstractChannel{I}
     cond_take::Threads.Condition                 # waiting for data to become available
     cond_put::Threads.Condition                  # waiting for a writeable slot
     cond_ret::Threads.Condition                  # waiting for a data to be returned
 
-    function BidirectionalChannel{I,O}() where {I,O}
+    function BidirectionalChannel{I, O}() where {I, O}
         lock = ReentrantLock()
         cond_put = Threads.Condition(lock)
         cond_take = Threads.Condition(lock)
@@ -61,9 +64,9 @@ struct BidirectionalChannel{I,O} <: AbstractChannel{I}
 end
 
 Base.put!(c::BidirectionalChannel{I}, v) where {I} = put!(c, convert(I, v))
-function Base.put!(c::BidirectionalChannel{I,O}, v::I) where {I,O}
+function Base.put!(c::BidirectionalChannel{I, O}, v::I) where {I, O}
     lock(c)
-    try
+    return try
         # wait for a slot to be available
         while isempty(c.cond_take)
             Base.wait(c.cond_put)
@@ -79,9 +82,9 @@ function Base.put!(c::BidirectionalChannel{I,O}, v::I) where {I,O}
     end
 end
 
-function Base.take!(f::Base.Callable, c::BidirectionalChannel{I,O}) where {I,O}
+function Base.take!(f::Base.Callable, c::BidirectionalChannel{I, O}) where {I, O}
     lock(c)
-    try
+    return try
         # notify the producer that we're ready to accept a value
         notify(c.cond_put, nothing, false, false)
 
@@ -131,7 +134,7 @@ end
 ## slow path: synchronize on a separate thread
 
 const MAX_SYNC_THREADS = 4
-const sync_channels = Array{BidirectionalChannel{SyncObject,ze_result_t}}(undef, MAX_SYNC_THREADS)
+const sync_channels = Array{BidirectionalChannel{SyncObject, ze_result_t}}(undef, MAX_SYNC_THREADS)
 const sync_channel_cursor = Threads.Atomic{UInt32}(1)
 const sync_channel_lock = Base.ReentrantLock()
 
@@ -143,6 +146,7 @@ function synchronization_worker(data)
         # wait for work
         take!(gcsafe_synchronize, chan)
     end
+    return
 end
 
 @noinline function create_synchronization_worker(i)
@@ -154,7 +158,7 @@ end
 
         # should be safe to assign before threads are running;
         # any user will just submit work that makes it block
-        sync_channels[i] = BidirectionalChannel{SyncObject,ze_result_t}()
+        sync_channels[i] = BidirectionalChannel{SyncObject, ze_result_t}()
 
         # we don't know what the size of uv_thread_t is, so reserve enough space
         tid = Ref{NTuple{32, UInt8}}(ntuple(i -> 0, 32))
diff --git a/src/compiler/compilation.jl b/src/compiler/compilation.jl
index a6a78b0..931079c 100644
--- a/src/compiler/compilation.jl
+++ b/src/compiler/compilation.jl
@@ -1,7 +1,7 @@
 ## gpucompiler interface implementation
 
 Base.@kwdef struct oneAPICompilerParams <: AbstractCompilerParams
-    sub_group_size::Union{Nothing,Int} = nothing
+    sub_group_size::Union{Nothing, Int} = nothing
 end
 
 const oneAPICompilerConfig = CompilerConfig{SPIRVCompilerTarget, oneAPICompilerParams}
@@ -281,8 +281,10 @@ function _driver_supports_bfloat16_spirv(dev=device())
     end
 end
 
-@noinline function _compiler_config(dev; kernel=true, name=nothing, always_inline=false,
-                                    sub_group_size=nothing, kwargs...)
+@noinline function _compiler_config(
+        dev; kernel = true, name = nothing, always_inline = false,
+        sub_group_size = nothing, kwargs...
+    )
     properties = oneL0.module_properties(dev)
     supports_fp16 = properties.fp16flags & oneL0.ZE_DEVICE_MODULE_FLAG_FP16 == oneL0.ZE_DEVICE_MODULE_FLAG_FP16
     supports_fp64 = properties.fp64flags & oneL0.ZE_DEVICE_MODULE_FLAG_FP64 == oneL0.ZE_DEVICE_MODULE_FLAG_FP64
diff --git a/src/context.jl b/src/context.jl
index 521adb3..282de5c 100644
--- a/src/context.jl
+++ b/src/context.jl
@@ -401,7 +401,7 @@ println("GPU work completed")
 
 See also: [`global_stream`](@ref), [`context`](@ref), [`device`](@ref)
 """
-function oneL0.synchronize(s::oneStream; blocking::Bool=false)
+function oneL0.synchronize(s::oneStream; blocking::Bool = false)
     sync = blocking ? oneL0.synchronize : oneL0.nonblocking_synchronize
     sync(s.list)
     q = s.queue
@@ -412,8 +412,8 @@ function oneL0.synchronize(s::oneStream; blocking::Bool=false)
     return
 end
 
-function oneL0.synchronize(; blocking::Bool=false)
-    oneL0.synchronize(global_stream(context(), device()); blocking)
+function oneL0.synchronize(; blocking::Bool = false)
+    return oneL0.synchronize(global_stream(context(), device()); blocking)
 end
 
 # Julia → MKL ordering: everything Julia appended to the task's immediate list must be
diff --git a/src/oneAPIKernels.jl b/src/oneAPIKernels.jl
index 8bc28db..b0b1254 100644
--- a/src/oneAPIKernels.jl
+++ b/src/oneAPIKernels.jl
@@ -98,7 +98,7 @@ struct oneAPIKernel{F, H <: oneAPI.HostKernel}
     host::H
 end
 
-function KI.kernel_function(backend::oneAPIBackend, f::F, tt::TT=Tuple{}; name = nothing, kwargs...) where {F,TT}
+function KI.kernel_function(backend::oneAPIBackend, f::F, tt::TT = Tuple{}; name = nothing, kwargs...) where {F, TT}
     # compile for the sub-group width that `KI.sub_group_size` promises
     sub_group_size = KI.sub_group_size(backend)
     if get(kwargs, :sub_group_size, sub_group_size) != sub_group_size
@@ -110,12 +110,14 @@ function KI.kernel_function(backend::oneAPIBackend, f::F, tt::TT=Tuple{}; name =
         zefunction(kernel_convert(f), tt; name, backend.always_inline, kwargs...)
     end
     kern = oneAPIKernel(f, host)
-    KI.Kernel{oneAPIBackend, typeof(kern)}(backend, kern)
+    return KI.Kernel{oneAPIBackend, typeof(kern)}(backend, kern)
 end
 
 # XXX: calling a `HostKernel` takes the arguments as varargs, so this splats them
-function KI.launch(obj::KI.Kernel{oneAPIBackend}, groups::Dims{3}, items::Dims{3},
-                   args::Tuple; kwargs...)
+function KI.launch(
+        obj::KI.Kernel{oneAPIBackend}, groups::Dims{3}, items::Dims{3},
+        args::Tuple; kwargs...
+    )
     # KernelInterface has validated the launch geometry
     if haskey(kwargs, :items) || haskey(kwargs, :groups)
         throw(ArgumentError("KernelInterface kernels take `numgroups`, `workgroupsize` or `ndrange`, not `items` or `groups`"))
@@ -171,19 +173,21 @@ function device_limits(dev::oneAPI.oneL0.ZeDevice = device())
     limits = get!(task_local_storage(), :oneAPIDeviceLimits) do
         Dict{oneAPI.oneL0.ZeDevice, DeviceLimits}()
     end::Dict{oneAPI.oneL0.ZeDevice, DeviceLimits}
-    get!(limits, dev) do
+    return get!(limits, dev) do
         props = oneAPI.oneL0.compute_properties(dev)
         module_props = oneAPI.oneL0.module_properties(dev)
         # the sub-group width that `kernel_function` compiles for: 32, like a CUDA warp, if
         # the device supports it, and 0 if the device has no sub-groups
         sg_sizes = props.subGroupSizes
         sub_group_size = 32 in sg_sizes ? 32 : maximum(sg_sizes; init = 0)
-        (; max_work_group_size = props.maxTotalGroupSize,
-           max_work_group_dims = (props.maxGroupSizeX, props.maxGroupSizeY, props.maxGroupSizeZ),
-           max_num_groups = (props.maxGroupCountX, props.maxGroupCountY, props.maxGroupCountZ),
-           sub_group_size,
-           supports_float16 = module_props.flags & oneAPI.oneL0.ZE_DEVICE_MODULE_FLAG_FP16 != 0,
-           supports_float64 = module_props.flags & oneAPI.oneL0.ZE_DEVICE_MODULE_FLAG_FP64 != 0)
+        (;
+            max_work_group_size = props.maxTotalGroupSize,
+            max_work_group_dims = (props.maxGroupSizeX, props.maxGroupSizeY, props.maxGroupSizeZ),
+            max_num_groups = (props.maxGroupCountX, props.maxGroupCountY, props.maxGroupCountZ),
+            sub_group_size,
+            supports_float16 = module_props.flags & oneAPI.oneL0.ZE_DEVICE_MODULE_FLAG_FP16 != 0,
+            supports_float64 = module_props.flags & oneAPI.oneL0.ZE_DEVICE_MODULE_FLAG_FP64 != 0,
+        )
     end
 end
 KI.max_work_group_size(::oneAPIBackend)::Int = device_limits().max_work_group_size
@@ -253,7 +257,7 @@ end
     sub_group_barrier(SPIRVIntrinsics.LOCAL_MEM_FENCE | SPIRVIntrinsics.GLOBAL_MEM_FENCE)
 end
 
-@device_override function KI.shfl_down(val::T, offset::Integer) where T
+@device_override function KI.shfl_down(val::T, offset::Integer) where {T}
     sub_group_shuffle(val, get_sub_group_local_id() + offset)
 end
 
@@ -282,11 +286,15 @@ function KI.record_event(::oneAPIBackend)
     ctx = oneAPI.context()
     dev = oneAPI.device()
 
-    pool = oneAPI.oneL0.ZeEventPool(ctx, 1;
-                                    flags = oneAPI.oneL0.ZE_EVENT_POOL_FLAG_HOST_VISIBLE)
-    ev = oneAPI.oneL0.ZeEvent(pool, 1;
-                              signal = oneAPI.oneL0.ZE_EVENT_SCOPE_FLAG_HOST,
-                              wait = oneAPI.oneL0.ZE_EVENT_SCOPE_FLAG_HOST)
+    pool = oneAPI.oneL0.ZeEventPool(
+        ctx, 1;
+        flags = oneAPI.oneL0.ZE_EVENT_POOL_FLAG_HOST_VISIBLE
+    )
+    ev = oneAPI.oneL0.ZeEvent(
+        pool, 1;
+        signal = oneAPI.oneL0.ZE_EVENT_SCOPE_FLAG_HOST,
+        wait = oneAPI.oneL0.ZE_EVENT_SCOPE_FLAG_HOST
+    )
 
     @lock pending_events_lock begin
         filter!(!Base.isdone, pending_events)
diff --git a/src/utils.jl b/src/utils.jl
index 05bceb8..ce3dcac 100644
--- a/src/utils.jl
+++ b/src/utils.jl
@@ -31,13 +31,15 @@ function versioninfo(io::IO=stdout)
     get_module(name::Symbol) = (name, getfield(oneAPI, name))
     function get_module(pkg::Tuple{String, String})
         id = Base.PkgId(Base.UUID(pkg[1]), pkg[2])
-        (pkg[2], get(Base.loaded_modules, id, nothing))
+        return (pkg[2], get(Base.loaded_modules, id, nothing))
     end
 
     println(io, "Julia packages:")
     println(io, "- oneAPI.jl: $(Base.pkgversion(oneAPI))")
-    for pkg in [:GPUArrays, :GPUCompiler, ("63c18a36-062a-441e-b654-da1e3ab1ce7c", "KernelAbstractions"),
-                 :KernelInterface, :LLVM, :SPIRVIntrinsics]
+    for pkg in [
+            :GPUArrays, :GPUCompiler, ("63c18a36-062a-441e-b654-da1e3ab1ce7c", "KernelAbstractions"),
+            :KernelInterface, :LLVM, :SPIRVIntrinsics,
+        ]
         name, mod = get_module(pkg)
         isnothing(mod) || println(io, "- $(name): $(Base.pkgversion(mod))")
     end
diff --git a/test/execution.jl b/test/execution.jl
index 2f943d7..f559365 100644
--- a/test/execution.jl
+++ b/test/execution.jl
@@ -41,20 +41,20 @@ end
 end
 
 
-function store_max_sub_group_size(a)
-    @inbounds a[1] = get_max_sub_group_size()
-    return
-end
+    function store_max_sub_group_size(a)
+        @inbounds a[1] = get_max_sub_group_size()
+        return
+    end
 
-@testset "sub-group size" begin
-    a = oneArray{UInt32}(undef, 1)
-    sizes = oneL0.compute_properties(device()).subGroupSizes
-    for sub_group_size in sizes
-        @oneapi items=64 sub_group_size store_max_sub_group_size(a)
-        @test Array(a)[1] == sub_group_size
+    @testset "sub-group size" begin
+        a = oneArray{UInt32}(undef, 1)
+        sizes = oneL0.compute_properties(device()).subGroupSizes
+        for sub_group_size in sizes
+            @oneapi items = 64 sub_group_size store_max_sub_group_size(a)
+            @test Array(a)[1] == sub_group_size
+        end
+        @test_throws ArgumentError @oneapi sub_group_size = maximum(sizes) + 1 dummy()
     end
-    @test_throws ArgumentError @oneapi sub_group_size=maximum(sizes)+1 dummy()
-end
 
 
 @testset "inference" begin
@@ -772,7 +772,7 @@ end
 
 @testset "cooperative synchronize" begin
     a = oneArray{UInt32}(undef, 64)
-    slow(iters) = @oneapi items=64 slow_kernel(a, UInt32(iters))
+    slow(iters) = @oneapi items = 64 slow_kernel(a, UInt32(iters))
     slow(1)
     synchronize()
     # warm up the slow path of `synchronize`: compiling it would end the calibration early
@@ -801,7 +801,7 @@ end
 
     # blocking synchronization is still available
     slow(1)
-    @test synchronize(; blocking=true) === nothing
+    @test synchronize(; blocking = true) === nothing
 end
 
 ############################################################################################
diff --git a/test/kernelabstractions.jl b/test/kernelabstractions.jl
index e3f3b56..68fa9dd 100644
--- a/test/kernelabstractions.jl
+++ b/test/kernelabstractions.jl
@@ -6,12 +6,12 @@ include(joinpath(dirname(pathof(KernelAbstractions)), "..", "test", "testsuite.j
 skip_tests=Set([
     "sparse",
     "Convert", # Need to opt out of i128
-    "Random", # oneAPI doesn't support Random's default RNG in kernels yet
-    # these run kernels on KernelAbstractions' POCL-based CPU back-end
-    "CPU synchronization",
-    "fallback test: callable types",
+        "Random", # oneAPI doesn't support Random's default RNG in kernels yet
+        # these run kernels on KernelAbstractions' POCL-based CPU back-end
+        "CPU synchronization",
+        "fallback test: callable types",
 ])
-Testsuite.testsuite(()->oneAPIBackend(), "oneAPI", oneAPI, oneArray, oneDeviceArray; skip_tests)
+Testsuite.testsuite(() -> oneAPIBackend(), "oneAPI", oneAPI, oneArray, oneDeviceArray; skip_tests)
 
 KA.@kernel function store_global_linear!(A)
     I = KA.@index(Global, Linear)
@@ -28,7 +28,7 @@ end
 
 @testset "launch configuration" begin
     backend = oneAPIBackend()
-    function select(kernel, ndrange, workgroupsize=nothing)
+    function select(kernel, ndrange, workgroupsize = nothing)
         ndrange, workgroupsize, iterspace, _ = KA.launch_config(kernel, ndrange, workgroupsize)
         KA.select_launch(kernel, workgroupsize, iterspace)
     end
@@ -40,20 +40,22 @@ end
 
     # which doesn't need divisions to compute the index of a dynamic N-d range
     A = oneAPI.zeros(Int, 64, 32, 16)
-    ir = sprint(io -> oneAPI.@device_code_llvm io=io kernel(A; ndrange=size(A)))
+    ir = sprint(io -> oneAPI.@device_code_llvm io = io kernel(A; ndrange = size(A)))
     @test !occursin(r"\b[us](div|rem) ", ir)
     @test Array(A) == LinearIndices(A)
 
     # tuning for more work-groups (the testsuite above uses the default)
-    Testsuite.launch_testsuite(()->oneAPIBackend(; prefer_blocks=true), oneArray)
+    Testsuite.launch_testsuite(() -> oneAPIBackend(; prefer_blocks = true), oneArray)
 
     # iteration spaces that don't fit 32 bits use 64-bit indices
     kernel = store_last_index!(backend)
     A = oneAPI.zeros(Int, 2)
-    for (dims, launch) in (((2^16 + 1, 2^15), KA.NDLaunch{Int}()),
-                           ((2^11 + 1, 2^10, 2^10, 1), KA.LinearLaunch{Int}()))
+    for (dims, launch) in (
+            ((2^16 + 1, 2^15), KA.NDLaunch{Int}()),
+            ((2^11 + 1, 2^10, 2^10, 1), KA.LinearLaunch{Int}()),
+        )
         @test select(kernel, dims) === launch
-        kernel(A; ndrange=dims)
+        kernel(A; ndrange = dims)
         @test Array(A) == [prod(dims), dims[2]]
     end
 end
@@ -69,14 +71,14 @@ end
 @testset "tuning" begin
     # tuning receives the number of work-items; prefer_blocks launches more, smaller groups
     A = oneAPI.zeros(Int, 1024)
-    kernel = KI.@launch oneAPIBackend() launch=false ki_store_index!(A)
-    items = KI.launch_configuration(kernel; nitems=1024).workgroupsize
+    kernel = KI.@launch oneAPIBackend() launch = false ki_store_index!(A)
+    items = KI.launch_configuration(kernel; nitems = 1024).workgroupsize
     @test items <= KI.max_work_group_size(kernel)
-    kernel = KI.@launch oneAPIBackend(; prefer_blocks=true) launch=false ki_store_index!(A)
-    fewer = KI.launch_configuration(kernel; nitems=1024).workgroupsize
+    kernel = KI.@launch oneAPIBackend(; prefer_blocks = true) launch = false ki_store_index!(A)
+    fewer = KI.launch_configuration(kernel; nitems = 1024).workgroupsize
     @test fewer < items
     # ... but not for a bound on the work-group size alone
-    @test KI.launch_configuration(kernel; max_work_group_size=1024).workgroupsize == items
-    KI.@launch oneAPIBackend(; prefer_blocks=true) ndrange=length(A) ki_store_index!(A)
+    @test KI.launch_configuration(kernel; max_work_group_size = 1024).workgroupsize == items
+    KI.@launch oneAPIBackend(; prefer_blocks = true) ndrange = length(A) ki_store_index!(A)
     @test Array(A) == 1:1024
 end
diff --git a/test/kernelinterface.jl b/test/kernelinterface.jl
index 27c0194..c42c9c9 100644
--- a/test/kernelinterface.jl
+++ b/test/kernelinterface.jl
@@ -30,15 +30,15 @@ end
     workgroupsize = (width + 1, 2)
     n = prod(workgroupsize)
     num, sizes, max, id, lane = (oneArray{UInt32}(undef, n) for _ in 1:5)
-    KI.@launch backend workgroupsize=workgroupsize ki_subgroup_kernel(num, sizes, max, id, lane)
+    KI.@launch backend workgroupsize = workgroupsize ki_subgroup_kernel(num, sizes, max, id, lane)
     @test all(==(3), Array(num))
     @test all(==(width), Array(max))
-    @test Array(sizes) == [i < 2width ? width : 2 for i in 0:n-1]
-    @test Array(id) == [div(i, width) + 1 for i in 0:n-1]
-    @test Array(lane) == [rem(i, width) + 1 for i in 0:n-1]
+    @test Array(sizes) == [i < 2width ? width : 2 for i in 0:(n - 1)]
+    @test Array(id) == [div(i, width) + 1 for i in 0:(n - 1)]
+    @test Array(lane) == [rem(i, width) + 1 for i in 0:(n - 1)]
 
     # the width is fixed
-    @test_throws ArgumentError KI.@launch backend launch=false sub_group_size=(width == 16 ? 8 : 16) ki_subgroup_kernel(num, sizes, max, id, lane)
+    @test_throws ArgumentError KI.@launch backend launch = false sub_group_size = (width == 16 ? 8 : 16) ki_subgroup_kernel(num, sizes, max, id, lane)
 end
 
 function ki_fill!(A)
@@ -51,15 +51,15 @@ end
 
 @testset "launch keywords" begin
     A = oneAPI.zeros(Int, 4)
-    kernel = KI.@launch oneAPIBackend() launch=false ki_fill!(A)
+    kernel = KI.@launch oneAPIBackend() launch = false ki_fill!(A)
 
     # oneAPI's launch options are passed on
-    kernel(A; ndrange=4, queue=oneAPI.global_stream(oneAPI.context(), oneAPI.device()))
+    kernel(A; ndrange = 4, queue = oneAPI.global_stream(oneAPI.context(), oneAPI.device()))
     @test Array(A) == 1:4
 
     # but not ones that would override the launch geometry
-    @test_throws ArgumentError kernel(A; ndrange=4, items=8)
-    @test_throws ArgumentError kernel(A; ndrange=4, groups=2)
+    @test_throws ArgumentError kernel(A; ndrange = 4, items = 8)
+    @test_throws ArgumentError kernel(A; ndrange = 4, groups = 2)
 end
 
 @testset "versioninfo" begin

…edKernels from their main branches

None of them is registered: KernelAbstractions 0.10 and KernelInterface 0.4
come from their development branch, and AcceleratedKernels from its main
branch, whose registered release excludes KernelAbstractions 0.10.
AcceleratedKernels depends on KernelAbstractions, so resolving oneAPI needs
KernelAbstractions from [sources] too, which Pkg only accepts for a regular
dependency: it is one until this commit is dropped. The test project takes
oneAPI from the checkout, so that it doesn't resolve the registered release.

Julia 1.10 ignores [sources], and Julia 1.11 doesn't pick them up for the
test environment, so CI develops all three explicitly there, and the
documentation is built with Julia 1.12.

Drop this commit once they are registered.

This branch has not been deployed

No deployments
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