Implement KernelInterface, and support KernelAbstractions 0.10

[maleadt]

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.

[github-actions]

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 5d52351..40e3996 100644
--- a/lib/level-zero/synchronization.jl
+++ b/lib/level-zero/synchronization.jl
@@ -36,7 +36,8 @@ gcsafe_synchronize(queue::ZeCommandQueue) =
         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
 
 """
     nonblocking_synchronize(list_or_queue)
diff --git a/src/compiler/compilation.jl b/src/compiler/compilation.jl
index 5e435cb..76ef39e 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}
@@ -272,8 +272,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/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 19d71ae..3369cac 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
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