diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc
index 3dfb74f..de7a0be 100644
--- a/src/tint/writer/hlsl/generator_impl.cc
+++ b/src/tint/writer/hlsl/generator_impl.cc
@@ -1779,8 +1779,8 @@
 
             {  // T compare_value = <compare_value>;
                 auto pre = line();
-                if (!EmitTypeAndName(pre, TypeOf(compare_value), ast::StorageClass::kNone,
-                                     ast::Access::kUndefined, compare)) {
+                if (!EmitTypeAndName(pre, TypeOf(compare_value)->UnwrapRef(),
+                                     ast::StorageClass::kNone, ast::Access::kUndefined, compare)) {
                     return false;
                 }
                 pre << " = ";
diff --git a/src/tint/writer/spirv/builder.cc b/src/tint/writer/spirv/builder.cc
index 4368e3c..c71f0fc 100644
--- a/src/tint/writer/spirv/builder.cc
+++ b/src/tint/writer/spirv/builder.cc
@@ -3272,7 +3272,8 @@
                                                                      value,
                                                                  });
         case sem::BuiltinType::kAtomicCompareExchangeWeak: {
-            auto comparator = GenerateExpression(call->Arguments()[1]->Declaration());
+            auto comparator =
+                GenerateExpressionWithLoadIfNeeded(call->Arguments()[1]->Declaration());
             if (comparator == 0) {
                 return false;
             }
diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.spvasm b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.spvasm
index 0c93134..0999969 100644
--- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.spvasm
@@ -1,9 +1,7 @@
-SKIP: FAILED
-
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 35
+; Bound: 36
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
@@ -45,32 +43,30 @@
      %uint_0 = OpConstant %uint 0
 %_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
 %_ptr_Function___atomic_compare_exchange_resulti32 = OpTypePointer Function %__atomic_compare_exchange_resulti32
-         %28 = OpConstantNull %__atomic_compare_exchange_resulti32
+         %29 = OpConstantNull %__atomic_compare_exchange_resulti32
 %atomicCompareExchangeWeak_1bd40a = OpFunction %void None %5
           %8 = OpLabel
       %arg_1 = OpVariable %_ptr_Function_int Function %12
       %arg_2 = OpVariable %_ptr_Function_int Function %12
-        %res = OpVariable %_ptr_Function___atomic_compare_exchange_resulti32 Function %28
+        %res = OpVariable %_ptr_Function___atomic_compare_exchange_resulti32 Function %29
                OpStore %arg_1 %int_1
                OpStore %arg_2 %int_1
          %22 = OpAccessChain %_ptr_StorageBuffer_int %sb_rw %uint_0
          %23 = OpLoad %int %arg_2
-         %24 = OpAtomicCompareExchange %int %22 %uint_1 %uint_0 %uint_0 %23 %arg_1
-         %25 = OpIEqual %bool %24 %23
-         %14 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %24 %25
+         %24 = OpLoad %int %arg_1
+         %25 = OpAtomicCompareExchange %int %22 %uint_1 %uint_0 %uint_0 %23 %24
+         %26 = OpIEqual %bool %25 %23
+         %14 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %25 %26
                OpStore %res %14
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %5
-         %30 = OpLabel
-         %31 = OpFunctionCall %void %atomicCompareExchangeWeak_1bd40a
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %atomicCompareExchangeWeak_1bd40a
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %5
-         %33 = OpLabel
-         %34 = OpFunctionCall %void %atomicCompareExchangeWeak_1bd40a
+         %34 = OpLabel
+         %35 = OpFunctionCall %void %atomicCompareExchangeWeak_1bd40a
                OpReturn
                OpFunctionEnd
-1:1: AtomicCompareExchange: expected Comparator to be of type Result Type
-  %24 = OpAtomicCompareExchange %int %22 %uint_1 %uint_0 %uint_0 %23 %arg_1
-
diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.spvasm b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.spvasm
index 62eac61..3723643 100644
--- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.spvasm
@@ -1,9 +1,7 @@
-SKIP: FAILED
-
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 33
+; Bound: 34
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
@@ -43,32 +41,30 @@
      %uint_0 = OpConstant %uint 0
 %_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
 %_ptr_Function___atomic_compare_exchange_resultu32 = OpTypePointer Function %__atomic_compare_exchange_resultu32
-         %26 = OpConstantNull %__atomic_compare_exchange_resultu32
+         %27 = OpConstantNull %__atomic_compare_exchange_resultu32
 %atomicCompareExchangeWeak_63d8e6 = OpFunction %void None %5
           %8 = OpLabel
       %arg_1 = OpVariable %_ptr_Function_uint Function %12
       %arg_2 = OpVariable %_ptr_Function_uint Function %12
-        %res = OpVariable %_ptr_Function___atomic_compare_exchange_resultu32 Function %26
+        %res = OpVariable %_ptr_Function___atomic_compare_exchange_resultu32 Function %27
                OpStore %arg_1 %uint_1
                OpStore %arg_2 %uint_1
          %20 = OpAccessChain %_ptr_StorageBuffer_uint %sb_rw %uint_0
          %21 = OpLoad %uint %arg_2
-         %22 = OpAtomicCompareExchange %uint %20 %uint_1 %uint_0 %uint_0 %21 %arg_1
-         %23 = OpIEqual %bool %22 %21
-         %14 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %22 %23
+         %22 = OpLoad %uint %arg_1
+         %23 = OpAtomicCompareExchange %uint %20 %uint_1 %uint_0 %uint_0 %21 %22
+         %24 = OpIEqual %bool %23 %21
+         %14 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %23 %24
                OpStore %res %14
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %5
-         %28 = OpLabel
-         %29 = OpFunctionCall %void %atomicCompareExchangeWeak_63d8e6
+         %29 = OpLabel
+         %30 = OpFunctionCall %void %atomicCompareExchangeWeak_63d8e6
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %5
-         %31 = OpLabel
-         %32 = OpFunctionCall %void %atomicCompareExchangeWeak_63d8e6
+         %32 = OpLabel
+         %33 = OpFunctionCall %void %atomicCompareExchangeWeak_63d8e6
                OpReturn
                OpFunctionEnd
-1:1: AtomicCompareExchange: expected Comparator to be of type Result Type
-  %22 = OpAtomicCompareExchange %uint %20 %uint_1 %uint_0 %uint_0 %21 %arg_1
-
diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.hlsl b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.hlsl
index 854c568..a445a1d 100644
--- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.hlsl
+++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.hlsl
@@ -1,17 +1,34 @@
-SKIP: FAILED
+struct atomic_compare_exchange_resultu32 {
+  uint old_value;
+  bool exchanged;
+};
+groupshared uint arg_0;
 
-
-var<workgroup> arg_0 : atomic<u32>;
-
-fn atomicCompareExchangeWeak_83580d() {
-  var arg_1 = 1u;
-  var arg_2 = 1u;
-  var res = atomicCompareExchangeWeak(&(arg_0), arg_1, arg_2);
+void atomicCompareExchangeWeak_83580d() {
+  uint arg_1 = 1u;
+  uint arg_2 = 1u;
+  atomic_compare_exchange_resultu32 atomic_result = (atomic_compare_exchange_resultu32)0;
+  uint atomic_compare_value = arg_1;
+  InterlockedCompareExchange(arg_0, atomic_compare_value, arg_2, atomic_result.old_value);
+  atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
+  atomic_compare_exchange_resultu32 res = atomic_result;
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
+struct tint_symbol_1 {
+  uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+  {
+    uint atomic_result_1 = 0u;
+    InterlockedExchange(arg_0, 0u, atomic_result_1);
+  }
+  GroupMemoryBarrierWithGroupSync();
   atomicCompareExchangeWeak_83580d();
 }
 
-Failed to generate: error: unknown type in EmitType
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+  compute_main_inner(tint_symbol.local_invocation_index);
+  return;
+}
diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.spvasm b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.spvasm
index 4f519e0..b163b6f 100644
--- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.spvasm
@@ -1,9 +1,7 @@
-SKIP: FAILED
-
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 40
+; Bound: 41
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
@@ -39,37 +37,35 @@
      %uint_2 = OpConstant %uint 2
      %uint_0 = OpConstant %uint 0
 %_ptr_Function___atomic_compare_exchange_resultu32 = OpTypePointer Function %__atomic_compare_exchange_resultu32
-         %26 = OpConstantNull %__atomic_compare_exchange_resultu32
-         %27 = OpTypeFunction %void %uint
+         %27 = OpConstantNull %__atomic_compare_exchange_resultu32
+         %28 = OpTypeFunction %void %uint
    %uint_264 = OpConstant %uint 264
 %atomicCompareExchangeWeak_83580d = OpFunction %void None %6
           %9 = OpLabel
       %arg_1 = OpVariable %_ptr_Function_uint Function %13
       %arg_2 = OpVariable %_ptr_Function_uint Function %13
-        %res = OpVariable %_ptr_Function___atomic_compare_exchange_resultu32 Function %26
+        %res = OpVariable %_ptr_Function___atomic_compare_exchange_resultu32 Function %27
                OpStore %arg_1 %uint_1
                OpStore %arg_2 %uint_1
          %21 = OpLoad %uint %arg_2
-         %22 = OpAtomicCompareExchange %uint %arg_0 %uint_2 %uint_0 %uint_0 %21 %arg_1
-         %23 = OpIEqual %bool %22 %21
-         %15 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %22 %23
+         %22 = OpLoad %uint %arg_1
+         %23 = OpAtomicCompareExchange %uint %arg_0 %uint_2 %uint_0 %uint_0 %21 %22
+         %24 = OpIEqual %bool %23 %21
+         %15 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %23 %24
                OpStore %res %15
                OpReturn
                OpFunctionEnd
-%compute_main_inner = OpFunction %void None %27
+%compute_main_inner = OpFunction %void None %28
 %local_invocation_index = OpFunctionParameter %uint
-         %30 = OpLabel
+         %31 = OpLabel
                OpAtomicStore %arg_0 %uint_2 %uint_0 %13
                OpControlBarrier %uint_2 %uint_2 %uint_264
-         %35 = OpFunctionCall %void %atomicCompareExchangeWeak_83580d
+         %36 = OpFunctionCall %void %atomicCompareExchangeWeak_83580d
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %6
-         %37 = OpLabel
-         %39 = OpLoad %uint %local_invocation_index_1
-         %38 = OpFunctionCall %void %compute_main_inner %39
+         %38 = OpLabel
+         %40 = OpLoad %uint %local_invocation_index_1
+         %39 = OpFunctionCall %void %compute_main_inner %40
                OpReturn
                OpFunctionEnd
-1:1: AtomicCompareExchange: expected Comparator to be of type Result Type
-  %22 = OpAtomicCompareExchange %uint %arg_0 %uint_2 %uint_0 %uint_0 %21 %arg_1
-
diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.hlsl b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.hlsl
index ca63be9..c14dade 100644
--- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.hlsl
+++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.hlsl
@@ -1,17 +1,34 @@
-SKIP: FAILED
+struct atomic_compare_exchange_resulti32 {
+  int old_value;
+  bool exchanged;
+};
+groupshared int arg_0;
 
-
-var<workgroup> arg_0 : atomic<i32>;
-
-fn atomicCompareExchangeWeak_e88938() {
-  var arg_1 = 1;
-  var arg_2 = 1;
-  var res = atomicCompareExchangeWeak(&(arg_0), arg_1, arg_2);
+void atomicCompareExchangeWeak_e88938() {
+  int arg_1 = 1;
+  int arg_2 = 1;
+  atomic_compare_exchange_resulti32 atomic_result = (atomic_compare_exchange_resulti32)0;
+  int atomic_compare_value = arg_1;
+  InterlockedCompareExchange(arg_0, atomic_compare_value, arg_2, atomic_result.old_value);
+  atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
+  atomic_compare_exchange_resulti32 res = atomic_result;
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
+struct tint_symbol_1 {
+  uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+  {
+    int atomic_result_1 = 0;
+    InterlockedExchange(arg_0, 0, atomic_result_1);
+  }
+  GroupMemoryBarrierWithGroupSync();
   atomicCompareExchangeWeak_e88938();
 }
 
-Failed to generate: error: unknown type in EmitType
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+  compute_main_inner(tint_symbol.local_invocation_index);
+  return;
+}
diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.spvasm b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.spvasm
index 9bf039b..b4e3102 100644
--- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.spvasm
@@ -1,9 +1,7 @@
-SKIP: FAILED
-
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 41
+; Bound: 42
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
@@ -40,37 +38,35 @@
      %uint_2 = OpConstant %uint 2
      %uint_0 = OpConstant %uint 0
 %_ptr_Function___atomic_compare_exchange_resulti32 = OpTypePointer Function %__atomic_compare_exchange_resulti32
-         %27 = OpConstantNull %__atomic_compare_exchange_resulti32
-         %28 = OpTypeFunction %void %uint
+         %28 = OpConstantNull %__atomic_compare_exchange_resulti32
+         %29 = OpTypeFunction %void %uint
    %uint_264 = OpConstant %uint 264
 %atomicCompareExchangeWeak_e88938 = OpFunction %void None %7
          %10 = OpLabel
       %arg_1 = OpVariable %_ptr_Function_int Function %14
       %arg_2 = OpVariable %_ptr_Function_int Function %14
-        %res = OpVariable %_ptr_Function___atomic_compare_exchange_resulti32 Function %27
+        %res = OpVariable %_ptr_Function___atomic_compare_exchange_resulti32 Function %28
                OpStore %arg_1 %int_1
                OpStore %arg_2 %int_1
          %22 = OpLoad %int %arg_2
-         %23 = OpAtomicCompareExchange %int %arg_0 %uint_2 %uint_0 %uint_0 %22 %arg_1
-         %24 = OpIEqual %bool %23 %22
-         %16 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %23 %24
+         %23 = OpLoad %int %arg_1
+         %24 = OpAtomicCompareExchange %int %arg_0 %uint_2 %uint_0 %uint_0 %22 %23
+         %25 = OpIEqual %bool %24 %22
+         %16 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %24 %25
                OpStore %res %16
                OpReturn
                OpFunctionEnd
-%compute_main_inner = OpFunction %void None %28
+%compute_main_inner = OpFunction %void None %29
 %local_invocation_index = OpFunctionParameter %uint
-         %31 = OpLabel
+         %32 = OpLabel
                OpAtomicStore %arg_0 %uint_2 %uint_0 %14
                OpControlBarrier %uint_2 %uint_2 %uint_264
-         %36 = OpFunctionCall %void %atomicCompareExchangeWeak_e88938
+         %37 = OpFunctionCall %void %atomicCompareExchangeWeak_e88938
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %7
-         %38 = OpLabel
-         %40 = OpLoad %uint %local_invocation_index_1
-         %39 = OpFunctionCall %void %compute_main_inner %40
+         %39 = OpLabel
+         %41 = OpLoad %uint %local_invocation_index_1
+         %40 = OpFunctionCall %void %compute_main_inner %41
                OpReturn
                OpFunctionEnd
-1:1: AtomicCompareExchange: expected Comparator to be of type Result Type
-  %23 = OpAtomicCompareExchange %int %arg_0 %uint_2 %uint_0 %uint_0 %22 %arg_1
-