From bde6ac1636497d459aa0fb1fff881478a9027b42 Mon Sep 17 00:00:00 2001
From: Michel Schanen <michel.schanen@gmail.com>
Date: Fri, 12 Dec 2025 11:04:49 -0600
Subject: [PATCH] Iron out atomics and add tests for Float32

---
 src/device/atomics.jl     | 22 ++++++++++++++++++++++
 src/oneAPI.jl             |  1 +
 test/device/intrinsics.jl | 16 ++++++++--------
 3 files changed, 31 insertions(+), 8 deletions(-)
 create mode 100644 src/device/atomics.jl

diff --git a/src/device/atomics.jl b/src/device/atomics.jl
new file mode 100644
index 00000000..dc04b532
--- /dev/null
+++ b/src/device/atomics.jl
@@ -0,0 +1,22 @@
+# Atomic operation device overrides and fallbacks
+
+# Fallback wrappers for Float32 atomic_inc!/atomic_dec!
+# Intel Level Zero doesn't support these directly for floating-point types,
+# so we implement them using atomic_add!/atomic_sub!
+
+@device_override @inline function SPIRVIntrinsics.atomic_inc!(p::LLVMPtr{Float32,AS}) where {AS}
+    SPIRVIntrinsics.atomic_add!(p, Float32(1))
+end
+
+@device_override @inline function SPIRVIntrinsics.atomic_dec!(p::LLVMPtr{Float32,AS}) where {AS}
+    SPIRVIntrinsics.atomic_sub!(p, Float32(1))
+end
+
+# Float64 fallbacks (if Float64 is supported on device)
+@device_override @inline function SPIRVIntrinsics.atomic_inc!(p::LLVMPtr{Float64,AS}) where {AS}
+    SPIRVIntrinsics.atomic_add!(p, Float64(1))
+end
+
+@device_override @inline function SPIRVIntrinsics.atomic_dec!(p::LLVMPtr{Float64,AS}) where {AS}
+    SPIRVIntrinsics.atomic_sub!(p, Float64(1))
+end
diff --git a/src/oneAPI.jl b/src/oneAPI.jl
index b9caa398..7cd2a26c 100644
--- a/src/oneAPI.jl
+++ b/src/oneAPI.jl
@@ -34,6 +34,7 @@ Base.Experimental.@MethodTable(method_table)
 include("device/runtime.jl")
 include("device/array.jl")
 include("device/quirks.jl")
+include("device/atomics.jl")
 
 # essential stuff
 include("context.jl")
diff --git a/test/device/intrinsics.jl b/test/device/intrinsics.jl
index 5e5605ef..f27698c2 100644
--- a/test/device/intrinsics.jl
+++ b/test/device/intrinsics.jl
@@ -276,7 +276,7 @@ end
 
 @testset "atomics (low level)" begin
 
-@testset "atomic_add($T)" for T in [Int32, UInt32]
+@testset "atomic_add($T)" for T in [Int32, UInt32, Float32]
     a = oneArray([zero(T)])
 
     function kernel(a, b)
@@ -288,7 +288,7 @@ end
     @test Array(a)[1] == T(256)
 end
 
-@testset "atomic_sub($T)" for T in [Int32, UInt32]
+@testset "atomic_sub($T)" for T in [Int32, UInt32, Float32]
     a = oneArray([T(256)])
 
     function kernel(a, b)
@@ -300,7 +300,7 @@ end
     @test Array(a)[1] == T(0)
 end
 
-@testset "atomic_inc($T)" for T in [Int32, UInt32]
+@testset "atomic_inc($T)" for T in [Int32, UInt32, Float32]
     a = oneArray([zero(T)])
 
     function kernel(a)
@@ -312,7 +312,7 @@ end
     @test Array(a)[1] == T(256)
 end
 
-@testset "atomic_dec($T)" for T in [Int32, UInt32]
+@testset "atomic_dec($T)" for T in [Int32, UInt32, Float32]
     a = oneArray([T(256)])
 
     function kernel(a)
@@ -324,12 +324,12 @@ end
     @test Array(a)[1] == T(0)
 end
 
-@testset "atomic_min($T)" for T in [Int32, UInt32]
+@testset "atomic_min($T)" for T in [Int32, UInt32, Float32]
     a = oneArray([T(256)])
 
     function kernel(a, T)
         i = get_global_id()
-        oneAPI.atomic_min!(pointer(a), i%T)
+        oneAPI.atomic_min!(pointer(a), T(i))
         return
     end
 
@@ -337,12 +337,12 @@ end
     @test Array(a)[1] == one(T)
 end
 
-@testset "atomic_max($T)" for T in [Int32, UInt32]
+@testset "atomic_max($T)" for T in [Int32, UInt32, Float32]
     a = oneArray([zero(T)])
 
     function kernel(a, T)
         i = get_global_id()
-        oneAPI.atomic_max!(pointer(a), i%T)
+        oneAPI.atomic_max!(pointer(a), T(i))
         return
     end