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:

Builds on [pocl] Enable float add and min/max atomics via cl_ext_float_atomics #737 (merged), which permits the SPV_EXT_shader_atomic_float_{add,min_max} extensions based on cl_ext_float_atomics, and extends the test/atomics.jl it introduced: add and min/max run for Float32 and, where supports_float64, Float64 through Atomix.
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"]).

CI status

Requires Atomix 1.2.1 (Complete the GPU extensions: atomic get/set, swap, float min/max JuliaConcurrent/Atomix.jl#80 and Create MultiEvent from tuple of empty MultiEvents #81, both released) on the GPU backends: the extensions previously implemented get/set! as error("not implemented") and had no swap. Rebased onto main, which includes [pocl] Enable float add and min/max atomics via cl_ext_float_atomics #737 and the Atomix = "1.2.1" compat.
The raw UnsafeAtomics tests run with integer element types only: atomicrmw fadd/fmin/fmax need SPIR-V extensions that NVIDIA's OpenCL driver and OpenCL.jl's C backend do not provide, and UnsafeAtomics has no fallback. Atomix covers floats. The syncscope test uses the system scope only, since NVPTX rejects single-thread scoped atomics and fences.

🤖 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	`a6a5221`...	main / `a6a5221`...
saxpy/default/Float32/1024	0.0854 ± 0.0075 ms	0.0861 ± 0.0093 ms	0.993 ± 0.14
saxpy/default/Float32/1048576	0.536 ± 0.029 ms	0.518 ± 0.026 ms	1.03 ± 0.077
saxpy/default/Float32/16384	0.0689 ± 0.033 ms	0.0692 ± 0.034 ms	0.995 ± 0.68
saxpy/default/Float32/2048	0.0857 ± 0.025 ms	0.0866 ± 0.023 ms	0.99 ± 0.38
saxpy/default/Float32/256	0.0859 ± 0.012 ms	0.0868 ± 0.011 ms	0.989 ± 0.19
saxpy/default/Float32/262144	0.181 ± 0.031 ms	0.182 ± 0.031 ms	0.999 ± 0.24
saxpy/default/Float32/32768	0.0758 ± 0.032 ms	0.0758 ± 0.032 ms	1 ± 0.59
saxpy/default/Float32/4096	0.0801 ± 0.032 ms	0.0846 ± 0.031 ms	0.946 ± 0.52
saxpy/default/Float32/512	0.085 ± 0.01 ms	0.0861 ± 0.0076 ms	0.987 ± 0.15
saxpy/default/Float32/64	0.0861 ± 0.013 ms	0.0867 ± 0.012 ms	0.993 ± 0.21
saxpy/default/Float32/65536	0.0931 ± 0.031 ms	0.0908 ± 0.031 ms	1.03 ± 0.49
saxpy/default/Float64/1024	0.0856 ± 0.022 ms	0.0855 ± 0.016 ms	1 ± 0.32
saxpy/default/Float64/1048576	0.595 ± 0.098 ms	0.62 ± 0.1 ms	0.959 ± 0.22
saxpy/default/Float64/16384	0.0689 ± 0.031 ms	0.0689 ± 0.03 ms	0.999 ± 0.62
saxpy/default/Float64/2048	0.0861 ± 0.028 ms	0.0857 ± 0.026 ms	1 ± 0.45
saxpy/default/Float64/256	0.086 ± 0.01 ms	0.0866 ± 0.01 ms	0.993 ± 0.17
saxpy/default/Float64/262144	0.192 ± 0.035 ms	0.188 ± 0.036 ms	1.02 ± 0.27
saxpy/default/Float64/32768	0.0798 ± 0.03 ms	0.0793 ± 0.03 ms	1.01 ± 0.54
saxpy/default/Float64/4096	0.0823 ± 0.03 ms	0.0837 ± 0.029 ms	0.984 ± 0.49
saxpy/default/Float64/512	0.0861 ± 0.0096 ms	0.0858 ± 0.012 ms	1 ± 0.18
saxpy/default/Float64/64	0.0865 ± 0.014 ms	0.0865 ± 0.012 ms	1 ± 0.22
saxpy/default/Float64/65536	0.1 ± 0.031 ms	0.0957 ± 0.032 ms	1.05 ± 0.47
saxpy/static workgroup=(1024,)/Float32/1024	0.0827 ± 0.011 ms	0.0834 ± 0.011 ms	0.992 ± 0.19
saxpy/static workgroup=(1024,)/Float32/1048576	0.438 ± 0.027 ms	0.438 ± 0.027 ms	1 ± 0.087
saxpy/static workgroup=(1024,)/Float32/16384	0.065 ± 0.031 ms	0.0642 ± 0.032 ms	1.01 ± 0.69
saxpy/static workgroup=(1024,)/Float32/2048	0.0826 ± 0.027 ms	0.0836 ± 0.023 ms	0.988 ± 0.42
saxpy/static workgroup=(1024,)/Float32/256	0.0834 ± 0.014 ms	0.084 ± 0.013 ms	0.993 ± 0.23
saxpy/static workgroup=(1024,)/Float32/262144	0.16 ± 0.031 ms	0.157 ± 0.031 ms	1.02 ± 0.29
saxpy/static workgroup=(1024,)/Float32/32768	0.0692 ± 0.029 ms	0.069 ± 0.03 ms	1 ± 0.6
saxpy/static workgroup=(1024,)/Float32/4096	0.0814 ± 0.03 ms	0.085 ± 0.029 ms	0.958 ± 0.48
saxpy/static workgroup=(1024,)/Float32/512	0.0828 ± 0.012 ms	0.0833 ± 0.01 ms	0.994 ± 0.19
saxpy/static workgroup=(1024,)/Float32/64	0.0836 ± 0.012 ms	0.0841 ± 0.013 ms	0.994 ± 0.21
saxpy/static workgroup=(1024,)/Float32/65536	0.0845 ± 0.031 ms	0.0829 ± 0.031 ms	1.02 ± 0.54
saxpy/static workgroup=(1024,)/Float64/1024	0.0825 ± 0.021 ms	0.0829 ± 0.017 ms	0.996 ± 0.32
saxpy/static workgroup=(1024,)/Float64/1048576	0.529 ± 0.084 ms	0.509 ± 0.085 ms	1.04 ± 0.24
saxpy/static workgroup=(1024,)/Float64/16384	0.0639 ± 0.028 ms	0.067 ± 0.029 ms	0.954 ± 0.58
saxpy/static workgroup=(1024,)/Float64/2048	0.0827 ± 0.029 ms	0.0841 ± 0.024 ms	0.984 ± 0.45
saxpy/static workgroup=(1024,)/Float64/256	0.0837 ± 0.011 ms	0.0833 ± 0.013 ms	1 ± 0.2
saxpy/static workgroup=(1024,)/Float64/262144	0.184 ± 0.032 ms	0.183 ± 0.035 ms	1 ± 0.26
saxpy/static workgroup=(1024,)/Float64/32768	0.0743 ± 0.029 ms	0.075 ± 0.028 ms	0.99 ± 0.53
saxpy/static workgroup=(1024,)/Float64/4096	0.0781 ± 0.029 ms	0.0797 ± 0.03 ms	0.98 ± 0.51
saxpy/static workgroup=(1024,)/Float64/512	0.0835 ± 0.0098 ms	0.0832 ± 0.014 ms	1 ± 0.2
saxpy/static workgroup=(1024,)/Float64/64	0.0837 ± 0.01 ms	0.0838 ± 0.014 ms	0.999 ± 0.21
saxpy/static workgroup=(1024,)/Float64/65536	0.0938 ± 0.029 ms	0.091 ± 0.03 ms	1.03 ± 0.47
time_to_load	0.969 ± 0.015 s	0.969 ± 0.016 s	1 ± 0.023

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.

christiangnrd · 2026-09-08T16:26:23Z

Does this close #308?

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)

UnsafeAtomics emits `atomicrmw fmin`/`fmax` directly and has no fallback where the target lacks the instruction; OpenCL.jl's OpenCL C program backend cannot enable the SPIR-V float min/max extension, so those cases fail there. Atomix's float min/max (with its compare-and-swap fallback) remains tested on every backend. Assisted-by: Claude Code (Fable 5.1)

`UnsafeAtomics.add!` on floats emits `atomicrmw fadd`, which needs the SPV_EXT_shader_atomic_float_add extension that NVIDIA's OpenCL driver does not provide, so the float cases fail on the OpenCL (CUDA) backend like the float min/max ones did on the OpenCL C backend. Raw UnsafeAtomics has no fallback; Atomix covers the float operations portably. Assisted-by: Claude Code (Fable 5.1)

NVPTX refuses to select it: `LLVM ERROR: Unsupported scope "Thread" for seq_cst fence` (Julia 1.12, LLVM 18). The single-thread scoped store and add stay. Assisted-by: Claude Code (Fable 5.1)

NVPTX rejects single-thread scoped atomics outright on LLVM 21 (`LLVM ERROR: Atomics need scope > "Thread"`), so the kernel now exercises the explicit `none` scope for store, fence and add. Assisted-by: Claude Code (Fable 5.1)

On Julia 1.10 and 1.13 UnsafeAtomics emits that fence through inline assembly, which the SPIR-V backend rejects (SPV_INTEL_inline_assembly). Release/acquire fences remain covered by the "fences" test. Assisted-by: Claude Code (Fable 5.1)

- Load UnsafeAtomics through its UUID, like Pkg in testsuite.jl, so backend packages need not add it to their test environments. - Fuse the Atomix max and min kernels into one, mirroring the UnsafeAtomics form; halves the kernel variants compiled per type. - One UnsafeAtomics add kernel exercises the default and every explicit ordering on separate histogram columns, replacing a second kernel and a loop that compiled five variants for the same coverage. - Drop a dead assignment, a redundant type parameter, and stale comments about restrictions that no longer exist. Assisted-by: Claude Code (Fable 5.1)

codecov · 2026-09-09T16:57:10Z

Codecov Report

✅ All modified and coverable lines are covered by tests.
✅ Project coverage is 63.72%. Comparing base (18c697d) to head (a6a5221).

Additional details and impacted files

@@           Coverage Diff           @@
##             main     #726   +/-   ##
=======================================
  Coverage   63.72%   63.72%           
=======================================
  Files          23       23           
  Lines        1935     1935           
=======================================
  Hits         1233     1233           
  Misses        702      702

☔ View full report in Codecov by Harness.
📢 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.

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 mentioned this pull request Sep 7, 2026

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

Merged

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

vchuravy mentioned this pull request Sep 8, 2026

Complete the GPU extensions: atomic get/set, swap, float min/max JuliaConcurrent/Atomix.jl#80

Merged

christiangnrd force-pushed the vc/atomics-tests branch from a6ba637 to 3a92bd8 Compare September 8, 2026 13:15

christiangnrd mentioned this pull request Sep 8, 2026

CI: CUDA and OpenCL(CUDA) jobs on Julia 1.12+ test KernelAbstractions main instead of the PR #753

Closed

vchuravy force-pushed the vc/atomics-tests branch 2 times, most recently from 0452642 to 0e8a63e Compare September 9, 2026 09:49

vchuravy and others added 9 commits September 9, 2026 12:09

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

2d490fa

…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…

cd2a44d

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

Smoke test

4a1a045

Drop the single-thread seq_cst fence from the syncscope test

40e1f26

NVPTX refuses to select it: `LLVM ERROR: Unsupported scope "Thread" for seq_cst fence` (Julia 1.12, LLVM 18). The single-thread scoped store and add stay. Assisted-by: Claude Code (Fable 5.1)

Drop the seq_cst fence from the syncscope test

a644d16

On Julia 1.10 and 1.13 UnsafeAtomics emits that fence through inline assembly, which the SPIR-V backend rejects (SPV_INTEL_inline_assembly). Release/acquire fences remain covered by the "fences" test. Assisted-by: Claude Code (Fable 5.1)

vchuravy force-pushed the vc/atomics-tests branch from 0e8a63e to a644d16 Compare September 9, 2026 10:10

vchuravy commented Sep 9, 2026

View reviewed changes

Comment thread Project.toml Outdated

vchuravy added 2 commits September 9, 2026 12:11

Apply suggestion from @vchuravy

246971e

vchuravy merged commit 901cd0e into main Sep 9, 2026
73 checks passed

vchuravy deleted the vc/atomics-tests branch September 9, 2026 17:14

Navigation Menu

Add Atomix and UnsafeAtomics atomics tests - #726

Add Atomix and UnsafeAtomics atomics tests#726
vchuravy merged 11 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!

christiangnrd commented Sep 8, 2026

Uh oh!

Uh oh!

codecov Bot commented Sep 9, 2026

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

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)

CI status

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!

christiangnrd commented Sep 8, 2026

Uh oh!

Uh oh!

codecov Bot commented Sep 9, 2026

Codecov Report

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 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