Add Atomix and UnsafeAtomics atomics tests

[vchuravy]

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
• @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
• max!/min!
• store!/modify!/cas!/xchg!/load sequence
• add! with explicit seq_cst ordering

Notes:

• 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).
• 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]

Suggested change

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

[vchuravy]

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]

Benchmark Results

Show table

	main	291db15...	main / 291db15...
saxpy/default/Float32/1024	0.0681 ± 0.024 ms	0.0638 ± 0.029 ms	1.07 ± 0.62
saxpy/default/Float32/1048576	0.489 ± 0.017 ms	0.488 ± 0.019 ms	1 ± 0.052
saxpy/default/Float32/16384	0.0533 ± 0.031 ms	0.0562 ± 0.03 ms	0.949 ± 0.75
saxpy/default/Float32/2048	0.0602 ± 0.028 ms	0.063 ± 0.027 ms	0.955 ± 0.6
saxpy/default/Float32/256	0.0653 ± 0.031 ms	0.0613 ± 0.033 ms	1.07 ± 0.77
saxpy/default/Float32/262144	0.177 ± 0.025 ms	0.175 ± 0.025 ms	1.01 ± 0.21
saxpy/default/Float32/32768	0.0705 ± 0.032 ms	0.0756 ± 0.031 ms	0.932 ± 0.57
saxpy/default/Float32/4096	0.066 ± 0.026 ms	0.0667 ± 0.023 ms	0.989 ± 0.52
saxpy/default/Float32/512	0.063 ± 0.034 ms	0.0626 ± 0.031 ms	1.01 ± 0.73
saxpy/default/Float32/64	0.0638 ± 0.033 ms	0.0632 ± 0.029 ms	1.01 ± 0.7
saxpy/default/Float32/65536	0.0905 ± 0.031 ms	0.0923 ± 0.03 ms	0.98 ± 0.46
saxpy/default/Float64/1024	0.0652 ± 0.027 ms	0.0636 ± 0.029 ms	1.03 ± 0.64
saxpy/default/Float64/1048576	0.591 ± 0.069 ms	0.602 ± 0.075 ms	0.982 ± 0.17
saxpy/default/Float64/16384	0.0712 ± 0.031 ms	0.0732 ± 0.03 ms	0.972 ± 0.58
saxpy/default/Float64/2048	0.0603 ± 0.028 ms	0.0629 ± 0.026 ms	0.958 ± 0.59
saxpy/default/Float64/256	0.0627 ± 0.034 ms	0.0627 ± 0.031 ms	0.999 ± 0.74
saxpy/default/Float64/262144	0.183 ± 0.023 ms	0.182 ± 0.023 ms	1.01 ± 0.18
saxpy/default/Float64/32768	0.0806 ± 0.031 ms	0.0813 ± 0.03 ms	0.991 ± 0.53
saxpy/default/Float64/4096	0.0667 ± 0.025 ms	0.067 ± 0.023 ms	0.996 ± 0.51
saxpy/default/Float64/512	0.064 ± 0.033 ms	0.0616 ± 0.032 ms	1.04 ± 0.75
saxpy/default/Float64/64	0.0628 ± 0.034 ms	0.0621 ± 0.032 ms	1.01 ± 0.76
saxpy/default/Float64/65536	0.0978 ± 0.029 ms	0.099 ± 0.028 ms	0.988 ± 0.41
saxpy/static workgroup=(1024,)/Float32/1024	0.0662 ± 0.027 ms	0.0639 ± 0.029 ms	1.04 ± 0.63
saxpy/static workgroup=(1024,)/Float32/1048576	0.459 ± 0.024 ms	0.458 ± 0.022 ms	1 ± 0.072
saxpy/static workgroup=(1024,)/Float32/16384	0.0544 ± 0.03 ms	0.0584 ± 0.029 ms	0.932 ± 0.69
saxpy/static workgroup=(1024,)/Float32/2048	0.059 ± 0.027 ms	0.0612 ± 0.027 ms	0.965 ± 0.61
saxpy/static workgroup=(1024,)/Float32/256	0.0655 ± 0.031 ms	0.0628 ± 0.03 ms	1.04 ± 0.7
saxpy/static workgroup=(1024,)/Float32/262144	0.169 ± 0.03 ms	0.167 ± 0.029 ms	1.02 ± 0.25
saxpy/static workgroup=(1024,)/Float32/32768	0.0648 ± 0.03 ms	0.0659 ± 0.03 ms	0.982 ± 0.64
saxpy/static workgroup=(1024,)/Float32/4096	0.0636 ± 0.027 ms	0.0642 ± 0.023 ms	0.991 ± 0.55
saxpy/static workgroup=(1024,)/Float32/512	0.0668 ± 0.028 ms	0.0651 ± 0.029 ms	1.03 ± 0.63
saxpy/static workgroup=(1024,)/Float32/64	0.0637 ± 0.031 ms	0.0652 ± 0.028 ms	0.978 ± 0.64
saxpy/static workgroup=(1024,)/Float32/65536	0.0874 ± 0.031 ms	0.0849 ± 0.03 ms	1.03 ± 0.51
saxpy/static workgroup=(1024,)/Float64/1024	0.0621 ± 0.029 ms	0.0641 ± 0.028 ms	0.969 ± 0.62
saxpy/static workgroup=(1024,)/Float64/1048576	0.485 ± 0.02 ms	0.484 ± 0.019 ms	1 ± 0.057
saxpy/static workgroup=(1024,)/Float64/16384	0.0686 ± 0.029 ms	0.0714 ± 0.029 ms	0.96 ± 0.56
saxpy/static workgroup=(1024,)/Float64/2048	0.0597 ± 0.028 ms	0.0604 ± 0.027 ms	0.989 ± 0.63
saxpy/static workgroup=(1024,)/Float64/256	0.0611 ± 0.033 ms	0.0608 ± 0.033 ms	1 ± 0.76
saxpy/static workgroup=(1024,)/Float64/262144	0.173 ± 0.03 ms	0.172 ± 0.026 ms	1.01 ± 0.23
saxpy/static workgroup=(1024,)/Float64/32768	0.0741 ± 0.03 ms	0.0771 ± 0.03 ms	0.961 ± 0.54
saxpy/static workgroup=(1024,)/Float64/4096	0.0656 ± 0.025 ms	0.0646 ± 0.023 ms	1.02 ± 0.53
saxpy/static workgroup=(1024,)/Float64/512	0.0607 ± 0.031 ms	0.0633 ± 0.03 ms	0.958 ± 0.67
saxpy/static workgroup=(1024,)/Float64/64	0.0568 ± 0.034 ms	0.0586 ± 0.035 ms	0.97 ± 0.81
saxpy/static workgroup=(1024,)/Float64/65536	0.094 ± 0.03 ms	0.0915 ± 0.03 ms	1.03 ± 0.47
time_to_load	1.11 ± 0.0024 s	1.13 ± 0.0012 s	0.986 ± 0.0023

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]

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

[vchuravy]

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

[simeonschaub]

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]

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.