Repository: gfx-rs/wgpu
Pull request #6387 (2024-10-09 17:05:28)

[naga spv-out] Expand LocalType to permit pointers to matrices.
In `back::spv`:

- Factor out the numeric variants of `LocalType` into a new enum, `NumericType`.

- Split the `Value` variant into `Numeric` and `LocalPointer` variants, and let `LocalPointer` point to any numeric type, including matrices.

  In subsequent PRs, we'll need to spill matrices out into temporary local variables. This means we'll need to generate SPIR-V pointer-to-matrix types, so `LocalType` needs to be able to represent that.

- Rename `make_local` to `LocalType::from_inner`.

  Change the free function `back::spv::make_local` into an associated function `LocalType::from_inner`.


jimblandy: This PR is big, but boring. All it's doing is rearranging some enums... that happen to be widely used :sob: 

Patch:
diff --git a/CHANGELOG.md b/CHANGELOG.md
index a9756151a..80b811e89 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -93,6 +93,10 @@ By @bradwerth [#6216](https://github.com/gfx-rs/wgpu/pull/6216).
 
 - Allow using [VK_GOOGLE_display_timing](https://registry.khronos.org/vulkan/specs/1.3-extensions/man/html/VK_GOOGLE_display_timing.html) unsafely with the `VULKAN_GOOGLE_DISPLAY_TIMING` feature. By @DJMcNab in [#6149](https://github.com/gfx-rs/wgpu/pull/6149)
 
+#### Metal
+
+- Implement `atomicCompareExchangeWeak`. By @AsherJingkongChen in [#6265](https://github.com/gfx-rs/wgpu/pull/6265)
+
 ### Bug Fixes
 
 - Fix incorrect hlsl image output type conversion. By @atlv24 in [#6123](https://github.com/gfx-rs/wgpu/pull/6123)
@@ -123,6 +127,9 @@ By @bradwerth [#6216](https://github.com/gfx-rs/wgpu/pull/6216).
 - Document `wgpu_hal` bounds-checking promises, and adapt `wgpu_core`'s lazy initialization logic to the slightly weaker-than-expected guarantees. By @jimblandy in [#6201](https://github.com/gfx-rs/wgpu/pull/6201)
 - Raise validation error instead of panicking in `{Render,Compute}Pipeline::get_bind_group_layout` on native / WebGL. By @bgr360 in [#6280](https://github.com/gfx-rs/wgpu/pull/6280).
 - **BREAKING**: Remove the last exposed C symbols in project, located in `wgpu_core::render::bundle::bundle_ffi`, to allow multiple versions of WGPU to compile together. By @ErichDonGubler in [#6272](https://github.com/gfx-rs/wgpu/pull/6272).
+- Call `flush_mapped_ranges` when unmapping write-mapped buffers. By @teoxoy in [#6089](https://github.com/gfx-rs/wgpu/pull/6089).
+- When mapping buffers for reading, mark buffers as initialized only when they have `MAP_WRITE` usage. By @teoxoy in [#6178](https://github.com/gfx-rs/wgpu/pull/6178).
+- Add a separate pipeline constants error. By @teoxoy in [#6094](https://github.com/gfx-rs/wgpu/pull/6094).
 
 #### GLES / OpenGL
 
@@ -143,6 +150,16 @@ By @bradwerth [#6216](https://github.com/gfx-rs/wgpu/pull/6216).
 - `wgpu_hal::gles::Adapter::new_external` now requires the context to be current when dropping the adapter and related objects. By @Imberflur in [#6114](https://github.com/gfx-rs/wgpu/pull/6114).
 - Reduce the amount of debug and trace logs emitted by wgpu-core and wgpu-hal. By @nical in [#6065](https://github.com/gfx-rs/wgpu/issues/6065)
 - `Rg11b10Float` is renamed to `Rg11b10Ufloat`. By @sagudev in [#6108](https://github.com/gfx-rs/wgpu/pull/6108)
+- Invalidate the device when we encounter driver-induced device loss or on unexpected errors. By @teoxoy in [#6229](https://github.com/gfx-rs/wgpu/pull/6229).
+- Make Vulkan error handling more robust. By @teoxoy in [#6119](https://github.com/gfx-rs/wgpu/pull/6119).
+
+#### Internal
+
+- Tracker simplifications. By @teoxoy in [#6073](https://github.com/gfx-rs/wgpu/pull/6073) & [#6088](https://github.com/gfx-rs/wgpu/pull/6088).
+- D3D12 cleanup. By @teoxoy in [#6200](https://github.com/gfx-rs/wgpu/pull/6200).
+- Use `ManuallyDrop` in remaining places. By @teoxoy in [#6092](https://github.com/gfx-rs/wgpu/pull/6092).
+- Move out invalidity from the `Registry`. By @teoxoy in [#6243](https://github.com/gfx-rs/wgpu/pull/6243).
+- Remove `backend` from ID. By @teoxoy in [#6263](https://github.com/gfx-rs/wgpu/pull/6263).
 
 #### HAL
 
diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs
index 7ab97f491..7fb30b4f7 100644
--- a/naga/src/back/msl/writer.rs
+++ b/naga/src/back/msl/writer.rs
@@ -33,6 +33,7 @@ const RAY_QUERY_FIELD_INTERSECTION: &str = "intersection";
 const RAY_QUERY_FIELD_READY: &str = "ready";
 const RAY_QUERY_FUN_MAP_INTERSECTION: &str = "_map_intersection_type";
 
+pub(crate) const ATOMIC_COMP_EXCH_FUNCTION: &str = "naga_atomic_compare_exchange_weak_explicit";
 pub(crate) const MODF_FUNCTION: &str = "naga_modf";
 pub(crate) const FREXP_FUNCTION: &str = "naga_frexp";
 
@@ -1279,42 +1280,6 @@ impl<W: Write> Writer<W> {
         Ok(())
     }
 
-    fn put_atomic_operation(
-        &mut self,
-        pointer: Handle<crate::Expression>,
-        key: &str,
-        value: Handle<crate::Expression>,
-        context: &ExpressionContext,
-    ) -> BackendResult {
-        // If the pointer we're passing to the atomic operation needs to be conditional
-        // for `ReadZeroSkipWrite`, the condition needs to *surround* the atomic op, and
-        // the pointer operand should be unchecked.
-        let policy = context.choose_bounds_check_policy(pointer);
-        let checked = policy == index::BoundsCheckPolicy::ReadZeroSkipWrite
-            && self.put_bounds_checks(pointer, context, back::Level(0), "")?;
-
-        // If requested and successfully put bounds checks, continue the ternary expression.
-        if checked {
-            write!(self.out, " ? ")?;
-        }
-
-        write!(
-            self.out,
-            "{NAMESPACE}::atomic_{key}_explicit({ATOMIC_REFERENCE}"
-        )?;
-        self.put_access_chain(pointer, policy, context)?;
-        write!(self.out, ", ")?;
-        self.put_expression(value, context, true)?;
-        write!(self.out, ", {NAMESPACE}::memory_order_relaxed)")?;
-
-        // Finish the ternary expression.
-        if checked {
-            write!(self.out, " : DefaultConstructible()")?;
-        }
-
-        Ok(())
-    }
-
     /// Emit code for the arithmetic expression of the dot product.
     ///
     fn put_dot_product(
@@ -3182,24 +3147,65 @@ impl<W: Write> Writer<W> {
                     value,
                     result,
                 } => {
+                    let context = &context.expression;
+
                     // This backend supports `SHADER_INT64_ATOMIC_MIN_MAX` but not
                     // `SHADER_INT64_ATOMIC_ALL_OPS`, so we can assume that if `result` is
                     // `Some`, we are not operating on a 64-bit value, and that if we are
                     // operating on a 64-bit value, `result` is `None`.
                     write!(self.out, "{level}")?;
-                    let fun_str = if let Some(result) = result {
+                    let fun_key = if let Some(result) = result {
                         let res_name = Baked(result).to_string();
-                        self.start_baking_expression(result, &context.expression, &res_name)?;
+                        self.start_baking_expression(result, context, &res_name)?;
                         self.named_expressions.insert(result, res_name);
-                        fun.to_msl()?
-                    } else if context.expression.resolve_type(value).scalar_width() == Some(8) {
+                        fun.to_msl()
+                    } else if context.resolve_type(value).scalar_width() == Some(8) {
                         fun.to_msl_64_bit()?
                     } else {
-                        fun.to_msl()?
+                        fun.to_msl()
                     };
 
-                    self.put_atomic_operation(pointer, fun_str, value, &context.expression)?;
-                    // done
+                    // If the pointer we're passing to the atomic operation needs to be conditional
+                    // for `ReadZeroSkipWrite`, the condition needs to *surround* the atomic op, and
+                    // the pointer operand should be unchecked.
+                    let policy = context.choose_bounds_check_policy(pointer);
+                    let checked = policy == index::BoundsCheckPolicy::ReadZeroSkipWrite
+                        && self.put_bounds_checks(pointer, context, back::Level(0), "")?;
+
+                    // If requested and successfully put bounds checks, continue the ternary expression.
+                    if checked {
+                        write!(self.out, " ? ")?;
+                    }
+
+                    // Put the atomic function invocation.
+                    match *fun {
+                        crate::AtomicFunction::Exchange { compare: Some(cmp) } => {
+                            write!(self.out, "{ATOMIC_COMP_EXCH_FUNCTION}({ATOMIC_REFERENCE}")?;
+                            self.put_access_chain(pointer, policy, context)?;
+                            write!(self.out, ", ")?;
+                            self.put_expression(cmp, context, true)?;
+                            write!(self.out, ", ")?;
+                            self.put_expression(value, context, true)?;
+                            write!(self.out, ")")?;
+                        }
+                        _ => {
+                            write!(
+                                self.out,
+                                "{NAMESPACE}::atomic_{fun_key}_explicit({ATOMIC_REFERENCE}"
+                            )?;
+                            self.put_access_chain(pointer, policy, context)?;
+                            write!(self.out, ", ")?;
+                            self.put_expression(value, context, true)?;
+                            write!(self.out, ", {NAMESPACE}::memory_order_relaxed)")?;
+                        }
+                    }
+
+                    // Finish the ternary expression.
+                    if checked {
+                        write!(self.out, " : DefaultConstructible()")?;
+                    }
+
+                    // Done
                     writeln!(self.out, ";")?;
                 }
                 crate::Statement::WorkGroupUniformLoad { pointer, result } => {
@@ -3827,7 +3833,33 @@ impl<W: Write> Writer<W> {
 }}"
                     )?;
                 }
-                &crate::PredeclaredType::AtomicCompareExchangeWeakResult { .. } => {}
+                &crate::PredeclaredType::AtomicCompareExchangeWeakResult(scalar) => {
+                    let arg_type_name = scalar.to_msl_name();
+                    let called_func_name = "atomic_compare_exchange_weak_explicit";
+                    let defined_func_name = ATOMIC_COMP_EXCH_FUNCTION;
+                    let struct_name = &self.names[&NameKey::Type(*struct_ty)];
+
+                    writeln!(self.out)?;
+
+                    for address_space_name in ["device", "threadgroup"] {
+                        writeln!(
+                            self.out,
+                            "\
+template <typename A>
+{struct_name} {defined_func_name}(
+    {address_space_name} A *atomic_ptr,
+    {arg_type_name} cmp,
+    {arg_type_name} v
+) {{
+    bool swapped = {NAMESPACE}::{called_func_name}(
+        atomic_ptr, &cmp, v,
+        metal::memory_order_relaxed, metal::memory_order_relaxed
+    );
+    return {struct_name}{{cmp, swapped}};
+}}"
+                        )?;
+                    }
+                }
             }
         }
 
@@ -6065,8 +6097,8 @@ fn test_stack_size() {
 }
 
 impl crate::AtomicFunction {
-    fn to_msl(self) -> Result<&'static str, Error> {
-        Ok(match self {
+    const fn to_msl(self) -> &'static str {
+        match self {
             Self::Add => "fetch_add",
             Self::Subtract => "fetch_sub",
             Self::And => "fetch_and",
@@ -6075,10 +6107,8 @@ impl crate::AtomicFunction {
             Self::Min => "fetch_min",
             Self::Max => "fetch_max",
             Self::Exchange { compare: None } => "exchange",
-            Self::Exchange { compare: Some(_) } => Err(Error::FeatureNotImplemented(
-                "atomic CompareExchange".to_string(),
-            ))?,
-        })
+            Self::Exchange { compare: Some(_) } => ATOMIC_COMP_EXCH_FUNCTION,
+        }
     }
 
     fn to_msl_64_bit(self) -> Result<&'static str, Error> {
diff --git a/naga/src/back/spv/block.rs b/naga/src/back/spv/block.rs
index 496ae4fad..d7c56d3f4 100644
--- a/naga/src/back/spv/block.rs
+++ b/naga/src/back/spv/block.rs
@@ -3,8 +3,8 @@ Implementations for `BlockContext` methods.
 */
 
 use super::{
-    helpers, index::BoundsCheckResult, make_local, selection::Selection, Block, BlockContext,
-    Dimension, Error, Instruction, LocalType, LookupType, ResultMember, Writer, WriterFlags,
+    helpers, index::BoundsCheckResult, selection::Selection, Block, BlockContext, Dimension, Error,
+    Instruction, LocalType, LookupType, NumericType, ResultMember, Writer, WriterFlags,
 };
 use crate::{arena::Handle, proc::TypeResolution, Statement};
 use spirv::Word;
@@ -105,10 +105,9 @@ impl Writer {
         position_id: Word,
         body: &mut Vec<Instruction>,
     ) -> Result<(), Error> {
-        let float_ptr_type_id = self.get_type_id(LookupType::Local(LocalType::Value {
-            vector_size: None,
-            scalar: crate::Scalar::F32,
-            pointer_space: Some(spirv::StorageClass::Output),
+        let float_ptr_type_id = self.get_type_id(LookupType::Local(LocalType::LocalPointer {
+            base: NumericType::Scalar(crate::Scalar::F32),
+            class: spirv::StorageClass::Output,
         }));
         let index_y_id = self.get_index_constant(1);
         let access_id = self.id_gen.next();
@@ -119,11 +118,9 @@ impl Writer {
             &[index_y_id],
         ));
 
-        let float_type_id = self.get_type_id(LookupType::Local(LocalType::Value {
-            vector_size: None,
-            scalar: crate::Scalar::F32,
-            pointer_space: None,
-        }));
+        let float_type_id = self.get_type_id(LookupType::Local(LocalType::Numeric(
+            NumericType::Scalar(crate::Scalar::F32),
+        )));
         let load_id = self.id_gen.next();
         body.push(Instruction::load(float_type_id, load_id, access_id, None));
 
@@ -145,11 +142,9 @@ impl Writer {
         frag_depth_id: Word,
         body: &mut Vec<Instruction>,
     ) -> Result<(), Error> {
-        let float_type_id = self.get_type_id(LookupType::Local(LocalType::Value {
-            vector_size: None,
-            scalar: crate::Scalar::F32,
-            pointer_space: None,
-        }));
+        let float_type_id = self.get_type_id(LookupType::Local(LocalType::Numeric(
+            NumericType::Scalar(crate::Scalar::F32),
+        )));
         let zero_scalar_id = self.get_constant_scalar(crate::Literal::F32(0.0));
         let one_scalar_id = self.get_constant_scalar(crate::Literal::F32(1.0));
 
@@ -830,12 +825,8 @@ impl<'w> BlockContext<'w> {
                         let mut arg2_id = self.writer.get_constant_scalar_with(1, scalar)?;
 
                         if let Some(size) = maybe_size {
-                            let ty = LocalType::Value {
-                                vector_size: Some(size),
-                                scalar,
-                                pointer_space: None,
-                            }
-                            .into();
+                            let ty =
+                                LocalType::Numeric(NumericType::Vector { size, scalar }).into();
 
                             self.temp_list.clear();
                             self.temp_list.resize(size as _, arg1_id);
@@ -950,12 +941,9 @@ impl<'w> BlockContext<'w> {
                                 &crate::TypeInner::Vector { size, .. },
                                 &crate::TypeInner::Scalar(scalar),
                             ) => {
-                                let selector_type_id =
-                                    self.get_type_id(LookupType::Local(LocalType::Value {
-                                        vector_size: Some(size),
-                                        scalar,
-                                        pointer_space: None,
-                                    }));
+                                let selector_type_id = self.get_type_id(LookupType::Local(
+                                    LocalType::Numeric(NumericType::Vector { size, scalar }),
+                                ));
                                 self.temp_list.clear();
                                 self.temp_list.resize(size as usize, arg2_id);
 
@@ -998,12 +986,8 @@ impl<'w> BlockContext<'w> {
                     Mf::CountTrailingZeros => {
                         let uint_id = match *arg_ty {
                             crate::TypeInner::Vector { size, scalar } => {
-                                let ty = LocalType::Value {
-                                    vector_size: Some(size),
-                                    scalar,
-                                    pointer_space: None,
-                                }
-                                .into();
+                                let ty =
+                                    LocalType::Numeric(NumericType::Vector { size, scalar }).into();
 
                                 self.temp_list.clear();
                                 self.temp_list.resize(
@@ -1040,12 +1024,8 @@ impl<'w> BlockContext<'w> {
                     Mf::CountLeadingZeros => {
                         let (int_type_id, int_id, width) = match *arg_ty {
                             crate::TypeInner::Vector { size, scalar } => {
-                                let ty = LocalType::Value {
-                                    vector_size: Some(size),
-                                    scalar,
-                                    pointer_space: None,
-                                }
-                                .into();
+                                let ty =
+                                    LocalType::Numeric(NumericType::Vector { size, scalar }).into();
 
                                 self.temp_list.clear();
                                 self.temp_list.resize(
@@ -1061,11 +1041,9 @@ impl<'w> BlockContext<'w> {
                                 )
                             }
                             crate::TypeInner::Scalar(scalar) => (
-                                self.get_type_id(LookupType::Local(LocalType::Value {
-                                    vector_size: None,
-                                    scalar,
-                                    pointer_space: None,
-                                })),
+                                self.get_type_id(LookupType::Local(LocalType::Numeric(
+                                    NumericType::Scalar(scalar),
+                                ))),
                                 self.writer
                                     .get_constant_scalar_with(scalar.width * 8 - 1, scalar)?,
                                 scalar.width,
@@ -1130,14 +1108,9 @@ impl<'w> BlockContext<'w> {
                             .writer
                             .get_constant_scalar(crate::Literal::U32(bit_width as u32));
 
-                        let u32_type = self.get_type_id(LookupType::Local(LocalType::Value {
-                            vector_size: None,
-                            scalar: crate::Scalar {
-                                kind: crate::ScalarKind::Uint,
-                                width: 4,
-                            },
-                            pointer_space: None,
-                        }));
+                        let u32_type = self.get_type_id(LookupType::Local(LocalType::Numeric(
+                            NumericType::Scalar(crate::Scalar::U32),
+                        )));
 
                         // o = min(offset, w)
                         let offset_id = self.gen_id();
@@ -1186,14 +1159,9 @@ impl<'w> BlockContext<'w> {
                             .writer
                             .get_constant_scalar(crate::Literal::U32(bit_width as u32));
 
-                        let u32_type = self.get_type_id(LookupType::Local(LocalType::Value {
-                            vector_size: None,
-                            scalar: crate::Scalar {
-                                kind: crate::ScalarKind::Uint,
-                                width: 4,
-                            },
-                            pointer_space: None,
-                        }));
+                        let u32_type = self.get_type_id(LookupType::Local(LocalType::Numeric(
+                            NumericType::Scalar(crate::Scalar::U32),
+                        )));
 
                         // o = min(offset, w)
                         let offset_id = self.gen_id();
@@ -1259,23 +1227,16 @@ impl<'w> BlockContext<'w> {
                             Mf::Pack4xU8 => (crate::ScalarKind::Uint, false),
                             _ => unreachable!(),
                         };
-                        let uint_type_id = self.get_type_id(LookupType::Local(LocalType::Value {
-                            vector_size: None,
-                            scalar: crate::Scalar {
-                                kind: crate::ScalarKind::Uint,
-                                width: 4,
-                            },
-                            pointer_space: None,
-                        }));
+                        let uint_type_id = self.get_type_id(LookupType::Local(LocalType::Numeric(
+                            NumericType::Scalar(crate::Scalar::U32),
+                        )));
 
-                        let int_type_id = self.get_type_id(LookupType::Local(LocalType::Value {
-                            vector_size: None,
-                            scalar: crate::Scalar {
+                        let int_type_id = self.get_type_id(LookupType::Local(LocalType::Numeric(
+                            NumericType::Scalar(crate::Scalar {
                                 kind: int_type,
                                 width: 4,
-                            },
-                            pointer_space: None,
-                        }));
+                            }),
+                        )));
 
                         let mut last_instruction = Instruction::new(spirv::Op::Nop);
 
@@ -1352,24 +1313,17 @@ impl<'w> BlockContext<'w> {
                             _ => unreachable!(),
                         };
 
-                        let sint_type_id = self.get_type_id(LookupType::Local(LocalType::Value {
-                            vector_size: None,
-                            scalar: crate::Scalar {
-                                kind: crate::ScalarKind::Sint,
-                                width: 4,
-                            },
-                            pointer_space: None,
-                        }));
+                        let sint_type_id = self.get_type_id(LookupType::Local(LocalType::Numeric(
+                            NumericType::Scalar(crate::Scalar::I32),
+                        )));
 
                         let eight = self.writer.get_constant_scalar(crate::Literal::U32(8));
-                        let int_type_id = self.get_type_id(LookupType::Local(LocalType::Value {
-                            vector_size: None,
-                            scalar: crate::Scalar {
+                        let int_type_id = self.get_type_id(LookupType::Local(LocalType::Numeric(
+                            NumericType::Scalar(crate::Scalar {
                                 kind: int_type,
                                 width: 4,
-                            },
-                            pointer_space: None,
-                        }));
+                            }),
+                        )));
                         block
                             .body
                             .reserve(usize::from(VEC_LENGTH) * 2 + usize::from(is_signed));
@@ -1533,11 +1487,10 @@ impl<'w> BlockContext<'w> {
                                 self.writer.get_constant_scalar_with(0, src_scalar)?;
                             let zero_id = match src_size {
                                 Some(size) => {
-                                    let ty = LocalType::Value {
-                                        vector_size: Some(size),
+                                    let ty = LocalType::Numeric(NumericType::Vector {
+                                        size,
                                         scalar: src_scalar,
-                                        pointer_space: None,
-                                    }
+                                    })
                                     .into();
 
                                     self.temp_list.clear();
@@ -1562,11 +1515,10 @@ impl<'w> BlockContext<'w> {
                                 self.writer.get_constant_scalar_with(1, dst_scalar)?;
                             let (accept_id, reject_id) = match src_size {
                                 Some(size) => {
-                                    let ty = LocalType::Value {
-                                        vector_size: Some(size),
+                                    let ty = LocalType::Numeric(NumericType::Vector {
+                                        size,
                                         scalar: dst_scalar,
-                                        pointer_space: None,
-                                    }
+                                    })
                                     .into();
 
                                     self.temp_list.clear();
@@ -1704,12 +1656,12 @@ impl<'w> BlockContext<'w> {
                     self.temp_list.clear();
                     self.temp_list.resize(size as usize, condition_id);
 
-                    let bool_vector_type_id =
-                        self.get_type_id(LookupType::Local(LocalType::Value {
-                            vector_size: Some(size),
+                    let bool_vector_type_id = self.get_type_id(LookupType::Local(
+                        LocalType::Numeric(NumericType::Vector {
+                            size,
                             scalar: condition_scalar,
-                            pointer_space: None,
-                        }));
+                        }),
+                    ));
 
                     let id = self.gen_id();
                     block.body.push(Instruction::composite_construct(
@@ -1809,7 +1761,9 @@ impl<'w> BlockContext<'w> {
                 Some(ty) => ty,
                 None => LookupType::Handle(ty_handle),
             },
-            TypeResolution::Value(ref inner) => LookupType::Local(make_local(inner).unwrap()),
+            TypeResolution::Value(ref inner) => {
+                LookupType::Local(LocalType::from_inner(inner).unwrap())
+            }
         };
         let result_type_id = self.get_type_id(result_lookup_ty);
 
@@ -2029,11 +1983,11 @@ impl<'w> BlockContext<'w> {
     ) {
         self.temp_list.clear();
 
-        let vector_type_id = self.get_type_id(LookupType::Local(LocalType::Value {
-            vector_size: Some(rows),
-            scalar: crate::Scalar::float(width),
-            pointer_space: None,
-        }));
+        let vector_type_id =
+            self.get_type_id(LookupType::Local(LocalType::Numeric(NumericType::Vector {
+                size: rows,
+                scalar: crate::Scalar::float(width),
+            })));
 
         for index in 0..columns as u32 {
             let column_id_left = self.gen_id();
@@ -2735,20 +2689,15 @@ impl<'w> BlockContext<'w> {
                         crate::AtomicFunction::Exchange { compare: Some(cmp) } => {
                             let scalar_type_id = match *value_inner {
                                 crate::TypeInner::Scalar(scalar) => {
-                                    self.get_type_id(LookupType::Local(LocalType::Value {
-                                        vector_size: None,
-                                        scalar,
-                                        pointer_space: None,
-                                    }))
+                                    self.get_type_id(LookupType::Local(LocalType::Numeric(
+                                        NumericType::Scalar(scalar),
+                                    )))
                                 }
                                 _ => unimplemented!(),
                             };
-                            let bool_type_id =
-                                self.get_type_id(LookupType::Local(LocalType::Value {
-                                    vector_size: None,
-                                    scalar: crate::Scalar::BOOL,
-                                    pointer_space: None,
-                                }));
+                            let bool_type_id = self.get_type_id(LookupType::Local(
+                                LocalType::Numeric(NumericType::Scalar(crate::Scalar::BOOL)),
+                            ));
 
                             let cas_result_id = self.gen_id();
                             let equality_result_id = self.gen_id();
diff --git a/naga/src/back/spv/image.rs b/naga/src/back/spv/image.rs
index 769971d13..a76d015f3 100644
--- a/naga/src/back/spv/image.rs
+++ b/naga/src/back/spv/image.rs
@@ -4,7 +4,7 @@ Generating SPIR-V for image operations.
 
 use super::{
     selection::{MergeTuple, Selection},
-    Block, BlockContext, Error, IdGenerator, Instruction, LocalType, LookupType,
+    Block, BlockContext, Error, IdGenerator, Instruction, LocalType, LookupType, NumericType,
 };
 use crate::arena::Handle;
 use spirv::Word;
@@ -126,11 +126,10 @@ impl Load {
         // the right SPIR-V type for the access instruction here.
         let type_id = match image_class {
             crate::ImageClass::Depth { .. } => {
-                ctx.get_type_id(LookupType::Local(LocalType::Value {
-                    vector_size: Some(crate::VectorSize::Quad),
+                ctx.get_type_id(LookupType::Local(LocalType::Numeric(NumericType::Vector {
+                    size: crate::VectorSize::Quad,
                     scalar: crate::Scalar::F32,
-                    pointer_space: None,
-                }))
+                })))
             }
             _ => result_type_id,
         };
@@ -292,15 +291,15 @@ impl<'w> BlockContext<'w> {
         // Find the component type of `coordinates`, and figure out the size the
         // combined coordinate vector will have.
         let (component_scalar, size) = match *inner_ty {
-            Ti::Scalar(scalar @ crate::Scalar { width: 4, .. }) => (scalar, Some(Vs::Bi)),
+            Ti::Scalar(scalar @ crate::Scalar { width: 4, .. }) => (scalar, Vs::Bi),
             Ti::Vector {
                 scalar: scalar @ crate::Scalar { width: 4, .. },
                 size: Vs::Bi,
-            } => (scalar, Some(Vs::Tri)),
+            } => (scalar, Vs::Tri),
             Ti::Vector {
                 scalar: scalar @ crate::Scalar { width: 4, .. },
                 size: Vs::Tri,
-            } => (scalar, Some(Vs::Quad)),
+            } => (scalar, Vs::Quad),
             Ti::Vector { size: Vs::Quad, .. } => {
                 return Err(Error::Validation("extending vec4 coordinate"));
             }
@@ -340,11 +339,9 @@ impl<'w> BlockContext<'w> {
             }
         };
         let reconciled_array_index_id = if let Some(cast) = cast {
-            let component_ty_id = self.get_type_id(LookupType::Local(LocalType::Value {
-                vector_size: None,
-                scalar: component_scalar,
-                pointer_space: None,
-            }));
+            let component_ty_id = self.get_type_id(LookupType::Local(LocalType::Numeric(
+                NumericType::Scalar(component_scalar),
+            )));
             let reconciled_id = self.gen_id();
             block.body.push(Instruction::unary(
                 cast,
@@ -358,11 +355,11 @@ impl<'w> BlockContext<'w> {
         };
 
         // Find the SPIR-V type for the combined coordinates/index vector.
-        let type_id = self.get_type_id(LookupType::Local(LocalType::Value {
-            vector_size: size,
-            scalar: component_scalar,
-            pointer_space: None,
-        }));
+        let type_id =
+            self.get_type_id(LookupType::Local(LocalType::Numeric(NumericType::Vector {
+                size,
+                scalar: component_scalar,
+            })));
 
         // Schmear the coordinates and index together.
         let value_id = self.gen_id();
@@ -374,7 +371,7 @@ impl<'w> BlockContext<'w> {
         Ok(ImageCoordinates {
             value_id,
             type_id,
-            size,
+            size: Some(size),
         })
     }
 
@@ -529,11 +526,9 @@ impl<'w> BlockContext<'w> {
             &[spirv::Capability::ImageQuery],
         )?;
 
-        let i32_type_id = self.get_type_id(LookupType::Local(LocalType::Value {
-            vector_size: None,
-            scalar: crate::Scalar::I32,
-            pointer_space: None,
-        }));
+        let i32_type_id = self.get_type_id(LookupType::Local(LocalType::Numeric(
+            NumericType::Scalar(crate::Scalar::I32),
+        )));
 
         // If `level` is `Some`, clamp it to fall within bounds. This must
         // happen first, because we'll use it to query the image size for
@@ -616,11 +611,9 @@ impl<'w> BlockContext<'w> {
         )?;
 
         let bool_type_id = self.writer.get_bool_type_id();
-        let i32_type_id = self.get_type_id(LookupType::Local(LocalType::Value {
-            vector_size: None,
-            scalar: crate::Scalar::I32,
-            pointer_space: None,
-        }));
+        let i32_type_id = self.get_type_id(LookupType::Local(LocalType::Numeric(
+            NumericType::Scalar(crate::Scalar::I32),
+        )));
 
         let null_id = access.out_of_bounds_value(self);
 
@@ -683,11 +676,15 @@ impl<'w> BlockContext<'w> {
         );
 
         // Compare the coordinates against the bounds.
-        let coords_bool_type_id = self.get_type_id(LookupType::Local(LocalType::Value {
-            vector_size: coordinates.size,
-            scalar: crate::Scalar::BOOL,
-            pointer_space: None,
-        }));
+        let coords_numeric_type = match coordinates.size {
+            Some(size) => NumericType::Vector {
+                size,
+                scalar: crate::Scalar::BOOL,
+            },
+            None => NumericType::Scalar(crate::Scalar::BOOL),
+        };
+        let coords_bool_type_id =
+            self.get_type_id(LookupType::Local(LocalType::Numeric(coords_numeric_type)));
         let coords_conds_id = self.gen_id();
         selection.block().body.push(Instruction::binary(
             spirv::Op::ULessThan,
@@ -838,11 +835,10 @@ impl<'w> BlockContext<'w> {
             _ => false,
         };
         let sample_result_type_id = if needs_sub_access {
-            self.get_type_id(LookupType::Local(LocalType::Value {
-                vector_size: Some(crate::VectorSize::Quad),
+            self.get_type_id(LookupType::Local(LocalType::Numeric(NumericType::Vector {
+                size: crate::VectorSize::Quad,
                 scalar: crate::Scalar::F32,
-                pointer_space: None,
-            }))
+            })))
         } else {
             result_type_id
         };
@@ -1038,11 +1034,16 @@ impl<'w> BlockContext<'w> {
                     4 => Some(crate::VectorSize::Quad),
                     _ => None,
                 };
-                let extended_size_type_id = self.get_type_id(LookupType::Local(LocalType::Value {
-                    vector_size,
-                    scalar: crate::Scalar::U32,
-                    pointer_space: None,
-                }));
+                let vector_numeric_type = match vector_size {
+                    Some(size) => NumericType::Vector {
+                        size,
+                        scalar: crate::Scalar::U32,
+                    },
+                    None => NumericType::Scalar(crate::Scalar::U32),
+                };
+
+                let extended_size_type_id =
+                    self.get_type_id(LookupType::Local(LocalType::Numeric(vector_numeric_type)));
 
                 let (query_op, level_id) = match class {
                     Ic::Sampled { multi: true, .. }
@@ -1108,11 +1109,11 @@ impl<'w> BlockContext<'w> {
                     Id::D2 | Id::Cube => crate::VectorSize::Tri,
                     Id::D3 => crate::VectorSize::Quad,
                 };
-                let extended_size_type_id = self.get_type_id(LookupType::Local(LocalType::Value {
-                    vector_size: Some(vec_size),
-                    scalar: crate::Scalar::U32,
-                    pointer_space: None,
-                }));
+                let extended_size_type_id =
+                    self.get_type_id(LookupType::Local(LocalType::Numeric(NumericType::Vector {
+                        size: vec_size,
+                        scalar: crate::Scalar::U32,
+                    })));
                 let id_extended = self.gen_id();
                 let mut inst = Instruction::image_query(
                     spirv::Op::ImageQuerySizeLod,
diff --git a/naga/src/back/spv/mod.rs b/naga/src/back/spv/mod.rs
index e6397017c..aa4de6846 100644
--- a/naga/src/back/spv/mod.rs
+++ b/naga/src/back/spv/mod.rs
@@ -231,6 +231,21 @@ impl LocalImageType {
     }
 }
 
+/// A numeric type, for use in [`LocalType`].
+#[derive(Debug, PartialEq, Hash, Eq, Copy, Clone)]
+enum NumericType {
+    Scalar(crate::Scalar),
+    Vector {
+        size: crate::VectorSize,
+        scalar: crate::Scalar,
+    },
+    Matrix {
+        columns: crate::VectorSize,
+        rows: crate::VectorSize,
+        scalar: crate::Scalar,
+    },
+}
+
 /// A SPIR-V type constructed during code generation.
 ///
 /// This is the variant of [`LookupType`] used to represent types that might not
@@ -246,9 +261,9 @@ impl LocalImageType {
 /// never synthesizes new struct types, so `LocalType` has nothing for that.
 ///
 /// Each `LocalType` variant should be handled identically to its analogous
-/// `TypeInner` variant. You can use the [`make_local`] function to help with
-/// this, by converting everything possible to a `LocalType` before inspecting
-/// it.
+/// `TypeInner` variant. You can use the [`LocalType::from_inner`] function to
+/// help with this, by converting everything possible to a `LocalType` before
+/// inspecting it.
 ///
 /// ## `LocalType` equality and SPIR-V `OpType` uniqueness
 ///
@@ -276,19 +291,11 @@ impl LocalImageType {
 /// [`TypeInner`]: crate::TypeInner
 #[derive(Debug, PartialEq, Hash, Eq, Copy, Clone)]
 enum LocalType {
-    /// A scalar, vector, or pointer to one of those.
-    Value {
-        /// If `None`, this represents a scalar type. If `Some`, this represents
-        /// a vector type of the given size.
-        vector_size: Option<crate::VectorSize>,
-        scalar: crate::Scalar,
-        pointer_space: Option<spirv::StorageClass>,
-    },
-    /// A matrix of floating-point values.
-    Matrix {
-        columns: crate::VectorSize,
-        rows: crate::VectorSize,
-        width: crate::Bytes,
+    /// A numeric type.
+    Numeric(NumericType),
+    LocalPointer {
+        base: NumericType,
+        class: spirv::StorageClass,
     },
     Pointer {
         base: Handle<crate::Type>,
@@ -357,52 +364,57 @@ struct LookupFunctionType {
     return_type_id: Word,
 }
 
-fn make_local(inner: &crate::TypeInner) -> Option<LocalType> {
-    Some(match *inner {
-        crate::TypeInner::Scalar(scalar) | crate::TypeInner::Atomic(scalar) => LocalType::Value {
-            vector_size: None,
-            scalar,
-            pointer_space: None,
-        },
-        crate::TypeInner::Vector { size, scalar } => LocalType::Value {
-            vector_size: Some(size),
-            scalar,
-            pointer_space: None,
-        },
-        crate::TypeInner::Matrix {
-            columns,
-            rows,
-            scalar,
-        } => LocalType::Matrix {
-            columns,
-            rows,
-            width: scalar.width,
-        },
-        crate::TypeInner::Pointer { base, space } => LocalType::Pointer {
-            base,
-            class: helpers::map_storage_class(space),
-        },
-        crate::TypeInner::ValuePointer {
-            size,
-            scalar,
-            space,
-        } => LocalType::Value {
-            vector_size: size,
-            scalar,
-            pointer_space: Some(helpers::map_storage_class(space)),
-        },
-        crate::TypeInner::Image {
-            dim,
-            arrayed,
-            class,
-        } => LocalType::Image(LocalImageType::from_inner(dim, arrayed, class)),
-        crate::TypeInner::Sampler { comparison: _ } => LocalType::Sampler,
-        crate::TypeInner::AccelerationStructure => LocalType::AccelerationStructure,
-        crate::TypeInner::RayQuery => LocalType::RayQuery,
-        crate::TypeInner::Array { .. }
-        | crate::TypeInner::Struct { .. }
-        | crate::TypeInner::BindingArray { .. } => return None,
-    })
+impl LocalType {
+    fn from_inner(inner: &crate::TypeInner) -> Option<Self> {
+        Some(match *inner {
+            crate::TypeInner::Scalar(scalar) | crate::TypeInner::Atomic(scalar) => {
+                LocalType::Numeric(NumericType::Scalar(scalar))
+            }
+            crate::TypeInner::Vector { size, scalar } => {
+                LocalType::Numeric(NumericType::Vector { size, scalar })
+            }
+            crate::TypeInner::Matrix {
+                columns,
+                rows,
+                scalar,
+            } => LocalType::Numeric(NumericType::Matrix {
+                columns,
+                rows,
+                scalar,
+            }),
+            crate::TypeInner::Pointer { base, space } => LocalType::Pointer {
+                base,
+                class: helpers::map_storage_class(space),
+            },
+            crate::TypeInner::ValuePointer {
+                size: Some(size),
+                scalar,
+                space,
+            } => LocalType::LocalPointer {
+                base: NumericType::Vector { size, scalar },
+                class: helpers::map_storage_class(space),
+            },
+            crate::TypeInner::ValuePointer {
+                size: None,
+                scalar,
+                space,
+            } => LocalType::LocalPointer {
+                base: NumericType::Scalar(scalar),
+                class: helpers::map_storage_class(space),
+            },
+            crate::TypeInner::Image {
+                dim,
+                arrayed,
+                class,
+            } => LocalType::Image(LocalImageType::from_inner(dim, arrayed, class)),
+            crate::TypeInner::Sampler { comparison: _ } => LocalType::Sampler,
+            crate::TypeInner::AccelerationStructure => LocalType::AccelerationStructure,
+            crate::TypeInner::RayQuery => LocalType::RayQuery,
+            crate::TypeInner::Array { .. }
+            | crate::TypeInner::Struct { .. }
+            | crate::TypeInner::BindingArray { .. } => return None,
+        })
+    }
 }
 
 #[derive(Debug)]
diff --git a/naga/src/back/spv/ray.rs b/naga/src/back/spv/ray.rs
index bc2c4ce3c..c2daf4b3f 100644
--- a/naga/src/back/spv/ray.rs
+++ b/naga/src/back/spv/ray.rs
@@ -2,7 +2,7 @@
 Generating SPIR-V for ray query operations.
 */
 
-use super::{Block, BlockContext, Instruction, LocalType, LookupType};
+use super::{Block, BlockContext, Instruction, LocalType, LookupType, NumericType};
 use crate::arena::Handle;
 
 impl<'w> BlockContext<'w> {
@@ -22,11 +22,9 @@ impl<'w> BlockContext<'w> {
                 let desc_id = self.cached[descriptor];
                 let acc_struct_id = self.get_handle_id(acceleration_structure);
 
-                let flag_type_id = self.get_type_id(LookupType::Local(LocalType::Value {
-                    vector_size: None,
-                    scalar: crate::Scalar::U32,
-                    pointer_space: None,
-                }));
+                let flag_type_id = self.get_type_id(LookupType::Local(LocalType::Numeric(
+                    NumericType::Scalar(crate::Scalar::U32),
+                )));
                 let ray_flags_id = self.gen_id();
                 block.body.push(Instruction::composite_extract(
                     flag_type_id,
@@ -42,11 +40,9 @@ impl<'w> BlockContext<'w> {
                     &[1],
                 ));
 
-                let scalar_type_id = self.get_type_id(LookupType::Local(LocalType::Value {
-                    vector_size: None,
-                    scalar: crate::Scalar::F32,
-                    pointer_space: None,
-                }));
+                let scalar_type_id = self.get_type_id(LookupType::Local(LocalType::Numeric(
+                    NumericType::Scalar(crate::Scalar::F32),
+                )));
                 let tmin_id = self.gen_id();
                 block.body.push(Instruction::composite_extract(
                     scalar_type_id,
@@ -62,11 +58,11 @@ impl<'w> BlockContext<'w> {
                     &[3],
                 ));
 
-                let vector_type_id = self.get_type_id(LookupType::Local(LocalType::Value {
-                    vector_size: Some(crate::VectorSize::Tri),
-                    scalar: crate::Scalar::F32,
-                    pointer_space: None,
-                }));
+                let vector_type_id =
+                    self.get_type_id(LookupType::Local(LocalType::Numeric(NumericType::Vector {
+                        size: crate::VectorSize::Tri,
+                        scalar: crate::Scalar::F32,
+                    })));
                 let ray_origin_id = self.gen_id();
                 block.body.push(Instruction::composite_extract(
                     vector_type_id,
@@ -116,11 +112,9 @@ impl<'w> BlockContext<'w> {
             spirv::RayQueryIntersection::RayQueryCommittedIntersectionKHR as _,
         ));
 
-        let flag_type_id = self.get_type_id(LookupType::Local(LocalType::Value {
-            vector_size: None,
-            scalar: crate::Scalar::U32,
-            pointer_space: None,
-        }));
+        let flag_type_id = self.get_type_id(LookupType::Local(LocalType::Numeric(
+            NumericType::Scalar(crate::Scalar::U32),
+        )));
         let kind_id = self.gen_id();
         block.body.push(Instruction::ray_query_get_intersection(
             spirv::Op::RayQueryGetIntersectionTypeKHR,
@@ -170,11 +164,9 @@ impl<'w> BlockContext<'w> {
             intersection_id,
         ));
 
-        let scalar_type_id = self.get_type_id(LookupType::Local(LocalType::Value {
-            vector_size: None,
-            scalar: crate::Scalar::F32,
-            pointer_space: None,
-        }));
+        let scalar_type_id = self.get_type_id(LookupType::Local(LocalType::Numeric(
+            NumericType::Scalar(crate::Scalar::F32),
+        )));
         let t_id = self.gen_id();
         block.body.push(Instruction::ray_query_get_intersection(
             spirv::Op::RayQueryGetIntersectionTKHR,
@@ -184,11 +176,11 @@ impl<'w> BlockContext<'w> {
             intersection_id,
         ));
 
-        let barycentrics_type_id = self.get_type_id(LookupType::Local(LocalType::Value {
-            vector_size: Some(crate::VectorSize::Bi),
-            scalar: crate::Scalar::F32,
-            pointer_space: None,
-        }));
+        let barycentrics_type_id =
+            self.get_type_id(LookupType::Local(LocalType::Numeric(NumericType::Vector {
+                size: crate::VectorSize::Bi,
+                scalar: crate::Scalar::F32,
+            })));
         let barycentrics_id = self.gen_id();
         block.body.push(Instruction::ray_query_get_intersection(
             spirv::Op::RayQueryGetIntersectionBarycentricsKHR,
@@ -198,11 +190,9 @@ impl<'w> BlockContext<'w> {
             intersection_id,
         ));
 
-        let bool_type_id = self.get_type_id(LookupType::Local(LocalType::Value {
-            vector_size: None,
-            scalar: crate::Scalar::BOOL,
-            pointer_space: None,
-        }));
+        let bool_type_id = self.get_type_id(LookupType::Local(LocalType::Numeric(
+            NumericType::Scalar(crate::Scalar::BOOL),
+        )));
         let front_face_id = self.gen_id();
         block.body.push(Instruction::ray_query_get_intersection(
             spirv::Op::RayQueryGetIntersectionFrontFaceKHR,
@@ -212,11 +202,12 @@ impl<'w> BlockContext<'w> {
             intersection_id,
         ));
 
-        let transform_type_id = self.get_type_id(LookupType::Local(LocalType::Matrix {
-            columns: crate::VectorSize::Quad,
-            rows: crate::VectorSize::Tri,
-            width: 4,
-        }));
+        let transform_type_id =
+            self.get_type_id(LookupType::Local(LocalType::Numeric(NumericType::Matrix {
+                columns: crate::VectorSize::Quad,
+                rows: crate::VectorSize::Tri,
+                scalar: crate::Scalar::F32,
+            })));
         let object_to_world_id = self.gen_id();
         block.body.push(Instruction::ray_query_get_intersection(
             spirv::Op::RayQueryGetIntersectionObjectToWorldKHR,
diff --git a/naga/src/back/spv/subgroup.rs b/naga/src/back/spv/subgroup.rs
index c952cb11a..eb273f5f1 100644
--- a/naga/src/back/spv/subgroup.rs
+++ b/naga/src/back/spv/subgroup.rs
@@ -1,4 +1,4 @@
-use super::{Block, BlockContext, Error, Instruction};
+use super::{Block, BlockContext, Error, Instruction, NumericType};
 use crate::{
     arena::Handle,
     back::spv::{LocalType, LookupType},
@@ -16,11 +16,11 @@ impl<'w> BlockContext<'w> {
             "GroupNonUniformBallot",
             &[spirv::Capability::GroupNonUniformBallot],
         )?;
-        let vec4_u32_type_id = self.get_type_id(LookupType::Local(LocalType::Value {
-            vector_size: Some(crate::VectorSize::Quad),
-            scalar: crate::Scalar::U32,
-            pointer_space: None,
-        }));
+        let vec4_u32_type_id =
+            self.get_type_id(LookupType::Local(LocalType::Numeric(NumericType::Vector {
+                size: crate::VectorSize::Quad,
+                scalar: crate::Scalar::U32,
+            })));
         let exec_scope_id = self.get_index_constant(spirv::Scope::Subgroup as u32);
         let predicate = if let Some(predicate) = *predicate {
             self.cached[predicate]
diff --git a/naga/src/back/spv/writer.rs b/naga/src/back/spv/writer.rs
index cd02c2219..cfff00be4 100644
--- a/naga/src/back/spv/writer.rs
+++ b/naga/src/back/spv/writer.rs
@@ -1,10 +1,10 @@
 use super::{
     block::DebugInfoInner,
     helpers::{contains_builtin, global_needs_wrapper, map_storage_class},
-    make_local, Block, BlockContext, CachedConstant, CachedExpressions, DebugInfo,
-    EntryPointContext, Error, Function, FunctionArgument, GlobalVariable, IdGenerator, Instruction,
-    LocalType, LocalVariable, LogicalLayout, LookupFunctionType, LookupType, Options,
-    PhysicalLayout, PipelineOptions, ResultMember, Writer, WriterFlags, BITS_PER_BYTE,
+    Block, BlockContext, CachedConstant, CachedExpressions, DebugInfo, EntryPointContext, Error,
+    Function, FunctionArgument, GlobalVariable, IdGenerator, Instruction, LocalType, LocalVariable,
+    LogicalLayout, LookupFunctionType, LookupType, NumericType, Options, PhysicalLayout,
+    PipelineOptions, ResultMember, Writer, WriterFlags, BITS_PER_BYTE,
 };
 use crate::{
     arena::{Handle, HandleVec, UniqueArena},
@@ -254,7 +254,9 @@ impl Writer {
     pub(super) fn get_expression_lookup_type(&mut self, tr: &TypeResolution) -> LookupType {
         match *tr {
             TypeResolution::Handle(ty_handle) => LookupType::Handle(ty_handle),
-            TypeResolution::Value(ref inner) => LookupType::Local(make_local(inner).unwrap()),
+            TypeResolution::Value(ref inner) => {
+                LookupType::Local(LocalType::from_inner(inner).unwrap())
+            }
         }
     }
 
@@ -289,83 +291,52 @@ impl Writer {
     }
 
     pub(super) fn get_uint_type_id(&mut self) -> Word {
-        let local_type = LocalType::Value {
-            vector_size: None,
-            scalar: crate::Scalar::U32,
-            pointer_space: None,
-        };
+        let local_type = LocalType::Numeric(NumericType::Scalar(crate::Scalar::U32));
         self.get_type_id(local_type.into())
     }
 
     pub(super) fn get_float_type_id(&mut self) -> Word {
-        let local_type = LocalType::Value {
-            vector_size: None,
-            scalar: crate::Scalar::F32,
-            pointer_space: None,
-        };
+        let local_type = LocalType::Numeric(NumericType::Scalar(crate::Scalar::F32));
         self.get_type_id(local_type.into())
     }
 
     pub(super) fn get_uint3_type_id(&mut self) -> Word {
-        let local_type = LocalType::Value {
-            vector_size: Some(crate::VectorSize::Tri),
+        let local_type = LocalType::Numeric(NumericType::Vector {
+            size: crate::VectorSize::Tri,
             scalar: crate::Scalar::U32,
-            pointer_space: None,
-        };
+        });
         self.get_type_id(local_type.into())
     }
 
     pub(super) fn get_float_pointer_type_id(&mut self, class: spirv::StorageClass) -> Word {
-        let lookup_type = LookupType::Local(LocalType::Value {
-            vector_size: None,
-            scalar: crate::Scalar::F32,
-            pointer_space: Some(class),
-        });
-        if let Some(&id) = self.lookup_type.get(&lookup_type) {
-            id
-        } else {
-            let id = self.id_gen.next();
-            let ty_id = self.get_float_type_id();
-            let instruction = Instruction::type_pointer(id, class, ty_id);
-            instruction.to_words(&mut self.logical_layout.declarations);
-            self.lookup_type.insert(lookup_type, id);
-            id
-        }
+        let local_type = LocalType::LocalPointer {
+            base: NumericType::Scalar(crate::Scalar::F32),
+            class,
+        };
+        self.get_type_id(local_type.into())
     }
 
     pub(super) fn get_uint3_pointer_type_id(&mut self, class: spirv::StorageClass) -> Word {
-        let lookup_type = LookupType::Local(LocalType::Value {
-            vector_size: Some(crate::VectorSize::Tri),
-            scalar: crate::Scalar::U32,
-            pointer_space: Some(class),
-        });
-        if let Some(&id) = self.lookup_type.get(&lookup_type) {
-            id
-        } else {
-            let id = self.id_gen.next();
-            let ty_id = self.get_uint3_type_id();
-            let instruction = Instruction::type_pointer(id, class, ty_id);
-            instruction.to_words(&mut self.logical_layout.declarations);
-            self.lookup_type.insert(lookup_type, id);
-            id
-        }
+        let local_type = LocalType::LocalPointer {
+            base: NumericType::Vector {
+                size: crate::VectorSize::Tri,
+                scalar: crate::Scalar::U32,
+            },
+            class,
+        };
+        self.get_type_id(local_type.into())
     }
 
     pub(super) fn get_bool_type_id(&mut self) -> Word {
-        let local_type = LocalType::Value {
-            vector_size: None,
-            scalar: crate::Scalar::BOOL,
-            pointer_space: None,
-        };
+        let local_type = LocalType::Numeric(NumericType::Scalar(crate::Scalar::BOOL));
         self.get_type_id(local_type.into())
     }
 
     pub(super) fn get_bool3_type_id(&mut self) -> Word {
-        let local_type = LocalType::Value {
-            vector_size: Some(crate::VectorSize::Tri),
+        let local_type = LocalType::Numeric(NumericType::Vector {
+            size: crate::VectorSize::Tri,
             scalar: crate::Scalar::BOOL,
-            pointer_space: None,
-        };
+        });
         self.get_type_id(local_type.into())
     }
 
@@ -933,62 +904,50 @@ impl Writer {
         Ok(())
     }
 
+    fn write_numeric_type_declaration_local(&mut self, id: Word, numeric: NumericType) {
+        let instruction =
+            match numeric {
+                NumericType::Scalar(scalar) => self.make_scalar(id, scalar),
+                NumericType::Vector { size, scalar } => {
+                    let scalar_id = self.get_type_id(LookupType::Local(LocalType::Numeric(
+                        NumericType::Scalar(scalar),
+                    )));
+                    Instruction::type_vector(id, scalar_id, size)
+                }
+                NumericType::Matrix {
+                    columns,
+                    rows,
+                    scalar,
+                } => {
+                    let column_id = self.get_type_id(LookupType::Local(LocalType::Numeric(
+                        NumericType::Vector { size: rows, scalar },
+                    )));
+                    Instruction::type_matrix(id, column_id, columns)
+                }
+            };
+
+        instruction.to_words(&mut self.logical_layout.declarations);
+    }
+
     fn write_type_declaration_local(&mut self, id: Word, local_ty: LocalType) {
         let instruction = match local_ty {
-            LocalType::Value {
-                vector_size: None,
-                scalar,
-                pointer_space: None,
-            } => self.make_scalar(id, scalar),
-            LocalType::Value {
-                vector_size: Some(size),
-                scalar,
-                pointer_space: None,
-            } => {
-                let scalar_id = self.get_type_id(LookupType::Local(LocalType::Value {
-                    vector_size: None,
-                    scalar,
-                    pointer_space: None,
-                }));
-                Instruction::type_vector(id, scalar_id, size)
+            LocalType::Numeric(numeric) => {
+                self.write_numeric_type_declaration_local(id, numeric);
+                return;
             }
-            LocalType::Matrix {
-                columns,
-                rows,
-                width,
-            } => {
-                let vector_id = self.get_type_id(LookupType::Local(LocalType::Value {
-                    vector_size: Some(rows),
-                    scalar: crate::Scalar::float(width),
-                    pointer_space: None,
-                }));
-                Instruction::type_matrix(id, vector_id, columns)
+            LocalType::LocalPointer { base, class } => {
+                let base_id = self.get_type_id(LookupType::Local(LocalType::Numeric(base)));
+                Instruction::type_pointer(id, class, base_id)
             }
             LocalType::Pointer { base, class } => {
                 let type_id = self.get_type_id(LookupType::Handle(base));
                 Instruction::type_pointer(id, class, type_id)
             }
-            LocalType::Value {
-                vector_size,
-                scalar,
-                pointer_space: Some(class),
-            } => {
-                let type_id = self.get_type_id(LookupType::Local(LocalType::Value {
-                    vector_size,
-                    scalar,
-                    pointer_space: None,
-                }));
-                Instruction::type_pointer(id, class, type_id)
-            }
             LocalType::Image(image) => {
-                let local_type = LocalType::Value {
-                    vector_size: None,
-                    scalar: crate::Scalar {
-                        kind: image.sampled_type,
-                        width: 4,
-                    },
-                    pointer_space: None,
-                };
+                let local_type = LocalType::Numeric(NumericType::Scalar(crate::Scalar {
+                    kind: image.sampled_type,
+                    width: 4,
+                }));
                 let type_id = self.get_type_id(LookupType::Local(local_type));
                 Instruction::type_image(id, type_id, image.dim, image.flags, image.image_format)
             }
@@ -1025,7 +984,7 @@ impl Writer {
         // because some types which map to the same LocalType have different
         // capability requirements. See https://github.com/gfx-rs/wgpu/issues/5569
         self.request_type_capabilities(&ty.inner)?;
-        let id = if let Some(local) = make_local(&ty.inner) {
+        let id = if let Some(local) = LocalType::from_inner(&ty.inner) {
             // This type can be represented as a `LocalType`, so check if we've
             // already written an instruction for it. If not, do so now, with
             // `write_type_declaration_local`.
@@ -1222,11 +1181,9 @@ impl Writer {
                 self.debugs.push(Instruction::name(id, name));
             }
         }
-        let type_id = self.get_type_id(LookupType::Local(LocalType::Value {
-            vector_size: None,
-            scalar: value.scalar(),
-            pointer_space: None,
-        }));
+        let type_id = self.get_type_id(LookupType::Local(LocalType::Numeric(NumericType::Scalar(
+            value.scalar(),
+        ))));
         let instruction = match *value {
             crate::Literal::F64(value) => {
                 let bits = value.to_bits();


Test patch:
diff --git a/naga/tests/out/msl/atomicCompareExchange.msl b/naga/tests/out/msl/atomicCompareExchange.msl
new file mode 100644
index 000000000..800b5b201
--- /dev/null
+++ b/naga/tests/out/msl/atomicCompareExchange.msl
@@ -0,0 +1,161 @@
+// language: metal1.0
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using metal::uint;
+
+struct type_2 {
+    metal::atomic_int inner[128];
+};
+struct type_4 {
+    metal::atomic_uint inner[128];
+};
+struct _atomic_compare_exchange_resultSint4_ {
+    int old_value;
+    bool exchanged;
+};
+struct _atomic_compare_exchange_resultUint4_ {
+    uint old_value;
+    bool exchanged;
+};
+
+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};
+}
+
+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};
+}
+constant uint SIZE = 128u;
+
+kernel void test_atomic_compare_exchange_i32_(
+  device type_2& arr_i32_ [[user(fake0)]]
+) {
+    uint i = 0u;
+    int old = {};
+    bool exchanged = {};
+#define LOOP_IS_REACHABLE if (volatile bool unpredictable_jump_over_loop = true; unpredictable_jump_over_loop)
+    bool loop_init = true;
+    LOOP_IS_REACHABLE 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 = metal::atomic_load_explicit(&arr_i32_.inner[_e6], metal::memory_order_relaxed);
+            old = _e8;
+            exchanged = false;
+            LOOP_IS_REACHABLE while(true) {
+                bool _e12 = exchanged;
+                if (!(_e12)) {
+                } else {
+                    break;
+                }
+                {
+                    int _e14 = old;
+                    int new_ = as_type<int>(as_type<float>(_e14) + 1.0);
+                    uint _e20 = i;
+                    int _e22 = old;
+                    _atomic_compare_exchange_resultSint4_ _e23 = naga_atomic_compare_exchange_weak_explicit(&arr_i32_.inner[_e20], _e22, new_);
+                    old = _e23.old_value;
+                    exchanged = _e23.exchanged;
+                }
+            }
+        }
+    }
+    return;
+}
+
+
+kernel void test_atomic_compare_exchange_u32_(
+  device type_4& arr_u32_ [[user(fake0)]]
+) {
+    uint i_1 = 0u;
+    uint old_1 = {};
+    bool exchanged_1 = {};
+    bool loop_init_1 = true;
+    LOOP_IS_REACHABLE while(true) {
+        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 = metal::atomic_load_explicit(&arr_u32_.inner[_e6], metal::memory_order_relaxed);
+            old_1 = _e8;
+            exchanged_1 = false;
+            LOOP_IS_REACHABLE while(true) {
+                bool _e12 = exchanged_1;
+                if (!(_e12)) {
+                } else {
+                    break;
+                }
+                {
+                    uint _e14 = old_1;
+                    uint new_1 = as_type<uint>(as_type<float>(_e14) + 1.0);
+                    uint _e20 = i_1;
+                    uint _e22 = old_1;
+                    _atomic_compare_exchange_resultUint4_ _e23 = naga_atomic_compare_exchange_weak_explicit(&arr_u32_.inner[_e20], _e22, new_1);
+                    old_1 = _e23.old_value;
+                    exchanged_1 = _e23.exchanged;
+                }
+            }
+        }
+    }
+    return;
+}
diff --git a/naga/tests/out/msl/overrides-atomicCompareExchangeWeak.msl b/naga/tests/out/msl/overrides-atomicCompareExchangeWeak.msl
new file mode 100644
index 000000000..d87190c59
--- /dev/null
+++ b/naga/tests/out/msl/overrides-atomicCompareExchangeWeak.msl
@@ -0,0 +1,48 @@
+// language: metal1.0
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using metal::uint;
+
+struct _atomic_compare_exchange_resultUint4_ {
+    uint 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};
+}
+constant int o = 2;
+
+kernel void f(
+  metal::uint3 __local_invocation_id [[thread_position_in_threadgroup]]
+, threadgroup metal::atomic_uint& a
+) {
+    if (metal::all(__local_invocation_id == metal::uint3(0u))) {
+        metal::atomic_store_explicit(&a, 0, metal::memory_order_relaxed);
+    }
+    metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
+    _atomic_compare_exchange_resultUint4_ _e5 = naga_atomic_compare_exchange_weak_explicit(&a, 2u, 1u);
+    return;
+}
diff --git a/naga/tests/snapshots.rs b/naga/tests/snapshots.rs
index adf67f833..78b2331dc 100644
--- a/naga/tests/snapshots.rs
+++ b/naga/tests/snapshots.rs
@@ -773,7 +773,10 @@ fn convert_wgsl() {
             "atomicOps",
             Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL,
         ),
-        ("atomicCompareExchange", Targets::SPIRV | Targets::WGSL),
+        (
+            "atomicCompareExchange",
+            Targets::SPIRV | Targets::METAL | Targets::WGSL,
+        ),
         (
             "padding",
             Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL,
@@ -917,7 +920,7 @@ fn convert_wgsl() {
         ),
         (
             "overrides-atomicCompareExchangeWeak",
-            Targets::IR | Targets::SPIRV,
+            Targets::IR | Targets::SPIRV | Targets::METAL,
         ),
         (
             "overrides-ray-query",