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 OpenCL.jl to that model in one step. It replaces #474, which added a KernelInterface back end next to the KernelAbstractions 0.9 one, and #475, which ported that to an earlier version of KernelAbstractions 0.10.

OpenCLBackend now implements KernelInterface, and OpenCL.jl depends on KernelInterface instead of KernelAbstractions. The KernelInterface layer is usable on its own, without KernelAbstractions' macros:

using OpenCL, pocl_jll
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 = CLArray(rand(Float32, 1000)); b = CLArray(rand(Float32, 1000)); c = similar(a)
KI.@launch OpenCLBackend() 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, and the Adapt rules behind @Const and for moving arrays back to the CPU. OpenCL.jl'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.

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 anymore; the test suite checks that in the generated code. With a KernelAbstractions kernel that copies a Float32 array, on PoCL 7.2 with 8 threads (POCL_MAX_PTHREAD_COUNT=8) on a Ryzen 9 9950X, best of six runs each (the machine was busy, so the numbers are noisy):

	main (KA 0.9)	this PR
256×256×256, @index(Global, Cartesian)	6.3 ms	4.3 ms
256×256×256, @index(Global, Linear)	5.8 ms	4.5 ms
255×257×129, @index(Global, Cartesian)	3.3 ms	1.8 ms
255×257×129, @index(Global, Linear)	3.0 ms	1.9 ms
launch of an empty kernel	4.4 µs, 3296 bytes	3.3 µs, 1376 bytes
launch of a kernel with 40 arguments	16.6 µs, 60480 bytes	5.1 µs, 1376 bytes

Launching kernels with many arguments also got a lot cheaper. KernelInterface 0.4 passes the arguments to the back end as a tuple (JuliaGPU/KernelAbstractions.jl#811), and the first commit makes OpenCL.jl's own launch path (the kernel object's call, cl.call and cl.set_args!) pass them along as a tuple too, instead of splatting them at every layer. That also benefits @opencl: launching a kernel with 40 arguments went from 10.3 µs and 21.5 kB to 5.2 µs and 976 bytes. KI.launch rejects global_size and local_size, which would override the launch geometry KernelInterface has validated, but passes OpenCL.jl's other launch options on:

kernel = KI.@launch OpenCLBackend() launch=false vadd(c, a, b)
kernel(c, a, b; ndrange=length(c), wait_on=events)   # fine
kernel(c, a, b; ndrange=length(c), local_size=256)   # ArgumentError

Kernels compiled through KernelInterface keep the callable they were compiled from, and convert it again at every launch, as @opencl launch=false already did: a closure that captures a CLArray converts to one that only holds a pointer, and converting it at launch is what registers the array's memory with the kernel and makes the launch's queue its owner.

An OpenCLBackend keeps its platform. Previously, allocating or launching with a backend for a platform other than the active one printed a warning, on every launch. Now, work for such a backend activates the default device of its platform first, as KI.device! would, so that the arrays it creates can be used afterwards; get_backend of an array returns a backend for the array's platform, and launching a kernel on a device it wasn't compiled for throws an error instead of failing in the driver.

The next commits implement parts of KernelInterface that are optional: sub-groups, ordering work across tasks with OpenCL events for KernelAbstractions.@spawn, and KI.versioninfo. Sub-groups are only reported for devices where kernels get a fixed sub-group width, i.e. with cl_intel_required_subgroup_size (which PoCL and Intel's CPU runtime have), because KernelInterface promises that width to every kernel; a sub_group_size compiler option asking for another width is rejected. Waiting for an event of another device happens on the host, because OpenCL.jl creates a context per device and a queue can only wait for events of its own context.

What remains:

• KI.synchronize (and waiting for another device's event) still blocks the thread, where KernelInterface asks for a cooperative wait. Wait for the device cooperatively #518 addresses that separately.
• KernelAbstractions converts the arguments twice, once to determine the argument types to compile for and once when launching. kernel_convert is supposed to be pure, so this only shows with a conversion that has side effects.
• Intel's CPU OpenCL runtime (2026.21) forms sub-groups per row of a multi-dimensional work-group, so a (33, 2) work-group has four sub-groups of 32 and 1 work-items. KernelInterface promised three, which KernelInterface: don't promise how work-items form sub-groups KernelAbstractions.jl#815 relaxes; with that, sub-groups are reported on this runtime too, and it passes KernelInterface's testsuite. Until #815 is merged, the testsuite on KernelAbstractions' main still fails there, but that runtime isn't tested in CI. The same runtime also gets KernelAbstractions' N-d launches with a partial last work-group along x wrong through its SPIR-V path (Intel's CPU runtime skips the partial last work-group along x in N-d launches #521). PoCL and NVIDIA's OpenCL pass everything.

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

[codecov]

Codecov Report

❌ Patch coverage is 98.57143% with 2 lines in your changes missing coverage. Please review.
✅ Project coverage is 86.05%. Comparing base (e3f48f3) to head (8c6f02f).
⚠️ Report is 1 commits behind head on main.

Files with missing lines	Patch %	Lines
src/OpenCLKernels.jl	99.21%	1 Missing ⚠️
src/compiler/execution.jl	75.00%	1 Missing ⚠️

Additional details and impacted files

@@            Coverage Diff             @@
##             main     #519      +/-   ##
==========================================
+ Coverage   85.53%   86.05%   +0.52%     
==========================================
  Files          19       20       +1     
  Lines        1680     1707      +27     
==========================================
+ Hits         1437     1469      +32     
+ Misses        243      238       -5     

☔ View full report in Codecov by Harness.
📢 Have feedback on the report? Share it here.

🚀 New features to boost your workflow:

• ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.