From dce94c99bf85fddc6c55c8b01abfb38da7c4a4f6 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Wed, 5 Nov 2025 18:53:25 -0400
Subject: [PATCH 01/10] Add KernelIntrinsics support

---
 Project.toml         |  2 +-
 src/OpenCLKernels.jl | 68 +++++++++++++++++++++++++++++++-------------
 2 files changed, 50 insertions(+), 20 deletions(-)

diff --git a/Project.toml b/Project.toml
index ae9646f1..f4aa07f7 100644
--- a/Project.toml
+++ b/Project.toml
@@ -32,7 +32,7 @@ SPIRVIntrinsics = {path = "lib/intrinsics"}
 Adapt = "4"
 GPUArrays = "11.2.1"
 GPUCompiler = "1.7.1"
-KernelAbstractions = "0.9.38"
+KernelAbstractions = "0.10"
 LLVM = "9.1"
 LinearAlgebra = "1"
 OpenCL_jll = "=2024.10.24"
diff --git a/src/OpenCLKernels.jl b/src/OpenCLKernels.jl
index e06102cf..626b7fa0 100644
--- a/src/OpenCLKernels.jl
+++ b/src/OpenCLKernels.jl
@@ -1,9 +1,10 @@
 module OpenCLKernels
 
 using ..OpenCL
-using ..OpenCL: @device_override, method_table
+using ..OpenCL: @device_override, method_table, kernel_convert, clfunction
 
 import KernelAbstractions as KA
+import KernelAbstractions.KernelIntrinsics as KI
 
 import StaticArrays
 
@@ -126,33 +127,62 @@ function (obj::KA.Kernel{OpenCLBackend})(args...; ndrange=nothing, workgroupsize
     return nothing
 end
 
+KI.argconvert(::OpenCLBackend, arg) = kernel_convert(arg)
+
+function KI.kernel_function(::OpenCLBackend, f::F, tt::TT=Tuple{}; name = nothing, kwargs...) where {F,TT}
+    kern = clfunction(f, tt; name, kwargs...)
+    KI.Kernel{OpenCLBackend, typeof(kern)}(OpenCLBackend(), kern)
+end
+
+function (obj::KI.Kernel{OpenCLBackend})(args...; numworkgroups = 1, workgroupsize = 1)
+    KI.check_launch_args(numworkgroups, workgroupsize)
+
+    local_size = (workgroupsize..., ntuple(_ -> 1, 3 - length(workgroupsize))...)
+
+    numworkgroups = (numworkgroups..., ntuple(_ -> 1, 3 - length(numworkgroups))...)
+    global_size = local_size .* numworkgroups
+
+    obj.kern(args...; local_size, global_size)
+    return nothing
+end
+
+
+function KI.kernel_max_work_group_size(kernel::KI.Kernel{<:OpenCLBackend}; max_work_items::Int=typemax(Int))::Int
+    wginfo = cl.work_group_info(kernel.kern.fun, cl.device())
+    Int(min(wginfo.size, max_work_items))
+end
+function KI.max_work_group_size(::OpenCLBackend)::Int
+    Int(cl.device().max_work_group_size)
+end
+function KI.multiprocessor_count(::OpenCLBackend)::Int
+    Int(cl.device().max_compute_units)
+end
 
 ## Indexing Functions
+## COV_EXCL_START
 
-@device_override @inline function KA.__index_Local_Linear(ctx)
-    return get_local_id(1)
+@device_override @inline function KI.get_local_id()
+    return (; x = Int(get_local_id(1)), y = Int(get_local_id(2)), z = Int(get_local_id(3)))
 end
 
-@device_override @inline function KA.__index_Group_Linear(ctx)
-    return get_group_id(1)
+@device_override @inline function KI.get_group_id()
+    return (; x = Int(get_group_id(1)), y = Int(get_group_id(2)), z = Int(get_group_id(3)))
 end
 
-@device_override @inline function KA.__index_Global_Linear(ctx)
-    #return get_global_id(1)    # JuliaGPU/OpenCL.jl#346
-    I = KA.__index_Global_Cartesian(ctx)
-    @inbounds LinearIndices(KA.__ndrange(ctx))[I]
+@device_override @inline function KI.get_global_id()
+    return (; x = Int(get_global_id(1)), y = Int(get_global_id(2)), z = Int(get_global_id(3)))
 end
 
-@device_override @inline function KA.__index_Local_Cartesian(ctx)
-    @inbounds KA.workitems(KA.__iterspace(ctx))[get_local_id(1)]
+@device_override @inline function KI.get_local_size()
+    return (; x = Int(get_local_size(1)), y = Int(get_local_size(2)), z = Int(get_local_size(3)))
 end
 
-@device_override @inline function KA.__index_Group_Cartesian(ctx)
-    @inbounds KA.blocks(KA.__iterspace(ctx))[get_group_id(1)]
+@device_override @inline function KI.get_num_groups()
+    return (; x = Int(get_num_groups(1)), y = Int(get_num_groups(2)), z = Int(get_num_groups(3)))
 end
 
-@device_override @inline function KA.__index_Global_Cartesian(ctx)
-    return @inbounds KA.expand(KA.__iterspace(ctx), get_group_id(1), get_local_id(1))
+@device_override @inline function KI.get_global_size()
+    return (; x = Int(get_global_size(1)), y = Int(get_global_size(2)), z = Int(get_global_size(3)))
 end
 
 @device_override @inline function KA.__validindex(ctx)
@@ -167,7 +197,7 @@ end
 
 ## Shared and Scratch Memory
 
-@device_override @inline function KA.SharedMemory(::Type{T}, ::Val{Dims}, ::Val{Id}) where {T, Dims, Id}
+@device_override @inline function KI.localmemory(::Type{T}, ::Val{Dims}) where {T, Dims}
     ptr = OpenCL.emit_localmemory(T, Val(prod(Dims)))
     CLDeviceArray(Dims, ptr)
 end
@@ -179,14 +209,14 @@ end
 
 ## Synchronization and Printing
 
-@device_override @inline function KA.__synchronize()
+@device_override @inline function KI.barrier()
     work_group_barrier(OpenCL.LOCAL_MEM_FENCE | OpenCL.GLOBAL_MEM_FENCE)
 end
 
-@device_override @inline function KA.__print(args...)
+@device_override @inline function KI._print(args...)
     OpenCL._print(args...)
 end
-
+## COV_EXCL_STOP
 
 ## Other
 

From 8e80290cb301c30e10188e785abd6504e40d36ee Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Mon, 17 Nov 2025 22:02:15 -0400
Subject: [PATCH 02/10] Temp CI

---
 Project.toml      | 3 ++-
 test/Project.toml | 1 +
 2 files changed, 3 insertions(+), 1 deletion(-)

diff --git a/Project.toml b/Project.toml
index f4aa07f7..7ae28129 100644
--- a/Project.toml
+++ b/Project.toml
@@ -27,12 +27,13 @@ spirv2clc_jll = "f0274c0c-8c8a-59f1-85b7-f7d60330c5fb"
 
 [sources]
 SPIRVIntrinsics = {path = "lib/intrinsics"}
+KernelAbstractions = {rev = "main", url = "https://github.com/JuliaGPU/KernelAbstractions.jl"}
 
 [compat]
 Adapt = "4"
 GPUArrays = "11.2.1"
 GPUCompiler = "1.7.1"
-KernelAbstractions = "0.10"
+KernelAbstractions = "0.9, 0.10"
 LLVM = "9.1"
 LinearAlgebra = "1"
 OpenCL_jll = "=2024.10.24"
diff --git a/test/Project.toml b/test/Project.toml
index 64eacfcf..cb13f2dc 100644
--- a/test/Project.toml
+++ b/test/Project.toml
@@ -10,6 +10,7 @@ KernelAbstractions = "63c18a36-062a-441e-b654-da1e3ab1ce7c"
 LinearAlgebra = "37e2e46d-f89d-539d-b4ee-838fcccc9c8e"
 OpenCL = "08131aa3-fb12-5dee-8b74-c09406e224a2"
 ParallelTestRunner = "d3525ed8-44d0-4b2c-a655-542cee43accc"
+Pkg = "44cfe95a-1eb2-52ea-b672-e2afdf69b78f"
 Preferences = "21216c6a-2e73-6563-6e65-726566657250"
 Printf = "de0858da-6303-5e67-8744-51eddeeeb8d7"
 REPL = "3fa0cd96-eef1-5676-8a61-b3b8758bbffb"

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

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

diff --git a/src/OpenCLKernels.jl b/src/OpenCLKernels.jl
index 626b7fa0..98d8f372 100644
--- a/src/OpenCLKernels.jl
+++ b/src/OpenCLKernels.jl
@@ -4,7 +4,7 @@ using ..OpenCL
 using ..OpenCL: @device_override, method_table, kernel_convert, clfunction
 
 import KernelAbstractions as KA
-import KernelAbstractions.KernelIntrinsics as KI
+import KernelAbstractions.KernelInterface as KI
 
 import StaticArrays
 

From 8e1937a4e615f330efe024f57dce50a762a3bcc5 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Thu, 2 Jul 2026 11:00:29 -0300
Subject: [PATCH 04/10] KA.versioninfo()

---
 src/OpenCLKernels.jl | 2 ++
 1 file changed, 2 insertions(+)

diff --git a/src/OpenCLKernels.jl b/src/OpenCLKernels.jl
index 98d8f372..3b84e180 100644
--- a/src/OpenCLKernels.jl
+++ b/src/OpenCLKernels.jl
@@ -18,6 +18,8 @@ export OpenCLBackend
 struct OpenCLBackend <: KA.GPU
 end
 
+KA.versioninfo(io::IO, ::OpenCLBackend) = OpenCL.versioninfo(io)
+
 function KA.allocate(::OpenCLBackend, ::Type{T}, dims::Tuple; unified::Bool = false) where T
     if unified
         memory_backend = cl.unified_memory_backend()

From 13cee50ae3aa66475b7e20c1c22e8a07b0a144c9 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Sat, 4 Jul 2026 16:18:12 -0300
Subject: [PATCH 05/10] Skip CPU-only tests

---
 test/kernelabstractions.jl | 2 ++
 1 file changed, 2 insertions(+)

diff --git a/test/kernelabstractions.jl b/test/kernelabstractions.jl
index 97e54154..e53b79c8 100644
--- a/test/kernelabstractions.jl
+++ b/test/kernelabstractions.jl
@@ -15,5 +15,7 @@ end
 skip_tests=Set([
     "sparse",
     "Convert", # Need to opt out of i128
+    "CPU synchronization",
+    "fallback test: callable types"
 ])
 KATestSuite.testsuite(OpenCLBackend, "OpenCL", OpenCL, CLArray, CLDeviceArray; skip_tests)

From a7ccee2a4a922e70a80ec43a73ea1385a4315f01 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Thu, 22 Jan 2026 11:01:40 -0400
Subject: [PATCH 06/10] KI sub-groups

---
 src/OpenCLKernels.jl | 17 +++++++++++++++++
 1 file changed, 17 insertions(+)

diff --git a/src/OpenCLKernels.jl b/src/OpenCLKernels.jl
index 3b84e180..b0be4504 100644
--- a/src/OpenCLKernels.jl
+++ b/src/OpenCLKernels.jl
@@ -156,6 +156,9 @@ end
 function KI.max_work_group_size(::OpenCLBackend)::Int
     Int(cl.device().max_work_group_size)
 end
+function KI.sub_group_size(::OpenCLBackend)::Int
+    cl.sub_group_size(cl.device())
+end
 function KI.multiprocessor_count(::OpenCLBackend)::Int
     Int(cl.device().max_compute_units)
 end
@@ -187,6 +190,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 = KA.__index_Global_Cartesian(ctx)
@@ -215,6 +228,10 @@ end
     work_group_barrier(OpenCL.LOCAL_MEM_FENCE | OpenCL.GLOBAL_MEM_FENCE)
 end
 
+@device_override @inline function KI.sub_group_barrier()
+    sub_group_barrier(OpenCL.LOCAL_MEM_FENCE | OpenCL.GLOBAL_MEM_FENCE)
+end
+
 @device_override @inline function KI._print(args...)
     OpenCL._print(args...)
 end

From 0042c603003566be3e7a5f124192055fbdd7e474 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Wed, 7 Jan 2026 12:33:14 -0400
Subject: [PATCH 07/10] Subgroup shuffle

---
 src/OpenCLKernels.jl | 20 ++++++++++++++++++++
 1 file changed, 20 insertions(+)

diff --git a/src/OpenCLKernels.jl b/src/OpenCLKernels.jl
index b0be4504..10a1c25d 100644
--- a/src/OpenCLKernels.jl
+++ b/src/OpenCLKernels.jl
@@ -6,6 +6,8 @@ using ..OpenCL: @device_override, method_table, kernel_convert, clfunction
 import KernelAbstractions as KA
 import KernelAbstractions.KernelInterface as KI
 
+import SPIRVIntrinsics
+
 import StaticArrays
 
 import Adapt
@@ -163,6 +165,20 @@ function KI.multiprocessor_count(::OpenCLBackend)::Int
     Int(cl.device().max_compute_units)
 end
 
+function KI.shfl_down_types(::OpenCLBackend)
+    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
 ## COV_EXCL_START
 
@@ -232,6 +248,10 @@ end
     sub_group_barrier(OpenCL.LOCAL_MEM_FENCE | OpenCL.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...)
     OpenCL._print(args...)
 end

From b97c072b85600df231ec51a9cb3a5183765c6123 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Thu, 2 Jul 2026 10:57:44 -0300
Subject: [PATCH 08/10] Up

---
 src/OpenCLKernels.jl | 1 +
 1 file changed, 1 insertion(+)

diff --git a/src/OpenCLKernels.jl b/src/OpenCLKernels.jl
index 10a1c25d..d9232511 100644
--- a/src/OpenCLKernels.jl
+++ b/src/OpenCLKernels.jl
@@ -38,6 +38,7 @@ function KA.allocate(::OpenCLBackend, ::Type{T}, dims::Tuple; unified::Bool = fa
 end
 
 KA.supports_unified(::OpenCLBackend) = cl.default_memory_backend(cl.device(); unified=true) !== nothing
+KA.supports_subgroups(::OpenCLBackend) = cl.sub_groups_supported(cl.device())
 
 KA.get_backend(::CLArray) = OpenCLBackend()
 # TODO should be non-blocking

From 11ec4500e09ea0f78511ba22ce665bba1fee78a6 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Wed, 8 Jul 2026 18:18:37 -0300
Subject: [PATCH 09/10] Remove support_subgroups

---
 src/OpenCLKernels.jl | 1 -
 1 file changed, 1 deletion(-)

diff --git a/src/OpenCLKernels.jl b/src/OpenCLKernels.jl
index d9232511..10a1c25d 100644
--- a/src/OpenCLKernels.jl
+++ b/src/OpenCLKernels.jl
@@ -38,7 +38,6 @@ function KA.allocate(::OpenCLBackend, ::Type{T}, dims::Tuple; unified::Bool = fa
 end
 
 KA.supports_unified(::OpenCLBackend) = cl.default_memory_backend(cl.device(); unified=true) !== nothing
-KA.supports_subgroups(::OpenCLBackend) = cl.sub_groups_supported(cl.device())
 
 KA.get_backend(::CLArray) = OpenCLBackend()
 # TODO should be non-blocking

From 4c8efce3e9096e81579435124f6076b1489a3899 Mon Sep 17 00:00:00 2001
From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com>
Date: Wed, 8 Jul 2026 19:02:13 -0300
Subject: [PATCH 10/10] Fix

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

diff --git a/src/OpenCLKernels.jl b/src/OpenCLKernels.jl
index 10a1c25d..8ed382e7 100644
--- a/src/OpenCLKernels.jl
+++ b/src/OpenCLKernels.jl
@@ -166,6 +166,9 @@ function KI.multiprocessor_count(::OpenCLBackend)::Int
 end
 
 function KI.shfl_down_types(::OpenCLBackend)
+    backend_extensions = cl.device().extensions
+    "cl_khr_subgroup_shuffle" in backend_extensions || return DataType[]
+
     res = copy(SPIRVIntrinsics.gentypes)
 
     backend_extensions = cl.device().extensions