Add Atomix and UnsafeAtomics atomics tests by vchuravy · Pull Request #726 · JuliaGPU/KernelAbstractions.jl

vchuravy · 2026-07-05T19:52:48Z

Adds an Atomics testset to the shared testsuite covering the atomics functionality KernelAbstractions re-exports from Atomix, plus raw-pointer atomics from UnsafeAtomics.

Atomix (all backends, gated on `supports_atomics(backend())`)

contended histogram @atomic += for Int32/UInt32/Float32/Float64
@atomic max/@atomic min reductions across the ndrange
atomic load (@atomic A[i]) and store (@atomic A[i] = v)
@atomicswap
@atomicreplace with both succeeding and failing CAS

UnsafeAtomics (CPU backend only, `@test_skip` elsewhere)

contended add! histogram for Int32/UInt32/Float32/Float64
max!/min!
store!/modify!/cas!/xchg!/load sequence
add! with explicit seq_cst ordering

Notes:

Depends on [pocl] Enable float add and min/max atomics via cl_ext_float_atomics #737, which permits the SPV_EXT_shader_atomic_float_{add,min_max} extensions based on cl_ext_float_atomics. This PR is rebased on top of it and extends the test/atomics.jl it introduces: add and min/max run for Float32 and, where supports_float64, Float64, for both Atomix and UnsafeAtomics. Merge [pocl] Enable float add and min/max atomics via cl_ext_float_atomics #737 first.
UnsafeAtomics is added as a direct test dependency, so backend packages that include this testsuite need it in their test environments once they pick up this change. Backends can also opt out via skip_tests = Set(["Atomics"]).

🤖 Generated with Claude Code

vchuravy · 2026-07-05T19:56:15Z

+        if !(backend() isa CPU)
+            @test_skip "UnsafeAtomics tests only run on the CPU backend"
+            return
+        end


Suggested change

if !(backend() isa CPU)

@test_skip "UnsafeAtomics tests only run on the CPU backend"

return

end

Done in 291db15 — the CPU-only gate is removed, so the UnsafeAtomics tests now run on every backend.

vchuravy

Missing:

fences
orderings
scopes

github-actions · 2026-07-05T20:11:46Z

Benchmark Results

Show table

	main	`f7698e2`...	main / `f7698e2`...
saxpy/default/Float32/1024	0.0666 ± 0.039 ms	0.0594 ± 0.034 ms	1.12 ± 0.92
saxpy/default/Float32/1048576	0.595 ± 0.037 ms	0.507 ± 0.022 ms	1.17 ± 0.088
saxpy/default/Float32/16384	0.0594 ± 0.031 ms	0.0558 ± 0.03 ms	1.06 ± 0.81
saxpy/default/Float32/2048	0.073 ± 0.031 ms	0.0627 ± 0.027 ms	1.16 ± 0.71
saxpy/default/Float32/256	0.0536 ± 0.039 ms	0.0552 ± 0.035 ms	0.971 ± 0.94
saxpy/default/Float32/262144	0.207 ± 0.028 ms	0.18 ± 0.024 ms	1.15 ± 0.22
saxpy/default/Float32/32768	0.0699 ± 0.035 ms	0.06 ± 0.031 ms	1.16 ± 0.83
saxpy/default/Float32/4096	0.0752 ± 0.033 ms	0.0661 ± 0.026 ms	1.14 ± 0.67
saxpy/default/Float32/512	0.0473 ± 0.039 ms	0.0521 ± 0.035 ms	0.908 ± 0.97
saxpy/default/Float32/64	0.0615 ± 0.04 ms	0.0586 ± 0.035 ms	1.05 ± 0.92
saxpy/default/Float32/65536	0.105 ± 0.034 ms	0.084 ± 0.03 ms	1.25 ± 0.6
saxpy/default/Float64/1024	0.0694 ± 0.037 ms	0.061 ± 0.03 ms	1.14 ± 0.83
saxpy/default/Float64/1048576	0.639 ± 0.13 ms	0.53 ± 0.27 ms	1.21 ± 0.67
saxpy/default/Float64/16384	0.0642 ± 0.034 ms	0.0712 ± 0.03 ms	0.902 ± 0.61
saxpy/default/Float64/2048	0.0742 ± 0.032 ms	0.0623 ± 0.03 ms	1.19 ± 0.77
saxpy/default/Float64/256	0.0503 ± 0.039 ms	0.0572 ± 0.035 ms	0.88 ± 0.86
saxpy/default/Float64/262144	0.21 ± 0.021 ms	0.18 ± 0.017 ms	1.17 ± 0.16
saxpy/default/Float64/32768	0.0889 ± 0.034 ms	0.0782 ± 0.03 ms	1.14 ± 0.61
saxpy/default/Float64/4096	0.0779 ± 0.027 ms	0.0668 ± 0.025 ms	1.17 ± 0.6
saxpy/default/Float64/512	0.0506 ± 0.039 ms	0.0558 ± 0.034 ms	0.906 ± 0.89
saxpy/default/Float64/64	0.0561 ± 0.04 ms	0.0576 ± 0.034 ms	0.975 ± 0.9
saxpy/default/Float64/65536	0.109 ± 0.031 ms	0.093 ± 0.027 ms	1.17 ± 0.48
saxpy/static workgroup=(1024,)/Float32/1024	0.0493 ± 0.039 ms	0.058 ± 0.034 ms	0.85 ± 0.84
saxpy/static workgroup=(1024,)/Float32/1048576	0.531 ± 0.027 ms	0.455 ± 0.024 ms	1.17 ± 0.085
saxpy/static workgroup=(1024,)/Float32/16384	0.0566 ± 0.03 ms	0.0524 ± 0.028 ms	1.08 ± 0.82
saxpy/static workgroup=(1024,)/Float32/2048	0.0716 ± 0.03 ms	0.0621 ± 0.028 ms	1.15 ± 0.71
saxpy/static workgroup=(1024,)/Float32/256	0.0504 ± 0.04 ms	0.0578 ± 0.034 ms	0.872 ± 0.85
saxpy/static workgroup=(1024,)/Float32/262144	0.191 ± 0.033 ms	0.165 ± 0.03 ms	1.16 ± 0.29
saxpy/static workgroup=(1024,)/Float32/32768	0.066 ± 0.034 ms	0.0578 ± 0.029 ms	1.14 ± 0.82
saxpy/static workgroup=(1024,)/Float32/4096	0.0743 ± 0.028 ms	0.0659 ± 0.023 ms	1.13 ± 0.59
saxpy/static workgroup=(1024,)/Float32/512	0.0482 ± 0.039 ms	0.0541 ± 0.035 ms	0.891 ± 0.92
saxpy/static workgroup=(1024,)/Float32/64	0.0659 ± 0.039 ms	0.0559 ± 0.034 ms	1.18 ± 1
saxpy/static workgroup=(1024,)/Float32/65536	0.0925 ± 0.034 ms	0.081 ± 0.03 ms	1.14 ± 0.6
saxpy/static workgroup=(1024,)/Float64/1024	0.0671 ± 0.037 ms	0.0602 ± 0.03 ms	1.11 ± 0.83
saxpy/static workgroup=(1024,)/Float64/1048576	0.552 ± 0.025 ms	0.474 ± 0.02 ms	1.16 ± 0.071
saxpy/static workgroup=(1024,)/Float64/16384	0.0618 ± 0.032 ms	0.0678 ± 0.028 ms	0.911 ± 0.61
saxpy/static workgroup=(1024,)/Float64/2048	0.0731 ± 0.031 ms	0.0634 ± 0.027 ms	1.15 ± 0.7
saxpy/static workgroup=(1024,)/Float64/256	0.0477 ± 0.039 ms	0.0607 ± 0.032 ms	0.786 ± 0.77
saxpy/static workgroup=(1024,)/Float64/262144	0.196 ± 0.032 ms	0.169 ± 0.028 ms	1.16 ± 0.27
saxpy/static workgroup=(1024,)/Float64/32768	0.0848 ± 0.034 ms	0.0728 ± 0.03 ms	1.16 ± 0.67
saxpy/static workgroup=(1024,)/Float64/4096	0.074 ± 0.027 ms	0.0651 ± 0.024 ms	1.14 ± 0.59
saxpy/static workgroup=(1024,)/Float64/512	0.055 ± 0.04 ms	0.0583 ± 0.034 ms	0.944 ± 0.87
saxpy/static workgroup=(1024,)/Float64/64	0.0485 ± 0.039 ms	0.0574 ± 0.033 ms	0.846 ± 0.85
saxpy/static workgroup=(1024,)/Float64/65536	0.103 ± 0.034 ms	0.0885 ± 0.03 ms	1.16 ± 0.54
time_to_load	1.08 ± 0.0077 s	0.93 ± 0.004 s	1.16 ± 0.0097

Benchmark Plots

A plot of the benchmark results have been uploaded as an artifact to the workflow run for this PR.
Go to "Actions"->"Benchmark a pull request"->[the most recent run]->"Artifacts" (at the bottom).

vchuravy · 2026-07-06T08:26:57Z

+    @testset "Atomix" begin
+        # Float32 is excluded since atomic float add requires the SPIR-V
+        # extension SPV_EXT_shader_atomic_float_add, unavailable with PoCL.
+        @testset "atomic add ($T)" for T in (Int32, UInt32)


We should implement fallbacks for this JuliaGPU/GPUCompiler.jl#652

Added a TODO comment referencing JuliaGPU/GPUCompiler.jl#652 next to the Float32 exclusion (291db15).

simeonschaub · 2026-07-06T13:27:15Z

Float atomic add is excluded: the CPU backend compiles through PoCL and atomic float add requires the SPV_EXT_shader_atomic_float_add SPIR-V extension, which PoCL does not provide (LLVM ERROR: The atomic float instruction requires the following SPIR-V extension: SPV_EXT_shader_atomic_float_add).

I just tested it, PoCL supports SPV_EXT_shader_atomic_float_min_max just fine, the extension just needs to be enabled:

julia> using OpenCL, KernelAbstractions, SPIRVIntrinsics, pocl_jll

julia> gentype = Float32; as = 1
1

julia> @eval SPIRVIntrinsics @device_function atomic_max!(p::LLVMPtr{$gentype,$as}, val::$gentype) =
           @builtin_ccall("__spirv_AtomicFMaxEXT", $gentype,
                          (LLVMPtr{$gentype,$as}, UInt32, UInt32, $gentype),
                          p, UInt32(atomic_scope),
                          UInt32(atomic_memory_semantics(Val($as))), val)

julia> function atomix_max!(A)
           @inbounds KernelAbstractions.@atomic max(A[1], eltype(A)(get_global_id()))
           return nothing
       end
atomix_max! (generic function with 4 methods)

julia> a = OpenCL.zeros(T)
0-dimensional CLArray{Float32, 0, OpenCL.cl.UnifiedDeviceMemory}:
0.0

julia> @opencl global_size = 1000 extensions = ["SPV_EXT_shader_atomic_float_min_max"] atomix_max!(a)
OpenCL.HostKernel{typeof(atomix_max!), Tuple{CLDeviceArray{Float32, 0, 1}}}(atomix_max!, OpenCL.Kernel("_Z11atomix_max_13CLDeviceArrayI7Float32Li0ELi1EE" nargs=2), false)

julia> a
0-dimensional CLArray{Float32, 0, OpenCL.cl.UnifiedDeviceMemory}:
1000.0

Probably makes sense to enable by default for PoCLBackend, or what do you think?

vchuravy · 2026-07-09T15:09:53Z

Addressed the review in 291db15:

UnsafeAtomics tests now run on all backends (CPU-only gate removed, per suggestion).
Orderings: Atomix kernel exercising :release/:acquire load & store, :monotonic/:acquire_release/:sequentially_consistent RMW, and ordered @atomicswap; UnsafeAtomics add parametrized over monotonic/acquire/release/acq_rel/seq_cst, plus acquire/release load & store.
Fences: non-blocking message-passing test — workitem 1 publishes data with fence(release) + monotonic flag store, observers check flag with monotonic load + fence(acquire); anyone who saw the flag must see the data.
Scopes: kernel using system-scope (none) contended adds and singlethread-scoped store/fence/add.

25 tests passing on the CPU backend.

The SPIR-V backend refuses to translate a module containing an atomic fadd unless the extension guarding it has been listed: LLVM ERROR: The atomic float instruction requires the following SPIR-V extension: SPV_EXT_shader_atomic_float_add Enzyme's reverse mode runs into this because it accumulates gradients with atomic fadd, so reverse-mode AD over a POCL kernel fails to compile outright. Derive the extension list from the device the way `supports_fp16`/`supports_fp64` already are, keyed off the corresponding OpenCL extension (`cl_ext_float_atomics`, which PoCL advertises). Listing an extension only permits it -- nothing is emitted unless a module needs those instructions -- so kernels that do not use float atomics are unaffected. An explicit `extensions=` keyword still wins, and `default_spirv_extensions` gives later extensions an obvious home. Verified end-to-end: with this change, reverse-mode Enzyme over a POCL kernel compiles, runs and produces correct gradients. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

…pabilities Atomix lowers `@atomic A[i] += x` and `@atomic max(A[i], x)` on floats to LLVM `atomicrmw fadd`/`fmin`/`fmax`, which the SPIR-V backend only translates when SPV_EXT_shader_atomic_float_add resp. SPV_EXT_shader_atomic_float_min_max is permitted. Derive both from the per-precision cl_ext_float_atomics capability bitfields instead of the bare extension string, requiring the global and local memory bits like OpenCL.jl does, and expose the bitfields as device properties. Adds an `Atomics` testsuite entry exercising add/min/max for Int32, UInt32, Float32 and (where supported) Float64 on all backends, and a POCL-specific check that both extensions are permitted on pocl's CPU device while an explicit `extensions=` still wins. Assisted-by: Claude Code (Fable 5.1)

Adds an Atomics testsuite exercising @atomic/@atomicswap/@atomicreplace (via Atomix) in kernels on all backends, plus UnsafeAtomics pointer-based atomics on the CPU backend. UnsafeAtomics becomes a direct test dependency. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

…gs, fences, and syncscopes Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

…nsions Run the Atomix and UnsafeAtomics add and min/max tests for Float32 and, where the backend supports it, Float64 in addition to the integer types. Assisted-by: Claude Code (Fable 5.1)

vchuravy requested a review from christiangnrd July 5, 2026 19:55

vchuravy commented Jul 5, 2026

View reviewed changes

vchuravy commented Jul 6, 2026

View reviewed changes

vchuravy added this to the 0.10.0 milestone Sep 7, 2026

vchuravy and others added 2 commits September 7, 2026 21:16

vchuravy mentioned this pull request Sep 7, 2026

[pocl] Enable float add and min/max atomics via cl_ext_float_atomics #737

Draft

vchuravy and others added 3 commits September 7, 2026 21:23

Address review: run UnsafeAtomics tests on all backends; test orderin…

a55b144

…gs, fences, and syncscopes Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

Test float atomics now that POCL permits the SPIR-V float atomic exte…

f7698e2

…nsions Run the Atomix and UnsafeAtomics add and min/max tests for Float32 and, where the backend supports it, Float64 in addition to the integer types. Assisted-by: Claude Code (Fable 5.1)

vchuravy force-pushed the vc/atomics-tests branch from 291db15 to f7698e2 Compare September 7, 2026 19:25

Navigation Menu

Uh oh!

Add Atomix and UnsafeAtomics atomics tests - #726

Add Atomix and UnsafeAtomics atomics tests#726
vchuravy wants to merge 5 commits into
mainfrom
vc/atomics-tests

vchuravy commented Jul 5, 2026 •

edited

Loading

Uh oh!

vchuravy Jul 5, 2026

Uh oh!

vchuravy Jul 9, 2026

Uh oh!

vchuravy left a comment

Uh oh!

github-actions Bot commented Jul 5, 2026 •

edited

Loading

Uh oh!

vchuravy Jul 6, 2026

Uh oh!

vchuravy Jul 9, 2026

Uh oh!

simeonschaub commented Jul 6, 2026

Uh oh!

vchuravy commented Jul 9, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Uh oh!

Conversation

vchuravy commented Jul 5, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Atomix (all backends, gated on supports_atomics(backend()))

UnsafeAtomics (CPU backend only, @test_skip elsewhere)

Uh oh!

vchuravy Jul 5, 2026

Choose a reason for hiding this comment

Uh oh!

vchuravy Jul 9, 2026

Choose a reason for hiding this comment

Uh oh!

vchuravy left a comment

Choose a reason for hiding this comment

Uh oh!

github-actions Bot commented Jul 5, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Benchmark Results

Benchmark Plots

Uh oh!

vchuravy Jul 6, 2026

Choose a reason for hiding this comment

Uh oh!

vchuravy Jul 9, 2026

Choose a reason for hiding this comment

Uh oh!

simeonschaub commented Jul 6, 2026

Uh oh!

vchuravy commented Jul 9, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

vchuravy commented Jul 5, 2026 •

edited

Loading

Atomix (all backends, gated on `supports_atomics(backend())`)

UnsafeAtomics (CPU backend only, `@test_skip` elsewhere)

github-actions Bot commented Jul 5, 2026 •

edited

Loading