Fill GPUCompiler's SPIR-V atomics descriptor from the device

[maleadt]

GPUCompiler 2.13 lowers atomics for SPIR-V itself and only selects floating-point addition instructions where SPIRVCompilerTarget.atomics allows them, adding the SPIR-V extensions they need. This fills that descriptor from ZE_extension_float_atomics (per precision, global and local memory), ignoring half- and double-precision bits when the device lacks that arithmetic (an Iris Xe reports fp64 atomics without fp64 support). 64-bit integer atomics come from ZE_DEVICE_MODULE_FLAG_INT64_ATOMICS. Without the extension, every float add uses compare-and-swap. SPV_EXT_shader_atomic_float_add is no longer declared for every device. Requires GPUCompiler 2.13 and SPIRVIntrinsics 1.4.

The descriptor can be overridden with the new atomics compiler keyword (@oneapi atomics=oneAPI.SPIRVAtomics(...), zefunction, code reflection), which replaces the device-derived one. Enabling operations the device or toolchain doesn't support can fail compilation or terminate the process: on Xe-LP, IGC rejects a half-precision atomic add and exits.

GPUCompiler implements 8- and 16-bit atomics on the containing aligned 32-bit word. So that word is always part of the allocation, device, shared and host allocations are now padded to a multiple of 4 bytes and aligned to at least 4 bytes; array dimensions are unchanged. The docs state the requirements for memory oneAPI.jl didn't allocate, and that adjacent elements must not be modified with plain stores while such an atomic may be in flight.

Also un-skips the low-level Float32 atomic tests on integrated GPUs, and adds tests for the atomics keyword and for odd-sized Int8/UInt16/Float16 arrays under neighbour contention.

Tested on an Intel Iris Xe (Xe-LP) only, NEO 26.18.38308, Julia 1.12.7, GPUCompiler 2.13.0 and a local SPIRVIntrinsics 1.4.0, with both the LLVM SPIR-V back-end and ONEAPI_LTS=1 (translator): device/intrinsics, kernelabstractions and level-zero pass (device/intrinsics again after the keyword change). IGC dumps show OpAtomicFAddEXT for Float32 add per the descriptor (Xe-LP itself compiles it to a compare-and-swap loop) and a 32-bit compare-and-swap for Float16. Not tested on Arc, Data Center GPU Max, or the actual LTS driver stack.

CI can't resolve until SPIRVIntrinsics 1.4.0 is registered (JuliaGPU/OpenCL.jl#542).

[github-actions]

Your PR requires formatting changes to meet the project's style guidelines.
Please consider running Runic (git runic main) to apply these changes.

Click here to view the suggested changes.

diff --git a/lib/level-zero/device.jl b/lib/level-zero/device.jl
index 8b1e9f7..b14b8a8 100644
--- a/lib/level-zero/device.jl
+++ b/lib/level-zero/device.jl
@@ -119,9 +119,9 @@ function float_atomic_properties(dev::ZeDevice)
 
     props = atomic_props_ref[]
     return (
-        fp16flags=props.fp16Flags,
-        fp32flags=props.fp32Flags,
-        fp64flags=props.fp64Flags,
+        fp16flags = props.fp16Flags,
+        fp32flags = props.fp32Flags,
+        fp64flags = props.fp64Flags,
     )
 end
 
diff --git a/src/compiler/compilation.jl b/src/compiler/compilation.jl
index 33b5537..f5d8418 100644
--- a/src/compiler/compilation.jl
+++ b/src/compiler/compilation.jl
@@ -268,7 +268,7 @@ end
 # implemented with compare-and-swap loops.
 function _atomics_capabilities(dev, properties; supports_fp16, supports_fp64)
     int64 = properties.flags & oneL0.ZE_DEVICE_MODULE_FLAG_INT64_ATOMICS ==
-            oneL0.ZE_DEVICE_MODULE_FLAG_INT64_ATOMICS
+        oneL0.ZE_DEVICE_MODULE_FLAG_INT64_ATOMICS
 
     fp_atomics = oneL0.float_atomic_properties(dev)
     fp_atomics === nothing && return SPIRVAtomics(; int64)
@@ -282,11 +282,11 @@ function _atomics_capabilities(dev, properties; supports_fp16, supports_fp64)
     return SPIRVAtomics(;
         int64,
         fadd_f16_global = supports_fp16 && has_add(fp_atomics.fp16flags, global_add),
-        fadd_f16_local  = supports_fp16 && has_add(fp_atomics.fp16flags, local_add),
+        fadd_f16_local = supports_fp16 && has_add(fp_atomics.fp16flags, local_add),
         fadd_f32_global = has_add(fp_atomics.fp32flags, global_add),
-        fadd_f32_local  = has_add(fp_atomics.fp32flags, local_add),
+        fadd_f32_local = has_add(fp_atomics.fp32flags, local_add),
         fadd_f64_global = supports_fp64 && has_add(fp_atomics.fp64flags, global_add),
-        fadd_f64_local  = supports_fp64 && has_add(fp_atomics.fp64flags, local_add),
+        fadd_f64_local = supports_fp64 && has_add(fp_atomics.fp64flags, local_add),
     )
 end
 
@@ -326,8 +326,9 @@ end
     atomics = _atomics_capabilities(dev, properties; supports_fp16, supports_fp64)
 
     # create GPUCompiler objects
-    target = SPIRVCompilerTarget(; backend, extensions = extensions_str, atomics,
-                                   supports_fp16, supports_fp64, supports_bfloat16,
+    target = SPIRVCompilerTarget(;
+        backend, extensions = extensions_str, atomics,
+        supports_fp16, supports_fp64, supports_bfloat16,
                                    driver = :intel, kwargs...)
     params = oneAPICompilerParams()
     CompilerConfig(target, params; kernel, name, always_inline)
diff --git a/test/device/intrinsics.jl b/test/device/intrinsics.jl
index 6130a87..ecf4436 100644
--- a/test/device/intrinsics.jl
+++ b/test/device/intrinsics.jl
@@ -278,7 +278,7 @@ end
 
 # @testset "atomics (low level)" begin
 
-    @testset "atomic_add($T)" for T in [Int32, UInt32, Float32]
+@testset "atomic_add($T)" for T in [Int32, UInt32, Float32]
         a = oneArray([zero(T)])
 
         function kernel(a, b)
@@ -290,7 +290,7 @@ end
         @test Array(a)[1] == T(256)
     end
 
-    @testset "atomic_sub($T)" for T in [Int32, UInt32, Float32]
+@testset "atomic_sub($T)" for T in [Int32, UInt32, Float32]
         a = oneArray([T(256)])
 
         function kernel(a, b)
@@ -326,7 +326,7 @@ end
         @test Array(a)[1] == T(0)
     end
 
-    @testset "atomic_min($T)" for T in [Int32, UInt32, Float32]
+@testset "atomic_min($T)" for T in [Int32, UInt32, Float32]
         a = oneArray([T(256)])
 
         function kernel(a, T)
@@ -339,7 +339,7 @@ end
         @test Array(a)[1] == one(T)
     end
 
-    @testset "atomic_max($T)" for T in [Int32, UInt32, Float32]
+@testset "atomic_max($T)" for T in [Int32, UInt32, Float32]
         a = oneArray([zero(T)])
 
         function kernel(a, T)
@@ -406,7 +406,7 @@ end
         @test Array(a)[1] == zero(T)
     end
 
-    @testset "atomic_xchg($T)" for T in [Int32, UInt32, Float32]
+@testset "atomic_xchg($T)" for T in [Int32, UInt32, Float32]
         a = oneArray([zero(T)])
 
         function kernel(a, b)
@@ -428,20 +428,20 @@ end
     a = oneArray(Float32[0])
     tt = Tuple{typeof(oneAPI.kernel_convert(a)), Float32}
 
-    spirv = sprint(io -> oneAPI.code_spirv(io, kernel, tt; kernel=true))
+    spirv = sprint(io -> oneAPI.code_spirv(io, kernel, tt; kernel = true))
     fp_atomics = oneL0.float_atomic_properties(device())
     if fp_atomics !== nothing &&
-       fp_atomics.fp32flags & oneL0.ZE_DEVICE_FP_ATOMIC_EXT_FLAG_GLOBAL_ADD != 0
+            fp_atomics.fp32flags & oneL0.ZE_DEVICE_FP_ATOMIC_EXT_FLAG_GLOBAL_ADD != 0
         @test occursin("OpAtomicFAddEXT", spirv)
     end
 
     # without the capability, a compare-and-swap loop is used
     atomics = oneAPI.SPIRVAtomics()
-    spirv = sprint(io -> oneAPI.code_spirv(io, kernel, tt; kernel=true, atomics))
+    spirv = sprint(io -> oneAPI.code_spirv(io, kernel, tt; kernel = true, atomics))
     @test !occursin("OpAtomicFAddEXT", spirv)
     @test occursin("OpAtomicCompareExchange", spirv)
 
-    @oneapi items=256 atomics=atomics kernel(a, 1f0)
+    @oneapi items = 256 atomics = atomics kernel(a, 1.0f0)
     @test Array(a)[1] == 256
 end
 
@@ -450,8 +450,10 @@ end
     # and the last element of an odd-sized array needs the allocation to be padded
     function kernel(a, x)
         i = (get_global_id() - 1) % length(a) + 1
-        UnsafeAtomics.modify!(pointer(a, i), +, x, UnsafeAtomics.monotonic,
-                              UnsafeAtomics.device)
+        UnsafeAtomics.modify!(
+            pointer(a, i), +, x, UnsafeAtomics.monotonic,
+            UnsafeAtomics.device
+        )
         return
     end
 
@@ -463,7 +465,7 @@ end
         @test sizeof(a.data[]) % 4 == 0
 
         # 1020 work-items, so that every element is incremented as often
-        @oneapi items=204 groups=5 kernel(a, one(T))
+        @oneapi items = 204 groups = 5 kernel(a, one(T))
         k = 1020 ÷ n
         @test Array(a) == fill(T <: Integer ? k % T : T(k), n)
     end
diff --git a/test/level-zero.jl b/test/level-zero.jl
index c210021..1f2bdfd 100644
--- a/test/level-zero.jl
+++ b/test/level-zero.jl
@@ -47,7 +47,7 @@ show(devnull, MIME("text/plain"), dev)
 properties(dev)
 compute_properties(dev)
 module_properties(dev)
-float_atomic_properties(dev)
+    float_atomic_properties(dev)
 memory_properties(dev)
 memory_access_properties(dev)
 cache_properties(dev)

[codecov]

Codecov Report

✅ All modified and coverable lines are covered by tests.
✅ Project coverage is 79.08%. Comparing base (89abc90) to head (333853b).
⚠️ Report is 1 commits behind head on main.

Additional details and impacted files

@@            Coverage Diff             @@
##             main     #668      +/-   ##
==========================================
+ Coverage   79.00%   79.08%   +0.07%     
==========================================
  Files          56       56              
  Lines        4063     4088      +25     
==========================================
+ Hits         3210     3233      +23     
- Misses        853      855       +2     

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