Repository: gfx-rs/wgpu (language: rust)

Issue:
Assigning an array of `vec3(0.0)` to an array of `vec3f(0.0)` fails
<!-- Thank you for filing this! Please read the [debugging tips](https://github.com/gfx-rs/wgpu/wiki/Debugging-wgpu-Applications).
That may let you investigate on your own, or provide additional information that helps us to assist.-->

**Description**
Assigning an array of `vec3(0.0)` to an array of `vec3f(0.0)` fails with an error about the types not matching.

**Repro steps**

```
var weird: array<vec3f, 1> = array( vec3(0.0) );
```


**Expected vs observed behavior**
They should be the same type, so it shouldn't fail.

**Extra materials**

The error:

```
Shader validation error: Function [0] 'strange' is invalid
   ┌─ LandscapeMeshPipeline Shader:31:1
   │
31 │ ╭ fn strange() {
32 │ │     var weird: array<vec3<f32>, 1> = array( vec3(0.0) );
   │ │     ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ naga::LocalVariable [0]
   │ ╰────────────────────────────────────────────────────────^ naga::Function [0]
   │
   = Local variable [0] 'weird' is invalid
   = Initializer doesn't match the variable type


      Function [0] 'strange' is invalid
        Local variable [0] 'weird' is invalid
          Initializer doesn't match the variable type
```

**Platform**
Using wgpu 24.0 and vulkan 1.4.303 on Linux.



Discussion / hints:
Fun fact: `vec3f` and `vec3(0.0)` _are_ actually different types. The former is a `vec3<f32>`, the latter is a `vec3<{AbstractInt}>`. We _should_ be converting them, though, via the `array<S,N>` of the conversion rank table in the WGSL spec. ([link](https://www.w3.org/TR/WGSL/#conversion-rank)).

I'm not positive, but it seems the OP's intent is that this seems like a bug, given that conversion from the `array`'s inner components of the OP convert just fine, i.e.:

```wgsl
var weird: vec3f = vec3(0.0);
```

Assigning to @jamienicol, who's been working on other type conversion issues in <https://github.com/gfx-rs/wgpu/issues/4400>.
Fun one. One the face of it of course we should be able to convert an `array<vec3<AbstactFloat>>` to `array<vec3<f32>>`. And it does indeed work fine for `vec2` or `vec4`. But the issue is the stride. 

[Here](https://github.com/gfx-rs/wgpu/blob/43eb6c99b047ff811f0b050b6ceddb1856683b4b/naga/src/front/wgsl/lower/mod.rs#L1473-L1487) is where we lower a local var declaration. `explicit_ty` in this testcase would be `Array { base: Scalar::F32, size: Constant(1), stride: 16 }`. The call to `try_automatic_conversions()` successfully converts this to a `Array { base: Scalar::F32, size: Constant(1), stride: 12 }`. These types are not equal, so validation complains.

[Here](https://github.com/gfx-rs/wgpu/blob/43eb6c99b047ff811f0b050b6ceddb1856683b4b/naga/src/proc/constant_evaluator.rs#L1721) is where we create the new type after casting the array. You can see we just use the size of the base type, which is `12` for `vec3<32>`.

Whereas [here](https://github.com/gfx-rs/wgpu/blob/43eb6c99b047ff811f0b050b6ceddb1856683b4b/naga/src/front/wgsl/lower/mod.rs#L3262) is where we choose the stride for the explict type. We use `proc::Layouter` to properly handle alignment.

I guess we should use a `Layouter` during constant evaluation, too.

Fix patch:
diff --git a/naga/src/back/pipeline_constants.rs b/naga/src/back/pipeline_constants.rs
index ef498b6554..7261c1d858 100644
--- a/naga/src/back/pipeline_constants.rs
+++ b/naga/src/back/pipeline_constants.rs
@@ -75,6 +75,7 @@ pub fn process_overrides<'a>(
     let mut adjusted_constant_initializers = HashSet::with_capacity(module.constants.len());
 
     let mut global_expression_kind_tracker = crate::proc::ExpressionKindTracker::new();
+    let mut layouter = crate::proc::Layouter::default();
 
     // An iterator through the original overrides table, consumed in
     // approximate tandem with the global expressions.
@@ -147,6 +148,7 @@ pub fn process_overrides<'a>(
         let mut evaluator = ConstantEvaluator::for_wgsl_module(
             &mut module,
             &mut global_expression_kind_tracker,
+            &mut layouter,
             false,
         );
         adjust_expr(&adjusted_global_expressions, &mut expr);
@@ -186,13 +188,13 @@ pub fn process_overrides<'a>(
 
     let mut functions = mem::take(&mut module.functions);
     for (_, function) in functions.iter_mut() {
-        process_function(&mut module, &override_map, function)?;
+        process_function(&mut module, &override_map, &mut layouter, function)?;
     }
     module.functions = functions;
 
     let mut entry_points = mem::take(&mut module.entry_points);
     for ep in entry_points.iter_mut() {
-        process_function(&mut module, &override_map, &mut ep.function)?;
+        process_function(&mut module, &override_map, &mut layouter, &mut ep.function)?;
         process_workgroup_size_override(&mut module, &adjusted_global_expressions, ep)?;
     }
     module.entry_points = entry_points;
@@ -365,6 +367,7 @@ fn process_override(
 fn process_function(
     module: &mut Module,
     override_map: &HandleVec<Override, Handle<Constant>>,
+    layouter: &mut crate::proc::Layouter,
     function: &mut Function,
 ) -> Result<(), ConstantEvaluatorError> {
     // A map from original local expression handles to
@@ -389,6 +392,7 @@ fn process_function(
         module,
         &mut function.expressions,
         &mut local_expression_kind_tracker,
+        layouter,
         &mut emitter,
         &mut block,
         false,
diff --git a/naga/src/front/glsl/context.rs b/naga/src/front/glsl/context.rs
index b4cb1c874e..04e89e8397 100644
--- a/naga/src/front/glsl/context.rs
+++ b/naga/src/front/glsl/context.rs
@@ -8,9 +8,9 @@ use super::{
     Frontend, Result,
 };
 use crate::{
-    front::Typifier, proc::Emitter, AddressSpace, Arena, BinaryOperator, Block, Expression,
-    FastHashMap, FunctionArgument, Handle, Literal, LocalVariable, RelationalFunction, Scalar,
-    Span, Statement, Type, TypeInner, VectorSize,
+    front::Typifier, proc::Emitter, proc::Layouter, AddressSpace, Arena, BinaryOperator, Block,
+    Expression, FastHashMap, FunctionArgument, Handle, Literal, LocalVariable, RelationalFunction,
+    Scalar, Span, Statement, Type, TypeInner, VectorSize,
 };
 use std::ops::Index;
 
@@ -72,6 +72,7 @@ pub struct Context<'a> {
 
     pub const_typifier: Typifier,
     pub typifier: Typifier,
+    layouter: Layouter,
     emitter: Emitter,
     stmt_ctx: Option<StmtContext>,
     pub body: Block,
@@ -103,6 +104,7 @@ impl<'a> Context<'a> {
 
             const_typifier: Typifier::new(),
             typifier: Typifier::new(),
+            layouter: Layouter::default(),
             emitter: Emitter::default(),
             stmt_ctx: Some(StmtContext::new()),
             body: Block::new(),
@@ -260,12 +262,14 @@ impl<'a> Context<'a> {
             crate::proc::ConstantEvaluator::for_glsl_module(
                 self.module,
                 self.global_expression_kind_tracker,
+                &mut self.layouter,
             )
         } else {
             crate::proc::ConstantEvaluator::for_glsl_function(
                 self.module,
                 &mut self.expressions,
                 &mut self.local_expression_kind_tracker,
+                &mut self.layouter,
                 &mut self.emitter,
                 &mut self.body,
             )
diff --git a/naga/src/front/wgsl/lower/construction.rs b/naga/src/front/wgsl/lower/construction.rs
index e52d4776ab..ee973fb9f1 100644
--- a/naga/src/front/wgsl/lower/construction.rs
+++ b/naga/src/front/wgsl/lower/construction.rs
@@ -490,8 +490,8 @@ impl<'source> Lowerer<'source, '_> {
                         NonZeroU32::new(u32::try_from(components.len()).unwrap()).unwrap(),
                     ),
                     stride: {
-                        self.layouter.update(ctx.module.to_ctx()).unwrap();
-                        self.layouter[base].to_stride()
+                        ctx.layouter.update(ctx.module.to_ctx()).unwrap();
+                        ctx.layouter[base].to_stride()
                     },
                 };
                 let ty = ctx.ensure_type_exists(inner);
@@ -616,8 +616,8 @@ impl<'source> Lowerer<'source, '_> {
                 let base = self.resolve_ast_type(base, &mut ctx.as_global())?;
                 let size = self.array_size(size, &mut ctx.as_global())?;
 
-                self.layouter.update(ctx.module.to_ctx()).unwrap();
-                let stride = self.layouter[base].to_stride();
+                ctx.layouter.update(ctx.module.to_ctx()).unwrap();
+                let stride = ctx.layouter[base].to_stride();
 
                 let ty = ctx.ensure_type_exists(crate::TypeInner::Array { base, size, stride });
                 Constructor::Type(ty)
diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs
index e405d3e56f..a1bb65533f 100644
--- a/naga/src/front/wgsl/lower/mod.rs
+++ b/naga/src/front/wgsl/lower/mod.rs
@@ -87,6 +87,8 @@ pub struct GlobalContext<'source, 'temp, 'out> {
 
     const_typifier: &'temp mut Typifier,
 
+    layouter: &'temp mut Layouter,
+
     global_expression_kind_tracker: &'temp mut crate::proc::ExpressionKindTracker,
 }
 
@@ -98,6 +100,7 @@ impl<'source> GlobalContext<'source, '_, '_> {
             types: self.types,
             module: self.module,
             const_typifier: self.const_typifier,
+            layouter: self.layouter,
             expr_type: ExpressionContextType::Constant(None),
             global_expression_kind_tracker: self.global_expression_kind_tracker,
         }
@@ -110,6 +113,7 @@ impl<'source> GlobalContext<'source, '_, '_> {
             types: self.types,
             module: self.module,
             const_typifier: self.const_typifier,
+            layouter: self.layouter,
             expr_type: ExpressionContextType::Override,
             global_expression_kind_tracker: self.global_expression_kind_tracker,
         }
@@ -165,6 +169,7 @@ pub struct StatementContext<'source, 'temp, 'out> {
 
     const_typifier: &'temp mut Typifier,
     typifier: &'temp mut Typifier,
+    layouter: &'temp mut Layouter,
     function: &'out mut crate::Function,
     /// Stores the names of expressions that are assigned in `let` statement
     /// Also stores the spans of the names, for use in errors.
@@ -198,6 +203,7 @@ impl<'a, 'temp> StatementContext<'a, 'temp, '_> {
             types: self.types,
             ast_expressions: self.ast_expressions,
             const_typifier: self.const_typifier,
+            layouter: self.layouter,
             global_expression_kind_tracker: self.global_expression_kind_tracker,
             module: self.module,
             expr_type: ExpressionContextType::Constant(Some(LocalExpressionContext {
@@ -224,6 +230,7 @@ impl<'a, 'temp> StatementContext<'a, 'temp, '_> {
             types: self.types,
             ast_expressions: self.ast_expressions,
             const_typifier: self.const_typifier,
+            layouter: self.layouter,
             global_expression_kind_tracker: self.global_expression_kind_tracker,
             module: self.module,
             expr_type: ExpressionContextType::Runtime(LocalExpressionContext {
@@ -244,6 +251,7 @@ impl<'a, 'temp> StatementContext<'a, 'temp, '_> {
             types: self.types,
             module: self.module,
             const_typifier: self.const_typifier,
+            layouter: self.layouter,
             global_expression_kind_tracker: self.global_expression_kind_tracker,
         }
     }
@@ -364,6 +372,7 @@ pub struct ExpressionContext<'source, 'temp, 'out> {
     ///
     /// [`module::global_expressions`]: crate::Module::global_expressions
     const_typifier: &'temp mut Typifier,
+    layouter: &'temp mut Layouter,
     global_expression_kind_tracker: &'temp mut crate::proc::ExpressionKindTracker,
 
     /// Whether we are lowering a constant expression or a general
@@ -379,6 +388,7 @@ impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> {
             types: self.types,
             ast_expressions: self.ast_expressions,
             const_typifier: self.const_typifier,
+            layouter: self.layouter,
             module: self.module,
             expr_type: ExpressionContextType::Constant(match self.expr_type {
                 ExpressionContextType::Runtime(ref mut local_expression_context)
@@ -406,6 +416,7 @@ impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> {
             types: self.types,
             module: self.module,
             const_typifier: self.const_typifier,
+            layouter: self.layouter,
             global_expression_kind_tracker: self.global_expression_kind_tracker,
         }
     }
@@ -416,6 +427,7 @@ impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> {
                 self.module,
                 &mut rctx.function.expressions,
                 rctx.local_expression_kind_tracker,
+                self.layouter,
                 rctx.emitter,
                 rctx.block,
                 false,
@@ -425,6 +437,7 @@ impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> {
                     self.module,
                     &mut rctx.function.expressions,
                     rctx.local_expression_kind_tracker,
+                    self.layouter,
                     rctx.emitter,
                     rctx.block,
                     true,
@@ -433,11 +446,13 @@ impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> {
             ExpressionContextType::Constant(None) => ConstantEvaluator::for_wgsl_module(
                 self.module,
                 self.global_expression_kind_tracker,
+                self.layouter,
                 false,
             ),
             ExpressionContextType::Override => ConstantEvaluator::for_wgsl_module(
                 self.module,
                 self.global_expression_kind_tracker,
+                self.layouter,
                 true,
             ),
         }
@@ -1016,15 +1031,11 @@ impl SubgroupGather {
 
 pub struct Lowerer<'source, 'temp> {
     index: &'temp Index<'source>,
-    layouter: Layouter,
 }
 
 impl<'source, 'temp> Lowerer<'source, 'temp> {
-    pub fn new(index: &'temp Index<'source>) -> Self {
-        Self {
-            index,
-            layouter: Layouter::default(),
-        }
+    pub const fn new(index: &'temp Index<'source>) -> Self {
+        Self { index }
     }
 
     pub fn lower(
@@ -1043,6 +1054,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
             types: &tu.types,
             module: &mut module,
             const_typifier: &mut Typifier::new(),
+            layouter: &mut Layouter::default(),
             global_expression_kind_tracker: &mut crate::proc::ExpressionKindTracker::new(),
         };
 
@@ -1299,6 +1311,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
             ast_expressions: ctx.ast_expressions,
             const_typifier: ctx.const_typifier,
             typifier: &mut typifier,
+            layouter: ctx.layouter,
             function: &mut function,
             named_expressions: &mut named_expressions,
             types: ctx.types,
@@ -3051,10 +3064,10 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
         for member in s.members.iter() {
             let ty = self.resolve_ast_type(member.ty, ctx)?;
 
-            self.layouter.update(ctx.module.to_ctx()).unwrap();
+            ctx.layouter.update(ctx.module.to_ctx()).unwrap();
 
-            let member_min_size = self.layouter[ty].size;
-            let member_min_alignment = self.layouter[ty].alignment;
+            let member_min_size = ctx.layouter[ty].size;
+            let member_min_alignment = ctx.layouter[ty].alignment;
 
             let member_size = if let Some(size_expr) = member.size {
                 let (size, span) = self.const_u32(size_expr, &mut ctx.as_const())?;
@@ -3258,8 +3271,8 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
                 let base = self.resolve_ast_type(base, ctx)?;
                 let size = self.array_size(size, ctx)?;
 
-                self.layouter.update(ctx.module.to_ctx()).unwrap();
-                let stride = self.layouter[base].to_stride();
+                ctx.layouter.update(ctx.module.to_ctx()).unwrap();
+                let stride = ctx.layouter[base].to_stride();
 
                 crate::TypeInner::Array { base, size, stride }
             }
diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs
index c8911077b7..bf97bedc34 100644
--- a/naga/src/proc/constant_evaluator.rs
+++ b/naga/src/proc/constant_evaluator.rs
@@ -312,6 +312,8 @@ pub struct ConstantEvaluator<'a> {
 
     /// Tracks the constness of expressions residing in [`Self::expressions`]
     expression_kind_tracker: &'a mut ExpressionKindTracker,
+
+    layouter: &'a mut crate::proc::Layouter,
 }
 
 #[derive(Debug)]
@@ -594,6 +596,7 @@ impl<'a> ConstantEvaluator<'a> {
     pub fn for_wgsl_module(
         module: &'a mut crate::Module,
         global_expression_kind_tracker: &'a mut ExpressionKindTracker,
+        layouter: &'a mut crate::proc::Layouter,
         in_override_ctx: bool,
     ) -> Self {
         Self::for_module(
@@ -604,6 +607,7 @@ impl<'a> ConstantEvaluator<'a> {
             }),
             module,
             global_expression_kind_tracker,
+            layouter,
         )
     }
 
@@ -614,11 +618,13 @@ impl<'a> ConstantEvaluator<'a> {
     pub fn for_glsl_module(
         module: &'a mut crate::Module,
         global_expression_kind_tracker: &'a mut ExpressionKindTracker,
+        layouter: &'a mut crate::proc::Layouter,
     ) -> Self {
         Self::for_module(
             Behavior::Glsl(GlslRestrictions::Const),
             module,
             global_expression_kind_tracker,
+            layouter,
         )
     }
 
@@ -626,6 +632,7 @@ impl<'a> ConstantEvaluator<'a> {
         behavior: Behavior<'a>,
         module: &'a mut crate::Module,
         global_expression_kind_tracker: &'a mut ExpressionKindTracker,
+        layouter: &'a mut crate::proc::Layouter,
     ) -> Self {
         Self {
             behavior,
@@ -634,6 +641,7 @@ impl<'a> ConstantEvaluator<'a> {
             overrides: &module.overrides,
             expressions: &mut module.global_expressions,
             expression_kind_tracker: global_expression_kind_tracker,
+            layouter,
         }
     }
 
@@ -645,6 +653,7 @@ impl<'a> ConstantEvaluator<'a> {
         module: &'a mut crate::Module,
         expressions: &'a mut Arena<Expression>,
         local_expression_kind_tracker: &'a mut ExpressionKindTracker,
+        layouter: &'a mut crate::proc::Layouter,
         emitter: &'a mut super::Emitter,
         block: &'a mut crate::Block,
         is_const: bool,
@@ -665,6 +674,7 @@ impl<'a> ConstantEvaluator<'a> {
             overrides: &module.overrides,
             expressions,
             expression_kind_tracker: local_expression_kind_tracker,
+            layouter,
         }
     }
 
@@ -676,6 +686,7 @@ impl<'a> ConstantEvaluator<'a> {
         module: &'a mut crate::Module,
         expressions: &'a mut Arena<Expression>,
         local_expression_kind_tracker: &'a mut ExpressionKindTracker,
+        layouter: &'a mut crate::proc::Layouter,
         emitter: &'a mut super::Emitter,
         block: &'a mut crate::Block,
     ) -> Self {
@@ -690,6 +701,7 @@ impl<'a> ConstantEvaluator<'a> {
             overrides: &module.overrides,
             expressions,
             expression_kind_tracker: local_expression_kind_tracker,
+            layouter,
         }
     }
 
@@ -1718,7 +1730,11 @@ impl<'a> ConstantEvaluator<'a> {
                 self.types.insert(Type { name: None, inner }, span)
             }
         };
-        let new_base_stride = self.types[new_base].inner.size(self.to_ctx());
+        let mut layouter = std::mem::take(self.layouter);
+        layouter.update(self.to_ctx()).unwrap();
+        *self.layouter = layouter;
+
+        let new_base_stride = self.layouter[new_base].to_stride();
         let new_array_ty = self.types.insert(
             Type {
                 name: None,
@@ -2567,6 +2583,7 @@ mod tests {
             overrides: &overrides,
             expressions: &mut global_expressions,
             expression_kind_tracker,
+            layouter: &mut crate::proc::Layouter::default(),
         };
 
         let res1 = solver
@@ -2653,6 +2670,7 @@ mod tests {
             overrides: &overrides,
             expressions: &mut global_expressions,
             expression_kind_tracker,
+            layouter: &mut crate::proc::Layouter::default(),
         };
 
         let res = solver
@@ -2771,6 +2789,7 @@ mod tests {
             overrides: &overrides,
             expressions: &mut global_expressions,
             expression_kind_tracker,
+            layouter: &mut crate::proc::Layouter::default(),
         };
 
         let root1 = Expression::AccessIndex { base, index: 1 };
@@ -2864,6 +2883,7 @@ mod tests {
             overrides: &overrides,
             expressions: &mut global_expressions,
             expression_kind_tracker,
+            layouter: &mut crate::proc::Layouter::default(),
         };
 
         let solved_compose = solver
@@ -2946,6 +2966,7 @@ mod tests {
             overrides: &overrides,
             expressions: &mut global_expressions,
             expression_kind_tracker,
+            layouter: &mut crate::proc::Layouter::default(),
         };
 
         let solved_compose = solver
@@ -3034,6 +3055,7 @@ mod tests {
             overrides: &overrides,
             expressions: &mut global_expressions,
             expression_kind_tracker,
+            layouter: &mut crate::proc::Layouter::default(),
         };
 
         let solved_add = solver


Test patch:
diff --git a/naga/tests/in/abstract-types-var.wgsl b/naga/tests/in/abstract-types-var.wgsl
index 4f0bfe44bc..e12a14b372 100644
--- a/naga/tests/in/abstract-types-var.wgsl
+++ b/naga/tests/in/abstract-types-var.wgsl
@@ -43,6 +43,10 @@ var<private> xafpaiaf: array<f32, 2> = array(1,   2.0);
 var<private> xafpafai: array<f32, 2> = array(1.0, 2);
 var<private> xafpafaf: array<f32, 2> = array(1.0, 2.0);
 
+var<private> xavipai: array<vec3<i32>, 1> = array(vec3(1));
+var<private> xavfpai: array<vec3<f32>, 1> = array(vec3(1));
+var<private> xavfpaf: array<vec3<f32>, 1> = array(vec3(1.0));
+
 var<private> ivispai = vec2(1);
 var<private> ivfspaf = vec2(1.0);
 var<private> ivis_ai = vec2<i32>(1);
@@ -58,6 +62,10 @@ var<private> iafpafaf = array(1.0, 2.0);
 var<private> iafpaiaf = array(1, 2.0);
 var<private> iafpafai = array(1.0, 2);
 
+var<private> iavipai = array(vec3(1));
+var<private> iavfpai = array(vec3(1));
+var<private> iavfpaf = array(vec3(1.0));
+
 fn all_constant_arguments() {
     var xvipaiai: vec2<i32> = vec2(42, 43);
     var xvupaiai: vec2<u32> = vec2(44, 45);
@@ -100,6 +108,10 @@ fn all_constant_arguments() {
     var xafpafai: array<f32, 2> = array(1.0, 2);
     var xafpafaf: array<f32, 2> = array(1.0, 2.0);
 
+    var xavipai: array<vec3<i32>, 1> = array(vec3(1));
+    var xavfpai: array<vec3<f32>, 1> = array(vec3(1));
+    var xavfpaf: array<vec3<f32>, 1> = array(vec3(1.0));
+
     var iaipaiai = array(1,   2);
     var iafpaiaf = array(1,   2.0);
     var iafpafai = array(1.0, 2);
@@ -147,6 +159,10 @@ fn all_constant_arguments() {
     xafpafai = array(1.0, 2);
     xafpafaf = array(1.0, 2.0);
 
+    xavipai = array(vec3(1));
+    xavfpai = array(vec3(1));
+    xavfpaf = array(vec3(1.0));
+
     iaipaiai = array(1,   2);
     iafpaiaf = array(1,   2.0);
     iafpafai = array(1.0, 2);
diff --git a/naga/tests/out/msl/abstract-types-var.msl b/naga/tests/out/msl/abstract-types-var.msl
index 1c9d6fcb25..da88a396e8 100644
--- a/naga/tests/out/msl/abstract-types-var.msl
+++ b/naga/tests/out/msl/abstract-types-var.msl
@@ -10,6 +10,12 @@ struct type_7 {
 struct type_8 {
     int inner[2];
 };
+struct type_10 {
+    metal::int3 inner[1];
+};
+struct type_12 {
+    metal::float3 inner[1];
+};
 
 void all_constant_arguments(
 ) {
@@ -46,6 +52,9 @@ void all_constant_arguments(
     type_7 xafpaiaf = type_7 {1.0, 2.0};
     type_7 xafpafai = type_7 {1.0, 2.0};
     type_7 xafpafaf = type_7 {1.0, 2.0};
+    type_10 xavipai = type_10 {metal::int3(1)};
+    type_12 xavfpai = type_12 {metal::float3(1.0)};
+    type_12 xavfpaf = type_12 {metal::float3(1.0)};
     type_8 iaipaiai = type_8 {1, 2};
     type_7 iafpaiaf = type_7 {1.0, 2.0};
     type_7 iafpafai = type_7 {1.0, 2.0};
@@ -83,6 +92,9 @@ void all_constant_arguments(
     xafpaiaf = type_7 {1.0, 2.0};
     xafpafai = type_7 {1.0, 2.0};
     xafpafaf = type_7 {1.0, 2.0};
+    xavipai = type_10 {metal::int3(1)};
+    xavfpai = type_12 {metal::float3(1.0)};
+    xavfpaf = type_12 {metal::float3(1.0)};
     iaipaiai = type_8 {1, 2};
     iafpaiaf = type_7 {1.0, 2.0};
     iafpafai = type_7 {1.0, 2.0};
diff --git a/naga/tests/out/spv/abstract-types-var.spvasm b/naga/tests/out/spv/abstract-types-var.spvasm
index 47249bf278..3c3f42b74b 100644
--- a/naga/tests/out/spv/abstract-types-var.spvasm
+++ b/naga/tests/out/spv/abstract-types-var.spvasm
@@ -1,13 +1,15 @@
 ; SPIR-V
 ; Version: 1.1
 ; Generator: rspirv
-; Bound: 269
+; Bound: 290
 OpCapability Shader
 OpCapability Linkage
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
 OpDecorate %10 ArrayStride 4
 OpDecorate %12 ArrayStride 4
+OpDecorate %14 ArrayStride 16
+OpDecorate %17 ArrayStride 16
 %2 = OpTypeVoid
 %3 = OpTypeInt 32 1
 %4 = OpTypeVector %3 2
@@ -19,342 +21,366 @@ OpDecorate %12 ArrayStride 4
 %11 = OpConstant  %5  2
 %10 = OpTypeArray %7 %11
 %12 = OpTypeArray %3 %11
-%13 = OpConstant  %3  42
-%14 = OpConstant  %3  43
-%15 = OpConstantComposite  %4  %13 %14
-%16 = OpConstant  %5  44
-%17 = OpConstant  %5  45
-%18 = OpConstantComposite  %6  %16 %17
-%19 = OpConstant  %7  46.0
-%20 = OpConstant  %7  47.0
-%21 = OpConstantComposite  %8  %19 %20
-%22 = OpConstant  %5  42
-%23 = OpConstant  %5  43
-%24 = OpConstantComposite  %6  %22 %23
-%25 = OpConstant  %7  1.0
-%26 = OpConstant  %7  2.0
-%27 = OpConstantComposite  %8  %25 %26
-%28 = OpConstant  %7  3.0
-%29 = OpConstant  %7  4.0
-%30 = OpConstantComposite  %8  %28 %29
-%31 = OpConstantComposite  %9  %27 %30
-%32 = OpConstant  %3  1
-%33 = OpConstantComposite  %4  %32 %32
-%34 = OpConstantComposite  %8  %25 %25
-%35 = OpConstant  %5  1
-%36 = OpConstantComposite  %6  %35 %35
-%37 = OpConstantComposite  %10  %25 %26
-%38 = OpConstant  %3  2
-%39 = OpConstantComposite  %12  %32 %38
-%41 = OpTypePointer Private %4
-%40 = OpVariable  %41  Private %15
-%43 = OpTypePointer Private %6
-%42 = OpVariable  %43  Private %18
-%45 = OpTypePointer Private %8
-%44 = OpVariable  %45  Private %21
-%46 = OpVariable  %43  Private %24
-%47 = OpVariable  %43  Private %24
-%48 = OpVariable  %43  Private %24
-%49 = OpVariable  %43  Private %24
-%51 = OpTypePointer Private %9
-%50 = OpVariable  %51  Private %31
-%52 = OpVariable  %51  Private %31
-%53 = OpVariable  %51  Private %31
-%54 = OpVariable  %51  Private %31
-%55 = OpVariable  %51  Private %31
-%56 = OpVariable  %41  Private %33
-%57 = OpVariable  %45  Private %34
-%58 = OpVariable  %41  Private %33
-%59 = OpVariable  %43  Private %36
-%60 = OpVariable  %45  Private %34
-%61 = OpVariable  %45  Private %34
-%63 = OpTypePointer Private %10
-%62 = OpVariable  %63  Private %37
-%64 = OpVariable  %63  Private %37
-%66 = OpTypePointer Private %12
-%65 = OpVariable  %66  Private %39
-%67 = OpVariable  %63  Private %37
-%68 = OpVariable  %63  Private %37
-%69 = OpVariable  %63  Private %37
-%70 = OpVariable  %41  Private %33
-%71 = OpVariable  %45  Private %34
-%72 = OpVariable  %41  Private %33
-%73 = OpVariable  %43  Private %36
-%74 = OpVariable  %45  Private %34
-%75 = OpVariable  %45  Private %34
-%76 = OpVariable  %63  Private %37
-%77 = OpVariable  %63  Private %37
-%78 = OpVariable  %66  Private %39
-%79 = OpVariable  %63  Private %37
-%80 = OpVariable  %63  Private %37
-%81 = OpVariable  %63  Private %37
-%84 = OpTypeFunction %2
-%86 = OpTypePointer Function %4
-%88 = OpTypePointer Function %6
-%90 = OpTypePointer Function %8
-%96 = OpTypePointer Function %9
-%112 = OpTypePointer Function %10
-%117 = OpTypePointer Function %12
-%132 = OpTypePointer Function %5
-%133 = OpConstantNull  %5
-%135 = OpTypePointer Function %3
-%136 = OpConstantNull  %3
-%138 = OpTypePointer Function %7
-%139 = OpConstantNull  %7
-%141 = OpConstantNull  %6
-%143 = OpConstantNull  %6
-%145 = OpConstantNull  %6
-%147 = OpConstantNull  %6
-%149 = OpConstantNull  %9
-%151 = OpConstantNull  %9
-%153 = OpConstantNull  %9
-%155 = OpConstantNull  %9
-%157 = OpConstantNull  %10
-%159 = OpConstantNull  %10
-%161 = OpConstantNull  %10
-%163 = OpConstantNull  %10
-%165 = OpConstantNull  %12
-%167 = OpConstantNull  %12
-%169 = OpConstantNull  %10
-%171 = OpConstantNull  %10
-%173 = OpConstantNull  %10
-%175 = OpConstantNull  %10
-%177 = OpConstantNull  %12
-%179 = OpConstantNull  %12
-%83 = OpFunction  %2  None %84
-%82 = OpLabel
-%127 = OpVariable  %112  Function %37
-%124 = OpVariable  %117  Function %39
-%121 = OpVariable  %112  Function %37
-%118 = OpVariable  %117  Function %39
-%114 = OpVariable  %112  Function %37
-%110 = OpVariable  %90  Function %34
-%107 = OpVariable  %86  Function %33
-%104 = OpVariable  %96  Function %31
-%101 = OpVariable  %96  Function %31
-%98 = OpVariable  %96  Function %31
-%94 = OpVariable  %88  Function %24
-%91 = OpVariable  %88  Function %24
-%85 = OpVariable  %86  Function %15
-%125 = OpVariable  %112  Function %37
-%122 = OpVariable  %112  Function %37
-%119 = OpVariable  %117  Function %39
-%115 = OpVariable  %112  Function %37
-%111 = OpVariable  %112  Function %37
-%108 = OpVariable  %88  Function %36
-%105 = OpVariable  %86  Function %33
-%102 = OpVariable  %96  Function %31
-%99 = OpVariable  %96  Function %31
-%95 = OpVariable  %96  Function %31
-%92 = OpVariable  %88  Function %24
-%87 = OpVariable  %88  Function %18
-%126 = OpVariable  %112  Function %37
-%123 = OpVariable  %112  Function %37
-%120 = OpVariable  %112  Function %37
-%116 = OpVariable  %117  Function %39
-%113 = OpVariable  %112  Function %37
-%109 = OpVariable  %90  Function %34
-%106 = OpVariable  %90  Function %34
-%103 = OpVariable  %96  Function %31
-%100 = OpVariable  %96  Function %31
-%97 = OpVariable  %96  Function %31
-%93 = OpVariable  %88  Function %24
-%89 = OpVariable  %90  Function %21
-OpBranch %128
-%128 = OpLabel
-OpStore %85 %15
-OpStore %87 %18
-OpStore %89 %21
-OpStore %91 %24
-OpStore %92 %24
-OpStore %93 %24
-OpStore %94 %24
-OpStore %95 %31
-OpStore %97 %31
-OpStore %98 %31
-OpStore %99 %31
-OpStore %100 %31
-OpStore %101 %31
-OpStore %102 %31
-OpStore %103 %31
-OpStore %104 %31
-OpStore %105 %33
-OpStore %106 %34
-OpStore %107 %33
-OpStore %108 %36
-OpStore %109 %34
-OpStore %110 %34
-OpStore %111 %37
-OpStore %113 %37
-OpStore %114 %37
-OpStore %115 %37
-OpStore %116 %39
-OpStore %118 %39
-OpStore %119 %39
-OpStore %120 %37
-OpStore %121 %37
-OpStore %122 %37
-OpStore %123 %37
-OpStore %124 %39
-OpStore %125 %37
-OpStore %126 %37
-OpStore %127 %37
+%13 = OpTypeVector %3 3
+%15 = OpConstant  %5  1
+%14 = OpTypeArray %13 %15
+%16 = OpTypeVector %7 3
+%17 = OpTypeArray %16 %15
+%18 = OpConstant  %3  42
+%19 = OpConstant  %3  43
+%20 = OpConstantComposite  %4  %18 %19
+%21 = OpConstant  %5  44
+%22 = OpConstant  %5  45
+%23 = OpConstantComposite  %6  %21 %22
+%24 = OpConstant  %7  46.0
+%25 = OpConstant  %7  47.0
+%26 = OpConstantComposite  %8  %24 %25
+%27 = OpConstant  %5  42
+%28 = OpConstant  %5  43
+%29 = OpConstantComposite  %6  %27 %28
+%30 = OpConstant  %7  1.0
+%31 = OpConstant  %7  2.0
+%32 = OpConstantComposite  %8  %30 %31
+%33 = OpConstant  %7  3.0
+%34 = OpConstant  %7  4.0
+%35 = OpConstantComposite  %8  %33 %34
+%36 = OpConstantComposite  %9  %32 %35
+%37 = OpConstant  %3  1
+%38 = OpConstantComposite  %4  %37 %37
+%39 = OpConstantComposite  %8  %30 %30
+%40 = OpConstantComposite  %6  %15 %15
+%41 = OpConstantComposite  %10  %30 %31
+%42 = OpConstant  %3  2
+%43 = OpConstantComposite  %12  %37 %42
+%44 = OpConstantComposite  %13  %37 %37 %37
+%45 = OpConstantComposite  %14  %44
+%46 = OpConstantComposite  %16  %30 %30 %30
+%47 = OpConstantComposite  %17  %46
+%49 = OpTypePointer Private %4
+%48 = OpVariable  %49  Private %20
+%51 = OpTypePointer Private %6
+%50 = OpVariable  %51  Private %23
+%53 = OpTypePointer Private %8
+%52 = OpVariable  %53  Private %26
+%54 = OpVariable  %51  Private %29
+%55 = OpVariable  %51  Private %29
+%56 = OpVariable  %51  Private %29
+%57 = OpVariable  %51  Private %29
+%59 = OpTypePointer Private %9
+%58 = OpVariable  %59  Private %36
+%60 = OpVariable  %59  Private %36
+%61 = OpVariable  %59  Private %36
+%62 = OpVariable  %59  Private %36
+%63 = OpVariable  %59  Private %36
+%64 = OpVariable  %49  Private %38
+%65 = OpVariable  %53  Private %39
+%66 = OpVariable  %49  Private %38
+%67 = OpVariable  %51  Private %40
+%68 = OpVariable  %53  Private %39
+%69 = OpVariable  %53  Private %39
+%71 = OpTypePointer Private %10
+%70 = OpVariable  %71  Private %41
+%72 = OpVariable  %71  Private %41
+%74 = OpTypePointer Private %12
+%73 = OpVariable  %74  Private %43
+%75 = OpVariable  %71  Private %41
+%76 = OpVariable  %71  Private %41
+%77 = OpVariable  %71  Private %41
+%79 = OpTypePointer Private %14
+%78 = OpVariable  %79  Private %45
+%81 = OpTypePointer Private %17
+%80 = OpVariable  %81  Private %47
+%82 = OpVariable  %81  Private %47
+%83 = OpVariable  %49  Private %38
+%84 = OpVariable  %53  Private %39
+%85 = OpVariable  %49  Private %38
+%86 = OpVariable  %51  Private %40
+%87 = OpVariable  %53  Private %39
+%88 = OpVariable  %53  Private %39
+%89 = OpVariable  %71  Private %41
+%90 = OpVariable  %71  Private %41
+%91 = OpVariable  %74  Private %43
+%92 = OpVariable  %71  Private %41
+%93 = OpVariable  %71  Private %41
+%94 = OpVariable  %71  Private %41
+%95 = OpVariable  %79  Private %45
+%96 = OpVariable  %79  Private %45
+%97 = OpVariable  %81  Private %47
+%100 = OpTypeFunction %2
+%102 = OpTypePointer Function %4
+%104 = OpTypePointer Function %6
+%106 = OpTypePointer Function %8
+%112 = OpTypePointer Function %9
+%128 = OpTypePointer Function %10
+%133 = OpTypePointer Function %12
+%141 = OpTypePointer Function %14
+%143 = OpTypePointer Function %17
+%153 = OpTypePointer Function %5
+%154 = OpConstantNull  %5
+%156 = OpTypePointer Function %3
+%157 = OpConstantNull  %3
+%159 = OpTypePointer Function %7
+%160 = OpConstantNull  %7
+%162 = OpConstantNull  %6
+%164 = OpConstantNull  %6
+%166 = OpConstantNull  %6
+%168 = OpConstantNull  %6
+%170 = OpConstantNull  %9
+%172 = OpConstantNull  %9
+%174 = OpConstantNull  %9
+%176 = OpConstantNull  %9
+%178 = OpConstantNull  %10
+%180 = OpConstantNull  %10
+%182 = OpConstantNull  %10
+%184 = OpConstantNull  %10
+%186 = OpConstantNull  %12
+%188 = OpConstantNull  %12
+%190 = OpConstantNull  %10
+%192 = OpConstantNull  %10
+%194 = OpConstantNull  %10
+%196 = OpConstantNull  %10
+%198 = OpConstantNull  %12
+%200 = OpConstantNull  %12
+%99 = OpFunction  %2  None %100
+%98 = OpLabel
+%148 = OpVariable  %128  Function %41
+%145 = OpVariable  %133  Function %43
+%140 = OpVariable  %141  Function %45
+%137 = OpVariable  %128  Function %41
+%134 = OpVariable  %133  Function %43
+%130 = OpVariable  %128  Function %41
+%126 = OpVariable  %106  Function %39
+%123 = OpVariable  %102  Function %38
+%120 = OpVariable  %112  Function %36
+%117 = OpVariable  %112  Function %36
+%114 = OpVariable  %112  Function %36
+%110 = OpVariable  %104  Function %29
+%107 = OpVariable  %104  Function %29
+%101 = OpVariable  %102  Function %20
+%146 = OpVariable  %128  Function %41
+%142 = OpVariable  %143  Function %47
+%138 = OpVariable  %128  Function %41
+%135 = OpVariable  %133  Function %43
+%131 = OpVariable  %128  Function %41
+%127 = OpVariable  %128  Function %41
+%124 = OpVariable  %104  Function %40
+%121 = OpVariable  %102  Function %38
+%118 = OpVariable  %112  Function %36
+%115 = OpVariable  %112  Function %36
+%111 = OpVariable  %112  Function %36
+%108 = OpVariable  %104  Function %29
+%103 = OpVariable  %104  Function %23
+%147 = OpVariable  %128  Function %41
+%144 = OpVariable  %143  Function %47
+%139 = OpVariable  %128  Function %41
+%136 = OpVariable  %128  Function %41
+%132 = OpVariable  %133  Function %43
+%129 = OpVariable  %128  Function %41
+%125 = OpVariable  %106  Function %39
+%122 = OpVariable  %106  Function %39
+%119 = OpVariable  %112  Function %36
+%116 = OpVariable  %112  Function %36
+%113 = OpVariable  %112  Function %36
+%109 = OpVariable  %104  Function %29
+%105 = OpVariable  %106  Function %26
+OpBranch %149
+%149 = OpLabel
+OpStore %101 %20
+OpStore %103 %23
+OpStore %105 %26
+OpStore %107 %29
+OpStore %108 %29
+OpStore %109 %29
+OpStore %110 %29
+OpStore %111 %36
+OpStore %113 %36
+OpStore %114 %36
+OpStore %115 %36
+OpStore %116 %36
+OpStore %117 %36
+OpStore %118 %36
+OpStore %119 %36
+OpStore %120 %36
+OpStore %121 %38
+OpStore %122 %39
+OpStore %123 %38
+OpStore %124 %40
+OpStore %125 %39
+OpStore %126 %39
+OpStore %127 %41
+OpStore %129 %41
+OpStore %130 %41
+OpStore %131 %41
+OpStore %132 %43
+OpStore %134 %43
+OpStore %135 %43
+OpStore %136 %41
+OpStore %137 %41
+OpStore %138 %41
+OpStore %139 %41
+OpStore %140 %45
+OpStore %142 %47
+OpStore %144 %47
+OpStore %145 %43
+OpStore %146 %41
+OpStore %147 %41
+OpStore %148 %41
 OpReturn
 OpFunctionEnd
-%130 = OpFunction  %2  None %84
-%129 = OpLabel
-%178 = OpVariable  %117  Function %179
-%172 = OpVariable  %112  Function %173
-%166 = OpVariable  %117  Function %167
-%160 = OpVariable  %112  Function %161
-%154 = OpVariable  %96  Function %155
-%148 = OpVariable  %96  Function %149
-%142 = OpVariable  %88  Function %143
-%134 = OpVariable  %135  Function %136
-%176 = OpVariable  %117  Function %177
-%170 = OpVariable  %112  Function %171
-%164 = OpVariable  %117  Function %165
-%158 = OpVariable  %112  Function %159
-%152 = OpVariable  %96  Function %153
-%146 = OpVariable  %88  Function %147
-%140 = OpVariable  %88  Function %141
-%131 = OpVariable  %132  Function %133
-%174 = OpVariable  %112  Function %175
-%168 = OpVariable  %112  Function %169
-%162 = OpVariable  %112  Function %163
-%156 = OpVariable  %112  Function %157
-%150 = OpVariable  %96  Function %151
-%144 = OpVariable  %88  Function %145
-%137 = OpVariable  %138  Function %139
-OpBranch %180
-%180 = OpLabel
-%181 = OpLoad  %5  %131
-%182 = OpCompositeConstruct  %6  %181 %23
-OpStore %140 %182
-%183 = OpLoad  %5  %131
-%184 = OpCompositeConstruct  %6  %22 %183
-OpStore %142 %184
-%185 = OpLoad  %5  %131
-%186 = OpCompositeConstruct  %6  %185 %23
-OpStore %144 %186
-%187 = OpLoad  %5  %131
-%188 = OpCompositeConstruct  %6  %22 %187
-OpStore %146 %188
-%189 = OpLoad  %7  %137
-%190 = OpCompositeConstruct  %8  %189 %26
-%191 = OpCompositeConstruct  %9  %190 %30
-OpStore %148 %191
-%192 = OpLoad  %7  %137
-%193 = OpCompositeConstruct  %8  %25 %192
-%194 = OpCompositeConstruct  %9  %193 %30
-OpStore %150 %194
-%195 = OpLoad  %7  %137
-%196 = OpCompositeConstruct  %8  %195 %29
-%197 = OpCompositeConstruct  %9  %27 %196
-OpStore %152 %197
-%198 = OpLoad  %7  %137
-%199 = OpCompositeConstruct  %8  %28 %198
-%200 = OpCompositeConstruct  %9  %27 %199
-OpStore %154 %200
-%201 = OpLoad  %7  %137
-%202 = OpCompositeConstruct  %10  %201 %26
-OpStore %156 %202
-%203 = OpLoad  %7  %137
-%204 = OpCompositeConstruct  %10  %25 %203
-OpStore %158 %204
-%205 = OpLoad  %7  %137
-%206 = OpCompositeConstruct  %10  %205 %26
-OpStore %160 %206
-%207 = OpLoad  %7  %137
-%208 = OpCompositeConstruct  %10  %25 %207
-OpStore %162 %208
-%209 = OpLoad  %3  %134
-%210 = OpCompositeConstruct  %12  %209 %38
-OpStore %164 %210
-%211 = OpLoad  %3  %134
-%212 = OpCompositeConstruct  %12  %32 %211
-OpStore %166 %212
-%213 = OpLoad  %7  %137
-%214 = OpCompositeConstruct  %10  %213 %26
-OpStore %168 %214
-%215 = OpLoad  %7  %137
-%216 = OpCompositeConstruct  %10  %25 %215
-OpStore %170 %216
-%217 = OpLoad  %7  %137
-%218 = OpCompositeConstruct  %10  %217 %26
-OpStore %172 %218
-%219 = OpLoad  %7  %137
-%220 = OpCompositeConstruct  %10  %25 %219
-OpStore %174 %220
-%221 = OpLoad  %3  %134
-%222 = OpCompositeConstruct  %12  %221 %38
-OpStore %176 %222
-%223 = OpLoad  %3  %134
-%224 = OpCompositeConstruct  %12  %32 %223
-OpStore %178 %224
-%225 = OpLoad  %5  %131
-%226 = OpCompositeConstruct  %6  %225 %23
-OpStore %140 %226
-%227 = OpLoad  %5  %131
-%228 = OpCompositeConstruct  %6  %22 %227
-OpStore %142 %228
-%229 = OpLoad  %5  %131
-%230 = OpCompositeConstruct  %6  %229 %23
-OpStore %144 %230
-%231 = OpLoad  %5  %131
-%232 = OpCompositeConstruct  %6  %22 %231
-OpStore %146 %232
-%233 = OpLoad  %7  %137
-%234 = OpCompositeConstruct  %8  %233 %26
-%235 = OpCompositeConstruct  %9  %234 %30
-OpStore %148 %235
-%236 = OpLoad  %7  %137
-%237 = OpCompositeConstruct  %8  %25 %236
-%238 = OpCompositeConstruct  %9  %237 %30
-OpStore %150 %238
-%239 = OpLoad  %7  %137
-%240 = OpCompositeConstruct  %8  %239 %29
-%241 = OpCompositeConstruct  %9  %27 %240
-OpStore %152 %241
-%242 = OpLoad  %7  %137
-%243 = OpCompositeConstruct  %8  %28 %242
-%244 = OpCompositeConstruct  %9  %27 %243
-OpStore %154 %244
-%245 = OpLoad  %7  %137
-%246 = OpCompositeConstruct  %10  %245 %26
-OpStore %156 %246
-%247 = OpLoad  %7  %137
-%248 = OpCompositeConstruct  %10  %25 %247
-OpStore %158 %248
-%249 = OpLoad  %7  %137
-%250 = OpCompositeConstruct  %10  %249 %26
-OpStore %160 %250
-%251 = OpLoad  %7  %137
-%252 = OpCompositeConstruct  %10  %25 %251
-OpStore %162 %252
-%253 = OpLoad  %3  %134
-%254 = OpCompositeConstruct  %12  %253 %38
-OpStore %164 %254
-%255 = OpLoad  %3  %134
-%256 = OpCompositeConstruct  %12  %32 %255
-OpStore %166 %256
-%257 = OpLoad  %7  %137
-%258 = OpCompositeConstruct  %10  %257 %26
-OpStore %168 %258
-%259 = OpLoad  %7  %137
-%260 = OpCompositeConstruct  %10  %25 %259
-OpStore %170 %260
-%261 = OpLoad  %7  %137
-%262 = OpCompositeConstruct  %10  %261 %26
-OpStore %172 %262
-%263 = OpLoad  %7  %137
-%264 = OpCompositeConstruct  %10  %25 %263
-OpStore %174 %264
-%265 = OpLoad  %3  %134
-%266 = OpCompositeConstruct  %12  %265 %38
-OpStore %176 %266
-%267 = OpLoad  %3  %134
-%268 = OpCompositeConstruct  %12  %32 %267
-OpStore %178 %268
+%151 = OpFunction  %2  None %100
+%150 = OpLabel
+%199 = OpVariable  %133  Function %200
+%193 = OpVariable  %128  Function %194
+%187 = OpVariable  %133  Function %188
+%181 = OpVariable  %128  Function %182
+%175 = OpVariable  %112  Function %176
+%169 = OpVariable  %112  Function %170
+%163 = OpVariable  %104  Function %164
+%155 = OpVariable  %156  Function %157
+%197 = OpVariable  %133  Function %198
+%191 = OpVariable  %128  Function %192
+%185 = OpVariable  %133  Function %186
+%179 = OpVariable  %128  Function %180
+%173 = OpVariable  %112  Function %174
+%167 = OpVariable  %104  Function %168
+%161 = OpVariable  %104  Function %162
+%152 = OpVariable  %153  Function %154
+%195 = OpVariable  %128  Function %196
+%189 = OpVariable  %128  Function %190
+%183 = OpVariable  %128  Function %184
+%177 = OpVariable  %128  Function %178
+%171 = OpVariable  %112  Function %172
+%165 = OpVariable  %104  Function %166
+%158 = OpVariable  %159  Function %160
+OpBranch %201
+%201 = OpLabel
+%202 = OpLoad  %5  %152
+%203 = OpCompositeConstruct  %6  %202 %28
+OpStore %161 %203
+%204 = OpLoad  %5  %152
+%205 = OpCompositeConstruct  %6  %27 %204
+OpStore %163 %205
+%206 = OpLoad  %5  %152
+%207 = OpCompositeConstruct  %6  %206 %28
+OpStore %165 %207
+%208 = OpLoad  %5  %152
+%209 = OpCompositeConstruct  %6  %27 %208
+OpStore %167 %209
+%210 = OpLoad  %7  %158
+%211 = OpCompositeConstruct  %8  %210 %31
+%212 = OpCompositeConstruct  %9  %211 %35
+OpStore %169 %212
+%213 = OpLoad  %7  %158
+%214 = OpCompositeConstruct  %8  %30 %213
+%215 = OpCompositeConstruct  %9  %214 %35
+OpStore %171 %215
+%216 = OpLoad  %7  %158
+%217 = OpCompositeConstruct  %8  %216 %34
+%218 = OpCompositeConstruct  %9  %32 %217
+OpStore %173 %218
+%219 = OpLoad  %7  %158
+%220 = OpCompositeConstruct  %8  %33 %219
+%221 = OpCompositeConstruct  %9  %32 %220
+OpStore %175 %221
+%222 = OpLoad  %7  %158
+%223 = OpCompositeConstruct  %10  %222 %31
+OpStore %177 %223
+%224 = OpLoad  %7  %158
+%225 = OpCompositeConstruct  %10  %30 %224
+OpStore %179 %225
+%226 = OpLoad  %7  %158
+%227 = OpCompositeConstruct  %10  %226 %31
+OpStore %181 %227
+%228 = OpLoad  %7  %158
+%229 = OpCompositeConstruct  %10  %30 %228
+OpStore %183 %229
+%230 = OpLoad  %3  %155
+%231 = OpCompositeConstruct  %12  %230 %42
+OpStore %185 %231
+%232 = OpLoad  %3  %155
+%233 = OpCompositeConstruct  %12  %37 %232
+OpStore %187 %233
+%234 = OpLoad  %7  %158
+%235 = OpCompositeConstruct  %10  %234 %31
+OpStore %189 %235
+%236 = OpLoad  %7  %158
+%237 = OpCompositeConstruct  %10  %30 %236
+OpStore %191 %237
+%238 = OpLoad  %7  %158
+%239 = OpCompositeConstruct  %10  %238 %31
+OpStore %193 %239
+%240 = OpLoad  %7  %158
+%241 = OpCompositeConstruct  %10  %30 %240
+OpStore %195 %241
+%242 = OpLoad  %3  %155
+%243 = OpCompositeConstruct  %12  %242 %42
+OpStore %197 %243
+%244 = OpLoad  %3  %155
+%245 = OpCompositeConstruct  %12  %37 %244
+OpStore %199 %245
+%246 = OpLoad  %5  %152
+%247 = OpCompositeConstruct  %6  %246 %28
+OpStore %161 %247
+%248 = OpLoad  %5  %152
+%249 = OpCompositeConstruct  %6  %27 %248
+OpStore %163 %249
+%250 = OpLoad  %5  %152
+%251 = OpCompositeConstruct  %6  %250 %28
+OpStore %165 %251
+%252 = OpLoad  %5  %152
+%253 = OpCompositeConstruct  %6  %27 %252
+OpStore %167 %253
+%254 = OpLoad  %7  %158
+%255 = OpCompositeConstruct  %8  %254 %31
+%256 = OpCompositeConstruct  %9  %255 %35
+OpStore %169 %256
+%257 = OpLoad  %7  %158
+%258 = OpCompositeConstruct  %8  %30 %257
+%259 = OpCompositeConstruct  %9  %258 %35
+OpStore %171 %259
+%260 = OpLoad  %7  %158
+%261 = OpCompositeConstruct  %8  %260 %34
+%262 = OpCompositeConstruct  %9  %32 %261
+OpStore %173 %262
+%263 = OpLoad  %7  %158
+%264 = OpCompositeConstruct  %8  %33 %263
+%265 = OpCompositeConstruct  %9  %32 %264
+OpStore %175 %265
+%266 = OpLoad  %7  %158
+%267 = OpCompositeConstruct  %10  %266 %31
+OpStore %177 %267
+%268 = OpLoad  %7  %158
+%269 = OpCompositeConstruct  %10  %30 %268
+OpStore %179 %269
+%270 = OpLoad  %7  %158
+%271 = OpCompositeConstruct  %10  %270 %31
+OpStore %181 %271
+%272 = OpLoad  %7  %158
+%273 = OpCompositeConstruct  %10  %30 %272
+OpStore %183 %273
+%274 = OpLoad  %3  %155
+%275 = OpCompositeConstruct  %12  %274 %42
+OpStore %185 %275
+%276 = OpLoad  %3  %155
+%277 = OpCompositeConstruct  %12  %37 %276
+OpStore %187 %277
+%278 = OpLoad  %7  %158
+%279 = OpCompositeConstruct  %10  %278 %31
+OpStore %189 %279
+%280 = OpLoad  %7  %158
+%281 = OpCompositeConstruct  %10  %30 %280
+OpStore %191 %281
+%282 = OpLoad  %7  %158
+%283 = OpCompositeConstruct  %10  %282 %31
+OpStore %193 %283
+%284 = OpLoad  %7  %158
+%285 = OpCompositeConstruct  %10  %30 %284
+OpStore %195 %285
+%286 = OpLoad  %3  %155
+%287 = OpCompositeConstruct  %12  %286 %42
+OpStore %197 %287
+%288 = OpLoad  %3  %155
+%289 = OpCompositeConstruct  %12  %37 %288
+OpStore %199 %289
 OpReturn
 OpFunctionEnd
\ No newline at end of file
diff --git a/naga/tests/out/wgsl/abstract-types-var.wgsl b/naga/tests/out/wgsl/abstract-types-var.wgsl
index cd35501732..58d5f414f2 100644
--- a/naga/tests/out/wgsl/abstract-types-var.wgsl
+++ b/naga/tests/out/wgsl/abstract-types-var.wgsl
@@ -22,6 +22,9 @@ var<private> xafpaiai_1: array<i32, 2> = array<i32, 2>(1i, 2i);
 var<private> xafpaiaf_1: array<f32, 2> = array<f32, 2>(1f, 2f);
 var<private> xafpafai_1: array<f32, 2> = array<f32, 2>(1f, 2f);
 var<private> xafpafaf_1: array<f32, 2> = array<f32, 2>(1f, 2f);
+var<private> xavipai_1: array<vec3<i32>, 1> = array<vec3<i32>, 1>(vec3(1i));
+var<private> xavfpai_1: array<vec3<f32>, 1> = array<vec3<f32>, 1>(vec3(1f));
+var<private> xavfpaf_1: array<vec3<f32>, 1> = array<vec3<f32>, 1>(vec3(1f));
 var<private> ivispai: vec2<i32> = vec2(1i);
 var<private> ivfspaf: vec2<f32> = vec2(1f);
 var<private> ivis_ai: vec2<i32> = vec2(1i);
@@ -34,6 +37,9 @@ var<private> iaipaiai_1: array<i32, 2> = array<i32, 2>(1i, 2i);
 var<private> iafpafaf_1: array<f32, 2> = array<f32, 2>(1f, 2f);
 var<private> iafpaiaf_1: array<f32, 2> = array<f32, 2>(1f, 2f);
 var<private> iafpafai_1: array<f32, 2> = array<f32, 2>(1f, 2f);
+var<private> iavipai: array<vec3<i32>, 1> = array<vec3<i32>, 1>(vec3(1i));
+var<private> iavfpai: array<vec3<i32>, 1> = array<vec3<i32>, 1>(vec3(1i));
+var<private> iavfpaf: array<vec3<f32>, 1> = array<vec3<f32>, 1>(vec3(1f));
 
 fn all_constant_arguments() {
     var xvipaiai: vec2<i32> = vec2<i32>(42i, 43i);
@@ -69,6 +75,9 @@ fn all_constant_arguments() {
     var xafpaiaf: array<f32, 2> = array<f32, 2>(1f, 2f);
     var xafpafai: array<f32, 2> = array<f32, 2>(1f, 2f);
     var xafpafaf: array<f32, 2> = array<f32, 2>(1f, 2f);
+    var xavipai: array<vec3<i32>, 1> = array<vec3<i32>, 1>(vec3(1i));
+    var xavfpai: array<vec3<f32>, 1> = array<vec3<f32>, 1>(vec3(1f));
+    var xavfpaf: array<vec3<f32>, 1> = array<vec3<f32>, 1>(vec3(1f));
     var iaipaiai: array<i32, 2> = array<i32, 2>(1i, 2i);
     var iafpaiaf: array<f32, 2> = array<f32, 2>(1f, 2f);
     var iafpafai: array<f32, 2> = array<f32, 2>(1f, 2f);
@@ -107,6 +116,9 @@ fn all_constant_arguments() {
     xafpaiaf = array<f32, 2>(1f, 2f);
     xafpafai = array<f32, 2>(1f, 2f);
     xafpafaf = array<f32, 2>(1f, 2f);
+    xavipai = array<vec3<i32>, 1>(vec3(1i));
+    xavfpai = array<vec3<f32>, 1>(vec3(1f));
+    xavfpaf = array<vec3<f32>, 1>(vec3(1f));
     iaipaiai = array<i32, 2>(1i, 2i);
     iafpaiaf = array<f32, 2>(1f, 2f);
     iafpafai = array<f32, 2>(1f, 2f);