diff --git a/.github/workflows/Test.yml b/.github/workflows/Test.yml
index aa691a2a8..a7b7560a1 100644
--- a/.github/workflows/Test.yml
+++ b/.github/workflows/Test.yml
@@ -17,7 +17,7 @@ jobs:
   test:
     name: Julia ${{ matrix.version }} - ${{ matrix.os }}
     runs-on: ${{ matrix.os }}
-    timeout-minutes: 60
+    timeout-minutes: 90
     strategy:
       fail-fast: false
       matrix:
diff --git a/Project.toml b/Project.toml
index 1bd6c0337..88649a38c 100644
--- a/Project.toml
+++ b/Project.toml
@@ -10,6 +10,7 @@ AcceleratedKernels = "6a4ca0a5-0e36-4168-a932-d9be78d558f1"
 Adapt = "79e6a3ab-5dfb-504d-930d-738a2a938a0e"
 GPUArraysCore = "46192b85-c4d5-4398-a991-12ede77f4527"
 KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c"
+KernelInterface = "4ee993da-d684-4d17-a7dd-4e58e78d92bf"
 LLVM = "929cbde3-209d-540e-8aea-75f648917ca0"
 LinearAlgebra = "37e2e46d-f89d-539d-b4ee-838fcccc9c8e"
 Printf = "de0858da-6303-5e67-8744-51eddeeeb8d7"
@@ -31,7 +32,8 @@ AcceleratedKernels = "0.5"
 Adapt = "4.7.3"
 GPUArraysCore = "= 0.2.1"
 JLD2 = "0.4, 0.5, 0.6"
-KernelAbstractions = "0.9.43, 0.10"
+KernelAbstractions = "0.10"
+KernelInterface = "0.4.1"
 LLVM = "3.9, 4, 5, 6, 7, 8, 9, 10"
 LinearAlgebra = "1"
 Printf = "1"
diff --git a/lib/JLArrays/Project.toml b/lib/JLArrays/Project.toml
index 698fb8ebd..dff9d5184 100644
--- a/lib/JLArrays/Project.toml
+++ b/lib/JLArrays/Project.toml
@@ -1,12 +1,13 @@
 name = "JLArrays"
 uuid = "27aeb0d3-9eb9-45fb-866b-73c2ecf80fcb"
-version = "0.4.2"
+version = "0.5"
 authors = ["Tim Besard <tim.besard@gmail.com>"]
 
 [deps]
 Adapt = "79e6a3ab-5dfb-504d-930d-738a2a938a0e"
 GPUArrays = "0c68f7d7-f131-5f86-a1c3-88cf8149b2d7"
 KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c"
+KernelInterface = "4ee993da-d684-4d17-a7dd-4e58e78d92bf"
 LinearAlgebra = "37e2e46d-f89d-539d-b4ee-838fcccc9c8e"
 Random = "9a3f8284-a2c9-5f02-9a11-845980a1fd5c"
 SparseArrays = "2f01184e-e22b-5df5-ae63-d93ebab69eaf"
@@ -14,7 +15,8 @@ SparseArrays = "2f01184e-e22b-5df5-ae63-d93ebab69eaf"
 [compat]
 Adapt = "2.0, 3.0, 4.0"
 GPUArrays = "12.1"
-KernelAbstractions = "0.9"
+KernelAbstractions = "0.10"
+KernelInterface = "0.4.2"
 LinearAlgebra = "1"
 Random = "1"
 SparseArrays = "1"
diff --git a/lib/JLArrays/src/JLArrays.jl b/lib/JLArrays/src/JLArrays.jl
index c6caf8524..f8695ddaa 100644
--- a/lib/JLArrays/src/JLArrays.jl
+++ b/lib/JLArrays/src/JLArrays.jl
@@ -16,11 +16,7 @@ using Random
 
 import GPUArrays: dense_array_type
 
-import KernelAbstractions
-
-@static if isdefined(JLArrays.KernelAbstractions, :POCL) # KA v0.10
-    import KernelAbstractions: POCL
-end
+import KernelAbstractions: KernelAbstractions, POCL
 
 module AS
 
@@ -44,4 +40,7 @@ include("findall.jl")
 include("JLKernels.jl")
 import .JLKernels: JLBackend
 
+# code reflection: JLArray kernels are compiled by POCL, for the device-side argument types
+import .POCL: code_typed, code_warntype, code_llvm, code_native
+
 end
diff --git a/lib/JLArrays/src/JLKernels.jl b/lib/JLArrays/src/JLKernels.jl
index ca4d84d0f..bee588748 100644
--- a/lib/JLArrays/src/JLKernels.jl
+++ b/lib/JLArrays/src/JLKernels.jl
@@ -3,8 +3,8 @@ module JLKernels
 using ..JLArrays
 using ..JLArrays: jlconvert
 
-import KernelAbstractions
-import KernelAbstractions: Kernel, StaticSize, DynamicSize, partition, launch_config
+import KernelAbstractions: KernelAbstractions, POCLBackend
+import KernelInterface as KI
 
 import Adapt
 
@@ -13,73 +13,65 @@ import Adapt
 
 export JLBackend
 
-const MAXTHREADS = 256
+struct JLBackend <: KI.Backend end
 
-struct JLBackend <: KernelAbstractions.GPU
-    static::Bool
-    JLBackend(;static::Bool=false) = new(static)
-end
+KI.get_backend(a::JLA) where JLA <: JLArray = JLBackend()
+KI.get_backend(a::JLA) where JLA <: Union{JLSparseMatrixCSC, JLSparseMatrixCSR, JLSparseVector} = JLBackend()
 
-KernelAbstractions.get_backend(a::JLA) where JLA <: JLArray = JLBackend()
-KernelAbstractions.get_backend(a::JLA) where JLA <: Union{JLSparseMatrixCSC, JLSparseMatrixCSR, JLSparseVector} = JLBackend()
+KI.allocate(::JLBackend, ::Type{T}, dims::Tuple) where T = JLArray{T}(undef, dims)
 
-function KernelAbstractions.mkcontext(kernel::Kernel{JLBackend}, I, _ndrange, iterspace, ::Dynamic) where Dynamic
-    return KernelAbstractions.CompilerMetadata{KernelAbstractions.ndrange(kernel), Dynamic}(I, _ndrange, iterspace)
-end
+## KernelInterface
+
+# JLArrays are executed by POCL: convert arguments to their POCL-compatible (CPU)
+# counterparts, and forward compilation and launches to the POCL backend. KernelAbstractions
+# kernels are launched through these methods too.
+
+KI.functional(::JLBackend) = true
 
-KernelAbstractions.allocate(::JLBackend, ::Type{T}, dims::Tuple) where T = JLArray{T}(undef, dims)
-
-@inline function launch_config(kernel::Kernel{JLBackend}, ndrange, workgroupsize)
-    if ndrange isa Integer
-        ndrange = (ndrange,)
-    end
-    if workgroupsize isa Integer
-        workgroupsize = (workgroupsize, )
-    end
-
-    if KernelAbstractions.workgroupsize(kernel) <: DynamicSize && workgroupsize === nothing
-        workgroupsize = (MAXTHREADS,) # Vectorization, 4x unrolling, minimal grain size
-    end
-    iterspace, dynamic = partition(kernel, ndrange, workgroupsize)
-    # partition checked that the ndrange's agreed
-    if KernelAbstractions.ndrange(kernel) <: StaticSize
-        ndrange = nothing
-    end
-
-    return ndrange, workgroupsize, iterspace, dynamic
+# POCL launches synchronously, and JLArrays live in host memory
+KI.synchronize(::JLBackend) = nothing
+
+function KI.copyto!(::JLBackend, A, B)
+    length(A) == length(B) ||
+        throw(ArgumentError("Arrays must match in length, got $(length(A)) and $(length(B))"))
+    Base.copyto!(A, B)
+    return A
 end
 
-@static if isdefined(KernelAbstractions, :isgpu) # KA v0.9
-    KernelAbstractions.isgpu(b::JLBackend) = false
+KI.argconvert(::JLBackend, arg) = KI.argconvert(POCLBackend(), jlconvert(arg))
+
+# the POCL kernel, and the callable it was compiled from, which has to be kept alive since
+# its converted form only holds pointers to the arrays it captures
+struct JLKernel{K, F}
+    kernel::K
+    f::F
 end
 
-@static if !isdefined(KernelAbstractions, :POCL) # KA v0.9
-    function convert_to_cpu(obj::Kernel{JLBackend, W, N, F}) where {W, N, F}
-        return Kernel{typeof(KernelAbstractions.CPU(; static = obj.backend.static)), W, N, F}(KernelAbstractions.CPU(; static = obj.backend.static), obj.f)
-    end
-else
-    function convert_to_cpu(obj::Kernel{JLBackend, W, N, F}) where {W, N, F}
-        return Kernel{typeof(KernelAbstractions.POCLBackend()), W, N, F}(KernelAbstractions.POCLBackend(), obj.f)
-    end
+function KI.kernel_function(backend::JLBackend, f::F, tt::TT = Tuple{}; kwargs...) where {F, TT}
+    kernel = KI.kernel_function(POCLBackend(), jlconvert(f), tt; kwargs...)
+    kern = JLKernel(kernel, f)
+    return KI.Kernel{JLBackend, typeof(kern)}(backend, kern)
 end
 
-function (obj::Kernel{JLBackend})(args...; ndrange=nothing, workgroupsize=nothing)
-    ndrange, workgroupsize, _, _ = launch_config(obj, ndrange, workgroupsize)
+function KI.launch(obj::KI.Kernel{JLBackend}, groups::Dims{3}, items::Dims{3}, args::Tuple; kwargs...)
     # Use `map` rather than `jlconvert.(args)` to skip the broadcast
     # machinery (broadcasted/materialize/ntuple) that would otherwise
     # be specialized per unique arg-tuple type.
-    device_args = map(jlconvert, args)
-    new_obj = convert_to_cpu(obj)
-    new_obj(device_args...; ndrange, workgroupsize)
+    f = obj.kern.f
+    GC.@preserve f args KI.launch(obj.kern.kernel, groups, items, map(jlconvert, args); kwargs...)
 end
 
-Adapt.adapt_storage(::JLBackend, a::Array) = Adapt.adapt(JLArrays.JLArray, a)
-Adapt.adapt_storage(::JLBackend, a::JLArrays.JLArray) = a
-
-@static if !isdefined(KernelAbstractions, :POCL) # KA v0.9
-    Adapt.adapt_storage(::KernelAbstractions.CPU, a::JLArrays.JLArray) = convert(Array, a)
-else
-    Adapt.adapt_storage(::KernelAbstractions.POCLBackend, a::JLArrays.JLArray) = convert(Array, a)
+KI.max_work_group_size(kernel::KI.Kernel{JLBackend}) = KI.max_work_group_size(kernel.kern.kernel)
+KI.launch_configuration(kernel::KI.Kernel{JLBackend}; kwargs...) =
+    KI.launch_configuration(kernel.kern.kernel; kwargs...)
+for f in (:max_work_group_size, :max_work_group_dims, :max_num_groups, :sub_group_size,
+          :multiprocessor_count, :supports_subgroups, :supports_linear_subgroups,
+          :supports_independent_subgroups, :supports_partitionable_subgroups,
+          :supports_float64, :supports_atomics)
+    @eval KI.$f(::JLBackend) = KI.$f(POCLBackend())
 end
+KI.supports_shuffle(::JLBackend, ::Type{T}) where {T} = KI.supports_shuffle(POCLBackend(), T)
+
+Adapt.adapt_storage(::JLBackend, x) = Adapt.adapt(JLArray, x)
 
 end
diff --git a/lib/JLArrays/src/array.jl b/lib/JLArrays/src/array.jl
index d47184a8b..31875c9d7 100644
--- a/lib/JLArrays/src/array.jl
+++ b/lib/JLArrays/src/array.jl
@@ -172,14 +172,9 @@ Adapt.adapt_storage(::Type{<:JLArray{T}}, xs::AbstractArray) where {T} =
 Adapt.adapt_storage(::Type{Array}, xs::JLArray) = convert(Array, xs)
 
 # adapt for the device
-@static if !isdefined(JLArrays.KernelAbstractions, :POCL) # KA v0.9
-    Adapt.adapt_storage(::Adaptor, x::JLArray{T,N}) where {T,N} =
-      JLDeviceArray{T,N}(x.data[], x.offset, x.dims)
-else
-    function Adapt.adapt_storage(::Adaptor, x::JLArray{T,N}) where {T,N}
-        arr = typed_data(x)
-        Adapt.adapt_storage(POCL.KernelAdaptor([pointer(arr)]), arr)
-    end
+function Adapt.adapt_storage(::Adaptor, x::JLArray{T,N}) where {T,N}
+    arr = typed_data(x)
+    Adapt.adapt_storage(POCL.KernelAdaptor([pointer(arr)]), arr)
 end
 
 ## conversions
diff --git a/lib/JLArrays/src/device/array.jl b/lib/JLArrays/src/device/array.jl
index 457697782..a6267d23c 100644
--- a/lib/JLArrays/src/device/array.jl
+++ b/lib/JLArrays/src/device/array.jl
@@ -9,30 +9,3 @@ struct JlRefValue{T} <: Ref{T}
 end
 Base.getindex(r::JlRefValue) = r.x
 Adapt.adapt_structure(to::Adaptor, r::Base.RefValue) = JlRefValue(adapt(to, r[]))
-
-## executed on-device
-
-# array type
-@static if !isdefined(JLArrays.KernelAbstractions, :POCL) # KA v0.9
-    struct JLDeviceArray{T, N} <: AbstractDeviceArray{T, N}
-        data::Vector{UInt8}
-        offset::Int
-        dims::Dims{N}
-    end
-
-    Base.elsize(::Type{<:JLDeviceArray{T}}) where {T} = sizeof(T)
-
-    Base.size(x::JLDeviceArray) = x.dims
-    Base.sizeof(x::JLDeviceArray) = Base.elsize(x) * length(x)
-
-    Base.unsafe_convert(::Type{Ptr{T}}, x::JLDeviceArray{T}) where {T} =
-        convert(Ptr{T}, pointer(x.data)) + x.offset
-
-    # conversion of untyped data to a typed Array
-    function typed_data(x::JLDeviceArray{T}) where {T}
-        unsafe_wrap(Array, pointer(x), x.dims)
-    end
-
-    @inline Base.getindex(A::JLDeviceArray, index::Integer) = getindex(typed_data(A), index)
-    @inline Base.setindex!(A::JLDeviceArray, x, index::Integer) = setindex!(typed_data(A), x, index)
-end
diff --git a/lib/JLArrays/src/sparse.jl b/lib/JLArrays/src/sparse.jl
index 032bd422f..8adfea701 100644
--- a/lib/JLArrays/src/sparse.jl
+++ b/lib/JLArrays/src/sparse.jl
@@ -153,6 +153,11 @@ function JLSparseMatrixCSC(xs::SparseMatrixCSC{Tv, Ti}) where {Ti, Tv}
     return JLSparseMatrixCSC{Tv, Ti}(colPtr, rowVal, nzVal, (xs.m, xs.n))
 end
 JLSparseMatrixCSC(xs::SparseVector) = JLSparseMatrixCSC(SparseMatrixCSC(xs))
+
+# dense to sparse, on the host
+SparseArrays.sparse(x::JLVector) = JLSparseVector(sparse(Array(x)))
+SparseArrays.sparse(x::JLMatrix) = JLSparseMatrixCSC(sparse(Array(x)))
+
 Base.length(x::JLSparseMatrixCSC) = prod(x.dims)
 Base.size(x::JLSparseMatrixCSC) = x.dims
 
@@ -209,9 +214,15 @@ end
 
 ## adapt for the device
 
-Adapt.adapt_structure(to::Adaptor, x::JLSparseMatrixCSC{Tv,Ti}) where {Tv,Ti} =
-GPUSparseDeviceMatrixCSC{Tv,Ti,JLDeviceArray{Ti, 1}, JLDeviceArray{Tv, 1}, AS.Generic}(adapt(to, x.colPtr), adapt(to, x.rowVal), adapt(to, x.nzVal), x.dims, x.nnz)
-Adapt.adapt_structure(to::Adaptor, x::JLSparseMatrixCSR{Tv,Ti}) where {Tv,Ti} =
-GPUSparseDeviceMatrixCSR{Tv,Ti,JLDeviceArray{Ti, 1}, JLDeviceArray{Tv, 1}, AS.Generic}(adapt(to, x.rowPtr), adapt(to, x.colVal), adapt(to, x.nzVal), x.dims, x.nnz)
-Adapt.adapt_structure(to::Adaptor, x::JLSparseVector{Tv,Ti}) where {Tv,Ti} =
-GPUSparseDeviceVector{Tv,Ti,JLDeviceArray{Ti, 1}, JLDeviceArray{Tv, 1}, AS.Generic}(adapt(to, x.iPtr), adapt(to, x.nzVal), x.len, x.nnz)
+function Adapt.adapt_structure(to::Adaptor, x::JLSparseMatrixCSC{Tv,Ti}) where {Tv,Ti}
+    colPtr, rowVal, nzVal = adapt(to, x.colPtr), adapt(to, x.rowVal), adapt(to, x.nzVal)
+    GPUSparseDeviceMatrixCSC{Tv,Ti,typeof(colPtr),typeof(nzVal),AS.Generic}(colPtr, rowVal, nzVal, x.dims, x.nnz)
+end
+function Adapt.adapt_structure(to::Adaptor, x::JLSparseMatrixCSR{Tv,Ti}) where {Tv,Ti}
+    rowPtr, colVal, nzVal = adapt(to, x.rowPtr), adapt(to, x.colVal), adapt(to, x.nzVal)
+    GPUSparseDeviceMatrixCSR{Tv,Ti,typeof(rowPtr),typeof(nzVal),AS.Generic}(rowPtr, colVal, nzVal, x.dims, x.nnz)
+end
+function Adapt.adapt_structure(to::Adaptor, x::JLSparseVector{Tv,Ti}) where {Tv,Ti}
+    iPtr, nzVal = adapt(to, x.iPtr), adapt(to, x.nzVal)
+    GPUSparseDeviceVector{Tv,Ti,typeof(iPtr),typeof(nzVal),AS.Generic}(iPtr, nzVal, x.len, x.nnz)
+end
diff --git a/src/host/gemm.jl b/src/host/gemm.jl
index d018806d6..cd37f7576 100644
--- a/src/host/gemm.jl
+++ b/src/host/gemm.jl
@@ -1,5 +1,5 @@
 using LinearAlgebra: MulAddMul, wrapper_char, _unwrap
-
+import KernelInterface as KI
 ## Matmul kernel
 
 # A[i, j] of a symmetric/hermitian matrix with only one triangle stored
@@ -32,43 +32,37 @@ end
     end
 end
 
+# additive identity to start accumulating from: -0.0, not 0.0, for IEEE floats, so that a
+# sum of -0.0 products stays -0.0. other element types needn't support unary minus.
+@inline accumulator_zero(::Type{T}) where {T <: Number} = -zero(T)
+@inline accumulator_zero(::Type{T}) where {T} = zero(T)
+
 const TILE_DIM = 16
-@kernel unsafe_indices = true function coalesced_matmul_kernel!(
+function coalesced_matmul_kernel!(
         C, A, B, add,
         M, N, K, ::Val{TA}, ::Val{TB},
         ::Val{TILE},
     ) where {TA, TB, TILE}
-    @uniform TAT = eltype(A); @uniform TBT = eltype(B); @uniform R = eltype(C)
-
-    # li, lj, _ = KI.get_local_id()
-    li, lj = @index(Local, NTuple)
+    TAT = eltype(A); TBT = eltype(B); R = eltype(C)
 
-    # grow, gcol, _ = KI.get_group_id()
-    grow, gcol = @index(Group, NTuple)
-    # Uncomment once moved to KernelInterface
-    # gi = (grow - 1) * TILE + li
-    # gj = (gcol - 1) * TILE + lj
+    li, lj, _ = KI.get_local_id()
+    grow, gcol, _ = KI.get_group_id()
+    gi = (grow - 1) * TILE + li
+    gj = (gcol - 1) * TILE + lj
 
-    # padded by one row to avoid bank conflicts on shared memory
-    tile1 = @localmem TAT (TILE + 1, TILE)
-    tile2 = @localmem TBT (TILE + 1, TILE)
+    # Revisit potentially padding if value of `TILE_DIM` ever changed
+    tile1 = KI.localmemory(TAT, (TILE, TILE))
+    tile2 = KI.localmemory(TBT, (TILE, TILE))
 
     # private variable for tile C
-    @uniform Tacc = promote_type(R, typeof(zero(TAT) * zero(TBT) + zero(TAT) * zero(TBT)))
-    # acc = -zero(Tacc) #Use this when KernelInterface
-    acc = @private Tacc 1
-    @inbounds acc[1] = zero(Tacc)
+    Tacc = promote_type(R, typeof(zero(TAT) * zero(TBT) + zero(TAT) * zero(TBT)))
+    acc = accumulator_zero(Tacc)
 
     # number of tiles depends on inner dimension
-    # NUM_TILES = div(K + TILE - 1, TILE)
-    @uniform NUM_TILES = cld(K, TILE)
+    NUM_TILES = cld(K, TILE)
 
     # loop over all tiles needed for this calculation
     for t in 0:(NUM_TILES - 1)
-        # Remove once moved to KernelInterface
-        gi = (grow - 1) * TILE + li
-        gj = (gcol - 1) * TILE + lj
-
         # load inputs into tiles, with bounds checking for non-square matrices
         k0 = t * TILE
         ac = k0 + lj
@@ -85,28 +79,21 @@ const TILE_DIM = 16
         end
 
         # wait for all tiles to be loaded
-         @synchronize()
-
-        # Remove once moved to KernelInterface
-         gi = (grow - 1) * TILE + li
-         gj = (gcol - 1) * TILE + lj
+        KI.barrier()
 
-        @inbounds for k in 1:TILE
-            acc[1] = muladd(tile1[li, k], tile2[k, lj], acc[1])
+        # skip the zero padding past K: adding its +0.0 products would turn a -0.0 into +0.0
+        @inbounds for k in 1:min(TILE, K - k0)
+            acc = muladd(tile1[li, k], tile2[k, lj], acc)
         end
 
-         @synchronize()
+         KI.barrier()
     end
 
-    # Remove once moved to KernelInterface
-    gi = (grow - 1) * TILE + li
-    gj = (gcol - 1) * TILE + lj
-
     # save if inbounds
     if gi <= M && gj <= N
-        @inbounds C[gi, gj] = add(acc[1], C[gi, gj])
+        @inbounds C[gi, gj] = add(acc, C[gi, gj])
     end
-    # return
+    return
 end
 
 @inline function check_matmul_shapes(C, A, B, cA='N', cB='N')
@@ -144,9 +131,7 @@ function generic_matmatmul!(C::AbstractArray{R}, cA::AbstractChar, cB::AbstractC
     workgroupsize=(TILE_DIM, TILE_DIM)
     numworkgroups=(cld(M, TILE_DIM), cld(N, TILE_DIM))
 
-    # KI.@kernel KI.get_backend(C) workgroupsize=workgroupsize numworkgroups=numworkgroups coalesced_matmul_kernel!(
-            # C, A, B, add, Int(M), Int(N), Int(K), Val(cA), Val(cB), Val(TILE_DIM))
-    coalesced_matmul_kernel!(get_backend(C), workgroupsize)(
-            C, A, B, add, Int(M), Int(N), Int(K), Val(cA), Val(cB), Val(TILE_DIM); ndrange=numworkgroups .* workgroupsize)
+    KI.@launch KI.get_backend(C) workgroupsize=workgroupsize numgroups=numworkgroups coalesced_matmul_kernel!(
+            C, A, B, add, Int(M), Int(N), Int(K), Val(cA), Val(cB), Val(TILE_DIM))
     C
 end
diff --git a/src/host/linalg.jl b/src/host/linalg.jl
index 3c6788e23..c80c44693 100644
--- a/src/host/linalg.jl
+++ b/src/host/linalg.jl
@@ -1182,7 +1182,7 @@ end
 ## reflect
 
 function LinearAlgebra.reflect!(x::AbstractGPUArray, y::AbstractGPUArray, c::Number, s::Number)
-    @kernel function  reflect_kernel!(x, y, c, s)
+    @kernel function reflect_kernel!(x, y, c, s)
         i = @index(Global, Linear)
         @inbounds xi = x[i]
         @inbounds yi = y[i]
diff --git a/test/Project.toml b/test/Project.toml
index 7e816ed54..ab94d87b8 100644
--- a/test/Project.toml
+++ b/test/Project.toml
@@ -1,13 +1,17 @@
 [deps]
 Adapt = "79e6a3ab-5dfb-504d-930d-738a2a938a0e"
 GPUArrays = "0c68f7d7-f131-5f86-a1c3-88cf8149b2d7"
+InteractiveUtils = "b77e0a4c-d291-57a0-90e8-8db25a27a240"
 JLArrays = "27aeb0d3-9eb9-45fb-866b-73c2ecf80fcb"
 JLD2 = "033835bb-8acc-5ee8-8aae-3f567f8a3819"
 KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c"
+KernelInterface = "4ee993da-d684-4d17-a7dd-4e58e78d92bf"
 LinearAlgebra = "37e2e46d-f89d-539d-b4ee-838fcccc9c8e"
 ParallelTestRunner = "d3525ed8-44d0-4b2c-a655-542cee43accc"
 Random = "9a3f8284-a2c9-5f02-9a11-845980a1fd5c"
 SparseArrays = "2f01184e-e22b-5df5-ae63-d93ebab69eaf"
+SpecialFunctions = "276daf66-3868-5448-9aa4-cd146d93841b"
+StaticArrays = "90137ffa-7385-5640-81b9-e52037218182"
 Statistics = "10745b16-79ce-11e8-11f9-7d13ad32a3b2"
 Test = "8dfed614-e22c-5e08-85e1-65c5234f0b40"
 
diff --git a/test/runtests.jl b/test/runtests.jl
index c2b9ea471..0daf898ca 100644
--- a/test/runtests.jl
+++ b/test/runtests.jl
@@ -11,6 +11,12 @@ const init_worker_code = quote
 
     TestSuite.sparse_types(::Type{<:JLArray}) = (JLSparseVector, JLSparseMatrixCSC, JLSparseMatrixCSR)
     TestSuite.sparse_types(::Type{<:Array})   = (SparseVector, SparseMatrixCSC)
+
+    # KA v0.10 runs JLArray kernels through POCL, whose SPIR-V back-end lacks 128-bit integers
+    if isdefined(JLArrays.KernelAbstractions, :POCL)
+        TestSuite.supported_eltypes(::Type{<:JLArray}) =
+            filter(!=(Int128), TestSuite.supported_eltypes())
+    end
 end
 
 const init_code = quote
@@ -26,6 +32,27 @@ for AT in (:JLArray, :Array), name in keys(TestSuite.tests)
     testsuite["$(AT)/$name"] = :(TestSuite.tests[$name]($AT))
 end
 
+testsuite["JLArray/kernelabstractions"] = quote
+    import KernelAbstractions
+    import KernelAbstractions as KA
+    using JLArrays.JLKernels
+
+    include(joinpath(dirname(pathof(KernelAbstractions)), "..", "test", "testsuite.jl"))
+
+    Testsuite.testsuite(JLBackend, "JLArrays", JLArrays, JLArray, KernelAbstractions.POCL.CLDeviceArray)
+end
+
+testsuite["JLArray/kernelinterface"] = quote
+    import KernelInterface
+    import KernelInterface as KI
+    using JLArrays.JLKernels
+    import Adapt
+
+    include(joinpath(dirname(pathof(KernelInterface)), "..", "test", "testsuite.jl"))
+
+    Testsuite.testsuite(JLBackend(), JLArray)
+end
+
 # captured variables that get boxed (`Core.Box`) lose their type
 if isdefined(Test, :detect_closure_boxes)
     testsuite["closure boxes"] = quote