From 4ef4772b155099570b285deacff896bf26786ded Mon Sep 17 00:00:00 2001 From: Valentin Churavy Date: Mon, 20 Jul 2026 14:30:13 +0200 Subject: [PATCH 1/5] Rework KernelInterface into a sibling package Extract the `KernelInterface` submodule (`src/interface.jl`) into a standalone subpackage under `lib/KernelInterface`, depended on by KernelAbstractions via a path source. KernelAbstractions keeps re-exporting `KernelInterface`/`KI` so existing references continue to work unchanged. KernelInterface no longer depends on GPUCompiler: the only two symbols used (`split_kwargs`, `assign_args!`) are small macro helpers, now vendored, leaving the interface package free of non-stdlib dependencies. Co-Authored-By: Claude Opus 4.8 --- Project.toml | 5 ++ lib/KernelInterface/Project.toml | 7 ++ .../KernelInterface/src/KernelInterface.jl | 68 ++++++++++++++++++- src/KernelAbstractions.jl | 4 +- 4 files changed, 79 insertions(+), 5 deletions(-) create mode 100644 lib/KernelInterface/Project.toml rename src/interface.jl => lib/KernelInterface/src/KernelInterface.jl (87%) diff --git a/Project.toml b/Project.toml index 5bd606365..c6b9bf465 100644 --- a/Project.toml +++ b/Project.toml @@ -8,6 +8,7 @@ Adapt = "79e6a3ab-5dfb-504d-930d-738a2a938a0e" Atomix = "a9b6321e-bd34-4604-b9c9-b65b8de01458" GPUCompiler = "61eb1bfa-7361-4325-ad38-22787b887f55" InteractiveUtils = "b77e0a4c-d291-57a0-90e8-8db25a27a240" +KernelInterface = "4ee993da-d684-4d17-a7dd-4e58e78d92bf" LLVM = "929cbde3-209d-540e-8aea-75f648917ca0" MacroTools = "1914dd2f-81c6-5fcd-8719-6d5c9610ff09" PrecompileTools = "aea7be01-6a6a-4083-8856-8a6e6704d82a" @@ -21,6 +22,9 @@ SPIRV_Tools_jll = "6ac6d60f-d740-5983-97d7-a4482c0689f4" UUIDs = "cf7118a7-6976-5b1a-9a39-7adc72f591a4" pocl_standalone_jll = "54f56a70-6062-5590-a942-1226658f6c83" +[sources] +KernelInterface = {path = "lib/KernelInterface"} + [weakdeps] EnzymeCore = "f151be2c-9106-41f4-ab19-57ee4f262869" LinearAlgebra = "37e2e46d-f89d-539d-b4ee-838fcccc9c8e" @@ -38,6 +42,7 @@ Atomix = "0.1, 1" EnzymeCore = "0.7, 0.8.1" GPUCompiler = "2" InteractiveUtils = "1.6" +KernelInterface = "0.1" LLVM = "9.9" LinearAlgebra = "1.6" MacroTools = "0.5" diff --git a/lib/KernelInterface/Project.toml b/lib/KernelInterface/Project.toml new file mode 100644 index 000000000..5051fddb7 --- /dev/null +++ b/lib/KernelInterface/Project.toml @@ -0,0 +1,7 @@ +name = "KernelInterface" +uuid = "4ee993da-d684-4d17-a7dd-4e58e78d92bf" +authors = ["Valentin Churavy and contributors"] +version = "0.1.0" + +[compat] +julia = "1.10" diff --git a/src/interface.jl b/lib/KernelInterface/src/KernelInterface.jl similarity index 87% rename from src/interface.jl rename to lib/KernelInterface/src/KernelInterface.jl index 6748c227a..a56a2f9d5 100644 --- a/src/interface.jl +++ b/lib/KernelInterface/src/KernelInterface.jl @@ -12,8 +12,70 @@ like allocating arrays on a backend. """ module KernelInterface -import ..KernelAbstractions: Backend -import GPUCompiler: split_kwargs, assign_args! +## macro tools +# Vendored from GPUCompiler to keep KernelInterface free of heavy dependencies. + +# split keyword arguments expressions into groups. returns vectors of keyword argument +# values, one more than the number of groups (unmatched keywords in the last vector). +# intended for use in macros; the resulting groups can be used in expressions. +# can be used at run time, but not in performance critical code. +function split_kwargs(kwargs, kw_groups...) + kwarg_groups = ntuple(_ -> [], length(kw_groups) + 1) + for kwarg in kwargs + # decode + if Meta.isexpr(kwarg, :(=)) + # use in macros + key, val = kwarg.args + elseif kwarg isa Pair{Symbol, <:Any} + # use in functions + key, val = kwarg + else + throw(ArgumentError("non-keyword argument like option '$kwarg'")) + end + isa(key, Symbol) || throw(ArgumentError("non-symbolic keyword '$key'")) + + # find a matching group + group = length(kwarg_groups) + for (i, kws) in enumerate(kw_groups) + if key in kws + group = i + break + end + end + push!(kwarg_groups[group], kwarg) + end + + return kwarg_groups +end + +# assign arguments to variables, handle splatting +function assign_args!(code, _args) + nargs = length(_args) + + # handle splatting + splats = Vector{Bool}(undef, nargs) + args = Vector{Any}(undef, nargs) + for i in 1:nargs + splats[i] = Meta.isexpr(_args[i], :(...)) + args[i] = splats[i] ? _args[i].args[1] : _args[i] + end + + # assign arguments to variables + vars = Vector{Symbol}(undef, nargs) + for i in 1:nargs + vars[i] = gensym() + push!(code.args, :($(vars[i]) = $(args[i]))) + end + + # convert the arguments, compile the function and call the kernel + # while keeping the original arguments alive + var_exprs = Vector{Any}(undef, nargs) + for i in 1:nargs + var_exprs[i] = splats[i] ? Expr(:(...), vars[i]) : vars[i] + end + + return vars, var_exprs +end """ get_global_size()::@NamedTuple{x::Int, y::Int, z::Int} @@ -219,7 +281,7 @@ Returns a vector of `DataType`s supported on `backend` Backend implementations **must** implement this function only if they support `shfl_down` for any types. """ -shfl_down_types(::Backend) = DataType[] +shfl_down_types(_) = DataType[] """ diff --git a/src/KernelAbstractions.jl b/src/KernelAbstractions.jl index 48329134d..0352d62f8 100644 --- a/src/KernelAbstractions.jl +++ b/src/KernelAbstractions.jl @@ -230,8 +230,8 @@ synchronize(backend) """ abstract type Backend end -include("interface.jl") -import .KernelInterface as KI +import KernelInterface +import KernelInterface as KI export KernelInterface ### From 3d4260e781d97b87322387144ff30a6ba0e73b55 Mon Sep 17 00:00:00 2001 From: Valentin Churavy Date: Mon, 27 Jul 2026 10:57:21 +0200 Subject: [PATCH 2/5] Move host `_print` fallback into KernelInterface Now that `KernelInterface` is a separate package, defining `KI._print(items...)` in KernelAbstractions is type piracy: neither the function nor any of the argument types are owned by KernelAbstractions, and Aqua's piracy check flags it. Move the generated host fallback next to the `_print` interface stub in KernelInterface, where it is a normal generic fallback. Backends keep overriding it with `@device_override _print(args...)`, unchanged. Co-Authored-By: Claude Opus 5 --- lib/KernelInterface/src/KernelInterface.jl | 21 ++++++++++++++++++++- src/KernelAbstractions.jl | 19 ------------------- 2 files changed, 20 insertions(+), 20 deletions(-) diff --git a/lib/KernelInterface/src/KernelInterface.jl b/lib/KernelInterface/src/KernelInterface.jl index a56a2f9d5..2415ef554 100644 --- a/lib/KernelInterface/src/KernelInterface.jl +++ b/lib/KernelInterface/src/KernelInterface.jl @@ -343,8 +343,27 @@ end ``` If the backend does not support printing, define it to return `nothing`. + +The generic fallback prints on the host, which keeps CPU backends working. +`Val` arguments are unwrapped, since `KernelAbstractions.@print` uses them to +pass literal strings through to backends that require compile-time format strings. """ -function _print end +@generated function _print(items...) + args = [] + + for i in 1:length(items) + item = :(items[$i]) + T = items[i] + if T <: Val + item = QuoteNode(T.parameters[1]) + end + push!(args, item) + end + + return quote + print($(args...)) + end +end """ diff --git a/src/KernelAbstractions.jl b/src/KernelAbstractions.jl index 0352d62f8..4866d4b29 100644 --- a/src/KernelAbstractions.jl +++ b/src/KernelAbstractions.jl @@ -402,25 +402,6 @@ macro context() return esc(:(__ctx__)) end -# Defined to keep cpu support for `__print` -@generated function KI._print(items...) - str = "" - args = [] - - for i in 1:length(items) - item = :(items[$i]) - T = items[i] - if T <: Val - item = QuoteNode(T.parameters[1]) - end - push!(args, item) - end - - return quote - print($(args...)) - end -end - """ @print(items...) From 908f82cb69e7a670c2d4bff8872f9af8386c3d25 Mon Sep 17 00:00:00 2001 From: Valentin Churavy Date: Mon, 27 Jul 2026 11:34:55 +0200 Subject: [PATCH 3/5] [KernelInterface] Terminate `localmemory` host fallback `localmemory(::Type{T}, dims) where {T}` forwards to `localmemory(T, Val(dims))`, but `dims` is untyped, so the `Val` call matches the same method again. Backends only supply the `Val` form via `@device_override`, which lives in the overlay method table, so in a kernel the override terminates the recursion -- but off device it recurses until the stack overflows and the process segfaults. This is reachable from the host: `KernelAbstractions.SharedMemory` passes a `Val` straight to `KI.localmemory`, so `@localmem` used outside a kernel segfaults rather than erroring. Add a terminating `Val` method that errors, matching `barrier()`. Co-Authored-By: Claude Opus 5 --- lib/KernelInterface/src/KernelInterface.jl | 17 ++++++++++++++++- 1 file changed, 16 insertions(+), 1 deletion(-) diff --git a/lib/KernelInterface/src/KernelInterface.jl b/lib/KernelInterface/src/KernelInterface.jl index 2415ef554..2d79b9b0e 100644 --- a/lib/KernelInterface/src/KernelInterface.jl +++ b/lib/KernelInterface/src/KernelInterface.jl @@ -250,6 +250,11 @@ Declare memory that is local to a workgroup. """ localmemory(::Type{T}, dims) where {T} = localmemory(T, Val(dims)) +# The `Val` form only exists in a backend's overlay method table, so off-device it +# would otherwise fall back to the forwarding method above and recurse forever. +localmemory(::Type{T}, ::Val) where {T} = + error("Local memory used outside kernel or not captured") + """ shfl_down(val::T, offset::Integer) where T @@ -389,6 +394,15 @@ struct Kernel{B, Kern} kern::Kern end +""" + check_launch_args(numworkgroups, workgroupsize) + +Validate the launch configuration passed to a [`Kernel`](@ref), throwing an +`ArgumentError` if either argument has more than 3 dimensions. + +Backends may call this from their kernel-launch method instead of writing their +own check. +""" function check_launch_args(numworkgroups, workgroupsize) length(numworkgroups) <= 3 || throw(ArgumentError("`numworkgroups` only accepts up to 3 dimensions")) @@ -479,7 +493,8 @@ function argconvert end KI.kernel_function(::NewBackend, f::F, tt::TT=Tuple{}; name=nothing, kwargs...) where {F,TT} Low-level interface to compile a function invocation for the currently-active GPU, returning -a callable kernel object. For a higher-level interface, use [`KI.@kernel`](@ref). +a callable kernel object. For a higher-level interface, use +[`KernelInterface.@kernel`](@ref). Currently, `kernel_function` only supports the `name` keyword argument as it is the only one by all backends. From 996da41198c376e384cb09a0f05eba9994680825 Mon Sep 17 00:00:00 2001 From: Valentin Churavy Date: Mon, 27 Jul 2026 11:35:05 +0200 Subject: [PATCH 4/5] [KernelInterface] Add a standalone test suite Tests the sibling package on its own, without KernelAbstractions or a backend, which also guards the claim that it stays dependency-free. Covers the vendored `split_kwargs`/`assign_args!` helpers, the host fallbacks (`barrier`, `shfl_down_types`, `multiprocessor_count`, `localmemory`, `_print`), `check_launch_args`, and drives `KI.@kernel` end to end against a mock backend, including its expansion-time error paths. Also asserts that the interface stubs have no methods, so a backend that forgets one gets a MethodError. Run in CI by a new `KernelInterface` job. Co-Authored-By: Claude Opus 5 --- .github/workflows/ci.yml | 24 ++++ lib/KernelInterface/test/Project.toml | 7 + lib/KernelInterface/test/runtests.jl | 187 ++++++++++++++++++++++++++ 3 files changed, 218 insertions(+) create mode 100644 lib/KernelInterface/test/Project.toml create mode 100644 lib/KernelInterface/test/runtests.jl diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 5a32aea13..50ce0764f 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -164,6 +164,30 @@ jobs: with: files: lcov.info + KernelInterface: + name: KernelInterface + runs-on: ${{ matrix.os }} + strategy: + fail-fast: false + matrix: + version: ['1.10', '1.11', '1.12', '1.13-nightly'] + os: [ubuntu-24.04, macOS-15, windows-2022] + steps: + - uses: actions/checkout@v7 + - uses: julia-actions/install-juliaup@v3 + with: + channel: ${{ matrix.version }} + - uses: julia-actions/cache@v3 + # Tested on its own, without KernelAbstractions, to keep the sibling + # package standalone and dependency-free. + - uses: julia-actions/julia-buildpkg@v1 + with: + project: lib/KernelInterface + - uses: julia-actions/julia-runtest@v1 + with: + project: lib/KernelInterface + annotate: true + OpenCL: name: OpenCL (POCL) runs-on: ubuntu-latest diff --git a/lib/KernelInterface/test/Project.toml b/lib/KernelInterface/test/Project.toml new file mode 100644 index 000000000..2f98bebc9 --- /dev/null +++ b/lib/KernelInterface/test/Project.toml @@ -0,0 +1,7 @@ +[deps] +Aqua = "4c88cf16-eb10-579e-8560-4a9242c79595" +KernelInterface = "4ee993da-d684-4d17-a7dd-4e58e78d92bf" +Test = "8dfed614-e22c-5e08-85e1-65c5234f0b40" + +[compat] +Aqua = "0.8" diff --git a/lib/KernelInterface/test/runtests.jl b/lib/KernelInterface/test/runtests.jl new file mode 100644 index 000000000..17aeb7e1d --- /dev/null +++ b/lib/KernelInterface/test/runtests.jl @@ -0,0 +1,187 @@ +using KernelInterface +using Aqua +using Test + +const KI = KernelInterface + +# `_print`'s host fallback writes to `stdout`, so capture it through a real file. +function capture_stdout(f) + return mktemp() do path, io + redirect_stdout(f, io) + flush(io) + return read(path, String) + end +end + +@testset "standalone" begin + # KernelInterface is what backends implement against, so it must stay loadable + # without dragging in KernelAbstractions or a compiler stack. + toml = read(joinpath(pkgdir(KernelInterface), "Project.toml"), String) + @test !occursin("[deps]", toml) + @test !occursin("[sources]", toml) +end + +# NOTE: this runs before the mock backend below defines methods on `argconvert` +# and `kernel_function`. +@testset "interface stubs" begin + # These have no fallback on purpose: a backend that forgets to `@device_override` + # them should get a MethodError rather than silently wrong behaviour. + stubs = [ + KI.get_global_size, KI.get_global_id, + KI.get_local_size, KI.get_local_id, + KI.get_num_groups, KI.get_group_id, + KI.get_sub_group_size, KI.get_max_sub_group_size, + KI.get_num_sub_groups, KI.get_sub_group_id, + KI.get_sub_group_local_id, + KI.shfl_down, + KI.kernel_max_work_group_size, KI.max_work_group_size, KI.sub_group_size, + KI.argconvert, KI.kernel_function, + ] + for stub in stubs + @test isempty(methods(stub)) + end +end + +@testset "host fallbacks" begin + # Barriers are meaningless off-device and must say so rather than no-op. + @test_throws "used outside kernel" KI.barrier() + @test_throws "used outside kernel" KI.sub_group_barrier() + + # Permissive defaults: a backend only implements these if it can do better. + @test KI.shfl_down_types(nothing) == DataType[] + @test KI.multiprocessor_count(nothing) == 0 + + # `localmemory` forwards the untyped `dims` to the `Val` form backends override. + # Off-device that form is unimplemented, and must error rather than recurse + # back into the forwarding method. + @test_throws "used outside kernel" KI.localmemory(Float32, (2, 2)) + @test_throws "used outside kernel" KI.localmemory(Float32, Val((2, 2))) +end + +@testset "_print" begin + # The host fallback keeps `KernelAbstractions.@print` working outside a kernel. + # `@print` wraps literals in `Val` so backends can use them as format strings; + # the fallback has to unwrap them again. + @test capture_stdout(() -> KI._print()) == "" + @test capture_stdout(() -> KI._print(Val(Symbol("hello\n")))) == "hello\n" + @test capture_stdout(() -> KI._print(1, 2)) == "12" + @test capture_stdout(() -> KI._print(Val(Symbol("x = ")), 42, Val(Symbol("\n")))) == + "x = 42\n" + @test capture_stdout(() -> KI._print(Val(3), " ", Val(:sym))) == "3 sym" +end + +@testset "check_launch_args" begin + @test KI.check_launch_args(1, 1) === nothing + @test KI.check_launch_args((1, 2, 3), (1, 2, 3)) === nothing + @test_throws ArgumentError KI.check_launch_args((1, 2, 3, 4), 1) + @test_throws ArgumentError KI.check_launch_args(1, (1, 2, 3, 4)) +end + +@testset "Kernel" begin + kernel = KI.Kernel(:backend, :kern) + @test kernel.backend === :backend + @test kernel.kern === :kern +end + +@testset "split_kwargs" begin + kwargs = [:(launch = false), :(name = "foo"), :(numworkgroups = 2)] + macro_kw, compiler_kw, launch_kw, other = KI.split_kwargs( + kwargs, KI.MACRO_KWARGS, KI.COMPILER_KWARGS, KI.LAUNCH_KWARGS + ) + @test macro_kw == [:(launch = false)] + @test compiler_kw == [:(name = "foo")] + @test launch_kw == [:(numworkgroups = 2)] + @test isempty(other) + + # Unmatched keywords land in the trailing group rather than erroring. + _, unmatched = KI.split_kwargs([:(bogus = 1)], [:launch]) + @test unmatched == [:(bogus = 1)] + + # Also usable at run time with pairs instead of expressions. + matched, _ = KI.split_kwargs([:launch => false], [:launch]) + @test matched == [:launch => false] + + @test_throws ArgumentError KI.split_kwargs([:(f(x))], [:launch]) + @test_throws ArgumentError KI.split_kwargs([Expr(:(=), 1, 2)], [:launch]) +end + +@testset "assign_args!" begin + code = Expr(:block) + vars, var_exprs = KI.assign_args!(code, [:a, :(b...)]) + @test length(vars) == 2 + # Arguments are hoisted into gensyms so the caller can `GC.@preserve` them. + @test code.args == [:($(vars[1]) = a), :($(vars[2]) = b)] + @test var_exprs[1] === vars[1] + @test var_exprs[2] == Expr(:..., vars[2]) +end + +# A minimal backend, exercising the contract `KI.@kernel` expects of one. +struct MockBackend end + +struct MockKernel + f::Any + tt::Any + name::Any + launches::Vector{Any} +end + +KI.argconvert(::MockBackend, arg) = arg +function KI.kernel_function(::MockBackend, f, tt = Tuple{}; name = nothing, kwargs...) + return MockKernel(f, tt, name, []) +end +function (kernel::MockKernel)(args...; kwargs...) + push!(kernel.launches, (args, Dict(kwargs))) + return nothing +end + +dummy(a, b) = nothing + +@testset "@kernel" begin + backend = MockBackend() + + kernel = KI.@kernel backend numworkgroups = 2 workgroupsize = 4 dummy(1, 2.0) + @test kernel isa MockKernel + @test kernel.f === dummy + @test kernel.tt == Tuple{Int, Float64} + args, launch_kwargs = only(kernel.launches) + @test args == (1, 2.0) + @test launch_kwargs == Dict(:numworkgroups => 2, :workgroupsize => 4) + + # `launch=false` compiles only; the caller launches later. + deferred = KI.@kernel backend launch = false dummy(1, 2.0) + @test isempty(deferred.launches) + + # Compiler kwargs reach `kernel_function` instead of the launch. + named = KI.@kernel backend launch = false name = "mykernel" dummy(1, 2.0) + @test named.name == "mykernel" + + # Splatted arguments are supported. + splatted = KI.@kernel backend launch = false dummy((1, 2.0)...) + @test splatted.tt == Tuple{Int, Float64} + + @testset "errors" begin + # These throw during macro expansion, so they cannot be written as a plain + # `@test_throws` call. `macroexpand` wraps such errors in a `LoadError`. + function expansion_error(ex) + try + macroexpand(@__MODULE__, ex) + catch err + return err isa LoadError ? err.error : err + end + return nothing + end + + @test expansion_error(:(KI.@kernel backend bogus = 1 dummy(1))) isa ArgumentError + @test expansion_error(:(KI.@kernel backend dummy)) isa ArgumentError + @test expansion_error(:(KI.@kernel backend launch = 1 dummy(1))) isa ArgumentError + @test expansion_error(:(KI.@kernel backend "notakwarg" dummy(1))) isa ArgumentError + # launch-time kwargs are meaningless when we are not launching + @test expansion_error( + :(KI.@kernel backend launch = false numworkgroups = 2 dummy(1)) + ) isa ErrorException + end +end + +@testset "Aqua" begin + Aqua.test_all(KernelInterface) +end From 417316bb5cbc17d5c6264be18035f0033562d13b Mon Sep 17 00:00:00 2001 From: Valentin Churavy Date: Mon, 27 Jul 2026 11:35:11 +0200 Subject: [PATCH 5/5] [KernelInterface] Document the interface Adds a KernelInterface manual page covering the device-side API, the host-side API, and what a backend has to implement, with `@docs` blocks for every KernelInterface docstring. Fixes #730. Co-Authored-By: Claude Opus 5 --- docs/Project.toml | 2 + docs/make.jl | 4 +- docs/src/kernelinterface.md | 143 ++++++++++++++++++++++++++++++++++++ 3 files changed, 148 insertions(+), 1 deletion(-) create mode 100644 docs/src/kernelinterface.md diff --git a/docs/Project.toml b/docs/Project.toml index 96dfe1047..24ef1b15c 100644 --- a/docs/Project.toml +++ b/docs/Project.toml @@ -1,9 +1,11 @@ [deps] Documenter = "e30172f5-a6a5-5a46-863b-614d45cd2de4" KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c" +KernelInterface = "4ee993da-d684-4d17-a7dd-4e58e78d92bf" [compat] Documenter = "1" [sources] KernelAbstractions = {path = ".."} +KernelInterface = {path = "../lib/KernelInterface"} diff --git a/docs/make.jl b/docs/make.jl index 57e427c1f..5300339ee 100644 --- a/docs/make.jl +++ b/docs/make.jl @@ -1,11 +1,12 @@ using KernelAbstractions +using KernelInterface using Documenter function main() ci = get(ENV, "CI", "") == "true" makedocs(; - modules = [KernelAbstractions], + modules = [KernelAbstractions, KernelInterface], authors = "JuliaGPU and contributors", repo = "https://github.com/JuliaGPU/KernelAbstractions.jl/blob/{commit}{path}#L{line}", sitename = "KernelAbstractions.jl", @@ -35,6 +36,7 @@ function main() "examples/atomix.md", ], # Examples "API" => "api.md", + "KernelInterface" => "kernelinterface.md", "Extras" => [ "extras/unrolling.md", "extras/pocl_debugging.md", diff --git a/docs/src/kernelinterface.md b/docs/src/kernelinterface.md new file mode 100644 index 000000000..e0e68b297 --- /dev/null +++ b/docs/src/kernelinterface.md @@ -0,0 +1,143 @@ +# [KernelInterface](@id kernelinterface) + +```@meta +CurrentModule = KernelInterface +``` + +`KernelInterface` (conventionally imported as `KI`) is the low-level API that +backends implement, and that `KernelAbstractions` builds its higher-level kernel +language on top of. + +It ships as a standalone package under `lib/KernelInterface` with **no +dependencies outside the standard library**, so a backend can implement the +interface without taking on `KernelAbstractions` or its compiler stack: + +```julia +using KernelInterface +const KI = KernelInterface +``` + +`KernelAbstractions` re-exports it, so `KernelAbstractions.KernelInterface` and +`KernelAbstractions.KI` refer to the same module. + +!!! note + Most of the functions below are stubs with no methods. They exist so that + backends can add device-side implementations with + `GPUCompiler.@device_override`, and so kernels can call them generically. + Calling one without a backend that implements it is a `MethodError`. + +```@docs +KernelInterface +``` + +## Device-side API + +These are called from inside a kernel. A backend provides each one with + +```julia +@device_override KI.get_global_id() = ... +``` + +along with the corresponding on-device functionality. + +### Indexing + +All index queries are **1-based** and return a named tuple of `x`, `y` and `z` +components. + +```@docs +get_global_size +get_global_id +get_local_size +get_local_id +get_num_groups +get_group_id +``` + +### Sub-groups + +```@docs +get_sub_group_size +get_max_sub_group_size +get_num_sub_groups +get_sub_group_id +get_sub_group_local_id +``` + +### Barriers + +```@docs +barrier +sub_group_barrier +``` + +### Memory + +```@docs +localmemory +``` + +### Communication + +```@docs +shfl_down +shfl_down_types +``` + +### Printing + +```@docs +KernelInterface._print +``` + +`_print` is the one device-side function with a working host fallback: it prints +its arguments with `Base.print`, unwrapping any `Val`-wrapped literals. That is +what makes [`KernelAbstractions.@print`](@ref) usable outside of a kernel. + +## Host-side API + +### Backend queries + +```@docs +max_work_group_size +sub_group_size +multiprocessor_count +``` + +### Compilation and launching + +```@docs +Kernel +kernel_function +kernel_max_work_group_size +check_launch_args +argconvert +KernelInterface.@kernel +``` + +!!! note + `KI.@kernel` is **not** `KernelAbstractions.@kernel`. `KI.@kernel` wraps a + backend's own compile-and-launch path — the equivalent of `@cuda` or + `@metal` — and prefixes a *call*. [`KernelAbstractions.@kernel`](@ref) + prefixes a *definition* and produces a kernel written in the higher-level + KernelAbstractions language. + +## Implementing a backend + +A backend must, at minimum: + +1. `@device_override` the device-side functions it supports. The indexing + queries and [`barrier`](@ref) are required; sub-group and + [`shfl_down`](@ref) support is optional. +2. Implement [`argconvert`](@ref) and [`kernel_function`](@ref) for its backend + type, returning a [`Kernel`](@ref). +3. Make that `Kernel` callable, accepting `numworkgroups` and `workgroupsize` as + a scalar `Integer` or a 1-, 2- or 3-element tuple. Use + [`check_launch_args`](@ref) to validate them, or check them directly. +4. Report its limits through [`kernel_max_work_group_size`](@ref) and, where + applicable, [`max_work_group_size`](@ref), [`sub_group_size`](@ref) and + [`multiprocessor_count`](@ref). + +The PoCL backend in `src/pocl/backend.jl` is a complete worked example. + +See also the [notes for backend implementations](@ref implementations_notes).