From 206843f0289295cb6d9a67b7eab83927d30f5741 Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Sun, 27 Sep 2026 13:24:28 +0200 Subject: [PATCH 1/5] Inline the untyped KernelInterface index queries KernelAbstractions' index functions call the zero-argument forms, which weren't inlined on CUDA: every work-item made two calls, returning the ids through local memory, making e.g. broadcasts up to 2x slower. --- lib/KernelInterface/src/device.jl | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/lib/KernelInterface/src/device.jl b/lib/KernelInterface/src/device.jl index 60a62fa09..ade52fb33 100644 --- a/lib/KernelInterface/src/device.jl +++ b/lib/KernelInterface/src/device.jl @@ -11,7 +11,7 @@ Return the number of global work-items specified as a tuple of type `T`. ``` The zero-argument form forwards to `get_global_size(Int)`. """ -get_global_size() = get_global_size(Int) +@inline get_global_size() = get_global_size(Int) """ get_global_id([::Type{T}=Int])::@NamedTuple{x::T, y::T, z::T} @@ -28,7 +28,7 @@ Returns the unique global work-item ID as a tuple of type `T`. `T` defaults to ` ``` The zero-argument form forwards to `get_global_id(Int)`. """ -get_global_id() = get_global_id(Int) +@inline get_global_id() = get_global_id(Int) """ get_local_size([::Type{T}=Int])::@NamedTuple{x::T, y::T, z::T} @@ -43,7 +43,7 @@ Return the number of local work-items specified as a tuple of type `T`. ``` The zero-argument form forwards to `get_local_size(Int)`. """ -get_local_size() = get_local_size(Int) +@inline get_local_size() = get_local_size(Int) """ get_local_id([::Type{T}=Int])::@NamedTuple{x::T, y::T, z::T} @@ -60,7 +60,7 @@ Returns the unique local work-item ID as a tuple of type `T`. `T` defaults to `I ``` The zero-argument form forwards to `get_local_id(Int)`. """ -get_local_id() = get_local_id(Int) +@inline get_local_id() = get_local_id(Int) """ get_num_groups([::Type{T}=Int])::@NamedTuple{x::T, y::T, z::T} @@ -74,7 +74,7 @@ Returns the number of groups as a tuple of type `T`. `T` defaults to `Int`. ``` The zero-argument form forwards to `get_num_groups(Int)`. """ -get_num_groups() = get_num_groups(Int) +@inline get_num_groups() = get_num_groups(Int) """ get_group_id([::Type{T}=Int])::@NamedTuple{x::T, y::T, z::T} @@ -91,7 +91,7 @@ Returns the unique group ID as a tuple of type `T`. `T` defaults to `Int`. ``` The zero-argument form forwards to `get_group_id(Int)`. """ -get_group_id() = get_group_id(Int) +@inline get_group_id() = get_group_id(Int) """ get_sub_group_size()::UInt32 From b2e9daaaf6ac80405961eb3d8e62f9502749e2da Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Sun, 27 Sep 2026 13:16:49 +0200 Subject: [PATCH 2/5] KernelInterface: add per-dimension launch limits Add `max_work_group_dims` and `max_num_groups`, and respect the former when distributing work-items over a workgroup, so that e.g. an ndrange of (1, 1, 5000) no longer gets a (1, 1, 1024) workgroup that CUDA can't launch. Also specify that the typed index queries compute in `T`. POCL reports its per-dimension workgroup limits, and the testsuite checks that back ends can launch at their reported limits and that automatically sized launches respect them. --- Project.toml | 2 +- docs/src/kernelinterface.md | 2 + lib/KernelInterface/Project.toml | 2 +- lib/KernelInterface/src/device.jl | 12 ++++++ lib/KernelInterface/src/launch.jl | 60 +++++++++++++++++++++++---- lib/KernelInterface/test/interface.jl | 40 ++++++++++++++++++ lib/KernelInterface/test/runtests.jl | 22 ++++++++++ src/pocl/backend.jl | 14 ++++++- 8 files changed, 141 insertions(+), 13 deletions(-) diff --git a/Project.toml b/Project.toml index 98d30cd1f..48c72ffb8 100644 --- a/Project.toml +++ b/Project.toml @@ -42,7 +42,7 @@ Adapt = "0.4, 1.0, 2.0, 3.0, 4" Atomix = "1.2.1" EnzymeCore = "0.7, 0.8.1" GPUCompiler = "2.7" -KernelInterface = "0.2" +KernelInterface = "0.2.3" LLVM = "9.9" LinearAlgebra = "1.6" MacroTools = "0.5" diff --git a/docs/src/kernelinterface.md b/docs/src/kernelinterface.md index 8714fcf0e..0c09aa088 100644 --- a/docs/src/kernelinterface.md +++ b/docs/src/kernelinterface.md @@ -171,6 +171,8 @@ supports_float64 ```@docs max_work_group_size +max_work_group_dims +max_num_groups sub_group_size multiprocessor_count ``` diff --git a/lib/KernelInterface/Project.toml b/lib/KernelInterface/Project.toml index 936a21459..e1660a954 100644 --- a/lib/KernelInterface/Project.toml +++ b/lib/KernelInterface/Project.toml @@ -1,7 +1,7 @@ name = "KernelInterface" uuid = "4ee993da-d684-4d17-a7dd-4e58e78d92bf" authors = ["Valentin Churavy and contributors"] -version = "0.2.2" +version = "0.2.3" [compat] julia = "1.10" diff --git a/lib/KernelInterface/src/device.jl b/lib/KernelInterface/src/device.jl index ade52fb33..8814544f1 100644 --- a/lib/KernelInterface/src/device.jl +++ b/lib/KernelInterface/src/device.jl @@ -4,6 +4,8 @@ Return the number of global work-items specified as a tuple of type `T`. `T` defaults to `Int`. +The value is computed in `T`, and is undefined when it does not fit in `T`. + !!! note Backend implementations **must** implement: ``` @@ -18,6 +20,8 @@ Return the number of global work-items specified as a tuple of type `T`. Returns the unique global work-item ID as a tuple of type `T`. `T` defaults to `Int`. +The value is computed in `T`, and is undefined when it does not fit in `T`. + !!! note 1-based. @@ -36,6 +40,8 @@ Returns the unique global work-item ID as a tuple of type `T`. `T` defaults to ` Return the number of local work-items specified as a tuple of type `T`. `T` defaults to `Int`. +The value is computed in `T`, and is undefined when it does not fit in `T`. + !!! note Backend implementations **must** implement: ``` @@ -50,6 +56,8 @@ Return the number of local work-items specified as a tuple of type `T`. Returns the unique local work-item ID as a tuple of type `T`. `T` defaults to `Int`. +The value is computed in `T`, and is undefined when it does not fit in `T`. + !!! note 1-based. @@ -67,6 +75,8 @@ Returns the unique local work-item ID as a tuple of type `T`. `T` defaults to `I Returns the number of groups as a tuple of type `T`. `T` defaults to `Int`. +The value is computed in `T`, and is undefined when it does not fit in `T`. + !!! note Backend implementations **must** implement: ``` @@ -81,6 +91,8 @@ Returns the number of groups as a tuple of type `T`. `T` defaults to `Int`. Returns the unique group ID as a tuple of type `T`. `T` defaults to `Int`. +The value is computed in `T`, and is undefined when it does not fit in `T`. + !!! note 1-based. diff --git a/lib/KernelInterface/src/launch.jl b/lib/KernelInterface/src/launch.jl index 5a79ae9d1..a5126af13 100644 --- a/lib/KernelInterface/src/launch.jl +++ b/lib/KernelInterface/src/launch.jl @@ -55,14 +55,26 @@ function check_launch_args(numworkgroups, workgroupsize, ndrange) return end -function threads_to_workgroupsize(threads, ndrange) - total = Ref(1) - return map(ndrange) do n - # each dimension must be at least 1, even for a zero-sized ndrange - x = max(min(div(threads, total[]), n), 1) - total[] *= x - return x - end +""" + threads_to_workgroupsize(threads, ndrange, [limits]) + +Distribute `threads` work-items over the dimensions of `ndrange`, filling the first +dimension first. Dimension `d` gets at most `limits[d]` work-items; dimensions past the end +of `limits` are only bounded by `threads`. + +Every dimension gets at least one work-item, even for a zero-sized `ndrange`. +""" +threads_to_workgroupsize(threads, ndrange::Tuple, limits = ()) = + _threads_to_workgroupsize(threads, 1, ndrange, limits) +threads_to_workgroupsize(threads, ndrange::Integer, limits = ()) = + only(threads_to_workgroupsize(threads, (ndrange,), limits)) +# written recursively, because a closure updating the running total would box it +_threads_to_workgroupsize(threads, total, ::Tuple{}, limits) = () +function _threads_to_workgroupsize(threads, total, ndrange::Tuple, limits) + limit = isempty(limits) ? typemax(Int) : first(limits) + x = max(min(div(threads, total), first(ndrange), limit), 1) + rest = isempty(limits) ? () : Base.tail(limits) + return (x, _threads_to_workgroupsize(threads, total * x, Base.tail(ndrange), rest)...) end """ @@ -86,7 +98,7 @@ writing their own heuristic for calculating launch size. else workgroupsize = if workgroupsize == () max_wgs = kernel_max_work_group_size(kernel; max_work_items = min(prod(ndrange), max_work_items)) - threads_to_workgroupsize(max_wgs, ndrange) + threads_to_workgroupsize(max_wgs, ndrange, max_work_group_dims(kernel.backend)) else workgroupsize end @@ -130,6 +142,36 @@ kernel launch with too big a workgroup is attempted. """ function max_work_group_size end +""" + max_work_group_dims(backend)::NTuple{3, Int} + +The maximum number of work-items along each dimension of a workgroup, for the currently +active device of `backend`. [`max_work_group_size`](@ref) bounds their product. + +!!! note + Backend implementations **should** implement: + ``` + max_work_group_dims(backend::NewBackend)::NTuple{3, Int} + ``` + The fallback does not limit individual dimensions. +""" +max_work_group_dims(::Backend) = (typemax(Int), typemax(Int), typemax(Int)) + +""" + max_num_groups(backend)::NTuple{3, Int} + +The maximum number of workgroups along each dimension of a launch, for the currently +active device of `backend`. + +!!! note + Backend implementations **should** implement: + ``` + max_num_groups(backend::NewBackend)::NTuple{3, Int} + ``` + The fallback does not limit the number of workgroups. +""" +max_num_groups(::Backend) = (typemax(Int), typemax(Int), typemax(Int)) + """ sub_group_size(backend)::Int diff --git a/lib/KernelInterface/test/interface.jl b/lib/KernelInterface/test/interface.jl index ff838302c..30172bf72 100644 --- a/lib/KernelInterface/test/interface.jl +++ b/lib/KernelInterface/test/interface.jl @@ -186,6 +186,46 @@ function interface_testsuite(backend, AT) @test_throws ArgumentError (KI.@kernel backend() numworkgroups = (2, 2, 2) workgroupsize = (2, 2, 2, 2) launch_kernel3d(arr3d)) end + @testset "Launch limits" begin + max_dims = KI.max_work_group_dims(backend()) + max_groups = KI.max_num_groups(backend()) + @test max_dims isa NTuple{3, Int} && all(>=(1), max_dims) + @test max_groups isa NTuple{3, Int} && all(>=(1), max_groups) + + function fill_kernel(arr) + i, j, k = KI.get_global_id() + if i <= size(arr, 1) && j <= size(arr, 2) && k <= size(arr, 3) + @inbounds arr[i, j, k] = 1.0f0 + end + return + end + kernel = KI.@kernel backend() launch = false fill_kernel(AT(zeros(Float32, 1, 1, 1))) + function fill_test(dims; kwargs...) + arr = AT(zeros(Float32, dims)) + kernel(arr; kwargs...) + KI.synchronize(backend()) + return all(Array(arr) .== 1) + end + + # automatically chosen workgroup sizes respect the per-dimension limits + @testset "ndrange = $ndrange" for ndrange in ((1, 1, 5000), (1, 5000, 1), (1, 3, 2000)) + @test fill_test(ndrange; ndrange) + end + + # the reported limits can be launched + max_items = KI.kernel_max_work_group_size(kernel) + @testset "dimension $d" for d in 1:3 + items = min(max_dims[d], max_items) + workgroupsize = ntuple(i -> i == d ? items : 1, 3) + @test fill_test(workgroupsize; workgroupsize, numworkgroups = (1, 1, 1)) + + # don't launch (practically) unlimited grids + groups = min(max_groups[d], 2^16) + numworkgroups = ntuple(i -> i == d ? groups : 1, 3) + @test fill_test(numworkgroups; workgroupsize = (1, 1, 1), numworkgroups) + end + end + @testset "Host return types" begin b = backend() diff --git a/lib/KernelInterface/test/runtests.jl b/lib/KernelInterface/test/runtests.jl index 53a8be54a..944e5648c 100644 --- a/lib/KernelInterface/test/runtests.jl +++ b/lib/KernelInterface/test/runtests.jl @@ -219,6 +219,18 @@ end @test KI.threads_to_workgroupsize(256, (0, 4)) == (1, 4) @test KI.threads_to_workgroupsize(256, (4, 0)) == (4, 1) @test KI.threads_to_workgroupsize(0, (5,)) == (1,) + + # Per-dimension limits, as for CUDA's (1024, 1024, 64) blocks. + @test KI.threads_to_workgroupsize(1024, (1, 1, 5000), (1024, 1024, 64)) == (1, 1, 64) + @test KI.threads_to_workgroupsize(1024, (2000, 3), (512, 1024, 64)) == (512, 2) + # dimensions past the limits are only bounded by the thread budget + @test KI.threads_to_workgroupsize(64, (1, 1, 1, 100), (1024, 1024, 64)) == (1, 1, 1, 64) +end + +@testset "per-dimension limits" begin + # Unlimited unless the backend says otherwise. + @test KI.max_work_group_dims(StubBackend()) == (typemax(Int), typemax(Int), typemax(Int)) + @test KI.max_num_groups(StubBackend()) == (typemax(Int), typemax(Int), typemax(Int)) end @testset "Kernel" begin @@ -236,7 +248,17 @@ function KI.kernel_max_work_group_size(k::KI.Kernel{SizedBackend}; max_work_item return min(k.backend.maxThreads, max_work_items) end +# ... and a fixed limit per workgroup dimension +struct DimsBackend <: KI.Backend end +KI.kernel_max_work_group_size(k::KI.Kernel{DimsBackend}; max_work_items::Int = typemax(Int)) = + min(1024, max_work_items) +KI.max_work_group_dims(::DimsBackend) = (1024, 1024, 64) + @testset "auto_launch_sizes" begin + # the per-dimension limit is respected + @test KI.auto_launch_sizes(KI.Kernel(DimsBackend(), nothing), (), (), (1, 1, 5000)) === + ((1, 1, 79), (1, 1, 64)) + kernel = KI.Kernel(SizedBackend(256), nothing) # Without an ndrange the sizes pass through, defaulting to 1. diff --git a/src/pocl/backend.jl b/src/pocl/backend.jl index 0c1296571..fc0d4b332 100644 --- a/src/pocl/backend.jl +++ b/src/pocl/backend.jl @@ -234,9 +234,19 @@ function KI.kernel_max_work_group_size(kernel::KI.Kernel{<:POCLBackend}; max_wor wginfo = cl.work_group_info(kernel.kern.fun, device()) return Int(min(wginfo.size, max_work_items)) end -function KI.max_work_group_size(::POCLBackend)::Int - return Int(device().max_work_group_size) +# querying the device allocates, so cache the limits that every launch needs +function device_limits() + return get!(task_local_storage(), :POCLLimits) do + dev = device() + sizes = dev.max_work_item_size + (; + max_work_group_size = Int(dev.max_work_group_size), + max_work_group_dims = ntuple(d -> d <= length(sizes) ? sizes[d] : 1, 3), + ) + end::@NamedTuple{max_work_group_size::Int, max_work_group_dims::NTuple{3, Int}} end +KI.max_work_group_size(::POCLBackend)::Int = device_limits().max_work_group_size +KI.max_work_group_dims(::POCLBackend)::NTuple{3, Int} = device_limits().max_work_group_dims function KI.sub_group_size(::POCLBackend)::Int # POCL can technically support any sub_group size. # Check for common values used on GPUs then From 370ed77c89fe2c7c408b4a7b48758dbeefabd0ad Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Sun, 27 Sep 2026 22:11:03 +0200 Subject: [PATCH 3/5] Run the KernelInterface testsuite on POCL Only CUDA.jl ran it so far. Skip the events tests, which need asynchronous launches. --- test/runtests.jl | 12 ++++++++++++ 1 file changed, 12 insertions(+) diff --git a/test/runtests.jl b/test/runtests.jl index 27c1334ca..728feff1c 100644 --- a/test/runtests.jl +++ b/test/runtests.jl @@ -25,6 +25,18 @@ KernelAbstractions.versioninfo(POCLBackend()) import KernelAbstractions.POCL: POCL, @opencl, @device_code_llvm +module KernelInterfaceTests + import KernelInterface + using Test + include(joinpath(pkgdir(KernelInterface), "test", "testsuite.jl")) +end +@testset "POCL KernelInterface" begin + # POCL launches synchronously, so there's nothing for the events tests to order + KernelInterfaceTests.Testsuite.testsuite( + POCLBackend, "POCL", POCL, Array, POCL.CLDeviceArray; skip_tests = Set(["Events"]) + ) +end + @testset "POCL float atomics" begin # pocl's CPU device natively supports float add and min/max atomics in both global # and local memory, so the SPIR-V extensions guarding them must be permitted From 074f7fa28f75f29a3ee888f2ee8245599063647a Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Sun, 27 Sep 2026 15:27:10 +0200 Subject: [PATCH 4/5] CI: develop KernelInterface before building the package On Julia 1.10, which ignores [sources], building the package already resolves the environment, which fails if KernelAbstractions requires a KernelInterface version that isn't registered yet. --- .github/workflows/ci.yml | 17 +++++++++++------ 1 file changed, 11 insertions(+), 6 deletions(-) diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 25231d318..5a2b0fee7 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -142,9 +142,6 @@ jobs: # end # end # ' - # Nightly tracks the next Julia release; don't fail CI on it - - uses: julia-actions/julia-buildpkg@v1 - continue-on-error: ${{ matrix.version == 'nightly' }} # Known limitation: `[sources]` is only supported from Julia 1.11 on, so on 1.10 # the `KernelInterface = {path = "lib/KernelInterface"}` entry in Project.toml is # silently ignored and Pkg resolves KernelInterface from the registry instead -- @@ -152,10 +149,15 @@ jobs: # explicitly so 1.10 tests the copy in this repo, like 1.11+ already do. # # Required on every platform. Windows runs its tests de-escalated below, but that - # applies to the test run only, not to `Pkg.develop`. + # applies to the test run only, not to `Pkg.develop`. Do this before building, which + # already resolves the environment (and fails if this repo requires an unregistered + # version of KernelInterface). - name: Dev KernelInterface shell: bash run: julia -e 'using Pkg; Pkg.activate("."); Pkg.develop(; path="lib/KernelInterface")' + # Nightly tracks the next Julia release; don't fail CI on it + - uses: julia-actions/julia-buildpkg@v1 + continue-on-error: ${{ matrix.version == 'nightly' }} - uses: julia-actions/julia-runtest@v1 if: runner.os != 'Windows' continue-on-error: ${{ matrix.version == 'nightly' }} @@ -228,10 +230,13 @@ jobs: with: channel: ${{ matrix.version }} - uses: julia-actions/cache@v3 - - uses: julia-actions/julia-buildpkg@v1 - continue-on-error: ${{ matrix.version == 'nightly' }} # KernelInterface is dev'ed explicitly here for the same reason as in the CI job # above: Julia 1.10 does not support `[sources]`. + - name: Dev KernelInterface + shell: bash + run: julia -e 'using Pkg; Pkg.activate("."); Pkg.develop(; path="lib/KernelInterface")' + - uses: julia-actions/julia-buildpkg@v1 + continue-on-error: ${{ matrix.version == 'nightly' }} - name: "Instantiating project" run: | julia -e 'println("--- :julia: Instantiating project") From 68c294da7b11c1fa80c28883e760e06718899023 Mon Sep 17 00:00:00 2001 From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com> Date: Mon, 28 Sep 2026 00:05:17 -0300 Subject: [PATCH 5/5] Oops --- .buildkite/pipeline.yml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.buildkite/pipeline.yml b/.buildkite/pipeline.yml index 2211c1bee..8b1b7b029 100644 --- a/.buildkite/pipeline.yml +++ b/.buildkite/pipeline.yml @@ -23,7 +23,7 @@ steps: PackageSpec(; name="KernelInterface", path=joinpath(pwd(), "lib", "KernelInterface"))])' || exit 3 julia -e 'println("--- :julia: Developing CUDA") using Pkg - url="https://github.com/christiangnrd/CUDA.jl" + url="https://github.com/JuliaGPU/CUDA.jl" rev="intrinsics" subdirs = ["CUDACore", "CUDATools", "lib/cusolver", "lib/curand", "lib/cublas", "lib/cufft", "lib/cusparse", "lib/cupti", "lib/nvml", "lib/custatevec", "lib/cudnn", "lib/cutensor", "lib/cutensornet"] Pkg.add([(;url, rev); [(; url, rev, subdir) for subdir in subdirs]])'