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 CUDA.jl to that model in one step. It replaces #3246, which added KernelInterface next to the existing KernelAbstractions 0.9 back end, and #3302/#3304, which then ported that to KernelAbstractions 0.10.

CUDABackend now implements KernelInterface, in CUDACore, which depends on KernelInterface instead of KernelAbstractions. That makes the KernelInterface layer usable on its own, without KernelAbstractions' macros:

using CUDA
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 = CUDA.rand(1000); b = CUDA.rand(1000); c = similar(a)
KI.@launch CUDABackend() 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 the maxthreads hint for kernels with a static workgroup size. CUDA's own copy of the KernelAbstractions launch path goes away: partitioning the ndrange, building the kernel's context, tuning the workgroup size, and calling the kernel. What remains specific to KernelAbstractions is a 37-line extension.

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 for the generated PTX. With a KA kernel that copies a Float32 array, on an RTX 5080:

	main (KA 0.9)	this PR
256×256×256, @index(Global, Cartesian)	180 µs	161 µs
256×256×256, @index(Global, Linear)	175 µs	149 µs
255×257×129, @index(Global, Cartesian)	86 µs	41 µs
255×257×129, @index(Global, Linear)	80 µs	37 µs
launch of an empty kernel	1.72 µs	1.89 µs
launch of a kernel with 40 arguments	1.95 µs, 336 bytes	1.98 µs, 304 bytes

CUDABackend(; prefer_blocks=true) still prefers more, smaller blocks. It now takes effect in KI.launch_configuration, which receives the number of work-items to cover.

Launching kernels with many arguments stays as cheap as with #3309. KernelInterface 0.4 passes the arguments to the back end as a tuple (JuliaGPU/KernelAbstractions.jl#811), and CUDA forwards that tuple to its own launch. KI.launch rejects threads and blocks, which would override the launch geometry KernelInterface has validated, but passes CUDA's other launch options on:

kernel = KI.@launch CUDABackend() launch=false vadd(c, a, b)
kernel(c, a, b; ndrange=length(c), stream=stream())   # fine
kernel(c, a, b; ndrange=length(c), threads=256)       # ArgumentError

The first commit is the port. The next ones implement parts of KernelInterface that are optional: sub-groups (warps, including a partial last warp), ordering work across tasks with CUDA events for KernelAbstractions.@spawn, and KI.versioninfo.

One regression remains. KernelAbstractions converts the arguments twice, once to determine the argument types to compile for and once when launching, where CUDA converted them once. cudaconvert is supposed to be pure, so this only shows with a conversion that has side effects. Two tests that count conversions are now @test_broken.

This needs a breaking release, since it requires KernelAbstractions 0.10, which isn't registered yet. Until then, the last commit takes KernelAbstractions and KernelInterface from their development branch, through [sources] and, for Julia 1.10 and 1.11, which don't pick those up, by developing them explicitly on Buildkite. That commit is dropped before merging.

[github-actions]

CUDA.jl Benchmarks

Details

Benchmark suite	Current: 19a22c6	Previous: a3202f2	Ratio
array/accumulate/Float32/1d	99292 ns	99028 ns	1.00
array/accumulate/Float32/dims=1	73540 ns	72921 ns	1.01
array/accumulate/Float32/dims=1L	1589481 ns	1588681 ns	1.00
array/accumulate/Float32/dims=2	139277 ns	138806 ns	1.00
array/accumulate/Float32/dims=2L	656037 ns	655618 ns	1.00
array/accumulate/Int64/1d	118658 ns	118327 ns	1.00
array/accumulate/Int64/dims=1	77559 ns	76942 ns	1.01
array/accumulate/Int64/dims=1L	1698768 ns	1697861 ns	1.00
array/accumulate/Int64/dims=2	151034 ns	151829 ns	0.99
array/accumulate/Int64/dims=2L	987438 ns	986767 ns	1.00
array/broadcast	10473 ns	15832 ns	0.66
array/broadcast launch	7604.25 ns	6879.4 ns	1.11
array/construct	918.0857142857143 ns	866.433962264151 ns	1.06
array/copy	16593 ns	16688 ns	0.99
array/copyto!/cpu_to_gpu	208249 ns	208406 ns	1.00
array/copyto!/gpu_to_cpu	240381 ns	241247 ns	1.00
array/copyto!/gpu_to_gpu	10216.333333333334 ns	8802.333333333334 ns	1.16
array/iteration/findall/bool	133260 ns	130199 ns	1.02
array/iteration/findall/int	144143 ns	138662 ns	1.04
array/iteration/findfirst/bool	70515 ns	67562 ns	1.04
array/iteration/findfirst/int	71870 ns	69121 ns	1.04
array/iteration/findmin/1d	64048 ns	59127 ns	1.08
array/iteration/findmin/2d	91988 ns	96956 ns	0.95
array/iteration/logical	185119 ns	180975 ns	1.02
array/iteration/scalar	61503 ns	58859 ns	1.04
array/permutedims/2d	45248 ns	46242 ns	0.98
array/permutedims/3d	42703 ns	46855 ns	0.91
array/permutedims/4d	45820 ns	48368 ns	0.95
array/random/rand/Float32	11673 ns	11809 ns	0.99
array/random/rand/Int64	20341 ns	18752 ns	1.08
array/random/rand!/Float32	7877.25 ns	7829.5 ns	1.01
array/random/rand!/Int64	17734 ns	15615 ns	1.14
array/random/randn/Float32	32328 ns	32695 ns	0.99
array/random/randn!/Float32	23654 ns	23566 ns	1.00
array/reductions/mapreduce/Float32/1d	33699 ns	31821 ns	1.06
array/reductions/mapreduce/Float32/dims=1	37954 ns	37330 ns	1.02
array/reductions/mapreduce/Float32/dims=1L	51448 ns	50978 ns	1.01
array/reductions/mapreduce/Float32/dims=2	55844 ns	54829 ns	1.02
array/reductions/mapreduce/Float32/dims=2L	68161 ns	67220 ns	1.01
array/reductions/mapreduce/Int64/1d	41254 ns	39331 ns	1.05
array/reductions/mapreduce/Int64/dims=1	41187 ns	40315 ns	1.02
array/reductions/mapreduce/Int64/dims=1L	88987 ns	88477 ns	1.01
array/reductions/mapreduce/Int64/dims=2	58196 ns	57574 ns	1.01
array/reductions/mapreduce/Int64/dims=2L	83802 ns	83461 ns	1.00
array/reductions/reduce/Float32/1d	33824 ns	32196 ns	1.05
array/reductions/reduce/Float32/dims=1	38049 ns	37272 ns	1.02
array/reductions/reduce/Float32/dims=1L	51046 ns	50697 ns	1.01
array/reductions/reduce/Float32/dims=2	55738 ns	55042 ns	1.01
array/reductions/reduce/Float32/dims=2L	68157 ns	67551 ns	1.01
array/reductions/reduce/Int64/1d	41342 ns	39026 ns	1.06
array/reductions/reduce/Int64/dims=1	40763 ns	40114 ns	1.02
array/reductions/reduce/Int64/dims=1L	89125 ns	88494 ns	1.01
array/reductions/reduce/Int64/dims=2	58181 ns	57350 ns	1.01
array/reductions/reduce/Int64/dims=2L	84287 ns	83169 ns	1.01
array/reverse/1d	17332 ns	17038 ns	1.02
array/reverse/1dL	70043 ns	69729 ns	1.00
array/reverse/1dL_inplace	67843 ns	67315 ns	1.01
array/reverse/1d_inplace	10699 ns	8467.666666666666 ns	1.26
array/reverse/2d	20732 ns	20236 ns	1.02
array/reverse/2dL	73963 ns	73577 ns	1.01
array/reverse/2dL_inplace	67586 ns	67185 ns	1.01
array/reverse/2d_inplace	12393 ns	9814 ns	1.26
array/sorting/1d	2658475 ns	2646705 ns	1.00
array/sorting/2d	1019293 ns	1017870 ns	1.00
array/sorting/by	3174280 ns	3174206 ns	1.00
cuda/synchronization/context/auto	1022.2 ns	6849.6 ns	0.15
cuda/synchronization/context/blocking	802.7142857142857 ns	815.9772727272727 ns	0.98
cuda/synchronization/context/nonblocking	5821.8 ns	6897.4 ns	0.84
cuda/synchronization/stream/auto	887.2727272727273 ns	700.5918367346939 ns	1.27
cuda/synchronization/stream/blocking	687.6190476190476 ns	892.3673469387755 ns	0.77
cuda/synchronization/stream/nonblocking	5699.666666666667 ns	7433 ns	0.77
integration/byval/reference	148398 ns	148108 ns	1.00
integration/byval/slices=1	149559 ns	149157 ns	1.00
integration/byval/slices=2	292214 ns	291918 ns	1.00
integration/byval/slices=3	435190 ns	434943 ns	1.00
integration/cudadevrt	105455 ns	105211 ns	1.00
integration/volumerhs	9150971 ns	9139091 ns	1.00
kernel/indexing	13329 ns	13078 ns	1.02
kernel/indexing_checked	14124 ns	13558 ns	1.04
kernel/launch	2486.3333333333335 ns	2086.222222222222 ns	1.19
kernel/occupancy	984.7857142857143 ns	684.6734693877551 ns	1.44
kernel/rand	14896 ns	16385 ns	0.91
latency/import	4291957078 ns	4265954564 ns	1.01
latency/precompile	5021407533 ns	5038240985 ns	1.00
latency/ttfp	4737418266 ns	4715207932 ns	1.00

This comment was automatically generated by workflow using github-action-benchmark.