diff --git a/CHANGELOG.md b/CHANGELOG.md
index b92c3f02df..39c7016313 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -49,6 +49,7 @@ Bottom level categories:
 #### Naga
 
 - When emitting GLSL, Uniform and Storage Buffer memory layouts are now emitted even if no explicit binding is given. By @cloone8 in [#7579](https://github.com/gfx-rs/wgpu/pull/7579).
+- Add support for `atomicCompareExchangeWeak` in HLSL and GLSL backends. By @cryvosh in [#7658](https://github.com/gfx-rs/wgpu/pull/7658)
 
 ### Bug Fixes
 
diff --git a/naga/src/back/glsl/mod.rs b/naga/src/back/glsl/mod.rs
index 52a47487ea..07cd6bc555 100644
--- a/naga/src/back/glsl/mod.rs
+++ b/naga/src/back/glsl/mod.rs
@@ -621,6 +621,9 @@ pub struct Writer<'a, W> {
     multiview: Option<core::num::NonZeroU32>,
     /// Mapping of varying variables to their location. Needed for reflections.
     varying: crate::FastHashMap<String, VaryingLocation>,
+
+    /// Set of special type names whose definitions have already been written. To prevent duplicates.
+    written_special_struct_names: crate::FastHashSet<String>,
 }
 
 impl<'a, W: Write> Writer<'a, W> {
@@ -688,6 +691,7 @@ impl<'a, W: Write> Writer<'a, W> {
             need_bake_expressions: Default::default(),
             continue_ctx: back::continue_forward::ContinueCtx::default(),
             varying: Default::default(),
+            written_special_struct_names: Default::default(),
         };
 
         // Find all features required to print this module
@@ -787,6 +791,12 @@ impl<'a, W: Write> Writer<'a, W> {
         // you can't make a struct without adding all of its members first.
         for (handle, ty) in self.module.types.iter() {
             if let TypeInner::Struct { ref members, .. } = ty.inner {
+                // Skip special atomic compare exchange result structs (generated in next loop)
+                let struct_name = &self.names[&NameKey::Type(handle)];
+                if struct_name.starts_with("_atomic_compare_exchange_result") {
+                    continue;
+                }
+
                 // Structures ending with runtime-sized arrays can only be
                 // rendered as shader storage blocks in GLSL, not stand-alone
                 // struct types.
@@ -794,16 +804,23 @@ impl<'a, W: Write> Writer<'a, W> {
                     .inner
                     .is_dynamically_sized(&self.module.types)
                 {
-                    let name = &self.names[&NameKey::Type(handle)];
-                    write!(self.out, "struct {name} ")?;
+                    write!(self.out, "struct {struct_name} ")?;
                     self.write_struct_body(handle, members)?;
                     writeln!(self.out, ";")?;
                 }
             }
         }
 
-        // Write functions to create special types.
+        // Write functions and struct definitions for special types.
         for (type_key, struct_ty) in self.module.special_types.predeclared_types.iter() {
+            let struct_name = &self.names[&NameKey::Type(*struct_ty)];
+            if !self
+                .written_special_struct_names
+                .insert(struct_name.clone())
+            {
+                continue;
+            }
+
             match type_key {
                 &crate::PredeclaredType::ModfResult { size, scalar }
                 | &crate::PredeclaredType::FrexpResult { size, scalar } => {
@@ -835,8 +852,6 @@ impl<'a, W: Write> Writer<'a, W> {
                             (FREXP_FUNCTION, "frexp", other_type_name)
                         };
 
-                    let struct_name = &self.names[&NameKey::Type(*struct_ty)];
-
                     writeln!(self.out)?;
                     if !self.options.version.supports_frexp_function()
                         && matches!(type_key, &crate::PredeclaredType::FrexpResult { .. })
@@ -860,7 +875,14 @@ impl<'a, W: Write> Writer<'a, W> {
                         )?;
                     }
                 }
-                &crate::PredeclaredType::AtomicCompareExchangeWeakResult { .. } => {}
+                &crate::PredeclaredType::AtomicCompareExchangeWeakResult(scalar) => {
+                    let scalar_str = glsl_scalar(scalar)?.full;
+                    writeln!(
+                        self.out,
+                        "struct {} {{\n    {} old_value;\n    bool exchanged;\n}};",
+                        struct_name, scalar_str
+                    )?;
+                }
             }
         }
 
@@ -1118,6 +1140,17 @@ impl<'a, W: Write> Writer<'a, W> {
     /// # Notes
     /// Adds no trailing or leading whitespace
     fn write_type(&mut self, ty: Handle<crate::Type>) -> BackendResult {
+        for (key, &handle) in self.module.special_types.predeclared_types.iter() {
+            if handle == ty {
+                if let crate::PredeclaredType::AtomicCompareExchangeWeakResult(_) = *key {
+                    let name = &self.names[&NameKey::Type(ty)];
+                    write!(self.out, "{name}")?;
+                    return Ok(());
+                }
+                break;
+            }
+        }
+
         match self.module.types[ty].inner {
             // glsl has no pointer types so just write types as normal and loads are skipped
             TypeInner::Pointer { base, .. } => self.write_type(base),
@@ -2572,33 +2605,49 @@ impl<'a, W: Write> Writer<'a, W> {
                 result,
             } => {
                 write!(self.out, "{level}")?;
-                if let Some(result) = result {
-                    let res_name = Baked(result).to_string();
-                    let res_ty = ctx.resolve_type(result, &self.module.types);
-                    self.write_value_type(res_ty)?;
-                    write!(self.out, " {res_name} = ")?;
-                    self.named_expressions.insert(result, res_name);
-                }
 
-                let fun_str = fun.to_glsl();
-                write!(self.out, "atomic{fun_str}(")?;
-                self.write_expr(pointer, ctx)?;
-                write!(self.out, ", ")?;
-                // handle the special cases
                 match *fun {
-                    crate::AtomicFunction::Subtract => {
-                        // we just wrote `InterlockedAdd`, so negate the argument
-                        write!(self.out, "-")?;
+                    crate::AtomicFunction::Exchange {
+                        compare: Some(compare_expr),
+                    } => {
+                        let result_handle = result.expect("CompareExchange must have a result");
+                        let res_name = Baked(result_handle).to_string();
+                        self.write_type(ctx.info[result_handle].ty.handle().unwrap())?;
+                        write!(self.out, " {res_name};")?;
+                        write!(self.out, " {res_name}.old_value = atomicCompSwap(")?;
+                        self.write_expr(pointer, ctx)?;
+                        write!(self.out, ", ")?;
+                        self.write_expr(compare_expr, ctx)?;
+                        write!(self.out, ", ")?;
+                        self.write_expr(value, ctx)?;
+                        writeln!(self.out, ");")?;
+
+                        write!(
+                            self.out,
+                            "{level}{res_name}.exchanged = ({res_name}.old_value == "
+                        )?;
+                        self.write_expr(compare_expr, ctx)?;
+                        writeln!(self.out, ");")?;
+                        self.named_expressions.insert(result_handle, res_name);
                     }
-                    crate::AtomicFunction::Exchange { compare: Some(_) } => {
-                        return Err(Error::Custom(
-                            "atomic CompareExchange is not implemented".to_string(),
-                        ));
+                    _ => {
+                        if let Some(result) = result {
+                            let res_name = Baked(result).to_string();
+                            self.write_type(ctx.info[result].ty.handle().unwrap())?;
+                            write!(self.out, " {res_name} = ")?;
+                            self.named_expressions.insert(result, res_name);
+                        }
+                        let fun_str = fun.to_glsl();
+                        write!(self.out, "atomic{fun_str}(")?;
+                        self.write_expr(pointer, ctx)?;
+                        write!(self.out, ", ")?;
+                        if let crate::AtomicFunction::Subtract = *fun {
+                            write!(self.out, "-")?;
+                        }
+                        self.write_expr(value, ctx)?;
+                        writeln!(self.out, ");")?;
                     }
-                    _ => {}
                 }
-                self.write_expr(value, ctx)?;
-                writeln!(self.out, ");")?;
             }
             // Stores a value into an image.
             Statement::ImageAtomic {
diff --git a/naga/src/back/hlsl/conv.rs b/naga/src/back/hlsl/conv.rs
index 3b8fb8fb3d..ed40cbe510 100644
--- a/naga/src/back/hlsl/conv.rs
+++ b/naga/src/back/hlsl/conv.rs
@@ -222,7 +222,7 @@ impl crate::AtomicFunction {
             Self::Min => "Min",
             Self::Max => "Max",
             Self::Exchange { compare: None } => "Exchange",
-            Self::Exchange { .. } => "", //TODO
+            Self::Exchange { .. } => "CompareExchange",
         }
     }
 }
diff --git a/naga/src/back/hlsl/writer.rs b/naga/src/back/hlsl/writer.rs
index bb90db7859..3660ab5c31 100644
--- a/naga/src/back/hlsl/writer.rs
+++ b/naga/src/back/hlsl/writer.rs
@@ -2358,49 +2358,60 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                 result,
             } => {
                 write!(self.out, "{level}")?;
-                let res_name = match result {
-                    None => None,
-                    Some(result) => {
-                        let name = Baked(result).to_string();
-                        match func_ctx.info[result].ty {
-                            proc::TypeResolution::Handle(handle) => {
-                                self.write_type(module, handle)?
-                            }
-                            proc::TypeResolution::Value(ref value) => {
-                                self.write_value_type(module, value)?
-                            }
-                        };
-                        write!(self.out, " {name}; ")?;
-                        Some((result, name))
-                    }
+                let res_var_info = if let Some(res_handle) = result {
+                    let name = Baked(res_handle).to_string();
+                    match func_ctx.info[res_handle].ty {
+                        proc::TypeResolution::Handle(handle) => self.write_type(module, handle)?,
+                        proc::TypeResolution::Value(ref value) => {
+                            self.write_value_type(module, value)?
+                        }
+                    };
+                    write!(self.out, " {name}; ")?;
+                    Some((res_handle, name))
+                } else {
+                    None
                 };
-
-                // Validation ensures that `pointer` has a `Pointer` type.
                 let pointer_space = func_ctx
                     .resolve_type(pointer, &module.types)
                     .pointer_space()
                     .unwrap();
-
                 let fun_str = fun.to_hlsl_suffix();
+                let compare_expr = match *fun {
+                    crate::AtomicFunction::Exchange { compare: Some(cmp) } => Some(cmp),
+                    _ => None,
+                };
                 match pointer_space {
                     crate::AddressSpace::WorkGroup => {
                         write!(self.out, "Interlocked{fun_str}(")?;
                         self.write_expr(module, pointer, func_ctx)?;
+                        self.emit_hlsl_atomic_tail(
+                            module,
+                            func_ctx,
+                            fun,
+                            compare_expr,
+                            value,
+                            &res_var_info,
+                        )?;
                     }
                     crate::AddressSpace::Storage { .. } => {
                         let var_handle = self.fill_access_chain(module, pointer, func_ctx)?;
-                        // The call to `self.write_storage_address` wants
-                        // mutable access to all of `self`, so temporarily take
-                        // ownership of our reusable access chain buffer.
-                        let chain = mem::take(&mut self.temp_access_chain);
                         let var_name = &self.names[&NameKey::GlobalVariable(var_handle)];
                         let width = match func_ctx.resolve_type(value, &module.types) {
                             &TypeInner::Scalar(Scalar { width: 8, .. }) => "64",
                             _ => "",
                         };
                         write!(self.out, "{var_name}.Interlocked{fun_str}{width}(")?;
+                        let chain = mem::take(&mut self.temp_access_chain);
                         self.write_storage_address(module, &chain, func_ctx)?;
                         self.temp_access_chain = chain;
+                        self.emit_hlsl_atomic_tail(
+                            module,
+                            func_ctx,
+                            fun,
+                            compare_expr,
+                            value,
+                            &res_var_info,
+                        )?;
                     }
                     ref other => {
                         return Err(Error::Custom(format!(
@@ -2408,29 +2419,17 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                         )))
                     }
                 }
-                write!(self.out, ", ")?;
-                // handle the special cases
-                match *fun {
-                    crate::AtomicFunction::Subtract => {
-                        // we just wrote `InterlockedAdd`, so negate the argument
-                        write!(self.out, "-")?;
-                    }
-                    crate::AtomicFunction::Exchange { compare: Some(_) } => {
-                        return Err(Error::Unimplemented("atomic CompareExchange".to_string()));
+                if let Some(cmp) = compare_expr {
+                    if let Some(&(res_handle, ref res_name)) = res_var_info.as_ref() {
+                        write!(
+                            self.out,
+                            "{level}{res_name}.exchanged = ({res_name}.old_value == "
+                        )?;
+                        self.write_expr(module, cmp, func_ctx)?;
+                        writeln!(self.out, ");")?;
+                        self.named_expressions.insert(res_handle, res_name.clone());
                     }
-                    _ => {}
-                }
-                self.write_expr(module, value, func_ctx)?;
-
-                // The `original_value` out parameter is optional for all the
-                // `Interlocked` functions we generate other than
-                // `InterlockedExchange`.
-                if let Some((result, name)) = res_name {
-                    write!(self.out, ", {name}")?;
-                    self.named_expressions.insert(result, name);
                 }
-
-                writeln!(self.out, ");")?;
             }
             Statement::ImageAtomic {
                 image,
@@ -4287,6 +4286,38 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
         }
         Ok(())
     }
+
+    /// Helper to emit the shared tail of an HLSL atomic call (arguments, value, result)
+    fn emit_hlsl_atomic_tail(
+        &mut self,
+        module: &Module,
+        func_ctx: &back::FunctionCtx<'_>,
+        fun: &crate::AtomicFunction,
+        compare_expr: Option<Handle<crate::Expression>>,
+        value: Handle<crate::Expression>,
+        res_var_info: &Option<(Handle<crate::Expression>, String)>,
+    ) -> BackendResult {
+        if let Some(cmp) = compare_expr {
+            write!(self.out, ", ")?;
+            self.write_expr(module, cmp, func_ctx)?;
+        }
+        write!(self.out, ", ")?;
+        if let crate::AtomicFunction::Subtract = *fun {
+            write!(self.out, "-")?;
+        }
+        self.write_expr(module, value, func_ctx)?;
+        if let Some(&(res_handle, ref res_name)) = res_var_info.as_ref() {
+            write!(self.out, ", ")?;
+            if compare_expr.is_some() {
+                write!(self.out, "{res_name}.old_value")?;
+            } else {
+                write!(self.out, "{res_name}")?;
+            }
+            self.named_expressions.insert(res_handle, res_name.clone());
+        }
+        writeln!(self.out, ");")?;
+        Ok(())
+    }
 }
 
 pub(super) struct MatrixType {
diff --git a/naga/tests/in/wgsl/atomicCompareExchange-int64.toml b/naga/tests/in/wgsl/atomicCompareExchange-int64.toml
index 6107dfc559..99227637f3 100644
--- a/naga/tests/in/wgsl/atomicCompareExchange-int64.toml
+++ b/naga/tests/in/wgsl/atomicCompareExchange-int64.toml
@@ -1,7 +1,8 @@
 god_mode = true
-targets = "SPIRV | WGSL"
+targets = "SPIRV | HLSL | WGSL"
 
 [hlsl]
+shader_model = "V6_6"
 fake_missing_bindings = true
 push_constants_target = { register = 0, space = 0 }
 restrict_indexing = true
diff --git a/naga/tests/in/wgsl/atomicCompareExchange.toml b/naga/tests/in/wgsl/atomicCompareExchange.toml
index 7d87498324..c9811b10d1 100644
--- a/naga/tests/in/wgsl/atomicCompareExchange.toml
+++ b/naga/tests/in/wgsl/atomicCompareExchange.toml
@@ -1 +1 @@
-targets = "SPIRV | METAL | WGSL"
+targets = "SPIRV | METAL | GLSL | HLSL | WGSL"
diff --git a/naga/tests/in/wgsl/atomicOps-int64.wgsl b/naga/tests/in/wgsl/atomicOps-int64.wgsl
index 42857d2fa4..42d5a88e78 100644
--- a/naga/tests/in/wgsl/atomicOps-int64.wgsl
+++ b/naga/tests/in/wgsl/atomicOps-int64.wgsl
@@ -129,13 +129,12 @@ fn cs_main(@builtin(local_invocation_id) id: vec3<u32>) {
     atomicExchange(&workgroup_struct.atomic_scalar, 1lu);
     atomicExchange(&workgroup_struct.atomic_arr[1], 1li);
 
-    // // TODO: https://github.com/gpuweb/gpuweb/issues/2021
-    // atomicCompareExchangeWeak(&storage_atomic_scalar, 1lu);
-    // atomicCompareExchangeWeak(&storage_atomic_arr[1], 1li);
-    // atomicCompareExchangeWeak(&storage_struct.atomic_scalar, 1lu);
-    // atomicCompareExchangeWeak(&storage_struct.atomic_arr[1], 1li);
-    // atomicCompareExchangeWeak(&workgroup_atomic_scalar, 1lu);
-    // atomicCompareExchangeWeak(&workgroup_atomic_arr[1], 1li);
-    // atomicCompareExchangeWeak(&workgroup_struct.atomic_scalar, 1lu);
-    // atomicCompareExchangeWeak(&workgroup_struct.atomic_arr[1], 1li);
+    let cas_res_0 = atomicCompareExchangeWeak(&storage_atomic_scalar, 1lu, 2lu);
+    let cas_res_1 = atomicCompareExchangeWeak(&storage_atomic_arr[1], 1li, 2li);
+    let cas_res_2 = atomicCompareExchangeWeak(&storage_struct.atomic_scalar, 1lu, 2lu);
+    let cas_res_3 = atomicCompareExchangeWeak(&storage_struct.atomic_arr[1], 1li, 2li);
+    let cas_res_4 = atomicCompareExchangeWeak(&workgroup_atomic_scalar, 1lu, 2lu);
+    let cas_res_5 = atomicCompareExchangeWeak(&workgroup_atomic_arr[1], 1li, 2li);
+    let cas_res_6 = atomicCompareExchangeWeak(&workgroup_struct.atomic_scalar, 1lu, 2lu);
+    let cas_res_7 = atomicCompareExchangeWeak(&workgroup_struct.atomic_arr[1], 1li, 2li);
 }
diff --git a/naga/tests/in/wgsl/atomicOps.wgsl b/naga/tests/in/wgsl/atomicOps.wgsl
index c1dd6b6326..3f0fa09441 100644
--- a/naga/tests/in/wgsl/atomicOps.wgsl
+++ b/naga/tests/in/wgsl/atomicOps.wgsl
@@ -129,13 +129,12 @@ fn cs_main(@builtin(local_invocation_id) id: vec3<u32>) {
     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);
+    let cas_res_0 = atomicCompareExchangeWeak(&storage_atomic_scalar, 1u, 2u);
+    let cas_res_1 = atomicCompareExchangeWeak(&storage_atomic_arr[1], 1i, 2i);
+    let cas_res_2 = atomicCompareExchangeWeak(&storage_struct.atomic_scalar, 1u, 2u);
+    let cas_res_3 = atomicCompareExchangeWeak(&storage_struct.atomic_arr[1], 1i, 2i);
+    let cas_res_4 = atomicCompareExchangeWeak(&workgroup_atomic_scalar, 1u, 2u);
+    let cas_res_5 = atomicCompareExchangeWeak(&workgroup_atomic_arr[1], 1i, 2i);
+    let cas_res_6 = atomicCompareExchangeWeak(&workgroup_struct.atomic_scalar, 1u, 2u);
+    let cas_res_7 = atomicCompareExchangeWeak(&workgroup_struct.atomic_arr[1], 1i, 2i);
 }
diff --git a/naga/tests/out/glsl/wgsl-atomicCompareExchange.test_atomic_compare_exchange_i32.Compute.glsl b/naga/tests/out/glsl/wgsl-atomicCompareExchange.test_atomic_compare_exchange_i32.Compute.glsl
new file mode 100644
index 0000000000..92d9c4355f
--- /dev/null
+++ b/naga/tests/out/glsl/wgsl-atomicCompareExchange.test_atomic_compare_exchange_i32.Compute.glsl
@@ -0,0 +1,63 @@
+#version 310 es
+
+precision highp float;
+precision highp int;
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+struct _atomic_compare_exchange_resultSint4_ {
+    int old_value;
+    bool exchanged;
+};
+struct _atomic_compare_exchange_resultUint4_ {
+    uint old_value;
+    bool exchanged;
+};
+const uint SIZE = 128u;
+
+layout(std430) buffer type_2_block_0Compute { int _group_0_binding_0_cs[128]; };
+
+
+void main() {
+    uint i = 0u;
+    int old = 0;
+    bool exchanged = false;
+    bool loop_init = true;
+    while(true) {
+        if (!loop_init) {
+            uint _e27 = i;
+            i = (_e27 + 1u);
+        }
+        loop_init = false;
+        uint _e2 = i;
+        if ((_e2 < SIZE)) {
+        } else {
+            break;
+        }
+        {
+            uint _e6 = i;
+            int _e8 = _group_0_binding_0_cs[_e6];
+            old = _e8;
+            exchanged = false;
+            while(true) {
+                bool _e12 = exchanged;
+                if (!(_e12)) {
+                } else {
+                    break;
+                }
+                {
+                    int _e14 = old;
+                    int new = floatBitsToInt((intBitsToFloat(_e14) + 1.0));
+                    uint _e20 = i;
+                    int _e22 = old;
+                    _atomic_compare_exchange_resultSint4_ _e23; _e23.old_value = atomicCompSwap(_group_0_binding_0_cs[_e20], _e22, new);
+                    _e23.exchanged = (_e23.old_value == _e22);
+                    old = _e23.old_value;
+                    exchanged = _e23.exchanged;
+                }
+            }
+        }
+    }
+    return;
+}
+
diff --git a/naga/tests/out/glsl/wgsl-atomicCompareExchange.test_atomic_compare_exchange_u32.Compute.glsl b/naga/tests/out/glsl/wgsl-atomicCompareExchange.test_atomic_compare_exchange_u32.Compute.glsl
new file mode 100644
index 0000000000..bbb4dea843
--- /dev/null
+++ b/naga/tests/out/glsl/wgsl-atomicCompareExchange.test_atomic_compare_exchange_u32.Compute.glsl
@@ -0,0 +1,63 @@
+#version 310 es
+
+precision highp float;
+precision highp int;
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+struct _atomic_compare_exchange_resultSint4_ {
+    int old_value;
+    bool exchanged;
+};
+struct _atomic_compare_exchange_resultUint4_ {
+    uint old_value;
+    bool exchanged;
+};
+const uint SIZE = 128u;
+
+layout(std430) buffer type_4_block_0Compute { uint _group_0_binding_1_cs[128]; };
+
+
+void main() {
+    uint i_1 = 0u;
+    uint old_1 = 0u;
+    bool exchanged_1 = false;
+    bool loop_init = true;
+    while(true) {
+        if (!loop_init) {
+            uint _e27 = i_1;
+            i_1 = (_e27 + 1u);
+        }
+        loop_init = false;
+        uint _e2 = i_1;
+        if ((_e2 < SIZE)) {
+        } else {
+            break;
+        }
+        {
+            uint _e6 = i_1;
+            uint _e8 = _group_0_binding_1_cs[_e6];
+            old_1 = _e8;
+            exchanged_1 = false;
+            while(true) {
+                bool _e12 = exchanged_1;
+                if (!(_e12)) {
+                } else {
+                    break;
+                }
+                {
+                    uint _e14 = old_1;
+                    uint new = floatBitsToUint((uintBitsToFloat(_e14) + 1.0));
+                    uint _e20 = i_1;
+                    uint _e22 = old_1;
+                    _atomic_compare_exchange_resultUint4_ _e23; _e23.old_value = atomicCompSwap(_group_0_binding_1_cs[_e20], _e22, new);
+                    _e23.exchanged = (_e23.old_value == _e22);
+                    old_1 = _e23.old_value;
+                    exchanged_1 = _e23.exchanged;
+                }
+            }
+        }
+    }
+    return;
+}
+
diff --git a/naga/tests/out/glsl/wgsl-atomicOps.cs_main.Compute.glsl b/naga/tests/out/glsl/wgsl-atomicOps.cs_main.Compute.glsl
index b69c5107ce..7e1b5061ba 100644
--- a/naga/tests/out/glsl/wgsl-atomicOps.cs_main.Compute.glsl
+++ b/naga/tests/out/glsl/wgsl-atomicOps.cs_main.Compute.glsl
@@ -9,6 +9,14 @@ struct Struct {
     uint atomic_scalar;
     int atomic_arr[2];
 };
+struct _atomic_compare_exchange_resultUint4_ {
+    uint old_value;
+    bool exchanged;
+};
+struct _atomic_compare_exchange_resultSint4_ {
+    int old_value;
+    bool exchanged;
+};
 layout(std430) buffer type_block_0Compute { uint _group_0_binding_0_cs; };
 
 layout(std430) buffer type_2_block_1Compute { int _group_0_binding_1_cs[2]; };
@@ -127,6 +135,22 @@ void main() {
     int _e295 = atomicExchange(workgroup_atomic_arr[1], 1);
     uint _e299 = atomicExchange(workgroup_struct.atomic_scalar, 1u);
     int _e304 = atomicExchange(workgroup_struct.atomic_arr[1], 1);
+    _atomic_compare_exchange_resultUint4_ _e308; _e308.old_value = atomicCompSwap(_group_0_binding_0_cs, 1u, 2u);
+    _e308.exchanged = (_e308.old_value == 1u);
+    _atomic_compare_exchange_resultSint4_ _e313; _e313.old_value = atomicCompSwap(_group_0_binding_1_cs[1], 1, 2);
+    _e313.exchanged = (_e313.old_value == 1);
+    _atomic_compare_exchange_resultUint4_ _e318; _e318.old_value = atomicCompSwap(_group_0_binding_2_cs.atomic_scalar, 1u, 2u);
+    _e318.exchanged = (_e318.old_value == 1u);
+    _atomic_compare_exchange_resultSint4_ _e324; _e324.old_value = atomicCompSwap(_group_0_binding_2_cs.atomic_arr[1], 1, 2);
+    _e324.exchanged = (_e324.old_value == 1);
+    _atomic_compare_exchange_resultUint4_ _e328; _e328.old_value = atomicCompSwap(workgroup_atomic_scalar, 1u, 2u);
+    _e328.exchanged = (_e328.old_value == 1u);
+    _atomic_compare_exchange_resultSint4_ _e333; _e333.old_value = atomicCompSwap(workgroup_atomic_arr[1], 1, 2);
+    _e333.exchanged = (_e333.old_value == 1);
+    _atomic_compare_exchange_resultUint4_ _e338; _e338.old_value = atomicCompSwap(workgroup_struct.atomic_scalar, 1u, 2u);
+    _e338.exchanged = (_e338.old_value == 1u);
+    _atomic_compare_exchange_resultSint4_ _e344; _e344.old_value = atomicCompSwap(workgroup_struct.atomic_arr[1], 1, 2);
+    _e344.exchanged = (_e344.old_value == 1);
     return;
 }
 
diff --git a/naga/tests/out/hlsl/wgsl-atomicCompareExchange-int64.hlsl b/naga/tests/out/hlsl/wgsl-atomicCompareExchange-int64.hlsl
new file mode 100644
index 0000000000..e0e4ca4021
--- /dev/null
+++ b/naga/tests/out/hlsl/wgsl-atomicCompareExchange-int64.hlsl
@@ -0,0 +1,127 @@
+struct NagaConstants {
+    int first_vertex;
+    int first_instance;
+    uint other;
+};
+ConstantBuffer<NagaConstants> _NagaConstants: register(b0, space1);
+
+struct _atomic_compare_exchange_resultSint8_ {
+    int64_t old_value;
+    bool exchanged;
+    int _end_pad_0;
+};
+
+struct _atomic_compare_exchange_resultUint8_ {
+    uint64_t old_value;
+    bool exchanged;
+    int _end_pad_0;
+};
+
+static const uint SIZE = 128u;
+
+RWByteAddressBuffer arr_i64_ : register(u0);
+RWByteAddressBuffer arr_u64_ : register(u1);
+
+[numthreads(1, 1, 1)]
+void test_atomic_compare_exchange_i64_()
+{
+    uint i = 0u;
+    int64_t old = (int64_t)0;
+    bool exchanged = (bool)0;
+
+    uint2 loop_bound = uint2(4294967295u, 4294967295u);
+    bool loop_init = true;
+    while(true) {
+        if (all(loop_bound == uint2(0u, 0u))) { break; }
+        loop_bound -= uint2(loop_bound.y == 0u, 1u);
+        if (!loop_init) {
+            uint _e26 = i;
+            i = (_e26 + 1u);
+        }
+        loop_init = false;
+        uint _e2 = i;
+        if ((_e2 < SIZE)) {
+        } else {
+            break;
+        }
+        {
+            uint _e6 = i;
+            int64_t _e8 = arr_i64_.Load<int64_t>(_e6*8);
+            old = _e8;
+            exchanged = false;
+            uint2 loop_bound_1 = uint2(4294967295u, 4294967295u);
+            while(true) {
+                if (all(loop_bound_1 == uint2(0u, 0u))) { break; }
+                loop_bound_1 -= uint2(loop_bound_1.y == 0u, 1u);
+                bool _e12 = exchanged;
+                if (!(_e12)) {
+                } else {
+                    break;
+                }
+                {
+                    int64_t _e14 = old;
+                    int64_t new_ = (_e14 + 10L);
+                    uint _e19 = i;
+                    int64_t _e21 = old;
+                    _atomic_compare_exchange_resultSint8_ _e22; arr_i64_.InterlockedCompareExchange64(_e19*8, _e21, new_, _e22.old_value);
+                    _e22.exchanged = (_e22.old_value == _e21);
+                    old = _e22.old_value;
+                    exchanged = _e22.exchanged;
+                }
+            }
+        }
+    }
+    return;
+}
+
+[numthreads(1, 1, 1)]
+void test_atomic_compare_exchange_u64_()
+{
+    uint i_1 = 0u;
+    uint64_t old_1 = (uint64_t)0;
+    bool exchanged_1 = (bool)0;
+
+    uint2 loop_bound_2 = uint2(4294967295u, 4294967295u);
+    bool loop_init_1 = true;
+    while(true) {
+        if (all(loop_bound_2 == uint2(0u, 0u))) { break; }
+        loop_bound_2 -= uint2(loop_bound_2.y == 0u, 1u);
+        if (!loop_init_1) {
+            uint _e26 = i_1;
+            i_1 = (_e26 + 1u);
+        }
+        loop_init_1 = false;
+        uint _e2 = i_1;
+        if ((_e2 < SIZE)) {
+        } else {
+            break;
+        }
+        {
+            uint _e6 = i_1;
+            uint64_t _e8 = arr_u64_.Load<uint64_t>(_e6*8);
+            old_1 = _e8;
+            exchanged_1 = false;
+            uint2 loop_bound_3 = uint2(4294967295u, 4294967295u);
+            while(true) {
+                if (all(loop_bound_3 == uint2(0u, 0u))) { break; }
+                loop_bound_3 -= uint2(loop_bound_3.y == 0u, 1u);
+                bool _e12 = exchanged_1;
+                if (!(_e12)) {
+                } else {
+                    break;
+                }
+                {
+                    uint64_t _e14 = old_1;
+                    uint64_t new_1 = (_e14 + 10uL);
+                    uint _e19 = i_1;
+                    uint64_t _e21 = old_1;
+                    _atomic_compare_exchange_resultUint8_ _e22; arr_u64_.InterlockedCompareExchange64(_e19*8, _e21, new_1, _e22.old_value);
+                    _e22.exchanged = (_e22.old_value == _e21);
+                    old_1 = _e22.old_value;
+                    exchanged_1 = _e22.exchanged;
+                }
+            }
+        }
+    }
+    return;
+}
diff --git a/naga/tests/out/hlsl/wgsl-atomicCompareExchange-int64.ron b/naga/tests/out/hlsl/wgsl-atomicCompareExchange-int64.ron
new file mode 100644
index 0000000000..b7059590b4
--- /dev/null
+++ b/naga/tests/out/hlsl/wgsl-atomicCompareExchange-int64.ron
@@ -0,0 +1,16 @@
+(
+    vertex:[
+    ],
+    fragment:[
+    ],
+    compute:[
+        (
+            entry_point:"test_atomic_compare_exchange_i64_",
+            target_profile:"cs_6_6",
+        ),
+        (
+            entry_point:"test_atomic_compare_exchange_u64_",
+            target_profile:"cs_6_6",
+        ),
+    ],
+)
diff --git a/naga/tests/out/hlsl/wgsl-atomicCompareExchange.hlsl b/naga/tests/out/hlsl/wgsl-atomicCompareExchange.hlsl
new file mode 100644
index 0000000000..d523f7a9d2
--- /dev/null
+++ b/naga/tests/out/hlsl/wgsl-atomicCompareExchange.hlsl
@@ -0,0 +1,118 @@
+struct _atomic_compare_exchange_resultSint4_ {
+    int old_value;
+    bool exchanged;
+};
+
+struct _atomic_compare_exchange_resultUint4_ {
+    uint old_value;
+    bool exchanged;
+};
+
+static const uint SIZE = 128u;
+
+RWByteAddressBuffer arr_i32_ : register(u0);
+RWByteAddressBuffer arr_u32_ : register(u1);
+
+[numthreads(1, 1, 1)]
+void test_atomic_compare_exchange_i32_()
+{
+    uint i = 0u;
+    int old = (int)0;
+    bool exchanged = (bool)0;
+
+    uint2 loop_bound = uint2(4294967295u, 4294967295u);
+    bool loop_init = true;
+    while(true) {
+        if (all(loop_bound == uint2(0u, 0u))) { break; }
+        loop_bound -= uint2(loop_bound.y == 0u, 1u);
+        if (!loop_init) {
+            uint _e27 = i;
+            i = (_e27 + 1u);
+        }
+        loop_init = false;
+        uint _e2 = i;
+        if ((_e2 < SIZE)) {
+        } else {
+            break;
+        }
+        {
+            uint _e6 = i;
+            int _e8 = asint(arr_i32_.Load(_e6*4));
+            old = _e8;
+            exchanged = false;
+            uint2 loop_bound_1 = uint2(4294967295u, 4294967295u);
+            while(true) {
+                if (all(loop_bound_1 == uint2(0u, 0u))) { break; }
+                loop_bound_1 -= uint2(loop_bound_1.y == 0u, 1u);
+                bool _e12 = exchanged;
+                if (!(_e12)) {
+                } else {
+                    break;
+                }
+                {
+                    int _e14 = old;
+                    int new_ = asint((asfloat(_e14) + 1.0));
+                    uint _e20 = i;
+                    int _e22 = old;
+                    _atomic_compare_exchange_resultSint4_ _e23; arr_i32_.InterlockedCompareExchange(_e20*4, _e22, new_, _e23.old_value);
+                    _e23.exchanged = (_e23.old_value == _e22);
+                    old = _e23.old_value;
+                    exchanged = _e23.exchanged;
+                }
+            }
+        }
+    }
+    return;
+}
+
+[numthreads(1, 1, 1)]
+void test_atomic_compare_exchange_u32_()
+{
+    uint i_1 = 0u;
+    uint old_1 = (uint)0;
+    bool exchanged_1 = (bool)0;
+
+    uint2 loop_bound_2 = uint2(4294967295u, 4294967295u);
+    bool loop_init_1 = true;
+    while(true) {
+        if (all(loop_bound_2 == uint2(0u, 0u))) { break; }
+        loop_bound_2 -= uint2(loop_bound_2.y == 0u, 1u);
+        if (!loop_init_1) {
+            uint _e27 = i_1;
+            i_1 = (_e27 + 1u);
+        }
+        loop_init_1 = false;
+        uint _e2 = i_1;
+        if ((_e2 < SIZE)) {
+        } else {
+            break;
+        }
+        {
+            uint _e6 = i_1;
+            uint _e8 = asuint(arr_u32_.Load(_e6*4));
+            old_1 = _e8;
+            exchanged_1 = false;
+            uint2 loop_bound_3 = uint2(4294967295u, 4294967295u);
+            while(true) {
+                if (all(loop_bound_3 == uint2(0u, 0u))) { break; }
+                loop_bound_3 -= uint2(loop_bound_3.y == 0u, 1u);
+                bool _e12 = exchanged_1;
+                if (!(_e12)) {
+                } else {
+                    break;
+                }
+                {
+                    uint _e14 = old_1;
+                    uint new_1 = asuint((asfloat(_e14) + 1.0));
+                    uint _e20 = i_1;
+                    uint _e22 = old_1;
+                    _atomic_compare_exchange_resultUint4_ _e23; arr_u32_.InterlockedCompareExchange(_e20*4, _e22, new_1, _e23.old_value);
+                    _e23.exchanged = (_e23.old_value == _e22);
+                    old_1 = _e23.old_value;
+                    exchanged_1 = _e23.exchanged;
+                }
+            }
+        }
+    }
+    return;
+}
diff --git a/naga/tests/out/hlsl/wgsl-atomicCompareExchange.ron b/naga/tests/out/hlsl/wgsl-atomicCompareExchange.ron
new file mode 100644
index 0000000000..70f4fcb3d1
--- /dev/null
+++ b/naga/tests/out/hlsl/wgsl-atomicCompareExchange.ron
@@ -0,0 +1,16 @@
+(
+    vertex:[
+    ],
+    fragment:[
+    ],
+    compute:[
+        (
+            entry_point:"test_atomic_compare_exchange_i32_",
+            target_profile:"cs_5_1",
+        ),
+        (
+            entry_point:"test_atomic_compare_exchange_u32_",
+            target_profile:"cs_5_1",
+        ),
+    ],
+)
diff --git a/naga/tests/out/hlsl/wgsl-atomicOps-int64.hlsl b/naga/tests/out/hlsl/wgsl-atomicOps-int64.hlsl
index ea88f81753..5c05b4ea9f 100644
--- a/naga/tests/out/hlsl/wgsl-atomicOps-int64.hlsl
+++ b/naga/tests/out/hlsl/wgsl-atomicOps-int64.hlsl
@@ -10,6 +10,18 @@ struct Struct {
     int64_t atomic_arr[2];
 };
 
+struct _atomic_compare_exchange_resultUint8_ {
+    uint64_t old_value;
+    bool exchanged;
+    int _end_pad_0;
+};
+
+struct _atomic_compare_exchange_resultSint8_ {
+    int64_t old_value;
+    bool exchanged;
+    int _end_pad_0;
+};
+
 RWByteAddressBuffer storage_atomic_scalar : register(u0);
 RWByteAddressBuffer storage_atomic_arr : register(u1);
 RWByteAddressBuffer storage_struct : register(u2);
@@ -114,5 +126,21 @@ void cs_main(uint3 id : SV_GroupThreadID, uint3 __local_invocation_id : SV_Group
     int64_t _e279; InterlockedExchange(workgroup_atomic_arr[1], 1L, _e279);
     uint64_t _e283; InterlockedExchange(workgroup_struct.atomic_scalar, 1uL, _e283);
     int64_t _e288; InterlockedExchange(workgroup_struct.atomic_arr[1], 1L, _e288);
+    _atomic_compare_exchange_resultUint8_ _e292; storage_atomic_scalar.InterlockedCompareExchange64(0, 1uL, 2uL, _e292.old_value);
+    _e292.exchanged = (_e292.old_value == 1uL);
+    _atomic_compare_exchange_resultSint8_ _e297; storage_atomic_arr.InterlockedCompareExchange64(8, 1L, 2L, _e297.old_value);
+    _e297.exchanged = (_e297.old_value == 1L);
+    _atomic_compare_exchange_resultUint8_ _e302; storage_struct.InterlockedCompareExchange64(0, 1uL, 2uL, _e302.old_value);
+    _e302.exchanged = (_e302.old_value == 1uL);
+    _atomic_compare_exchange_resultSint8_ _e308; storage_struct.InterlockedCompareExchange64(8+8, 1L, 2L, _e308.old_value);
+    _e308.exchanged = (_e308.old_value == 1L);
+    _atomic_compare_exchange_resultUint8_ _e312; InterlockedCompareExchange(workgroup_atomic_scalar, 1uL, 2uL, _e312.old_value);
+    _e312.exchanged = (_e312.old_value == 1uL);
+    _atomic_compare_exchange_resultSint8_ _e317; InterlockedCompareExchange(workgroup_atomic_arr[1], 1L, 2L, _e317.old_value);
+    _e317.exchanged = (_e317.old_value == 1L);
+    _atomic_compare_exchange_resultUint8_ _e322; InterlockedCompareExchange(workgroup_struct.atomic_scalar, 1uL, 2uL, _e322.old_value);
+    _e322.exchanged = (_e322.old_value == 1uL);
+    _atomic_compare_exchange_resultSint8_ _e328; InterlockedCompareExchange(workgroup_struct.atomic_arr[1], 1L, 2L, _e328.old_value);
+    _e328.exchanged = (_e328.old_value == 1L);
     return;
 }
diff --git a/naga/tests/out/hlsl/wgsl-atomicOps.hlsl b/naga/tests/out/hlsl/wgsl-atomicOps.hlsl
index 52d11aa325..5771c898d9 100644
--- a/naga/tests/out/hlsl/wgsl-atomicOps.hlsl
+++ b/naga/tests/out/hlsl/wgsl-atomicOps.hlsl
@@ -3,6 +3,16 @@ struct Struct {
     int atomic_arr[2];
 };
 
+struct _atomic_compare_exchange_resultUint4_ {
+    uint old_value;
+    bool exchanged;
+};
+
+struct _atomic_compare_exchange_resultSint4_ {
+    int old_value;
+    bool exchanged;
+};
+
 RWByteAddressBuffer storage_atomic_scalar : register(u0);
 RWByteAddressBuffer storage_atomic_arr : register(u1);
 RWByteAddressBuffer storage_struct : register(u2);
@@ -107,5 +117,21 @@ void cs_main(uint3 id : SV_GroupThreadID, uint3 __local_invocation_id : SV_Group
     int _e295; InterlockedExchange(workgroup_atomic_arr[1], int(1), _e295);
     uint _e299; InterlockedExchange(workgroup_struct.atomic_scalar, 1u, _e299);
     int _e304; InterlockedExchange(workgroup_struct.atomic_arr[1], int(1), _e304);
+    _atomic_compare_exchange_resultUint4_ _e308; storage_atomic_scalar.InterlockedCompareExchange(0, 1u, 2u, _e308.old_value);
+    _e308.exchanged = (_e308.old_value == 1u);
+    _atomic_compare_exchange_resultSint4_ _e313; storage_atomic_arr.InterlockedCompareExchange(4, int(1), int(2), _e313.old_value);
+    _e313.exchanged = (_e313.old_value == int(1));
+    _atomic_compare_exchange_resultUint4_ _e318; storage_struct.InterlockedCompareExchange(0, 1u, 2u, _e318.old_value);
+    _e318.exchanged = (_e318.old_value == 1u);
+    _atomic_compare_exchange_resultSint4_ _e324; storage_struct.InterlockedCompareExchange(4+4, int(1), int(2), _e324.old_value);
+    _e324.exchanged = (_e324.old_value == int(1));
+    _atomic_compare_exchange_resultUint4_ _e328; InterlockedCompareExchange(workgroup_atomic_scalar, 1u, 2u, _e328.old_value);
+    _e328.exchanged = (_e328.old_value == 1u);
+    _atomic_compare_exchange_resultSint4_ _e333; InterlockedCompareExchange(workgroup_atomic_arr[1], int(1), int(2), _e333.old_value);
+    _e333.exchanged = (_e333.old_value == int(1));
+    _atomic_compare_exchange_resultUint4_ _e338; InterlockedCompareExchange(workgroup_struct.atomic_scalar, 1u, 2u, _e338.old_value);
+    _e338.exchanged = (_e338.old_value == 1u);
+    _atomic_compare_exchange_resultSint4_ _e344; InterlockedCompareExchange(workgroup_struct.atomic_arr[1], int(1), int(2), _e344.old_value);
+    _e344.exchanged = (_e344.old_value == int(1));
     return;
 }
diff --git a/naga/tests/out/msl/wgsl-atomicOps.msl b/naga/tests/out/msl/wgsl-atomicOps.msl
index 4732b4a32d..860a5a3e5a 100644
--- a/naga/tests/out/msl/wgsl-atomicOps.msl
+++ b/naga/tests/out/msl/wgsl-atomicOps.msl
@@ -11,6 +11,64 @@ struct Struct {
     metal::atomic_uint atomic_scalar;
     type_2 atomic_arr;
 };
+struct _atomic_compare_exchange_resultUint4_ {
+    uint old_value;
+    bool exchanged;
+};
+struct _atomic_compare_exchange_resultSint4_ {
+    int old_value;
+    bool exchanged;
+};
+
+template <typename A>
+_atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit(
+    device A *atomic_ptr,
+    uint cmp,
+    uint v
+) {
+    bool swapped = metal::atomic_compare_exchange_weak_explicit(
+        atomic_ptr, &cmp, v,
+        metal::memory_order_relaxed, metal::memory_order_relaxed
+    );
+    return _atomic_compare_exchange_resultUint4_{cmp, swapped};
+}
+template <typename A>
+_atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit(
+    threadgroup A *atomic_ptr,
+    uint cmp,
+    uint v
+) {
+    bool swapped = metal::atomic_compare_exchange_weak_explicit(
+        atomic_ptr, &cmp, v,
+        metal::memory_order_relaxed, metal::memory_order_relaxed
+    );
+    return _atomic_compare_exchange_resultUint4_{cmp, swapped};
+}
+
+template <typename A>
+_atomic_compare_exchange_resultSint4_ naga_atomic_compare_exchange_weak_explicit(
+    device A *atomic_ptr,
+    int cmp,
+    int v
+) {
+    bool swapped = metal::atomic_compare_exchange_weak_explicit(
+        atomic_ptr, &cmp, v,
+        metal::memory_order_relaxed, metal::memory_order_relaxed
+    );
+    return _atomic_compare_exchange_resultSint4_{cmp, swapped};
+}
+template <typename A>
+_atomic_compare_exchange_resultSint4_ naga_atomic_compare_exchange_weak_explicit(
+    threadgroup A *atomic_ptr,
+    int cmp,
+    int v
+) {
+    bool swapped = metal::atomic_compare_exchange_weak_explicit(
+        atomic_ptr, &cmp, v,
+        metal::memory_order_relaxed, metal::memory_order_relaxed
+    );
+    return _atomic_compare_exchange_resultSint4_{cmp, swapped};
+}
 
 struct cs_mainInput {
 };
@@ -122,5 +180,13 @@ kernel void cs_main(
     int _e295 = metal::atomic_exchange_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed);
     uint _e299 = metal::atomic_exchange_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
     int _e304 = metal::atomic_exchange_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed);
+    _atomic_compare_exchange_resultUint4_ _e308 = naga_atomic_compare_exchange_weak_explicit(&storage_atomic_scalar, 1u, 2u);
+    _atomic_compare_exchange_resultSint4_ _e313 = naga_atomic_compare_exchange_weak_explicit(&storage_atomic_arr.inner[1], 1, 2);
+    _atomic_compare_exchange_resultUint4_ _e318 = naga_atomic_compare_exchange_weak_explicit(&storage_struct.atomic_scalar, 1u, 2u);
+    _atomic_compare_exchange_resultSint4_ _e324 = naga_atomic_compare_exchange_weak_explicit(&storage_struct.atomic_arr.inner[1], 1, 2);
+    _atomic_compare_exchange_resultUint4_ _e328 = naga_atomic_compare_exchange_weak_explicit(&workgroup_atomic_scalar, 1u, 2u);
+    _atomic_compare_exchange_resultSint4_ _e333 = naga_atomic_compare_exchange_weak_explicit(&workgroup_atomic_arr.inner[1], 1, 2);
+    _atomic_compare_exchange_resultUint4_ _e338 = naga_atomic_compare_exchange_weak_explicit(&workgroup_struct.atomic_scalar, 1u, 2u);
+    _atomic_compare_exchange_resultSint4_ _e344 = naga_atomic_compare_exchange_weak_explicit(&workgroup_struct.atomic_arr.inner[1], 1, 2);
     return;
 }
diff --git a/naga/tests/out/spv/wgsl-atomicOps-int64.spvasm b/naga/tests/out/spv/wgsl-atomicOps-int64.spvasm
index 943107f500..806ce47991 100644
--- a/naga/tests/out/spv/wgsl-atomicOps-int64.spvasm
+++ b/naga/tests/out/spv/wgsl-atomicOps-int64.spvasm
@@ -1,31 +1,35 @@
 ; SPIR-V
 ; Version: 1.0
 ; Generator: rspirv
-; Bound: 193
+; Bound: 227
 OpCapability Shader
 OpCapability Int64Atomics
 OpCapability Int64
 OpExtension "SPV_KHR_storage_buffer_storage_class"
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
-OpEntryPoint GLCompute %29 "cs_main" %26
-OpExecutionMode %29 LocalSize 2 1 1
+OpEntryPoint GLCompute %32 "cs_main" %29
+OpExecutionMode %32 LocalSize 2 1 1
 OpDecorate %5 ArrayStride 8
 OpMemberDecorate %8 0 Offset 0
 OpMemberDecorate %8 1 Offset 8
-OpDecorate %10 DescriptorSet 0
-OpDecorate %10 Binding 0
-OpDecorate %11 Block
 OpMemberDecorate %11 0 Offset 0
+OpMemberDecorate %11 1 Offset 8
+OpMemberDecorate %12 0 Offset 0
+OpMemberDecorate %12 1 Offset 8
 OpDecorate %13 DescriptorSet 0
-OpDecorate %13 Binding 1
+OpDecorate %13 Binding 0
 OpDecorate %14 Block
 OpMemberDecorate %14 0 Offset 0
 OpDecorate %16 DescriptorSet 0
-OpDecorate %16 Binding 2
+OpDecorate %16 Binding 1
 OpDecorate %17 Block
 OpMemberDecorate %17 0 Offset 0
-OpDecorate %26 BuiltIn LocalInvocationId
+OpDecorate %19 DescriptorSet 0
+OpDecorate %19 Binding 2
+OpDecorate %20 Block
+OpMemberDecorate %20 0 Offset 0
+OpDecorate %29 BuiltIn LocalInvocationId
 %2 = OpTypeVoid
 %3 = OpTypeInt 64 0
 %4 = OpTypeInt 64 1
@@ -34,213 +38,247 @@ OpDecorate %26 BuiltIn LocalInvocationId
 %5 = OpTypeArray %4 %6
 %8 = OpTypeStruct %3 %5
 %9 = OpTypeVector %7 3
-%11 = OpTypeStruct %3
-%12 = OpTypePointer StorageBuffer %11
-%10 = OpVariable  %12  StorageBuffer
-%14 = OpTypeStruct %5
+%10 = OpTypeBool
+%11 = OpTypeStruct %3 %10
+%12 = OpTypeStruct %4 %10
+%14 = OpTypeStruct %3
 %15 = OpTypePointer StorageBuffer %14
 %13 = OpVariable  %15  StorageBuffer
-%17 = OpTypeStruct %8
+%17 = OpTypeStruct %5
 %18 = OpTypePointer StorageBuffer %17
 %16 = OpVariable  %18  StorageBuffer
-%20 = OpTypePointer Workgroup %3
-%19 = OpVariable  %20  Workgroup
-%22 = OpTypePointer Workgroup %5
-%21 = OpVariable  %22  Workgroup
-%24 = OpTypePointer Workgroup %8
-%23 = OpVariable  %24  Workgroup
-%27 = OpTypePointer Input %9
-%26 = OpVariable  %27  Input
-%30 = OpTypeFunction %2
-%31 = OpTypePointer StorageBuffer %3
-%32 = OpConstant  %7  0
-%34 = OpTypePointer StorageBuffer %5
-%36 = OpTypePointer StorageBuffer %8
-%38 = OpConstant  %3  1
-%39 = OpConstant  %4  1
-%41 = OpConstantNull  %3
-%42 = OpConstantNull  %5
-%43 = OpConstantNull  %8
-%44 = OpConstantNull  %9
-%46 = OpTypeBool
-%45 = OpTypeVector %46 3
-%51 = OpConstant  %7  264
-%54 = OpTypeInt 32 1
-%53 = OpConstant  %54  1
-%55 = OpConstant  %7  64
-%56 = OpTypePointer StorageBuffer %4
-%57 = OpConstant  %7  1
-%61 = OpConstant  %54  2
-%62 = OpConstant  %7  256
-%63 = OpTypePointer Workgroup %4
-%29 = OpFunction  %2  None %30
-%25 = OpLabel
-%28 = OpLoad  %9  %26
-%33 = OpAccessChain  %31  %10 %32
-%35 = OpAccessChain  %34  %13 %32
-%37 = OpAccessChain  %36  %16 %32
-OpBranch %40
-%40 = OpLabel
-%47 = OpIEqual  %45  %28 %44
-%48 = OpAll  %46  %47
-OpSelectionMerge %49 None
-OpBranchConditional %48 %50 %49
-%50 = OpLabel
-OpStore %19 %41
-OpStore %21 %42
-OpStore %23 %43
-OpBranch %49
-%49 = OpLabel
-OpControlBarrier %6 %6 %51
-OpBranch %52
-%52 = OpLabel
-OpAtomicStore %33 %53 %55 %38
-%58 = OpAccessChain  %56  %35 %57
-OpAtomicStore %58 %53 %55 %39
-%59 = OpAccessChain  %31  %37 %32
-OpAtomicStore %59 %53 %55 %38
-%60 = OpAccessChain  %56  %37 %57 %57
-OpAtomicStore %60 %53 %55 %39
-OpAtomicStore %19 %61 %62 %38
-%64 = OpAccessChain  %63  %21 %57
-OpAtomicStore %64 %61 %62 %39
-%65 = OpAccessChain  %20  %23 %32
-OpAtomicStore %65 %61 %62 %38
-%66 = OpAccessChain  %63  %23 %57 %57
-OpAtomicStore %66 %61 %62 %39
-OpControlBarrier %6 %6 %51
-%67 = OpAtomicLoad  %3  %33 %53 %55
-%68 = OpAccessChain  %56  %35 %57
-%69 = OpAtomicLoad  %4  %68 %53 %55
-%70 = OpAccessChain  %31  %37 %32
-%71 = OpAtomicLoad  %3  %70 %53 %55
-%72 = OpAccessChain  %56  %37 %57 %57
-%73 = OpAtomicLoad  %4  %72 %53 %55
-%74 = OpAtomicLoad  %3  %19 %61 %62
-%75 = OpAccessChain  %63  %21 %57
-%76 = OpAtomicLoad  %4  %75 %61 %62
-%77 = OpAccessChain  %20  %23 %32
-%78 = OpAtomicLoad  %3  %77 %61 %62
-%79 = OpAccessChain  %63  %23 %57 %57
-%80 = OpAtomicLoad  %4  %79 %61 %62
-OpControlBarrier %6 %6 %51
-%81 = OpAtomicIAdd  %3  %33 %53 %55 %38
-%83 = OpAccessChain  %56  %35 %57
-%82 = OpAtomicIAdd  %4  %83 %53 %55 %39
-%85 = OpAccessChain  %31  %37 %32
-%84 = OpAtomicIAdd  %3  %85 %53 %55 %38
-%87 = OpAccessChain  %56  %37 %57 %57
-%86 = OpAtomicIAdd  %4  %87 %53 %55 %39
-%88 = OpAtomicIAdd  %3  %19 %61 %62 %38
-%90 = OpAccessChain  %63  %21 %57
-%89 = OpAtomicIAdd  %4  %90 %61 %62 %39
-%92 = OpAccessChain  %20  %23 %32
-%91 = OpAtomicIAdd  %3  %92 %61 %62 %38
-%94 = OpAccessChain  %63  %23 %57 %57
-%93 = OpAtomicIAdd  %4  %94 %61 %62 %39
-OpControlBarrier %6 %6 %51
-%95 = OpAtomicISub  %3  %33 %53 %55 %38
-%97 = OpAccessChain  %56  %35 %57
-%96 = OpAtomicISub  %4  %97 %53 %55 %39
-%99 = OpAccessChain  %31  %37 %32
-%98 = OpAtomicISub  %3  %99 %53 %55 %38
-%101 = OpAccessChain  %56  %37 %57 %57
-%100 = OpAtomicISub  %4  %101 %53 %55 %39
-%102 = OpAtomicISub  %3  %19 %61 %62 %38
-%104 = OpAccessChain  %63  %21 %57
-%103 = OpAtomicISub  %4  %104 %61 %62 %39
-%106 = OpAccessChain  %20  %23 %32
-%105 = OpAtomicISub  %3  %106 %61 %62 %38
-%108 = OpAccessChain  %63  %23 %57 %57
-%107 = OpAtomicISub  %4  %108 %61 %62 %39
-OpControlBarrier %6 %6 %51
-%109 = OpAtomicUMax  %3  %33 %53 %55 %38
-%111 = OpAccessChain  %56  %35 %57
-%110 = OpAtomicSMax  %4  %111 %53 %55 %39
-%113 = OpAccessChain  %31  %37 %32
-%112 = OpAtomicUMax  %3  %113 %53 %55 %38
-%115 = OpAccessChain  %56  %37 %57 %57
-%114 = OpAtomicSMax  %4  %115 %53 %55 %39
-%116 = OpAtomicUMax  %3  %19 %61 %62 %38
-%118 = OpAccessChain  %63  %21 %57
-%117 = OpAtomicSMax  %4  %118 %61 %62 %39
-%120 = OpAccessChain  %20  %23 %32
-%119 = OpAtomicUMax  %3  %120 %61 %62 %38
-%122 = OpAccessChain  %63  %23 %57 %57
-%121 = OpAtomicSMax  %4  %122 %61 %62 %39
-OpControlBarrier %6 %6 %51
-%123 = OpAtomicUMin  %3  %33 %53 %55 %38
-%125 = OpAccessChain  %56  %35 %57
-%124 = OpAtomicSMin  %4  %125 %53 %55 %39
-%127 = OpAccessChain  %31  %37 %32
-%126 = OpAtomicUMin  %3  %127 %53 %55 %38
-%129 = OpAccessChain  %56  %37 %57 %57
-%128 = OpAtomicSMin  %4  %129 %53 %55 %39
-%130 = OpAtomicUMin  %3  %19 %61 %62 %38
-%132 = OpAccessChain  %63  %21 %57
-%131 = OpAtomicSMin  %4  %132 %61 %62 %39
-%134 = OpAccessChain  %20  %23 %32
-%133 = OpAtomicUMin  %3  %134 %61 %62 %38
-%136 = OpAccessChain  %63  %23 %57 %57
-%135 = OpAtomicSMin  %4  %136 %61 %62 %39
-OpControlBarrier %6 %6 %51
-%137 = OpAtomicAnd  %3  %33 %53 %55 %38
-%139 = OpAccessChain  %56  %35 %57
-%138 = OpAtomicAnd  %4  %139 %53 %55 %39
-%141 = OpAccessChain  %31  %37 %32
-%140 = OpAtomicAnd  %3  %141 %53 %55 %38
-%143 = OpAccessChain  %56  %37 %57 %57
-%142 = OpAtomicAnd  %4  %143 %53 %55 %39
-%144 = OpAtomicAnd  %3  %19 %61 %62 %38
-%146 = OpAccessChain  %63  %21 %57
-%145 = OpAtomicAnd  %4  %146 %61 %62 %39
-%148 = OpAccessChain  %20  %23 %32
-%147 = OpAtomicAnd  %3  %148 %61 %62 %38
-%150 = OpAccessChain  %63  %23 %57 %57
-%149 = OpAtomicAnd  %4  %150 %61 %62 %39
-OpControlBarrier %6 %6 %51
-%151 = OpAtomicOr  %3  %33 %53 %55 %38
-%153 = OpAccessChain  %56  %35 %57
-%152 = OpAtomicOr  %4  %153 %53 %55 %39
-%155 = OpAccessChain  %31  %37 %32
-%154 = OpAtomicOr  %3  %155 %53 %55 %38
-%157 = OpAccessChain  %56  %37 %57 %57
-%156 = OpAtomicOr  %4  %157 %53 %55 %39
-%158 = OpAtomicOr  %3  %19 %61 %62 %38
-%160 = OpAccessChain  %63  %21 %57
-%159 = OpAtomicOr  %4  %160 %61 %62 %39
-%162 = OpAccessChain  %20  %23 %32
-%161 = OpAtomicOr  %3  %162 %61 %62 %38
-%164 = OpAccessChain  %63  %23 %57 %57
-%163 = OpAtomicOr  %4  %164 %61 %62 %39
-OpControlBarrier %6 %6 %51
-%165 = OpAtomicXor  %3  %33 %53 %55 %38
-%167 = OpAccessChain  %56  %35 %57
-%166 = OpAtomicXor  %4  %167 %53 %55 %39
-%169 = OpAccessChain  %31  %37 %32
-%168 = OpAtomicXor  %3  %169 %53 %55 %38
-%171 = OpAccessChain  %56  %37 %57 %57
-%170 = OpAtomicXor  %4  %171 %53 %55 %39
-%172 = OpAtomicXor  %3  %19 %61 %62 %38
-%174 = OpAccessChain  %63  %21 %57
-%173 = OpAtomicXor  %4  %174 %61 %62 %39
-%176 = OpAccessChain  %20  %23 %32
-%175 = OpAtomicXor  %3  %176 %61 %62 %38
-%178 = OpAccessChain  %63  %23 %57 %57
-%177 = OpAtomicXor  %4  %178 %61 %62 %39
-%179 = OpAtomicExchange  %3  %33 %53 %55 %38
-%181 = OpAccessChain  %56  %35 %57
-%180 = OpAtomicExchange  %4  %181 %53 %55 %39
-%183 = OpAccessChain  %31  %37 %32
-%182 = OpAtomicExchange  %3  %183 %53 %55 %38
-%185 = OpAccessChain  %56  %37 %57 %57
-%184 = OpAtomicExchange  %4  %185 %53 %55 %39
-%186 = OpAtomicExchange  %3  %19 %61 %62 %38
-%188 = OpAccessChain  %63  %21 %57
-%187 = OpAtomicExchange  %4  %188 %61 %62 %39
-%190 = OpAccessChain  %20  %23 %32
-%189 = OpAtomicExchange  %3  %190 %61 %62 %38
-%192 = OpAccessChain  %63  %23 %57 %57
-%191 = OpAtomicExchange  %4  %192 %61 %62 %39
+%20 = OpTypeStruct %8
+%21 = OpTypePointer StorageBuffer %20
+%19 = OpVariable  %21  StorageBuffer
+%23 = OpTypePointer Workgroup %3
+%22 = OpVariable  %23  Workgroup
+%25 = OpTypePointer Workgroup %5
+%24 = OpVariable  %25  Workgroup
+%27 = OpTypePointer Workgroup %8
+%26 = OpVariable  %27  Workgroup
+%30 = OpTypePointer Input %9
+%29 = OpVariable  %30  Input
+%33 = OpTypeFunction %2
+%34 = OpTypePointer StorageBuffer %3
+%35 = OpConstant  %7  0
+%37 = OpTypePointer StorageBuffer %5
+%39 = OpTypePointer StorageBuffer %8
+%41 = OpConstant  %3  1
+%42 = OpConstant  %4  1
+%43 = OpConstant  %3  2
+%44 = OpConstant  %4  2
+%46 = OpConstantNull  %3
+%47 = OpConstantNull  %5
+%48 = OpConstantNull  %8
+%49 = OpConstantNull  %9
+%50 = OpTypeVector %10 3
+%55 = OpConstant  %7  264
+%58 = OpTypeInt 32 1
+%57 = OpConstant  %58  1
+%59 = OpConstant  %7  64
+%60 = OpTypePointer StorageBuffer %4
+%61 = OpConstant  %7  1
+%65 = OpConstant  %58  2
+%66 = OpConstant  %7  256
+%67 = OpTypePointer Workgroup %4
+%32 = OpFunction  %2  None %33
+%28 = OpLabel
+%31 = OpLoad  %9  %29
+%36 = OpAccessChain  %34  %13 %35
+%38 = OpAccessChain  %37  %16 %35
+%40 = OpAccessChain  %39  %19 %35
+OpBranch %45
+%45 = OpLabel
+%51 = OpIEqual  %50  %31 %49
+%52 = OpAll  %10  %51
+OpSelectionMerge %53 None
+OpBranchConditional %52 %54 %53
+%54 = OpLabel
+OpStore %22 %46
+OpStore %24 %47
+OpStore %26 %48
+OpBranch %53
+%53 = OpLabel
+OpControlBarrier %6 %6 %55
+OpBranch %56
+%56 = OpLabel
+OpAtomicStore %36 %57 %59 %41
+%62 = OpAccessChain  %60  %38 %61
+OpAtomicStore %62 %57 %59 %42
+%63 = OpAccessChain  %34  %40 %35
+OpAtomicStore %63 %57 %59 %41
+%64 = OpAccessChain  %60  %40 %61 %61
+OpAtomicStore %64 %57 %59 %42
+OpAtomicStore %22 %65 %66 %41
+%68 = OpAccessChain  %67  %24 %61
+OpAtomicStore %68 %65 %66 %42
+%69 = OpAccessChain  %23  %26 %35
+OpAtomicStore %69 %65 %66 %41
+%70 = OpAccessChain  %67  %26 %61 %61
+OpAtomicStore %70 %65 %66 %42
+OpControlBarrier %6 %6 %55
+%71 = OpAtomicLoad  %3  %36 %57 %59
+%72 = OpAccessChain  %60  %38 %61
+%73 = OpAtomicLoad  %4  %72 %57 %59
+%74 = OpAccessChain  %34  %40 %35
+%75 = OpAtomicLoad  %3  %74 %57 %59
+%76 = OpAccessChain  %60  %40 %61 %61
+%77 = OpAtomicLoad  %4  %76 %57 %59
+%78 = OpAtomicLoad  %3  %22 %65 %66
+%79 = OpAccessChain  %67  %24 %61
+%80 = OpAtomicLoad  %4  %79 %65 %66
+%81 = OpAccessChain  %23  %26 %35
+%82 = OpAtomicLoad  %3  %81 %65 %66
+%83 = OpAccessChain  %67  %26 %61 %61
+%84 = OpAtomicLoad  %4  %83 %65 %66
+OpControlBarrier %6 %6 %55
+%85 = OpAtomicIAdd  %3  %36 %57 %59 %41
+%87 = OpAccessChain  %60  %38 %61
+%86 = OpAtomicIAdd  %4  %87 %57 %59 %42
+%89 = OpAccessChain  %34  %40 %35
+%88 = OpAtomicIAdd  %3  %89 %57 %59 %41
+%91 = OpAccessChain  %60  %40 %61 %61
+%90 = OpAtomicIAdd  %4  %91 %57 %59 %42
+%92 = OpAtomicIAdd  %3  %22 %65 %66 %41
+%94 = OpAccessChain  %67  %24 %61
+%93 = OpAtomicIAdd  %4  %94 %65 %66 %42
+%96 = OpAccessChain  %23  %26 %35
+%95 = OpAtomicIAdd  %3  %96 %65 %66 %41
+%98 = OpAccessChain  %67  %26 %61 %61
+%97 = OpAtomicIAdd  %4  %98 %65 %66 %42
+OpControlBarrier %6 %6 %55
+%99 = OpAtomicISub  %3  %36 %57 %59 %41
+%101 = OpAccessChain  %60  %38 %61
+%100 = OpAtomicISub  %4  %101 %57 %59 %42
+%103 = OpAccessChain  %34  %40 %35
+%102 = OpAtomicISub  %3  %103 %57 %59 %41
+%105 = OpAccessChain  %60  %40 %61 %61
+%104 = OpAtomicISub  %4  %105 %57 %59 %42
+%106 = OpAtomicISub  %3  %22 %65 %66 %41
+%108 = OpAccessChain  %67  %24 %61
+%107 = OpAtomicISub  %4  %108 %65 %66 %42
+%110 = OpAccessChain  %23  %26 %35
+%109 = OpAtomicISub  %3  %110 %65 %66 %41
+%112 = OpAccessChain  %67  %26 %61 %61
+%111 = OpAtomicISub  %4  %112 %65 %66 %42
+OpControlBarrier %6 %6 %55
+%113 = OpAtomicUMax  %3  %36 %57 %59 %41
+%115 = OpAccessChain  %60  %38 %61
+%114 = OpAtomicSMax  %4  %115 %57 %59 %42
+%117 = OpAccessChain  %34  %40 %35
+%116 = OpAtomicUMax  %3  %117 %57 %59 %41
+%119 = OpAccessChain  %60  %40 %61 %61
+%118 = OpAtomicSMax  %4  %119 %57 %59 %42
+%120 = OpAtomicUMax  %3  %22 %65 %66 %41
+%122 = OpAccessChain  %67  %24 %61
+%121 = OpAtomicSMax  %4  %122 %65 %66 %42
+%124 = OpAccessChain  %23  %26 %35
+%123 = OpAtomicUMax  %3  %124 %65 %66 %41
+%126 = OpAccessChain  %67  %26 %61 %61
+%125 = OpAtomicSMax  %4  %126 %65 %66 %42
+OpControlBarrier %6 %6 %55
+%127 = OpAtomicUMin  %3  %36 %57 %59 %41
+%129 = OpAccessChain  %60  %38 %61
+%128 = OpAtomicSMin  %4  %129 %57 %59 %42
+%131 = OpAccessChain  %34  %40 %35
+%130 = OpAtomicUMin  %3  %131 %57 %59 %41
+%133 = OpAccessChain  %60  %40 %61 %61
+%132 = OpAtomicSMin  %4  %133 %57 %59 %42
+%134 = OpAtomicUMin  %3  %22 %65 %66 %41
+%136 = OpAccessChain  %67  %24 %61
+%135 = OpAtomicSMin  %4  %136 %65 %66 %42
+%138 = OpAccessChain  %23  %26 %35
+%137 = OpAtomicUMin  %3  %138 %65 %66 %41
+%140 = OpAccessChain  %67  %26 %61 %61
+%139 = OpAtomicSMin  %4  %140 %65 %66 %42
+OpControlBarrier %6 %6 %55
+%141 = OpAtomicAnd  %3  %36 %57 %59 %41
+%143 = OpAccessChain  %60  %38 %61
+%142 = OpAtomicAnd  %4  %143 %57 %59 %42
+%145 = OpAccessChain  %34  %40 %35
+%144 = OpAtomicAnd  %3  %145 %57 %59 %41
+%147 = OpAccessChain  %60  %40 %61 %61
+%146 = OpAtomicAnd  %4  %147 %57 %59 %42
+%148 = OpAtomicAnd  %3  %22 %65 %66 %41
+%150 = OpAccessChain  %67  %24 %61
+%149 = OpAtomicAnd  %4  %150 %65 %66 %42
+%152 = OpAccessChain  %23  %26 %35
+%151 = OpAtomicAnd  %3  %152 %65 %66 %41
+%154 = OpAccessChain  %67  %26 %61 %61
+%153 = OpAtomicAnd  %4  %154 %65 %66 %42
+OpControlBarrier %6 %6 %55
+%155 = OpAtomicOr  %3  %36 %57 %59 %41
+%157 = OpAccessChain  %60  %38 %61
+%156 = OpAtomicOr  %4  %157 %57 %59 %42
+%159 = OpAccessChain  %34  %40 %35
+%158 = OpAtomicOr  %3  %159 %57 %59 %41
+%161 = OpAccessChain  %60  %40 %61 %61
+%160 = OpAtomicOr  %4  %161 %57 %59 %42
+%162 = OpAtomicOr  %3  %22 %65 %66 %41
+%164 = OpAccessChain  %67  %24 %61
+%163 = OpAtomicOr  %4  %164 %65 %66 %42
+%166 = OpAccessChain  %23  %26 %35
+%165 = OpAtomicOr  %3  %166 %65 %66 %41
+%168 = OpAccessChain  %67  %26 %61 %61
+%167 = OpAtomicOr  %4  %168 %65 %66 %42
+OpControlBarrier %6 %6 %55
+%169 = OpAtomicXor  %3  %36 %57 %59 %41
+%171 = OpAccessChain  %60  %38 %61
+%170 = OpAtomicXor  %4  %171 %57 %59 %42
+%173 = OpAccessChain  %34  %40 %35
+%172 = OpAtomicXor  %3  %173 %57 %59 %41
+%175 = OpAccessChain  %60  %40 %61 %61
+%174 = OpAtomicXor  %4  %175 %57 %59 %42
+%176 = OpAtomicXor  %3  %22 %65 %66 %41
+%178 = OpAccessChain  %67  %24 %61
+%177 = OpAtomicXor  %4  %178 %65 %66 %42
+%180 = OpAccessChain  %23  %26 %35
+%179 = OpAtomicXor  %3  %180 %65 %66 %41
+%182 = OpAccessChain  %67  %26 %61 %61
+%181 = OpAtomicXor  %4  %182 %65 %66 %42
+%183 = OpAtomicExchange  %3  %36 %57 %59 %41
+%185 = OpAccessChain  %60  %38 %61
+%184 = OpAtomicExchange  %4  %185 %57 %59 %42
+%187 = OpAccessChain  %34  %40 %35
+%186 = OpAtomicExchange  %3  %187 %57 %59 %41
+%189 = OpAccessChain  %60  %40 %61 %61
+%188 = OpAtomicExchange  %4  %189 %57 %59 %42
+%190 = OpAtomicExchange  %3  %22 %65 %66 %41
+%192 = OpAccessChain  %67  %24 %61
+%191 = OpAtomicExchange  %4  %192 %65 %66 %42
+%194 = OpAccessChain  %23  %26 %35
+%193 = OpAtomicExchange  %3  %194 %65 %66 %41
+%196 = OpAccessChain  %67  %26 %61 %61
+%195 = OpAtomicExchange  %4  %196 %65 %66 %42
+%198 = OpAtomicCompareExchange  %3  %36 %57 %59 %59 %43 %41
+%199 = OpIEqual  %10  %198 %41
+%197 = OpCompositeConstruct  %11  %198 %199
+%201 = OpAccessChain  %60  %38 %61
+%202 = OpAtomicCompareExchange  %4  %201 %57 %59 %59 %44 %42
+%203 = OpIEqual  %10  %202 %42
+%200 = OpCompositeConstruct  %12  %202 %203
+%205 = OpAccessChain  %34  %40 %35
+%206 = OpAtomicCompareExchange  %3  %205 %57 %59 %59 %43 %41
+%207 = OpIEqual  %10  %206 %41
+%204 = OpCompositeConstruct  %11  %206 %207
+%209 = OpAccessChain  %60  %40 %61 %61
+%210 = OpAtomicCompareExchange  %4  %209 %57 %59 %59 %44 %42
+%211 = OpIEqual  %10  %210 %42
+%208 = OpCompositeConstruct  %12  %210 %211
+%213 = OpAtomicCompareExchange  %3  %22 %65 %66 %66 %43 %41
+%214 = OpIEqual  %10  %213 %41
+%212 = OpCompositeConstruct  %11  %213 %214
+%216 = OpAccessChain  %67  %24 %61
+%217 = OpAtomicCompareExchange  %4  %216 %65 %66 %66 %44 %42
+%218 = OpIEqual  %10  %217 %42
+%215 = OpCompositeConstruct  %12  %217 %218
+%220 = OpAccessChain  %23  %26 %35
+%221 = OpAtomicCompareExchange  %3  %220 %65 %66 %66 %43 %41
+%222 = OpIEqual  %10  %221 %41
+%219 = OpCompositeConstruct  %11  %221 %222
+%224 = OpAccessChain  %67  %26 %61 %61
+%225 = OpAtomicCompareExchange  %4  %224 %65 %66 %66 %44 %42
+%226 = OpIEqual  %10  %225 %42
+%223 = OpCompositeConstruct  %12  %225 %226
 OpReturn
 OpFunctionEnd
\ No newline at end of file
diff --git a/naga/tests/out/spv/wgsl-atomicOps.spvasm b/naga/tests/out/spv/wgsl-atomicOps.spvasm
index de4d687824..f000dc21fd 100644
--- a/naga/tests/out/spv/wgsl-atomicOps.spvasm
+++ b/naga/tests/out/spv/wgsl-atomicOps.spvasm
@@ -1,29 +1,33 @@
 ; SPIR-V
 ; Version: 1.1
 ; Generator: rspirv
-; Bound: 189
+; Bound: 221
 OpCapability Shader
 OpExtension "SPV_KHR_storage_buffer_storage_class"
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
-OpEntryPoint GLCompute %28 "cs_main" %25
-OpExecutionMode %28 LocalSize 2 1 1
+OpEntryPoint GLCompute %31 "cs_main" %28
+OpExecutionMode %31 LocalSize 2 1 1
 OpDecorate %5 ArrayStride 4
 OpMemberDecorate %7 0 Offset 0
 OpMemberDecorate %7 1 Offset 4
-OpDecorate %9 DescriptorSet 0
-OpDecorate %9 Binding 0
-OpDecorate %10 Block
 OpMemberDecorate %10 0 Offset 0
+OpMemberDecorate %10 1 Offset 4
+OpMemberDecorate %11 0 Offset 0
+OpMemberDecorate %11 1 Offset 4
 OpDecorate %12 DescriptorSet 0
-OpDecorate %12 Binding 1
+OpDecorate %12 Binding 0
 OpDecorate %13 Block
 OpMemberDecorate %13 0 Offset 0
 OpDecorate %15 DescriptorSet 0
-OpDecorate %15 Binding 2
+OpDecorate %15 Binding 1
 OpDecorate %16 Block
 OpMemberDecorate %16 0 Offset 0
-OpDecorate %25 BuiltIn LocalInvocationId
+OpDecorate %18 DescriptorSet 0
+OpDecorate %18 Binding 2
+OpDecorate %19 Block
+OpMemberDecorate %19 0 Offset 0
+OpDecorate %28 BuiltIn LocalInvocationId
 %2 = OpTypeVoid
 %3 = OpTypeInt 32 0
 %4 = OpTypeInt 32 1
@@ -31,210 +35,242 @@ OpDecorate %25 BuiltIn LocalInvocationId
 %5 = OpTypeArray %4 %6
 %7 = OpTypeStruct %3 %5
 %8 = OpTypeVector %3 3
-%10 = OpTypeStruct %3
-%11 = OpTypePointer StorageBuffer %10
-%9 = OpVariable  %11  StorageBuffer
-%13 = OpTypeStruct %5
+%9 = OpTypeBool
+%10 = OpTypeStruct %3 %9
+%11 = OpTypeStruct %4 %9
+%13 = OpTypeStruct %3
 %14 = OpTypePointer StorageBuffer %13
 %12 = OpVariable  %14  StorageBuffer
-%16 = OpTypeStruct %7
+%16 = OpTypeStruct %5
 %17 = OpTypePointer StorageBuffer %16
 %15 = OpVariable  %17  StorageBuffer
-%19 = OpTypePointer Workgroup %3
-%18 = OpVariable  %19  Workgroup
-%21 = OpTypePointer Workgroup %5
-%20 = OpVariable  %21  Workgroup
-%23 = OpTypePointer Workgroup %7
-%22 = OpVariable  %23  Workgroup
-%26 = OpTypePointer Input %8
-%25 = OpVariable  %26  Input
-%29 = OpTypeFunction %2
-%30 = OpTypePointer StorageBuffer %3
-%31 = OpConstant  %3  0
-%33 = OpTypePointer StorageBuffer %5
-%35 = OpTypePointer StorageBuffer %7
-%37 = OpConstant  %3  1
-%38 = OpConstant  %4  1
-%40 = OpConstantNull  %3
-%41 = OpConstantNull  %5
-%42 = OpConstantNull  %7
-%43 = OpConstantNull  %8
-%45 = OpTypeBool
-%44 = OpTypeVector %45 3
-%50 = OpConstant  %3  264
-%52 = OpConstant  %3  64
-%53 = OpTypePointer StorageBuffer %4
-%57 = OpConstant  %4  2
-%58 = OpConstant  %3  256
-%59 = OpTypePointer Workgroup %4
-%28 = OpFunction  %2  None %29
-%24 = OpLabel
-%27 = OpLoad  %8  %25
-%32 = OpAccessChain  %30  %9 %31
-%34 = OpAccessChain  %33  %12 %31
-%36 = OpAccessChain  %35  %15 %31
-OpBranch %39
-%39 = OpLabel
-%46 = OpIEqual  %44  %27 %43
-%47 = OpAll  %45  %46
-OpSelectionMerge %48 None
-OpBranchConditional %47 %49 %48
-%49 = OpLabel
-OpStore %18 %40
-OpStore %20 %41
-OpStore %22 %42
-OpBranch %48
-%48 = OpLabel
-OpControlBarrier %6 %6 %50
+%19 = OpTypeStruct %7
+%20 = OpTypePointer StorageBuffer %19
+%18 = OpVariable  %20  StorageBuffer
+%22 = OpTypePointer Workgroup %3
+%21 = OpVariable  %22  Workgroup
+%24 = OpTypePointer Workgroup %5
+%23 = OpVariable  %24  Workgroup
+%26 = OpTypePointer Workgroup %7
+%25 = OpVariable  %26  Workgroup
+%29 = OpTypePointer Input %8
+%28 = OpVariable  %29  Input
+%32 = OpTypeFunction %2
+%33 = OpTypePointer StorageBuffer %3
+%34 = OpConstant  %3  0
+%36 = OpTypePointer StorageBuffer %5
+%38 = OpTypePointer StorageBuffer %7
+%40 = OpConstant  %3  1
+%41 = OpConstant  %4  1
+%42 = OpConstant  %4  2
+%44 = OpConstantNull  %3
+%45 = OpConstantNull  %5
+%46 = OpConstantNull  %7
+%47 = OpConstantNull  %8
+%48 = OpTypeVector %9 3
+%53 = OpConstant  %3  264
+%55 = OpConstant  %3  64
+%56 = OpTypePointer StorageBuffer %4
+%60 = OpConstant  %3  256
+%61 = OpTypePointer Workgroup %4
+%31 = OpFunction  %2  None %32
+%27 = OpLabel
+%30 = OpLoad  %8  %28
+%35 = OpAccessChain  %33  %12 %34
+%37 = OpAccessChain  %36  %15 %34
+%39 = OpAccessChain  %38  %18 %34
+OpBranch %43
+%43 = OpLabel
+%49 = OpIEqual  %48  %30 %47
+%50 = OpAll  %9  %49
+OpSelectionMerge %51 None
+OpBranchConditional %50 %52 %51
+%52 = OpLabel
+OpStore %21 %44
+OpStore %23 %45
+OpStore %25 %46
 OpBranch %51
 %51 = OpLabel
-OpAtomicStore %32 %38 %52 %37
-%54 = OpAccessChain  %53  %34 %37
-OpAtomicStore %54 %38 %52 %38
-%55 = OpAccessChain  %30  %36 %31
-OpAtomicStore %55 %38 %52 %37
-%56 = OpAccessChain  %53  %36 %37 %37
-OpAtomicStore %56 %38 %52 %38
-OpAtomicStore %18 %57 %58 %37
-%60 = OpAccessChain  %59  %20 %37
-OpAtomicStore %60 %57 %58 %38
-%61 = OpAccessChain  %19  %22 %31
-OpAtomicStore %61 %57 %58 %37
-%62 = OpAccessChain  %59  %22 %37 %37
-OpAtomicStore %62 %57 %58 %38
-OpControlBarrier %6 %6 %50
-%63 = OpAtomicLoad  %3  %32 %38 %52
-%64 = OpAccessChain  %53  %34 %37
-%65 = OpAtomicLoad  %4  %64 %38 %52
-%66 = OpAccessChain  %30  %36 %31
-%67 = OpAtomicLoad  %3  %66 %38 %52
-%68 = OpAccessChain  %53  %36 %37 %37
-%69 = OpAtomicLoad  %4  %68 %38 %52
-%70 = OpAtomicLoad  %3  %18 %57 %58
-%71 = OpAccessChain  %59  %20 %37
-%72 = OpAtomicLoad  %4  %71 %57 %58
-%73 = OpAccessChain  %19  %22 %31
-%74 = OpAtomicLoad  %3  %73 %57 %58
-%75 = OpAccessChain  %59  %22 %37 %37
-%76 = OpAtomicLoad  %4  %75 %57 %58
-OpControlBarrier %6 %6 %50
-%77 = OpAtomicIAdd  %3  %32 %38 %52 %37
-%79 = OpAccessChain  %53  %34 %37
-%78 = OpAtomicIAdd  %4  %79 %38 %52 %38
-%81 = OpAccessChain  %30  %36 %31
-%80 = OpAtomicIAdd  %3  %81 %38 %52 %37
-%83 = OpAccessChain  %53  %36 %37 %37
-%82 = OpAtomicIAdd  %4  %83 %38 %52 %38
-%84 = OpAtomicIAdd  %3  %18 %57 %58 %37
-%86 = OpAccessChain  %59  %20 %37
-%85 = OpAtomicIAdd  %4  %86 %57 %58 %38
-%88 = OpAccessChain  %19  %22 %31
-%87 = OpAtomicIAdd  %3  %88 %57 %58 %37
-%90 = OpAccessChain  %59  %22 %37 %37
-%89 = OpAtomicIAdd  %4  %90 %57 %58 %38
-OpControlBarrier %6 %6 %50
-%91 = OpAtomicISub  %3  %32 %38 %52 %37
-%93 = OpAccessChain  %53  %34 %37
-%92 = OpAtomicISub  %4  %93 %38 %52 %38
-%95 = OpAccessChain  %30  %36 %31
-%94 = OpAtomicISub  %3  %95 %38 %52 %37
-%97 = OpAccessChain  %53  %36 %37 %37
-%96 = OpAtomicISub  %4  %97 %38 %52 %38
-%98 = OpAtomicISub  %3  %18 %57 %58 %37
-%100 = OpAccessChain  %59  %20 %37
-%99 = OpAtomicISub  %4  %100 %57 %58 %38
-%102 = OpAccessChain  %19  %22 %31
-%101 = OpAtomicISub  %3  %102 %57 %58 %37
-%104 = OpAccessChain  %59  %22 %37 %37
-%103 = OpAtomicISub  %4  %104 %57 %58 %38
-OpControlBarrier %6 %6 %50
-%105 = OpAtomicUMax  %3  %32 %38 %52 %37
-%107 = OpAccessChain  %53  %34 %37
-%106 = OpAtomicSMax  %4  %107 %38 %52 %38
-%109 = OpAccessChain  %30  %36 %31
-%108 = OpAtomicUMax  %3  %109 %38 %52 %37
-%111 = OpAccessChain  %53  %36 %37 %37
-%110 = OpAtomicSMax  %4  %111 %38 %52 %38
-%112 = OpAtomicUMax  %3  %18 %57 %58 %37
-%114 = OpAccessChain  %59  %20 %37
-%113 = OpAtomicSMax  %4  %114 %57 %58 %38
-%116 = OpAccessChain  %19  %22 %31
-%115 = OpAtomicUMax  %3  %116 %57 %58 %37
-%118 = OpAccessChain  %59  %22 %37 %37
-%117 = OpAtomicSMax  %4  %118 %57 %58 %38
-OpControlBarrier %6 %6 %50
-%119 = OpAtomicUMin  %3  %32 %38 %52 %37
-%121 = OpAccessChain  %53  %34 %37
-%120 = OpAtomicSMin  %4  %121 %38 %52 %38
-%123 = OpAccessChain  %30  %36 %31
-%122 = OpAtomicUMin  %3  %123 %38 %52 %37
-%125 = OpAccessChain  %53  %36 %37 %37
-%124 = OpAtomicSMin  %4  %125 %38 %52 %38
-%126 = OpAtomicUMin  %3  %18 %57 %58 %37
-%128 = OpAccessChain  %59  %20 %37
-%127 = OpAtomicSMin  %4  %128 %57 %58 %38
-%130 = OpAccessChain  %19  %22 %31
-%129 = OpAtomicUMin  %3  %130 %57 %58 %37
-%132 = OpAccessChain  %59  %22 %37 %37
-%131 = OpAtomicSMin  %4  %132 %57 %58 %38
-OpControlBarrier %6 %6 %50
-%133 = OpAtomicAnd  %3  %32 %38 %52 %37
-%135 = OpAccessChain  %53  %34 %37
-%134 = OpAtomicAnd  %4  %135 %38 %52 %38
-%137 = OpAccessChain  %30  %36 %31
-%136 = OpAtomicAnd  %3  %137 %38 %52 %37
-%139 = OpAccessChain  %53  %36 %37 %37
-%138 = OpAtomicAnd  %4  %139 %38 %52 %38
-%140 = OpAtomicAnd  %3  %18 %57 %58 %37
-%142 = OpAccessChain  %59  %20 %37
-%141 = OpAtomicAnd  %4  %142 %57 %58 %38
-%144 = OpAccessChain  %19  %22 %31
-%143 = OpAtomicAnd  %3  %144 %57 %58 %37
-%146 = OpAccessChain  %59  %22 %37 %37
-%145 = OpAtomicAnd  %4  %146 %57 %58 %38
-OpControlBarrier %6 %6 %50
-%147 = OpAtomicOr  %3  %32 %38 %52 %37
-%149 = OpAccessChain  %53  %34 %37
-%148 = OpAtomicOr  %4  %149 %38 %52 %38
-%151 = OpAccessChain  %30  %36 %31
-%150 = OpAtomicOr  %3  %151 %38 %52 %37
-%153 = OpAccessChain  %53  %36 %37 %37
-%152 = OpAtomicOr  %4  %153 %38 %52 %38
-%154 = OpAtomicOr  %3  %18 %57 %58 %37
-%156 = OpAccessChain  %59  %20 %37
-%155 = OpAtomicOr  %4  %156 %57 %58 %38
-%158 = OpAccessChain  %19  %22 %31
-%157 = OpAtomicOr  %3  %158 %57 %58 %37
-%160 = OpAccessChain  %59  %22 %37 %37
-%159 = OpAtomicOr  %4  %160 %57 %58 %38
-OpControlBarrier %6 %6 %50
-%161 = OpAtomicXor  %3  %32 %38 %52 %37
-%163 = OpAccessChain  %53  %34 %37
-%162 = OpAtomicXor  %4  %163 %38 %52 %38
-%165 = OpAccessChain  %30  %36 %31
-%164 = OpAtomicXor  %3  %165 %38 %52 %37
-%167 = OpAccessChain  %53  %36 %37 %37
-%166 = OpAtomicXor  %4  %167 %38 %52 %38
-%168 = OpAtomicXor  %3  %18 %57 %58 %37
-%170 = OpAccessChain  %59  %20 %37
-%169 = OpAtomicXor  %4  %170 %57 %58 %38
-%172 = OpAccessChain  %19  %22 %31
-%171 = OpAtomicXor  %3  %172 %57 %58 %37
-%174 = OpAccessChain  %59  %22 %37 %37
-%173 = OpAtomicXor  %4  %174 %57 %58 %38
-%175 = OpAtomicExchange  %3  %32 %38 %52 %37
-%177 = OpAccessChain  %53  %34 %37
-%176 = OpAtomicExchange  %4  %177 %38 %52 %38
-%179 = OpAccessChain  %30  %36 %31
-%178 = OpAtomicExchange  %3  %179 %38 %52 %37
-%181 = OpAccessChain  %53  %36 %37 %37
-%180 = OpAtomicExchange  %4  %181 %38 %52 %38
-%182 = OpAtomicExchange  %3  %18 %57 %58 %37
-%184 = OpAccessChain  %59  %20 %37
-%183 = OpAtomicExchange  %4  %184 %57 %58 %38
-%186 = OpAccessChain  %19  %22 %31
-%185 = OpAtomicExchange  %3  %186 %57 %58 %37
-%188 = OpAccessChain  %59  %22 %37 %37
-%187 = OpAtomicExchange  %4  %188 %57 %58 %38
+OpControlBarrier %6 %6 %53
+OpBranch %54
+%54 = OpLabel
+OpAtomicStore %35 %41 %55 %40
+%57 = OpAccessChain  %56  %37 %40
+OpAtomicStore %57 %41 %55 %41
+%58 = OpAccessChain  %33  %39 %34
+OpAtomicStore %58 %41 %55 %40
+%59 = OpAccessChain  %56  %39 %40 %40
+OpAtomicStore %59 %41 %55 %41
+OpAtomicStore %21 %42 %60 %40
+%62 = OpAccessChain  %61  %23 %40
+OpAtomicStore %62 %42 %60 %41
+%63 = OpAccessChain  %22  %25 %34
+OpAtomicStore %63 %42 %60 %40
+%64 = OpAccessChain  %61  %25 %40 %40
+OpAtomicStore %64 %42 %60 %41
+OpControlBarrier %6 %6 %53
+%65 = OpAtomicLoad  %3  %35 %41 %55
+%66 = OpAccessChain  %56  %37 %40
+%67 = OpAtomicLoad  %4  %66 %41 %55
+%68 = OpAccessChain  %33  %39 %34
+%69 = OpAtomicLoad  %3  %68 %41 %55
+%70 = OpAccessChain  %56  %39 %40 %40
+%71 = OpAtomicLoad  %4  %70 %41 %55
+%72 = OpAtomicLoad  %3  %21 %42 %60
+%73 = OpAccessChain  %61  %23 %40
+%74 = OpAtomicLoad  %4  %73 %42 %60
+%75 = OpAccessChain  %22  %25 %34
+%76 = OpAtomicLoad  %3  %75 %42 %60
+%77 = OpAccessChain  %61  %25 %40 %40
+%78 = OpAtomicLoad  %4  %77 %42 %60
+OpControlBarrier %6 %6 %53
+%79 = OpAtomicIAdd  %3  %35 %41 %55 %40
+%81 = OpAccessChain  %56  %37 %40
+%80 = OpAtomicIAdd  %4  %81 %41 %55 %41
+%83 = OpAccessChain  %33  %39 %34
+%82 = OpAtomicIAdd  %3  %83 %41 %55 %40
+%85 = OpAccessChain  %56  %39 %40 %40
+%84 = OpAtomicIAdd  %4  %85 %41 %55 %41
+%86 = OpAtomicIAdd  %3  %21 %42 %60 %40
+%88 = OpAccessChain  %61  %23 %40
+%87 = OpAtomicIAdd  %4  %88 %42 %60 %41
+%90 = OpAccessChain  %22  %25 %34
+%89 = OpAtomicIAdd  %3  %90 %42 %60 %40
+%92 = OpAccessChain  %61  %25 %40 %40
+%91 = OpAtomicIAdd  %4  %92 %42 %60 %41
+OpControlBarrier %6 %6 %53
+%93 = OpAtomicISub  %3  %35 %41 %55 %40
+%95 = OpAccessChain  %56  %37 %40
+%94 = OpAtomicISub  %4  %95 %41 %55 %41
+%97 = OpAccessChain  %33  %39 %34
+%96 = OpAtomicISub  %3  %97 %41 %55 %40
+%99 = OpAccessChain  %56  %39 %40 %40
+%98 = OpAtomicISub  %4  %99 %41 %55 %41
+%100 = OpAtomicISub  %3  %21 %42 %60 %40
+%102 = OpAccessChain  %61  %23 %40
+%101 = OpAtomicISub  %4  %102 %42 %60 %41
+%104 = OpAccessChain  %22  %25 %34
+%103 = OpAtomicISub  %3  %104 %42 %60 %40
+%106 = OpAccessChain  %61  %25 %40 %40
+%105 = OpAtomicISub  %4  %106 %42 %60 %41
+OpControlBarrier %6 %6 %53
+%107 = OpAtomicUMax  %3  %35 %41 %55 %40
+%109 = OpAccessChain  %56  %37 %40
+%108 = OpAtomicSMax  %4  %109 %41 %55 %41
+%111 = OpAccessChain  %33  %39 %34
+%110 = OpAtomicUMax  %3  %111 %41 %55 %40
+%113 = OpAccessChain  %56  %39 %40 %40
+%112 = OpAtomicSMax  %4  %113 %41 %55 %41
+%114 = OpAtomicUMax  %3  %21 %42 %60 %40
+%116 = OpAccessChain  %61  %23 %40
+%115 = OpAtomicSMax  %4  %116 %42 %60 %41
+%118 = OpAccessChain  %22  %25 %34
+%117 = OpAtomicUMax  %3  %118 %42 %60 %40
+%120 = OpAccessChain  %61  %25 %40 %40
+%119 = OpAtomicSMax  %4  %120 %42 %60 %41
+OpControlBarrier %6 %6 %53
+%121 = OpAtomicUMin  %3  %35 %41 %55 %40
+%123 = OpAccessChain  %56  %37 %40
+%122 = OpAtomicSMin  %4  %123 %41 %55 %41
+%125 = OpAccessChain  %33  %39 %34
+%124 = OpAtomicUMin  %3  %125 %41 %55 %40
+%127 = OpAccessChain  %56  %39 %40 %40
+%126 = OpAtomicSMin  %4  %127 %41 %55 %41
+%128 = OpAtomicUMin  %3  %21 %42 %60 %40
+%130 = OpAccessChain  %61  %23 %40
+%129 = OpAtomicSMin  %4  %130 %42 %60 %41
+%132 = OpAccessChain  %22  %25 %34
+%131 = OpAtomicUMin  %3  %132 %42 %60 %40
+%134 = OpAccessChain  %61  %25 %40 %40
+%133 = OpAtomicSMin  %4  %134 %42 %60 %41
+OpControlBarrier %6 %6 %53
+%135 = OpAtomicAnd  %3  %35 %41 %55 %40
+%137 = OpAccessChain  %56  %37 %40
+%136 = OpAtomicAnd  %4  %137 %41 %55 %41
+%139 = OpAccessChain  %33  %39 %34
+%138 = OpAtomicAnd  %3  %139 %41 %55 %40
+%141 = OpAccessChain  %56  %39 %40 %40
+%140 = OpAtomicAnd  %4  %141 %41 %55 %41
+%142 = OpAtomicAnd  %3  %21 %42 %60 %40
+%144 = OpAccessChain  %61  %23 %40
+%143 = OpAtomicAnd  %4  %144 %42 %60 %41
+%146 = OpAccessChain  %22  %25 %34
+%145 = OpAtomicAnd  %3  %146 %42 %60 %40
+%148 = OpAccessChain  %61  %25 %40 %40
+%147 = OpAtomicAnd  %4  %148 %42 %60 %41
+OpControlBarrier %6 %6 %53
+%149 = OpAtomicOr  %3  %35 %41 %55 %40
+%151 = OpAccessChain  %56  %37 %40
+%150 = OpAtomicOr  %4  %151 %41 %55 %41
+%153 = OpAccessChain  %33  %39 %34
+%152 = OpAtomicOr  %3  %153 %41 %55 %40
+%155 = OpAccessChain  %56  %39 %40 %40
+%154 = OpAtomicOr  %4  %155 %41 %55 %41
+%156 = OpAtomicOr  %3  %21 %42 %60 %40
+%158 = OpAccessChain  %61  %23 %40
+%157 = OpAtomicOr  %4  %158 %42 %60 %41
+%160 = OpAccessChain  %22  %25 %34
+%159 = OpAtomicOr  %3  %160 %42 %60 %40
+%162 = OpAccessChain  %61  %25 %40 %40
+%161 = OpAtomicOr  %4  %162 %42 %60 %41
+OpControlBarrier %6 %6 %53
+%163 = OpAtomicXor  %3  %35 %41 %55 %40
+%165 = OpAccessChain  %56  %37 %40
+%164 = OpAtomicXor  %4  %165 %41 %55 %41
+%167 = OpAccessChain  %33  %39 %34
+%166 = OpAtomicXor  %3  %167 %41 %55 %40
+%169 = OpAccessChain  %56  %39 %40 %40
+%168 = OpAtomicXor  %4  %169 %41 %55 %41
+%170 = OpAtomicXor  %3  %21 %42 %60 %40
+%172 = OpAccessChain  %61  %23 %40
+%171 = OpAtomicXor  %4  %172 %42 %60 %41
+%174 = OpAccessChain  %22  %25 %34
+%173 = OpAtomicXor  %3  %174 %42 %60 %40
+%176 = OpAccessChain  %61  %25 %40 %40
+%175 = OpAtomicXor  %4  %176 %42 %60 %41
+%177 = OpAtomicExchange  %3  %35 %41 %55 %40
+%179 = OpAccessChain  %56  %37 %40
+%178 = OpAtomicExchange  %4  %179 %41 %55 %41
+%181 = OpAccessChain  %33  %39 %34
+%180 = OpAtomicExchange  %3  %181 %41 %55 %40
+%183 = OpAccessChain  %56  %39 %40 %40
+%182 = OpAtomicExchange  %4  %183 %41 %55 %41
+%184 = OpAtomicExchange  %3  %21 %42 %60 %40
+%186 = OpAccessChain  %61  %23 %40
+%185 = OpAtomicExchange  %4  %186 %42 %60 %41
+%188 = OpAccessChain  %22  %25 %34
+%187 = OpAtomicExchange  %3  %188 %42 %60 %40
+%190 = OpAccessChain  %61  %25 %40 %40
+%189 = OpAtomicExchange  %4  %190 %42 %60 %41
+%192 = OpAtomicCompareExchange  %3  %35 %41 %55 %55 %6 %40
+%193 = OpIEqual  %9  %192 %40
+%191 = OpCompositeConstruct  %10  %192 %193
+%195 = OpAccessChain  %56  %37 %40
+%196 = OpAtomicCompareExchange  %4  %195 %41 %55 %55 %42 %41
+%197 = OpIEqual  %9  %196 %41
+%194 = OpCompositeConstruct  %11  %196 %197
+%199 = OpAccessChain  %33  %39 %34
+%200 = OpAtomicCompareExchange  %3  %199 %41 %55 %55 %6 %40
+%201 = OpIEqual  %9  %200 %40
+%198 = OpCompositeConstruct  %10  %200 %201
+%203 = OpAccessChain  %56  %39 %40 %40
+%204 = OpAtomicCompareExchange  %4  %203 %41 %55 %55 %42 %41
+%205 = OpIEqual  %9  %204 %41
+%202 = OpCompositeConstruct  %11  %204 %205
+%207 = OpAtomicCompareExchange  %3  %21 %42 %60 %60 %6 %40
+%208 = OpIEqual  %9  %207 %40
+%206 = OpCompositeConstruct  %10  %207 %208
+%210 = OpAccessChain  %61  %23 %40
+%211 = OpAtomicCompareExchange  %4  %210 %42 %60 %60 %42 %41
+%212 = OpIEqual  %9  %211 %41
+%209 = OpCompositeConstruct  %11  %211 %212
+%214 = OpAccessChain  %22  %25 %34
+%215 = OpAtomicCompareExchange  %3  %214 %42 %60 %60 %6 %40
+%216 = OpIEqual  %9  %215 %40
+%213 = OpCompositeConstruct  %10  %215 %216
+%218 = OpAccessChain  %61  %25 %40 %40
+%219 = OpAtomicCompareExchange  %4  %218 %42 %60 %60 %42 %41
+%220 = OpIEqual  %9  %219 %41
+%217 = OpCompositeConstruct  %11  %219 %220
 OpReturn
 OpFunctionEnd
\ No newline at end of file
diff --git a/naga/tests/out/wgsl/wgsl-atomicOps-int64.wgsl b/naga/tests/out/wgsl/wgsl-atomicOps-int64.wgsl
index 364108c646..7b23ac02a9 100644
--- a/naga/tests/out/wgsl/wgsl-atomicOps-int64.wgsl
+++ b/naga/tests/out/wgsl/wgsl-atomicOps-int64.wgsl
@@ -103,5 +103,13 @@ fn cs_main(@builtin(local_invocation_id) id: vec3<u32>) {
     let _e279 = atomicExchange((&workgroup_atomic_arr[1]), 1li);
     let _e283 = atomicExchange((&workgroup_struct.atomic_scalar), 1lu);
     let _e288 = atomicExchange((&workgroup_struct.atomic_arr[1]), 1li);
+    let _e292 = atomicCompareExchangeWeak((&storage_atomic_scalar), 1lu, 2lu);
+    let _e297 = atomicCompareExchangeWeak((&storage_atomic_arr[1]), 1li, 2li);
+    let _e302 = atomicCompareExchangeWeak((&storage_struct.atomic_scalar), 1lu, 2lu);
+    let _e308 = atomicCompareExchangeWeak((&storage_struct.atomic_arr[1]), 1li, 2li);
+    let _e312 = atomicCompareExchangeWeak((&workgroup_atomic_scalar), 1lu, 2lu);
+    let _e317 = atomicCompareExchangeWeak((&workgroup_atomic_arr[1]), 1li, 2li);
+    let _e322 = atomicCompareExchangeWeak((&workgroup_struct.atomic_scalar), 1lu, 2lu);
+    let _e328 = atomicCompareExchangeWeak((&workgroup_struct.atomic_arr[1]), 1li, 2li);
     return;
 }
diff --git a/naga/tests/out/wgsl/wgsl-atomicOps.wgsl b/naga/tests/out/wgsl/wgsl-atomicOps.wgsl
index be102e6833..f35acc4e72 100644
--- a/naga/tests/out/wgsl/wgsl-atomicOps.wgsl
+++ b/naga/tests/out/wgsl/wgsl-atomicOps.wgsl
@@ -103,5 +103,13 @@ fn cs_main(@builtin(local_invocation_id) id: vec3<u32>) {
     let _e295 = atomicExchange((&workgroup_atomic_arr[1]), 1i);
     let _e299 = atomicExchange((&workgroup_struct.atomic_scalar), 1u);
     let _e304 = atomicExchange((&workgroup_struct.atomic_arr[1]), 1i);
+    let _e308 = atomicCompareExchangeWeak((&storage_atomic_scalar), 1u, 2u);
+    let _e313 = atomicCompareExchangeWeak((&storage_atomic_arr[1]), 1i, 2i);
+    let _e318 = atomicCompareExchangeWeak((&storage_struct.atomic_scalar), 1u, 2u);
+    let _e324 = atomicCompareExchangeWeak((&storage_struct.atomic_arr[1]), 1i, 2i);
+    let _e328 = atomicCompareExchangeWeak((&workgroup_atomic_scalar), 1u, 2u);
+    let _e333 = atomicCompareExchangeWeak((&workgroup_atomic_arr[1]), 1i, 2i);
+    let _e338 = atomicCompareExchangeWeak((&workgroup_struct.atomic_scalar), 1u, 2u);
+    let _e344 = atomicCompareExchangeWeak((&workgroup_struct.atomic_arr[1]), 1i, 2i);
     return;
 }