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 8bc9180ef..c37eb75e7 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"
@@ -48,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"
@@ -63,6 +67,7 @@ GPUArrays = "11.5.14"
 GPUCompiler = "2.9"
 GPUToolbox = "3"
 KernelAbstractions = "0.9.2"
+KernelInterface = "0.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/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 636b3bf27..53d4e6dfb 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")
@@ -82,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.
@@ -102,7 +104,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
@@ -147,10 +149,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..1c65272e9 100644
--- a/src/ROCKernels.jl
+++ b/src/ROCKernels.jl
@@ -1,201 +1,216 @@
-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
 
 """
-    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 <: KA.GPU end
+struct ROCBackend <: KI.Backend end
+
+KI.versioninfo(io::IO, ::ROCBackend) = AMDGPU.versioninfo(io)
 
-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.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)
+function KI.record_event(::ROCBackend)
+    return HIP.HIPEvent(AMDGPU.stream())
+end
 
-function KA.priority!(::ROCBackend, priority::Symbol)
+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)
+
+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)
+    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 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
+KI.device(::ROCBackend, A::AMDGPU.ROCArray) = AMDGPU.device_id(AMDGPU.device(A))
 
-    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.supports_float64(::ROCBackend) = true
+KI.supports_atomics(::ROCBackend) = true
 
-    return ndrange, workgroupsize, iterspace, dynamic
-end
+KI.argconvert(::ROCBackend, arg) = rocconvert(arg)
 
-function threads_to_workgroupsize(threads, ndrange)
-    total = 1
-    return map(ndrange) do n
-        x = min(div(threads, total), n)
-        total *= x
-        return x
+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
 
-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)
+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
 
-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(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}; 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)))
+end
+function KI.max_work_group_size(::ROCBackend)::Int
+    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())
+# 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 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
+    Int(HIP.wavefrontsize(AMDGPU.device()))
 end
+function KI.multiprocessor_count(::ROCBackend)::Int
+    Int(HIP.attribute(AMDGPU.device(), HIP.hipDeviceAttributeMultiprocessorCount))
+end
+
+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
 
-@device_override @inline function KA.__index_Local_Linear(ctx)
-    return AMDGPU.Device.threadIdx().x
+# 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 = Device.workitemIdx().x % T, y = Device.workitemIdx().y % T, z = Device.workitemIdx().z % T)
 end
 
-@device_override @inline function KA.__index_Group_Linear(ctx)
-    return AMDGPU.Device.blockIdx().x
+@device_override @inline function KI.get_group_id(::Type{T}) where {T}
+    return (; x = Device.workgroupIdx().x % T, y = Device.workgroupIdx().y % T, z = Device.workgroupIdx().z % T)
 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_local_size(::Type{T}) where {T}
+    return (; x = Device.workgroupDim().x % T, y = Device.workgroupDim().y % T, z = Device.workgroupDim().z % T)
 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_num_groups(::Type{T}) where {T}
+    return (; x = Device.gridGroupDim().x % T, y = Device.gridGroupDim().y % T, z = Device.gridGroupDim().z % T)
 end
 
-@device_override @inline function KA.__index_Group_Cartesian(ctx)
-    @inbounds KA.blocks(KA.__iterspace(ctx))[AMDGPU.Device.blockIdx().x]
+# 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
 
-@device_override @inline function KA.__index_Global_Cartesian(ctx)
-    return @inbounds KA.expand(KA.__iterspace(ctx), AMDGPU.Device.blockIdx().x, AMDGPU.Device.threadIdx().x)
+@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
 
-@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
+# 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} = cld(workgroup_items(), Device.wavefrontsize()) % 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} = (hardware_lane() + 0x1) % T
+
 # 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...)
-    # TODO
+@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, offset % Cint)
 end
 
+# not supported, see the `ROCBackend` docstring
+@device_override @inline KI._print(args...) = nothing
+## 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/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)
 
diff --git a/src/device/gcn/synchronization.jl b/src/device/gcn/synchronization.jl
index 6d8b682dc..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(
                     $("""
@@ -29,6 +29,22 @@ 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 lanes of the wavefront have reached this call, and makes their memory
+accesses before it visible to the other lanes of the wavefront.
+"""
+@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
 
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
 
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
diff --git a/test/Project.toml b/test/Project.toml
index 8a24df04d..ba8e5bb93 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"
@@ -27,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
new file mode 100644
index 000000000..a22ba45ba
--- /dev/null
+++ b/test/kernelinterface_tests.jl
@@ -0,0 +1,81 @@
+using Test
+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
+
+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
+
+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)
+    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