Changelog and more tests

Author: cryvosh

CHANGELOG.md

@@ -57,6 +57,9 @@ Bottom level categories:
 - `DisplayHandle` should now be passed to `InstanceDescriptor` for correct EGL initialization on Wayland. By @MarijnS95 in [#8012](https://github.com/gfx-rs/wgpu/pull/8012)
   Note that the existing workaround to create surfaces before the adapter is no longer valid.
 
+#### naga
+- Fixed `workgroupUniformLoad` incorrectly returning an atomic when called on an atomic, it now returns the inner `T` as per the spec. By @cryvosh in [#8791](https://github.com/gfx-rs/wgpu/pull/8791).
+
 ### Documentation
 
 #### General

naga/tests/in/wgsl/workgroup-uniform-load-atomic.wgsl

@@ -1,8 +1,13 @@
 // Test workgroupUniformLoad specialization for atomic<T> -> T
-// Issue: https://github.com/gfx-rs/wgpu/issues/8785
+
+struct AtomicStruct {
+    atomic_scalar: atomic<u32>,
+    atomic_arr: array<atomic<i32>, 2>,
+}
 
 var<workgroup> wg_scalar: atomic<u32>;
 var<workgroup> wg_signed: atomic<i32>;
+var<workgroup> wg_struct: AtomicStruct;
 
 @compute @workgroup_size(64)
 fn test_atomic_workgroup_uniform_load(
@@ -11,9 +16,11 @@ fn test_atomic_workgroup_uniform_load(
 ) {
     let active_tile_index = workgroup_id.x + workgroup_id.y * 32768;
 
-    // Each thread may set the atomic
+    // Each thread may set the atomics
     atomicOr(&wg_scalar, u32(active_tile_index >= 64));
     atomicAdd(&wg_signed, 1i);
+    atomicStore(&wg_struct.atomic_scalar, 1u);
+    atomicAdd(&wg_struct.atomic_arr[0], 1i);
 
     workgroupBarrier();
 
@@ -23,8 +30,14 @@ fn test_atomic_workgroup_uniform_load(
     // workgroupUniformLoad on atomic<i32> should return i32
     let signed_val: i32 = workgroupUniformLoad(&wg_signed);
 
-    // Should be able to use the result in comparisons
-    if scalar_val == 0u && signed_val > 0i {
+    // workgroupUniformLoad on struct.atomic_scalar should return u32
+    let struct_scalar: u32 = workgroupUniformLoad(&wg_struct.atomic_scalar);
+    
+    // workgroupUniformLoad on struct.atomic_arr[i] should return i32
+    let struct_arr_val: i32 = workgroupUniformLoad(&wg_struct.atomic_arr[0]);
+    
+    // Should be able to use all results in comparisons
+    if scalar_val == 0u && signed_val > 0i && struct_scalar > 0u && struct_arr_val > 0i {
         return;
     }
 }

naga/tests/out/glsl/wgsl-workgroup-uniform-load-atomic.test_atomic_workgroup_uniform_load.Compute.glsl

@@ -5,43 +5,76 @@ precision highp int;
 
 layout(local_size_x = 64, local_size_y = 1, local_size_z = 1) in;
 
+struct AtomicStruct {
+    uint atomic_scalar;
+    int atomic_arr[2];
+};
 shared uint wg_scalar;
 
 shared int wg_signed;
 
+shared AtomicStruct wg_struct;
+
 
 void main() {
     if (gl_LocalInvocationID == uvec3(0u)) {
         wg_scalar = 0u;
         wg_signed = 0;
+        wg_struct = AtomicStruct(0u, int[2](0, 0));
     }
     memoryBarrierShared();
     barrier();
     uvec3 workgroup_id = gl_WorkGroupID;
     uvec3 local_id = gl_LocalInvocationID;
     bool local = false;
+    bool local_1 = false;
+    bool local_2 = false;
     uint active_tile_index = (workgroup_id.x + (workgroup_id.y * 32768u));
     uint _e11 = atomicOr(wg_scalar, uint((active_tile_index >= 64u)));
     int _e14 = atomicAdd(wg_signed, 1);
+    wg_struct.atomic_scalar = 1u;
+    int _e22 = atomicAdd(wg_struct.atomic_arr[0], 1);
+    memoryBarrierShared();
+    barrier();
+    memoryBarrierShared();
+    barrier();
+    uint _e24 = wg_scalar;
+    memoryBarrierShared();
+    barrier();
     memoryBarrierShared();
     barrier();
+    int _e26 = wg_signed;
     memoryBarrierShared();
     barrier();
-    uint _e16 = wg_scalar;
     memoryBarrierShared();
     barrier();
+    uint _e29 = wg_struct.atomic_scalar;
     memoryBarrierShared();
     barrier();
-    int _e18 = wg_signed;
     memoryBarrierShared();
     barrier();
-    if ((_e16 == 0u)) {
-        local = (_e18 > 0);
+    int _e33 = wg_struct.atomic_arr[0];
+    memoryBarrierShared();
+    barrier();
+    if ((_e24 == 0u)) {
+        local = (_e26 > 0);
     } else {
         local = false;
     }
-    bool _e26 = local;
-    if (_e26) {
+    bool _e41 = local;
+    if (_e41) {
+        local_1 = (_e29 > 0u);
+    } else {
+        local_1 = false;
+    }
+    bool _e47 = local_1;
+    if (_e47) {
+        local_2 = (_e33 > 0);
+    } else {
+        local_2 = false;
+    }
+    bool _e53 = local_2;
+    if (_e53) {
         return;
     } else {
         return;

naga/tests/out/hlsl/wgsl-workgroup-uniform-load-atomic.hlsl

@@ -1,33 +1,62 @@
+struct AtomicStruct {
+    uint atomic_scalar;
+    int atomic_arr[2];
+};
+
 groupshared uint wg_scalar;
 groupshared int wg_signed;
+groupshared AtomicStruct wg_struct;
 
 [numthreads(64, 1, 1)]
 void test_atomic_workgroup_uniform_load(uint3 workgroup_id : SV_GroupID, uint3 local_id : SV_GroupThreadID, uint3 __local_invocation_id : SV_GroupThreadID)
 {
     if (all(__local_invocation_id == uint3(0u, 0u, 0u))) {
         wg_scalar = (uint)0;
         wg_signed = (int)0;
+        wg_struct = (AtomicStruct)0;
     }
     GroupMemoryBarrierWithGroupSync();
     bool local = (bool)0;
+    bool local_1 = (bool)0;
+    bool local_2 = (bool)0;
 
     uint active_tile_index = (workgroup_id.x + (workgroup_id.y * 32768u));
     uint _e11; InterlockedOr(wg_scalar, uint((active_tile_index >= 64u)), _e11);
     int _e14; InterlockedAdd(wg_signed, int(1), _e14);
+    wg_struct.atomic_scalar = 1u;
+    int _e22; InterlockedAdd(wg_struct.atomic_arr[0], int(1), _e22);
+    GroupMemoryBarrierWithGroupSync();
+    GroupMemoryBarrierWithGroupSync();
+    uint _e24 = wg_scalar;
+    GroupMemoryBarrierWithGroupSync();
+    GroupMemoryBarrierWithGroupSync();
+    int _e26 = wg_signed;
     GroupMemoryBarrierWithGroupSync();
     GroupMemoryBarrierWithGroupSync();
-    uint _e16 = wg_scalar;
+    uint _e29 = wg_struct.atomic_scalar;
     GroupMemoryBarrierWithGroupSync();
     GroupMemoryBarrierWithGroupSync();
-    int _e18 = wg_signed;
+    int _e33 = wg_struct.atomic_arr[0];
     GroupMemoryBarrierWithGroupSync();
-    if ((_e16 == 0u)) {
-        local = (_e18 > int(0));
+    if ((_e24 == 0u)) {
+        local = (_e26 > int(0));
     } else {
         local = false;
     }
-    bool _e26 = local;
-    if (_e26) {
+    bool _e41 = local;
+    if (_e41) {
+        local_1 = (_e29 > 0u);
+    } else {
+        local_1 = false;
+    }
+    bool _e47 = local_1;
+    if (_e47) {
+        local_2 = (_e33 > int(0));
+    } else {
+        local_2 = false;
+    }
+    bool _e53 = local_2;
+    if (_e53) {
         return;
     } else {
         return;

naga/tests/out/msl/wgsl-workgroup-uniform-load-atomic.msl

@@ -4,6 +4,13 @@
 
 using metal::uint;
 
+struct type_2 {
+    metal::atomic_int inner[2];
+};
+struct AtomicStruct {
+    metal::atomic_uint atomic_scalar;
+    type_2 atomic_arr;
+};
 
 struct test_atomic_workgroup_uniform_loadInput {
 };
@@ -12,30 +19,57 @@ kernel void test_atomic_workgroup_uniform_load(
 , metal::uint3 local_id [[thread_position_in_threadgroup]]
 , threadgroup metal::atomic_uint& wg_scalar
 , threadgroup metal::atomic_int& wg_signed
+, threadgroup AtomicStruct& wg_struct
 ) {
     if (metal::all(local_id == metal::uint3(0u))) {
         metal::atomic_store_explicit(&wg_scalar, 0, metal::memory_order_relaxed);
         metal::atomic_store_explicit(&wg_signed, 0, metal::memory_order_relaxed);
+        metal::atomic_store_explicit(&wg_struct.atomic_scalar, 0, metal::memory_order_relaxed);
+        for (int __i0 = 0; __i0 < 2; __i0++) {
+            metal::atomic_store_explicit(&wg_struct.atomic_arr.inner[__i0], 0, metal::memory_order_relaxed);
+        }
     }
     metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
     bool local = {};
+    bool local_1 = {};
+    bool local_2 = {};
     uint active_tile_index = workgroup_id.x + (workgroup_id.y * 32768u);
     uint _e11 = metal::atomic_fetch_or_explicit(&wg_scalar, static_cast<uint>(active_tile_index >= 64u), metal::memory_order_relaxed);
     int _e14 = metal::atomic_fetch_add_explicit(&wg_signed, 1, metal::memory_order_relaxed);
+    metal::atomic_store_explicit(&wg_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
+    int _e22 = metal::atomic_fetch_add_explicit(&wg_struct.atomic_arr.inner[0], 1, metal::memory_order_relaxed);
     metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
     metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
     uint unnamed = metal::atomic_load_explicit(&wg_scalar, metal::memory_order_relaxed);
     metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
     metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
     int unnamed_1 = metal::atomic_load_explicit(&wg_signed, metal::memory_order_relaxed);
     metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
+    metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
+    uint unnamed_2 = metal::atomic_load_explicit(&wg_struct.atomic_scalar, metal::memory_order_relaxed);
+    metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
+    metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
+    int unnamed_3 = metal::atomic_load_explicit(&wg_struct.atomic_arr.inner[0], metal::memory_order_relaxed);
+    metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
     if (unnamed == 0u) {
         local = unnamed_1 > 0;
     } else {
         local = false;
     }
-    bool _e26 = local;
-    if (_e26) {
+    bool _e41 = local;
+    if (_e41) {
+        local_1 = unnamed_2 > 0u;
+    } else {
+        local_1 = false;
+    }
+    bool _e47 = local_1;
+    if (_e47) {
+        local_2 = unnamed_3 > 0;
+    } else {
+        local_2 = false;
+    }
+    bool _e53 = local_2;
+    if (_e53) {
         return;
     } else {
         return;