Add atomic float support #399

michel2323 · 2025-10-31T17:07:02Z

JuliaGPU/oneAPI.jl#544 Adding the SPV_EXT_shader_atomic_float_add extension in oneAPI.jl seems to work.

github-actions · 2025-10-31T17:08:03Z

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/lib/intrinsics/src/atomic.jl b/lib/intrinsics/src/atomic.jl
index a93f926..d627741 100644
--- a/lib/intrinsics/src/atomic.jl
+++ b/lib/intrinsics/src/atomic.jl
@@ -47,17 +47,21 @@ for gentype in atomic_types, as in atomic_memory_types
 @device_function atomic_xor!(p::LLVMPtr{$gentype,$as}, val::$gentype) =
     @builtin_ccall("atomic_xor", $gentype,
                    (LLVMPtr{$gentype,$as}, $gentype), p, val)
-end
-if gentype in atomic_integer_types
-    @eval begin
-    @device_function atomic_xchg!(p::LLVMPtr{$gentype,$as}, val::$gentype) =
-        @builtin_ccall("atomic_xchg", $gentype,
-                    (LLVMPtr{$gentype,$as}, $gentype), p, val)
-
-    @device_function atomic_cmpxchg!(p::LLVMPtr{$gentype,$as}, cmp::$gentype, val::$gentype) =
-        @builtin_ccall("atomic_cmpxchg", $gentype,
-                    (LLVMPtr{$gentype,$as}, $gentype, $gentype), p, cmp, val)
     end
+    if gentype in atomic_integer_types
+        @eval begin
+            @device_function atomic_xchg!(p::LLVMPtr{$gentype, $as}, val::$gentype) =
+                @builtin_ccall(
+                "atomic_xchg", $gentype,
+                (LLVMPtr{$gentype, $as}, $gentype), p, val
+            )
+
+            @device_function atomic_cmpxchg!(p::LLVMPtr{$gentype, $as}, cmp::$gentype, val::$gentype) =
+                @builtin_ccall(
+                "atomic_cmpxchg", $gentype,
+                (LLVMPtr{$gentype, $as}, $gentype, $gentype), p, cmp, val
+            )
+        end
 end
 end
 
diff --git a/test/atomics.jl b/test/atomics.jl
index 71fefa8..9a7799e 100644
--- a/test/atomics.jl
+++ b/test/atomics.jl
@@ -8,43 +8,43 @@ all_types = vcat(integer_types, float_types)
 dev = OpenCL.cl.device()
 
 # Arithmetic operations
-function test_atomic_add(counter::AbstractArray{T}) where T
+function test_atomic_add(counter::AbstractArray{T}) where {T}
     OpenCL.@atomic counter[] += one(T)
     return
 end
-function test_atomic_sub(counter::AbstractArray{T}) where T
+function test_atomic_sub(counter::AbstractArray{T}) where {T}
     OpenCL.@atomic counter[] -= one(T)
     return
 end
 # Bitwise operations
-function test_atomic_and(counter::AbstractArray{T}) where T
+function test_atomic_and(counter::AbstractArray{T}) where {T}
     OpenCL.@atomic counter[] &= ~(one(T) << (get_global_id() - 1))
     return
 end
-function test_atomic_or(counter::AbstractArray{T}) where T
+function test_atomic_or(counter::AbstractArray{T}) where {T}
     OpenCL.@atomic counter[] |= one(T) << (get_global_id() - 1)
     return
 end
-function test_atomic_xor(counter::AbstractArray{T}) where T
+function test_atomic_xor(counter::AbstractArray{T}) where {T}
     OpenCL.@atomic counter[] ⊻= one(T) << ((get_global_id() - 1) % 32)
     return
 end
 # Min/max operations - use low-level API directly
-function test_atomic_max(counter::AbstractArray{T}) where T
+function test_atomic_max(counter::AbstractArray{T}) where {T}
     OpenCL.atomic_max!(pointer(counter), T(get_global_id()))
     return
 end
-function test_atomic_min(counter::AbstractArray{T}) where T
+function test_atomic_min(counter::AbstractArray{T}) where {T}
     OpenCL.atomic_min!(pointer(counter), T(get_global_id()))
     return
 end
 # Exchange operation - use low-level API directly
-function test_atomic_xchg(counter::AbstractArray{T}) where T
+function test_atomic_xchg(counter::AbstractArray{T}) where {T}
     OpenCL.atomic_xchg!(pointer(counter), one(T))
     return
 end
 # Compare-and-swap operation - use low-level API directly
-function test_atomic_cas(counter::AbstractArray{T}) where T
+function test_atomic_cas(counter::AbstractArray{T}) where {T}
     OpenCL.atomic_cmpxchg!(pointer(counter), zero(T), one(T))
     return
 end
@@ -63,61 +63,61 @@ atomic_operations = [
     (test_atomic_cas, 0, 1),
 ]
 @testset "atomics" begin
-@testset "$kernel_func - $T" for (kernel_func, init_val, expected_val) in atomic_operations, T in all_types
-    # Skip Int64/UInt64 if not supported
-    if sizeof(T) == 8 && T <: Integer && !("cl_khr_int64_extended_atomics" in dev.extensions)
-        continue
-    end
+    @testset "$kernel_func - $T" for (kernel_func, init_val, expected_val) in atomic_operations, T in all_types
+        # Skip Int64/UInt64 if not supported
+        if sizeof(T) == 8 && T <: Integer && !("cl_khr_int64_extended_atomics" in dev.extensions)
+            continue
+        end
 
-    # Skip Float64 if not supported
-    if T == Float64 && !("cl_khr_fp64" in dev.extensions)
-        continue
-    end
+        # Skip Float64 if not supported
+        if T == Float64 && !("cl_khr_fp64" in dev.extensions)
+            continue
+        end
 
-    # Bitwise operations (only valid for integers)
-    if kernel_func in [test_atomic_and, test_atomic_or, test_atomic_xor] && T <: AbstractFloat
-        continue
+        # Bitwise operations (only valid for integers)
+        if kernel_func in [test_atomic_and, test_atomic_or, test_atomic_xor] && T <: AbstractFloat
+            continue
     end
 
-    # Min/max operations (only supported for 32-bit integers in OpenCL)
-    if kernel_func in [test_atomic_min, test_atomic_max] && !(T in [Int32, UInt32])
-        continue
-    end
+        # Min/max operations (only supported for 32-bit integers in OpenCL)
+        if kernel_func in [test_atomic_min, test_atomic_max] && !(T in [Int32, UInt32])
+            continue
+        end
 
-    if T <: Integer
-        init_val %= T
-        expected_val %= T
-    end
+        if T <: Integer
+            init_val %= T
+            expected_val %= T
+        end
 
-    a = OpenCL.fill(T(init_val))
-    @opencl global_size=1000 kernel_func(a)
-    result_val = OpenCL.@allowscalar a[]
-    @test result_val === T(expected_val)
+        a = OpenCL.fill(T(init_val))
+        @opencl global_size = 1000 kernel_func(a)
+        result_val = OpenCL.@allowscalar a[]
+        @test result_val === T(expected_val)
 end
 
 
-@testset "atomic_add! ($T)" for T in [Float32, Float64]
-    # Float64 requires cl_khr_fp64 extension
-    if T == Float64 && !("cl_khr_fp64" in cl.device().extensions)
-        continue
-    end
+    @testset "atomic_add! ($T)" for T in [Float32, Float64]
+        # Float64 requires cl_khr_fp64 extension
+        if T == Float64 && !("cl_khr_fp64" in cl.device().extensions)
+            continue
+        end
     if "cl_ext_float_atomics" in cl.device().extensions
-        @eval function atomic_float_add(counter, val::$T)
+            @eval function atomic_float_add(counter, val::$T)
             @builtin_ccall(
-                "atomic_add", $T,
-                (LLVMPtr{$T, AS.CrossWorkgroup}, $T),
+                    "atomic_add", $T,
+                    (LLVMPtr{$T, AS.CrossWorkgroup}, $T),
                 pointer(counter), val,
             )
             return
         end
 
         @testset "SPV_EXT_shader_atomic_float_add extension" begin
-            a = OpenCL.zeros(T)
-            @opencl global_size = 1000 extensions = ["SPV_EXT_shader_atomic_float_add"] atomic_float_add(a, one(T))
-            @test OpenCL.@allowscalar a[] == T(1000.0)
+                a = OpenCL.zeros(T)
+                @opencl global_size = 1000 extensions = ["SPV_EXT_shader_atomic_float_add"] atomic_float_add(a, one(T))
+                @test OpenCL.@allowscalar a[] == T(1000.0)
 
             spv = sprint() do io
-                OpenCL.code_native(io, atomic_float_add, Tuple{CLDeviceArray{T, 0, 1}, T}; extensions = ["SPV_EXT_shader_atomic_float_add"])
+                    OpenCL.code_native(io, atomic_float_add, Tuple{CLDeviceArray{T, 0, 1}, T}; extensions = ["SPV_EXT_shader_atomic_float_add"])
             end
             @test occursin("OpExtension \"SPV_EXT_shader_atomic_float_add\"", spv)
             @test occursin("OpAtomicFAddEXT", spv)

codecov · 2025-10-31T17:09:47Z

Codecov Report

✅ All modified and coverable lines are covered by tests.
✅ Project coverage is 80.68%. Comparing base (71cf159) to head (e12f093).
⚠️ Report is 1 commits behind head on master.

Additional details and impacted files

@@            Coverage Diff             @@
##           master     #399      +/-   ##
==========================================
+ Coverage   80.27%   80.68%   +0.41%     
==========================================
  Files          12       12              
  Lines         730      730              
==========================================
+ Hits          586      589       +3     
+ Misses        144      141       -3

☔ 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.

michel2323 · 2025-11-06T18:24:39Z

Bump. Can someone take a look? @vchuravy ? I don't see how the failing tests are connected to the changes here.

maleadt · 2025-11-07T10:32:50Z

Thanks. Can you add a test that works here too?
CI failures look unrelated.

michel2323 · 2025-11-18T14:32:18Z

@maleadt Added tests and rebased.

simeonschaub

The test can now be modified to actually call the atomic_add! method from SPIRVIntrinsics instead of defining its own function. It is also not exercising the Float64 path ATM, since Float32 is hard coded inside the test.

Ideally, we would also be testing the other atomic intrinsics that have been defined for Float32 and Float64 now. (Are the bitwise intrinsics like atomic_or! and friends even defined for floats?)

michel2323 · 2025-11-19T19:03:27Z

@simeonschaub I tried to add more tests, but except for atomic_add, they all seem to fail, and I don't understand why. I see something about a fallback in the definition of @atomic. I've added atomic_sub that fails.

Can we remove the @builtin_ccall based tests?

michel2323 · 2025-11-20T16:26:06Z

A restful night resolved the issues. This should pass now. A few comments:

OpenCL only supports Int32 for min, max. Feels weird.
Bitwise atomics only for integers

michel2323 · 2025-11-25T15:53:50Z

@simeonschaub @maleadt @vchuravy Bump. I don't think the errors are related to my code changes.

simeonschaub

Sorry for taking so long to review! I took the liberty of improving some of the tests a little, hope that's ok

michel2323 · 2025-11-26T14:14:38Z

Of course! Who can merge? @vchuravy?

michel2323 mentioned this pull request Oct 31, 2025

Issues with floating-point atomics JuliaGPU/oneAPI.jl#508

Closed

michel2323 force-pushed the ms/atomic_float branch from 1ac8383 to e962893 Compare November 18, 2025 14:31

simeonschaub requested changes Nov 18, 2025

View reviewed changes

michel2323 requested a review from simeonschaub November 20, 2025 16:32

michel2323 and others added 7 commits November 25, 2025 09:56

Add atomic float support

1900e8b

Add tests

5a8199f

More tests

c1b2f5f

Compilation works

a9a6373

Check results

1fff397

Using 1000 items and checking type

3c58e28

Fix

58f3b27

michel2323 force-pushed the ms/atomic_float branch from 8d36d26 to 58f3b27 Compare November 25, 2025 15:56

avoid use of eval, improve some of the tests

5a1a66b

simeonschaub approved these changes Nov 26, 2025

View reviewed changes

bump SPIRVIntrinsics version

e12f093

simeonschaub enabled auto-merge (squash) November 27, 2025 12:44

simeonschaub merged commit db69fa8 into JuliaGPU:master Nov 27, 2025
29 of 37 checks passed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

Add atomic float support #399

Add atomic float support #399

Uh oh!

michel2323 commented Oct 31, 2025

Uh oh!

github-actions bot commented Oct 31, 2025 •

edited

Loading

Uh oh!

codecov bot commented Oct 31, 2025 •

edited

Loading

Uh oh!

michel2323 commented Nov 6, 2025

Uh oh!

maleadt commented Nov 7, 2025

Uh oh!

michel2323 commented Nov 18, 2025

Uh oh!

simeonschaub left a comment

Uh oh!

michel2323 commented Nov 19, 2025 •

edited

Loading

Uh oh!

michel2323 commented Nov 20, 2025 •

edited

Loading

Uh oh!

michel2323 commented Nov 25, 2025 •

edited

Loading

Uh oh!

simeonschaub left a comment

Uh oh!

michel2323 commented Nov 26, 2025

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Add atomic float support #399

Add atomic float support #399

Uh oh!

Conversation

michel2323 commented Oct 31, 2025

Uh oh!

github-actions bot commented Oct 31, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

codecov bot commented Oct 31, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Codecov Report

Uh oh!

michel2323 commented Nov 6, 2025

Uh oh!

maleadt commented Nov 7, 2025

Uh oh!

michel2323 commented Nov 18, 2025

Uh oh!

simeonschaub left a comment

Choose a reason for hiding this comment

Uh oh!

michel2323 commented Nov 19, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

michel2323 commented Nov 20, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

michel2323 commented Nov 25, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

simeonschaub left a comment

Choose a reason for hiding this comment

Uh oh!

michel2323 commented Nov 26, 2025

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

github-actions bot commented Oct 31, 2025 •

edited

Loading

codecov bot commented Oct 31, 2025 •

edited

Loading

michel2323 commented Nov 19, 2025 •

edited

Loading

michel2323 commented Nov 20, 2025 •

edited

Loading

michel2323 commented Nov 25, 2025 •

edited

Loading