From efb42e932886fa649d3d7336b1dab0ec1c1a7f6b Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Mon, 22 Dec 2025 16:04:53 -0400
Subject: [PATCH 01/17] Initial subgroups support

---
 src/interface.jl  | 111 ++++++++++++++++++++++++++++++++++++++++++++++
 test/interface.jl |  55 +++++++++++++++++++++++
 2 files changed, 166 insertions(+)

diff --git a/src/interface.jl b/src/interface.jl
index f2c21d2cc..7853ad1fc 100644
--- a/src/interface.jl
+++ b/src/interface.jl
@@ -102,6 +102,78 @@ Returns the unique group ID.
 """
 function get_group_id end
 
+"""
+    get_sub_group_size()::UInt32
+
+Returns the number of work-items in the sub-group.
+
+!!! note
+    Backend implementations **must** implement:
+    ```
+    @device_override get_sub_group_size()::UInt32
+    ```
+"""
+function get_sub_group_size end
+
+"""
+    get_max_sub_group_size()::UInt32
+
+Returns the maximum sub-group size for sub-groups in the current workgroup.
+
+!!! note
+    Backend implementations **must** implement:
+    ```
+    @device_override get_max_sub_group_size()::UInt32
+    ```
+"""
+function get_max_sub_group_size end
+
+"""
+    get_num_sub_groups()::UInt32
+
+Returns the number of sub-groups in the current workgroup.
+
+!!! note
+    Backend implementations **must** implement:
+    ```
+    @device_override get_num_sub_groups()::UInt32
+    ```
+"""
+function get_num_sub_groups end
+
+"""
+    get_sub_group_id()::UInt32
+
+Returns the sub-group ID within the work-group.
+
+!!! note
+    1-based.
+
+!!! note
+    Backend implementations **must** implement:
+    ```
+    @device_override get_sub_group_id()::UInt32
+    ```
+"""
+function get_sub_group_id end
+
+"""
+    get_sub_group_local_id()::UInt32
+
+Returns the work-item ID within the current sub-group.
+
+!!! note
+    1-based.
+
+!!! note
+    Backend implementations **must** implement:
+    ```
+    @device_override get_sub_group_local_id()::UInt32
+    ```
+"""
+function get_sub_group_local_id end
+
+
 """
     localmemory(::Type{T}, dims)
 
@@ -139,6 +211,29 @@ function barrier()
     error("Group barrier used outside kernel or not captured")
 end
 
+"""
+    sub_group_barrier()
+
+After a `sub_group_barrier()` call, all read and writes to global and local memory
+from each thread in the sub-group are visible in from all other threads in the
+sub-group.
+
+This does **not** guarantee that a write from a thread in a certain sub-group will
+be visible to a thread in a different sub-group.
+
+!!! note
+    `sub_group_barrier()` must be encountered by all workitems of a sub-group executing the kernel or by none at all.
+
+!!! note
+    Backend implementations **must** implement:
+    ```
+    @device_override sub_group_barrier()
+    ```
+"""
+function sub_group_barrier()
+    error("Sub-group barrier used outside kernel or not captured")
+end
+
 """
     _print(args...)
 
@@ -220,6 +315,22 @@ kernel launch with too big a workgroup is attempted.
 """
 function max_work_group_size end
 
+"""
+    sub_group_size(backend)::Int
+
+Returns a reasonable sub-group size supported by the currently
+active device for the specified backend. This would typically
+be 32, or 64 for devices that don't support 32.
+
+!!! note
+    Backend implementations **must** implement:
+    ```
+    sub_group_size(backend::NewBackend)::Int
+    ```
+    As well as the on-device functionality.
+"""
+function sub_group_size end
+
 """
     multiprocessor_count(backend::NewBackend)::Int
 
diff --git a/test/interface.jl b/test/interface.jl
index a6fdd05bc..cb8c59f71 100644
--- a/test/interface.jl
+++ b/test/interface.jl
@@ -23,6 +23,27 @@ function test_interface_kernel(results)
     end
     return
 end
+struct SubgroupData
+    sub_group_size::UInt32
+    max_sub_group_size::UInt32
+    num_sub_groups::UInt32
+    sub_group_id::UInt32
+    sub_group_local_id::UInt32
+end
+function test_subgroup_kernel(results)
+    i = KI.get_global_id().x
+
+    if i <= length(results)
+        @inbounds results[i] = SubgroupData(
+            KI.get_sub_group_size(),
+            KI.get_max_sub_group_size(),
+            KI.get_num_sub_groups(),
+            KI.get_sub_group_id(),
+            KI.get_sub_group_local_id()
+        )
+    end
+    return
+end
 
 function interface_testsuite(backend, AT)
     @testset "KernelInterface Tests" begin
@@ -122,6 +143,40 @@ function interface_testsuite(backend, AT)
                 @test k_data.local_id == expected_local
             end
         end
+
+        @testset "Sub-groups" begin
+            @test KI.sub_group_size(backend()) isa Int
+
+            # Test with small kernel
+            sg_size = KI.sub_group_size(backend())
+            sg_n = 2
+            workgroupsize = sg_size * sg_n
+            numworkgroups = 2
+            N = workgroupsize * numworkgroups
+
+            results = AT(Vector{SubgroupData}(undef, N))
+            kernel = KI.@kernel backend() launch = false test_subgroup_kernel(results)
+
+            kernel(results; workgroupsize, numworkgroups)
+            KernelAbstractions.synchronize(backend())
+
+            host_results = Array(results)
+
+            # Verify results make sense
+            for (i, sg_data) in enumerate(host_results)
+                @test sg_data.sub_group_size == sg_size
+                @test sg_data.max_sub_group_size == sg_size
+                @test sg_data.num_sub_groups == sg_n
+
+                # Group ID should be 1-based
+                expected_sub_group = div(((i - 1) % workgroupsize), sg_size) + 1
+                @test sg_data.sub_group_id == expected_sub_group
+
+                # Local ID should be 1-based within group
+                expected_sg_local = ((i - 1) % sg_size) + 1
+                @test sg_data.sub_group_local_id == expected_sg_local
+            end
+        end
     end
     return nothing
 end

From 0d8d42954332404a31e0680a631e0defb36bd0b4 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Fri, 21 Nov 2025 21:32:27 -0400
Subject: [PATCH 02/17] `shfl_down` intrinsics

Co-Authored-By: Anton Smirnov <tonysmn97@gmail.com>
---
 src/interface.jl  | 26 ++++++++++++++++++++++
 test/interface.jl | 55 +++++++++++++++++++++++++++++++++++++++++++++++
 2 files changed, 81 insertions(+)

diff --git a/src/interface.jl b/src/interface.jl
index 7853ad1fc..638ecb37e 100644
--- a/src/interface.jl
+++ b/src/interface.jl
@@ -188,6 +188,32 @@ Declare memory that is local to a workgroup.
 """
 localmemory(::Type{T}, dims) where {T} = localmemory(T, Val(dims))
 
+"""
+    shfl_down(val::T, offset::Integer) where T
+
+Read `val` from a lane with higher id given by `offset`.
+
+!!! note
+    Backend implementations **must** implement:
+    ```
+    @device_override shfl_down(val::T, offset::Integer) where T
+    ```
+    As well as the on-device functionality.
+"""
+function shfl_down end
+
+"""
+    shfl_down_types(::Backend)::Vector{DataType}
+
+Returns a vector of `DataType`s supported on `backend`
+
+!!! note
+    Backend implementations **must** implement this function
+    only if they support `shfl_down` for any types.
+"""
+shfl_down_types(::Backend) = DataType[]
+
+
 """
     barrier()
 
diff --git a/test/interface.jl b/test/interface.jl
index cb8c59f71..13037b16a 100644
--- a/test/interface.jl
+++ b/test/interface.jl
@@ -45,6 +45,44 @@ function test_subgroup_kernel(results)
     return
 end
 
+function shfl_down_test_kernel(a, b, ::Val{N}) where {N}
+    idx = KI.get_sub_group_local_id()
+
+    val = a[idx]
+
+    offset = 0x00000001
+    while offset < N
+        val += KI.shfl_down(val, offset)
+        offset <<= 1
+    end
+
+    KI.sub_group_barrier()
+
+    if idx == 1
+        b[idx] = val
+    end
+    return
+end
+
+function shfl_down_test_kernel(a, b, ::Val{N}) where {N}
+    idx = KI.get_sub_group_local_id()
+
+    val = a[idx]
+
+    offset = 0x00000001
+    while offset < N
+        val += KI.shfl_down(val, offset)
+        offset <<= 1
+    end
+
+    KI.sub_group_barrier()
+
+    if idx == 1
+        b[idx] = val
+    end
+    return
+end
+
 function interface_testsuite(backend, AT)
     @testset "KernelInterface Tests" begin
         @testset "Launch parameters" begin
@@ -177,6 +215,23 @@ function interface_testsuite(backend, AT)
                 @test sg_data.sub_group_local_id == expected_sg_local
             end
         end
+        @testset "shfl_down" begin
+            @test !isempty(KI.shfl_down_types(backend()))
+            types_to_test = setdiff(KI.shfl_down_types(backend()), [Bool])
+            @testset "$T" for T in types_to_test
+                N = KI.sub_group_size(backend())
+                a = zeros(T, N)
+                rand!(a, (0:1))
+
+                dev_a = AT(a)
+                dev_b = AT(zeros(T, N))
+
+                KI.@kernel backend() workgroupsize = N shfl_down_test_kernel(dev_a, dev_b, Val(N))
+
+                b = Array(dev_b)
+                @test sum(a) ≈ b[1]
+            end
+        end
     end
     return nothing
 end

From dda81e01e94382921059f875eed39a94cb74cf28 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Fri, 2 Jan 2026 00:20:06 -0400
Subject: [PATCH 03/17] All backends are synchronizing

---
 src/interface.jl | 8 ++++++++
 1 file changed, 8 insertions(+)

diff --git a/src/interface.jl b/src/interface.jl
index 638ecb37e..6748c227a 100644
--- a/src/interface.jl
+++ b/src/interface.jl
@@ -193,12 +193,20 @@ localmemory(::Type{T}, dims) where {T} = localmemory(T, Val(dims))
 
 Read `val` from a lane with higher id given by `offset`.
 
+!!! note
+    `shfl_down` must be encountered by all workitems of a sub-group executing the kernel or by none at all.
+
 !!! note
     Backend implementations **must** implement:
     ```
     @device_override shfl_down(val::T, offset::Integer) where T
     ```
     As well as the on-device functionality.
+
+    This implementation **must** be synchronizing.
+    That is, kernels using this function can safely assume that
+    they do **not** need a `sub_group_barrier` before calling
+    this function.
 """
 function shfl_down end
 

From a39f7f3feeac4f5248e3f0c8acaa5a02cf4db4a4 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Sat, 3 Jan 2026 14:07:42 -0400
Subject: [PATCH 04/17] Local backend support

---
 src/pocl/backend.jl              | 46 ++++++++++++++++++++++++++++++++
 src/pocl/compiler/compilation.jl | 14 +++++++---
 src/pocl/compiler/execution.jl   |  2 +-
 src/pocl/nanoOpenCL.jl           | 10 +++++++
 src/pocl/pocl.jl                 |  2 +-
 5 files changed, 69 insertions(+), 5 deletions(-)

diff --git a/src/pocl/backend.jl b/src/pocl/backend.jl
index 3b0fdb3fa..d8bad41d2 100644
--- a/src/pocl/backend.jl
+++ b/src/pocl/backend.jl
@@ -9,6 +9,8 @@ using SPIRV_LLVM_Backend_jll, SPIRV_Tools_jll
 import KernelAbstractions as KA
 import KernelAbstractions.KernelInterface as KI
 
+import SPIRVIntrinsics
+
 import StaticArrays
 
 import Adapt
@@ -227,10 +229,36 @@ end
 function KI.max_work_group_size(::POCLBackend)::Int
     return Int(device().max_work_group_size)
 end
+function KI.sub_group_size(::POCLBackend)::Int
+    sg_sizes = cl.device().sub_group_sizes
+    if 32 in sg_sizes
+        return 32
+    elseif 64 in sg_sizes
+        return 64
+    elseif 16 in sg_sizes
+        return 16
+    else
+        return 1
+    end
+end
 function KI.multiprocessor_count(::POCLBackend)::Int
     return Int(device().max_compute_units)
 end
 
+function KI.shfl_down_types(::POCLBackend)
+    res = copy(SPIRVIntrinsics.gentypes)
+
+    backend_extensions = cl.device().extensions
+    if "cl_khr_fp64" ∉ backend_extensions
+        res = setdiff(res, [Float64])
+    end
+    if "cl_khr_fp16" ∉ backend_extensions
+        res = setdiff(res, [Float16])
+    end
+
+    return res
+end
+
 ## Indexing Functions
 
 @device_override @inline function KI.get_local_id()
@@ -257,6 +285,16 @@ end
     return (; x = Int(get_global_size(1)), y = Int(get_global_size(2)), z = Int(get_global_size(3)))
 end
 
+@device_override KI.get_sub_group_size() = get_sub_group_size()
+
+@device_override KI.get_max_sub_group_size() = get_max_sub_group_size()
+
+@device_override KI.get_num_sub_groups() = get_num_sub_groups()
+
+@device_override KI.get_sub_group_id() = get_sub_group_id()
+
+@device_override KI.get_sub_group_local_id() = get_sub_group_local_id()
+
 @device_override @inline function KA.__validindex(ctx)
     if KA.__dynamic_checkbounds(ctx)
         I = @inbounds KA.expand(KA.__iterspace(ctx), get_group_id(1), get_local_id(1))
@@ -285,6 +323,14 @@ end
     work_group_barrier(POCL.LOCAL_MEM_FENCE | POCL.GLOBAL_MEM_FENCE)
 end
 
+@device_override @inline function KI.sub_group_barrier()
+    sub_group_barrier(POCL.LOCAL_MEM_FENCE | POCL.GLOBAL_MEM_FENCE)
+end
+
+@device_override function KI.shfl_down(val::T, offset::Integer) where {T}
+    sub_group_shuffle(val, get_sub_group_local_id() + offset)
+end
+
 @device_override @inline function KI._print(args...)
     POCL._print(args...)
 end
diff --git a/src/pocl/compiler/compilation.jl b/src/pocl/compiler/compilation.jl
index 8aee85447..f951de556 100644
--- a/src/pocl/compiler/compilation.jl
+++ b/src/pocl/compiler/compilation.jl
@@ -1,6 +1,9 @@
 ## gpucompiler interface
 
-struct OpenCLCompilerParams <: AbstractCompilerParams end
+Base.@kwdef struct OpenCLCompilerParams <: AbstractCompilerParams
+    sub_group_size::Int
+end
+
 const OpenCLCompilerConfig = CompilerConfig{SPIRVCompilerTarget, OpenCLCompilerParams}
 const OpenCLCompilerJob = CompilerJob{SPIRVCompilerTarget, OpenCLCompilerParams}
 
@@ -31,6 +34,8 @@ function GPUCompiler.finish_module!(
         job, mod, entry
     )
 
+    metadata(entry)["intel_reqd_sub_group_size"] = MDNode([ConstantInt(Int32(job.config.params.sub_group_size))])
+
     # if this kernel uses our RNG, we should prime the shared state.
     # XXX: these transformations should really happen at the Julia IR level...
     if haskey(functions(mod), "julia.opencl.random_keys") && job.config.kernel
@@ -130,14 +135,17 @@ function compiler_config(dev::cl.Device; kwargs...)
     end
     return config
 end
-@noinline function _compiler_config(dev; kernel = true, name = nothing, always_inline = false, kwargs...)
+@noinline function _compiler_config(dev; kernel = true, name = nothing, always_inline = false, sub_group_size = 32, kwargs...)
     supports_fp16 = "cl_khr_fp16" in dev.extensions
     supports_fp64 = "cl_khr_fp64" in dev.extensions
 
+    if sub_group_size ∉ dev.sub_group_sizes
+        @error("$sub_group_size is not a valid sub-group size for this device.")
+    end
 
     # create GPUCompiler objects
     target = SPIRVCompilerTarget(; supports_fp16, supports_fp64, validate = true, kwargs...)
-    params = OpenCLCompilerParams()
+    params = OpenCLCompilerParams(; sub_group_size)
     return CompilerConfig(target, params; kernel, name, always_inline)
 end
 
diff --git a/src/pocl/compiler/execution.jl b/src/pocl/compiler/execution.jl
index 595e7ff1a..fd22b57c8 100644
--- a/src/pocl/compiler/execution.jl
+++ b/src/pocl/compiler/execution.jl
@@ -4,7 +4,7 @@ export @opencl, clfunction, clconvert
 ## high-level @opencl interface
 
 const MACRO_KWARGS = [:launch]
-const COMPILER_KWARGS = [:kernel, :name, :always_inline, :validate]
+const COMPILER_KWARGS = [:kernel, :name, :always_inline, :validate, :sub_group_size]
 const LAUNCH_KWARGS = [:global_size, :local_size, :queue]
 
 macro opencl(ex...)
diff --git a/src/pocl/nanoOpenCL.jl b/src/pocl/nanoOpenCL.jl
index e7901ad07..eec955990 100644
--- a/src/pocl/nanoOpenCL.jl
+++ b/src/pocl/nanoOpenCL.jl
@@ -389,6 +389,8 @@ const CL_KERNEL_EXEC_INFO_SVM_PTRS = 0x11b6
 
 const CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM = 0x11b7
 
+const CL_DEVICE_SUB_GROUP_SIZES_INTEL = 0x4108
+
 struct CLError <: Exception
     code::Cint
 end
@@ -951,6 +953,14 @@ devices(p::Platform) = devices(p, CL_DEVICE_TYPE_ALL)
         return tuple([Int(r) for r in result]...)
     end
 
+    if s == :sub_group_sizes
+        res_size = Ref{Csize_t}()
+        clGetDeviceInfo(d, CL_DEVICE_SUB_GROUP_SIZES_INTEL, C_NULL, C_NULL, res_size)
+        result = Vector{Csize_t}(undef, res_size[] ÷ sizeof(Csize_t))
+        clGetDeviceInfo(d, CL_DEVICE_SUB_GROUP_SIZES_INTEL, sizeof(result), result, C_NULL)
+        return tuple([Int(r) for r in result]...)
+    end
+
     if s == :max_image2d_shape
         width = Ref{Csize_t}()
         height = Ref{Csize_t}()
diff --git a/src/pocl/pocl.jl b/src/pocl/pocl.jl
index de1bd4b32..007a491cc 100644
--- a/src/pocl/pocl.jl
+++ b/src/pocl/pocl.jl
@@ -41,7 +41,7 @@ function queue()
 end
 
 using GPUCompiler
-using LLVM, LLVM.Interop
+import LLVM: LLVM, MDNode, ConstantInt, metadata, Interop
 using SPIRV_LLVM_Backend_jll, SPIRV_Tools_jll
 using Adapt
 

From aa9f6ee7a0ceda6b16f6c1f3cc19e942440af3af Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Mon, 22 Dec 2025 21:43:21 -0400
Subject: [PATCH 05/17] [Temp] Use new intrinsics feature branches in CI

---
 .buildkite/pipeline.yml | 10 +++++-----
 1 file changed, 5 insertions(+), 5 deletions(-)

diff --git a/.buildkite/pipeline.yml b/.buildkite/pipeline.yml
index 6e229c6ea..04766f59a 100644
--- a/.buildkite/pipeline.yml
+++ b/.buildkite/pipeline.yml
@@ -16,7 +16,7 @@ steps:
       julia -e 'println("--- :julia: Developing CUDA")
                 using Pkg
                 url="https://github.com/christiangnrd/CUDA.jl"
-                rev="intrinsics"
+                rev="intrinsicsnew"
                 subdirs = ["CUDACore", "CUDATools", "lib/cusolver", "lib/curand", "lib/cublas", "lib/cufft", "lib/cusparse", "lib/cupti", "lib/nvml", "lib/custatevec", "lib/cudnn", "lib/cutensor", "lib/cutensornet"]
                 Pkg.add([(;url, rev); [(; url, rev, subdir) for subdir in subdirs]])'
       julia -e 'println("--- :julia: Instantiating project")
@@ -81,7 +81,7 @@ steps:
     command: |
       julia -e 'println("--- :julia: Developing Metal")
                 using Pkg
-                Pkg.add(url="https://github.com/JuliaGPU/Metal.jl", rev="kaintr")'
+                Pkg.add(url="https://github.com/JuliaGPU/Metal.jl", rev="kaintrnew")'
       julia -e 'println("--- :julia: Instantiating project")
                 using Pkg
                 Pkg.develop(; path=pwd())' || exit 3
@@ -113,7 +113,7 @@ steps:
     command: |
       julia -e 'println("--- :julia: Developing oneAPI")
                 using Pkg
-                Pkg.add(url="https://github.com/christiangnrd/oneAPI.jl", rev="intrinsics")
+                Pkg.add(url="https://github.com/christiangnrd/oneAPI.jl", rev="intrinsicsnew")
                 Pkg.add(url="https://github.com/christiangnrd/AcceleratedKernels.jl", rev="ka0.10simple")'
       julia -e 'println("--- :julia: Instantiating project")
                 using Pkg
@@ -149,7 +149,7 @@ steps:
                 Pkg.add(url="https://github.com/christiangnrd/AcceleratedKernels.jl", rev="ka0.10simple")'
       julia -e '
                 using Pkg
-                Pkg.add(url="https://github.com/christiangnrd/AMDGPU.jl", rev="intrinsics")
+                Pkg.add(url="https://github.com/christiangnrd/AMDGPU.jl", rev="intrinsicsnew")
                 println("--- :julia: Instantiating project")
                 Pkg.develop(; path=pwd())' || exit 3
 
@@ -181,7 +181,7 @@ steps:
     command: |
       julia -e 'println("--- :julia: Developing OpenCL")
                 using Pkg
-                Pkg.add(url="https://github.com/christiangnrd/OpenCL.jl", rev="intrinsics")
+                Pkg.add(url="https://github.com/christiangnrd/OpenCL.jl", rev="intrinsicsnew")
                 Pkg.develop(; name="SPIRVIntrinsics")'
       julia -e 'println("--- :julia: Instantiating project")
                 using Pkg

From fafb1164060edb833733ceed92e16db74588dfc6 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Thu, 28 May 2026 19:06:37 -0300
Subject: [PATCH 06/17] Fixes to pocl backend in line with OpenCL fixes

---
 src/pocl/compiler/compilation.jl | 14 +++++++++-----
 1 file changed, 9 insertions(+), 5 deletions(-)

diff --git a/src/pocl/compiler/compilation.jl b/src/pocl/compiler/compilation.jl
index f951de556..ded8367e3 100644
--- a/src/pocl/compiler/compilation.jl
+++ b/src/pocl/compiler/compilation.jl
@@ -1,7 +1,8 @@
 ## gpucompiler interface
 
 Base.@kwdef struct OpenCLCompilerParams <: AbstractCompilerParams
-    sub_group_size::Int
+    # request a fixed sub-group width via `intel_reqd_sub_group_size`
+    sub_group_size::Union{Nothing, Int} = nothing
 end
 
 const OpenCLCompilerConfig = CompilerConfig{SPIRVCompilerTarget, OpenCLCompilerParams}
@@ -34,7 +35,10 @@ function GPUCompiler.finish_module!(
         job, mod, entry
     )
 
-    metadata(entry)["intel_reqd_sub_group_size"] = MDNode([ConstantInt(Int32(job.config.params.sub_group_size))])
+    sg_size = job.config.params.sub_group_size
+    if sg_size !== nothing
+        metadata(entry)["intel_reqd_sub_group_size"] = MDNode([ConstantInt(Int32(sg_size))])
+    end
 
     # if this kernel uses our RNG, we should prime the shared state.
     # XXX: these transformations should really happen at the Julia IR level...
@@ -135,12 +139,12 @@ function compiler_config(dev::cl.Device; kwargs...)
     end
     return config
 end
-@noinline function _compiler_config(dev; kernel = true, name = nothing, always_inline = false, sub_group_size = 32, kwargs...)
+@noinline function _compiler_config(dev; kernel = true, name = nothing, always_inline = false, sub_group_size::Union{Nothing, Int} = nothing, kwargs...)
     supports_fp16 = "cl_khr_fp16" in dev.extensions
     supports_fp64 = "cl_khr_fp64" in dev.extensions
 
-    if sub_group_size ∉ dev.sub_group_sizes
-        @error("$sub_group_size is not a valid sub-group size for this device.")
+    if sub_group_size !== nothing && sub_group_size ∉ dev.sub_group_sizes
+        error("$sub_group_size is not a valid sub-group size for this device.")
     end
 
     # create GPUCompiler objects

From 7256803d23e59b180569b072074237b67b2dce24 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Fri, 29 May 2026 17:00:10 -0300
Subject: [PATCH 07/17] Temp?

---
 src/pocl/compiler/compilation.jl | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/src/pocl/compiler/compilation.jl b/src/pocl/compiler/compilation.jl
index ded8367e3..e5eefd31c 100644
--- a/src/pocl/compiler/compilation.jl
+++ b/src/pocl/compiler/compilation.jl
@@ -139,7 +139,7 @@ function compiler_config(dev::cl.Device; kwargs...)
     end
     return config
 end
-@noinline function _compiler_config(dev; kernel = true, name = nothing, always_inline = false, sub_group_size::Union{Nothing, Int} = nothing, kwargs...)
+@noinline function _compiler_config(dev; kernel = true, name = nothing, always_inline = false, sub_group_size::Union{Nothing, Int} = 32, kwargs...)
     supports_fp16 = "cl_khr_fp16" in dev.extensions
     supports_fp64 = "cl_khr_fp64" in dev.extensions
 

From 14e3cc770e5cc06ca6a6db99cb05fa70c832666a Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Mon, 15 Jun 2026 12:56:35 -0300
Subject: [PATCH 08/17] Fix

---
 src/pocl/pocl.jl | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

diff --git a/src/pocl/pocl.jl b/src/pocl/pocl.jl
index 007a491cc..681683e70 100644
--- a/src/pocl/pocl.jl
+++ b/src/pocl/pocl.jl
@@ -41,7 +41,8 @@ function queue()
 end
 
 using GPUCompiler
-import LLVM: LLVM, MDNode, ConstantInt, metadata, Interop
+using LLVM, LLVM.Interop
+import LLVM: LLVM, MDNode, ConstantInt, metadata
 using SPIRV_LLVM_Backend_jll, SPIRV_Tools_jll
 using Adapt
 

From bf926ff85c4ff739376367f53db3dafc60e8d214 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Thu, 2 Jul 2026 10:14:55 -0300
Subject: [PATCH 09/17] Correct branch

---
 .github/workflows/ci.yml | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml
index 64c294c97..229d78001 100644
--- a/.github/workflows/ci.yml
+++ b/.github/workflows/ci.yml
@@ -182,7 +182,7 @@ jobs:
         run: |
           julia -e 'println("--- :julia: Developing OpenCL")
                 using Pkg
-                Pkg.add(url="https://github.com/christiangnrd/OpenCL.jl", rev="intrinsics")
+                Pkg.add(url="https://github.com/christiangnrd/OpenCL.jl", rev="intrinsicsnew")
                 Pkg.develop(; name="SPIRVIntrinsics")'
       - name: "Instantiating project"
         run: |

From 018e68be9f914b0c45a670a6db3f13bdc2611b58 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Thu, 2 Jul 2026 10:57:17 -0300
Subject: [PATCH 10/17] Add `supports_subgroups`

---
 src/KernelAbstractions.jl | 11 +++++++++++
 test/interface.jl         |  4 +++-
 2 files changed, 14 insertions(+), 1 deletion(-)

diff --git a/src/KernelAbstractions.jl b/src/KernelAbstractions.jl
index 1b6962cbd..92c72a91d 100644
--- a/src/KernelAbstractions.jl
+++ b/src/KernelAbstractions.jl
@@ -647,6 +647,17 @@ function ones(backend::Backend, ::Type{T}, dims::Tuple; kwargs...) where {T}
     return data
 end
 
+"""
+    supports_subgroups(::Backend)::Bool
+
+Returns whether subgroups are supported by the backend.
+
+!!! note
+    Backend implementations **should** implement this function
+    only if they **do not** support subgroups.
+"""
+supports_subgroups(::Backend) = true
+
 """
     supports_unified(::Backend)::Bool
 
diff --git a/test/interface.jl b/test/interface.jl
index 13037b16a..2dc60802d 100644
--- a/test/interface.jl
+++ b/test/interface.jl
@@ -182,6 +182,7 @@ function interface_testsuite(backend, AT)
             end
         end
 
+    if !KI.supports_subgroups(backend())
         @testset "Sub-groups" begin
             @test KI.sub_group_size(backend()) isa Int
 
@@ -232,6 +233,7 @@ function interface_testsuite(backend, AT)
                 @test sum(a) ≈ b[1]
             end
         end
-    end
+    end # if !KI.supports_subgroups(backend())
+    end # @testset "KernelInterface Tests" begin
     return nothing
 end

From a23b891127c56a34366bd88bfb6358e0d96a39e0 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Thu, 2 Jul 2026 11:13:03 -0300
Subject: [PATCH 11/17] Fix

---
 test/interface.jl | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/test/interface.jl b/test/interface.jl
index 2dc60802d..5da304164 100644
--- a/test/interface.jl
+++ b/test/interface.jl
@@ -182,7 +182,7 @@ function interface_testsuite(backend, AT)
             end
         end
 
-    if !KI.supports_subgroups(backend())
+    if KernelAbstractions.supports_subgroups(backend())
         @testset "Sub-groups" begin
             @test KI.sub_group_size(backend()) isa Int
 
@@ -233,7 +233,7 @@ function interface_testsuite(backend, AT)
                 @test sum(a) ≈ b[1]
             end
         end
-    end # if !KI.supports_subgroups(backend())
+    end # if !KernelAbstractions.supports_subgroups(backend())
     end # @testset "KernelInterface Tests" begin
     return nothing
 end

From 45f4a819a9b2128d64fe01adabc48cb0b469a516 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Thu, 2 Jul 2026 11:49:35 -0300
Subject: [PATCH 12/17] Format

---
 test/interface.jl | 100 +++++++++++++++++++++++-----------------------
 1 file changed, 50 insertions(+), 50 deletions(-)

diff --git a/test/interface.jl b/test/interface.jl
index 5da304164..39e07cab2 100644
--- a/test/interface.jl
+++ b/test/interface.jl
@@ -182,58 +182,58 @@ function interface_testsuite(backend, AT)
             end
         end
 
-    if KernelAbstractions.supports_subgroups(backend())
-        @testset "Sub-groups" begin
-            @test KI.sub_group_size(backend()) isa Int
-
-            # Test with small kernel
-            sg_size = KI.sub_group_size(backend())
-            sg_n = 2
-            workgroupsize = sg_size * sg_n
-            numworkgroups = 2
-            N = workgroupsize * numworkgroups
-
-            results = AT(Vector{SubgroupData}(undef, N))
-            kernel = KI.@kernel backend() launch = false test_subgroup_kernel(results)
-
-            kernel(results; workgroupsize, numworkgroups)
-            KernelAbstractions.synchronize(backend())
-
-            host_results = Array(results)
-
-            # Verify results make sense
-            for (i, sg_data) in enumerate(host_results)
-                @test sg_data.sub_group_size == sg_size
-                @test sg_data.max_sub_group_size == sg_size
-                @test sg_data.num_sub_groups == sg_n
-
-                # Group ID should be 1-based
-                expected_sub_group = div(((i - 1) % workgroupsize), sg_size) + 1
-                @test sg_data.sub_group_id == expected_sub_group
-
-                # Local ID should be 1-based within group
-                expected_sg_local = ((i - 1) % sg_size) + 1
-                @test sg_data.sub_group_local_id == expected_sg_local
+        if KernelAbstractions.supports_subgroups(backend())
+            @testset "Sub-groups" begin
+                @test KI.sub_group_size(backend()) isa Int
+
+                # Test with small kernel
+                sg_size = KI.sub_group_size(backend())
+                sg_n = 2
+                workgroupsize = sg_size * sg_n
+                numworkgroups = 2
+                N = workgroupsize * numworkgroups
+
+                results = AT(Vector{SubgroupData}(undef, N))
+                kernel = KI.@kernel backend() launch = false test_subgroup_kernel(results)
+
+                kernel(results; workgroupsize, numworkgroups)
+                KernelAbstractions.synchronize(backend())
+
+                host_results = Array(results)
+
+                # Verify results make sense
+                for (i, sg_data) in enumerate(host_results)
+                    @test sg_data.sub_group_size == sg_size
+                    @test sg_data.max_sub_group_size == sg_size
+                    @test sg_data.num_sub_groups == sg_n
+
+                    # Group ID should be 1-based
+                    expected_sub_group = div(((i - 1) % workgroupsize), sg_size) + 1
+                    @test sg_data.sub_group_id == expected_sub_group
+
+                    # Local ID should be 1-based within group
+                    expected_sg_local = ((i - 1) % sg_size) + 1
+                    @test sg_data.sub_group_local_id == expected_sg_local
+                end
             end
-        end
-        @testset "shfl_down" begin
-            @test !isempty(KI.shfl_down_types(backend()))
-            types_to_test = setdiff(KI.shfl_down_types(backend()), [Bool])
-            @testset "$T" for T in types_to_test
-                N = KI.sub_group_size(backend())
-                a = zeros(T, N)
-                rand!(a, (0:1))
-
-                dev_a = AT(a)
-                dev_b = AT(zeros(T, N))
-
-                KI.@kernel backend() workgroupsize = N shfl_down_test_kernel(dev_a, dev_b, Val(N))
-
-                b = Array(dev_b)
-                @test sum(a) ≈ b[1]
+            @testset "shfl_down" begin
+                @test !isempty(KI.shfl_down_types(backend()))
+                types_to_test = setdiff(KI.shfl_down_types(backend()), [Bool])
+                @testset "$T" for T in types_to_test
+                    N = KI.sub_group_size(backend())
+                    a = zeros(T, N)
+                    rand!(a, (0:1))
+
+                    dev_a = AT(a)
+                    dev_b = AT(zeros(T, N))
+
+                    KI.@kernel backend() workgroupsize = N shfl_down_test_kernel(dev_a, dev_b, Val(N))
+
+                    b = Array(dev_b)
+                    @test sum(a) ≈ b[1]
+                end
             end
-        end
-    end # if !KernelAbstractions.supports_subgroups(backend())
+        end # if !KernelAbstractions.supports_subgroups(backend())
     end # @testset "KernelInterface Tests" begin
     return nothing
 end

From d2368c06c1976dcaa075f3e94c19ea61e8cd8ae3 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Sat, 4 Jul 2026 13:47:10 -0300
Subject: [PATCH 13/17] Add supports_subgroups to docs

---
 docs/src/api.md | 1 +
 1 file changed, 1 insertion(+)

diff --git a/docs/src/api.md b/docs/src/api.md
index 90e8af186..2881e6bf3 100644
--- a/docs/src/api.md
+++ b/docs/src/api.md
@@ -33,6 +33,7 @@ KernelAbstractions.pagelock!
 KernelAbstractions.unsafe_free!
 KernelAbstractions.functional
 KernelAbstractions.versioninfo
+KernelAbstractions.supports_subgroups
 KernelAbstractions.supports_unified
 KernelAbstractions.supports_atomics
 KernelAbstractions.supports_float64

From 632a6aa09aaa6377d1122599cc4631e94dc9018c Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Sat, 4 Jul 2026 16:19:51 -0300
Subject: [PATCH 14/17] Updaete AK branch

---
 .buildkite/pipeline.yml | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/.buildkite/pipeline.yml b/.buildkite/pipeline.yml
index 04766f59a..7e2269723 100644
--- a/.buildkite/pipeline.yml
+++ b/.buildkite/pipeline.yml
@@ -114,7 +114,7 @@ steps:
       julia -e 'println("--- :julia: Developing oneAPI")
                 using Pkg
                 Pkg.add(url="https://github.com/christiangnrd/oneAPI.jl", rev="intrinsicsnew")
-                Pkg.add(url="https://github.com/christiangnrd/AcceleratedKernels.jl", rev="ka0.10simple")'
+                Pkg.add(url="https://github.com/JuliaGPU/AcceleratedKernels.jl", rev="main")'
       julia -e 'println("--- :julia: Instantiating project")
                 using Pkg
                 Pkg.develop(; path=pwd())' || exit 3
@@ -146,7 +146,7 @@ steps:
     command: |
       julia -e 'println("--- :julia: Developing AMDGPU")
                 using Pkg
-                Pkg.add(url="https://github.com/christiangnrd/AcceleratedKernels.jl", rev="ka0.10simple")'
+                Pkg.add(url="https://github.com/JuliaGPU/AcceleratedKernels.jl", rev="main")'
       julia -e '
                 using Pkg
                 Pkg.add(url="https://github.com/christiangnrd/AMDGPU.jl", rev="intrinsicsnew")

From 25723c18048e696d2e99943f1b0dff0a1981c14a Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Sat, 4 Jul 2026 16:22:12 -0300
Subject: [PATCH 15/17] Readd comment

---
 src/pocl/backend.jl | 3 +++
 1 file changed, 3 insertions(+)

diff --git a/src/pocl/backend.jl b/src/pocl/backend.jl
index d8bad41d2..a72eae309 100644
--- a/src/pocl/backend.jl
+++ b/src/pocl/backend.jl
@@ -230,6 +230,9 @@ function KI.max_work_group_size(::POCLBackend)::Int
     return Int(device().max_work_group_size)
 end
 function KI.sub_group_size(::POCLBackend)::Int
+    # POCL can technically support any sub_group size.
+    #  Check for common values used on GPUs then
+    #  return 1 otherwise
     sg_sizes = cl.device().sub_group_sizes
     if 32 in sg_sizes
         return 32

From 80ba3cf52f2bb38f8a681b19a74732d5193d3823 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Wed, 8 Jul 2026 18:17:38 -0300
Subject: [PATCH 16/17] Remove `supports_subgroups`

All backends must support them
---
 docs/src/api.md           |  1 -
 src/KernelAbstractions.jl | 11 -----
 test/interface.jl         | 98 +++++++++++++++++++--------------------
 3 files changed, 48 insertions(+), 62 deletions(-)

diff --git a/docs/src/api.md b/docs/src/api.md
index 2881e6bf3..90e8af186 100644
--- a/docs/src/api.md
+++ b/docs/src/api.md
@@ -33,7 +33,6 @@ KernelAbstractions.pagelock!
 KernelAbstractions.unsafe_free!
 KernelAbstractions.functional
 KernelAbstractions.versioninfo
-KernelAbstractions.supports_subgroups
 KernelAbstractions.supports_unified
 KernelAbstractions.supports_atomics
 KernelAbstractions.supports_float64
diff --git a/src/KernelAbstractions.jl b/src/KernelAbstractions.jl
index 92c72a91d..1b6962cbd 100644
--- a/src/KernelAbstractions.jl
+++ b/src/KernelAbstractions.jl
@@ -647,17 +647,6 @@ function ones(backend::Backend, ::Type{T}, dims::Tuple; kwargs...) where {T}
     return data
 end
 
-"""
-    supports_subgroups(::Backend)::Bool
-
-Returns whether subgroups are supported by the backend.
-
-!!! note
-    Backend implementations **should** implement this function
-    only if they **do not** support subgroups.
-"""
-supports_subgroups(::Backend) = true
-
 """
     supports_unified(::Backend)::Bool
 
diff --git a/test/interface.jl b/test/interface.jl
index 39e07cab2..0deb73fec 100644
--- a/test/interface.jl
+++ b/test/interface.jl
@@ -182,58 +182,56 @@ function interface_testsuite(backend, AT)
             end
         end
 
-        if KernelAbstractions.supports_subgroups(backend())
-            @testset "Sub-groups" begin
-                @test KI.sub_group_size(backend()) isa Int
-
-                # Test with small kernel
-                sg_size = KI.sub_group_size(backend())
-                sg_n = 2
-                workgroupsize = sg_size * sg_n
-                numworkgroups = 2
-                N = workgroupsize * numworkgroups
-
-                results = AT(Vector{SubgroupData}(undef, N))
-                kernel = KI.@kernel backend() launch = false test_subgroup_kernel(results)
-
-                kernel(results; workgroupsize, numworkgroups)
-                KernelAbstractions.synchronize(backend())
-
-                host_results = Array(results)
-
-                # Verify results make sense
-                for (i, sg_data) in enumerate(host_results)
-                    @test sg_data.sub_group_size == sg_size
-                    @test sg_data.max_sub_group_size == sg_size
-                    @test sg_data.num_sub_groups == sg_n
-
-                    # Group ID should be 1-based
-                    expected_sub_group = div(((i - 1) % workgroupsize), sg_size) + 1
-                    @test sg_data.sub_group_id == expected_sub_group
-
-                    # Local ID should be 1-based within group
-                    expected_sg_local = ((i - 1) % sg_size) + 1
-                    @test sg_data.sub_group_local_id == expected_sg_local
-                end
+        @testset "Sub-groups" begin
+            @test KI.sub_group_size(backend()) isa Int
+
+            # Test with small kernel
+            sg_size = KI.sub_group_size(backend())
+            sg_n = 2
+            workgroupsize = sg_size * sg_n
+            numworkgroups = 2
+            N = workgroupsize * numworkgroups
+
+            results = AT(Vector{SubgroupData}(undef, N))
+            kernel = KI.@kernel backend() launch = false test_subgroup_kernel(results)
+
+            kernel(results; workgroupsize, numworkgroups)
+            KernelAbstractions.synchronize(backend())
+
+            host_results = Array(results)
+
+            # Verify results make sense
+            for (i, sg_data) in enumerate(host_results)
+                @test sg_data.sub_group_size == sg_size
+                @test sg_data.max_sub_group_size == sg_size
+                @test sg_data.num_sub_groups == sg_n
+
+                # Group ID should be 1-based
+                expected_sub_group = div(((i - 1) % workgroupsize), sg_size) + 1
+                @test sg_data.sub_group_id == expected_sub_group
+
+                # Local ID should be 1-based within group
+                expected_sg_local = ((i - 1) % sg_size) + 1
+                @test sg_data.sub_group_local_id == expected_sg_local
             end
-            @testset "shfl_down" begin
-                @test !isempty(KI.shfl_down_types(backend()))
-                types_to_test = setdiff(KI.shfl_down_types(backend()), [Bool])
-                @testset "$T" for T in types_to_test
-                    N = KI.sub_group_size(backend())
-                    a = zeros(T, N)
-                    rand!(a, (0:1))
-
-                    dev_a = AT(a)
-                    dev_b = AT(zeros(T, N))
-
-                    KI.@kernel backend() workgroupsize = N shfl_down_test_kernel(dev_a, dev_b, Val(N))
-
-                    b = Array(dev_b)
-                    @test sum(a) ≈ b[1]
-                end
+        end
+        @testset "shfl_down" begin
+            @test !isempty(KI.shfl_down_types(backend()))
+            types_to_test = setdiff(KI.shfl_down_types(backend()), [Bool])
+            @testset "$T" for T in types_to_test
+                N = KI.sub_group_size(backend())
+                a = zeros(T, N)
+                rand!(a, (0:1))
+
+                dev_a = AT(a)
+                dev_b = AT(zeros(T, N))
+
+                KI.@kernel backend() workgroupsize = N shfl_down_test_kernel(dev_a, dev_b, Val(N))
+
+                b = Array(dev_b)
+                @test sum(a) ≈ b[1]
             end
-        end # if !KernelAbstractions.supports_subgroups(backend())
+        end
     end # @testset "KernelInterface Tests" begin
     return nothing
 end

From 01b4b4b6418fa3ea5b38a74866619fd5373b92a7 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Wed, 8 Jul 2026 19:05:21 -0300
Subject: [PATCH 17/17] Fix again

---
 test/interface.jl | 97 ++++++++++++++++++++++++-----------------------
 1 file changed, 50 insertions(+), 47 deletions(-)

diff --git a/test/interface.jl b/test/interface.jl
index 0deb73fec..416ddcd1e 100644
--- a/test/interface.jl
+++ b/test/interface.jl
@@ -182,54 +182,57 @@ function interface_testsuite(backend, AT)
             end
         end
 
-        @testset "Sub-groups" begin
-            @test KI.sub_group_size(backend()) isa Int
-
-            # Test with small kernel
-            sg_size = KI.sub_group_size(backend())
-            sg_n = 2
-            workgroupsize = sg_size * sg_n
-            numworkgroups = 2
-            N = workgroupsize * numworkgroups
-
-            results = AT(Vector{SubgroupData}(undef, N))
-            kernel = KI.@kernel backend() launch = false test_subgroup_kernel(results)
-
-            kernel(results; workgroupsize, numworkgroups)
-            KernelAbstractions.synchronize(backend())
-
-            host_results = Array(results)
-
-            # Verify results make sense
-            for (i, sg_data) in enumerate(host_results)
-                @test sg_data.sub_group_size == sg_size
-                @test sg_data.max_sub_group_size == sg_size
-                @test sg_data.num_sub_groups == sg_n
-
-                # Group ID should be 1-based
-                expected_sub_group = div(((i - 1) % workgroupsize), sg_size) + 1
-                @test sg_data.sub_group_id == expected_sub_group
-
-                # Local ID should be 1-based within group
-                expected_sg_local = ((i - 1) % sg_size) + 1
-                @test sg_data.sub_group_local_id == expected_sg_local
+        # Used as a proxy for sub-group support
+        if !isempty(KI.shfl_down_types(backend()))
+            @testset "Sub-groups" begin
+                @test KI.sub_group_size(backend()) isa Int
+
+                # Test with small kernel
+                sg_size = KI.sub_group_size(backend())
+                sg_n = 2
+                workgroupsize = sg_size * sg_n
+                numworkgroups = 2
+                N = workgroupsize * numworkgroups
+
+                results = AT(Vector{SubgroupData}(undef, N))
+                kernel = KI.@kernel backend() launch = false test_subgroup_kernel(results)
+
+                kernel(results; workgroupsize, numworkgroups)
+                KernelAbstractions.synchronize(backend())
+
+                host_results = Array(results)
+
+                # Verify results make sense
+                for (i, sg_data) in enumerate(host_results)
+                    @test sg_data.sub_group_size == sg_size
+                    @test sg_data.max_sub_group_size == sg_size
+                    @test sg_data.num_sub_groups == sg_n
+
+                    # Group ID should be 1-based
+                    expected_sub_group = div(((i - 1) % workgroupsize), sg_size) + 1
+                    @test sg_data.sub_group_id == expected_sub_group
+
+                    # Local ID should be 1-based within group
+                    expected_sg_local = ((i - 1) % sg_size) + 1
+                    @test sg_data.sub_group_local_id == expected_sg_local
+                end
             end
-        end
-        @testset "shfl_down" begin
-            @test !isempty(KI.shfl_down_types(backend()))
-            types_to_test = setdiff(KI.shfl_down_types(backend()), [Bool])
-            @testset "$T" for T in types_to_test
-                N = KI.sub_group_size(backend())
-                a = zeros(T, N)
-                rand!(a, (0:1))
-
-                dev_a = AT(a)
-                dev_b = AT(zeros(T, N))
-
-                KI.@kernel backend() workgroupsize = N shfl_down_test_kernel(dev_a, dev_b, Val(N))
-
-                b = Array(dev_b)
-                @test sum(a) ≈ b[1]
+            @testset "shfl_down" begin
+                @test !isempty(KI.shfl_down_types(backend()))
+                types_to_test = setdiff(KI.shfl_down_types(backend()), [Bool])
+                @testset "$T" for T in types_to_test
+                    N = KI.sub_group_size(backend())
+                    a = zeros(T, N)
+                    rand!(a, (0:1))
+
+                    dev_a = AT(a)
+                    dev_b = AT(zeros(T, N))
+
+                    KI.@kernel backend() workgroupsize = N shfl_down_test_kernel(dev_a, dev_b, Val(N))
+
+                    b = Array(dev_b)
+                    @test sum(a) ≈ b[1]
+                end
             end
         end
     end # @testset "KernelInterface Tests" begin