tint/spirv: Fix atomicCompareExchangeWeak

We were missing an `UnwrapRef` when generating the return type, and
were generating invalid SPIR-V when the value being stored was a
reference. The auto-generated builtin tests only test with literal
values.

Fixed: tint:1573
Change-Id: If42280b3cc8ad3fba7355d333e02400c6db843fa
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/92144
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Auto-Submit: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: James Price <jrprice@google.com>

src/tint/writer/spirv/builder.cc

@@ -3161,7 +3161,7 @@ bool Builder::GenerateAtomicBuiltin(const sem::Call* call,
                 return false;
             }
 
-            auto* value_sem_type = TypeOf(call->Arguments()[2]->Declaration());
+            auto* value_sem_type = call->Target()->Signature().parameters[2]->Type();
 
             auto value_type = GenerateTypeIfNeeded(value_sem_type);
             if (value_type == 0) {

test/tint/bug/tint/1573.wgsl

@@ -0,0 +1,8 @@
+@group(0) @binding(0)
+var<storage, read_write> a : atomic<u32>;
+
+@stage(compute) @workgroup_size(16)
+fn main() {
+  var value = 42u;
+  let result = atomicCompareExchangeWeak(&a, 0u, value);
+}

test/tint/bug/tint/1573.wgsl.expected.glsl

@@ -0,0 +1,28 @@
+#version 310 es
+
+struct atomic_compare_exchange_resultu32 {
+  uint old_value;
+  bool exchanged;
+};
+
+
+struct a_block {
+  uint inner;
+};
+
+layout(binding = 0, std430) buffer a_block_1 {
+  uint inner;
+} a;
+void tint_symbol() {
+  uint value = 42u;
+  atomic_compare_exchange_resultu32 atomic_compare_result;
+  atomic_compare_result.old_value = atomicCompSwap(a.inner, 0u, value);
+  atomic_compare_result.exchanged = atomic_compare_result.old_value == 0u;
+  atomic_compare_exchange_resultu32 result = atomic_compare_result;
+}
+
+layout(local_size_x = 16, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}

test/tint/bug/tint/1573.wgsl.expected.hlsl

@@ -0,0 +1,21 @@
+RWByteAddressBuffer a : register(u0, space0);
+
+struct atomic_compare_exchange_weak_ret_type {
+  uint old_value;
+  bool exchanged;
+};
+
+atomic_compare_exchange_weak_ret_type tint_atomicCompareExchangeWeak(RWByteAddressBuffer buffer, uint offset, uint compare, uint value) {
+  atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0;
+  buffer.InterlockedCompareExchange(offset, compare, value, result.old_value);
+  result.exchanged = result.old_value == compare;
+  return result;
+}
+
+
+[numthreads(16, 1, 1)]
+void main() {
+  uint value = 42u;
+  const atomic_compare_exchange_weak_ret_type result = tint_atomicCompareExchangeWeak(a, 0u, 0u, value);
+  return;
+}

test/tint/bug/tint/1573.wgsl.expected.msl

@@ -0,0 +1,21 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+struct atomic_compare_exchange_resultu32 {
+  uint old_value;
+  bool exchanged;
+};
+template <typename A, typename T>
+atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(device A* atomic, T compare, T value) {
+  T old_value = compare;
+  bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
+  return {old_value, exchanged};
+}
+
+kernel void tint_symbol(device atomic_uint* tint_symbol_1 [[buffer(0)]]) {
+  uint value = 42u;
+  atomic_compare_exchange_resultu32 const result = atomicCompareExchangeWeak_1(tint_symbol_1, 0u, value);
+  return;
+}
+

test/tint/bug/tint/1573.wgsl.expected.spvasm

@@ -0,0 +1,48 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 24
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 16 1 1
+               OpName %a_block "a_block"
+               OpMemberName %a_block 0 "inner"
+               OpName %a "a"
+               OpName %main "main"
+               OpName %value "value"
+               OpName %__atomic_compare_exchange_resultu32 "__atomic_compare_exchange_resultu32"
+               OpMemberName %__atomic_compare_exchange_resultu32 0 "old_value"
+               OpMemberName %__atomic_compare_exchange_resultu32 1 "exchanged"
+               OpDecorate %a_block Block
+               OpMemberDecorate %a_block 0 Offset 0
+               OpDecorate %a DescriptorSet 0
+               OpDecorate %a Binding 0
+               OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0
+               OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4
+       %uint = OpTypeInt 32 0
+    %a_block = OpTypeStruct %uint
+%_ptr_StorageBuffer_a_block = OpTypePointer StorageBuffer %a_block
+          %a = OpVariable %_ptr_StorageBuffer_a_block StorageBuffer
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+    %uint_42 = OpConstant %uint 42
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %12 = OpConstantNull %uint
+       %bool = OpTypeBool
+%__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool
+     %uint_1 = OpConstant %uint 1
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+       %main = OpFunction %void None %5
+          %8 = OpLabel
+      %value = OpVariable %_ptr_Function_uint Function %12
+               OpStore %value %uint_42
+         %20 = OpAccessChain %_ptr_StorageBuffer_uint %a %uint_0
+         %21 = OpLoad %uint %value
+         %22 = OpAtomicCompareExchange %uint %20 %uint_1 %uint_0 %uint_0 %21 %uint_0
+         %23 = OpIEqual %bool %22 %21
+         %13 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %22 %23
+               OpReturn
+               OpFunctionEnd

test/tint/bug/tint/1573.wgsl.expected.wgsl

@@ -0,0 +1,7 @@
+@group(0) @binding(0) var<storage, read_write> a : atomic<u32>;
+
+@stage(compute) @workgroup_size(16)
+fn main() {
+  var value = 42u;
+  let result = atomicCompareExchangeWeak(&(a), 0u, value);
+}