diff --git a/Project.toml b/Project.toml index 569f884ff..f6ed01d91 100644 --- a/Project.toml +++ b/Project.toml @@ -11,7 +11,7 @@ ExprTools = "e2ba6199-217a-4e67-a87a-7c52f15ade04" GPUArrays = "0c68f7d7-f131-5f86-a1c3-88cf8149b2d7" GPUCompiler = "61eb1bfa-7361-4325-ad38-22787b887f55" GPUToolbox = "096a3bc2-3ced-46d0-87f4-dd12716f4bfc" -KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c" +KernelInterface = "4ee993da-d684-4d17-a7dd-4e58e78d92bf" LLVM = "929cbde3-209d-540e-8aea-75f648917ca0" Libdl = "8f399da3-3557-5675-b5ff-fb832c97cbdb" LinearAlgebra = "37e2e46d-f89d-539d-b4ee-838fcccc9c8e" @@ -26,11 +26,16 @@ SPIRV_LLVM_Translator_jll = "4a5d46fc-d8cf-5151-a261-86b458210efb" SPIRV_Tools_jll = "6ac6d60f-d740-5983-97d7-a4482c0689f4" SparseArrays = "2f01184e-e22b-5df5-ae63-d93ebab69eaf" SpecialFunctions = "276daf66-3868-5448-9aa4-cd146d93841b" -StaticArrays = "90137ffa-7385-5640-81b9-e52037218182" oneAPI_Level_Zero_Headers_jll = "f4bc562b-d309-54f8-9efb-476e56f0410d" oneAPI_Level_Zero_Loader_jll = "13eca655-d68d-5b81-8367-6d99d727ab01" oneAPI_Support_jll = "b049733a-a71d-5ed3-8eba-7d323ac00b36" +[weakdeps] +KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c" + +[extensions] +KernelAbstractionsExt = "KernelAbstractions" + [compat] AbstractFFTs = "1.5.0" Adapt = "4" @@ -39,7 +44,8 @@ ExprTools = "0.1" GPUArrays = "12.1" GPUCompiler = "2.13" GPUToolbox = "3.3.1" -KernelAbstractions = "0.9.39" +KernelAbstractions = "0.10" +KernelInterface = "0.4" LLVM = "10" NEO_jll = "=26.18.38308" PrecompileTools = "1" @@ -49,7 +55,6 @@ SPIRV_LLVM_Backend_jll = "23" SPIRV_LLVM_Translator_jll = "23" SPIRV_Tools_jll = "2025.4.0" SpecialFunctions = "1.3, 2" -StaticArrays = "1" julia = "1.10" oneAPI_Level_Zero_Loader_jll = "1.29" oneAPI_Support_jll = "0.10.0" diff --git a/ext/KernelAbstractionsExt.jl b/ext/KernelAbstractionsExt.jl new file mode 100644 index 000000000..f9a0d9fcd --- /dev/null +++ b/ext/KernelAbstractionsExt.jl @@ -0,0 +1,22 @@ +module KernelAbstractionsExt + +using oneAPI + +import KernelAbstractions as KA + +import Adapt + +Adapt.adapt_storage(::KA.CPU, a::oneArray) = convert(Array, a) + +# sparse arrays (oneMKL is only available on Linux) +@static if Sys.islinux() + import SparseArrays + Adapt.adapt_storage(::KA.CPU, a::oneAPI.oneMKL.oneAbstractSparseMatrix) = SparseArrays.SparseMatrixCSC(a) +end + + +## other + +Adapt.adapt_storage(to::KA.ConstAdaptor, a::oneDeviceArray) = Base.Experimental.Const(a) + +end diff --git a/lib/level-zero/event.jl b/lib/level-zero/event.jl index 92460f2b2..d805cf129 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) - desc_ref = Ref(ze_event_desc_t(; index=index-1)) + 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/module.jl b/lib/level-zero/module.jl index 91cc0dfcf..6e0635df6 100644 --- a/lib/level-zero/module.jl +++ b/lib/level-zero/module.jl @@ -85,13 +85,17 @@ mutable struct ZeKernel # Read on every launch by the scratch hedge, so it must not cost an API call. spill::Int + # cached maxGroupSize, seeded by `properties`; -1 while unqueried, and 0 without the + # MAX_GROUP_SIZE extension. Read on every KernelInterface launch. + max_group_size::Int + function ZeKernel(mod, name) GC.@preserve name begin desc_ref = Ref(ze_kernel_desc_t(; pKernelName=pointer(name))) handle_ref = Ref{ze_kernel_handle_t}() zeKernelCreate(mod, desc_ref, handle_ref) end - obj = new(mod, handle_ref[], ReentrantLock(), -1) + obj = new(mod, handle_ref[], ReentrantLock(), -1, -1) finalizer(obj) do obj zeKernelDestroy(obj) @@ -268,6 +272,8 @@ function properties(kernel::ZeKernel) props = props_ref[] kernel.spill = Int(props.spillMemSize) + kernel.max_group_size = max_group_size_props_ref === nothing ? 0 : + Int(max_group_size_props_ref[].maxGroupSize) return ( numKernelArgs=Int(props.numKernelArgs), requiredGroupSize=ZeDim3(props.requiredGroupSizeX, @@ -295,6 +301,13 @@ function spill_mem_size(kernel::ZeKernel) return s >= 0 ? s : Int(properties(kernel).spillMemSize) end +# Cached access to a kernel's maxGroupSize, `missing` without the MAX_GROUP_SIZE extension. +function max_group_size(kernel::ZeKernel) + s = kernel.max_group_size + s < 0 && (s = coalesce(properties(kernel).maxGroupSize, 0)) + return s > 0 ? s : missing +end + ## execution diff --git a/lib/level-zero/oneL0.jl b/lib/level-zero/oneL0.jl index c36b9eee3..f8e32e8ad 100644 --- a/lib/level-zero/oneL0.jl +++ b/lib/level-zero/oneL0.jl @@ -124,9 +124,9 @@ end include("context.jl") include("cmdqueue.jl") include("cmdlist.jl") -include("synchronization.jl") include("fence.jl") include("event.jl") +include("synchronization.jl") include("barrier.jl") include("module.jl") include("memory.jl") diff --git a/lib/level-zero/synchronization.jl b/lib/level-zero/synchronization.jl index e3a3a4e20..5d5235180 100644 --- a/lib/level-zero/synchronization.jl +++ b/lib/level-zero/synchronization.jl @@ -1,16 +1,16 @@ # cooperative synchronization # -# `zeCommandListHostSynchronize` and `zeCommandQueueSynchronize` block the calling thread -# until the work has completed, so no other task can run on it in the meantime. Instead, wait -# using GPUToolbox's `cooperative_wait`: first poll, which keeps the latency of short -# operations low, and then block in the driver on a separate thread, while the calling task -# yields. +# `zeCommandListHostSynchronize`, `zeCommandQueueSynchronize` and `zeEventHostSynchronize` +# block the calling thread until the work has completed, so no other task can run on it in +# the meantime. Instead, wait using GPUToolbox's `cooperative_wait`: first poll, which keeps +# the latency of short operations low, and then block in the driver on a separate thread, +# while the calling task yields. using GPUToolbox: cooperative_wait export nonblocking_synchronize -const SyncObject = Union{ZeImmediateCommandList, ZeCommandQueue} +const SyncObject = Union{ZeImmediateCommandList, ZeCommandQueue, ZeEvent} # with a zero timeout, a synchronization is a query function check_done(res::ze_result_t) @@ -34,13 +34,17 @@ gcsafe_synchronize(list::ZeImmediateCommandList) = gcsafe_synchronize(queue::ZeCommandQueue) = @gcsafe_ccall libze_loader.zeCommandQueueSynchronize( 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 """ nonblocking_synchronize(list_or_queue) + nonblocking_synchronize(event) -Wait for the work on an immediate command list or command queue to complete, like -[`synchronize`](@ref), but without blocking the calling thread: other tasks keep running -while this one waits. +Wait for the work on an immediate command list or command queue to complete, or for an +event to be signaled, like [`synchronize`](@ref) or `wait`, but without blocking the calling +thread: other tasks keep running while this one waits. """ function nonblocking_synchronize(obj::SyncObject) # when polling found the work to be done, synchronize again to check for errors diff --git a/src/compiler/compilation.jl b/src/compiler/compilation.jl index 33b5537b8..f8aa0cb0c 100644 --- a/src/compiler/compilation.jl +++ b/src/compiler/compilation.jl @@ -1,6 +1,9 @@ ## gpucompiler interface implementation -struct oneAPICompilerParams <: AbstractCompilerParams end +Base.@kwdef struct oneAPICompilerParams <: AbstractCompilerParams + sub_group_size::Union{Nothing,Int} = nothing +end + const oneAPICompilerConfig = CompilerConfig{SPIRVCompilerTarget, oneAPICompilerParams} const oneAPICompilerJob = CompilerJob{SPIRVCompilerTarget,oneAPICompilerParams} @@ -56,6 +59,11 @@ function GPUCompiler.finish_module!(job::oneAPICompilerJob, mod::LLVM.Module, Tuple{CompilerJob{SPIRVCompilerTarget}, typeof(mod), typeof(entry)}, job, mod, entry) + # require a sub-group size, instead of letting the compiler choose one + if job.config.params.sub_group_size !== nothing + entry.metadata["intel_reqd_sub_group_size"] = MDNode([ConstantInt(Int32(job.config.params.sub_group_size))]) + end + # OpenCL 2.0 push!( get!(mod.metadata, "opencl.ocl.version").operands, @@ -290,11 +298,18 @@ function _atomics_capabilities(dev, properties; supports_fp16, supports_fp64) ) end -@noinline function _compiler_config(dev; kernel=true, name=nothing, always_inline=false, 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 + if sub_group_size !== nothing + sizes = oneL0.compute_properties(dev).subGroupSizes + sub_group_size in sizes || + throw(ArgumentError("Sub-group size $sub_group_size is not supported by this device, which supports $(join(sizes, ", "))")) + end + # SPIR-V codegen path. The Aurora LTS NEO/IGC runtime only accepts SPIR-V from the # Khronos translator; the rolling stack uses the LLVM SPIR-V back-end. GPUCompiler picks # the tool from the target's `backend` field and loads the JLL lazily, so both can be @@ -329,7 +344,7 @@ end target = SPIRVCompilerTarget(; backend, extensions = extensions_str, atomics, supports_fp16, supports_fp64, supports_bfloat16, driver = :intel, kwargs...) - params = oneAPICompilerParams() + params = oneAPICompilerParams(; sub_group_size) CompilerConfig(target, params; kernel, name, always_inline) end diff --git a/src/compiler/execution.jl b/src/compiler/execution.jl index c5fb4b927..6dd470934 100644 --- a/src/compiler/execution.jl +++ b/src/compiler/execution.jl @@ -4,7 +4,7 @@ export @oneapi, zefunction, kernel_convert ## high-level @oneapi interface const MACRO_KWARGS = [:launch] -const COMPILER_KWARGS = [:kernel, :name, :always_inline, :atomics] +const COMPILER_KWARGS = [:kernel, :name, :always_inline, :atomics, :sub_group_size] const LAUNCH_KWARGS = [:groups, :items, :queue] """ @@ -33,6 +33,9 @@ launches the kernel on the GPU. additions on Xe-LP). Disabled floating-point operations are implemented with integer compare-and-swap loops instead, and disabling `int64` makes 64-bit atomic operations an error. +- `sub_group_size::Union{Int,Nothing}=nothing`: The sub-group size the kernel has to be + compiled for, one of the device's `oneL0.compute_properties(dev).subGroupSizes`. By + default, the compiler chooses one. ## Launch Keywords (runtime) - `groups`: Number of workgroups (required). Can be an integer or tuple. @@ -249,9 +252,9 @@ function launch_configuration(kernel::HostKernel{F,TT}) where {F,TT} # configurations, so roll our own version that behaves like CUDA's # occupancy API and assumes the kernel still does bounds checking. - kernel_props = oneL0.properties(kernel.fun) - group_size = if kernel_props.maxGroupSize !== missing - kernel_props.maxGroupSize + max_group_size = oneL0.max_group_size(kernel.fun) + group_size = if max_group_size !== missing + max_group_size else # without the MAX_GROUP_SIZE extension, we need to be conservative dev = kernel.fun.mod.device @@ -269,7 +272,7 @@ function launch_configuration(kernel::HostKernel{F,TT}) where {F,TT} # size but does not fold it into `maxGroupSize`, so account for it here. Rounded down to # a power of two, both because group sizes want to be anyway and to stay clear of the # limit rather than right at it. - spill = kernel_props.spillMemSize + spill = oneL0.spill_mem_size(kernel.fun) if spill > 0 && group_size * spill > MAX_GROUP_SCRATCH group_size = max(1, prevpow(2, max(1, MAX_GROUP_SCRATCH รท spill))) end diff --git a/src/context.jl b/src/context.jl index 6c9e188fa..521adb3ec 100644 --- a/src/context.jl +++ b/src/context.jl @@ -284,7 +284,7 @@ global_queue(ctx::ZeContext, dev::ZeDevice) = stream_queue(global_stream(ctx, de # Register `stream` as a stream targeting (ctx, dev) so `synchronize_all_streams`/ # `release` can find and drain it before freeing buffers whose in-flight work it may # still reference. EVERY stream that becomes a task's active stream must go through here -# โ€” not just the one `global_stream` creates but also the replacement `KA.priority!` +# โ€” not just the one `global_stream` creates but also the replacement `KI.priority!` # installs โ€” or the unregistered stream's in-flight work can outlive a freed buffer (a # use-after-free that faults and bans the context on the LTS NEO stack). Only the LTS # stack maintains the registry; on the rolling stack this is a no-op. Returns `stream`. diff --git a/src/oneAPI.jl b/src/oneAPI.jl index 1f098f680..f6489f6c2 100644 --- a/src/oneAPI.jl +++ b/src/oneAPI.jl @@ -12,7 +12,7 @@ using SpecialFunctions import Preferences -import KernelAbstractions: KernelAbstractions +import KernelInterface using LLVM, LLVM.IR, LLVM.Build, LLVM.Interop using Core: LLVMPtr diff --git a/src/oneAPIKernels.jl b/src/oneAPIKernels.jl index ae693caac..664aed0ac 100644 --- a/src/oneAPIKernels.jl +++ b/src/oneAPIKernels.jl @@ -1,11 +1,9 @@ module oneAPIKernels using ..oneAPI -using ..oneAPI: @device_override, SPIRVIntrinsics, method_table +using ..oneAPI: @device_override, SPIRVIntrinsics, method_table, kernel_convert, zefunction -import KernelAbstractions as KA - -import StaticArrays +import KernelInterface as KI import Adapt @@ -14,215 +12,250 @@ import Adapt export oneAPIBackend -struct oneAPIBackend <: KA.GPU +""" + oneAPIBackend(; prefer_blocks=false, always_inline=false) + +KernelInterface back end for oneAPI, which KernelAbstractions uses to launch kernels on +Intel GPUs. `prefer_blocks=true` spreads small launches over more, smaller work-groups. +""" +struct oneAPIBackend <: KI.Backend prefer_blocks::Bool always_inline::Bool end oneAPIBackend(; prefer_blocks = false, always_inline = false) = oneAPIBackend(prefer_blocks, always_inline) -@inline KA.allocate(::oneAPIBackend, ::Type{T}, dims::Tuple; unified::Bool = false) where {T} = oneArray{T, length(dims), unified ? oneAPI.oneL0.SharedBuffer : oneAPI.oneL0.DeviceBuffer}(undef, dims) -@inline KA.zeros(::oneAPIBackend, ::Type{T}, dims::Tuple; unified::Bool = false) where {T} = fill!(oneArray{T, length(dims), unified ? oneAPI.oneL0.SharedBuffer : oneAPI.oneL0.DeviceBuffer}(undef, dims), zero(T)) -@inline KA.ones(::oneAPIBackend, ::Type{T}, dims::Tuple; unified::Bool = false) where {T} = fill!(oneArray{T, length(dims), unified ? oneAPI.oneL0.SharedBuffer : oneAPI.oneL0.DeviceBuffer}(undef, dims), one(T)) - -KA.get_backend(::oneArray) = oneAPIBackend() -KA.synchronize(::oneAPIBackend) = oneAPI.oneL0.synchronize() -KA.supports_float64(::oneAPIBackend) = false # TODO: Check if this is device dependent -KA.supports_unified(::oneAPIBackend) = true +@inline KI.allocate(::oneAPIBackend, ::Type{T}, dims::Tuple; unified::Bool = false) where {T} = oneArray{T, length(dims), unified ? oneAPI.oneL0.SharedBuffer : oneAPI.oneL0.DeviceBuffer}(undef, dims) + +KI.get_backend(::oneArray) = oneAPIBackend() +KI.synchronize(::oneAPIBackend) = oneAPI.oneL0.synchronize() +KI.supports_float64(::oneAPIBackend) = device_limits().supports_float64 +KI.supports_unified(::oneAPIBackend) = true +KI.supports_atomics(::oneAPIBackend) = true +KI.supports_subgroups(::oneAPIBackend) = device_limits().sub_group_size > 0 +# Intel's GPU compiler forms sub-groups from consecutive work-items (which the KernelInterface +# testsuite checks), and they execute independently +KI.supports_linear_subgroups(backend::oneAPIBackend) = KI.supports_subgroups(backend) +KI.supports_independent_subgroups(backend::oneAPIBackend) = KI.supports_subgroups(backend) +# A sub-group is the SIMD lanes of one EU thread, so a segment of it executes the functions +# with a `width` together without the rest of the sub-group: the shuffles and the ballot are +# `OpGroupNonUniform*` instructions over the active lanes, and `KI.sub_group_barrier(width)` +# only has to order memory (see below). +KI.supports_partitionable_subgroups(backend::oneAPIBackend) = KI.supports_subgroups(backend) + +# The types `sub_group_shuffle` supports; other types are shuffled as words. `Float16` and +# `Float64` are reported as unsupported on devices without them, rather than shuffled as +# words: the overrides below are selected by type, not by device. +const ShuffleTypes = Union{SPIRVIntrinsics.gentypes...} + +function KI.supports_shuffle(backend::oneAPIBackend, ::Type{T}) where {T <: ShuffleTypes} + KI.supports_subgroups(backend) || return false + T === Float64 && return device_limits().supports_float64 + T === Float16 && return device_limits().supports_float16 + return true +end -KA.functional(::oneAPIBackend) = oneAPI.functional() +KI.functional(::oneAPIBackend) = oneAPI.functional() Adapt.adapt_storage(::oneAPIBackend, a::AbstractArray) = Adapt.adapt(oneArray, a) Adapt.adapt_storage(::oneAPIBackend, a::oneArray) = a -Adapt.adapt_storage(::KA.CPU, a::oneArray) = convert(Array, a) # sparse arrays (oneMKL is only available on Linux) @static if Sys.islinux() - import GPUArrays, SparseArrays - KA.get_backend(::oneAPI.oneMKL.oneAbstractSparseMatrix) = oneAPIBackend() + import GPUArrays + KI.get_backend(::oneAPI.oneMKL.oneAbstractSparseMatrix) = oneAPIBackend() # without this, `adapt_storage(::oneAPIBackend, ::AbstractArray)` would densify sparse arrays Adapt.adapt_storage(::oneAPIBackend, a::GPUArrays.AbstractGPUSparseArray) = a - Adapt.adapt_storage(::KA.CPU, a::oneAPI.oneMKL.oneAbstractSparseMatrix) = SparseArrays.SparseMatrixCSC(a) end ## Memory Operations -function KA.copyto!(::oneAPIBackend, A, B) +function KI.copyto!(::oneAPIBackend, A, B) + length(A) == length(B) || + throw(ArgumentError("Arrays must have the same length, got $(length(A)) and $(length(B))")) copyto!(A, B) # TODO: Address device to host copies in jl being synchronizing + return A end +KI.unsafe_free!(A::oneArray) = oneAPI.unsafe_free!(A) + ## Device Operations -function KA.ndevices(::oneAPIBackend) +function KI.ndevices(::oneAPIBackend) return length(oneAPI.devices()) end -function KA.device(::oneAPIBackend)::Int - dev = oneAPI.device() +function device_index(dev)::Int devs = oneAPI.devices() idx = findfirst(==(dev), devs) return idx === nothing ? 1 : idx end +KI.device(::oneAPIBackend)::Int = device_index(oneAPI.device()) +KI.device(::oneAPIBackend, A::oneArray)::Int = device_index(oneAPI.device(A)) -function KA.device!(backend::oneAPIBackend, id::Int) - return oneAPI.device!(id) +function KI.device!(backend::oneAPIBackend, id::Int) + oneAPI.device!(id) + return end ## Kernel Launch -function KA.mkcontext(kernel::KA.Kernel{oneAPIBackend}, _ndrange, iterspace) - KA.CompilerMetadata{KA.ndrange(kernel), KA.DynamicCheck}(_ndrange, iterspace) -end -function KA.mkcontext(kernel::KA.Kernel{oneAPIBackend}, I, _ndrange, iterspace, - ::Dynamic) where Dynamic - KA.CompilerMetadata{KA.ndrange(kernel), Dynamic}(I, _ndrange, iterspace) -end +KI.argconvert(::oneAPIBackend, arg) = kernel_convert(arg) -function KA.launch_config(kernel::KA.Kernel{oneAPIBackend}, ndrange, workgroupsize) - if ndrange isa Integer - ndrange = (ndrange,) - end - if workgroupsize isa Integer - workgroupsize = (workgroupsize, ) - end +# a compiled kernel, and the callable it was compiled from: the converted callable only holds +# pointers to the arrays a closure captures, so the kernel has to keep the original alive +struct oneAPIKernel{F, H <: oneAPI.HostKernel} + f::F + host::H +end - # partition checked that the ndrange's agreed - if KA.ndrange(kernel) <: KA.StaticSize - ndrange = nothing +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 + throw(ArgumentError("KernelInterface kernels are compiled for a sub-group size of $sub_group_size, got $(kwargs[:sub_group_size])")) end - - iterspace, dynamic = if KA.workgroupsize(kernel) <: KA.DynamicSize && - workgroupsize === nothing - # use ndrange as preliminary workgroupsize for autotuning - # (clamped to 1, since an empty ndrange cannot serve as a workgroup size) - KA.partition(kernel, ndrange, max.(ndrange, 1)) + host = if sub_group_size > 0 + zefunction(kernel_convert(f), tt; name, backend.always_inline, kwargs..., sub_group_size) else - KA.partition(kernel, ndrange, workgroupsize) + zefunction(kernel_convert(f), tt; name, backend.always_inline, kwargs...) end - - return ndrange, workgroupsize, iterspace, dynamic + kern = oneAPIKernel(f, host) + KI.Kernel{oneAPIBackend, typeof(kern)}(backend, kern) end -function threads_to_workgroupsize(threads, ndrange) - total = 1 - return map(ndrange) do n - x = max(1, min(div(threads, total), n)) - total *= x - return x +# 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...) + # 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`")) end + # kernels are compiled for a context and device, and by default launched on the task's + # stream for the active ones + (; f, host) = obj.kern + check_target(host.fun.mod, get(kwargs, :queue, nothing)) + GC.@preserve f host(args...; items, groups, kwargs...) + return end -function (obj::KA.Kernel{oneAPIBackend})(args...; ndrange=nothing, workgroupsize=nothing) - backend = KA.backend(obj) - - ndrange, workgroupsize, iterspace, dynamic = KA.launch_config(obj, ndrange, workgroupsize) - # this might not be the final context, since we may tune the workgroupsize - ctx = KA.mkcontext(obj, ndrange, iterspace) - - # If the kernel is statically sized we can tell the compiler about that - if KA.workgroupsize(obj) <: KA.StaticSize - # TODO: maxthreads - # maxthreads = prod(KA.get(KA.workgroupsize(obj))) +function check_target(mod::oneAPI.oneL0.ZeModule, queue) + ctx, dev = if queue === nothing + context(), device() + elseif queue isa oneAPI.oneStream + queue.ctx, queue.dev else - # maxthreads = nothing + queue.context, queue.device end + mod.context == ctx && mod.device == dev || + throw(ArgumentError("Cannot launch a kernel on another context or device than the one it was compiled for")) + return +end - kernel = @oneapi launch = false always_inline = backend.always_inline obj.f(ctx, args...) - - # figure out the optimal workgroupsize automatically - if KA.workgroupsize(obj) <: KA.DynamicSize && workgroupsize === nothing - items = oneAPI.launch_configuration(kernel) - - if backend.prefer_blocks - # Prefer blocks over threads: - # Reducing the workgroup size (items) increases the number of workgroups (blocks). - # We use a simple heuristic here since we lack full occupancy info (max_blocks) from launch_configuration. - - # If the total range is large enough, full workgroups are fine. - # If the range is small, we might want to reduce 'items' to create more blocks to fill the GPU. - # (Simplified logic compared to CUDA.jl which uses explicit occupancy calculators) - total_items = prod(ndrange) - if total_items < items * 16 # Heuristic factor - # Force at least a few blocks if possible by reducing items per block - target_blocks = 16 # Target at least 16 blocks - items = max(1, min(items, cld(total_items, target_blocks))) - end +function KI.max_work_group_size(kernel::KI.Kernel{oneAPIBackend})::Int + fun = kernel.kern.host.fun + max_group_size = oneAPI.oneL0.max_group_size(fun) + # without the MAX_GROUP_SIZE extension, the device limit is all we know + return coalesce(max_group_size, device_limits(fun.mod.device).max_work_group_size) +end +function KI.launch_configuration( + kernel::KI.Kernel{oneAPIBackend}; nitems::Union{Integer, Nothing} = nothing, + max_work_group_size::Integer = typemax(Int) + ) + items = min(oneAPI.launch_configuration(kernel.kern.host), max_work_group_size) + if kernel.backend.prefer_blocks && nitems !== nothing + # prefer more, smaller work-groups: without an occupancy API that tells how many + # work-groups fill the device, make sure that a small launch uses at least a few + target_groups = 16 + if nitems < items * target_groups + items = max(1, min(items, cld(nitems, target_groups))) end - - workgroupsize = threads_to_workgroupsize(items, ndrange) - iterspace, dynamic = KA.partition(obj, ndrange, workgroupsize) - ctx = KA.mkcontext(obj, ndrange, iterspace) end - - groups = length(KA.blocks(iterspace)) - items = length(KA.workitems(iterspace)) - - if groups == 0 - return nothing + return (; workgroupsize = Int(items)) +end +# querying the device allocates, so cache what every launch needs +const DeviceLimits = @NamedTuple{ + max_work_group_size::Int, max_work_group_dims::NTuple{3, Int}, max_num_groups::NTuple{3, Int}, + sub_group_size::Int, supports_float16::Bool, supports_float64::Bool, +} +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 + 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. Only widths of up to + # 64, so that `KI.sub_group_ballot` (a `UInt64` mask) covers every lane. + sg_sizes = filter(<=(64), 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) end - - # Launch kernel - kernel(ctx, args...; items, groups) - - return nothing +end +KI.max_work_group_size(::oneAPIBackend)::Int = device_limits().max_work_group_size +KI.max_work_group_dims(::oneAPIBackend)::NTuple{3, Int} = device_limits().max_work_group_dims +KI.max_num_groups(::oneAPIBackend)::NTuple{3, Int} = device_limits().max_num_groups +KI.sub_group_size(::oneAPIBackend)::Int = device_limits().sub_group_size +# Level Zero calls Xe cores sub-slices +function KI.multiprocessor_count(::oneAPIBackend)::Int + props = oneAPI.oneL0.properties(device()) + return props.numSlices * props.numSubslicesPerSlice end ## Indexing Functions +## COV_EXCL_START -@device_override @inline function KA.__index_Local_Linear(ctx) - return get_local_id() -end +# computed with `% T`, which unlike `T(x)` has no error path -@device_override @inline function KA.__index_Group_Linear(ctx) - return get_group_id() +@device_override @inline function KI.get_local_id(::Type{T}) where {T} + return (; x = get_local_id(1) % T, y = get_local_id(2) % T, z = get_local_id(3) % T) end -@device_override @inline function KA.__index_Global_Linear(ctx) - return get_global_id() +@device_override @inline function KI.get_group_id(::Type{T}) where {T} + return (; x = get_group_id(1) % T, y = get_group_id(2) % T, z = get_group_id(3) % T) end -@device_override @inline function KA.__index_Local_Cartesian(ctx) - @inbounds KA.workitems(KA.__iterspace(ctx))[get_local_id()] +@device_override @inline function KI.get_local_size(::Type{T}) where {T} + return (; x = get_local_size(1) % T, y = get_local_size(2) % T, z = get_local_size(3) % T) end -@device_override @inline function KA.__index_Group_Cartesian(ctx) - @inbounds KA.blocks(KA.__iterspace(ctx))[get_group_id()] +@device_override @inline function KI.get_num_groups(::Type{T}) where {T} + return (; x = get_num_groups(1) % T, y = get_num_groups(2) % T, z = get_num_groups(3) % T) end -@device_override @inline function KA.__index_Global_Cartesian(ctx) - return @inbounds KA.expand(KA.__iterspace(ctx), get_group_id(), get_local_id()) -end +@device_override KI.get_sub_group_size(::Type{T}) where {T} = get_sub_group_size() % T -@device_override @inline function KA.__validindex(ctx) - if KA.__dynamic_checkbounds(ctx) - I = @inbounds KA.expand(KA.__iterspace(ctx), get_group_id(), get_local_id()) - return I in KA.__ndrange(ctx) - else - return true - end -end +@device_override KI.get_max_sub_group_size(::Type{T}) where {T} = get_max_sub_group_size() % T + +@device_override KI.get_num_sub_groups(::Type{T}) where {T} = get_num_sub_groups() % T + +@device_override KI.get_sub_group_id(::Type{T}) where {T} = get_sub_group_id() % T +@device_override KI.get_sub_group_local_id(::Type{T}) where {T} = get_sub_group_local_id() % T -## Shared and Scratch Memory -@device_override @inline function KA.SharedMemory(::Type{T}, ::Val{Dims}, ::Val{Id}) where {T, Dims, Id} +## Shared Memory + +@device_override @inline function KI.localmemory(::Type{T}, ::Val{Dims}) where {T, Dims} ptr = oneAPI.emit_localmemory(T, Val(prod(Dims))) oneDeviceArray(Dims, ptr) end -@device_override @inline function KA.Scratchpad(ctx, ::Type{T}, ::Val{Dims}) where {T, Dims} - StaticArrays.MArray{KA.__size(Dims), T}(undef) -end - ## Synchronization and Printing -@device_override @inline function KA.__synchronize() +@device_override @inline function KI.barrier() # Fence both local and global memory across the workgroup barrier, matching CUDA # `__syncthreads` semantics. `barrier(0)` lowers to `OpControlBarrier` with # `SequentiallyConsistent` but WITHOUT any storage-class bit, which the SPIR-V spec @@ -232,18 +265,141 @@ end barrier(SPIRVIntrinsics.LOCAL_MEM_FENCE | SPIRVIntrinsics.GLOBAL_MEM_FENCE) end -@device_override @inline function KA.__print(args...) +@device_override @inline function KI.sub_group_barrier() + sub_group_barrier(SPIRVIntrinsics.LOCAL_MEM_FENCE | SPIRVIntrinsics.GLOBAL_MEM_FENCE) +end + +# A sub-group `OpControlBarrier` has to be executed by the whole sub-group, so a narrower +# segment, whose lanes execute together anyway, only fences its memory accesses +@device_override @inline function KI.sub_group_barrier(width::Integer) + if width == get_max_sub_group_size() + KI.sub_group_barrier() + else + SPIRVIntrinsics.atomic_work_item_fence( + SPIRVIntrinsics.LOCAL_MEM_FENCE | SPIRVIntrinsics.GLOBAL_MEM_FENCE, + SPIRVIntrinsics.memory_order_acq_rel, SPIRVIntrinsics.memory_scope_sub_group) + end +end + +@device_override KI.shfl(val::T, lane::Integer) where {T <: ShuffleTypes} = + sub_group_shuffle(val, lane) + +# past the sub-group width, `shfl_down` and `shfl_up` return the work-item's own value, which +# `sub_group_shuffle` (like SPIR-V's `OpGroupNonUniformShuffleDown`) leaves undefined +@device_override function KI.shfl_down(val::T, offset::Integer) where {T <: ShuffleTypes} + lane = get_sub_group_local_id() + # compared before adding, so that large offsets don't overflow + inside = offset <= get_max_sub_group_size() - lane + return sub_group_shuffle(val, ifelse(inside, lane + offset, lane)) +end + +@device_override function KI.shfl_up(val::T, offset::Integer) where {T <: ShuffleTypes} + lane = get_sub_group_local_id() + return sub_group_shuffle(val, ifelse(lane > offset, lane - offset, lane)) +end + +@device_override KI.shfl_xor(val::T, mask::Integer) where {T <: ShuffleTypes} = + sub_group_shuffle_xor(val, mask) + +@device_override KI.sub_group_any(pred::Bool) = SPIRVIntrinsics.sub_group_any(pred) + +@device_override KI.sub_group_all(pred::Bool) = SPIRVIntrinsics.sub_group_all(pred) + +@device_override function KI.sub_group_ballot(pred::Bool) + mask = SPIRVIntrinsics.sub_group_ballot(pred) + return UInt64(mask[1].value) | (UInt64(mask[2].value) << 32) +end + +# The native reductions and scans, for `+` on 32- and 64-bit integers and floats, and +# `min`/`max` on integers (OpenCL's `min` and `max` treat NaN and the sign of zero differently +# from Julia's). +const CollectiveIntTypes = Union{Int32, UInt32, Int64, UInt64} +const CollectiveTypes = Union{CollectiveIntTypes, Float16, Float32, Float64} + +@device_override KI.sub_group_reduce(::typeof(+), val::CollectiveTypes) = + SPIRVIntrinsics.sub_group_reduce_add(val) +@device_override KI.sub_group_reduce(::typeof(min), val::CollectiveIntTypes) = + SPIRVIntrinsics.sub_group_reduce_min(val) +@device_override KI.sub_group_reduce(::typeof(max), val::CollectiveIntTypes) = + SPIRVIntrinsics.sub_group_reduce_max(val) + +@device_override KI.sub_group_scan(::typeof(+), val::CollectiveTypes) = + SPIRVIntrinsics.sub_group_scan_inclusive_add(val) +@device_override KI.sub_group_scan(::typeof(min), val::CollectiveIntTypes) = + SPIRVIntrinsics.sub_group_scan_inclusive_min(val) +@device_override KI.sub_group_scan(::typeof(max), val::CollectiveIntTypes) = + SPIRVIntrinsics.sub_group_scan_inclusive_max(val) + +# the native exclusive scans start from the identity, not from `init` +@device_override function KI.sub_group_exclusive_scan(::typeof(+), val::T, init::T) where {T <: CollectiveTypes} + prefix = SPIRVIntrinsics.sub_group_scan_exclusive_add(val) + return ifelse(get_sub_group_local_id() == 1, init, init + prefix) +end + +@device_override @inline function KI._print(args...) oneAPI._print(args...) end +## COV_EXCL_STOP -## Other -Adapt.adapt_storage(to::KA.ConstAdaptor, a::oneDeviceArray) = Base.Experimental.Const(a) +## Events + +# Level Zero has no reusable timeline events, so every `record_event` creates a fresh +# one-shot event in its own host-visible pool. Its signal has `HOST` scope, the broadest: +# it flushes caches far enough for the host *and* other devices. +# +# `wait_event` waits for the event on the host, cooperatively: making the waiting task's +# stream wait for it on the device would require keeping the event alive until that stream +# has executed the wait. Recorded events are kept alive until they have been signaled, as +# destroying an event that the device still has to signal is undefined; not by the +# recording task, which may end before that. +const pending_events = oneAPI.oneL0.ZeEvent[] +const pending_events_lock = ReentrantLock() + +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) + + @lock pending_events_lock begin + filter!(!Base.isdone, pending_events) + push!(pending_events, ev) + end + + # the task's stream is an in-order immediate command list, so the signal fires once + # everything appended before it has completed. `execute!` first drains any pending + # oneMKL work on the companion queue (`mkl_wait!`), so that is captured as well. + try + oneAPI.oneL0.execute!(oneAPI.global_stream(ctx, dev)) do list + oneAPI.oneL0.append_signal!(list, ev) + end + catch + # the signal wasn't submitted, so nothing references the event + @lock pending_events_lock filter!(!=(ev), pending_events) + rethrow() + end + + return ev +end + +# the event can come from any device, also one of another driver +function KI.wait_event(::oneAPIBackend, ev::oneAPI.oneL0.ZeEvent) + oneAPI.oneL0.nonblocking_synchronize(ev) + return +end + + +## Other -KA.argconvert(::KA.Kernel{oneAPIBackend}, arg) = kernel_convert(arg) +KI.versioninfo(io::IO, ::oneAPIBackend) = oneAPI.versioninfo(io) -function KA.priority!(::oneAPIBackend, prio::Symbol) +function KI.priority!(::oneAPIBackend, prio::Symbol) if !(prio in (:high, :normal, :low)) error("priority must be one of :high, :normal, :low") end diff --git a/src/utils.jl b/src/utils.jl index ddf4b1f91..05bceb8af 100644 --- a/src/utils.jl +++ b/src/utils.jl @@ -27,11 +27,19 @@ function versioninfo(io::IO=stdout) println(io, "- LLVM: $(LLVM.version())") println(io) + + 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)) + end + println(io, "Julia packages:") println(io, "- oneAPI.jl: $(Base.pkgversion(oneAPI))") - for name in [:GPUArrays, :GPUCompiler, :KernelAbstractions, :LLVM, :SPIRVIntrinsics] - mod = getfield(oneAPI, name) - println(io, "- $(name): $(Base.pkgversion(mod))") + 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 println(io) diff --git a/test/Project.toml b/test/Project.toml index beab523b5..af23a69ba 100644 --- a/test/Project.toml +++ b/test/Project.toml @@ -9,6 +9,7 @@ GPUArrays = "0c68f7d7-f131-5f86-a1c3-88cf8149b2d7" InteractiveUtils = "b77e0a4c-d291-57a0-90e8-8db25a27a240" JLD2 = "033835bb-8acc-5ee8-8aae-3f567f8a3819" KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c" +KernelInterface = "4ee993da-d684-4d17-a7dd-4e58e78d92bf" LinearAlgebra = "37e2e46d-f89d-539d-b4ee-838fcccc9c8e" NEO_jll = "700fe977-ac61-5f37-bbc8-c6c4b2b6a9fd" ParallelTestRunner = "d3525ed8-44d0-4b2c-a655-542cee43accc" diff --git a/test/execution.jl b/test/execution.jl index 00af8e3a5..19d71ae65 100644 --- a/test/execution.jl +++ b/test/execution.jl @@ -41,6 +41,22 @@ end 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 + end + @test_throws ArgumentError @oneapi sub_group_size=maximum(sizes)+1 dummy() +end + + @testset "inference" begin foo() = @oneapi dummy() @inferred foo() diff --git a/test/kernelabstractions.jl b/test/kernelabstractions.jl index 48bbe1cb7..e3f3b56cb 100644 --- a/test/kernelabstractions.jl +++ b/test/kernelabstractions.jl @@ -1,8 +1,82 @@ import KernelAbstractions +import KernelAbstractions as KA +import KernelInterface as KI include(joinpath(dirname(pathof(KernelAbstractions)), "..", "test", "testsuite.jl")) 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", ]) -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) + @inbounds A[I] = I +end + +KA.@kernel function store_last_index!(A) + I = KA.@index(Global, Linear) + if I == prod(KA.@ndrange()) + @inbounds A[1] = I + @inbounds A[2] = KA.@index(Global, Cartesian)[2] + end +end + +@testset "launch configuration" begin + backend = oneAPIBackend() + function select(kernel, ndrange, workgroupsize=nothing) + ndrange, workgroupsize, iterspace, _ = KA.launch_config(kernel, ndrange, workgroupsize) + KA.select_launch(kernel, workgroupsize, iterspace) + end + + # kernels are launched on an N-d grid, computing indices in 32 bits + kernel = store_global_linear!(backend) + @test select(kernel, (64, 32, 16)) === KA.NDLaunch{Int32}() + @test select(kernel, (4, 4, 4, 4)) === KA.LinearLaunch{Int32}() + + # 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))) + @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) + + # 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}())) + @test select(kernel, dims) === launch + kernel(A; ndrange=dims) + @test Array(A) == [prod(dims), dims[2]] + end +end + +function ki_store_index!(A) + i = KI.get_global_id().x + if i <= length(A) + @inbounds A[i] = i + end + return +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 + @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 + @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 Array(A) == 1:1024 +end diff --git a/test/kernelinterface.jl b/test/kernelinterface.jl new file mode 100644 index 000000000..3acfec5dc --- /dev/null +++ b/test/kernelinterface.jl @@ -0,0 +1,70 @@ +import KernelInterface +import KernelInterface as KI + +include(joinpath(dirname(pathof(KernelInterface)), "..", "test", "testsuite.jl")) + +# synchronizing every submission makes a launch wait for the device, so the timing check that +# an event wait is ordered after earlier work can't distinguish a missing wait +skip_tests = oneL0.sync_each_submission() ? Set(["Events"]) : Set{String}() +Testsuite.testsuite(oneAPIBackend(), oneArray; skip_tests) + +function ki_subgroup_kernel(num, sizes, max, id, lane) + l = KI.get_local_id() + s = KI.get_local_size() + i = l.x + (l.y - 1) * s.x + @inbounds begin + num[i] = KI.get_num_sub_groups() + sizes[i] = KI.get_sub_group_size() + max[i] = KI.get_max_sub_group_size() + id[i] = KI.get_sub_group_id() + lane[i] = KI.get_sub_group_local_id() + end + return +end + +# KernelInterface leaves the formation of sub-groups unspecified; Intel GPUs form them from +# consecutive linear work-item indices +@testset "partial sub-groups" begin + backend = oneAPIBackend() + @test KI.supports_subgroups(backend) + width = KI.sub_group_size(backend) + + # a (width+1)x2 work-group is made up of 3 sub-groups, the last one only partially filled + 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) + @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] + + # 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) +end + +function ki_fill!(A) + i = KI.get_global_id().x + if i <= length(A) + @inbounds A[i] = i + end + return +end + +@testset "launch keywords" begin + A = oneAPI.zeros(Int, 4) + 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())) + @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) +end + +@testset "versioninfo" begin + @test occursin("oneAPI.jl", sprint(KI.versioninfo, oneAPIBackend())) +end