From 86ff07e306262f284636b389c8449f00f2b754f9 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Wed, 17 Dec 2025 12:56:50 -0400
Subject: [PATCH 01/13] Thread-local memory without `id` argument

---
 src/device/gcn/memory_static.jl | 8 ++++----
 1 file changed, 4 insertions(+), 4 deletions(-)

diff --git a/src/device/gcn/memory_static.jl b/src/device/gcn/memory_static.jl
index b57937911..be0eb4df3 100644
--- a/src/device/gcn/memory_static.jl
+++ b/src/device/gcn/memory_static.jl
@@ -2,7 +2,7 @@
 @generated function alloc_special(
     ::Val{id}, ::Type{T}, ::Val{as}, ::Val{len}, ::Val{zeroinit} = Val{false}(),
 ) where {id,T,as,len,zeroinit}
-    @dispose ctx=Context() begin
+    Context() do ctx
         eltyp = convert(LLVMType, T)
 
         # old versions of GPUArrays invoke _shmem with an integer id; make sure those are unique
@@ -24,8 +24,8 @@
         gv = GlobalVariable(mod, gv_typ, string(id), as)
         if len > 0
             if as == AS.Local
-                linkage!(gv, LLVM.API.LLVMExternalLinkage)
-                # NOTE: Backend doesn't support initializer for local AS
+                linkage!(gv, LLVM.API.LLVMInternalLinkage)
+                initializer!(gv, UndefValue(gv_typ))
             elseif as == AS.Private
                 linkage!(gv, LLVM.API.LLVMInternalLinkage)
                 initializer!(gv, null(gv_typ))
@@ -38,7 +38,7 @@
         alignment!(gv, Base.max(32, Base.datatype_alignment(T)))
 
         # generate IR
-        @dispose builder=IRBuilder() begin
+        IRBuilder() do builder
             entry = BasicBlock(llvm_f, "entry")
             position!(builder, entry)
 

From 972aa4a8298ad6430dfe30f69cb2a2efd43306c1 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Sat, 22 Aug 2026 15:43:52 -0300
Subject: [PATCH 02/13] sync_wavefront

---
 src/AMDGPU.jl                     |  2 +-
 src/device/gcn/synchronization.jl | 11 +++++++++++
 2 files changed, 12 insertions(+), 1 deletion(-)

diff --git a/src/AMDGPU.jl b/src/AMDGPU.jl
index 636b3bf27..94a6147f4 100644
--- a/src/AMDGPU.jl
+++ b/src/AMDGPU.jl
@@ -102,7 +102,7 @@ import .Device: ROCDeviceArray, AS, HostCall, HostCallHolder, hostcall!
 import .Device: @ROCDynamicLocalArray, @ROCStaticLocalArray
 import .Device: workitemIdx, workgroupIdx, workgroupDim, gridItemDim, gridGroupDim
 import .Device: threadIdx, blockIdx, blockDim
-import .Device: sync_workgroup, sync_workgroup_count, sync_workgroup_and, sync_workgroup_or
+import .Device: sync_workgroup, sync_wavefront, sync_workgroup_count, sync_workgroup_and, sync_workgroup_or
 import .Device: @rocprint, @rocprintln, @rocprintf
 
 export ROCDeviceArray, @ROCDynamicLocalArray, @ROCStaticLocalArray
diff --git a/src/device/gcn/synchronization.jl b/src/device/gcn/synchronization.jl
index 6d8b682dc..5e40d9cac 100644
--- a/src/device/gcn/synchronization.jl
+++ b/src/device/gcn/synchronization.jl
@@ -29,6 +29,17 @@ Waits until all wavefronts in a workgroup have reached this call and that their
     UnsafeAtomics.fence(UnsafeAtomics.seq_cst, AMDGPU.syncscope_workgroup)
 end
 
+"""
+    sync_wavefront()
+
+Waits until all wavefronts in a workgroup have reached this call and that their memory accesses are visible to other threads in the workgroup.
+"""
+@inline function sync_wavefront()
+    # This is a no-op https://github.com/llvm/llvm-project/blob/88b77d5eaa66747538a12c9876eeffdce31ddb71/openmp/device/src/Synchronization.cpp#L136-L140
+    ccall("llvm.amdgcn.wave.barrier", llvmcall, Cvoid, ())
+end
+
+
 """
     sync_workgroup_count(predicate::Cint)::Cint
 

From 122af31f7c99e073a5ef55cc921fe0574e6ecd9a Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Sat, 22 Aug 2026 15:57:39 -0300
Subject: [PATCH 03/13] KernelInterface

---
 Project.toml                  |   2 +
 src/AMDGPU.jl                 |   7 +-
 src/ROCKernels.jl             | 205 ++++++++++++++--------------------
 src/ROCKernelsOld.jl          | 201 +++++++++++++++++++++++++++++++++
 test/Project.toml             |   1 +
 test/kernelinterface_tests.jl |  15 +++
 6 files changed, 311 insertions(+), 120 deletions(-)
 create mode 100644 src/ROCKernelsOld.jl
 create mode 100644 test/kernelinterface_tests.jl

diff --git a/Project.toml b/Project.toml
index 8bc9180ef..0d22a3edd 100644
--- a/Project.toml
+++ b/Project.toml
@@ -19,6 +19,7 @@ GPUArrays = "0c68f7d7-f131-5f86-a1c3-88cf8149b2d7"
 GPUCompiler = "61eb1bfa-7361-4325-ad38-22787b887f55"
 GPUToolbox = "096a3bc2-3ced-46d0-87f4-dd12716f4bfc"
 KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c"
+KernelInterface = "4ee993da-d684-4d17-a7dd-4e58e78d92bf"
 LLVM = "929cbde3-209d-540e-8aea-75f648917ca0"
 LLVMDowngrader_jll = "f52de702-fb25-5922-94ba-81dd59b07444"
 Libdl = "8f399da3-3557-5675-b5ff-fb832c97cbdb"
@@ -63,6 +64,7 @@ GPUArrays = "11.5.14"
 GPUCompiler = "2.9"
 GPUToolbox = "3"
 KernelAbstractions = "0.9.2"
+KernelInterface = "0.1.0"
 LLVM = "9"
 LLVMDowngrader_jll = "0.11"
 Libdl = "1"
diff --git a/src/AMDGPU.jl b/src/AMDGPU.jl
index 94a6147f4..79fda4125 100644
--- a/src/AMDGPU.jl
+++ b/src/AMDGPU.jl
@@ -147,10 +147,15 @@ function Atomix.modify!(ref::ROCIndexableRef, op::OP, x, ord) where OP <: Union{
     GC.@preserve root UnsafeAtomics.modify!(ptr, op, x, ord, syncscope_agent)
 end
 
-include("ROCKernels.jl")
+# KernelAbstractions
+include("ROCKernelsOld.jl")
 import .ROCKernels: ROCBackend
 export ROCBackend
 
+# KernelInterface
+include("ROCKernels.jl")
+import .ROCInterface
+
 include("precompile.jl")
 
 function __init__()
diff --git a/src/ROCKernels.jl b/src/ROCKernels.jl
index 2315222fe..f9d1ad037 100644
--- a/src/ROCKernels.jl
+++ b/src/ROCKernels.jl
@@ -1,13 +1,14 @@
-module ROCKernels
+module ROCInterface
 
 export ROCBackend
 
 import AMDGPU
+import AMDGPU: rocconvert, hipfunction
 import AMDGPU.Device: @device_override
-using AMDGPU: GPUArrays, rocSPARSE
+using AMDGPU: GPUArrays, rocSPARSE, HIP, Device
 
 import Adapt
-import KernelAbstractions as KA
+import KernelInterface as KI
 import LLVM
 
 using StaticArraysCore: MArray
@@ -19,183 +20,149 @@ KernelAbstractions backend that executes kernels on an AMD GPU via AMDGPU.jl.
 Pass `ROCBackend()` to a KernelAbstractions kernel to run it on the GPU, or
 obtain it from an array with `KernelAbstractions.get_backend(::ROCArray)`.
 """
-struct ROCBackend <: KA.GPU end
+struct ROCBackend <: KI.GPU end
 
-KA.functional(::ROCBackend) = AMDGPU.functional()
-KA.ndevices(::ROCBackend) = AMDGPU.HIP.ndevices()
-KA.device(::ROCBackend) = AMDGPU.device_id()
-function KA.device!(kab::ROCBackend, id::Int)
-    (0 < id <= KA.ndevices(kab)) || throw(ArgumentError("Device id $id out of bounds."))
+KI.versioninfo(io::IO, ::ROCBackend) = AMDGPU.versioninfo(io)
+
+KI.functional(::ROCBackend) = AMDGPU.functional()
+KI.ndevices(::ROCBackend) = AMDGPU.HIP.ndevices()
+KI.device(::ROCBackend) = AMDGPU.device_id()
+function KI.device!(kab::ROCBackend, id::Int)
+    (0 < id <= KI.ndevices(kab)) || throw(ArgumentError("Device id $id out of bounds."))
     AMDGPU.device_id!(id)
     return
 end
 
 Adapt.adapt_storage(::ROCBackend, a::AbstractArray) = Adapt.adapt(AMDGPU.ROCArray, a)
 Adapt.adapt_storage(::ROCBackend, a::Union{AMDGPU.ROCArray, GPUArrays.AbstractGPUSparseArray}) = a
-Adapt.adapt_storage(::KA.CPU, a::Union{AMDGPU.ROCArray, GPUArrays.AbstractGPUSparseArray}) =
-    Adapt.adapt(Array, a)
-function Adapt.adapt_storage(::KA.ConstAdaptor, a::AMDGPU.ROCDeviceArray{T}) where T
-    ptr = LLVM.Interop.addrspacecast(Core.LLVMPtr{T,AMDGPU.Device.AS.Constant}, a.ptr)
-    AMDGPU.ROCDeviceArray(a.dims, ptr)
-end
 
-KA.get_backend(::AMDGPU.ROCArray) = ROCBackend()
-KA.get_backend(::AMDGPU.rocSPARSE.ROCSparseVector) = ROCBackend()
-KA.get_backend(::AMDGPU.rocSPARSE.ROCSparseMatrixCSC) = ROCBackend()
-KA.get_backend(::AMDGPU.rocSPARSE.ROCSparseMatrixCSR) = ROCBackend()
+KI.get_backend(::AMDGPU.ROCArray) = ROCBackend()
+KI.get_backend(::AMDGPU.rocSPARSE.ROCSparseVector) = ROCBackend()
+KI.get_backend(::AMDGPU.rocSPARSE.ROCSparseMatrixCSC) = ROCBackend()
+KI.get_backend(::AMDGPU.rocSPARSE.ROCSparseMatrixCSR) = ROCBackend()
 
-KA.argconvert(::KA.Kernel{ROCBackend}, arg) = AMDGPU.rocconvert(arg)
-KA.synchronize(::ROCBackend) = AMDGPU.synchronize()
+KI.synchronize(::ROCBackend) = AMDGPU.synchronize()
 
-KA.unsafe_free!(x::AMDGPU.ROCArray) = AMDGPU.unsafe_free!(x)
-KA.allocate(::ROCBackend, ::Type{T}, dims::Tuple) where T = AMDGPU.ROCArray{T}(undef, dims)
-KA.zeros(::ROCBackend, ::Type{T}, dims::Tuple) where T = AMDGPU.zeros(T, dims)
-KA.ones(::ROCBackend, ::Type{T}, dims::Tuple) where T = AMDGPU.ones(T, dims)
+KI.unsafe_free!(x::AMDGPU.ROCArray) = AMDGPU.unsafe_free!(x)
+KI.allocate(::ROCBackend, ::Type{T}, dims::Tuple) where T = AMDGPU.ROCArray{T}(undef, dims)
+KI.zeros(::ROCBackend, ::Type{T}, dims::Tuple) where T = AMDGPU.zeros(T, dims)
+KI.ones(::ROCBackend, ::Type{T}, dims::Tuple) where T = AMDGPU.ones(T, dims)
 
-function KA.priority!(::ROCBackend, priority::Symbol)
+function KI.priority!(::ROCBackend, priority::Symbol)
     priority ∉ (:high, :normal, :low) && error(
         "Priority `$priority` must be one of `:high`, `:normal`, `:low`.")
     AMDGPU.priority!(priority)
 end
 
-function KA.copyto!(::ROCBackend, A, B)
+function KI.copyto!(::ROCBackend, A, B)
     GC.@preserve A B begin
         copyto!(A, 1, B, 1, length(A))
     end
     return
 end
 
-function KA.pagelock!(::ROCBackend, x::Array)
+function KI.pagelock!(::ROCBackend, x::Array)
     AMDGPU.Mem.pin(pointer(x), sizeof(x))
     return
 end
 
-function KA.launch_config(kernel::KA.Kernel{ROCBackend}, ndrange, workgroupsize)
-    if ndrange isa Integer
-        ndrange = (ndrange,)
-    end
-    if workgroupsize isa Integer
-        workgroupsize = (workgroupsize, )
-    end
-
-    # partition checked that the ndrange's agreed
-    if KA.ndrange(kernel) <: KA.StaticSize
-        ndrange = nothing
-    end
-
-    iterspace, dynamic = if KA.workgroupsize(kernel) <: KA.DynamicSize && workgroupsize === nothing
-        workgroupsize = ntuple(
-            i -> i == 1 ? min(prod(ndrange), AMDGPU.Device._max_group_size) : 1,
-            length(ndrange))
-        KA.partition(kernel, ndrange, workgroupsize)
-    else
-        KA.partition(kernel, ndrange, workgroupsize)
-    end
+KI.argconvert(::ROCBackend, arg) = rocconvert(arg)
 
-    return ndrange, workgroupsize, iterspace, dynamic
+function KI.kernel_function(::ROCBackend, f::F, tt::TT=Tuple{}; name=nothing, kwargs...) where {F,TT}
+    kern = hipfunction(f, tt; name, kwargs...)
+    KI.Kernel{ROCBackend, typeof(kern)}(ROCBackend(), kern)
 end
 
-function threads_to_workgroupsize(threads, ndrange)
-    total = 1
-    return map(ndrange) do n
-        x = min(div(threads, total), n)
-        total *= x
-        return x
-    end
-end
+function (obj::KI.Kernel{ROCBackend})(args...; numworkgroups=(), workgroupsize=(), ndrange=(), max_work_group_size=typemax(Int))
+    KI.check_launch_args(numworkgroups, workgroupsize, ndrange)
+    prod(ndrange) == 0 && return nothing
 
-function (obj::KA.Kernel{ROCBackend})(args...; ndrange=nothing, workgroupsize=nothing)
-    ndrange, new_workgroupsize, iterspace, dynamic = KA.launch_config(obj, ndrange, workgroupsize)
-    ctx = KA.mkcontext(obj, ndrange, iterspace)
-    if KA.workgroupsize(obj) <: KA.StaticSize
-        maxthreads = prod(KA.get(KA.workgroupsize(obj)))
-    else
-        maxthreads = nothing
-    end
-    kernel = AMDGPU.@roc launch=false maxthreads=maxthreads obj.f(ctx, args...)
-
-    # If dynamic, figure out the optimal groupsize automatically.
-    is_dynamic =
-        KA.workgroupsize(obj) <: KA.DynamicSize &&
-        isnothing(workgroupsize)
-    if is_dynamic
-        (; groupsize) = AMDGPU.launch_configuration(kernel)
-        new_workgroupsize = threads_to_workgroupsize(groupsize, ndrange)
-        iterspace, dynamic = KA.partition(obj, ndrange, new_workgroupsize)
-        ctx = KA.mkcontext(obj, ndrange, iterspace)
-    end
+    numworkgroups, workgroupsize = KI.auto_launch_sizes(obj, numworkgroups, workgroupsize, ndrange, max_work_group_size)
+    obj.kern(args...; groupsize = workgroupsize, gridsize = numworkgroups)
+    return nothing
+end
 
-    nblocks = length(KA.blocks(iterspace))
-    nthreads = length(KA.workitems(iterspace))
-    nblocks == 0 && return
+function KI.kernel_max_work_group_size(kikern::KI.Kernel{<:ROCBackend}; max_work_items::Int=Int(typemax(Int32)))::Int
+    (; groupsize) = AMDGPU.launch_configuration(kikern.kern; max_block_size = max_work_items)
 
-    kernel(ctx, args...; groupsize=nthreads, gridsize=nblocks)
-    return
+    return Int(min(max_work_items, groupsize))
 end
-
-function KA.mkcontext(kernel::KA.Kernel{ROCBackend}, _ndrange, iterspace)
-    metadata = KA.CompilerMetadata{KA.ndrange(kernel), KA.DynamicCheck}(_ndrange, iterspace)
+function KI.max_work_group_size(::ROCBackend)::Int
+    Int(HIP.attribute(AMDGPU.HIP.device(), AMDGPU.HIP.hipDeviceAttributeMaxThreadsPerBlock))
 end
-function KA.mkcontext(kernel::KA.Kernel{ROCBackend}, I, _ndrange, iterspace, ::Dynamic) where Dynamic
-    metadata = KA.CompilerMetadata{KA.ndrange(kernel), Dynamic}(I, _ndrange, iterspace)
+function KI.sub_group_size(::ROCBackend)::Int
+    HIP.wavefrontsize(HIP.device())
+end
+function KI.multiprocessor_count(::ROCBackend)::Int
+    Int(HIP.attribute(AMDGPU.HIP.device(), AMDGPU.HIP.hipDeviceAttributeMultiprocessorCount))
 end
 
-# Indexing.
+KI.shfl_down_types(::ROCBackend) = DataType[Bool,
+                                             UInt8, UInt16, UInt32, UInt64, UInt128,
+                                             Int8, Int16, Int32, Int64, Int128,
+                                             Float16, Float32, Float64,
+                                             ComplexF16, ComplexF32, ComplexF64]
 
-@device_override @inline function KA.__index_Local_Linear(ctx)
-    return AMDGPU.Device.threadIdx().x
+# Indexing.
+## COV_EXCL_START
+@device_override @inline function KI.get_local_id()
+    return (; x = Int(AMDGPU.Device.workitemIdx().x), y = Int(AMDGPU.Device.workitemIdx().y), z = Int(AMDGPU.Device.workitemIdx().z))
 end
 
-@device_override @inline function KA.__index_Group_Linear(ctx)
-    return AMDGPU.Device.blockIdx().x
+@device_override @inline function KI.get_group_id()
+    return (; x = Int(AMDGPU.Device.workgroupIdx().x), y = Int(AMDGPU.Device.workgroupIdx().y), z = Int(AMDGPU.Device.workgroupIdx().z))
 end
 
-@device_override @inline function KA.__index_Global_Linear(ctx)
-    I =  @inbounds KA.expand(KA.__iterspace(ctx), AMDGPU.Device.blockIdx().x, AMDGPU.Device.threadIdx().x)
-    # TODO: This is unfortunate, can we get the linear index cheaper
-    @inbounds LinearIndices(KA.__ndrange(ctx))[I]
+@device_override @inline function KI.get_global_id()
+    return (; x = Int((AMDGPU.Device.workgroupIdx().x-1)*AMDGPU.Device.blockDim().x + AMDGPU.Device.workitemIdx().x), y = Int((AMDGPU.Device.workgroupIdx().y-1)*AMDGPU.Device.blockDim().y + AMDGPU.Device.workitemIdx().y), z = Int((AMDGPU.Device.workgroupIdx().z-1)*AMDGPU.Device.blockDim().z + AMDGPU.Device.workitemIdx().z))
 end
 
-@device_override @inline function KA.__index_Local_Cartesian(ctx)
-    @inbounds KA.workitems(KA.__iterspace(ctx))[AMDGPU.Device.threadIdx().x]
+@device_override @inline function KI.get_local_size()
+    return (; x = Int(AMDGPU.Device.workgroupDim().x), y = Int(AMDGPU.Device.workgroupDim().y), z = Int(AMDGPU.Device.workgroupDim().z))
 end
 
-@device_override @inline function KA.__index_Group_Cartesian(ctx)
-    @inbounds KA.blocks(KA.__iterspace(ctx))[AMDGPU.Device.blockIdx().x]
+@device_override @inline function KI.get_num_groups()
+    return (; x = Int(AMDGPU.Device.gridGroupDim().x), y = Int(AMDGPU.Device.gridGroupDim().y), z = Int(AMDGPU.Device.gridGroupDim().z))
 end
 
-@device_override @inline function KA.__index_Global_Cartesian(ctx)
-    return @inbounds KA.expand(KA.__iterspace(ctx), AMDGPU.Device.blockIdx().x, AMDGPU.Device.threadIdx().x)
+@device_override @inline function KI.get_global_size()
+    return (; x = Int(AMDGPU.Device.gridItemDim().x), y = Int(AMDGPU.Device.gridItemDim().y), z = Int(AMDGPU.Device.gridItemDim().z))
 end
 
-@device_override @inline function KA.__validindex(ctx)
-    if KA.__dynamic_checkbounds(ctx)
-        I = @inbounds KA.expand(KA.__iterspace(ctx), AMDGPU.Device.blockIdx().x, AMDGPU.Device.threadIdx().x)
-        return I in KA.__ndrange(ctx)
-    else
-        return true
-    end
-end
+@device_override KI.get_sub_group_size() = UInt32(Device.wavefrontsize())
+
+@device_override KI.get_max_sub_group_size() = UInt32(Device.wavefrontsize())
+
+@device_override KI.get_num_sub_groups() = UInt32(prod(Device.blockDim()) ÷ Device.wavefrontsize())
+
+@device_override KI.get_sub_group_id() = UInt32(((Device.threadIdx().x - 1) + Device.blockDim().x * (Device.threadIdx().y - 1) + Device.blockDim().x * Device.blockDim().y * (Device.threadIdx().z - 1)) ÷ Device.wavefrontsize()) + 0x1
+
+@device_override KI.get_sub_group_local_id() = UInt32(Device.activelane() + 0x1)
 
 # Shared memory.
 
-@device_override @inline function KA.SharedMemory(::Type{T}, ::Val{Dims}, ::Val{Id}) where {T, Dims, Id}
-    ptr = AMDGPU.Device.alloc_special(Val(Id), T, Val(AMDGPU.AS.Local), Val(prod(Dims)))
+@device_override @inline function KI.localmemory(::Type{T}, ::Val{Dims}) where {T, Dims}
+    ptr = AMDGPU.Device.alloc_special(Val(:shmem), T, Val(AMDGPU.AS.Local), Val(prod(Dims)))
     AMDGPU.ROCDeviceArray(Dims, ptr)
 end
 
-@device_override @inline function KA.Scratchpad(ctx, ::Type{T}, ::Val{Dims}) where {T, Dims}
-    MArray{KA.__size(Dims), T}(undef)
-end
-
 # Other.
 
-@device_override @inline function KA.__synchronize()
+@device_override @inline function KI.barrier()
     AMDGPU.Device.sync_workgroup()
 end
 
-@device_override @inline function KA.__print(args...)
+@device_override @inline function KI.sub_group_barrier()
+    AMDGPU.Device.sync_wavefront()
+end
+
+@device_override function KI.shfl_down(val::T, offset::Integer) where T
+    @inline AMDGPU.Device.shfl_down(val, Cint(offset))
+end
+
+@device_override @inline function KI._print(args...)
     # TODO
 end
+## COV_EXCL_STOP
 
 end
diff --git a/src/ROCKernelsOld.jl b/src/ROCKernelsOld.jl
new file mode 100644
index 000000000..2315222fe
--- /dev/null
+++ b/src/ROCKernelsOld.jl
@@ -0,0 +1,201 @@
+module ROCKernels
+
+export ROCBackend
+
+import AMDGPU
+import AMDGPU.Device: @device_override
+using AMDGPU: GPUArrays, rocSPARSE
+
+import Adapt
+import KernelAbstractions as KA
+import LLVM
+
+using StaticArraysCore: MArray
+
+"""
+    ROCBackend <: KernelAbstractions.GPU
+
+KernelAbstractions backend that executes kernels on an AMD GPU via AMDGPU.jl.
+Pass `ROCBackend()` to a KernelAbstractions kernel to run it on the GPU, or
+obtain it from an array with `KernelAbstractions.get_backend(::ROCArray)`.
+"""
+struct ROCBackend <: KA.GPU end
+
+KA.functional(::ROCBackend) = AMDGPU.functional()
+KA.ndevices(::ROCBackend) = AMDGPU.HIP.ndevices()
+KA.device(::ROCBackend) = AMDGPU.device_id()
+function KA.device!(kab::ROCBackend, id::Int)
+    (0 < id <= KA.ndevices(kab)) || throw(ArgumentError("Device id $id out of bounds."))
+    AMDGPU.device_id!(id)
+    return
+end
+
+Adapt.adapt_storage(::ROCBackend, a::AbstractArray) = Adapt.adapt(AMDGPU.ROCArray, a)
+Adapt.adapt_storage(::ROCBackend, a::Union{AMDGPU.ROCArray, GPUArrays.AbstractGPUSparseArray}) = a
+Adapt.adapt_storage(::KA.CPU, a::Union{AMDGPU.ROCArray, GPUArrays.AbstractGPUSparseArray}) =
+    Adapt.adapt(Array, a)
+function Adapt.adapt_storage(::KA.ConstAdaptor, a::AMDGPU.ROCDeviceArray{T}) where T
+    ptr = LLVM.Interop.addrspacecast(Core.LLVMPtr{T,AMDGPU.Device.AS.Constant}, a.ptr)
+    AMDGPU.ROCDeviceArray(a.dims, ptr)
+end
+
+KA.get_backend(::AMDGPU.ROCArray) = ROCBackend()
+KA.get_backend(::AMDGPU.rocSPARSE.ROCSparseVector) = ROCBackend()
+KA.get_backend(::AMDGPU.rocSPARSE.ROCSparseMatrixCSC) = ROCBackend()
+KA.get_backend(::AMDGPU.rocSPARSE.ROCSparseMatrixCSR) = ROCBackend()
+
+KA.argconvert(::KA.Kernel{ROCBackend}, arg) = AMDGPU.rocconvert(arg)
+KA.synchronize(::ROCBackend) = AMDGPU.synchronize()
+
+KA.unsafe_free!(x::AMDGPU.ROCArray) = AMDGPU.unsafe_free!(x)
+KA.allocate(::ROCBackend, ::Type{T}, dims::Tuple) where T = AMDGPU.ROCArray{T}(undef, dims)
+KA.zeros(::ROCBackend, ::Type{T}, dims::Tuple) where T = AMDGPU.zeros(T, dims)
+KA.ones(::ROCBackend, ::Type{T}, dims::Tuple) where T = AMDGPU.ones(T, dims)
+
+function KA.priority!(::ROCBackend, priority::Symbol)
+    priority ∉ (:high, :normal, :low) && error(
+        "Priority `$priority` must be one of `:high`, `:normal`, `:low`.")
+    AMDGPU.priority!(priority)
+end
+
+function KA.copyto!(::ROCBackend, A, B)
+    GC.@preserve A B begin
+        copyto!(A, 1, B, 1, length(A))
+    end
+    return
+end
+
+function KA.pagelock!(::ROCBackend, x::Array)
+    AMDGPU.Mem.pin(pointer(x), sizeof(x))
+    return
+end
+
+function KA.launch_config(kernel::KA.Kernel{ROCBackend}, ndrange, workgroupsize)
+    if ndrange isa Integer
+        ndrange = (ndrange,)
+    end
+    if workgroupsize isa Integer
+        workgroupsize = (workgroupsize, )
+    end
+
+    # partition checked that the ndrange's agreed
+    if KA.ndrange(kernel) <: KA.StaticSize
+        ndrange = nothing
+    end
+
+    iterspace, dynamic = if KA.workgroupsize(kernel) <: KA.DynamicSize && workgroupsize === nothing
+        workgroupsize = ntuple(
+            i -> i == 1 ? min(prod(ndrange), AMDGPU.Device._max_group_size) : 1,
+            length(ndrange))
+        KA.partition(kernel, ndrange, workgroupsize)
+    else
+        KA.partition(kernel, ndrange, workgroupsize)
+    end
+
+    return ndrange, workgroupsize, iterspace, dynamic
+end
+
+function threads_to_workgroupsize(threads, ndrange)
+    total = 1
+    return map(ndrange) do n
+        x = min(div(threads, total), n)
+        total *= x
+        return x
+    end
+end
+
+function (obj::KA.Kernel{ROCBackend})(args...; ndrange=nothing, workgroupsize=nothing)
+    ndrange, new_workgroupsize, iterspace, dynamic = KA.launch_config(obj, ndrange, workgroupsize)
+    ctx = KA.mkcontext(obj, ndrange, iterspace)
+    if KA.workgroupsize(obj) <: KA.StaticSize
+        maxthreads = prod(KA.get(KA.workgroupsize(obj)))
+    else
+        maxthreads = nothing
+    end
+    kernel = AMDGPU.@roc launch=false maxthreads=maxthreads obj.f(ctx, args...)
+
+    # If dynamic, figure out the optimal groupsize automatically.
+    is_dynamic =
+        KA.workgroupsize(obj) <: KA.DynamicSize &&
+        isnothing(workgroupsize)
+    if is_dynamic
+        (; groupsize) = AMDGPU.launch_configuration(kernel)
+        new_workgroupsize = threads_to_workgroupsize(groupsize, ndrange)
+        iterspace, dynamic = KA.partition(obj, ndrange, new_workgroupsize)
+        ctx = KA.mkcontext(obj, ndrange, iterspace)
+    end
+
+    nblocks = length(KA.blocks(iterspace))
+    nthreads = length(KA.workitems(iterspace))
+    nblocks == 0 && return
+
+    kernel(ctx, args...; groupsize=nthreads, gridsize=nblocks)
+    return
+end
+
+function KA.mkcontext(kernel::KA.Kernel{ROCBackend}, _ndrange, iterspace)
+    metadata = KA.CompilerMetadata{KA.ndrange(kernel), KA.DynamicCheck}(_ndrange, iterspace)
+end
+function KA.mkcontext(kernel::KA.Kernel{ROCBackend}, I, _ndrange, iterspace, ::Dynamic) where Dynamic
+    metadata = KA.CompilerMetadata{KA.ndrange(kernel), Dynamic}(I, _ndrange, iterspace)
+end
+
+# Indexing.
+
+@device_override @inline function KA.__index_Local_Linear(ctx)
+    return AMDGPU.Device.threadIdx().x
+end
+
+@device_override @inline function KA.__index_Group_Linear(ctx)
+    return AMDGPU.Device.blockIdx().x
+end
+
+@device_override @inline function KA.__index_Global_Linear(ctx)
+    I =  @inbounds KA.expand(KA.__iterspace(ctx), AMDGPU.Device.blockIdx().x, AMDGPU.Device.threadIdx().x)
+    # TODO: This is unfortunate, can we get the linear index cheaper
+    @inbounds LinearIndices(KA.__ndrange(ctx))[I]
+end
+
+@device_override @inline function KA.__index_Local_Cartesian(ctx)
+    @inbounds KA.workitems(KA.__iterspace(ctx))[AMDGPU.Device.threadIdx().x]
+end
+
+@device_override @inline function KA.__index_Group_Cartesian(ctx)
+    @inbounds KA.blocks(KA.__iterspace(ctx))[AMDGPU.Device.blockIdx().x]
+end
+
+@device_override @inline function KA.__index_Global_Cartesian(ctx)
+    return @inbounds KA.expand(KA.__iterspace(ctx), AMDGPU.Device.blockIdx().x, AMDGPU.Device.threadIdx().x)
+end
+
+@device_override @inline function KA.__validindex(ctx)
+    if KA.__dynamic_checkbounds(ctx)
+        I = @inbounds KA.expand(KA.__iterspace(ctx), AMDGPU.Device.blockIdx().x, AMDGPU.Device.threadIdx().x)
+        return I in KA.__ndrange(ctx)
+    else
+        return true
+    end
+end
+
+# Shared memory.
+
+@device_override @inline function KA.SharedMemory(::Type{T}, ::Val{Dims}, ::Val{Id}) where {T, Dims, Id}
+    ptr = AMDGPU.Device.alloc_special(Val(Id), T, Val(AMDGPU.AS.Local), Val(prod(Dims)))
+    AMDGPU.ROCDeviceArray(Dims, ptr)
+end
+
+@device_override @inline function KA.Scratchpad(ctx, ::Type{T}, ::Val{Dims}) where {T, Dims}
+    MArray{KA.__size(Dims), T}(undef)
+end
+
+# Other.
+
+@device_override @inline function KA.__synchronize()
+    AMDGPU.Device.sync_workgroup()
+end
+
+@device_override @inline function KA.__print(args...)
+    # TODO
+end
+
+end
diff --git a/test/Project.toml b/test/Project.toml
index 8a24df04d..af3919c3a 100644
--- a/test/Project.toml
+++ b/test/Project.toml
@@ -11,6 +11,7 @@ GPUCompiler = "61eb1bfa-7361-4325-ad38-22787b887f55"
 InteractiveUtils = "b77e0a4c-d291-57a0-90e8-8db25a27a240"
 JLD2 = "033835bb-8acc-5ee8-8aae-3f567f8a3819"
 KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c"
+KernelInterface = "4ee993da-d684-4d17-a7dd-4e58e78d92bf"
 LLVM = "929cbde3-209d-540e-8aea-75f648917ca0"
 LinearAlgebra = "37e2e46d-f89d-539d-b4ee-838fcccc9c8e"
 ParallelTestRunner = "d3525ed8-44d0-4b2c-a655-542cee43accc"
diff --git a/test/kernelinterface_tests.jl b/test/kernelinterface_tests.jl
new file mode 100644
index 000000000..2eba9d15f
--- /dev/null
+++ b/test/kernelinterface_tests.jl
@@ -0,0 +1,15 @@
+using Test
+using AMDGPU
+using AMDGPU: ROCInterface
+
+import KernelInterface
+include(joinpath(pkgdir(KernelInterface), "test", "testsuite.jl"))
+
+AMDGPU.allowscalar(false)
+
+@testset "kernelinterface" begin
+
+Testsuite.testsuite(
+    ROCInterface.ROCBackend, "ROCM", AMDGPU, ROCArray, AMDGPU.ROCDeviceArray)
+
+end

From b3d70de49196de4391ba5decebb534f33ae646e3 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Fri, 11 Sep 2026 16:50:15 -0300
Subject: [PATCH 04/13] Support KI 0.2

---
 Project.toml      |  2 +-
 src/ROCKernels.jl | 24 ++++++++++++------------
 2 files changed, 13 insertions(+), 13 deletions(-)

diff --git a/Project.toml b/Project.toml
index 0d22a3edd..da7b821ed 100644
--- a/Project.toml
+++ b/Project.toml
@@ -64,7 +64,7 @@ GPUArrays = "11.5.14"
 GPUCompiler = "2.9"
 GPUToolbox = "3"
 KernelAbstractions = "0.9.2"
-KernelInterface = "0.1.0"
+KernelInterface = "0.2"
 LLVM = "9"
 LLVMDowngrader_jll = "0.11"
 Libdl = "1"
diff --git a/src/ROCKernels.jl b/src/ROCKernels.jl
index f9d1ad037..fa4d70e7b 100644
--- a/src/ROCKernels.jl
+++ b/src/ROCKernels.jl
@@ -105,28 +105,28 @@ KI.shfl_down_types(::ROCBackend) = DataType[Bool,
 
 # Indexing.
 ## COV_EXCL_START
-@device_override @inline function KI.get_local_id()
-    return (; x = Int(AMDGPU.Device.workitemIdx().x), y = Int(AMDGPU.Device.workitemIdx().y), z = Int(AMDGPU.Device.workitemIdx().z))
+@device_override @inline function KI.get_local_id(::Type{T}) where {T}
+    return (; x = T(AMDGPU.Device.workitemIdx().x), y = T(AMDGPU.Device.workitemIdx().y), z = T(AMDGPU.Device.workitemIdx().z))
 end
 
-@device_override @inline function KI.get_group_id()
-    return (; x = Int(AMDGPU.Device.workgroupIdx().x), y = Int(AMDGPU.Device.workgroupIdx().y), z = Int(AMDGPU.Device.workgroupIdx().z))
+@device_override @inline function KI.get_group_id(::Type{T}) where {T}
+    return (; x = T(AMDGPU.Device.workgroupIdx().x), y = T(AMDGPU.Device.workgroupIdx().y), z = T(AMDGPU.Device.workgroupIdx().z))
 end
 
-@device_override @inline function KI.get_global_id()
-    return (; x = Int((AMDGPU.Device.workgroupIdx().x-1)*AMDGPU.Device.blockDim().x + AMDGPU.Device.workitemIdx().x), y = Int((AMDGPU.Device.workgroupIdx().y-1)*AMDGPU.Device.blockDim().y + AMDGPU.Device.workitemIdx().y), z = Int((AMDGPU.Device.workgroupIdx().z-1)*AMDGPU.Device.blockDim().z + AMDGPU.Device.workitemIdx().z))
+@device_override @inline function KI.get_global_id(::Type{T}) where {T}
+    return (; x = T((AMDGPU.Device.workgroupIdx().x-1)*AMDGPU.Device.blockDim().x + AMDGPU.Device.workitemIdx().x), y = T((AMDGPU.Device.workgroupIdx().y-1)*AMDGPU.Device.blockDim().y + AMDGPU.Device.workitemIdx().y), z = T((AMDGPU.Device.workgroupIdx().z-1)*AMDGPU.Device.blockDim().z + AMDGPU.Device.workitemIdx().z))
 end
 
-@device_override @inline function KI.get_local_size()
-    return (; x = Int(AMDGPU.Device.workgroupDim().x), y = Int(AMDGPU.Device.workgroupDim().y), z = Int(AMDGPU.Device.workgroupDim().z))
+@device_override @inline function KI.get_local_size(::Type{T}) where {T}
+    return (; x = T(AMDGPU.Device.workgroupDim().x), y = T(AMDGPU.Device.workgroupDim().y), z = T(AMDGPU.Device.workgroupDim().z))
 end
 
-@device_override @inline function KI.get_num_groups()
-    return (; x = Int(AMDGPU.Device.gridGroupDim().x), y = Int(AMDGPU.Device.gridGroupDim().y), z = Int(AMDGPU.Device.gridGroupDim().z))
+@device_override @inline function KI.get_num_groups(::Type{T}) where {T}
+    return (; x = T(AMDGPU.Device.gridGroupDim().x), y = T(AMDGPU.Device.gridGroupDim().y), z = T(AMDGPU.Device.gridGroupDim().z))
 end
 
-@device_override @inline function KI.get_global_size()
-    return (; x = Int(AMDGPU.Device.gridItemDim().x), y = Int(AMDGPU.Device.gridItemDim().y), z = Int(AMDGPU.Device.gridItemDim().z))
+@device_override @inline function KI.get_global_size(::Type{T}) where {T}
+    return (; x = T(AMDGPU.Device.gridItemDim().x), y = T(AMDGPU.Device.gridItemDim().y), z = T(AMDGPU.Device.gridItemDim().z))
 end
 
 @device_override KI.get_sub_group_size() = UInt32(Device.wavefrontsize())

From e33ef3b0443752597560f99ff175c4dd4f51def7 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Mon, 21 Sep 2026 16:23:12 -0300
Subject: [PATCH 05/13] KI 0.2.1

---
 Project.toml      | 2 +-
 src/ROCKernels.jl | 9 +++++++++
 2 files changed, 10 insertions(+), 1 deletion(-)

diff --git a/Project.toml b/Project.toml
index da7b821ed..d7d144f4b 100644
--- a/Project.toml
+++ b/Project.toml
@@ -64,7 +64,7 @@ GPUArrays = "11.5.14"
 GPUCompiler = "2.9"
 GPUToolbox = "3"
 KernelAbstractions = "0.9.2"
-KernelInterface = "0.2"
+KernelInterface = "0.2.1"
 LLVM = "9"
 LLVMDowngrader_jll = "0.11"
 Libdl = "1"
diff --git a/src/ROCKernels.jl b/src/ROCKernels.jl
index fa4d70e7b..137849c55 100644
--- a/src/ROCKernels.jl
+++ b/src/ROCKernels.jl
@@ -43,6 +43,15 @@ KI.get_backend(::AMDGPU.rocSPARSE.ROCSparseMatrixCSR) = ROCBackend()
 
 KI.synchronize(::ROCBackend) = AMDGPU.synchronize()
 
+function KI.record_event(::ROCBackend)
+    return HIP.HIPEvent(AMDGPU.stream())
+end
+
+function KI.wait_event(::ROCBackend, ev::HIP.HIPEvent)
+    HIP.hipStreamWaitEvent(AMDGPU.stream(), ev, 0)
+    return
+end
+
 KI.unsafe_free!(x::AMDGPU.ROCArray) = AMDGPU.unsafe_free!(x)
 KI.allocate(::ROCBackend, ::Type{T}, dims::Tuple) where T = AMDGPU.ROCArray{T}(undef, dims)
 KI.zeros(::ROCBackend, ::Type{T}, dims::Tuple) where T = AMDGPU.zeros(T, dims)

From 9aef27bb506a799db779d43579d876d4359386e8 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Mon, 28 Sep 2026 00:24:01 -0300
Subject: [PATCH 06/13] Add KernelInterface to versioninfo

---
 src/AMDGPU.jl | 1 +
 src/utils.jl  | 2 +-
 2 files changed, 2 insertions(+), 1 deletion(-)

diff --git a/src/AMDGPU.jl b/src/AMDGPU.jl
index 79fda4125..9401b175e 100644
--- a/src/AMDGPU.jl
+++ b/src/AMDGPU.jl
@@ -60,6 +60,7 @@ include("discovery/discovery.jl")
 using .ROCmDiscovery
 using .ROCmDiscovery: AMDGPU_LLVM_Backend_jll, LLVMDowngrader_jll
 
+using KernelInterface
 include("utils.jl")
 
 include("hsa/HSA.jl")
diff --git a/src/utils.jl b/src/utils.jl
index 11a92547d..799380df8 100644
--- a/src/utils.jl
+++ b/src/utils.jl
@@ -97,7 +97,7 @@ function versioninfo(io::IO=stdout)
     println(io, "Julia packages: ")
     println(io, "- AMDGPU.jl: $(Base.pkgversion(AMDGPU))")
     for pkg in [:GPUArrays, :GPUCompiler, ("63c18a36-062a-441e-b654-da1e3ab1ce7c", "KernelAbstractions"),
-                 :LLVM, :AMDGPU_LLVM_Backend_jll, :LLVMDowngrader_jll]
+                 :KernelInterface, :LLVM, :AMDGPU_LLVM_Backend_jll, :LLVMDowngrader_jll]
         name, mod = get_module(pkg)
         isnothing(mod) || println(io, "- $(name): $(Base.pkgversion(mod))")
     end

From 7f91c42e3e5760451c53db91c984888ce8bac819 Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard@gmail.com>
Date: Mon, 28 Sep 2026 07:09:09 +0200
Subject: [PATCH 07/13] KernelInterface: report per-dimension launch limits

Implement KI.max_work_group_dims from the device's maxThreadsDim, cached in
HIPDevice since it is queried on every automatically-sized launch, and
KI.max_num_groups as the number of workgroups that fits the work-item grid
limit for any valid workgroup size. Requires KernelInterface 0.2.3.
---
 Project.toml            |  2 +-
 docs/src/api/devices.md |  1 +
 src/ROCKernels.jl       | 13 +++++++++++++
 src/hip/device.jl       | 11 ++++++++++-
 4 files changed, 25 insertions(+), 2 deletions(-)

diff --git a/Project.toml b/Project.toml
index d7d144f4b..84b46df90 100644
--- a/Project.toml
+++ b/Project.toml
@@ -64,7 +64,7 @@ GPUArrays = "11.5.14"
 GPUCompiler = "2.9"
 GPUToolbox = "3"
 KernelAbstractions = "0.9.2"
-KernelInterface = "0.2.1"
+KernelInterface = "0.2.3"
 LLVM = "9"
 LLVMDowngrader_jll = "0.11"
 Libdl = "1"
diff --git a/docs/src/api/devices.md b/docs/src/api/devices.md
index 20194f57f..39447b2b1 100644
--- a/docs/src/api/devices.md
+++ b/docs/src/api/devices.md
@@ -42,6 +42,7 @@ AMDGPU.device_id!
 AMDGPU.HIP.name
 AMDGPU.HIP.wavefrontsize
 AMDGPU.HIP.gcn_arch
+AMDGPU.HIP.max_workgroup_dims
 AMDGPU.HIP.device_id
 AMDGPU.HIP.properties
 ```
diff --git a/src/ROCKernels.jl b/src/ROCKernels.jl
index 137849c55..5f0a4108d 100644
--- a/src/ROCKernels.jl
+++ b/src/ROCKernels.jl
@@ -99,6 +99,19 @@ end
 function KI.max_work_group_size(::ROCBackend)::Int
     Int(HIP.attribute(AMDGPU.HIP.device(), AMDGPU.HIP.hipDeviceAttributeMaxThreadsPerBlock))
 end
+# queried on every automatically-sized launch, so use the limits cached in the device
+KI.max_work_group_dims(::ROCBackend)::NTuple{3, Int} = HIP.max_workgroup_dims(AMDGPU.device())
+# HIP takes the grid size in workgroups, but the dispatch packet holds it in work-items
+# (as a UInt32 per dimension, which HIP checks), and the device code assumes workgroup
+# indices fit in an Int32 (see `Device._max_groups`). Report the number of workgroups
+# that can be launched with any valid workgroup size. HIP's `maxGridSize` isn't usable:
+# depending on the ROCm version it holds CUDA's block limits or the work-item limits.
+function KI.max_num_groups(backend::ROCBackend)::NTuple{3, Int}
+    dims = KI.max_work_group_dims(backend)
+    return ntuple(Val(3)) do d
+        Int(min(Device._max_groups[d], Device._max_grid_size[d] ÷ dims[d]))
+    end
+end
 function KI.sub_group_size(::ROCBackend)::Int
     HIP.wavefrontsize(HIP.device())
 end
diff --git a/src/hip/device.jl b/src/hip/device.jl
index 61b9408c3..479aea449 100644
--- a/src/hip/device.jl
+++ b/src/hip/device.jl
@@ -3,6 +3,7 @@ struct HIPDevice
     device_id::Cint
     gcn_arch::String
     wavefrontsize::Cint
+    max_workgroup_dims::NTuple{3, Int}
 end
 
 const DEFAULT_DEVICE = Ref{Union{Nothing, HIPDevice}}(nothing)
@@ -18,7 +19,8 @@ function HIPDevice(device_id::Integer)
     gcn_arch = unsafe_string(pointer([props.gcnArchName...]))
 
     wavefrontsize = props.warpSize
-    HIPDevice(device_ref[], device_id, gcn_arch, wavefrontsize)
+    max_workgroup_dims = Int.(props.maxThreadsDim)
+    HIPDevice(device_ref[], device_id, gcn_arch, wavefrontsize, max_workgroup_dims)
 end
 
 """
@@ -36,6 +38,13 @@ Get size of the wavefront. AMD GPUs support either 32 or 64.
 """
 wavefrontsize(d::HIPDevice)::Cint = d.wavefrontsize
 
+"""
+    max_workgroup_dims(d::HIPDevice)::NTuple{3, Int}
+
+Get the maximum number of work-items along each dimension of a workgroup.
+"""
+max_workgroup_dims(d::HIPDevice)::NTuple{3, Int} = d.max_workgroup_dims
+
 """
     gcn_arch(d::HIPDevice)::String
 

From 5d52333fd645a4602650386008dff49c66025cce Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard@gmail.com>
Date: Mon, 28 Sep 2026 10:01:30 +0200
Subject: [PATCH 08/13] Port to KernelInterface 0.3

- subtype `KI.Backend`, and implement `KI.launch` instead of calling the kernel;
- implement the four primitive index queries with `% T`, and let KernelInterface derive
  the global ones;
- typed sub-group queries, `supports_subgroups` and `supports_shuffle`;
- `max_work_group_size(kernel)` is the kernel's limit, `launch_configuration` the
  occupancy-based recommendation;
- report `Float64` and atomics support, and the device of an array;
- `kernel_function` keeps the backend it was given;
- `copyto!` checks the lengths and returns the destination;
- use the generic `zeros` and `ones`;
- `_print` is documented as unsupported.

KernelInterface 0.3 isn't registered yet, so get it from its branch, and develop it from a
clone on Julia 1.10, which ignores `[sources]`.
---
 .buildkite/pipeline.yml       |  35 +++++++++++-
 Project.toml                  |   5 +-
 src/ROCKernels.jl             | 100 +++++++++++++++++-----------------
 test/Project.toml             |   1 +
 test/kernelinterface_tests.jl |   3 +-
 5 files changed, 88 insertions(+), 56 deletions(-)

diff --git a/.buildkite/pipeline.yml b/.buildkite/pipeline.yml
index 6a5489c49..b75b3e285 100644
--- a/.buildkite/pipeline.yml
+++ b/.buildkite/pipeline.yml
@@ -74,7 +74,6 @@ steps:
     matrix:
       setup:
         version:
-          - "1.10"
           - "1.11"
           - "1.12"
         test_args:
@@ -84,6 +83,32 @@ steps:
       (build.message =~ /\[only [^\]]*(tests|julia)/ ||
        build.message !~ /\[only / && !build.pull_request.draft)
 
+  # XXX: KernelInterface 0.3 isn't registered yet, and Julia 1.10 ignores [sources]: develop
+  #      it from a clone before instantiating, which the julia-test plugin would do first
+  - <<: *julia
+    matrix:
+      setup:
+        version:
+          - "1.10"
+        test_args:
+          - "--verbose"
+    plugins:
+      - JuliaCI/julia#v1:
+          version: "{{matrix.version}}"
+      - JuliaCI/julia-coverage#v1:
+          codecov: true
+    command: |
+      git clone --depth 1 --branch tb/ki-0.3 https://github.com/JuliaGPU/KernelAbstractions.jl ka
+      julia --project -e '
+        using Pkg
+        Pkg.develop(path="ka/lib/KernelInterface")
+        Pkg.instantiate()
+        Pkg.test(; coverage=true, test_args=`{{matrix.test_args}}`)'
+    if: |
+      build.message !~ /\[skip [^\]]*(tests|julia)/ &&
+      (build.message =~ /\[only [^\]]*(tests|julia)/ ||
+       build.message !~ /\[only / && !build.pull_request.draft)
+
   - <<: *julia
     matrix:
       setup:
@@ -127,10 +152,14 @@ steps:
     plugins:
       - JuliaCI/julia#v1:
           version: "1.10"
-      - JuliaCI/julia-test#v1:
-          run_tests: false
+    # XXX: KernelInterface 0.3 isn't registered yet, and Julia 1.10 ignores [sources]
     command: |
+      git clone --depth 1 --branch tb/ki-0.3 https://github.com/JuliaGPU/KernelAbstractions.jl ka
       julia --project -e '
+        using Pkg
+        Pkg.develop(path="ka/lib/KernelInterface")
+        Pkg.instantiate()
+
         using AMDGPU
         @assert !AMDGPU.functional()'
     agents:
diff --git a/Project.toml b/Project.toml
index 84b46df90..c37eb75e7 100644
--- a/Project.toml
+++ b/Project.toml
@@ -49,6 +49,9 @@ AMDGPUEnzymeCoreExt = "EnzymeCore"
 AMDGPUSparseMatricesCSRExt = "SparseMatricesCSR"
 AMDGPUSpecialFunctionsExt = "SpecialFunctions"
 
+[sources]
+KernelInterface = {url = "https://github.com/JuliaGPU/KernelAbstractions.jl", rev = "tb/ki-0.3", subdir = "lib/KernelInterface"}
+
 [compat]
 AMDGPU_LLVM_Backend_jll = "23"
 AbstractFFTs = "1.0"
@@ -64,7 +67,7 @@ GPUArrays = "11.5.14"
 GPUCompiler = "2.9"
 GPUToolbox = "3"
 KernelAbstractions = "0.9.2"
-KernelInterface = "0.2.3"
+KernelInterface = "0.3"
 LLVM = "9"
 LLVMDowngrader_jll = "0.11"
 Libdl = "1"
diff --git a/src/ROCKernels.jl b/src/ROCKernels.jl
index 5f0a4108d..1b4900d71 100644
--- a/src/ROCKernels.jl
+++ b/src/ROCKernels.jl
@@ -14,13 +14,14 @@ import LLVM
 using StaticArraysCore: MArray
 
 """
-    ROCBackend <: KernelAbstractions.GPU
+    ROCBackend <: KernelInterface.Backend
 
-KernelAbstractions backend that executes kernels on an AMD GPU via AMDGPU.jl.
-Pass `ROCBackend()` to a KernelAbstractions kernel to run it on the GPU, or
-obtain it from an array with `KernelAbstractions.get_backend(::ROCArray)`.
+KernelInterface backend that executes kernels on an AMD GPU via AMDGPU.jl.
+Obtain it from an array with `KernelInterface.get_backend(::ROCArray)`.
+
+Printing from a kernel (`KernelInterface._print`) is not supported: it does nothing.
 """
-struct ROCBackend <: KI.GPU end
+struct ROCBackend <: KI.Backend end
 
 KI.versioninfo(io::IO, ::ROCBackend) = AMDGPU.versioninfo(io)
 
@@ -54,8 +55,6 @@ end
 
 KI.unsafe_free!(x::AMDGPU.ROCArray) = AMDGPU.unsafe_free!(x)
 KI.allocate(::ROCBackend, ::Type{T}, dims::Tuple) where T = AMDGPU.ROCArray{T}(undef, dims)
-KI.zeros(::ROCBackend, ::Type{T}, dims::Tuple) where T = AMDGPU.zeros(T, dims)
-KI.ones(::ROCBackend, ::Type{T}, dims::Tuple) where T = AMDGPU.ones(T, dims)
 
 function KI.priority!(::ROCBackend, priority::Symbol)
     priority ∉ (:high, :normal, :low) && error(
@@ -64,10 +63,12 @@ function KI.priority!(::ROCBackend, priority::Symbol)
 end
 
 function KI.copyto!(::ROCBackend, A, B)
+    length(A) == length(B) ||
+        throw(ArgumentError("Arrays must have the same length, got $(length(A)) and $(length(B))"))
     GC.@preserve A B begin
         copyto!(A, 1, B, 1, length(A))
     end
-    return
+    return A
 end
 
 function KI.pagelock!(::ROCBackend, x::Array)
@@ -75,29 +76,35 @@ function KI.pagelock!(::ROCBackend, x::Array)
     return
 end
 
+KI.device(::ROCBackend, A::AMDGPU.ROCArray) = AMDGPU.device_id(AMDGPU.device(A))
+
+KI.supports_float64(::ROCBackend) = true
+KI.supports_atomics(::ROCBackend) = true
+
 KI.argconvert(::ROCBackend, arg) = rocconvert(arg)
 
-function KI.kernel_function(::ROCBackend, f::F, tt::TT=Tuple{}; name=nothing, kwargs...) where {F,TT}
+function KI.kernel_function(backend::ROCBackend, f::F, tt::TT=Tuple{}; name=nothing, kwargs...) where {F,TT}
     kern = hipfunction(f, tt; name, kwargs...)
-    KI.Kernel{ROCBackend, typeof(kern)}(ROCBackend(), kern)
+    KI.Kernel{ROCBackend, typeof(kern)}(backend, kern)
 end
 
-function (obj::KI.Kernel{ROCBackend})(args...; numworkgroups=(), workgroupsize=(), ndrange=(), max_work_group_size=typemax(Int))
-    KI.check_launch_args(numworkgroups, workgroupsize, ndrange)
-    prod(ndrange) == 0 && return nothing
-
-    numworkgroups, workgroupsize = KI.auto_launch_sizes(obj, numworkgroups, workgroupsize, ndrange, max_work_group_size)
-    obj.kern(args...; groupsize = workgroupsize, gridsize = numworkgroups)
-    return nothing
+function KI.launch(obj::KI.Kernel{ROCBackend}, groups::Dims{3}, items::Dims{3}, args...; kwargs...)
+    obj.kern(args...; groupsize = items, gridsize = groups, kwargs...)
+    return
 end
 
-function KI.kernel_max_work_group_size(kikern::KI.Kernel{<:ROCBackend}; max_work_items::Int=Int(typemax(Int32)))::Int
-    (; groupsize) = AMDGPU.launch_configuration(kikern.kern; max_block_size = max_work_items)
-
-    return Int(min(max_work_items, groupsize))
+function KI.max_work_group_size(kernel::KI.Kernel{ROCBackend})::Int
+    max_items = Ref{Cint}()
+    HIP.hipFuncGetAttribute(max_items, HIP.HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, kernel.kern.fun)
+    return Int(max_items[])
+end
+function KI.launch_configuration(kernel::KI.Kernel{ROCBackend}; max_work_group_size::Integer=typemax(Int))
+    max_items = min(max_work_group_size, KI.max_work_group_size(kernel))
+    (; groupsize) = AMDGPU.launch_configuration(kernel.kern; max_block_size = max_items)
+    return (; workgroupsize = Int(min(groupsize, max_items)))
 end
 function KI.max_work_group_size(::ROCBackend)::Int
-    Int(HIP.attribute(AMDGPU.HIP.device(), AMDGPU.HIP.hipDeviceAttributeMaxThreadsPerBlock))
+    Int(HIP.attribute(AMDGPU.device(), HIP.hipDeviceAttributeMaxThreadsPerBlock))
 end
 # queried on every automatically-sized launch, so use the limits cached in the device
 KI.max_work_group_dims(::ROCBackend)::NTuple{3, Int} = HIP.max_workgroup_dims(AMDGPU.device())
@@ -113,53 +120,47 @@ function KI.max_num_groups(backend::ROCBackend)::NTuple{3, Int}
     end
 end
 function KI.sub_group_size(::ROCBackend)::Int
-    HIP.wavefrontsize(HIP.device())
+    Int(HIP.wavefrontsize(AMDGPU.device()))
 end
 function KI.multiprocessor_count(::ROCBackend)::Int
-    Int(HIP.attribute(AMDGPU.HIP.device(), AMDGPU.HIP.hipDeviceAttributeMultiprocessorCount))
+    Int(HIP.attribute(AMDGPU.device(), HIP.hipDeviceAttributeMultiprocessorCount))
 end
 
-KI.shfl_down_types(::ROCBackend) = DataType[Bool,
-                                             UInt8, UInt16, UInt32, UInt64, UInt128,
-                                             Int8, Int16, Int32, Int64, Int128,
-                                             Float16, Float32, Float64,
-                                             ComplexF16, ComplexF32, ComplexF64]
+KI.supports_subgroups(::ROCBackend) = true
+# `shfl_down` decomposes other types into 32-bit shuffles
+KI.supports_shuffle(::ROCBackend, ::Type{T}) where {T} =
+    T <: Union{Bool, Base.BitInteger, Base.IEEEFloat, Complex{<:Union{Base.BitInteger, Base.IEEEFloat}}}
 
 # Indexing.
 ## COV_EXCL_START
+
+# computed with `% T`, which unlike `T(x)` has no error path
+
 @device_override @inline function KI.get_local_id(::Type{T}) where {T}
-    return (; x = T(AMDGPU.Device.workitemIdx().x), y = T(AMDGPU.Device.workitemIdx().y), z = T(AMDGPU.Device.workitemIdx().z))
+    return (; x = Device.workitemIdx().x % T, y = Device.workitemIdx().y % T, z = Device.workitemIdx().z % T)
 end
 
 @device_override @inline function KI.get_group_id(::Type{T}) where {T}
-    return (; x = T(AMDGPU.Device.workgroupIdx().x), y = T(AMDGPU.Device.workgroupIdx().y), z = T(AMDGPU.Device.workgroupIdx().z))
-end
-
-@device_override @inline function KI.get_global_id(::Type{T}) where {T}
-    return (; x = T((AMDGPU.Device.workgroupIdx().x-1)*AMDGPU.Device.blockDim().x + AMDGPU.Device.workitemIdx().x), y = T((AMDGPU.Device.workgroupIdx().y-1)*AMDGPU.Device.blockDim().y + AMDGPU.Device.workitemIdx().y), z = T((AMDGPU.Device.workgroupIdx().z-1)*AMDGPU.Device.blockDim().z + AMDGPU.Device.workitemIdx().z))
+    return (; x = Device.workgroupIdx().x % T, y = Device.workgroupIdx().y % T, z = Device.workgroupIdx().z % T)
 end
 
 @device_override @inline function KI.get_local_size(::Type{T}) where {T}
-    return (; x = T(AMDGPU.Device.workgroupDim().x), y = T(AMDGPU.Device.workgroupDim().y), z = T(AMDGPU.Device.workgroupDim().z))
+    return (; x = Device.workgroupDim().x % T, y = Device.workgroupDim().y % T, z = Device.workgroupDim().z % T)
 end
 
 @device_override @inline function KI.get_num_groups(::Type{T}) where {T}
-    return (; x = T(AMDGPU.Device.gridGroupDim().x), y = T(AMDGPU.Device.gridGroupDim().y), z = T(AMDGPU.Device.gridGroupDim().z))
+    return (; x = Device.gridGroupDim().x % T, y = Device.gridGroupDim().y % T, z = Device.gridGroupDim().z % T)
 end
 
-@device_override @inline function KI.get_global_size(::Type{T}) where {T}
-    return (; x = T(AMDGPU.Device.gridItemDim().x), y = T(AMDGPU.Device.gridItemDim().y), z = T(AMDGPU.Device.gridItemDim().z))
-end
-
-@device_override KI.get_sub_group_size() = UInt32(Device.wavefrontsize())
+@device_override KI.get_sub_group_size(::Type{T}) where {T} = Device.wavefrontsize() % T
 
-@device_override KI.get_max_sub_group_size() = UInt32(Device.wavefrontsize())
+@device_override KI.get_max_sub_group_size(::Type{T}) where {T} = Device.wavefrontsize() % T
 
-@device_override KI.get_num_sub_groups() = UInt32(prod(Device.blockDim()) ÷ Device.wavefrontsize())
+@device_override KI.get_num_sub_groups(::Type{T}) where {T} = (prod(Device.workgroupDim()) ÷ Device.wavefrontsize()) % T
 
-@device_override KI.get_sub_group_id() = UInt32(((Device.threadIdx().x - 1) + Device.blockDim().x * (Device.threadIdx().y - 1) + Device.blockDim().x * Device.blockDim().y * (Device.threadIdx().z - 1)) ÷ Device.wavefrontsize()) + 0x1
+@device_override KI.get_sub_group_id(::Type{T}) where {T} = (((Device.workitemIdx().x - 0x1) + Device.workgroupDim().x * (Device.workitemIdx().y - 0x1) + Device.workgroupDim().x * Device.workgroupDim().y * (Device.workitemIdx().z - 0x1)) ÷ Device.wavefrontsize() + 0x1) % T
 
-@device_override KI.get_sub_group_local_id() = UInt32(Device.activelane() + 0x1)
+@device_override KI.get_sub_group_local_id(::Type{T}) where {T} = (Device.activelane() + 0x1) % T
 
 # Shared memory.
 
@@ -179,12 +180,11 @@ end
 end
 
 @device_override function KI.shfl_down(val::T, offset::Integer) where T
-    @inline AMDGPU.Device.shfl_down(val, Cint(offset))
+    @inline AMDGPU.Device.shfl_down(val, offset % Cint)
 end
 
-@device_override @inline function KI._print(args...)
-    # TODO
-end
+# not supported, see the `ROCBackend` docstring
+@device_override @inline KI._print(args...) = nothing
 ## COV_EXCL_STOP
 
 end
diff --git a/test/Project.toml b/test/Project.toml
index af3919c3a..ba8e5bb93 100644
--- a/test/Project.toml
+++ b/test/Project.toml
@@ -28,3 +28,4 @@ Test = "8dfed614-e22c-5e08-85e1-65c5234f0b40"
 
 [sources]
 AMDGPU = {path = ".."}
+KernelInterface = {url = "https://github.com/JuliaGPU/KernelAbstractions.jl", rev = "tb/ki-0.3", subdir = "lib/KernelInterface"}
diff --git a/test/kernelinterface_tests.jl b/test/kernelinterface_tests.jl
index 2eba9d15f..a56967ac4 100644
--- a/test/kernelinterface_tests.jl
+++ b/test/kernelinterface_tests.jl
@@ -9,7 +9,6 @@ AMDGPU.allowscalar(false)
 
 @testset "kernelinterface" begin
 
-Testsuite.testsuite(
-    ROCInterface.ROCBackend, "ROCM", AMDGPU, ROCArray, AMDGPU.ROCDeviceArray)
+Testsuite.testsuite(ROCInterface.ROCBackend(), ROCArray)
 
 end

From e589274e7ead049c0abe135f1613adf46bc5d7f9 Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard@gmail.com>
Date: Mon, 28 Sep 2026 10:03:00 +0200
Subject: [PATCH 09/13] KernelInterface: fix the sub-group queries for partial
 and divergent wavefronts
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

- `get_num_sub_groups` counts a partial last wavefront (`cld` instead of `÷`);
- `get_sub_group_size` is the number of work-items in the wavefront, which is smaller than
  the wavefront size for the partial one;
- `get_sub_group_local_id` is the hardware lane (`mbcnt`), not the index among the active
  lanes (`activelane`), which changes in divergent code.
---
 src/ROCKernels.jl             | 28 ++++++++++++++++---
 test/kernelinterface_tests.jl | 51 ++++++++++++++++++++++++++++++++++-
 2 files changed, 74 insertions(+), 5 deletions(-)

diff --git a/src/ROCKernels.jl b/src/ROCKernels.jl
index 1b4900d71..2b2abb3b3 100644
--- a/src/ROCKernels.jl
+++ b/src/ROCKernels.jl
@@ -152,15 +152,35 @@ end
     return (; x = Device.gridGroupDim().x % T, y = Device.gridGroupDim().y % T, z = Device.gridGroupDim().z % T)
 end
 
-@device_override KI.get_sub_group_size(::Type{T}) where {T} = Device.wavefrontsize() % T
+# wavefronts are formed from consecutive linear work-item indices
+@inline function linear_workitem_id()
+    return (Device.workitemIdx().x - 0x1) +
+           (Device.workitemIdx().y - 0x1) * Device.workgroupDim().x +
+           (Device.workitemIdx().z - 0x1) * Device.workgroupDim().x * Device.workgroupDim().y
+end
+
+@inline workgroup_items() = Device.workgroupDim().x * Device.workgroupDim().y * Device.workgroupDim().z
+
+# the 0-based lane of the work-item in its wavefront. unlike `Device.activelane()`, this
+# doesn't depend on which lanes are active. `mbcnt.hi` adds nothing in wave32 mode.
+@inline function hardware_lane()
+    lo = ccall("llvm.amdgcn.mbcnt.lo", llvmcall, UInt32, (UInt32, UInt32), typemax(UInt32), 0x0)
+    return ccall("llvm.amdgcn.mbcnt.hi", llvmcall, UInt32, (UInt32, UInt32), typemax(UInt32), lo)
+end
+
+# the last wavefront of a workgroup can be partial
+@device_override @inline function KI.get_sub_group_size(::Type{T}) where {T}
+    ws = Device.wavefrontsize()
+    return min(ws, workgroup_items() - (linear_workitem_id() ÷ ws) * ws) % T
+end
 
 @device_override KI.get_max_sub_group_size(::Type{T}) where {T} = Device.wavefrontsize() % T
 
-@device_override KI.get_num_sub_groups(::Type{T}) where {T} = (prod(Device.workgroupDim()) ÷ Device.wavefrontsize()) % T
+@device_override KI.get_num_sub_groups(::Type{T}) where {T} = cld(workgroup_items(), Device.wavefrontsize()) % T
 
-@device_override KI.get_sub_group_id(::Type{T}) where {T} = (((Device.workitemIdx().x - 0x1) + Device.workgroupDim().x * (Device.workitemIdx().y - 0x1) + Device.workgroupDim().x * Device.workgroupDim().y * (Device.workitemIdx().z - 0x1)) ÷ Device.wavefrontsize() + 0x1) % T
+@device_override KI.get_sub_group_id(::Type{T}) where {T} = (linear_workitem_id() ÷ Device.wavefrontsize() + 0x1) % T
 
-@device_override KI.get_sub_group_local_id(::Type{T}) where {T} = (Device.activelane() + 0x1) % T
+@device_override KI.get_sub_group_local_id(::Type{T}) where {T} = (hardware_lane() + 0x1) % T
 
 # Shared memory.
 
diff --git a/test/kernelinterface_tests.jl b/test/kernelinterface_tests.jl
index a56967ac4..59397ac20 100644
--- a/test/kernelinterface_tests.jl
+++ b/test/kernelinterface_tests.jl
@@ -3,12 +3,61 @@ using AMDGPU
 using AMDGPU: ROCInterface
 
 import KernelInterface
+import KernelInterface as KI
 include(joinpath(pkgdir(KernelInterface), "test", "testsuite.jl"))
 
 AMDGPU.allowscalar(false)
 
+function ki_subgroup_kernel(num, sizes, id, lane)
+    l = KI.get_local_id()
+    s = KI.get_local_size()
+    i = l.x + (l.y - 1) * s.x
+    @inbounds begin
+        num[i] = KI.get_num_sub_groups()
+        sizes[i] = KI.get_sub_group_size()
+        id[i] = KI.get_sub_group_id()
+        lane[i] = KI.get_sub_group_local_id()
+    end
+    return
+end
+
+# only the odd lanes are active when querying the lane id
+function ki_divergent_lane_kernel(lane)
+    i = KI.get_local_id().x
+    if isodd(i)
+        @inbounds lane[i] = KI.get_sub_group_local_id()
+    end
+    return
+end
+
 @testset "kernelinterface" begin
 
-Testsuite.testsuite(ROCInterface.ROCBackend(), ROCArray)
+backend = ROCInterface.ROCBackend()
+Testsuite.testsuite(backend, ROCArray)
+
+# KernelInterface leaves the formation of sub-groups unspecified; AMD GPUs form wavefronts
+# from consecutive linear work-item indices
+@testset "partial sub-groups" begin
+    ws = KI.sub_group_size(backend)
+    # a (ws + 1)x2 workgroup is made up of 3 wavefronts, the last one only partially filled
+    workgroupsize = (ws + 1, 2)
+    n = prod(workgroupsize)
+    num = ROCArray{UInt32}(undef, n)
+    sizes = ROCArray{UInt32}(undef, n)
+    id = ROCArray{UInt32}(undef, n)
+    lane = ROCArray{UInt32}(undef, n)
+    KI.@launch backend workgroupsize=workgroupsize ki_subgroup_kernel(num, sizes, id, lane)
+    @test all(==(3), Array(num))
+    @test Array(sizes) == [i < 2ws ? ws : 2 for i in 0:n-1]
+    @test Array(id) == [div(i, ws) + 1 for i in 0:n-1]
+    @test Array(lane) == [rem(i, ws) + 1 for i in 0:n-1]
+end
+
+@testset "lane ids under divergence" begin
+    ws = KI.sub_group_size(backend)
+    lane = AMDGPU.zeros(UInt32, 2ws)
+    KI.@launch backend workgroupsize=2ws ki_divergent_lane_kernel(lane)
+    @test Array(lane) == [isodd(i) ? mod1(i, ws) : 0 for i in 1:2ws]
+end
 
 end

From e65181240a5cb46ea72bb216f64ca2d18a53d1d6 Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard@gmail.com>
Date: Mon, 28 Sep 2026 10:04:23 +0200
Subject: [PATCH 10/13] Order memory in `sync_wavefront`

`llvm.amdgcn.wave.barrier` only keeps the compiler from moving code across it, so writes
before it weren't guaranteed to be visible to the other lanes afterwards. Fence it at
wavefront scope, like `sync_workgroup` does at workgroup scope. This makes
`KernelInterface.sub_group_barrier` fence memory, as KernelInterface 0.3 requires.
---
 docs/src/api/intrinsics.md        |  1 +
 src/AMDGPU.jl                     |  1 +
 src/device/gcn/synchronization.jl | 15 ++++++++++-----
 3 files changed, 12 insertions(+), 5 deletions(-)

diff --git a/docs/src/api/intrinsics.md b/docs/src/api/intrinsics.md
index 5a957110b..07b91f39f 100644
--- a/docs/src/api/intrinsics.md
+++ b/docs/src/api/intrinsics.md
@@ -23,6 +23,7 @@ AMDGPU.Device.blockDim
 
 ```@docs
 AMDGPU.sync_workgroup
+AMDGPU.sync_wavefront
 AMDGPU.sync_workgroup_count
 AMDGPU.sync_workgroup_and
 AMDGPU.sync_workgroup_or
diff --git a/src/AMDGPU.jl b/src/AMDGPU.jl
index 9401b175e..53d4e6dfb 100644
--- a/src/AMDGPU.jl
+++ b/src/AMDGPU.jl
@@ -83,6 +83,7 @@ Base.Experimental.@MethodTable(method_table)
 #needs to be before Device since sync uses this
 const syncscope_agent = UnsafeAtomics.Internal.LLVMSyncScope{:agent}()
 const syncscope_workgroup = UnsafeAtomics.Internal.LLVMSyncScope{:workgroup}()
+const syncscope_wavefront = UnsafeAtomics.Internal.LLVMSyncScope{:wavefront}()
 
 # Referenced by the generated `kernel_state()`, and generators run in the world they are
 # defined in, so this has to precede the device code.
diff --git a/src/device/gcn/synchronization.jl b/src/device/gcn/synchronization.jl
index 5e40d9cac..bdad1c110 100644
--- a/src/device/gcn/synchronization.jl
+++ b/src/device/gcn/synchronization.jl
@@ -1,6 +1,6 @@
 
 for ord in UnsafeAtomics.Internal.orderings
-    for sync in (AMDGPU.syncscope_agent, AMDGPU.syncscope_workgroup)
+    for sync in (AMDGPU.syncscope_agent, AMDGPU.syncscope_workgroup, AMDGPU.syncscope_wavefront)
         @eval @device_function function UnsafeAtomics.fence(::$(typeof(ord)), ::$(typeof(sync)))
             Base.llvmcall(
                     $("""
@@ -32,14 +32,19 @@ end
 """
     sync_wavefront()
 
-Waits until all wavefronts in a workgroup have reached this call and that their memory accesses are visible to other threads in the workgroup.
+Waits until all lanes of the wavefront have reached this call, and makes their memory
+accesses before it visible to the other lanes of the wavefront.
 """
-@inline function sync_wavefront()
-    # This is a no-op https://github.com/llvm/llvm-project/blob/88b77d5eaa66747538a12c9876eeffdce31ddb71/openmp/device/src/Synchronization.cpp#L136-L140
+@device_function @inline function sync_wavefront()
+    # the lanes of a wavefront execute in lockstep, so the barrier doesn't generate any
+    # code (https://github.com/llvm/llvm-project/blob/88b77d5eaa66747538a12c9876eeffdce31ddb71/openmp/device/src/Synchronization.cpp#L136-L140),
+    # but it keeps the compiler from moving code across it. like `sync_workgroup`, it needs
+    # fences to order memory.
+    UnsafeAtomics.fence(UnsafeAtomics.seq_cst, AMDGPU.syncscope_wavefront)
     ccall("llvm.amdgcn.wave.barrier", llvmcall, Cvoid, ())
+    UnsafeAtomics.fence(UnsafeAtomics.seq_cst, AMDGPU.syncscope_wavefront)
 end
 
-
 """
     sync_workgroup_count(predicate::Cint)::Cint
 

From 20a179e60280041cc40afe5b647187281a2ef728 Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard@gmail.com>
Date: Mon, 28 Sep 2026 10:05:01 +0200
Subject: [PATCH 11/13] KernelInterface: guarantee the wavefront size of
 compiled kernels

`KI.sub_group_size` promises that kernels from `KI.kernel_function` execute with that
sub-group width. `hipfunction` compiles for the device's wavefront size by default, but
`wavefrontsize64` could override it; reject a conflicting value.
---
 src/ROCKernels.jl             |  5 +++++
 test/kernelinterface_tests.jl | 18 ++++++++++++++++++
 2 files changed, 23 insertions(+)

diff --git a/src/ROCKernels.jl b/src/ROCKernels.jl
index 2b2abb3b3..0bdb8dbbc 100644
--- a/src/ROCKernels.jl
+++ b/src/ROCKernels.jl
@@ -84,6 +84,11 @@ KI.supports_atomics(::ROCBackend) = true
 KI.argconvert(::ROCBackend, arg) = rocconvert(arg)
 
 function KI.kernel_function(backend::ROCBackend, f::F, tt::TT=Tuple{}; name=nothing, kwargs...) where {F,TT}
+    # kernels have to execute with the wavefront size that `KI.sub_group_size` reports,
+    # the device's, which is also what `hipfunction` compiles for by default
+    if haskey(kwargs, :wavefrontsize64) && kwargs[:wavefrontsize64] != (KI.sub_group_size(backend) == 64)
+        throw(ArgumentError("`wavefrontsize64=$(kwargs[:wavefrontsize64])` conflicts with the wavefront size of the device, $(KI.sub_group_size(backend))"))
+    end
     kern = hipfunction(f, tt; name, kwargs...)
     KI.Kernel{ROCBackend, typeof(kern)}(backend, kern)
 end
diff --git a/test/kernelinterface_tests.jl b/test/kernelinterface_tests.jl
index 59397ac20..a22ba45ba 100644
--- a/test/kernelinterface_tests.jl
+++ b/test/kernelinterface_tests.jl
@@ -53,6 +53,24 @@ Testsuite.testsuite(backend, ROCArray)
     @test Array(lane) == [rem(i, ws) + 1 for i in 0:n-1]
 end
 
+function ki_wavefront_size_kernel(ws)
+    @inbounds ws[1] = KI.get_max_sub_group_size()
+    return
+end
+
+@testset "wavefront size" begin
+    ws = KI.sub_group_size(backend)
+    @test ws in (32, 64)
+    out = AMDGPU.zeros(Int, 1)
+    KI.@launch backend ki_wavefront_size_kernel(out)
+    @test Array(out)[1] == ws
+
+    # compiling for another wavefront size would break the guarantee
+    tt = Tuple{typeof(KI.argconvert(backend, out))}
+    @test_throws ArgumentError KI.kernel_function(backend, ki_wavefront_size_kernel, tt; wavefrontsize64 = ws == 32)
+    @test KI.kernel_function(backend, ki_wavefront_size_kernel, tt; wavefrontsize64 = ws == 64) isa KI.Kernel
+end
+
 @testset "lane ids under divergence" begin
     ws = KI.sub_group_size(backend)
     lane = AMDGPU.zeros(UInt32, 2ws)

From 21fc97c08a2ffaeb09bb74f733cffcd4400636d2 Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard@gmail.com>
Date: Mon, 28 Sep 2026 12:29:44 +0200
Subject: [PATCH 12/13] KernelInterface: specialize KI.launch on its arguments
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

Julia doesn't specialize a method on `args...` that it only passes through, which
made every launch through KernelInterface's generic launch dispatch dynamically
(+1.2 µs and +1 kB per launch on CUDA).
---
 src/ROCKernels.jl | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/src/ROCKernels.jl b/src/ROCKernels.jl
index 0bdb8dbbc..b6d5a7106 100644
--- a/src/ROCKernels.jl
+++ b/src/ROCKernels.jl
@@ -93,7 +93,7 @@ function KI.kernel_function(backend::ROCBackend, f::F, tt::TT=Tuple{}; name=noth
     KI.Kernel{ROCBackend, typeof(kern)}(backend, kern)
 end
 
-function KI.launch(obj::KI.Kernel{ROCBackend}, groups::Dims{3}, items::Dims{3}, args...; kwargs...)
+function KI.launch(obj::KI.Kernel{ROCBackend}, groups::Dims{3}, items::Dims{3}, args::Vararg{Any, N}; kwargs...) where {N}
     obj.kern(args...; groupsize = items, gridsize = groups, kwargs...)
     return
 end

From f915ef135a2af4cbc3dcb274ce8680535b0afb38 Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard@gmail.com>
Date: Wed, 30 Sep 2026 09:45:41 +0200
Subject: [PATCH 13/13] Accept the problem size in KI.launch_configuration

KernelInterface 0.3 passes the number of work-items to launch as `nitems`,
separately from the bound on the work-group size.
---
 src/ROCKernels.jl | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

diff --git a/src/ROCKernels.jl b/src/ROCKernels.jl
index b6d5a7106..1c65272e9 100644
--- a/src/ROCKernels.jl
+++ b/src/ROCKernels.jl
@@ -103,7 +103,8 @@ function KI.max_work_group_size(kernel::KI.Kernel{ROCBackend})::Int
     HIP.hipFuncGetAttribute(max_items, HIP.HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, kernel.kern.fun)
     return Int(max_items[])
 end
-function KI.launch_configuration(kernel::KI.Kernel{ROCBackend}; max_work_group_size::Integer=typemax(Int))
+function KI.launch_configuration(kernel::KI.Kernel{ROCBackend}; nitems::Union{Integer,Nothing}=nothing,
+                                 max_work_group_size::Integer=typemax(Int))
     max_items = min(max_work_group_size, KI.max_work_group_size(kernel))
     (; groupsize) = AMDGPU.launch_configuration(kernel.kern; max_block_size = max_items)
     return (; workgroupsize = Int(min(groupsize, max_items)))