Add atomic float support

[michel2323]

Enable SPV_EXT_shader_atomic_float_add
Fixes #508

[github-actions]

Your PR requires formatting changes to meet the project's style guidelines.
Please consider running Runic (git runic master) to apply these changes.

Click here to view the suggested changes.

diff --git a/src/device/atomics.jl b/src/device/atomics.jl
index dc04b53..346db57 100644
--- a/src/device/atomics.jl
+++ b/src/device/atomics.jl
@@ -4,19 +4,19 @@
 # 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}
+@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}
+@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}
+@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}
+@device_override @inline function SPIRVIntrinsics.atomic_dec!(p::LLVMPtr{Float64, AS}) where {AS}
     SPIRVIntrinsics.atomic_sub!(p, Float64(1))
 end
diff --git a/test/device/intrinsics.jl b/test/device/intrinsics.jl
index f27698c..32045e0 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, Float32]
+    @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, Float32]
+    @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, Float32]
+    @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, Float32]
+    @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, Float32]
+    @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), T(i))
+            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, Float32]
+    @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), T(i))
+            oneAPI.atomic_max!(pointer(a), T(i))
         return
     end
 

[michel2323]

Needs tests.

[codecov]

Codecov Report

✅ All modified and coverable lines are covered by tests.
✅ Project coverage is 79.10%. Comparing base (0c9c8b7) to head (fa63ecb).
⚠️ Report is 1 commits behind head on master.

Additional details and impacted files

@@           Coverage Diff           @@
##           master     #544   +/-   ##
=======================================
  Coverage   79.10%   79.10%           
=======================================
  Files          47       47           
  Lines        3001     3001           
=======================================
  Hits         2374     2374           
  Misses        627      627           

☔ View full report in Codecov by Sentry.
📢 Have feedback on the report? Share it here.

🚀 New features to boost your workflow:

• ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.