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 AMDGPU.jl to that model in one step, as JuliaGPU/CUDA.jl#3314 does for CUDA. It supersedes #1047, which added KernelInterface next to the KernelAbstractions 0.9 back end, and Christian's intrinsics branch, which ported an earlier version of it to KernelAbstractions 0.10.

ROCBackend now implements KernelInterface, and AMDGPU depends on KernelInterface instead of KernelAbstractions. That makes the KernelInterface layer usable on its own, without KernelAbstractions' macros:

using AMDGPU
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 = AMDGPU.rand(1000); b = AMDGPU.rand(1000); c = similar(a)
KI.@launch ROCBackend() 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 (which still reads through the constant address space), and the maxthreads hint for kernels with a static workgroup size. AMDGPU's own copy of the KernelAbstractions launch path goes away: partitioning the ndrange, building the kernel's context, tuning the workgroup size with the occupancy API, and calling the kernel. In practice KernelAbstractions is still always loaded, because AcceleratedKernels depends on it, so the extension is about where the code lives rather than about load time.

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 integer divisions to compute its index anymore, which matter more on AMD GPUs than elsewhere because they have no hardware integer division. The test suite checks that for the generated LLVM IR. I have no AMD GPU to measure the effect on, so unlike #3314 there are no timings here.

Some smaller changes came along:

• Each call site of KI.localmemory (and thus @localmem) needs its own memory, and KernelInterface doesn't pass an id to tell them apart. alloc_special now defines local memory with internal linkage, so that allocations don't alias because they share a name.
• AMDGPU.sync_wavefront() is a new barrier for the lanes of a wavefront, fenced at wavefront scope. It backs KI.sub_group_barrier.
• HIP.max_workgroup_dims(dev) reports the per-dimension workgroup limits, cached in the device, since KernelInterface queries them on every launch whose workgroup size it picks.
• shfl_down of a Complex{Int32} compiles now.

KI.launch rejects groupsize and gridsize, which would override the launch geometry KernelInterface has validated, but passes AMDGPU's other launch options on:

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

KI.kernel_function receives the kernel's callable unconverted. It compiles the converted form, which only holds pointers to the arrays a closure captures, so the returned kernel keeps the original callable alive. Every launch converts it again for the launch's stream, like the arguments, so that the arrays it captures are handed to that stream the same way.

KernelInterface 0.4 passes the kernel arguments to the back end as one tuple (JuliaGPU/KernelAbstractions.jl#811), so that kernels with many arguments stay cheap to launch. A HIPKernel can't be launched with a tuple yet, so KI.launch splats it for now, and launching a kernel with more than about 32 arguments still goes through dynamic dispatch, as it does on master. #1115 adds the tuple launch; once it is in, KI.launch should forward the tuple to it.

The first two commits prepare local memory and the workgroup limits, and the third is the port. The next ones implement parts of KernelInterface that are optional: sync_wavefront and sub-groups (wavefronts, including a partial last wavefront and lane ids that don't change in divergent code), ordering work across tasks with HIP events for KernelAbstractions.@spawn, and KI.versioninfo.

What changes for users: ROCBackend is a KernelInterface.Backend rather than a KernelAbstractions.GPU, which no longer exists. Code that uses KernelAbstractions.get_backend, allocate, synchronize and so on keeps working, since KernelAbstractions re-exports them from KernelInterface. @print in a kernel still does nothing on AMDGPU.

I could not run any of this on AMD hardware. Locally, without a GPU, I checked that AMDGPU and its KernelAbstractions extension precompile and load on Julia 1.10, 1.11, 1.12 and 1.13, that ROCBackend has methods for everything KernelInterface requires, and, compiling for gfx1030 and gfx90a through GPUCompiler, that KernelAbstractions kernels using local memory, @private, @Const, atomics and @print, the sub-group queries, and shfl_down for every type KI.supports_shuffle claims all compile. Whether they run correctly depends on AMDGPU's CI.

This needs a breaking release, since it requires KernelAbstractions 0.10, which isn't registered yet, as well as an AcceleratedKernels release that allows it (the registered 0.4.3 caps KernelAbstractions at 0.9). Until then, the last commit takes KernelAbstractions and KernelInterface from their development branch and AcceleratedKernels from its main branch, through [sources]. Julia 1.10 doesn't support [sources], and without workspaces Julia 1.11 only picks up KernelAbstractions from the test project's, so Buildkite runs the tests in the test project on those versions, and the Enzyme and GPU-less jobs move to Julia 1.12. That commit is dropped before merging, and the AcceleratedKernels compat bound should then require the release that supports KernelAbstractions 0.10.

[gbaraldi]

Tested on MI300A + MI250 (ROCm 7.2.4):

1. CI failure. KA.Scratchpad builds a StaticArraysCore.MArray(undef), whose constructor and indexing live in StaticArrays, and KA 0.10 no longer loads that. GPUArrays' matmul (@private) then fails with InvalidIRError (jl_f_throw_methoderror): 524 failures and 67 errors across 9 files on both GPUs, while main passes. Wrapping GPUCompiler.alloca(T, Val(prod(dims)), Val(AS.Private)) in a ROCDeviceArray (like POCL) fixes it; depending on StaticArrays would too.
2. Device switch. Calling a KI.Kernel after device! to another GPU segfaults in hipFuncGetAttribute (via KI.max_work_group_size(kernel)). Store the device in ROCKernel and re-resolve or throw.
3. [sources]: AcceleratedKernels main is now 0.5.0, outside compat, so instantiate fails. AK 0.5 also makes backend a kwarg, which breaks accumulate/cumsum/cumprod in src/kernels/accumulate.jl.
4. Nit: each launch does 2–3 uncached attribute queries (~0.4–0.5 µs each); they could be cached like max_workgroup_dims.

Perf on MI300A, main → PR: async launch 11.8 → 4.4 µs, launch+sync 28.1 → 17.8 µs, 3-D scale 0.334 → 0.222 ms. div/rem are gone from 3-D indexing (ISA 188 → 36 instructions).