fixup! test: add atomicOps shader test for all but hlsl-out

tests/in/atomicOps.wgsl

@@ -119,4 +119,23 @@ fn cs_main(@builtin(local_invocation_id) id: vec3<u32>) {
  atomicXor(&workgroup_atomic_arr[1], 1i);
  atomicXor(&workgroup_struct.atomic_scalar, 1u);
  atomicXor(&workgroup_struct.atomic_arr[1], 1i);
+
+ atomicExchange(&storage_atomic_scalar, 1u);
+ atomicExchange(&storage_atomic_arr[1], 1i);
+ atomicExchange(&storage_struct.atomic_scalar, 1u);
+ atomicExchange(&storage_struct.atomic_arr[1], 1i);
+ atomicExchange(&workgroup_atomic_scalar, 1u);
+ atomicExchange(&workgroup_atomic_arr[1], 1i);
+ atomicExchange(&workgroup_struct.atomic_scalar, 1u);
+ atomicExchange(&workgroup_struct.atomic_arr[1], 1i);
+
+ // // TODO: https://github.com/gpuweb/gpuweb/issues/2021
+ // atomicCompareExchangeWeak(&storage_atomic_scalar, 1u);
+ // atomicCompareExchangeWeak(&storage_atomic_arr[1], 1i);
+ // atomicCompareExchangeWeak(&storage_struct.atomic_scalar, 1u);
+ // atomicCompareExchangeWeak(&storage_struct.atomic_arr[1], 1i);
+ // atomicCompareExchangeWeak(&workgroup_atomic_scalar, 1u);
+ // atomicCompareExchangeWeak(&workgroup_atomic_arr[1], 1i);
+ // atomicCompareExchangeWeak(&workgroup_struct.atomic_scalar, 1u);
+ // atomicCompareExchangeWeak(&workgroup_struct.atomic_arr[1], 1i);
 }

tests/out/glsl/atomicOps.cs_main.Compute.glsl

@@ -119,6 +119,14 @@ void main() {
  int _e298 = atomicXor(workgroup_atomic_arr[1], 1);
  uint _e302 = atomicXor(workgroup_struct.atomic_scalar, 1u);
  int _e308 = atomicXor(workgroup_struct.atomic_arr[1], 1);
+ uint _e311 = atomicExchange(_group_0_binding_0_cs, 1u);
+ int _e316 = atomicExchange(_group_0_binding_1_cs[1], 1);
+ uint _e320 = atomicExchange(_group_0_binding_2_cs.atomic_scalar, 1u);
+ int _e326 = atomicExchange(_group_0_binding_2_cs.atomic_arr[1], 1);
+ uint _e329 = atomicExchange(workgroup_atomic_scalar, 1u);
+ int _e334 = atomicExchange(workgroup_atomic_arr[1], 1);
+ uint _e338 = atomicExchange(workgroup_struct.atomic_scalar, 1u);
+ int _e344 = atomicExchange(workgroup_struct.atomic_arr[1], 1);
  return;
 }
 

tests/out/msl/atomicOps.msl

@@ -114,5 +114,13 @@ kernel void cs_main(
  int _e298 = metal::atomic_fetch_xor_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed);
  uint _e302 = metal::atomic_fetch_xor_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
  int _e308 = metal::atomic_fetch_xor_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed);
+ uint _e311 = metal::atomic_exchange_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed);
+ int _e316 = metal::atomic_exchange_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed);
+ uint _e320 = metal::atomic_exchange_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
+ int _e326 = metal::atomic_exchange_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed);
+ uint _e329 = metal::atomic_exchange_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed);
+ int _e334 = metal::atomic_exchange_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed);
+ uint _e338 = metal::atomic_exchange_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
+ int _e344 = metal::atomic_exchange_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed);
  return;
 }

tests/out/spv/atomicOps.spvasm

@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.1
 ; Generator: rspirv
-; Bound: 175
+; Bound: 189
 OpCapability Shader
 OpExtension "SPV_KHR_storage_buffer_storage_class"
 %1 = OpExtInstImport "GLSL.std.450"
@@ -222,5 +222,19 @@ OpControlBarrier %50 %50 %51
 %171 = OpAtomicXor %6 %172 %3 %58 %5
 %174 = OpAccessChain %59 %24 %5 %5
 %173 = OpAtomicXor %4 %174 %3 %58 %7
+%175 = OpAtomicExchange %6 %34 %7 %53 %5
+%177 = OpAccessChain %54 %36 %5
+%176 = OpAtomicExchange %4 %177 %7 %53 %7
+%179 = OpAccessChain %32 %38 %33
+%178 = OpAtomicExchange %6 %179 %7 %53 %5
+%181 = OpAccessChain %54 %38 %5 %5
+%180 = OpAtomicExchange %4 %181 %7 %53 %7
+%182 = OpAtomicExchange %6 %20 %3 %58 %5
+%184 = OpAccessChain %59 %22 %5
+%183 = OpAtomicExchange %4 %184 %3 %58 %7
+%186 = OpAccessChain %21 %24 %33
+%185 = OpAtomicExchange %6 %186 %3 %58 %5
+%188 = OpAccessChain %59 %24 %5 %5
+%187 = OpAtomicExchange %4 %188 %3 %58 %7
 OpReturn
 OpFunctionEnd

tests/out/wgsl/atomicOps.wgsl

@@ -95,5 +95,13 @@ fn cs_main(@builtin(local_invocation_id) id: vec3<u32>) {
  let _e298 = atomicXor((&workgroup_atomic_arr[1]), 1);
  let _e302 = atomicXor((&workgroup_struct.atomic_scalar), 1u);
  let _e308 = atomicXor((&workgroup_struct.atomic_arr[1]), 1);
+ let _e311 = atomicExchange((&storage_atomic_scalar), 1u);
+ let _e316 = atomicExchange((&storage_atomic_arr[1]), 1);
+ let _e320 = atomicExchange((&storage_struct.atomic_scalar), 1u);
+ let _e326 = atomicExchange((&storage_struct.atomic_arr[1]), 1);
+ let _e329 = atomicExchange((&workgroup_atomic_scalar), 1u);
+ let _e334 = atomicExchange((&workgroup_atomic_arr[1]), 1);
+ let _e338 = atomicExchange((&workgroup_struct.atomic_scalar), 1u);
+ let _e344 = atomicExchange((&workgroup_struct.atomic_arr[1]), 1);
  return;
 }