Iron out atomics and add tests for Float32

Author: michel2323

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

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")

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,25 +324,25 @@ 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
 
     @oneapi items=256 kernel(a, T)
     @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