diff --git a/naga/src/back/spv/block.rs b/naga/src/back/spv/block.rs
index 7758d86c414..218f760ce4c 100644
--- a/naga/src/back/spv/block.rs
+++ b/naga/src/back/spv/block.rs
@@ -8,11 +8,14 @@ use arrayvec::ArrayVec;
 use spirv::Word;
 
 use super::{
-    index::BoundsCheckResult, selection::Selection, Block, BlockContext, Dimension, Error,
-    Instruction, LocalType, LookupType, NumericType, ResultMember, WrappedFunction, Writer,
-    WriterFlags,
+    helpers::map_storage_class, index::BoundsCheckResult, selection::Selection, Block,
+    BlockContext, Dimension, Error, IdGenerator, Instruction, LocalType, LookupType, NumericType,
+    ResultMember, WrappedFunction, Writer, WriterFlags,
+};
+use crate::{
+    arena::Handle, back::spv::helpers::is_uniform_matcx2_struct_member_access,
+    proc::index::GuardedIndex, Statement,
 };
-use crate::{arena::Handle, proc::index::GuardedIndex, Statement};
 
 fn get_dimension(type_inner: &crate::TypeInner) -> Dimension {
     match *type_inner {
@@ -34,6 +37,7 @@ fn get_dimension(type_inner: &crate::TypeInner) -> Dimension {
 /// the type of the given Naga IR [`Expression`] it's generating code for.
 ///
 /// [`Expression`]: crate::Expression
+#[derive(Copy, Clone)]
 enum AccessTypeAdjustment {
     /// No adjustment needed: the SPIR-V type should be the direct
     /// analog of the Naga IR expression type.
@@ -69,6 +73,17 @@ enum AccessTypeAdjustment {
     /// [`Access`]: crate::Expression::Access
     /// [`AccessIndex`]: crate::Expression::AccessIndex
     IntroducePointer(spirv::StorageClass),
+
+    /// The SPIR-V type should be an `OpPointer` to the std140 layout
+    /// compatible variant of the Naga IR expression's base type.
+    ///
+    /// This is used when accessing a type through an [`AddressSpace::Uniform`]
+    /// pointer in cases where the original type is incompatible with std140
+    /// layout requirements and we have therefore declared the uniform to be of
+    /// an alternative std140 compliant type.
+    ///
+    /// [`AddressSpace::Uniform`]: crate::AddressSpace::Uniform
+    UseStd140CompatType,
 }
 
 /// The results of emitting code for a left-hand-side expression.
@@ -409,6 +424,276 @@ impl BlockContext<'_> {
         block
     }
 
+    /// If `pointer` refers to an access chain that contains a dynamic indexing
+    /// of a two-row matrix in the [`Uniform`] address space, write code to
+    /// access the value returning the ID of the result. Else return None.
+    ///
+    /// Two-row matrices in the uniform address space will have been declared
+    /// using a alternative std140 layout compatible type, where each column is
+    /// a member of a containing struct. As a result, SPIR-V is unable to access
+    /// its columns with a non-constant index. To work around this limitation
+    /// this function will call [`Self::write_checked_load()`] to load the
+    /// matrix itself, which handles conversion from the std140 compatible type
+    /// to the real matrix type. It then calls a [`wrapper function`] to obtain
+    /// the correct column from the matrix, and possibly extracts a component
+    /// from the vector too.
+    ///
+    /// [`Uniform`]: crate::AddressSpace::Uniform
+    /// [`wrapper function`]: super::Writer::write_wrapped_matcx2_get_column
+    fn maybe_write_uniform_matcx2_dynamic_access(
+        &mut self,
+        pointer: Handle<crate::Expression>,
+        block: &mut Block,
+    ) -> Result<Option<Word>, Error> {
+        // If this access chain contains a dynamic matrix access, `pointer` is
+        // either a pointer to a vector (the column) or a scalar (a component
+        // within the column). In either case grab the pointer to the column,
+        // and remember the component index if there is one. If `pointer`
+        // points to any other type we're not interested.
+        let (column_pointer, component_index) = match self.fun_info[pointer]
+            .ty
+            .inner_with(&self.ir_module.types)
+            .pointer_base_type()
+        {
+            Some(resolution) => match *resolution.inner_with(&self.ir_module.types) {
+                crate::TypeInner::Scalar(_) => match self.ir_function.expressions[pointer] {
+                    crate::Expression::Access { base, index } => {
+                        (base, Some(GuardedIndex::Expression(index)))
+                    }
+                    crate::Expression::AccessIndex { base, index } => {
+                        (base, Some(GuardedIndex::Known(index)))
+                    }
+                    _ => return Ok(None),
+                },
+                crate::TypeInner::Vector { .. } => (pointer, None),
+                _ => return Ok(None),
+            },
+            None => return Ok(None),
+        };
+
+        // Ensure the column is accessed with a dynamic index (i.e.
+        // `Expression::Access`), and grab the pointer to the matrix.
+        let crate::Expression::Access {
+            base: matrix_pointer,
+            index: column_index,
+        } = self.ir_function.expressions[column_pointer]
+        else {
+            return Ok(None);
+        };
+
+        // Ensure the matrix pointer is in the uniform address space.
+        let crate::TypeInner::Pointer {
+            base: matrix_pointer_base_type,
+            space: crate::AddressSpace::Uniform,
+        } = *self.fun_info[matrix_pointer]
+            .ty
+            .inner_with(&self.ir_module.types)
+        else {
+            return Ok(None);
+        };
+
+        // Ensure the matrix pointer actually points to a Cx2 matrix.
+        let crate::TypeInner::Matrix {
+            columns,
+            rows: rows @ crate::VectorSize::Bi,
+            scalar,
+        } = self.ir_module.types[matrix_pointer_base_type].inner
+        else {
+            return Ok(None);
+        };
+
+        let matrix_type_id = self.get_numeric_type_id(NumericType::Matrix {
+            columns,
+            rows,
+            scalar,
+        });
+        let column_type_id = self.get_numeric_type_id(NumericType::Vector { size: rows, scalar });
+        let component_type_id = self.get_numeric_type_id(NumericType::Scalar(scalar));
+        let get_column_function_id = self.writer.wrapped_functions
+            [&WrappedFunction::MatCx2GetColumn {
+                r#type: matrix_pointer_base_type,
+            }];
+
+        let matrix_load_id = self.write_checked_load(
+            matrix_pointer,
+            block,
+            AccessTypeAdjustment::None,
+            matrix_type_id,
+        )?;
+
+        // Naga IR allows the index to be either an I32 or U32 but our wrapper
+        // function expects a U32 argument, so convert it if required.
+        let column_index_id = match *self.fun_info[column_index]
+            .ty
+            .inner_with(&self.ir_module.types)
+        {
+            crate::TypeInner::Scalar(crate::Scalar {
+                kind: crate::ScalarKind::Uint,
+                ..
+            }) => self.cached[column_index],
+            crate::TypeInner::Scalar(crate::Scalar {
+                kind: crate::ScalarKind::Sint,
+                ..
+            }) => {
+                let cast_id = self.gen_id();
+                let u32_type_id = self.writer.get_u32_type_id();
+                block.body.push(Instruction::unary(
+                    spirv::Op::Bitcast,
+                    u32_type_id,
+                    cast_id,
+                    self.cached[column_index],
+                ));
+                cast_id
+            }
+            _ => return Err(Error::Validation("Matrix access index must be u32 or i32")),
+        };
+        let column_id = self.gen_id();
+        block.body.push(Instruction::function_call(
+            column_type_id,
+            column_id,
+            get_column_function_id,
+            &[matrix_load_id, column_index_id],
+        ));
+        let result_id = match component_index {
+            Some(index) => self.write_vector_access(
+                component_type_id,
+                column_pointer,
+                Some(column_id),
+                index,
+                block,
+            )?,
+            None => column_id,
+        };
+
+        Ok(Some(result_id))
+    }
+
+    /// If `pointer` refers to two-row matrix that is a member of a struct in
+    /// the [`Uniform`] address space, write code to load the matrix returning
+    /// the ID of the result. Else return None.
+    ///
+    /// Two-row matrices that are struct members in the uniform address space
+    /// will have been decomposed such that the struct contains a separate
+    /// vector member for each column of the matrix. This function will load
+    /// each column separately from the containing struct, then composite them
+    /// into the real matrix type.
+    ///
+    /// [`Uniform`]: crate::AddressSpace::Uniform
+    fn maybe_write_load_uniform_matcx2_struct_member(
+        &mut self,
+        pointer: Handle<crate::Expression>,
+        block: &mut Block,
+    ) -> Result<Option<Word>, Error> {
+        // Check this is a uniform address space pointer to a two-row matrix.
+        let crate::TypeInner::Pointer {
+            base: matrix_type,
+            space: space @ crate::AddressSpace::Uniform,
+        } = *self.fun_info[pointer].ty.inner_with(&self.ir_module.types)
+        else {
+            return Ok(None);
+        };
+
+        let crate::TypeInner::Matrix {
+            columns,
+            rows: rows @ crate::VectorSize::Bi,
+            scalar,
+        } = self.ir_module.types[matrix_type].inner
+        else {
+            return Ok(None);
+        };
+
+        // Check this is a struct member. Note struct members can only be
+        // accessed with `AccessIndex`.
+        let crate::Expression::AccessIndex {
+            base: struct_pointer,
+            index: member_index,
+        } = self.ir_function.expressions[pointer]
+        else {
+            return Ok(None);
+        };
+
+        let crate::TypeInner::Pointer {
+            base: struct_type, ..
+        } = *self.fun_info[struct_pointer]
+            .ty
+            .inner_with(&self.ir_module.types)
+        else {
+            return Ok(None);
+        };
+
+        let crate::TypeInner::Struct { .. } = self.ir_module.types[struct_type].inner else {
+            return Ok(None);
+        };
+
+        let matrix_type_id = self.get_numeric_type_id(NumericType::Matrix {
+            columns,
+            rows,
+            scalar,
+        });
+        let column_type_id = self.get_numeric_type_id(NumericType::Vector { size: rows, scalar });
+        let column_pointer_type_id =
+            self.get_pointer_type_id(column_type_id, map_storage_class(space));
+        let column0_index = self.writer.std140_compat_uniform_types[&struct_type].member_indices
+            [member_index as usize];
+        let column_indices = (0..columns as u32)
+            .map(|c| self.get_index_constant(column0_index + c))
+            .collect::<ArrayVec<_, 4>>();
+
+        // Load each column from the struct, then composite into the real
+        // matrix type.
+        let load_mat_from_struct =
+            |struct_pointer_id: Word, id_gen: &mut IdGenerator, block: &mut Block| -> Word {
+                let mut column_ids: ArrayVec<Word, 4> = ArrayVec::new();
+                for index in &column_indices {
+                    let column_pointer_id = id_gen.next();
+                    block.body.push(Instruction::access_chain(
+                        column_pointer_type_id,
+                        column_pointer_id,
+                        struct_pointer_id,
+                        &[*index],
+                    ));
+                    let column_id = id_gen.next();
+                    block.body.push(Instruction::load(
+                        column_type_id,
+                        column_id,
+                        column_pointer_id,
+                        None,
+                    ));
+                    column_ids.push(column_id);
+                }
+                let result_id = id_gen.next();
+                block.body.push(Instruction::composite_construct(
+                    matrix_type_id,
+                    result_id,
+                    &column_ids,
+                ));
+                result_id
+            };
+
+        let result_id = match self.write_access_chain(
+            struct_pointer,
+            block,
+            AccessTypeAdjustment::UseStd140CompatType,
+        )? {
+            ExpressionPointer::Ready { pointer_id } => {
+                load_mat_from_struct(pointer_id, &mut self.writer.id_gen, block)
+            }
+            ExpressionPointer::Conditional { condition, access } => self
+                .write_conditional_indexed_load(
+                    matrix_type_id,
+                    condition,
+                    block,
+                    |id_gen, block| {
+                        let pointer_id = access.result_id.unwrap();
+                        block.body.push(access);
+                        load_mat_from_struct(pointer_id, id_gen, block)
+                    },
+                ),
+        };
+
+        Ok(Some(result_id))
+    }
+
     /// Cache an expression for a value.
     pub(super) fn cache_expression_value(
         &mut self,
@@ -501,9 +786,13 @@ impl BlockContext<'_> {
                         self.function.spilled_accesses.insert(expr_handle);
                         self.maybe_access_spilled_composite(expr_handle, block, result_type_id)?
                     }
-                    crate::TypeInner::Vector { .. } => {
-                        self.write_vector_access(expr_handle, base, index, block)?
-                    }
+                    crate::TypeInner::Vector { .. } => self.write_vector_access(
+                        result_type_id,
+                        base,
+                        None,
+                        GuardedIndex::Expression(index),
+                        block,
+                    )?,
                     crate::TypeInner::Array { .. } | crate::TypeInner::Matrix { .. } => {
                         // See if `index` is known at compile time.
                         match GuardedIndex::from_expression(
@@ -2092,6 +2381,20 @@ impl BlockContext<'_> {
                 AccessTypeAdjustment::IntroducePointer(class) => {
                     self.writer.get_resolution_pointer_id(resolution, class)
                 }
+                AccessTypeAdjustment::UseStd140CompatType => {
+                    match *resolution.inner_with(&self.ir_module.types) {
+                        crate::TypeInner::Pointer {
+                            base,
+                            space: space @ crate::AddressSpace::Uniform,
+                        } => self.writer.get_pointer_type_id(
+                            self.writer.std140_compat_uniform_types[&base].type_id,
+                            map_storage_class(space),
+                        ),
+                        _ => unreachable!(
+                            "`UseStd140CompatType` must only be used with uniform pointer types"
+                        ),
+                    }
+                }
             }
         };
 
@@ -2103,6 +2406,13 @@ impl BlockContext<'_> {
         // Is true if we are accessing into a binding array with a non-uniform index.
         let mut is_non_uniform_binding_array = false;
 
+        // The index value if the previously encountered expression was an
+        // `AccessIndex` of a matrix which has been decomposed into individual
+        // column vectors directly in the containing struct. The subsequent
+        // iteration will append the correct index to the list for accessing
+        // said column from the containing struct.
+        let mut prev_decomposed_matrix_index = None;
+
         self.temp_list.clear();
         let root_id = loop {
             // If `expr_handle` was spilled, then the temporary variable has exactly
@@ -2129,27 +2439,67 @@ impl BlockContext<'_> {
                     // Decide whether we're indexing a struct (bounds checks
                     // forbidden) or anything else (bounds checks required).
                     let mut base_ty = self.fun_info[base].ty.inner_with(&self.ir_module.types);
-                    if let crate::TypeInner::Pointer { base, .. } = *base_ty {
+                    let mut base_ty_handle = self.fun_info[base].ty.handle();
+                    let mut pointer_space = None;
+                    if let crate::TypeInner::Pointer { base, space } = *base_ty {
                         base_ty = &self.ir_module.types[base].inner;
+                        base_ty_handle = Some(base);
+                        pointer_space = Some(space);
                     }
-                    let index_id = if let crate::TypeInner::Struct { .. } = *base_ty {
-                        self.get_index_constant(index)
-                    } else {
-                        // `index` is constant, so this can't possibly require
-                        // setting `is_nonuniform_binding_array_access`.
-
-                        // Even though the index value is statically known, `base`
-                        // may be a runtime-sized array, so we still need to go
-                        // through the bounds check process.
-                        self.write_access_chain_index(
+                    match *base_ty {
+                        // When indexing a struct bounds checks are forbidden. If accessing the
+                        // struct through a uniform address space pointer, where the struct has
+                        // been declared with an alternative std140 compatible layout, we must use
+                        // the remapped member index. Additionally if the previous iteration was
+                        // accessing a column of a matrix member which has been decomposed directly
+                        // into the struct, we must ensure we access the correct column.
+                        crate::TypeInner::Struct { .. } => {
+                            let index = match base_ty_handle.and_then(|handle| {
+                                self.writer.std140_compat_uniform_types.get(&handle)
+                            }) {
+                                Some(std140_type_info)
+                                    if pointer_space == Some(crate::AddressSpace::Uniform) =>
+                                {
+                                    std140_type_info.member_indices[index as usize]
+                                        + prev_decomposed_matrix_index.take().unwrap_or(0)
+                                }
+                                _ => index,
+                            };
+                            let index_id = self.get_index_constant(index);
+                            self.temp_list.push(index_id);
+                        }
+                        // Bounds checks are not required when indexing a matrix. If indexing a
+                        // two-row matrix contained within a struct through a uniform address space
+                        // pointer then the matrix' columns will have been decomposed directly into
+                        // the containing struct. We skip adding an index to the list on this
+                        // iteration and instead adjust the index on the next iteration when
+                        // accessing the struct member.
+                        _ if is_uniform_matcx2_struct_member_access(
+                            self.ir_function,
+                            self.fun_info,
+                            self.ir_module,
                             base,
-                            GuardedIndex::Known(index),
-                            &mut accumulated_checks,
-                            block,
-                        )?
-                    };
-
-                    self.temp_list.push(index_id);
+                        ) =>
+                        {
+                            assert!(prev_decomposed_matrix_index.is_none());
+                            prev_decomposed_matrix_index = Some(index);
+                        }
+                        _ => {
+                            // `index` is constant, so this can't possibly require
+                            // setting `is_nonuniform_binding_array_access`.
+
+                            // Even though the index value is statically known, `base`
+                            // may be a runtime-sized array, so we still need to go
+                            // through the bounds check process.
+                            let index_id = self.write_access_chain_index(
+                                base,
+                                GuardedIndex::Known(index),
+                                &mut accumulated_checks,
+                                block,
+                            )?;
+                            self.temp_list.push(index_id);
+                        }
+                    }
                     base
                 }
                 crate::Expression::GlobalVariable(handle) => {
@@ -2310,57 +2660,119 @@ impl BlockContext<'_> {
         access_type_adjustment: AccessTypeAdjustment,
         result_type_id: Word,
     ) -> Result<Word, Error> {
-        match self.write_access_chain(pointer, block, access_type_adjustment)? {
-            ExpressionPointer::Ready { pointer_id } => {
-                let id = self.gen_id();
-                let atomic_space =
-                    match *self.fun_info[pointer].ty.inner_with(&self.ir_module.types) {
-                        crate::TypeInner::Pointer { base, space } => {
-                            match self.ir_module.types[base].inner {
-                                crate::TypeInner::Atomic { .. } => Some(space),
-                                _ => None,
-                            }
-                        }
-                        _ => None,
-                    };
-                let instruction = if let Some(space) = atomic_space {
-                    let (semantics, scope) = space.to_spirv_semantics_and_scope();
-                    let scope_constant_id = self.get_scope_constant(scope as u32);
-                    let semantics_id = self.get_index_constant(semantics.bits());
-                    Instruction::atomic_load(
-                        result_type_id,
-                        id,
-                        pointer_id,
-                        scope_constant_id,
-                        semantics_id,
-                    )
-                } else {
-                    Instruction::load(result_type_id, id, pointer_id, None)
-                };
-                block.body.push(instruction);
-                Ok(id)
+        if let Some(result_id) = self.maybe_write_uniform_matcx2_dynamic_access(pointer, block)? {
+            Ok(result_id)
+        } else if let Some(result_id) =
+            self.maybe_write_load_uniform_matcx2_struct_member(pointer, block)?
+        {
+            Ok(result_id)
+        } else {
+            // If `pointer` refers to a uniform address space pointer to a type
+            // which was declared using a std140 compatible type variant (i.e.
+            // is a two-row matrix, or a struct or array containing such a
+            // matrix) we must ensure the access chain and the type of the load
+            // instruction use the std140 compatible type variant.
+            struct WrappedLoad {
+                access_type_adjustment: AccessTypeAdjustment,
+                r#type: Handle<crate::Type>,
             }
-            ExpressionPointer::Conditional { condition, access } => {
-                //TODO: support atomics?
-                let value = self.write_conditional_indexed_load(
-                    result_type_id,
-                    condition,
-                    block,
-                    move |id_gen, block| {
-                        // The in-bounds path. Perform the access and the load.
-                        let pointer_id = access.result_id.unwrap();
-                        let value_id = id_gen.next();
-                        block.body.push(access);
-                        block.body.push(Instruction::load(
+            let mut wrapped_load = None;
+            if let crate::TypeInner::Pointer {
+                base: pointer_base_type,
+                space: crate::AddressSpace::Uniform,
+            } = *self.fun_info[pointer].ty.inner_with(&self.ir_module.types)
+            {
+                if self
+                    .writer
+                    .std140_compat_uniform_types
+                    .contains_key(&pointer_base_type)
+                {
+                    wrapped_load = Some(WrappedLoad {
+                        access_type_adjustment: AccessTypeAdjustment::UseStd140CompatType,
+                        r#type: pointer_base_type,
+                    });
+                };
+            };
+
+            let (load_type_id, access_type_adjustment) = match wrapped_load {
+                Some(ref wrapped_load) => (
+                    self.writer.std140_compat_uniform_types[&wrapped_load.r#type].type_id,
+                    wrapped_load.access_type_adjustment,
+                ),
+                None => (result_type_id, access_type_adjustment),
+            };
+
+            let load_id = match self.write_access_chain(pointer, block, access_type_adjustment)? {
+                ExpressionPointer::Ready { pointer_id } => {
+                    let id = self.gen_id();
+                    let atomic_space =
+                        match *self.fun_info[pointer].ty.inner_with(&self.ir_module.types) {
+                            crate::TypeInner::Pointer { base, space } => {
+                                match self.ir_module.types[base].inner {
+                                    crate::TypeInner::Atomic { .. } => Some(space),
+                                    _ => None,
+                                }
+                            }
+                            _ => None,
+                        };
+                    let instruction = if let Some(space) = atomic_space {
+                        let (semantics, scope) = space.to_spirv_semantics_and_scope();
+                        let scope_constant_id = self.get_scope_constant(scope as u32);
+                        let semantics_id = self.get_index_constant(semantics.bits());
+                        Instruction::atomic_load(
                             result_type_id,
-                            value_id,
+                            id,
                             pointer_id,
-                            None,
-                        ));
-                        value_id
-                    },
-                );
-                Ok(value)
+                            scope_constant_id,
+                            semantics_id,
+                        )
+                    } else {
+                        Instruction::load(load_type_id, id, pointer_id, None)
+                    };
+                    block.body.push(instruction);
+                    id
+                }
+                ExpressionPointer::Conditional { condition, access } => {
+                    //TODO: support atomics?
+                    self.write_conditional_indexed_load(
+                        load_type_id,
+                        condition,
+                        block,
+                        move |id_gen, block| {
+                            // The in-bounds path. Perform the access and the load.
+                            let pointer_id = access.result_id.unwrap();
+                            let value_id = id_gen.next();
+                            block.body.push(access);
+                            block.body.push(Instruction::load(
+                                load_type_id,
+                                value_id,
+                                pointer_id,
+                                None,
+                            ));
+                            value_id
+                        },
+                    )
+                }
+            };
+
+            match wrapped_load {
+                Some(ref wrapped_load) => {
+                    // If we loaded a std140 compat type then we must call the
+                    // function to convert the loaded value to the regular type.
+                    let result_id = self.gen_id();
+                    let function_id = self.writer.wrapped_functions
+                        [&WrappedFunction::ConvertFromStd140CompatType {
+                            r#type: wrapped_load.r#type,
+                        }];
+                    block.body.push(Instruction::function_call(
+                        result_type_id,
+                        result_id,
+                        function_id,
+                        &[load_id],
+                    ));
+                    Ok(result_id)
+                }
+                None => Ok(load_id),
             }
         }
     }
diff --git a/naga/src/back/spv/helpers.rs b/naga/src/back/spv/helpers.rs
index 84e130efaa3..acc52dfa566 100644
--- a/naga/src/back/spv/helpers.rs
+++ b/naga/src/back/spv/helpers.rs
@@ -122,6 +122,44 @@ pub fn global_needs_wrapper(ir_module: &crate::Module, var: &crate::GlobalVariab
     }
 }
 
+/// Returns true if `pointer` refers to two-row matrix which is a member of a
+/// struct in the [`crate::AddressSpace::Uniform`] address space.
+pub fn is_uniform_matcx2_struct_member_access(
+    ir_function: &crate::Function,
+    fun_info: &crate::valid::FunctionInfo,
+    ir_module: &crate::Module,
+    pointer: Handle<crate::Expression>,
+) -> bool {
+    if let crate::TypeInner::Pointer {
+        base: pointer_base_type,
+        space: crate::AddressSpace::Uniform,
+    } = *fun_info[pointer].ty.inner_with(&ir_module.types)
+    {
+        if let crate::TypeInner::Matrix {
+            rows: crate::VectorSize::Bi,
+            ..
+        } = ir_module.types[pointer_base_type].inner
+        {
+            if let crate::Expression::AccessIndex {
+                base: parent_pointer,
+                ..
+            } = ir_function.expressions[pointer]
+            {
+                if let crate::TypeInner::Pointer {
+                    base: parent_type, ..
+                } = *fun_info[parent_pointer].ty.inner_with(&ir_module.types)
+                {
+                    if let crate::TypeInner::Struct { .. } = ir_module.types[parent_type].inner {
+                        return true;
+                    }
+                }
+            }
+        }
+    }
+
+    false
+}
+
 ///HACK: this is taken from std unstable, remove it when std's floor_char_boundary is stable
 trait U8Internal {
     fn is_utf8_char_boundary(&self) -> bool;
diff --git a/naga/src/back/spv/index.rs b/naga/src/back/spv/index.rs
index 3a15ee88060..3ea448bbfbf 100644
--- a/naga/src/back/spv/index.rs
+++ b/naga/src/back/spv/index.rs
@@ -536,17 +536,18 @@ impl BlockContext<'_> {
     /// Emit code to subscript a vector by value with a computed index.
     ///
     /// Return the id of the element value.
+    ///
+    /// If `base_id_override` is provided, it is used as the vector expression
+    /// to be subscripted into, rather than the cached value of `base`.
     pub(super) fn write_vector_access(
         &mut self,
-        expr_handle: Handle<crate::Expression>,
+        result_type_id: Word,
         base: Handle<crate::Expression>,
-        index: Handle<crate::Expression>,
+        base_id_override: Option<Word>,
+        index: GuardedIndex,
         block: &mut Block,
     ) -> Result<Word, Error> {
-        let result_type_id = self.get_expression_type_id(&self.fun_info[expr_handle].ty);
-
-        let base_id = self.cached[base];
-        let index = GuardedIndex::Expression(index);
+        let base_id = base_id_override.unwrap_or_else(|| self.cached[base]);
 
         let result_id = match self.write_bounds_check(base, index, block)? {
             BoundsCheckResult::KnownInBounds(known_index) => {
diff --git a/naga/src/back/spv/mod.rs b/naga/src/back/spv/mod.rs
index 4690dc71951..6e3783242cc 100644
--- a/naga/src/back/spv/mod.rs
+++ b/naga/src/back/spv/mod.rs
@@ -1,7 +1,91 @@
 /*!
 Backend for [SPIR-V][spv] (Standard Portable Intermediate Representation).
 
+# Layout of values in `uniform` buffers
+
+WGSL's ["Internal Layout of Values"][ilov] rules specify how each WGSL type
+should be stored in `uniform` and `storage` buffers, and Naga IR adheres to
+these rules. The SPIR-V we generate must access values in that form, even when
+it is not what Vulkan would use normally. Fortunately the rules for `storage`
+buffers match Vulkan's, but some adjustments must be made when emitting SPIR-V
+for `uniform` buffers.
+
+## Padding in two-row matrices
+
+In Vulkan's ["extended layout"][extended-layout] (also known as std140) used
+for `uniform` buffers, matrices are defined in terms of arrays of their vector
+type, and arrays are defined to have an alignment equal to the alignment of
+their element type rounded up to a multiple of 16. This means that each column
+of the vector has a minimum alignment of 16. WGSL, and consequently Naga IR, on
+the other hand defines each column to have an alignment equal to the alignment
+of the vector type, without being rounded up to 16.
+
+To compensate for this, for any `struct` used as a `uniform` buffer which
+contains a two-row matrix, we declare an additional "std140 compatible" type
+in which each column of the matrix has been decomposed into the containing
+struct. For example, the following WGSL struct type:
+
+```ignore
+struct Baz {
+    m: mat3x2<f32>,
+}
+```
+
+is rendered as the SPIR-V struct type:
+
+```ignore
+OpTypeStruct %v2float %v2float %v2float
+```
+
+This has the effect that struct indices in Naga IR for such types do not
+correspond to the struct indices used in SPIR-V. A mapping of struct indices
+for these types is maintained in [`Std140CompatTypeInfo`].
+
+Additionally, any two-row matrices that are declared directly as uniform
+buffers without being wrapped in a struct are declared as a struct containing a
+vector member for each column. Any array of a two-row matrix in a uniform
+buffer is declared as an array of a struct containing a vector member for each
+column. Any struct or array within a uniform buffer which contains a member or
+whose base type requires a std140 compatible type declaration, itself requires a
+std140 compatible type declaration.
+
+Whenever a value of such a type is [`loaded`] we insert code to convert the
+loaded value from the std140 compatible type to the regular type. This occurs
+in `BlockContext::write_checked_load`, making use of the wrapper function
+defined by `Writer::write_wrapped_convert_from_std140_compat_type`. For matrices
+that have been decomposed as separate columns in the containing struct, we load
+each column separately then composite the matrix type in
+`BlockContext::maybe_write_load_uniform_matcx2_struct_member`.
+
+Whenever a column of a matrix that has been decomposed into its containing
+struct is [`accessed`] with a constant index we adjust the emitted access chain
+to access from the containing struct instead, in `BlockContext::write_access_chain`.
+
+Whenever a column of a uniform buffer two-row matrix is [`dynamically accessed`]
+we must first load the matrix type, converting it from its std140 compatible
+type as described above, then access the column using the wrapper function
+defined by `Writer::write_wrapped_matcx2_get_column`. This is handled by
+`BlockContext::maybe_write_uniform_matcx2_dynamic_access`.
+
+Note that this approach differs somewhat from the equivalent code in the HLSL
+backend. For HLSL all structs containing two-row matrices (or arrays of such)
+have their declarations modified, not just those used as uniform buffers.
+Two-row matrices and arrays of such only use modified type declarations when
+used as uniform buffers, or additionally when used as struct member in any
+context. This avoids the need to convert struct values when loading from uniform
+buffers, but when loading arrays and matrices from uniform buffers or from any
+struct the conversion is still required. In contrast, the approach used here
+always requires converting *any* affected type when loading from a uniform
+buffer, but consistently *only* when loading from a uniform buffer. As a result
+this also means we only have to handle loads and not stores, as uniform buffers
+are read-only.
+
 [spv]: https://www.khronos.org/registry/SPIR-V/
+[ilov]: https://gpuweb.github.io/gpuweb/wgsl/#internal-value-layout
+[extended-layout]: https://docs.vulkan.org/spec/latest/chapters/interfaces.html#interfaces-resources-layout
+[`loaded`]: crate::Expression::Load
+[`accessed`]: crate::Expression::AccessIndex
+[`dynamically accessed`]: crate::Expression::Access
 */
 
 mod block;
@@ -462,6 +546,12 @@ enum WrappedFunction {
         left_type_id: Word,
         right_type_id: Word,
     },
+    ConvertFromStd140CompatType {
+        r#type: Handle<crate::Type>,
+    },
+    MatCx2GetColumn {
+        r#type: Handle<crate::Type>,
+    },
 }
 
 /// A map from evaluated [`Expression`](crate::Expression)s to their SPIR-V ids.
@@ -722,6 +812,20 @@ impl BlockContext<'_> {
     }
 }
 
+/// Information about a type for which we have declared a std140 layout
+/// compatible variant, because the type is used in a uniform but does not
+/// adhere to std140 requirements. The uniform will be declared using the
+/// type `type_id`, and the result of any `Load` will be immediately converted
+/// to the base type. This is used for matrices with 2 rows, as well as any
+/// arrays or structs containing such matrices.
+pub struct Std140CompatTypeInfo {
+    /// ID of the std140 compatible type declaration.
+    type_id: Word,
+    /// For structs, a mapping of Naga IR struct member indices to the indices
+    /// used in the generated SPIR-V. For non-struct types this will be empty.
+    member_indices: Vec<u32>,
+}
+
 pub struct Writer {
     physical_layout: PhysicalLayout,
     logical_layout: LogicalLayout,
@@ -761,6 +865,7 @@ pub struct Writer {
     constant_ids: HandleVec<crate::Expression, Word>,
     cached_constants: crate::FastHashMap<CachedConstant, Word>,
     global_variables: HandleVec<crate::GlobalVariable, GlobalVariable>,
+    std140_compat_uniform_types: crate::FastHashMap<Handle<crate::Type>, Std140CompatTypeInfo>,
     fake_missing_bindings: bool,
     binding_map: BindingMap,
 
diff --git a/naga/src/back/spv/writer.rs b/naga/src/back/spv/writer.rs
index c86a53c6ef8..0f654318020 100644
--- a/naga/src/back/spv/writer.rs
+++ b/naga/src/back/spv/writer.rs
@@ -1,5 +1,6 @@
-use alloc::{string::String, vec, vec::Vec};
+use alloc::{format, string::String, vec, vec::Vec};
 
+use arrayvec::ArrayVec;
 use hashbrown::hash_map::Entry;
 use spirv::Word;
 
@@ -13,7 +14,11 @@ use super::{
 };
 use crate::{
     arena::{Handle, HandleVec, UniqueArena},
-    back::spv::{BindingInfo, WrappedFunction},
+    back::spv::{
+        helpers::is_uniform_matcx2_struct_member_access, BindingInfo, Std140CompatTypeInfo,
+        WrappedFunction,
+    },
+    common::ForDebugWithTypes as _,
     proc::{Alignment, TypeResolution},
     valid::{FunctionInfo, ModuleInfo},
 };
@@ -86,6 +91,7 @@ impl Writer {
             constant_ids: HandleVec::new(),
             cached_constants: crate::FastHashMap::default(),
             global_variables: HandleVec::new(),
+            std140_compat_uniform_types: crate::FastHashMap::default(),
             fake_missing_bindings: options.fake_missing_bindings,
             binding_map: options.binding_map.clone(),
             saved_cached: CachedExpressions::default(),
@@ -171,6 +177,7 @@ impl Writer {
             constant_ids: take(&mut self.constant_ids).recycle(),
             cached_constants: take(&mut self.cached_constants).recycle(),
             global_variables: take(&mut self.global_variables).recycle(),
+            std140_compat_uniform_types: take(&mut self.std140_compat_uniform_types).recycle(),
             saved_cached: take(&mut self.saved_cached).recycle(),
             temp_list: take(&mut self.temp_list).recycle(),
             ray_get_candidate_intersection_function: None,
@@ -526,6 +533,52 @@ impl Writer {
                         }
                     }
                 }
+                crate::Expression::Load { pointer } => {
+                    if let crate::TypeInner::Pointer {
+                        base: pointer_type,
+                        space: crate::AddressSpace::Uniform,
+                    } = *info[pointer].ty.inner_with(&ir_module.types)
+                    {
+                        if self.std140_compat_uniform_types.contains_key(&pointer_type) {
+                            // Loading a std140 compat type requires the wrapper function
+                            // to convert to the regular type.
+                            self.write_wrapped_convert_from_std140_compat_type(
+                                ir_module,
+                                pointer_type,
+                            )?;
+                        }
+                    }
+                }
+                crate::Expression::Access { base, .. } => {
+                    if let crate::TypeInner::Pointer {
+                        base: base_type,
+                        space: crate::AddressSpace::Uniform,
+                    } = *info[base].ty.inner_with(&ir_module.types)
+                    {
+                        // Dynamic accesses of a two-row matrix's columns require a
+                        // wrapper function.
+                        if let crate::TypeInner::Matrix {
+                            rows: crate::VectorSize::Bi,
+                            ..
+                        } = ir_module.types[base_type].inner
+                        {
+                            self.write_wrapped_matcx2_get_column(ir_module, base_type)?;
+                            // If the matrix is *not* directly a member of a struct, then
+                            // we additionally require a wrapper function to convert from
+                            // the std140 compat type to the regular type.
+                            if !is_uniform_matcx2_struct_member_access(
+                                ir_function,
+                                info,
+                                ir_module,
+                                base,
+                            ) {
+                                self.write_wrapped_convert_from_std140_compat_type(
+                                    ir_module, base_type,
+                                )?;
+                            }
+                        }
+                    }
+                }
                 _ => {}
             }
         }
@@ -726,6 +779,379 @@ impl Writer {
         Ok(())
     }
 
+    /// Writes a wrapper function to convert from a std140 compat type to its
+    /// corresponding regular type.
+    ///
+    /// See [`Self::write_std140_compat_type_declaration`] for more details.
+    fn write_wrapped_convert_from_std140_compat_type(
+        &mut self,
+        ir_module: &crate::Module,
+        r#type: Handle<crate::Type>,
+    ) -> Result<(), Error> {
+        // Check if we've already emitted this function.
+        let wrapped = WrappedFunction::ConvertFromStd140CompatType { r#type };
+        let function_id = match self.wrapped_functions.entry(wrapped) {
+            Entry::Occupied(_) => return Ok(()),
+            Entry::Vacant(e) => *e.insert(self.id_gen.next()),
+        };
+        if self.flags.contains(WriterFlags::DEBUG) {
+            self.debugs.push(Instruction::name(
+                function_id,
+                &format!("{:?}_from_std140", r#type.for_debug(&ir_module.types)),
+            ));
+        }
+        let param_type_id = self.std140_compat_uniform_types[&r#type].type_id;
+        let return_type_id = self.get_handle_type_id(r#type);
+
+        let mut function = Function::default();
+        let function_type_id = self.get_function_type(LookupFunctionType {
+            parameter_type_ids: vec![param_type_id],
+            return_type_id,
+        });
+        function.signature = Some(Instruction::function(
+            return_type_id,
+            function_id,
+            spirv::FunctionControl::empty(),
+            function_type_id,
+        ));
+        let param_id = self.id_gen.next();
+        function.parameters.push(FunctionArgument {
+            instruction: Instruction::function_parameter(param_type_id, param_id),
+            handle_id: 0,
+        });
+
+        let label_id = self.id_gen.next();
+        let mut block = Block::new(label_id);
+
+        let result_id = match ir_module.types[r#type].inner {
+            // Param is struct containing a vector member for each of the
+            // matrix's columns. Extract each column from the struct then
+            // composite into a matrix.
+            crate::TypeInner::Matrix {
+                columns,
+                rows: rows @ crate::VectorSize::Bi,
+                scalar,
+            } => {
+                let column_type_id =
+                    self.get_numeric_type_id(NumericType::Vector { size: rows, scalar });
+
+                let mut column_ids: ArrayVec<Word, 4> = ArrayVec::new();
+                for column in 0..columns as u32 {
+                    let column_id = self.id_gen.next();
+                    block.body.push(Instruction::composite_extract(
+                        column_type_id,
+                        column_id,
+                        param_id,
+                        &[column],
+                    ));
+                    column_ids.push(column_id);
+                }
+                let result_id = self.id_gen.next();
+                block.body.push(Instruction::composite_construct(
+                    return_type_id,
+                    result_id,
+                    &column_ids,
+                ));
+                result_id
+            }
+            // Param is an array where the base type is the std140 compatible
+            // type corresponding to `base`. Iterate through each element and
+            // call its conversion function, then composite into a new array.
+            crate::TypeInner::Array { base, size, .. } => {
+                // Ensure the conversion function for the array's base type is
+                // declared.
+                self.write_wrapped_convert_from_std140_compat_type(ir_module, base)?;
+
+                let element_type_id = self.get_handle_type_id(base);
+                let std140_element_type_id = self.std140_compat_uniform_types[&base].type_id;
+                let element_conversion_function_id = self.wrapped_functions
+                    [&WrappedFunction::ConvertFromStd140CompatType { r#type: base }];
+                let mut element_ids = Vec::new();
+                let size = match size.resolve(ir_module.to_ctx())? {
+                    crate::proc::IndexableLength::Known(size) => size,
+                    crate::proc::IndexableLength::Dynamic => {
+                        return Err(Error::Validation(
+                            "Uniform buffers cannot contain dynamic arrays",
+                        ))
+                    }
+                };
+                for i in 0..size {
+                    let std140_element_id = self.id_gen.next();
+                    block.body.push(Instruction::composite_extract(
+                        std140_element_type_id,
+                        std140_element_id,
+                        param_id,
+                        &[i],
+                    ));
+                    let element_id = self.id_gen.next();
+                    block.body.push(Instruction::function_call(
+                        element_type_id,
+                        element_id,
+                        element_conversion_function_id,
+                        &[std140_element_id],
+                    ));
+                    element_ids.push(element_id);
+                }
+                let result_id = self.id_gen.next();
+                block.body.push(Instruction::composite_construct(
+                    return_type_id,
+                    result_id,
+                    &element_ids,
+                ));
+                result_id
+            }
+            // Param is a struct where each two-row matrix member has been
+            // decomposed in to separate vector members for each column.
+            // Other members use their std140 compatible type if one exists, or
+            // else their regular type. Iterate through each member, converting
+            // or composing any matrices if required, then finally compose into
+            // the struct.
+            crate::TypeInner::Struct { ref members, .. } => {
+                let mut member_ids = Vec::new();
+                let mut next_index = 0;
+                for member in members {
+                    let member_id = self.id_gen.next();
+                    let member_type_id = self.get_handle_type_id(member.ty);
+                    match ir_module.types[member.ty].inner {
+                        crate::TypeInner::Matrix {
+                            columns,
+                            rows: rows @ crate::VectorSize::Bi,
+                            scalar,
+                        } => {
+                            let mut column_ids: ArrayVec<Word, 4> = ArrayVec::new();
+                            let column_type_id = self
+                                .get_numeric_type_id(NumericType::Vector { size: rows, scalar });
+                            for _ in 0..columns as u32 {
+                                let column_id = self.id_gen.next();
+                                block.body.push(Instruction::composite_extract(
+                                    column_type_id,
+                                    column_id,
+                                    param_id,
+                                    &[next_index],
+                                ));
+                                column_ids.push(column_id);
+                                next_index += 1;
+                            }
+                            block.body.push(Instruction::composite_construct(
+                                member_type_id,
+                                member_id,
+                                &column_ids,
+                            ));
+                        }
+                        _ => {
+                            // Ensure the conversion function for the member's
+                            // type is declared.
+                            self.write_wrapped_convert_from_std140_compat_type(
+                                ir_module, member.ty,
+                            )?;
+                            match self.std140_compat_uniform_types.get(&member.ty) {
+                                Some(std140_type_info) => {
+                                    let std140_member_id = self.id_gen.next();
+                                    block.body.push(Instruction::composite_extract(
+                                        std140_type_info.type_id,
+                                        std140_member_id,
+                                        param_id,
+                                        &[next_index],
+                                    ));
+                                    let function_id = self.wrapped_functions
+                                        [&WrappedFunction::ConvertFromStd140CompatType {
+                                            r#type: member.ty,
+                                        }];
+                                    block.body.push(Instruction::function_call(
+                                        member_type_id,
+                                        member_id,
+                                        function_id,
+                                        &[std140_member_id],
+                                    ));
+                                    next_index += 1;
+                                }
+                                None => {
+                                    let member_id = self.id_gen.next();
+                                    block.body.push(Instruction::composite_extract(
+                                        member_type_id,
+                                        member_id,
+                                        param_id,
+                                        &[next_index],
+                                    ));
+                                    next_index += 1;
+                                }
+                            }
+                        }
+                    }
+                    member_ids.push(member_id);
+                }
+                let result_id = self.id_gen.next();
+                block.body.push(Instruction::composite_construct(
+                    return_type_id,
+                    result_id,
+                    &member_ids,
+                ));
+                result_id
+            }
+            _ => unreachable!(),
+        };
+
+        function.consume(block, Instruction::return_value(result_id));
+        function.to_words(&mut self.logical_layout.function_definitions);
+        Ok(())
+    }
+
+    /// Writes a wrapper function to get an `OpTypeVector` column from an
+    /// `OpTypeMatrix` with a dynamic index.
+    ///
+    /// This is used when accessing a column of a [`TypeInner::Matrix`] through
+    /// a [`Uniform`] address space pointer. In such cases, the matrix will have
+    /// been declared in SPIR-V using an alternative type where each column is a
+    /// member of a containing struct. SPIR-V is unable to dynamically access
+    /// struct members, so instead we load the matrix then call this function to
+    /// access a column from the loaded value.
+    ///
+    /// [`TypeInner::Matrix`]: crate::TypeInner::Matrix
+    /// [`Uniform`]: crate::AddressSpace::Uniform
+    fn write_wrapped_matcx2_get_column(
+        &mut self,
+        ir_module: &crate::Module,
+        r#type: Handle<crate::Type>,
+    ) -> Result<(), Error> {
+        let wrapped = WrappedFunction::MatCx2GetColumn { r#type };
+        let function_id = match self.wrapped_functions.entry(wrapped) {
+            Entry::Occupied(_) => return Ok(()),
+            Entry::Vacant(e) => *e.insert(self.id_gen.next()),
+        };
+        if self.flags.contains(WriterFlags::DEBUG) {
+            self.debugs.push(Instruction::name(
+                function_id,
+                &format!("{:?}_get_column", r#type.for_debug(&ir_module.types)),
+            ));
+        }
+
+        let crate::TypeInner::Matrix {
+            columns,
+            rows: rows @ crate::VectorSize::Bi,
+            scalar,
+        } = ir_module.types[r#type].inner
+        else {
+            unreachable!();
+        };
+
+        let mut function = Function::default();
+        let matrix_type_id = self.get_handle_type_id(r#type);
+        let column_index_type_id = self.get_u32_type_id();
+        let column_type_id = self.get_numeric_type_id(NumericType::Vector { size: rows, scalar });
+        let matrix_param_id = self.id_gen.next();
+        let column_index_param_id = self.id_gen.next();
+        function.parameters.push(FunctionArgument {
+            instruction: Instruction::function_parameter(matrix_type_id, matrix_param_id),
+            handle_id: 0,
+        });
+        function.parameters.push(FunctionArgument {
+            instruction: Instruction::function_parameter(
+                column_index_type_id,
+                column_index_param_id,
+            ),
+            handle_id: 0,
+        });
+        let function_type_id = self.get_function_type(LookupFunctionType {
+            parameter_type_ids: vec![matrix_type_id, column_index_type_id],
+            return_type_id: column_type_id,
+        });
+        function.signature = Some(Instruction::function(
+            column_type_id,
+            function_id,
+            spirv::FunctionControl::empty(),
+            function_type_id,
+        ));
+
+        let label_id = self.id_gen.next();
+        let mut block = Block::new(label_id);
+
+        // Create a switch case for each column in the matrix, where each case
+        // extracts its column from the matrix. Finally we use OpPhi to return
+        // the correct column.
+        let merge_id = self.id_gen.next();
+        block.body.push(Instruction::selection_merge(
+            merge_id,
+            spirv::SelectionControl::NONE,
+        ));
+        let cases = (0..columns as u32)
+            .map(|i| super::instructions::Case {
+                value: i,
+                label_id: self.id_gen.next(),
+            })
+            .collect::<ArrayVec<_, 4>>();
+
+        // Which label we branch to in the default (column index out-of-bounds)
+        // case depends on our bounds check policy.
+        let default_id = match self.bounds_check_policies.index {
+            // For `Restrict`, treat the same as the final column.
+            crate::proc::BoundsCheckPolicy::Restrict => cases.last().unwrap().label_id,
+            // For `ReadZeroSkipWrite`, branch directly to the merge block. This
+            // will be handled in the `OpPhi` below to produce a zero value.
+            crate::proc::BoundsCheckPolicy::ReadZeroSkipWrite => merge_id,
+            // For `Unchecked` we create a new block containing an
+            // `OpUnreachable`.
+            crate::proc::BoundsCheckPolicy::Unchecked => self.id_gen.next(),
+        };
+        function.consume(
+            block,
+            Instruction::switch(column_index_param_id, default_id, &cases),
+        );
+
+        // Emit a block for each case, and produce a list of variable and parent
+        // block IDs that will be used in an `OpPhi` below to select the right
+        // value.
+        let mut var_parent_pairs = cases
+            .into_iter()
+            .map(|case| {
+                let mut block = Block::new(case.label_id);
+                let column_id = self.id_gen.next();
+                block.body.push(Instruction::composite_extract(
+                    column_type_id,
+                    column_id,
+                    matrix_param_id,
+                    &[case.value],
+                ));
+                function.consume(block, Instruction::branch(merge_id));
+                (column_id, case.label_id)
+            })
+            // Need capacity for up to 4 columns plus possibly a default case.
+            .collect::<ArrayVec<_, 5>>();
+
+        // Emit a block or append the variable and parent `OpPhi` pair for the
+        // column index out-of-bounds case, if required.
+        match self.bounds_check_policies.index {
+            // Don't need to do anything for `Restrict` as we have branched from
+            // the final column case's block.
+            crate::proc::BoundsCheckPolicy::Restrict => {}
+            // For `ReadZeroSkipWrite` we have branched directly from the block
+            // containing the `OpSwitch`. The `OpPhi` should produce a zero
+            // value.
+            crate::proc::BoundsCheckPolicy::ReadZeroSkipWrite => {
+                var_parent_pairs.push((self.get_constant_null(column_type_id), label_id));
+            }
+            // For `Unchecked` create a new block containing `OpUnreachable`.
+            // This does not need to be handled by the `OpPhi`.
+            crate::proc::BoundsCheckPolicy::Unchecked => {
+                function.consume(
+                    Block::new(default_id),
+                    Instruction::new(spirv::Op::Unreachable),
+                );
+            }
+        }
+
+        let mut block = Block::new(merge_id);
+        let result_id = self.id_gen.next();
+        block.body.push(Instruction::phi(
+            column_type_id,
+            result_id,
+            &var_parent_pairs,
+        ));
+
+        function.consume(block, Instruction::return_value(result_id));
+        function.to_words(&mut self.logical_layout.function_definitions);
+        Ok(())
+    }
+
     fn write_function(
         &mut self,
         ir_function: &crate::Function,
@@ -984,7 +1410,12 @@ impl Writer {
                         gv.handle_id = id;
                     } else if global_needs_wrapper(ir_module, var) {
                         let class = map_storage_class(var.space);
-                        let pointer_type_id = self.get_handle_pointer_type_id(var.ty, class);
+                        let pointer_type_id = match self.std140_compat_uniform_types.get(&var.ty) {
+                            Some(std140_type_info) if var.space == crate::AddressSpace::Uniform => {
+                                self.get_pointer_type_id(std140_type_info.type_id, class)
+                            }
+                            _ => self.get_handle_pointer_type_id(var.ty, class),
+                        };
                         let index_id = self.get_index_constant(0);
                         let id = self.id_gen.next();
                         prelude.body.push(Instruction::access_chain(
@@ -1523,6 +1954,247 @@ impl Writer {
         Ok(id)
     }
 
+    /// Writes a std140 layout compatible type declaration for a type. Returns
+    /// the ID of the declared type, or None if no declaration is required.
+    ///
+    /// This should be called for any type for which there exists a
+    /// [`GlobalVariable`] in the [`Uniform`] address space. If the type already
+    /// adheres to std140 layout rules it will return without declaring any
+    /// types. If the type contains another type which requires a std140
+    /// compatible type declaration, it will recursively call itself.
+    ///
+    /// When `handle` refers to a [`TypeInner::Matrix`] with 2 rows, the
+    /// declared type will be an `OpTypeStruct` containing an `OpVector` for
+    /// each of the matrix's columns.
+    ///
+    /// When `handle` refers to a [`TypeInner::Array`] whose base type is a
+    /// matrix with 2 rows, this will declare an `OpTypeArray` whose element
+    /// type is the matrix's corresponding std140 compatible type.
+    ///
+    /// When `handle` refers to a [`TypeInner::Struct`] and any of its members
+    /// require a std140 compatible type declaration, this will declare a new
+    /// struct with the following rules:
+    /// * Struct or array members will be declared with their std140 compatible
+    ///   type declaration, if one is required.
+    /// * Two-row matrix members will have each of their columns hoisted
+    ///   directly into the struct as 2-component vector members.
+    /// * All other members will be declared with their normal type.
+    ///
+    /// Note that this means the Naga IR index of a struct member may not match
+    /// the index in the generated SPIR-V. The mapping can be obtained via
+    /// `Std140TypeInfo::member_indices`.
+    ///
+    /// [`GlobalVariable`]: crate::GlobalVariable
+    /// [`Uniform`]: crate::AddressSpace::Uniform
+    /// [`TypeInner::Matrix`]: crate::TypeInner::Matrix
+    /// [`TypeInner::Array`]: crate::TypeInner::Array
+    /// [`TypeInner::Struct`]: crate::TypeInner::Struct
+    fn write_std140_compat_type_declaration(
+        &mut self,
+        module: &crate::Module,
+        handle: Handle<crate::Type>,
+    ) -> Result<Option<Word>, Error> {
+        if let Some(std140_type_info) = self.std140_compat_uniform_types.get(&handle) {
+            return Ok(Some(std140_type_info.type_id));
+        }
+
+        let type_inner = &module.types[handle].inner;
+        let std140_type_id = match *type_inner {
+            crate::TypeInner::Matrix {
+                columns,
+                rows: rows @ crate::VectorSize::Bi,
+                scalar,
+            } => {
+                let std140_type_id = self.id_gen.next();
+                let mut member_type_ids: ArrayVec<Word, 4> = ArrayVec::new();
+                let column_type_id =
+                    self.get_numeric_type_id(NumericType::Vector { size: rows, scalar });
+                for column in 0..columns as u32 {
+                    member_type_ids.push(column_type_id);
+                    self.annotations.push(Instruction::member_decorate(
+                        std140_type_id,
+                        column,
+                        spirv::Decoration::Offset,
+                        &[column * rows as u32 * scalar.width as u32],
+                    ));
+                    if self.flags.contains(WriterFlags::DEBUG) {
+                        self.debugs.push(Instruction::member_name(
+                            std140_type_id,
+                            column,
+                            &format!("col{column}"),
+                        ));
+                    }
+                }
+                Instruction::type_struct(std140_type_id, &member_type_ids)
+                    .to_words(&mut self.logical_layout.declarations);
+                self.std140_compat_uniform_types.insert(
+                    handle,
+                    Std140CompatTypeInfo {
+                        type_id: std140_type_id,
+                        member_indices: Vec::new(),
+                    },
+                );
+                Some(std140_type_id)
+            }
+            crate::TypeInner::Array { base, size, stride } => {
+                match self.write_std140_compat_type_declaration(module, base)? {
+                    Some(std140_base_type_id) => {
+                        let std140_type_id = self.id_gen.next();
+                        self.decorate(std140_type_id, spirv::Decoration::ArrayStride, &[stride]);
+                        let instruction = match size.resolve(module.to_ctx())? {
+                            crate::proc::IndexableLength::Known(length) => {
+                                let length_id = self.get_index_constant(length);
+                                Instruction::type_array(
+                                    std140_type_id,
+                                    std140_base_type_id,
+                                    length_id,
+                                )
+                            }
+                            crate::proc::IndexableLength::Dynamic => {
+                                unreachable!()
+                            }
+                        };
+                        instruction.to_words(&mut self.logical_layout.declarations);
+                        self.std140_compat_uniform_types.insert(
+                            handle,
+                            Std140CompatTypeInfo {
+                                type_id: std140_type_id,
+                                member_indices: Vec::new(),
+                            },
+                        );
+                        Some(std140_type_id)
+                    }
+                    None => None,
+                }
+            }
+            crate::TypeInner::Struct { ref members, .. } => {
+                let mut needs_std140_type = false;
+                for member in members {
+                    match module.types[member.ty].inner {
+                        // We don't need to write a std140 type for the matrix itself as
+                        // it will be decomposed into the parent struct. As a result, the
+                        // struct does need a std140 type, however.
+                        crate::TypeInner::Matrix {
+                            rows: crate::VectorSize::Bi,
+                            ..
+                        } => needs_std140_type = true,
+                        // If an array member needs a std140 type, because it is an array
+                        // (of an array, etc) of `matCx2`s, then the struct also needs
+                        // a std140 type which uses the std140 type for this member.
+                        crate::TypeInner::Array { .. }
+                            if self
+                                .write_std140_compat_type_declaration(module, member.ty)?
+                                .is_some() =>
+                        {
+                            needs_std140_type = true;
+                        }
+                        _ => {}
+                    }
+                }
+
+                if needs_std140_type {
+                    let std140_type_id = self.id_gen.next();
+                    let mut member_ids = Vec::new();
+                    let mut member_indices = Vec::new();
+                    let mut next_index = 0;
+
+                    for member in members {
+                        member_indices.push(next_index);
+                        match module.types[member.ty].inner {
+                            crate::TypeInner::Matrix {
+                                columns,
+                                rows: rows @ crate::VectorSize::Bi,
+                                scalar,
+                            } => {
+                                let vector_type_id =
+                                    self.get_numeric_type_id(NumericType::Vector {
+                                        size: rows,
+                                        scalar,
+                                    });
+                                for column in 0..columns as u32 {
+                                    self.annotations.push(Instruction::member_decorate(
+                                        std140_type_id,
+                                        next_index,
+                                        spirv::Decoration::Offset,
+                                        &[member.offset
+                                            + column * rows as u32 * scalar.width as u32],
+                                    ));
+                                    if self.flags.contains(WriterFlags::DEBUG) {
+                                        if let Some(ref name) = member.name {
+                                            self.debugs.push(Instruction::member_name(
+                                                std140_type_id,
+                                                next_index,
+                                                &format!("{name}_col{column}"),
+                                            ));
+                                        }
+                                    }
+                                    member_ids.push(vector_type_id);
+                                    next_index += 1;
+                                }
+                            }
+                            _ => {
+                                let member_id =
+                                    match self.std140_compat_uniform_types.get(&member.ty) {
+                                        Some(std140_member_type_info) => {
+                                            self.annotations.push(Instruction::member_decorate(
+                                                std140_type_id,
+                                                next_index,
+                                                spirv::Decoration::Offset,
+                                                &[member.offset],
+                                            ));
+                                            if self.flags.contains(WriterFlags::DEBUG) {
+                                                if let Some(ref name) = member.name {
+                                                    self.debugs.push(Instruction::member_name(
+                                                        std140_type_id,
+                                                        next_index,
+                                                        name,
+                                                    ));
+                                                }
+                                            }
+                                            std140_member_type_info.type_id
+                                        }
+                                        None => {
+                                            self.decorate_struct_member(
+                                                std140_type_id,
+                                                next_index as usize,
+                                                member,
+                                                &module.types,
+                                            )?;
+                                            self.get_handle_type_id(member.ty)
+                                        }
+                                    };
+                                member_ids.push(member_id);
+                                next_index += 1;
+                            }
+                        }
+                    }
+
+                    Instruction::type_struct(std140_type_id, &member_ids)
+                        .to_words(&mut self.logical_layout.declarations);
+                    self.std140_compat_uniform_types.insert(
+                        handle,
+                        Std140CompatTypeInfo {
+                            type_id: std140_type_id,
+                            member_indices,
+                        },
+                    );
+                    Some(std140_type_id)
+                } else {
+                    None
+                }
+            }
+            _ => None,
+        };
+
+        if let Some(std140_type_id) = std140_type_id {
+            if self.flags.contains(WriterFlags::DEBUG) {
+                let name = format!("std140_{:?}", handle.for_debug(&module.types));
+                self.debugs.push(Instruction::name(std140_type_id, &name));
+            }
+        }
+        Ok(std140_type_id)
+    }
+
     fn request_image_format_capabilities(
         &mut self,
         format: spirv::ImageFormat,
@@ -2295,16 +2967,31 @@ impl Writer {
             let wrapper_type_id = self.id_gen.next();
 
             self.decorate(wrapper_type_id, Decoration::Block, &[]);
-            let member = crate::StructMember {
-                name: None,
-                ty: global_variable.ty,
-                binding: None,
-                offset: 0,
-            };
-            self.decorate_struct_member(wrapper_type_id, 0, &member, &ir_module.types)?;
 
-            Instruction::type_struct(wrapper_type_id, &[inner_type_id])
-                .to_words(&mut self.logical_layout.declarations);
+            match self.std140_compat_uniform_types.get(&global_variable.ty) {
+                Some(std140_type_info) if global_variable.space == crate::AddressSpace::Uniform => {
+                    self.annotations.push(Instruction::member_decorate(
+                        wrapper_type_id,
+                        0,
+                        Decoration::Offset,
+                        &[0],
+                    ));
+                    Instruction::type_struct(wrapper_type_id, &[std140_type_info.type_id])
+                        .to_words(&mut self.logical_layout.declarations);
+                }
+                _ => {
+                    let member = crate::StructMember {
+                        name: None,
+                        ty: global_variable.ty,
+                        binding: None,
+                        offset: 0,
+                    };
+                    self.decorate_struct_member(wrapper_type_id, 0, &member, &ir_module.types)?;
+
+                    Instruction::type_struct(wrapper_type_id, &[inner_type_id])
+                        .to_words(&mut self.logical_layout.declarations);
+                }
+            }
 
             let pointer_type_id = self.id_gen.next();
             Instruction::type_pointer(pointer_type_id, class, wrapper_type_id)
@@ -2538,6 +3225,13 @@ impl Writer {
             self.write_type_declaration_arena(ir_module, handle)?;
         }
 
+        // write std140 layout compatible types required by uniforms
+        for (_, var) in ir_module.global_variables.iter() {
+            if var.space == crate::AddressSpace::Uniform {
+                self.write_std140_compat_type_declaration(ir_module, var.ty)?;
+            }
+        }
+
         // write all const-expressions as constants
         self.constant_ids
             .resize(ir_module.global_expressions.len(), 0);
diff --git a/naga/tests/in/wgsl/access.wgsl b/naga/tests/in/wgsl/access.wgsl
index 0e89bcfb0eb..65d959656b8 100644
--- a/naga/tests/in/wgsl/access.wgsl
+++ b/naga/tests/in/wgsl/access.wgsl
@@ -35,8 +35,8 @@ var<uniform> baz: Baz;
 var<storage,read_write> qux: vec2<i32>;
 
 fn test_matrix_within_struct_accesses() {
-    // Test HLSL accesses to Cx2 matrices. There are additional tests
-    // in `hlsl_mat_cx2.wgsl`.
+    // Test accesses to Cx2 matrices. There are additional tests in
+    // `mat_cx2.wgsl`.
 
     var idx = 1;
 
diff --git a/naga/tests/in/wgsl/hlsl_mat_cx2.toml b/naga/tests/in/wgsl/hlsl_mat_cx2.toml
deleted file mode 100644
index 3ca0b52f4e5..00000000000
--- a/naga/tests/in/wgsl/hlsl_mat_cx2.toml
+++ /dev/null
@@ -1 +0,0 @@
-targets = "HLSL"
diff --git a/naga/tests/in/wgsl/hlsl_mat_cx3.toml b/naga/tests/in/wgsl/hlsl_mat_cx3.toml
deleted file mode 100644
index 3ca0b52f4e5..00000000000
--- a/naga/tests/in/wgsl/hlsl_mat_cx3.toml
+++ /dev/null
@@ -1 +0,0 @@
-targets = "HLSL"
diff --git a/naga/tests/in/wgsl/mat_cx2.toml b/naga/tests/in/wgsl/mat_cx2.toml
new file mode 100644
index 00000000000..cc4a177e5f4
--- /dev/null
+++ b/naga/tests/in/wgsl/mat_cx2.toml
@@ -0,0 +1,4 @@
+targets = "HLSL | SPIRV"
+
+[spv]
+debug = true
diff --git a/naga/tests/in/wgsl/hlsl_mat_cx2.wgsl b/naga/tests/in/wgsl/mat_cx2.wgsl
similarity index 97%
rename from naga/tests/in/wgsl/hlsl_mat_cx2.wgsl
rename to naga/tests/in/wgsl/mat_cx2.wgsl
index 50bc188794a..6a2a3a6e63d 100644
--- a/naga/tests/in/wgsl/hlsl_mat_cx2.wgsl
+++ b/naga/tests/in/wgsl/mat_cx2.wgsl
@@ -1,9 +1,9 @@
-// Test HLSL handling of N-by-2 matrices.
-// See the doc comment on `naga::back::hlsl` for details.
+// Test handling of N-by-2 matrices.
+// See the doc comments on `naga::back::hlsl` and `naga::back::spv` for details.
 //
 // There are additional tests in `access.wgsl`.
 //
-// Tests that we don't apply this handling to other sizes are in hlsl_mat_cx3.wgsl.
+// Tests that we don't apply this handling to other sizes are in mat_cx3.wgsl.
 
 // Access type (3rd item in variable names)
 // S = Struct
diff --git a/naga/tests/in/wgsl/mat_cx3.toml b/naga/tests/in/wgsl/mat_cx3.toml
new file mode 100644
index 00000000000..a4553fd2f78
--- /dev/null
+++ b/naga/tests/in/wgsl/mat_cx3.toml
@@ -0,0 +1 @@
+targets = "HLSL | SPIRV"
diff --git a/naga/tests/in/wgsl/hlsl_mat_cx3.wgsl b/naga/tests/in/wgsl/mat_cx3.wgsl
similarity index 98%
rename from naga/tests/in/wgsl/hlsl_mat_cx3.wgsl
rename to naga/tests/in/wgsl/mat_cx3.wgsl
index e33f10fc9c5..bb17a437132 100644
--- a/naga/tests/in/wgsl/hlsl_mat_cx3.wgsl
+++ b/naga/tests/in/wgsl/mat_cx3.wgsl
@@ -1,5 +1,5 @@
 // Test HLSL handling of N-by-3 matrices. These should not receive the special
-// treatment that N-by-2 matrices receive (which is tested in hlsl_mat_cx2).
+// treatment that N-by-2 matrices receive (which is tested in mat_cx2.wgsl).
 
 // Access type (3rd item in variable names)
 // S = Struct
diff --git a/naga/tests/out/hlsl/wgsl-hlsl_mat_cx2.hlsl b/naga/tests/out/hlsl/wgsl-mat_cx2.hlsl
similarity index 100%
rename from naga/tests/out/hlsl/wgsl-hlsl_mat_cx2.hlsl
rename to naga/tests/out/hlsl/wgsl-mat_cx2.hlsl
diff --git a/naga/tests/out/hlsl/wgsl-hlsl_mat_cx2.ron b/naga/tests/out/hlsl/wgsl-mat_cx2.ron
similarity index 100%
rename from naga/tests/out/hlsl/wgsl-hlsl_mat_cx2.ron
rename to naga/tests/out/hlsl/wgsl-mat_cx2.ron
diff --git a/naga/tests/out/hlsl/wgsl-hlsl_mat_cx3.hlsl b/naga/tests/out/hlsl/wgsl-mat_cx3.hlsl
similarity index 100%
rename from naga/tests/out/hlsl/wgsl-hlsl_mat_cx3.hlsl
rename to naga/tests/out/hlsl/wgsl-mat_cx3.hlsl
diff --git a/naga/tests/out/hlsl/wgsl-hlsl_mat_cx3.ron b/naga/tests/out/hlsl/wgsl-mat_cx3.ron
similarity index 100%
rename from naga/tests/out/hlsl/wgsl-hlsl_mat_cx3.ron
rename to naga/tests/out/hlsl/wgsl-mat_cx3.ron
diff --git a/naga/tests/out/spv/wgsl-access.spvasm b/naga/tests/out/spv/wgsl-access.spvasm
index 31e8e5d4c0b..1678c50d436 100644
--- a/naga/tests/out/spv/wgsl-access.spvasm
+++ b/naga/tests/out/spv/wgsl-access.spvasm
@@ -1,16 +1,16 @@
 ; SPIR-V
 ; Version: 1.1
 ; Generator: rspirv
-; Bound: 414
+; Bound: 511
 OpCapability Shader
 OpExtension "SPV_KHR_storage_buffer_storage_class"
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
-OpEntryPoint Vertex %331 "foo_vert" %326 %329
-OpEntryPoint Fragment %387 "foo_frag" %386
-OpEntryPoint GLCompute %405 "foo_compute"
-OpExecutionMode %387 OriginUpperLeft
-OpExecutionMode %405 LocalSize 1 1 1
+OpEntryPoint Vertex %428 "foo_vert" %423 %426
+OpEntryPoint Fragment %484 "foo_frag" %483
+OpEntryPoint GLCompute %502 "foo_compute"
+OpExecutionMode %484 OriginUpperLeft
+OpExecutionMode %502 LocalSize 1 1 1
 %3 = OpString "access.wgsl"
 OpSource Unknown 0 %3 "// This snapshot tests accessing various containers, dereferencing pointers.
 
@@ -49,8 +49,8 @@ var<uniform> baz: Baz;
 var<storage,read_write> qux: vec2<i32>;
 
 fn test_matrix_within_struct_accesses() {
-    // Test HLSL accesses to Cx2 matrices. There are additional tests
-    // in `hlsl_mat_cx2.wgsl`.
+    // Test accesses to Cx2 matrices. There are additional tests in
+    // `mat_cx2.wgsl`.
 
     var idx = 1;
 
@@ -295,55 +295,71 @@ OpName %46 "Inner"
 OpMemberName %47 0 "om_nom_nom"
 OpMemberName %47 1 "thing"
 OpName %47 "Outer"
-OpName %52 "msl_padding_global_const"
-OpName %54 "bar"
-OpName %56 "baz"
-OpName %59 "qux"
-OpName %62 "nested_mat_cx2"
-OpName %66 "test_matrix_within_struct_accesses"
-OpName %94 "idx"
-OpName %96 "t"
-OpName %140 "test_matrix_within_array_within_struct_accesses"
-OpName %150 "idx"
-OpName %151 "t"
-OpName %197 "foo"
-OpName %198 "read_from_private"
-OpName %203 "a"
-OpName %204 "test_arr_as_arg"
-OpName %210 "p"
-OpName %211 "assign_through_ptr_fn"
-OpName %216 "foo"
-OpName %217 "assign_array_through_ptr_fn"
-OpName %224 "assign_through_ptr"
-OpName %229 "val"
-OpName %230 "arr"
-OpName %235 "p"
-OpName %236 "fetch_arg_ptr_member"
-OpName %242 "p"
-OpName %243 "assign_to_arg_ptr_member"
-OpName %248 "p"
-OpName %249 "fetch_arg_ptr_array_element"
-OpName %255 "p"
-OpName %256 "assign_to_arg_ptr_array_element"
-OpName %261 "assign_to_ptr_components"
-OpName %262 "s1"
-OpName %264 "a1"
-OpName %272 "value"
-OpName %273 "index_ptr"
-OpName %275 "a"
-OpName %284 "member_ptr"
-OpName %288 "s"
-OpName %294 "let_members_of_members"
-OpName %305 "var_members_of_members"
-OpName %306 "thing"
-OpName %308 "inner"
-OpName %311 "delishus"
-OpName %326 "vi"
-OpName %331 "foo_vert"
-OpName %342 "foo"
-OpName %343 "c2"
-OpName %387 "foo_frag"
-OpName %405 "foo_compute"
+OpMemberName %48 0 "m_col0"
+OpMemberName %48 1 "m_col1"
+OpMemberName %48 2 "m_col2"
+OpName %48 "std140_Baz"
+OpMemberName %49 0 "col0"
+OpMemberName %49 1 "col1"
+OpMemberName %49 2 "col2"
+OpMemberName %49 3 "col3"
+OpName %49 "std140_mat4x2<f32>"
+OpName %50 "std140_array<mat4x2<f32>, 2>"
+OpMemberName %51 0 "am"
+OpName %51 "std140_MatCx2InArray"
+OpName %56 "msl_padding_global_const"
+OpName %58 "bar"
+OpName %60 "baz"
+OpName %63 "qux"
+OpName %66 "nested_mat_cx2"
+OpName %69 "mat3x2<f32>_get_column"
+OpName %84 "test_matrix_within_struct_accesses"
+OpName %113 "idx"
+OpName %115 "t"
+OpName %186 "array<mat4x2<f32>, 2>_from_std140"
+OpName %190 "mat4x2<f32>_from_std140"
+OpName %204 "mat4x2<f32>_get_column"
+OpName %221 "test_matrix_within_array_within_struct_accesses"
+OpName %232 "idx"
+OpName %233 "t"
+OpName %294 "foo"
+OpName %295 "read_from_private"
+OpName %300 "a"
+OpName %301 "test_arr_as_arg"
+OpName %307 "p"
+OpName %308 "assign_through_ptr_fn"
+OpName %313 "foo"
+OpName %314 "assign_array_through_ptr_fn"
+OpName %321 "assign_through_ptr"
+OpName %326 "val"
+OpName %327 "arr"
+OpName %332 "p"
+OpName %333 "fetch_arg_ptr_member"
+OpName %339 "p"
+OpName %340 "assign_to_arg_ptr_member"
+OpName %345 "p"
+OpName %346 "fetch_arg_ptr_array_element"
+OpName %352 "p"
+OpName %353 "assign_to_arg_ptr_array_element"
+OpName %358 "assign_to_ptr_components"
+OpName %359 "s1"
+OpName %361 "a1"
+OpName %369 "value"
+OpName %370 "index_ptr"
+OpName %372 "a"
+OpName %381 "member_ptr"
+OpName %385 "s"
+OpName %391 "let_members_of_members"
+OpName %402 "var_members_of_members"
+OpName %403 "thing"
+OpName %405 "inner"
+OpName %408 "delishus"
+OpName %423 "vi"
+OpName %428 "foo_vert"
+OpName %439 "foo"
+OpName %440 "c2"
+OpName %484 "foo_frag"
+OpName %502 "foo_compute"
 OpMemberDecorate %7 0 Offset 0
 OpMemberDecorate %7 1 Offset 16
 OpMemberDecorate %7 2 Offset 28
@@ -381,23 +397,32 @@ OpMemberDecorate %45 0 Offset 0
 OpMemberDecorate %46 0 Offset 0
 OpMemberDecorate %47 0 Offset 0
 OpMemberDecorate %47 1 Offset 4
-OpDecorate %54 DescriptorSet 0
-OpDecorate %54 Binding 0
-OpDecorate %56 DescriptorSet 0
-OpDecorate %56 Binding 1
-OpDecorate %57 Block
-OpMemberDecorate %57 0 Offset 0
-OpDecorate %59 DescriptorSet 0
-OpDecorate %59 Binding 2
-OpDecorate %60 Block
-OpMemberDecorate %60 0 Offset 0
-OpDecorate %62 DescriptorSet 0
-OpDecorate %62 Binding 3
-OpDecorate %63 Block
-OpMemberDecorate %63 0 Offset 0
-OpDecorate %326 BuiltIn VertexIndex
-OpDecorate %329 BuiltIn Position
-OpDecorate %386 Location 0
+OpMemberDecorate %48 0 Offset 0
+OpMemberDecorate %48 1 Offset 8
+OpMemberDecorate %48 2 Offset 16
+OpMemberDecorate %49 0 Offset 0
+OpMemberDecorate %49 1 Offset 8
+OpMemberDecorate %49 2 Offset 16
+OpMemberDecorate %49 3 Offset 24
+OpDecorate %50 ArrayStride 32
+OpMemberDecorate %51 0 Offset 0
+OpDecorate %58 DescriptorSet 0
+OpDecorate %58 Binding 0
+OpDecorate %60 DescriptorSet 0
+OpDecorate %60 Binding 1
+OpDecorate %61 Block
+OpMemberDecorate %61 0 Offset 0
+OpDecorate %63 DescriptorSet 0
+OpDecorate %63 Binding 2
+OpDecorate %64 Block
+OpMemberDecorate %64 0 Offset 0
+OpDecorate %66 DescriptorSet 0
+OpDecorate %66 Binding 3
+OpDecorate %67 Block
+OpMemberDecorate %67 0 Offset 0
+OpDecorate %423 BuiltIn VertexIndex
+OpDecorate %426 BuiltIn Position
+OpDecorate %483 Location 0
 %2 = OpTypeVoid
 %4 = OpTypeInt 32 0
 %5 = OpTypeVector %4 3
@@ -443,294 +468,412 @@ OpDecorate %386 Location 0
 %45 = OpTypeStruct %6
 %46 = OpTypeStruct %6
 %47 = OpTypeStruct %46 %4
-%48 = OpConstant  %4  0
-%49 = OpConstantComposite  %5  %48 %48 %48
-%50 = OpConstant  %6  0
-%51 = OpConstantComposite  %7  %48 %49 %50
-%53 = OpTypePointer Private %7
-%52 = OpVariable  %53  Private %51
-%55 = OpTypePointer StorageBuffer %21
-%54 = OpVariable  %55  StorageBuffer
-%57 = OpTypeStruct %23
-%58 = OpTypePointer Uniform %57
-%56 = OpVariable  %58  Uniform
-%60 = OpTypeStruct %24
-%61 = OpTypePointer StorageBuffer %60
-%59 = OpVariable  %61  StorageBuffer
-%63 = OpTypeStruct %27
-%64 = OpTypePointer Uniform %63
-%62 = OpVariable  %64  Uniform
-%67 = OpTypeFunction %2
-%68 = OpTypePointer Uniform %23
-%70 = OpConstant  %6  1
-%71 = OpConstant  %9  1
-%72 = OpConstantComposite  %13  %71 %71
-%73 = OpConstant  %9  2
-%74 = OpConstantComposite  %13  %73 %73
-%75 = OpConstant  %9  3
-%76 = OpConstantComposite  %13  %75 %75
-%77 = OpConstantComposite  %22  %72 %74 %76
-%78 = OpConstantComposite  %23  %77
-%79 = OpConstant  %9  6
-%80 = OpConstantComposite  %13  %79 %79
-%81 = OpConstant  %9  5
-%82 = OpConstantComposite  %13  %81 %81
-%83 = OpConstant  %9  4
-%84 = OpConstantComposite  %13  %83 %83
-%85 = OpConstantComposite  %22  %80 %82 %84
-%86 = OpConstant  %9  9
-%87 = OpConstantComposite  %13  %86 %86
-%88 = OpConstant  %9  90
-%89 = OpConstantComposite  %13  %88 %88
-%90 = OpConstant  %9  10
-%91 = OpConstant  %9  20
-%92 = OpConstant  %9  30
-%93 = OpConstant  %9  40
-%95 = OpTypePointer Function %6
-%97 = OpTypePointer Function %23
-%101 = OpTypePointer Uniform %22
-%104 = OpTypePointer Uniform %13
-%110 = OpTypePointer Uniform %9
-%125 = OpTypePointer Function %22
-%127 = OpTypePointer Function %13
-%141 = OpTypePointer Uniform %27
-%143 = OpConstantNull  %26
-%144 = OpConstantComposite  %27  %143
-%145 = OpConstant  %9  8
-%146 = OpConstantComposite  %13  %145 %145
-%147 = OpConstant  %9  7
-%148 = OpConstantComposite  %13  %147 %147
-%149 = OpConstantComposite  %25  %146 %148 %80 %82
-%152 = OpTypePointer Function %27
-%156 = OpTypePointer Uniform %26
-%159 = OpTypePointer Uniform %25
-%181 = OpTypePointer Function %26
-%183 = OpTypePointer Function %25
-%199 = OpTypeFunction %9 %28
-%205 = OpTypeFunction %9 %30
-%212 = OpTypeFunction %2 %34
-%213 = OpConstant  %4  42
-%218 = OpTypeFunction %2 %36
-%219 = OpConstantComposite  %32  %71 %71 %71 %71
-%220 = OpConstantComposite  %32  %73 %73 %73 %73
-%221 = OpConstantComposite  %35  %219 %220
-%225 = OpConstant  %4  33
-%226 = OpConstantComposite  %32  %79 %79 %79 %79
-%227 = OpConstantComposite  %32  %147 %147 %147 %147
-%228 = OpConstantComposite  %35  %226 %227
-%237 = OpTypeFunction %4 %38
-%244 = OpTypeFunction %2 %38
-%250 = OpTypeFunction %4 %41
-%257 = OpTypeFunction %2 %41
-%263 = OpConstantNull  %37
-%265 = OpConstantNull  %39
-%274 = OpTypeFunction %42 %42
-%276 = OpTypePointer Function %43
-%277 = OpConstantNull  %43
-%280 = OpTypePointer Function %42
-%285 = OpTypeFunction %6
-%286 = OpConstant  %6  42
-%287 = OpConstantComposite  %45  %286
-%289 = OpTypePointer Function %45
-%295 = OpConstantNull  %47
-%307 = OpTypePointer Function %47
-%309 = OpTypePointer Function %46
-%310 = OpConstantNull  %46
-%312 = OpConstantNull  %6
-%327 = OpTypePointer Input %4
-%326 = OpVariable  %327  Input
-%330 = OpTypePointer Output %32
-%329 = OpVariable  %330  Output
-%333 = OpTypePointer StorageBuffer %24
-%336 = OpConstant  %9  0
-%337 = OpConstant  %4  3
-%338 = OpConstant  %6  3
-%339 = OpConstant  %6  4
-%340 = OpConstant  %6  5
-%341 = OpConstantNull  %30
-%344 = OpTypePointer Function %33
-%345 = OpConstantNull  %33
-%351 = OpTypePointer StorageBuffer %10
-%354 = OpTypePointer StorageBuffer %19
-%357 = OpTypePointer StorageBuffer %11
-%358 = OpTypePointer StorageBuffer %9
-%361 = OpTypePointer StorageBuffer %20
-%364 = OpTypePointer StorageBuffer %8
-%365 = OpTypePointer StorageBuffer %6
-%370 = OpConstant  %9  -2147483600
-%371 = OpConstant  %9  2147483500
-%380 = OpTypeVector %6 4
-%386 = OpVariable  %330  Output
-%389 = OpConstantComposite  %11  %336 %336 %336
-%390 = OpConstantComposite  %11  %71 %71 %71
-%391 = OpConstantComposite  %11  %73 %73 %73
-%392 = OpConstantComposite  %11  %75 %75 %75
-%393 = OpConstantComposite  %10  %389 %390 %391 %392
-%394 = OpConstantComposite  %18  %48 %48
-%395 = OpConstantComposite  %18  %44 %44
-%396 = OpConstantComposite  %19  %394 %395
-%397 = OpConstantNull  %24
-%398 = OpConstantComposite  %32  %336 %336 %336 %336
-%406 = OpConstantTrue  %42
-%66 = OpFunction  %2  None %67
-%65 = OpLabel
-%94 = OpVariable  %95  Function %70
-%96 = OpVariable  %97  Function %78
-%69 = OpAccessChain  %68  %56 %48
-OpBranch %98
-%98 = OpLabel
+%48 = OpTypeStruct %13 %13 %13
+%49 = OpTypeStruct %13 %13 %13 %13
+%50 = OpTypeArray %49 %15
+%51 = OpTypeStruct %50
+%52 = OpConstant  %4  0
+%53 = OpConstantComposite  %5  %52 %52 %52
+%54 = OpConstant  %6  0
+%55 = OpConstantComposite  %7  %52 %53 %54
+%57 = OpTypePointer Private %7
+%56 = OpVariable  %57  Private %55
+%59 = OpTypePointer StorageBuffer %21
+%58 = OpVariable  %59  StorageBuffer
+%61 = OpTypeStruct %48
+%62 = OpTypePointer Uniform %61
+%60 = OpVariable  %62  Uniform
+%64 = OpTypeStruct %24
+%65 = OpTypePointer StorageBuffer %64
+%63 = OpVariable  %65  StorageBuffer
+%67 = OpTypeStruct %51
+%68 = OpTypePointer Uniform %67
+%66 = OpVariable  %68  Uniform
+%72 = OpTypeFunction %13 %22 %4
+%85 = OpTypeFunction %2
+%86 = OpTypePointer Uniform %48
+%88 = OpConstant  %6  1
+%89 = OpTypePointer Uniform %23
+%90 = OpConstant  %9  1
+%91 = OpConstantComposite  %13  %90 %90
+%92 = OpConstant  %9  2
+%93 = OpConstantComposite  %13  %92 %92
+%94 = OpConstant  %9  3
+%95 = OpConstantComposite  %13  %94 %94
+%96 = OpConstantComposite  %22  %91 %93 %95
+%97 = OpConstantComposite  %23  %96
+%98 = OpConstant  %9  6
+%99 = OpConstantComposite  %13  %98 %98
+%100 = OpConstant  %9  5
+%101 = OpConstantComposite  %13  %100 %100
+%102 = OpConstant  %9  4
+%103 = OpConstantComposite  %13  %102 %102
+%104 = OpConstantComposite  %22  %99 %101 %103
+%105 = OpConstant  %9  9
+%106 = OpConstantComposite  %13  %105 %105
+%107 = OpConstant  %9  90
+%108 = OpConstantComposite  %13  %107 %107
+%109 = OpConstant  %9  10
+%110 = OpConstant  %9  20
+%111 = OpConstant  %9  30
+%112 = OpConstant  %9  40
+%114 = OpTypePointer Function %6
+%116 = OpTypePointer Function %23
+%120 = OpTypePointer Uniform %22
+%121 = OpTypePointer Uniform %13
+%141 = OpTypePointer Uniform %9
+%172 = OpTypePointer Function %22
+%174 = OpTypePointer Function %13
+%187 = OpTypeFunction %26 %50
+%191 = OpTypeFunction %25 %49
+%207 = OpTypeFunction %13 %25 %4
+%222 = OpTypePointer Uniform %51
+%224 = OpTypePointer Uniform %27
+%225 = OpConstantNull  %26
+%226 = OpConstantComposite  %27  %225
+%227 = OpConstant  %9  8
+%228 = OpConstantComposite  %13  %227 %227
+%229 = OpConstant  %9  7
+%230 = OpConstantComposite  %13  %229 %229
+%231 = OpConstantComposite  %25  %228 %230 %99 %101
+%234 = OpTypePointer Function %27
+%238 = OpTypePointer Uniform %26
+%239 = OpTypePointer Uniform %50
+%243 = OpTypePointer Uniform %25
+%244 = OpTypePointer Uniform %49
+%278 = OpTypePointer Function %26
+%280 = OpTypePointer Function %25
+%296 = OpTypeFunction %9 %28
+%302 = OpTypeFunction %9 %30
+%309 = OpTypeFunction %2 %34
+%310 = OpConstant  %4  42
+%315 = OpTypeFunction %2 %36
+%316 = OpConstantComposite  %32  %90 %90 %90 %90
+%317 = OpConstantComposite  %32  %92 %92 %92 %92
+%318 = OpConstantComposite  %35  %316 %317
+%322 = OpConstant  %4  33
+%323 = OpConstantComposite  %32  %98 %98 %98 %98
+%324 = OpConstantComposite  %32  %229 %229 %229 %229
+%325 = OpConstantComposite  %35  %323 %324
+%334 = OpTypeFunction %4 %38
+%341 = OpTypeFunction %2 %38
+%347 = OpTypeFunction %4 %41
+%354 = OpTypeFunction %2 %41
+%360 = OpConstantNull  %37
+%362 = OpConstantNull  %39
+%371 = OpTypeFunction %42 %42
+%373 = OpTypePointer Function %43
+%374 = OpConstantNull  %43
+%377 = OpTypePointer Function %42
+%382 = OpTypeFunction %6
+%383 = OpConstant  %6  42
+%384 = OpConstantComposite  %45  %383
+%386 = OpTypePointer Function %45
+%392 = OpConstantNull  %47
+%404 = OpTypePointer Function %47
+%406 = OpTypePointer Function %46
+%407 = OpConstantNull  %46
+%409 = OpConstantNull  %6
+%424 = OpTypePointer Input %4
+%423 = OpVariable  %424  Input
+%427 = OpTypePointer Output %32
+%426 = OpVariable  %427  Output
+%430 = OpTypePointer StorageBuffer %24
+%433 = OpConstant  %9  0
+%434 = OpConstant  %4  3
+%435 = OpConstant  %6  3
+%436 = OpConstant  %6  4
+%437 = OpConstant  %6  5
+%438 = OpConstantNull  %30
+%441 = OpTypePointer Function %33
+%442 = OpConstantNull  %33
+%448 = OpTypePointer StorageBuffer %10
+%451 = OpTypePointer StorageBuffer %19
+%454 = OpTypePointer StorageBuffer %11
+%455 = OpTypePointer StorageBuffer %9
+%458 = OpTypePointer StorageBuffer %20
+%461 = OpTypePointer StorageBuffer %8
+%462 = OpTypePointer StorageBuffer %6
+%467 = OpConstant  %9  -2147483600
+%468 = OpConstant  %9  2147483500
+%477 = OpTypeVector %6 4
+%483 = OpVariable  %427  Output
+%486 = OpConstantComposite  %11  %433 %433 %433
+%487 = OpConstantComposite  %11  %90 %90 %90
+%488 = OpConstantComposite  %11  %92 %92 %92
+%489 = OpConstantComposite  %11  %94 %94 %94
+%490 = OpConstantComposite  %10  %486 %487 %488 %489
+%491 = OpConstantComposite  %18  %52 %52
+%492 = OpConstantComposite  %18  %44 %44
+%493 = OpConstantComposite  %19  %491 %492
+%494 = OpConstantNull  %24
+%495 = OpConstantComposite  %32  %433 %433 %433 %433
+%503 = OpConstantTrue  %42
+%69 = OpFunction  %13  None %72
+%70 = OpFunctionParameter  %22
+%71 = OpFunctionParameter  %4
+%73 = OpLabel
+OpSelectionMerge %74 None
+OpSwitch %71 %78 0 %75 1 %76 2 %77
+%75 = OpLabel
+%79 = OpCompositeExtract  %13  %70 0
+OpBranch %74
+%76 = OpLabel
+%80 = OpCompositeExtract  %13  %70 1
+OpBranch %74
+%77 = OpLabel
+%81 = OpCompositeExtract  %13  %70 2
+OpBranch %74
+%78 = OpLabel
+OpUnreachable
+%74 = OpLabel
+%82 = OpPhi  %13  %79 %75 %80 %76 %81 %77
+OpReturnValue %82
+OpFunctionEnd
+%84 = OpFunction  %2  None %85
+%83 = OpLabel
+%113 = OpVariable  %114  Function %88
+%115 = OpVariable  %116  Function %97
+%87 = OpAccessChain  %86  %60 %52
+OpBranch %117
+%117 = OpLabel
 OpLine %3 43 5
-%99 = OpLoad  %6  %94
-%100 = OpISub  %6  %99 %70
+%118 = OpLoad  %6  %113
+%119 = OpISub  %6  %118 %88
 OpLine %3 43 5
-OpStore %94 %100
+OpStore %113 %119
 OpLine %3 46 14
-%102 = OpAccessChain  %101  %69 %48
-%103 = OpLoad  %22  %102
+%122 = OpAccessChain  %121  %87 %52
+%123 = OpLoad  %13  %122
+%124 = OpAccessChain  %121  %87 %44
+%125 = OpLoad  %13  %124
+%126 = OpAccessChain  %121  %87 %15
+%127 = OpLoad  %13  %126
+%128 = OpCompositeConstruct  %22  %123 %125 %127
 OpLine %3 47 14
 OpLine %3 47 14
-%105 = OpAccessChain  %104  %69 %48 %48
-%106 = OpLoad  %13  %105
+%129 = OpAccessChain  %121  %87 %52
+%130 = OpLoad  %13  %129
 OpLine %3 48 14
-%107 = OpLoad  %6  %94
-%108 = OpAccessChain  %104  %69 %48 %107
-%109 = OpLoad  %13  %108
+%131 = OpLoad  %6  %113
+%132 = OpAccessChain  %121  %87 %52
+%133 = OpLoad  %13  %132
+%134 = OpAccessChain  %121  %87 %44
+%135 = OpLoad  %13  %134
+%136 = OpAccessChain  %121  %87 %15
+%137 = OpLoad  %13  %136
+%138 = OpCompositeConstruct  %22  %133 %135 %137
+%139 = OpBitcast  %4  %131
+%140 = OpFunctionCall  %13  %69 %138 %139
 OpLine %3 49 14
 OpLine %3 49 14
 OpLine %3 49 14
-%111 = OpAccessChain  %110  %69 %48 %48 %44
-%112 = OpLoad  %9  %111
+%142 = OpAccessChain  %141  %87 %52 %44
+%143 = OpLoad  %9  %142
 OpLine %3 50 14
 OpLine %3 50 14
-%113 = OpLoad  %6  %94
-%114 = OpAccessChain  %110  %69 %48 %48 %113
-%115 = OpLoad  %9  %114
+%144 = OpLoad  %6  %113
+%145 = OpAccessChain  %141  %87 %52 %144
+%146 = OpLoad  %9  %145
 OpLine %3 51 14
-%116 = OpLoad  %6  %94
+%147 = OpLoad  %6  %113
 OpLine %3 51 14
-%117 = OpAccessChain  %110  %69 %48 %116 %44
-%118 = OpLoad  %9  %117
+%148 = OpAccessChain  %121  %87 %52
+%149 = OpLoad  %13  %148
+%150 = OpAccessChain  %121  %87 %44
+%151 = OpLoad  %13  %150
+%152 = OpAccessChain  %121  %87 %15
+%153 = OpLoad  %13  %152
+%154 = OpCompositeConstruct  %22  %149 %151 %153
+%155 = OpBitcast  %4  %147
+%156 = OpFunctionCall  %13  %69 %154 %155
+%157 = OpCompositeExtract  %9  %156 1
 OpLine %3 52 14
-%119 = OpLoad  %6  %94
-%120 = OpLoad  %6  %94
-%121 = OpAccessChain  %110  %69 %48 %119 %120
-%122 = OpLoad  %9  %121
+%158 = OpLoad  %6  %113
+%159 = OpLoad  %6  %113
+%160 = OpAccessChain  %121  %87 %52
+%161 = OpLoad  %13  %160
+%162 = OpAccessChain  %121  %87 %44
+%163 = OpLoad  %13  %162
+%164 = OpAccessChain  %121  %87 %15
+%165 = OpLoad  %13  %164
+%166 = OpCompositeConstruct  %22  %161 %163 %165
+%167 = OpBitcast  %4  %158
+%168 = OpFunctionCall  %13  %69 %166 %167
+%169 = OpVectorExtractDynamic  %9  %168 %159
 OpLine %3 54 29
 OpLine %3 54 45
 OpLine %3 54 13
 OpLine %3 56 5
-%123 = OpLoad  %6  %94
-%124 = OpIAdd  %6  %123 %70
+%170 = OpLoad  %6  %113
+%171 = OpIAdd  %6  %170 %88
 OpLine %3 56 5
-OpStore %94 %124
+OpStore %113 %171
 OpLine %3 59 5
 OpLine %3 59 23
 OpLine %3 59 39
 OpLine %3 59 11
 OpLine %3 59 5
-%126 = OpAccessChain  %125  %96 %48
-OpStore %126 %85
+%173 = OpAccessChain  %172  %115 %52
+OpStore %173 %104
 OpLine %3 60 5
 OpLine %3 60 5
 OpLine %3 60 14
 OpLine %3 60 5
-%128 = OpAccessChain  %127  %96 %48 %48
-OpStore %128 %87
+%175 = OpAccessChain  %174  %115 %52 %52
+OpStore %175 %106
 OpLine %3 61 5
-%129 = OpLoad  %6  %94
+%176 = OpLoad  %6  %113
 OpLine %3 61 16
 OpLine %3 61 5
-%130 = OpAccessChain  %127  %96 %48 %129
-OpStore %130 %89
+%177 = OpAccessChain  %174  %115 %52 %176
+OpStore %177 %108
 OpLine %3 62 5
 OpLine %3 62 5
 OpLine %3 62 5
 OpLine %3 62 5
-%131 = OpAccessChain  %28  %96 %48 %48 %44
-OpStore %131 %90
+%178 = OpAccessChain  %28  %115 %52 %52 %44
+OpStore %178 %109
 OpLine %3 63 5
 OpLine %3 63 5
-%132 = OpLoad  %6  %94
+%179 = OpLoad  %6  %113
 OpLine %3 63 5
-%133 = OpAccessChain  %28  %96 %48 %48 %132
-OpStore %133 %91
+%180 = OpAccessChain  %28  %115 %52 %52 %179
+OpStore %180 %110
 OpLine %3 64 5
-%134 = OpLoad  %6  %94
+%181 = OpLoad  %6  %113
 OpLine %3 64 5
 OpLine %3 64 5
-%135 = OpAccessChain  %28  %96 %48 %134 %44
-OpStore %135 %92
+%182 = OpAccessChain  %28  %115 %52 %181 %44
+OpStore %182 %111
 OpLine %3 65 5
-%136 = OpLoad  %6  %94
-%137 = OpLoad  %6  %94
+%183 = OpLoad  %6  %113
+%184 = OpLoad  %6  %113
 OpLine %3 65 5
-%138 = OpAccessChain  %28  %96 %48 %136 %137
-OpStore %138 %93
+%185 = OpAccessChain  %28  %115 %52 %183 %184
+OpStore %185 %112
 OpReturn
 OpFunctionEnd
-%140 = OpFunction  %2  None %67
-%139 = OpLabel
-%150 = OpVariable  %95  Function %70
-%151 = OpVariable  %152  Function %144
-%142 = OpAccessChain  %141  %62 %48
-OpBranch %153
-%153 = OpLabel
+%190 = OpFunction  %25  None %191
+%192 = OpFunctionParameter  %49
+%193 = OpLabel
+%194 = OpCompositeExtract  %13  %192 0
+%195 = OpCompositeExtract  %13  %192 1
+%196 = OpCompositeExtract  %13  %192 2
+%197 = OpCompositeExtract  %13  %192 3
+%198 = OpCompositeConstruct  %25  %194 %195 %196 %197
+OpReturnValue %198
+OpFunctionEnd
+%186 = OpFunction  %26  None %187
+%188 = OpFunctionParameter  %50
+%189 = OpLabel
+%199 = OpCompositeExtract  %49  %188 0
+%200 = OpFunctionCall  %25  %190 %199
+%201 = OpCompositeExtract  %49  %188 1
+%202 = OpFunctionCall  %25  %190 %201
+%203 = OpCompositeConstruct  %26  %200 %202
+OpReturnValue %203
+OpFunctionEnd
+%204 = OpFunction  %13  None %207
+%205 = OpFunctionParameter  %25
+%206 = OpFunctionParameter  %4
+%208 = OpLabel
+OpSelectionMerge %209 None
+OpSwitch %206 %214 0 %210 1 %211 2 %212 3 %213
+%210 = OpLabel
+%215 = OpCompositeExtract  %13  %205 0
+OpBranch %209
+%211 = OpLabel
+%216 = OpCompositeExtract  %13  %205 1
+OpBranch %209
+%212 = OpLabel
+%217 = OpCompositeExtract  %13  %205 2
+OpBranch %209
+%213 = OpLabel
+%218 = OpCompositeExtract  %13  %205 3
+OpBranch %209
+%214 = OpLabel
+OpUnreachable
+%209 = OpLabel
+%219 = OpPhi  %13  %215 %210 %216 %211 %217 %212 %218 %213
+OpReturnValue %219
+OpFunctionEnd
+%221 = OpFunction  %2  None %85
+%220 = OpLabel
+%232 = OpVariable  %114  Function %88
+%233 = OpVariable  %234  Function %226
+%223 = OpAccessChain  %222  %66 %52
+OpBranch %235
+%235 = OpLabel
 OpLine %3 78 5
-%154 = OpLoad  %6  %150
-%155 = OpISub  %6  %154 %70
+%236 = OpLoad  %6  %232
+%237 = OpISub  %6  %236 %88
 OpLine %3 78 5
-OpStore %150 %155
+OpStore %232 %237
 OpLine %3 81 14
-%157 = OpAccessChain  %156  %142 %48
-%158 = OpLoad  %26  %157
+%240 = OpAccessChain  %239  %223 %52
+%241 = OpLoad  %50  %240
+%242 = OpFunctionCall  %26  %186 %241
 OpLine %3 82 14
 OpLine %3 82 14
-%160 = OpAccessChain  %159  %142 %48 %48
-%161 = OpLoad  %25  %160
+%245 = OpAccessChain  %244  %223 %52 %52
+%246 = OpLoad  %49  %245
+%247 = OpFunctionCall  %25  %190 %246
 OpLine %3 83 14
 OpLine %3 83 14
 OpLine %3 83 14
-%162 = OpAccessChain  %104  %142 %48 %48 %48
-%163 = OpLoad  %13  %162
+%248 = OpAccessChain  %121  %223 %52 %52 %52
+%249 = OpLoad  %13  %248
 OpLine %3 84 14
 OpLine %3 84 14
-%164 = OpLoad  %6  %150
-%165 = OpAccessChain  %104  %142 %48 %48 %164
-%166 = OpLoad  %13  %165
+%250 = OpLoad  %6  %232
+%251 = OpAccessChain  %244  %223 %52 %52
+%252 = OpLoad  %49  %251
+%253 = OpFunctionCall  %25  %190 %252
+%254 = OpBitcast  %4  %250
+%255 = OpFunctionCall  %13  %204 %253 %254
 OpLine %3 85 14
 OpLine %3 85 14
 OpLine %3 85 14
 OpLine %3 85 14
-%167 = OpAccessChain  %110  %142 %48 %48 %48 %44
-%168 = OpLoad  %9  %167
+%256 = OpAccessChain  %141  %223 %52 %52 %52 %44
+%257 = OpLoad  %9  %256
 OpLine %3 86 14
 OpLine %3 86 14
 OpLine %3 86 14
-%169 = OpLoad  %6  %150
-%170 = OpAccessChain  %110  %142 %48 %48 %48 %169
-%171 = OpLoad  %9  %170
+%258 = OpLoad  %6  %232
+%259 = OpAccessChain  %141  %223 %52 %52 %52 %258
+%260 = OpLoad  %9  %259
 OpLine %3 87 14
 OpLine %3 87 14
-%172 = OpLoad  %6  %150
+%261 = OpLoad  %6  %232
 OpLine %3 87 14
-%173 = OpAccessChain  %110  %142 %48 %48 %172 %44
-%174 = OpLoad  %9  %173
+%262 = OpAccessChain  %244  %223 %52 %52
+%263 = OpLoad  %49  %262
+%264 = OpFunctionCall  %25  %190 %263
+%265 = OpBitcast  %4  %261
+%266 = OpFunctionCall  %13  %204 %264 %265
+%267 = OpCompositeExtract  %9  %266 1
 OpLine %3 88 14
 OpLine %3 88 14
-%175 = OpLoad  %6  %150
-%176 = OpLoad  %6  %150
-%177 = OpAccessChain  %110  %142 %48 %48 %175 %176
-%178 = OpLoad  %9  %177
+%268 = OpLoad  %6  %232
+%269 = OpLoad  %6  %232
+%270 = OpAccessChain  %244  %223 %52 %52
+%271 = OpLoad  %49  %270
+%272 = OpFunctionCall  %25  %190 %271
+%273 = OpBitcast  %4  %268
+%274 = OpFunctionCall  %13  %204 %272 %273
+%275 = OpVectorExtractDynamic  %9  %274 %269
 OpLine %3 90 13
 OpLine %3 92 5
-%179 = OpLoad  %6  %150
-%180 = OpIAdd  %6  %179 %70
+%276 = OpLoad  %6  %232
+%277 = OpIAdd  %6  %276 %88
 OpLine %3 92 5
-OpStore %150 %180
+OpStore %232 %277
 OpLine %3 95 5
 OpLine %3 95 5
-%182 = OpAccessChain  %181  %151 %48
-OpStore %182 %143
+%279 = OpAccessChain  %278  %233 %52
+OpStore %279 %225
 OpLine %3 96 5
 OpLine %3 96 5
 OpLine %3 96 27
@@ -738,359 +881,359 @@ OpLine %3 96 43
 OpLine %3 96 59
 OpLine %3 96 15
 OpLine %3 96 5
-%184 = OpAccessChain  %183  %151 %48 %48
-OpStore %184 %149
+%281 = OpAccessChain  %280  %233 %52 %52
+OpStore %281 %231
 OpLine %3 97 5
 OpLine %3 97 5
 OpLine %3 97 5
 OpLine %3 97 18
 OpLine %3 97 5
-%185 = OpAccessChain  %127  %151 %48 %48 %48
-OpStore %185 %87
+%282 = OpAccessChain  %174  %233 %52 %52 %52
+OpStore %282 %106
 OpLine %3 98 5
 OpLine %3 98 5
-%186 = OpLoad  %6  %150
+%283 = OpLoad  %6  %232
 OpLine %3 98 20
 OpLine %3 98 5
-%187 = OpAccessChain  %127  %151 %48 %48 %186
-OpStore %187 %89
+%284 = OpAccessChain  %174  %233 %52 %52 %283
+OpStore %284 %108
 OpLine %3 99 5
 OpLine %3 99 5
 OpLine %3 99 5
 OpLine %3 99 5
 OpLine %3 99 5
-%188 = OpAccessChain  %28  %151 %48 %48 %48 %44
-OpStore %188 %90
+%285 = OpAccessChain  %28  %233 %52 %52 %52 %44
+OpStore %285 %109
 OpLine %3 100 5
 OpLine %3 100 5
 OpLine %3 100 5
-%189 = OpLoad  %6  %150
+%286 = OpLoad  %6  %232
 OpLine %3 100 5
-%190 = OpAccessChain  %28  %151 %48 %48 %48 %189
-OpStore %190 %91
+%287 = OpAccessChain  %28  %233 %52 %52 %52 %286
+OpStore %287 %110
 OpLine %3 101 5
 OpLine %3 101 5
-%191 = OpLoad  %6  %150
+%288 = OpLoad  %6  %232
 OpLine %3 101 5
 OpLine %3 101 5
-%192 = OpAccessChain  %28  %151 %48 %48 %191 %44
-OpStore %192 %92
+%289 = OpAccessChain  %28  %233 %52 %52 %288 %44
+OpStore %289 %111
 OpLine %3 102 5
 OpLine %3 102 5
-%193 = OpLoad  %6  %150
-%194 = OpLoad  %6  %150
+%290 = OpLoad  %6  %232
+%291 = OpLoad  %6  %232
 OpLine %3 102 5
-%195 = OpAccessChain  %28  %151 %48 %48 %193 %194
-OpStore %195 %93
+%292 = OpAccessChain  %28  %233 %52 %52 %290 %291
+OpStore %292 %112
 OpReturn
 OpFunctionEnd
-%198 = OpFunction  %9  None %199
-%197 = OpFunctionParameter  %28
-%196 = OpLabel
-OpBranch %200
-%200 = OpLabel
+%295 = OpFunction  %9  None %296
+%294 = OpFunctionParameter  %28
+%293 = OpLabel
+OpBranch %297
+%297 = OpLabel
 OpLine %3 105 22
-%201 = OpLoad  %9  %197
-OpReturnValue %201
+%298 = OpLoad  %9  %294
+OpReturnValue %298
 OpFunctionEnd
-%204 = OpFunction  %9  None %205
-%203 = OpFunctionParameter  %30
-%202 = OpLabel
-OpBranch %206
-%206 = OpLabel
+%301 = OpFunction  %9  None %302
+%300 = OpFunctionParameter  %30
+%299 = OpLabel
+OpBranch %303
+%303 = OpLabel
 OpLine %3 110 12
-%207 = OpCompositeExtract  %29  %203 4
+%304 = OpCompositeExtract  %29  %300 4
 OpLine %3 110 12
-%208 = OpCompositeExtract  %9  %207 9
-OpReturnValue %208
+%305 = OpCompositeExtract  %9  %304 9
+OpReturnValue %305
 OpFunctionEnd
-%211 = OpFunction  %2  None %212
-%210 = OpFunctionParameter  %34
-%209 = OpLabel
-OpBranch %214
-%214 = OpLabel
+%308 = OpFunction  %2  None %309
+%307 = OpFunctionParameter  %34
+%306 = OpLabel
+OpBranch %311
+%311 = OpLabel
 OpLine %3 159 5
-OpStore %210 %213
+OpStore %307 %310
 OpReturn
 OpFunctionEnd
-%217 = OpFunction  %2  None %218
-%216 = OpFunctionParameter  %36
-%215 = OpLabel
-OpBranch %222
-%222 = OpLabel
+%314 = OpFunction  %2  None %315
+%313 = OpFunctionParameter  %36
+%312 = OpLabel
+OpBranch %319
+%319 = OpLabel
 OpLine %3 163 32
 OpLine %3 163 43
 OpLine %3 163 32
 OpLine %3 163 12
 OpLine %3 163 5
-OpStore %216 %221
+OpStore %313 %318
 OpReturn
 OpFunctionEnd
-%224 = OpFunction  %2  None %67
-%223 = OpLabel
-%229 = OpVariable  %34  Function %225
-%230 = OpVariable  %36  Function %228
-OpBranch %231
-%231 = OpLabel
+%321 = OpFunction  %2  None %85
+%320 = OpLabel
+%326 = OpVariable  %34  Function %322
+%327 = OpVariable  %36  Function %325
+OpBranch %328
+%328 = OpLabel
 OpLine %3 168 5
-%232 = OpFunctionCall  %2  %211 %229
+%329 = OpFunctionCall  %2  %308 %326
 OpLine %3 170 35
 OpLine %3 170 46
 OpLine %3 170 35
 OpLine %3 170 15
 OpLine %3 171 5
-%233 = OpFunctionCall  %2  %217 %230
+%330 = OpFunctionCall  %2  %314 %327
 OpReturn
 OpFunctionEnd
-%236 = OpFunction  %4  None %237
-%235 = OpFunctionParameter  %38
-%234 = OpLabel
-OpBranch %238
-%238 = OpLabel
+%333 = OpFunction  %4  None %334
+%332 = OpFunctionParameter  %38
+%331 = OpLabel
+OpBranch %335
+%335 = OpLabel
 OpLine %3 179 10
-%239 = OpAccessChain  %34  %235 %48
-%240 = OpLoad  %4  %239
-OpReturnValue %240
+%336 = OpAccessChain  %34  %332 %52
+%337 = OpLoad  %4  %336
+OpReturnValue %337
 OpFunctionEnd
-%243 = OpFunction  %2  None %244
-%242 = OpFunctionParameter  %38
-%241 = OpLabel
-OpBranch %245
-%245 = OpLabel
+%340 = OpFunction  %2  None %341
+%339 = OpFunctionParameter  %38
+%338 = OpLabel
+OpBranch %342
+%342 = OpLabel
 OpLine %3 183 3
 OpLine %3 183 3
-%246 = OpAccessChain  %34  %242 %48
-OpStore %246 %17
+%343 = OpAccessChain  %34  %339 %52
+OpStore %343 %17
 OpReturn
 OpFunctionEnd
-%249 = OpFunction  %4  None %250
-%248 = OpFunctionParameter  %41
-%247 = OpLabel
-OpBranch %251
-%251 = OpLabel
+%346 = OpFunction  %4  None %347
+%345 = OpFunctionParameter  %41
+%344 = OpLabel
+OpBranch %348
+%348 = OpLabel
 OpLine %3 187 10
-%252 = OpAccessChain  %34  %248 %44
-%253 = OpLoad  %4  %252
-OpReturnValue %253
+%349 = OpAccessChain  %34  %345 %44
+%350 = OpLoad  %4  %349
+OpReturnValue %350
 OpFunctionEnd
-%256 = OpFunction  %2  None %257
-%255 = OpFunctionParameter  %41
-%254 = OpLabel
-OpBranch %258
-%258 = OpLabel
+%353 = OpFunction  %2  None %354
+%352 = OpFunctionParameter  %41
+%351 = OpLabel
+OpBranch %355
+%355 = OpLabel
 OpLine %3 191 3
 OpLine %3 191 3
-%259 = OpAccessChain  %34  %255 %44
-OpStore %259 %17
+%356 = OpAccessChain  %34  %352 %44
+OpStore %356 %17
 OpReturn
 OpFunctionEnd
-%261 = OpFunction  %2  None %67
-%260 = OpLabel
-%262 = OpVariable  %38  Function %263
-%264 = OpVariable  %41  Function %265
-OpBranch %266
-%266 = OpLabel
+%358 = OpFunction  %2  None %85
+%357 = OpLabel
+%359 = OpVariable  %38  Function %360
+%361 = OpVariable  %41  Function %362
+OpBranch %363
+%363 = OpLabel
 OpLine %3 196 4
-%267 = OpFunctionCall  %2  %243 %262
+%364 = OpFunctionCall  %2  %340 %359
 OpLine %3 197 4
-%268 = OpFunctionCall  %4  %236 %262
+%365 = OpFunctionCall  %4  %333 %359
 OpLine %3 200 4
-%269 = OpFunctionCall  %2  %256 %264
+%366 = OpFunctionCall  %2  %353 %361
 OpLine %3 201 4
-%270 = OpFunctionCall  %4  %249 %264
+%367 = OpFunctionCall  %4  %346 %361
 OpReturn
 OpFunctionEnd
-%273 = OpFunction  %42  None %274
-%272 = OpFunctionParameter  %42
-%271 = OpLabel
-%275 = OpVariable  %276  Function %277
-OpBranch %278
-%278 = OpLabel
+%370 = OpFunction  %42  None %371
+%369 = OpFunctionParameter  %42
+%368 = OpLabel
+%372 = OpVariable  %373  Function %374
+OpBranch %375
+%375 = OpLabel
 OpLine %3 205 13
-%279 = OpCompositeConstruct  %43  %272
+%376 = OpCompositeConstruct  %43  %369
 OpLine %3 205 5
-OpStore %275 %279
+OpStore %372 %376
 OpLine %3 207 12
-%281 = OpAccessChain  %280  %275 %48
-%282 = OpLoad  %42  %281
-OpReturnValue %282
+%378 = OpAccessChain  %377  %372 %52
+%379 = OpLoad  %42  %378
+OpReturnValue %379
 OpFunctionEnd
-%284 = OpFunction  %6  None %285
-%283 = OpLabel
-%288 = OpVariable  %289  Function %287
-OpBranch %290
-%290 = OpLabel
+%381 = OpFunction  %6  None %382
+%380 = OpLabel
+%385 = OpVariable  %386  Function %384
+OpBranch %387
+%387 = OpLabel
 OpLine %3 213 16
 OpLine %3 215 12
-%291 = OpAccessChain  %95  %288 %48
-%292 = OpLoad  %6  %291
-OpReturnValue %292
+%388 = OpAccessChain  %114  %385 %52
+%389 = OpLoad  %6  %388
+OpReturnValue %389
 OpFunctionEnd
-%294 = OpFunction  %6  None %285
-%293 = OpLabel
-OpBranch %296
-%296 = OpLabel
+%391 = OpFunction  %6  None %382
+%390 = OpLabel
+OpBranch %393
+%393 = OpLabel
 OpLine %3 225 17
-%297 = OpCompositeExtract  %46  %295 0
+%394 = OpCompositeExtract  %46  %392 0
 OpLine %3 226 20
-%298 = OpCompositeExtract  %6  %297 0
+%395 = OpCompositeExtract  %6  %394 0
 OpLine %3 228 9
-%299 = OpCompositeExtract  %4  %295 1
-%300 = OpBitcast  %4  %298
-%301 = OpINotEqual  %42  %299 %300
+%396 = OpCompositeExtract  %4  %392 1
+%397 = OpBitcast  %4  %395
+%398 = OpINotEqual  %42  %396 %397
 OpLine %3 228 5
 OpLine %3 232 12
-%302 = OpCompositeExtract  %46  %295 0
-%303 = OpCompositeExtract  %6  %302 0
-OpReturnValue %303
+%399 = OpCompositeExtract  %46  %392 0
+%400 = OpCompositeExtract  %6  %399 0
+OpReturnValue %400
 OpFunctionEnd
-%305 = OpFunction  %6  None %285
-%304 = OpLabel
-%306 = OpVariable  %307  Function %295
-%308 = OpVariable  %309  Function %310
-%311 = OpVariable  %95  Function %312
-OpBranch %313
-%313 = OpLabel
+%402 = OpFunction  %6  None %382
+%401 = OpLabel
+%403 = OpVariable  %404  Function %392
+%405 = OpVariable  %406  Function %407
+%408 = OpVariable  %114  Function %409
+OpBranch %410
+%410 = OpLabel
 OpLine %3 238 17
-%314 = OpAccessChain  %309  %306 %48
-%315 = OpLoad  %46  %314
+%411 = OpAccessChain  %406  %403 %52
+%412 = OpLoad  %46  %411
 OpLine %3 238 5
-OpStore %308 %315
+OpStore %405 %412
 OpLine %3 239 20
-%316 = OpAccessChain  %95  %308 %48
-%317 = OpLoad  %6  %316
+%413 = OpAccessChain  %114  %405 %52
+%414 = OpLoad  %6  %413
 OpLine %3 239 5
-OpStore %311 %317
+OpStore %408 %414
 OpLine %3 241 9
-%318 = OpAccessChain  %34  %306 %44
-%319 = OpLoad  %4  %318
-%320 = OpLoad  %6  %311
-%321 = OpBitcast  %4  %320
-%322 = OpINotEqual  %42  %319 %321
+%415 = OpAccessChain  %34  %403 %44
+%416 = OpLoad  %4  %415
+%417 = OpLoad  %6  %408
+%418 = OpBitcast  %4  %417
+%419 = OpINotEqual  %42  %416 %418
 OpLine %3 241 5
 OpLine %3 245 12
-%323 = OpAccessChain  %95  %306 %48 %48
-%324 = OpLoad  %6  %323
-OpReturnValue %324
+%420 = OpAccessChain  %114  %403 %52 %52
+%421 = OpLoad  %6  %420
+OpReturnValue %421
 OpFunctionEnd
-%331 = OpFunction  %2  None %67
-%325 = OpLabel
-%342 = OpVariable  %28  Function %336
-%343 = OpVariable  %344  Function %345
-%328 = OpLoad  %4  %326
-%332 = OpAccessChain  %68  %56 %48
-%334 = OpAccessChain  %333  %59 %48
-%335 = OpAccessChain  %141  %62 %48
-OpBranch %346
-%346 = OpLabel
+%428 = OpFunction  %2  None %85
+%422 = OpLabel
+%439 = OpVariable  %28  Function %433
+%440 = OpVariable  %441  Function %442
+%425 = OpLoad  %4  %423
+%429 = OpAccessChain  %86  %60 %52
+%431 = OpAccessChain  %430  %63 %52
+%432 = OpAccessChain  %222  %66 %52
+OpBranch %443
+%443 = OpLabel
 OpLine %3 1 1
-%347 = OpLoad  %9  %342
+%444 = OpLoad  %9  %439
 OpLine %3 118 5
-OpStore %342 %71
+OpStore %439 %90
 OpLine %3 120 9
-%348 = OpLoad  %7  %52
+%445 = OpLoad  %7  %56
 OpLine %3 121 5
-%349 = OpFunctionCall  %2  %66
+%446 = OpFunctionCall  %2  %84
 OpLine %3 122 5
-%350 = OpFunctionCall  %2  %140
+%447 = OpFunctionCall  %2  %221
 OpLine %3 125 19
-%352 = OpAccessChain  %351  %54 %48
-%353 = OpLoad  %10  %352
+%449 = OpAccessChain  %448  %58 %52
+%450 = OpLoad  %10  %449
 OpLine %3 126 15
-%355 = OpAccessChain  %354  %54 %40
-%356 = OpLoad  %19  %355
+%452 = OpAccessChain  %451  %58 %40
+%453 = OpLoad  %19  %452
 OpLine %3 128 13
-%359 = OpAccessChain  %358  %54 %48 %337 %48
-%360 = OpLoad  %9  %359
+%456 = OpAccessChain  %455  %58 %52 %434 %52
+%457 = OpLoad  %9  %456
 OpLine %3 129 13
 OpLine %3 129 22
-%362 = OpArrayLength  %4  %54 5
+%459 = OpArrayLength  %4  %58 5
 OpLine %3 129 13
-%363 = OpISub  %4  %362 %15
-%366 = OpAccessChain  %365  %54 %31 %363 %48
-%367 = OpLoad  %6  %366
+%460 = OpISub  %4  %459 %15
+%463 = OpAccessChain  %462  %58 %31 %460 %52
+%464 = OpLoad  %6  %463
 OpLine %3 130 13
-%368 = OpLoad  %24  %334
+%465 = OpLoad  %24  %431
 OpLine %3 133 56
 OpLine %3 133 56
 OpLine %3 134 21
-%369 = OpFunctionCall  %9  %198 %342
+%466 = OpFunctionCall  %9  %295 %439
 OpLine %3 137 31
-%372 = OpExtInst  %9  %1 FClamp %360 %370 %371
-%373 = OpConvertFToS  %6  %372
+%469 = OpExtInst  %9  %1 FClamp %457 %467 %468
+%470 = OpConvertFToS  %6  %469
 OpLine %3 137 14
-%374 = OpCompositeConstruct  %33  %367 %373 %338 %339 %340
+%471 = OpCompositeConstruct  %33  %464 %470 %435 %436 %437
 OpLine %3 137 5
-OpStore %343 %374
+OpStore %440 %471
 OpLine %3 138 5
-%375 = OpIAdd  %4  %328 %44
+%472 = OpIAdd  %4  %425 %44
 OpLine %3 138 5
-%376 = OpAccessChain  %95  %343 %375
-OpStore %376 %286
+%473 = OpAccessChain  %114  %440 %472
+OpStore %473 %383
 OpLine %3 139 17
-%377 = OpAccessChain  %95  %343 %328
-%378 = OpLoad  %6  %377
+%474 = OpAccessChain  %114  %440 %425
+%475 = OpLoad  %6  %474
 OpLine %3 141 5
-%379 = OpFunctionCall  %9  %204 %341
+%476 = OpFunctionCall  %9  %301 %438
 OpLine %3 143 22
-%381 = OpCompositeConstruct  %380  %378 %378 %378 %378
-%382 = OpConvertSToF  %32  %381
-%383 = OpMatrixTimesVector  %11  %353 %382
+%478 = OpCompositeConstruct  %477  %475 %475 %475 %475
+%479 = OpConvertSToF  %32  %478
+%480 = OpMatrixTimesVector  %11  %450 %479
 OpLine %3 143 12
-%384 = OpCompositeConstruct  %32  %383 %73
-OpStore %329 %384
+%481 = OpCompositeConstruct  %32  %480 %92
+OpStore %426 %481
 OpReturn
 OpFunctionEnd
-%387 = OpFunction  %2  None %67
-%385 = OpLabel
-%388 = OpAccessChain  %333  %59 %48
-OpBranch %399
-%399 = OpLabel
+%484 = OpFunction  %2  None %85
+%482 = OpLabel
+%485 = OpAccessChain  %430  %63 %52
+OpBranch %496
+%496 = OpLabel
 OpLine %3 149 5
 OpLine %3 149 5
 OpLine %3 149 5
-%400 = OpAccessChain  %358  %54 %48 %44 %15
-OpStore %400 %71
+%497 = OpAccessChain  %455  %58 %52 %44 %15
+OpStore %497 %90
 OpLine %3 150 5
 OpLine %3 150 31
 OpLine %3 150 47
 OpLine %3 150 63
 OpLine %3 150 19
 OpLine %3 150 5
-%401 = OpAccessChain  %351  %54 %48
-OpStore %401 %393
+%498 = OpAccessChain  %448  %58 %52
+OpStore %498 %490
 OpLine %3 151 5
 OpLine %3 151 35
 OpLine %3 151 15
 OpLine %3 151 5
-%402 = OpAccessChain  %354  %54 %40
-OpStore %402 %396
+%499 = OpAccessChain  %451  %58 %40
+OpStore %499 %493
 OpLine %3 152 5
 OpLine %3 152 5
 OpLine %3 152 5
-%403 = OpAccessChain  %365  %54 %31 %44 %48
-OpStore %403 %70
+%500 = OpAccessChain  %462  %58 %31 %44 %52
+OpStore %500 %88
 OpLine %3 153 5
-OpStore %388 %397
+OpStore %485 %494
 OpLine %3 155 12
-OpStore %386 %398
+OpStore %483 %495
 OpReturn
 OpFunctionEnd
-%405 = OpFunction  %2  None %67
-%404 = OpLabel
-OpBranch %407
-%407 = OpLabel
+%502 = OpFunction  %2  None %85
+%501 = OpLabel
+OpBranch %504
+%504 = OpLabel
 OpLine %3 250 5
-%408 = OpFunctionCall  %2  %224
+%505 = OpFunctionCall  %2  %321
 OpLine %3 251 5
-%409 = OpFunctionCall  %2  %261
+%506 = OpFunctionCall  %2  %358
 OpLine %3 252 5
-%410 = OpFunctionCall  %42  %273 %406
+%507 = OpFunctionCall  %42  %370 %503
 OpLine %3 253 5
-%411 = OpFunctionCall  %6  %284
+%508 = OpFunctionCall  %6  %381
 OpLine %3 254 5
-%412 = OpFunctionCall  %6  %294
+%509 = OpFunctionCall  %6  %391
 OpLine %3 255 5
-%413 = OpFunctionCall  %6  %305
+%510 = OpFunctionCall  %6  %402
 OpReturn
 OpFunctionEnd
\ No newline at end of file
diff --git a/naga/tests/out/spv/wgsl-f16.spvasm b/naga/tests/out/spv/wgsl-f16.spvasm
index 4fc6963ce87..b280637aa2b 100644
--- a/naga/tests/out/spv/wgsl-f16.spvasm
+++ b/naga/tests/out/spv/wgsl-f16.spvasm
@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.0
 ; Generator: rspirv
-; Bound: 530
+; Bound: 568
 OpCapability Shader
 OpCapability Float16
 OpCapability StorageBuffer16BitAccess
@@ -11,8 +11,8 @@ OpExtension "SPV_KHR_storage_buffer_storage_class"
 OpExtension "SPV_KHR_16bit_storage"
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
-OpEntryPoint GLCompute %518 "main"
-OpExecutionMode %518 LocalSize 1 1 1
+OpEntryPoint GLCompute %556 "main"
+OpExecutionMode %556 LocalSize 1 1 1
 OpMemberDecorate %19 0 Offset 0
 OpMemberDecorate %19 1 Offset 4
 OpMemberDecorate %19 2 Offset 8
@@ -56,28 +56,63 @@ OpMemberDecorate %23 2 Offset 8
 OpMemberDecorate %23 3 Offset 14
 OpMemberDecorate %23 4 Offset 16
 OpMemberDecorate %23 5 Offset 20
-OpDecorate %28 DescriptorSet 0
-OpDecorate %28 Binding 0
-OpDecorate %29 Block
-OpMemberDecorate %29 0 Offset 0
-OpDecorate %31 NonWritable
-OpDecorate %31 DescriptorSet 0
-OpDecorate %31 Binding 1
-OpDecorate %32 Block
-OpMemberDecorate %32 0 Offset 0
-OpDecorate %34 NonWritable
-OpDecorate %34 DescriptorSet 0
-OpDecorate %34 Binding 2
-OpDecorate %35 Block
-OpMemberDecorate %35 0 Offset 0
-OpDecorate %37 DescriptorSet 0
-OpDecorate %37 Binding 3
-OpDecorate %38 Block
-OpMemberDecorate %38 0 Offset 0
-OpDecorate %40 DescriptorSet 0
-OpDecorate %40 Binding 4
-OpDecorate %41 Block
-OpMemberDecorate %41 0 Offset 0
+OpMemberDecorate %24 0 Offset 0
+OpMemberDecorate %24 1 Offset 4
+OpMemberDecorate %24 2 Offset 8
+OpMemberDecorate %24 3 Offset 12
+OpMemberDecorate %24 4 Offset 16
+OpMemberDecorate %24 5 Offset 24
+OpMemberDecorate %24 6 Offset 32
+OpMemberDecorate %24 7 Offset 40
+OpMemberDecorate %24 8 Offset 44
+OpMemberDecorate %24 9 Offset 48
+OpMemberDecorate %24 10 Offset 56
+OpMemberDecorate %24 10 ColMajor
+OpMemberDecorate %24 10 MatrixStride 8
+OpMemberDecorate %24 11 Offset 72
+OpMemberDecorate %24 11 ColMajor
+OpMemberDecorate %24 11 MatrixStride 8
+OpMemberDecorate %24 12 Offset 88
+OpMemberDecorate %24 13 Offset 92
+OpMemberDecorate %24 14 Offset 96
+OpMemberDecorate %24 15 Offset 104
+OpMemberDecorate %24 15 ColMajor
+OpMemberDecorate %24 15 MatrixStride 8
+OpMemberDecorate %24 16 Offset 128
+OpMemberDecorate %24 16 ColMajor
+OpMemberDecorate %24 16 MatrixStride 8
+OpMemberDecorate %24 17 Offset 152
+OpMemberDecorate %24 18 Offset 156
+OpMemberDecorate %24 19 Offset 160
+OpMemberDecorate %24 20 Offset 164
+OpMemberDecorate %24 21 Offset 168
+OpMemberDecorate %24 21 ColMajor
+OpMemberDecorate %24 21 MatrixStride 8
+OpMemberDecorate %24 22 Offset 200
+OpMemberDecorate %24 22 ColMajor
+OpMemberDecorate %24 22 MatrixStride 8
+OpDecorate %29 DescriptorSet 0
+OpDecorate %29 Binding 0
+OpDecorate %30 Block
+OpMemberDecorate %30 0 Offset 0
+OpDecorate %32 NonWritable
+OpDecorate %32 DescriptorSet 0
+OpDecorate %32 Binding 1
+OpDecorate %33 Block
+OpMemberDecorate %33 0 Offset 0
+OpDecorate %35 NonWritable
+OpDecorate %35 DescriptorSet 0
+OpDecorate %35 Binding 2
+OpDecorate %36 Block
+OpMemberDecorate %36 0 Offset 0
+OpDecorate %38 DescriptorSet 0
+OpDecorate %38 Binding 3
+OpDecorate %39 Block
+OpMemberDecorate %39 0 Offset 0
+OpDecorate %41 DescriptorSet 0
+OpDecorate %41 Binding 4
+OpDecorate %42 Block
+OpMemberDecorate %42 0 Offset 0
 %2 = OpTypeVoid
 %3 = OpTypeFloat 16
 %4 = OpTypeInt 32 0
@@ -100,559 +135,597 @@ OpMemberDecorate %41 0 Offset 0
 %20 = OpTypeArray %3 %21
 %22 = OpTypeStruct %20
 %23 = OpTypeStruct %3 %3 %8 %3 %3 %4
-%24 = OpConstant  %3  0.000000000000000000000000000000000000000021524
-%25 = OpConstant  %3  0.000000000000000000000000000000000000000027121
-%27 = OpTypePointer Private %3
-%26 = OpVariable  %27  Private %24
-%29 = OpTypeStruct %19
-%30 = OpTypePointer Uniform %29
-%28 = OpVariable  %30  Uniform
-%32 = OpTypeStruct %19
-%33 = OpTypePointer StorageBuffer %32
-%31 = OpVariable  %33  StorageBuffer
-%35 = OpTypeStruct %22
-%36 = OpTypePointer StorageBuffer %35
-%34 = OpVariable  %36  StorageBuffer
-%38 = OpTypeStruct %19
-%39 = OpTypePointer StorageBuffer %38
-%37 = OpVariable  %39  StorageBuffer
-%41 = OpTypeStruct %22
-%42 = OpTypePointer StorageBuffer %41
-%40 = OpVariable  %42  StorageBuffer
-%46 = OpTypeFunction %3 %3
-%47 = OpTypePointer Uniform %19
-%48 = OpConstant  %4  0
-%50 = OpTypePointer StorageBuffer %19
-%52 = OpTypePointer StorageBuffer %22
-%56 = OpConstant  %3  0.000000000000000000000000000000000000000088991
-%57 = OpConstant  %3  0.000000000000000000000000000000000000000024753
-%58 = OpConstant  %5  65504
-%59 = OpConstant  %5  -65504
-%60 = OpConstant  %4  65504
-%61 = OpConstant  %6  65504
-%62 = OpConstant  %6  -65504
-%64 = OpTypePointer Function %23
-%65 = OpConstantNull  %23
-%67 = OpTypePointer Function %3
-%76 = OpTypePointer Uniform %6
-%85 = OpTypePointer Uniform %3
-%86 = OpConstant  %4  3
-%93 = OpTypePointer StorageBuffer %5
-%94 = OpConstant  %4  1
-%97 = OpTypePointer StorageBuffer %4
-%100 = OpTypePointer StorageBuffer %6
-%103 = OpTypePointer StorageBuffer %3
-%110 = OpTypePointer StorageBuffer %7
-%111 = OpTypePointer Uniform %7
-%112 = OpConstant  %4  4
-%119 = OpTypePointer StorageBuffer %8
-%120 = OpTypePointer Uniform %8
-%121 = OpConstant  %4  5
-%128 = OpTypePointer StorageBuffer %9
-%129 = OpTypePointer Uniform %9
-%130 = OpConstant  %4  6
-%137 = OpTypePointer StorageBuffer %10
-%138 = OpTypePointer Uniform %10
-%139 = OpConstant  %4  8
-%152 = OpTypePointer StorageBuffer %11
-%153 = OpTypePointer Uniform %11
-%154 = OpConstant  %4  9
-%167 = OpTypePointer StorageBuffer %12
-%168 = OpTypePointer Uniform %12
-%169 = OpConstant  %4  10
-%182 = OpTypePointer StorageBuffer %13
-%183 = OpTypePointer Uniform %13
-%184 = OpConstant  %4  11
-%200 = OpTypePointer StorageBuffer %14
-%201 = OpTypePointer Uniform %14
-%202 = OpConstant  %4  12
-%218 = OpTypePointer StorageBuffer %15
-%219 = OpTypePointer Uniform %15
-%220 = OpConstant  %4  13
-%236 = OpTypePointer StorageBuffer %16
-%237 = OpTypePointer Uniform %16
-%238 = OpConstant  %4  14
-%257 = OpTypePointer StorageBuffer %17
-%258 = OpTypePointer Uniform %17
-%259 = OpConstant  %4  15
-%278 = OpTypePointer StorageBuffer %18
-%279 = OpTypePointer Uniform %18
-%280 = OpConstant  %4  16
-%299 = OpTypePointer StorageBuffer %20
-%338 = OpTypeVector %6 2
-%344 = OpTypeVector %6 3
-%350 = OpTypeVector %6 4
-%356 = OpTypeMatrix %338 2
-%370 = OpTypeMatrix %344 2
-%384 = OpTypeMatrix %350 2
-%398 = OpTypeMatrix %338 3
-%416 = OpTypeMatrix %344 3
-%434 = OpTypeMatrix %350 3
-%452 = OpTypeMatrix %338 4
-%474 = OpTypeMatrix %344 4
-%496 = OpTypeMatrix %350 4
-%519 = OpTypeFunction %2
-%525 = OpConstant  %3  0.000000000000000000000000000000000000000022959
-%528 = OpConstant  %4  7
-%45 = OpFunction  %3  None %46
-%44 = OpFunctionParameter  %3
-%43 = OpLabel
-%63 = OpVariable  %64  Function %65
-%66 = OpVariable  %67  Function %25
-%49 = OpAccessChain  %47  %28 %48
-%51 = OpAccessChain  %50  %31 %48
-%53 = OpAccessChain  %52  %34 %48
-%54 = OpAccessChain  %50  %37 %48
-%55 = OpAccessChain  %52  %40 %48
-OpBranch %68
-%68 = OpLabel
-%69 = OpLoad  %3  %26
-%70 = OpLoad  %3  %66
-%71 = OpFAdd  %3  %70 %56
-OpStore %66 %71
-%72 = OpLoad  %3  %66
+%24 = OpTypeStruct %4 %5 %6 %3 %7 %8 %9 %3 %7 %7 %11 %12 %7 %7 %7 %14 %15 %7 %7 %7 %7 %17 %18
+%25 = OpConstant  %3  0.000000000000000000000000000000000000000021524
+%26 = OpConstant  %3  0.000000000000000000000000000000000000000027121
+%28 = OpTypePointer Private %3
+%27 = OpVariable  %28  Private %25
+%30 = OpTypeStruct %24
+%31 = OpTypePointer Uniform %30
+%29 = OpVariable  %31  Uniform
+%33 = OpTypeStruct %19
+%34 = OpTypePointer StorageBuffer %33
+%32 = OpVariable  %34  StorageBuffer
+%36 = OpTypeStruct %22
+%37 = OpTypePointer StorageBuffer %36
+%35 = OpVariable  %37  StorageBuffer
+%39 = OpTypeStruct %19
+%40 = OpTypePointer StorageBuffer %39
+%38 = OpVariable  %40  StorageBuffer
+%42 = OpTypeStruct %22
+%43 = OpTypePointer StorageBuffer %42
+%41 = OpVariable  %43  StorageBuffer
+%47 = OpTypeFunction %3 %3
+%48 = OpTypePointer Uniform %24
+%49 = OpConstant  %4  0
+%51 = OpTypePointer StorageBuffer %19
+%53 = OpTypePointer StorageBuffer %22
+%57 = OpConstant  %3  0.000000000000000000000000000000000000000088991
+%58 = OpConstant  %3  0.000000000000000000000000000000000000000024753
+%59 = OpTypePointer Uniform %19
+%60 = OpConstant  %5  65504
+%61 = OpConstant  %5  -65504
+%62 = OpConstant  %4  65504
+%63 = OpConstant  %6  65504
+%64 = OpConstant  %6  -65504
+%66 = OpTypePointer Function %23
+%67 = OpConstantNull  %23
+%69 = OpTypePointer Function %3
+%78 = OpTypePointer Uniform %6
+%87 = OpTypePointer Uniform %3
+%88 = OpConstant  %4  3
+%95 = OpTypePointer StorageBuffer %5
+%96 = OpConstant  %4  1
+%99 = OpTypePointer StorageBuffer %4
+%102 = OpTypePointer StorageBuffer %6
+%105 = OpTypePointer StorageBuffer %3
+%112 = OpTypePointer StorageBuffer %7
+%113 = OpTypePointer Uniform %7
+%114 = OpConstant  %4  4
+%121 = OpTypePointer StorageBuffer %8
+%122 = OpTypePointer Uniform %8
+%123 = OpConstant  %4  5
+%130 = OpTypePointer StorageBuffer %9
+%131 = OpTypePointer Uniform %9
+%132 = OpConstant  %4  6
+%139 = OpTypePointer StorageBuffer %10
+%140 = OpTypePointer Uniform %10
+%141 = OpConstant  %4  8
+%142 = OpConstant  %4  9
+%158 = OpTypePointer StorageBuffer %11
+%159 = OpTypePointer Uniform %11
+%160 = OpConstant  %4  10
+%173 = OpTypePointer StorageBuffer %12
+%174 = OpTypePointer Uniform %12
+%175 = OpConstant  %4  11
+%188 = OpTypePointer StorageBuffer %13
+%189 = OpTypePointer Uniform %13
+%190 = OpConstant  %4  12
+%191 = OpConstant  %4  13
+%192 = OpConstant  %4  14
+%213 = OpTypePointer StorageBuffer %14
+%214 = OpTypePointer Uniform %14
+%215 = OpConstant  %4  15
+%231 = OpTypePointer StorageBuffer %15
+%232 = OpTypePointer Uniform %15
+%233 = OpConstant  %4  16
+%249 = OpTypePointer StorageBuffer %16
+%250 = OpTypePointer Uniform %16
+%251 = OpConstant  %4  17
+%252 = OpConstant  %4  18
+%253 = OpConstant  %4  19
+%254 = OpConstant  %4  20
+%280 = OpTypePointer StorageBuffer %17
+%281 = OpTypePointer Uniform %17
+%282 = OpConstant  %4  21
+%301 = OpTypePointer StorageBuffer %18
+%302 = OpTypePointer Uniform %18
+%303 = OpConstant  %4  22
+%322 = OpTypePointer StorageBuffer %20
+%361 = OpTypeVector %6 2
+%367 = OpTypeVector %6 3
+%373 = OpTypeVector %6 4
+%382 = OpTypeMatrix %361 2
+%396 = OpTypeMatrix %367 2
+%410 = OpTypeMatrix %373 2
+%429 = OpTypeMatrix %361 3
+%447 = OpTypeMatrix %367 3
+%465 = OpTypeMatrix %373 3
+%490 = OpTypeMatrix %361 4
+%512 = OpTypeMatrix %367 4
+%534 = OpTypeMatrix %373 4
+%557 = OpTypeFunction %2
+%563 = OpConstant  %3  0.000000000000000000000000000000000000000022959
+%566 = OpConstant  %4  7
+%46 = OpFunction  %3  None %47
+%45 = OpFunctionParameter  %3
+%44 = OpLabel
+%65 = OpVariable  %66  Function %67
+%68 = OpVariable  %69  Function %26
+%50 = OpAccessChain  %48  %29 %49
+%52 = OpAccessChain  %51  %32 %49
+%54 = OpAccessChain  %53  %35 %49
+%55 = OpAccessChain  %51  %38 %49
+%56 = OpAccessChain  %53  %41 %49
+OpBranch %70
+%70 = OpLabel
+%71 = OpLoad  %3  %27
+%72 = OpLoad  %3  %68
 %73 = OpFAdd  %3  %72 %57
-%74 = OpLoad  %3  %66
-%75 = OpFAdd  %3  %74 %73
-OpStore %66 %75
-%77 = OpAccessChain  %76  %49 %21
-%78 = OpLoad  %6  %77
-%79 = OpLoad  %3  %66
-%80 = OpFConvert  %6  %79
-%81 = OpFAdd  %6  %78 %80
-%82 = OpFConvert  %3  %81
-%83 = OpLoad  %3  %66
-%84 = OpFAdd  %3  %83 %82
-OpStore %66 %84
-%87 = OpAccessChain  %85  %49 %86
-%88 = OpLoad  %3  %87
-%89 = OpCompositeConstruct  %8  %88 %88 %88
-%90 = OpCompositeExtract  %3  %89 2
-%91 = OpLoad  %3  %66
-%92 = OpFAdd  %3  %91 %90
-OpStore %66 %92
-%95 = OpAccessChain  %93  %54 %94
-OpStore %95 %58
-%96 = OpAccessChain  %93  %54 %94
-OpStore %96 %59
-%98 = OpAccessChain  %97  %54 %48
-OpStore %98 %60
-%99 = OpAccessChain  %97  %54 %48
-OpStore %99 %48
-%101 = OpAccessChain  %100  %54 %21
-OpStore %101 %61
-%102 = OpAccessChain  %100  %54 %21
-OpStore %102 %62
-%104 = OpAccessChain  %85  %49 %86
-%105 = OpLoad  %3  %104
-%106 = OpAccessChain  %103  %51 %86
+OpStore %68 %73
+%74 = OpLoad  %3  %68
+%75 = OpFAdd  %3  %74 %58
+%76 = OpLoad  %3  %68
+%77 = OpFAdd  %3  %76 %75
+OpStore %68 %77
+%79 = OpAccessChain  %78  %50 %21
+%80 = OpLoad  %6  %79
+%81 = OpLoad  %3  %68
+%82 = OpFConvert  %6  %81
+%83 = OpFAdd  %6  %80 %82
+%84 = OpFConvert  %3  %83
+%85 = OpLoad  %3  %68
+%86 = OpFAdd  %3  %85 %84
+OpStore %68 %86
+%89 = OpAccessChain  %87  %50 %88
+%90 = OpLoad  %3  %89
+%91 = OpCompositeConstruct  %8  %90 %90 %90
+%92 = OpCompositeExtract  %3  %91 2
+%93 = OpLoad  %3  %68
+%94 = OpFAdd  %3  %93 %92
+OpStore %68 %94
+%97 = OpAccessChain  %95  %55 %96
+OpStore %97 %60
+%98 = OpAccessChain  %95  %55 %96
+OpStore %98 %61
+%100 = OpAccessChain  %99  %55 %49
+OpStore %100 %62
+%101 = OpAccessChain  %99  %55 %49
+OpStore %101 %49
+%103 = OpAccessChain  %102  %55 %21
+OpStore %103 %63
+%104 = OpAccessChain  %102  %55 %21
+OpStore %104 %64
+%106 = OpAccessChain  %87  %50 %88
 %107 = OpLoad  %3  %106
-%108 = OpFAdd  %3  %105 %107
-%109 = OpAccessChain  %103  %54 %86
-OpStore %109 %108
-%113 = OpAccessChain  %111  %49 %112
-%114 = OpLoad  %7  %113
-%115 = OpAccessChain  %110  %51 %112
+%108 = OpAccessChain  %105  %52 %88
+%109 = OpLoad  %3  %108
+%110 = OpFAdd  %3  %107 %109
+%111 = OpAccessChain  %105  %55 %88
+OpStore %111 %110
+%115 = OpAccessChain  %113  %50 %114
 %116 = OpLoad  %7  %115
-%117 = OpFAdd  %7  %114 %116
-%118 = OpAccessChain  %110  %54 %112
-OpStore %118 %117
-%122 = OpAccessChain  %120  %49 %121
-%123 = OpLoad  %8  %122
-%124 = OpAccessChain  %119  %51 %121
+%117 = OpAccessChain  %112  %52 %114
+%118 = OpLoad  %7  %117
+%119 = OpFAdd  %7  %116 %118
+%120 = OpAccessChain  %112  %55 %114
+OpStore %120 %119
+%124 = OpAccessChain  %122  %50 %123
 %125 = OpLoad  %8  %124
-%126 = OpFAdd  %8  %123 %125
-%127 = OpAccessChain  %119  %54 %121
-OpStore %127 %126
-%131 = OpAccessChain  %129  %49 %130
-%132 = OpLoad  %9  %131
-%133 = OpAccessChain  %128  %51 %130
+%126 = OpAccessChain  %121  %52 %123
+%127 = OpLoad  %8  %126
+%128 = OpFAdd  %8  %125 %127
+%129 = OpAccessChain  %121  %55 %123
+OpStore %129 %128
+%133 = OpAccessChain  %131  %50 %132
 %134 = OpLoad  %9  %133
-%135 = OpFAdd  %9  %132 %134
-%136 = OpAccessChain  %128  %54 %130
-OpStore %136 %135
-%140 = OpAccessChain  %138  %49 %139
-%141 = OpLoad  %10  %140
-%142 = OpAccessChain  %137  %51 %139
-%143 = OpLoad  %10  %142
-%145 = OpCompositeExtract  %7  %141 0
-%146 = OpCompositeExtract  %7  %143 0
-%147 = OpFAdd  %7  %145 %146
-%148 = OpCompositeExtract  %7  %141 1
-%149 = OpCompositeExtract  %7  %143 1
-%150 = OpFAdd  %7  %148 %149
-%144 = OpCompositeConstruct  %10  %147 %150
-%151 = OpAccessChain  %137  %54 %139
-OpStore %151 %144
-%155 = OpAccessChain  %153  %49 %154
-%156 = OpLoad  %11  %155
-%157 = OpAccessChain  %152  %51 %154
-%158 = OpLoad  %11  %157
-%160 = OpCompositeExtract  %8  %156 0
-%161 = OpCompositeExtract  %8  %158 0
-%162 = OpFAdd  %8  %160 %161
-%163 = OpCompositeExtract  %8  %156 1
-%164 = OpCompositeExtract  %8  %158 1
-%165 = OpFAdd  %8  %163 %164
-%159 = OpCompositeConstruct  %11  %162 %165
-%166 = OpAccessChain  %152  %54 %154
-OpStore %166 %159
-%170 = OpAccessChain  %168  %49 %169
-%171 = OpLoad  %12  %170
-%172 = OpAccessChain  %167  %51 %169
-%173 = OpLoad  %12  %172
-%175 = OpCompositeExtract  %9  %171 0
-%176 = OpCompositeExtract  %9  %173 0
-%177 = OpFAdd  %9  %175 %176
-%178 = OpCompositeExtract  %9  %171 1
-%179 = OpCompositeExtract  %9  %173 1
-%180 = OpFAdd  %9  %178 %179
-%174 = OpCompositeConstruct  %12  %177 %180
-%181 = OpAccessChain  %167  %54 %169
-OpStore %181 %174
-%185 = OpAccessChain  %183  %49 %184
-%186 = OpLoad  %13  %185
-%187 = OpAccessChain  %182  %51 %184
-%188 = OpLoad  %13  %187
-%190 = OpCompositeExtract  %7  %186 0
-%191 = OpCompositeExtract  %7  %188 0
-%192 = OpFAdd  %7  %190 %191
-%193 = OpCompositeExtract  %7  %186 1
-%194 = OpCompositeExtract  %7  %188 1
-%195 = OpFAdd  %7  %193 %194
-%196 = OpCompositeExtract  %7  %186 2
-%197 = OpCompositeExtract  %7  %188 2
-%198 = OpFAdd  %7  %196 %197
-%189 = OpCompositeConstruct  %13  %192 %195 %198
-%199 = OpAccessChain  %182  %54 %184
-OpStore %199 %189
-%203 = OpAccessChain  %201  %49 %202
-%204 = OpLoad  %14  %203
-%205 = OpAccessChain  %200  %51 %202
-%206 = OpLoad  %14  %205
-%208 = OpCompositeExtract  %8  %204 0
-%209 = OpCompositeExtract  %8  %206 0
-%210 = OpFAdd  %8  %208 %209
-%211 = OpCompositeExtract  %8  %204 1
-%212 = OpCompositeExtract  %8  %206 1
-%213 = OpFAdd  %8  %211 %212
-%214 = OpCompositeExtract  %8  %204 2
-%215 = OpCompositeExtract  %8  %206 2
-%216 = OpFAdd  %8  %214 %215
-%207 = OpCompositeConstruct  %14  %210 %213 %216
-%217 = OpAccessChain  %200  %54 %202
-OpStore %217 %207
-%221 = OpAccessChain  %219  %49 %220
-%222 = OpLoad  %15  %221
-%223 = OpAccessChain  %218  %51 %220
-%224 = OpLoad  %15  %223
-%226 = OpCompositeExtract  %9  %222 0
-%227 = OpCompositeExtract  %9  %224 0
-%228 = OpFAdd  %9  %226 %227
-%229 = OpCompositeExtract  %9  %222 1
-%230 = OpCompositeExtract  %9  %224 1
-%231 = OpFAdd  %9  %229 %230
-%232 = OpCompositeExtract  %9  %222 2
-%233 = OpCompositeExtract  %9  %224 2
-%234 = OpFAdd  %9  %232 %233
-%225 = OpCompositeConstruct  %15  %228 %231 %234
-%235 = OpAccessChain  %218  %54 %220
-OpStore %235 %225
-%239 = OpAccessChain  %237  %49 %238
-%240 = OpLoad  %16  %239
-%241 = OpAccessChain  %236  %51 %238
-%242 = OpLoad  %16  %241
-%244 = OpCompositeExtract  %7  %240 0
-%245 = OpCompositeExtract  %7  %242 0
-%246 = OpFAdd  %7  %244 %245
-%247 = OpCompositeExtract  %7  %240 1
-%248 = OpCompositeExtract  %7  %242 1
-%249 = OpFAdd  %7  %247 %248
-%250 = OpCompositeExtract  %7  %240 2
-%251 = OpCompositeExtract  %7  %242 2
-%252 = OpFAdd  %7  %250 %251
-%253 = OpCompositeExtract  %7  %240 3
-%254 = OpCompositeExtract  %7  %242 3
-%255 = OpFAdd  %7  %253 %254
-%243 = OpCompositeConstruct  %16  %246 %249 %252 %255
-%256 = OpAccessChain  %236  %54 %238
-OpStore %256 %243
-%260 = OpAccessChain  %258  %49 %259
-%261 = OpLoad  %17  %260
-%262 = OpAccessChain  %257  %51 %259
-%263 = OpLoad  %17  %262
-%265 = OpCompositeExtract  %8  %261 0
-%266 = OpCompositeExtract  %8  %263 0
-%267 = OpFAdd  %8  %265 %266
-%268 = OpCompositeExtract  %8  %261 1
-%269 = OpCompositeExtract  %8  %263 1
-%270 = OpFAdd  %8  %268 %269
-%271 = OpCompositeExtract  %8  %261 2
-%272 = OpCompositeExtract  %8  %263 2
-%273 = OpFAdd  %8  %271 %272
-%274 = OpCompositeExtract  %8  %261 3
-%275 = OpCompositeExtract  %8  %263 3
-%276 = OpFAdd  %8  %274 %275
-%264 = OpCompositeConstruct  %17  %267 %270 %273 %276
-%277 = OpAccessChain  %257  %54 %259
-OpStore %277 %264
-%281 = OpAccessChain  %279  %49 %280
-%282 = OpLoad  %18  %281
-%283 = OpAccessChain  %278  %51 %280
-%284 = OpLoad  %18  %283
-%286 = OpCompositeExtract  %9  %282 0
-%287 = OpCompositeExtract  %9  %284 0
-%288 = OpFAdd  %9  %286 %287
-%289 = OpCompositeExtract  %9  %282 1
-%290 = OpCompositeExtract  %9  %284 1
-%291 = OpFAdd  %9  %289 %290
-%292 = OpCompositeExtract  %9  %282 2
-%293 = OpCompositeExtract  %9  %284 2
-%294 = OpFAdd  %9  %292 %293
-%295 = OpCompositeExtract  %9  %282 3
-%296 = OpCompositeExtract  %9  %284 3
-%297 = OpFAdd  %9  %295 %296
-%285 = OpCompositeConstruct  %18  %288 %291 %294 %297
-%298 = OpAccessChain  %278  %54 %280
-OpStore %298 %285
-%300 = OpAccessChain  %299  %53 %48
-%301 = OpLoad  %20  %300
-%302 = OpAccessChain  %299  %55 %48
-OpStore %302 %301
-%303 = OpLoad  %3  %66
-%304 = OpExtInst  %3  %1 FAbs %303
-%305 = OpLoad  %3  %66
-%306 = OpFAdd  %3  %305 %304
-OpStore %66 %306
-%307 = OpLoad  %3  %66
-%308 = OpLoad  %3  %66
-%309 = OpLoad  %3  %66
-%310 = OpExtInst  %3  %1 FClamp %307 %308 %309
-%311 = OpLoad  %3  %66
-%312 = OpFAdd  %3  %311 %310
-OpStore %66 %312
-%313 = OpLoad  %3  %66
-%314 = OpCompositeConstruct  %7  %313 %313
-%315 = OpLoad  %3  %66
-%316 = OpCompositeConstruct  %7  %315 %315
-%317 = OpDot  %3  %314 %316
-%318 = OpLoad  %3  %66
-%319 = OpFAdd  %3  %318 %317
-OpStore %66 %319
-%320 = OpLoad  %3  %66
-%321 = OpLoad  %3  %66
-%322 = OpExtInst  %3  %1 FMax %320 %321
-%323 = OpLoad  %3  %66
-%324 = OpFAdd  %3  %323 %322
-OpStore %66 %324
-%325 = OpLoad  %3  %66
-%326 = OpLoad  %3  %66
-%327 = OpExtInst  %3  %1 FMin %325 %326
-%328 = OpLoad  %3  %66
+%135 = OpAccessChain  %130  %52 %132
+%136 = OpLoad  %9  %135
+%137 = OpFAdd  %9  %134 %136
+%138 = OpAccessChain  %130  %55 %132
+OpStore %138 %137
+%143 = OpAccessChain  %113  %50 %141
+%144 = OpLoad  %7  %143
+%145 = OpAccessChain  %113  %50 %142
+%146 = OpLoad  %7  %145
+%147 = OpCompositeConstruct  %10  %144 %146
+%148 = OpAccessChain  %139  %52 %141
+%149 = OpLoad  %10  %148
+%151 = OpCompositeExtract  %7  %147 0
+%152 = OpCompositeExtract  %7  %149 0
+%153 = OpFAdd  %7  %151 %152
+%154 = OpCompositeExtract  %7  %147 1
+%155 = OpCompositeExtract  %7  %149 1
+%156 = OpFAdd  %7  %154 %155
+%150 = OpCompositeConstruct  %10  %153 %156
+%157 = OpAccessChain  %139  %55 %141
+OpStore %157 %150
+%161 = OpAccessChain  %159  %50 %160
+%162 = OpLoad  %11  %161
+%163 = OpAccessChain  %158  %52 %142
+%164 = OpLoad  %11  %163
+%166 = OpCompositeExtract  %8  %162 0
+%167 = OpCompositeExtract  %8  %164 0
+%168 = OpFAdd  %8  %166 %167
+%169 = OpCompositeExtract  %8  %162 1
+%170 = OpCompositeExtract  %8  %164 1
+%171 = OpFAdd  %8  %169 %170
+%165 = OpCompositeConstruct  %11  %168 %171
+%172 = OpAccessChain  %158  %55 %142
+OpStore %172 %165
+%176 = OpAccessChain  %174  %50 %175
+%177 = OpLoad  %12  %176
+%178 = OpAccessChain  %173  %52 %160
+%179 = OpLoad  %12  %178
+%181 = OpCompositeExtract  %9  %177 0
+%182 = OpCompositeExtract  %9  %179 0
+%183 = OpFAdd  %9  %181 %182
+%184 = OpCompositeExtract  %9  %177 1
+%185 = OpCompositeExtract  %9  %179 1
+%186 = OpFAdd  %9  %184 %185
+%180 = OpCompositeConstruct  %12  %183 %186
+%187 = OpAccessChain  %173  %55 %160
+OpStore %187 %180
+%193 = OpAccessChain  %113  %50 %190
+%194 = OpLoad  %7  %193
+%195 = OpAccessChain  %113  %50 %191
+%196 = OpLoad  %7  %195
+%197 = OpAccessChain  %113  %50 %192
+%198 = OpLoad  %7  %197
+%199 = OpCompositeConstruct  %13  %194 %196 %198
+%200 = OpAccessChain  %188  %52 %175
+%201 = OpLoad  %13  %200
+%203 = OpCompositeExtract  %7  %199 0
+%204 = OpCompositeExtract  %7  %201 0
+%205 = OpFAdd  %7  %203 %204
+%206 = OpCompositeExtract  %7  %199 1
+%207 = OpCompositeExtract  %7  %201 1
+%208 = OpFAdd  %7  %206 %207
+%209 = OpCompositeExtract  %7  %199 2
+%210 = OpCompositeExtract  %7  %201 2
+%211 = OpFAdd  %7  %209 %210
+%202 = OpCompositeConstruct  %13  %205 %208 %211
+%212 = OpAccessChain  %188  %55 %175
+OpStore %212 %202
+%216 = OpAccessChain  %214  %50 %215
+%217 = OpLoad  %14  %216
+%218 = OpAccessChain  %213  %52 %190
+%219 = OpLoad  %14  %218
+%221 = OpCompositeExtract  %8  %217 0
+%222 = OpCompositeExtract  %8  %219 0
+%223 = OpFAdd  %8  %221 %222
+%224 = OpCompositeExtract  %8  %217 1
+%225 = OpCompositeExtract  %8  %219 1
+%226 = OpFAdd  %8  %224 %225
+%227 = OpCompositeExtract  %8  %217 2
+%228 = OpCompositeExtract  %8  %219 2
+%229 = OpFAdd  %8  %227 %228
+%220 = OpCompositeConstruct  %14  %223 %226 %229
+%230 = OpAccessChain  %213  %55 %190
+OpStore %230 %220
+%234 = OpAccessChain  %232  %50 %233
+%235 = OpLoad  %15  %234
+%236 = OpAccessChain  %231  %52 %191
+%237 = OpLoad  %15  %236
+%239 = OpCompositeExtract  %9  %235 0
+%240 = OpCompositeExtract  %9  %237 0
+%241 = OpFAdd  %9  %239 %240
+%242 = OpCompositeExtract  %9  %235 1
+%243 = OpCompositeExtract  %9  %237 1
+%244 = OpFAdd  %9  %242 %243
+%245 = OpCompositeExtract  %9  %235 2
+%246 = OpCompositeExtract  %9  %237 2
+%247 = OpFAdd  %9  %245 %246
+%238 = OpCompositeConstruct  %15  %241 %244 %247
+%248 = OpAccessChain  %231  %55 %191
+OpStore %248 %238
+%255 = OpAccessChain  %113  %50 %251
+%256 = OpLoad  %7  %255
+%257 = OpAccessChain  %113  %50 %252
+%258 = OpLoad  %7  %257
+%259 = OpAccessChain  %113  %50 %253
+%260 = OpLoad  %7  %259
+%261 = OpAccessChain  %113  %50 %254
+%262 = OpLoad  %7  %261
+%263 = OpCompositeConstruct  %16  %256 %258 %260 %262
+%264 = OpAccessChain  %249  %52 %192
+%265 = OpLoad  %16  %264
+%267 = OpCompositeExtract  %7  %263 0
+%268 = OpCompositeExtract  %7  %265 0
+%269 = OpFAdd  %7  %267 %268
+%270 = OpCompositeExtract  %7  %263 1
+%271 = OpCompositeExtract  %7  %265 1
+%272 = OpFAdd  %7  %270 %271
+%273 = OpCompositeExtract  %7  %263 2
+%274 = OpCompositeExtract  %7  %265 2
+%275 = OpFAdd  %7  %273 %274
+%276 = OpCompositeExtract  %7  %263 3
+%277 = OpCompositeExtract  %7  %265 3
+%278 = OpFAdd  %7  %276 %277
+%266 = OpCompositeConstruct  %16  %269 %272 %275 %278
+%279 = OpAccessChain  %249  %55 %192
+OpStore %279 %266
+%283 = OpAccessChain  %281  %50 %282
+%284 = OpLoad  %17  %283
+%285 = OpAccessChain  %280  %52 %215
+%286 = OpLoad  %17  %285
+%288 = OpCompositeExtract  %8  %284 0
+%289 = OpCompositeExtract  %8  %286 0
+%290 = OpFAdd  %8  %288 %289
+%291 = OpCompositeExtract  %8  %284 1
+%292 = OpCompositeExtract  %8  %286 1
+%293 = OpFAdd  %8  %291 %292
+%294 = OpCompositeExtract  %8  %284 2
+%295 = OpCompositeExtract  %8  %286 2
+%296 = OpFAdd  %8  %294 %295
+%297 = OpCompositeExtract  %8  %284 3
+%298 = OpCompositeExtract  %8  %286 3
+%299 = OpFAdd  %8  %297 %298
+%287 = OpCompositeConstruct  %17  %290 %293 %296 %299
+%300 = OpAccessChain  %280  %55 %215
+OpStore %300 %287
+%304 = OpAccessChain  %302  %50 %303
+%305 = OpLoad  %18  %304
+%306 = OpAccessChain  %301  %52 %233
+%307 = OpLoad  %18  %306
+%309 = OpCompositeExtract  %9  %305 0
+%310 = OpCompositeExtract  %9  %307 0
+%311 = OpFAdd  %9  %309 %310
+%312 = OpCompositeExtract  %9  %305 1
+%313 = OpCompositeExtract  %9  %307 1
+%314 = OpFAdd  %9  %312 %313
+%315 = OpCompositeExtract  %9  %305 2
+%316 = OpCompositeExtract  %9  %307 2
+%317 = OpFAdd  %9  %315 %316
+%318 = OpCompositeExtract  %9  %305 3
+%319 = OpCompositeExtract  %9  %307 3
+%320 = OpFAdd  %9  %318 %319
+%308 = OpCompositeConstruct  %18  %311 %314 %317 %320
+%321 = OpAccessChain  %301  %55 %233
+OpStore %321 %308
+%323 = OpAccessChain  %322  %54 %49
+%324 = OpLoad  %20  %323
+%325 = OpAccessChain  %322  %56 %49
+OpStore %325 %324
+%326 = OpLoad  %3  %68
+%327 = OpExtInst  %3  %1 FAbs %326
+%328 = OpLoad  %3  %68
 %329 = OpFAdd  %3  %328 %327
-OpStore %66 %329
-%330 = OpLoad  %3  %66
-%331 = OpExtInst  %3  %1 FSign %330
-%332 = OpLoad  %3  %66
-%333 = OpFAdd  %3  %332 %331
-OpStore %66 %333
-%334 = OpLoad  %3  %66
-%335 = OpFAdd  %3  %334 %24
-OpStore %66 %335
-%336 = OpAccessChain  %111  %49 %112
-%337 = OpLoad  %7  %336
-%339 = OpFConvert  %338  %337
-%340 = OpFConvert  %7  %339
-%341 = OpAccessChain  %110  %54 %112
-OpStore %341 %340
-%342 = OpAccessChain  %120  %49 %121
-%343 = OpLoad  %8  %342
-%345 = OpFConvert  %344  %343
-%346 = OpFConvert  %8  %345
-%347 = OpAccessChain  %119  %54 %121
-OpStore %347 %346
-%348 = OpAccessChain  %129  %49 %130
-%349 = OpLoad  %9  %348
-%351 = OpFConvert  %350  %349
-%352 = OpFConvert  %9  %351
-%353 = OpAccessChain  %128  %54 %130
-OpStore %353 %352
-%354 = OpAccessChain  %138  %49 %139
-%355 = OpLoad  %10  %354
-%357 = OpCompositeExtract  %7  %355 0
-%358 = OpFConvert  %338  %357
-%359 = OpCompositeExtract  %7  %355 1
-%360 = OpFConvert  %338  %359
-%361 = OpCompositeConstruct  %356  %358 %360
-%362 = OpCompositeExtract  %338  %361 0
+OpStore %68 %329
+%330 = OpLoad  %3  %68
+%331 = OpLoad  %3  %68
+%332 = OpLoad  %3  %68
+%333 = OpExtInst  %3  %1 FClamp %330 %331 %332
+%334 = OpLoad  %3  %68
+%335 = OpFAdd  %3  %334 %333
+OpStore %68 %335
+%336 = OpLoad  %3  %68
+%337 = OpCompositeConstruct  %7  %336 %336
+%338 = OpLoad  %3  %68
+%339 = OpCompositeConstruct  %7  %338 %338
+%340 = OpDot  %3  %337 %339
+%341 = OpLoad  %3  %68
+%342 = OpFAdd  %3  %341 %340
+OpStore %68 %342
+%343 = OpLoad  %3  %68
+%344 = OpLoad  %3  %68
+%345 = OpExtInst  %3  %1 FMax %343 %344
+%346 = OpLoad  %3  %68
+%347 = OpFAdd  %3  %346 %345
+OpStore %68 %347
+%348 = OpLoad  %3  %68
+%349 = OpLoad  %3  %68
+%350 = OpExtInst  %3  %1 FMin %348 %349
+%351 = OpLoad  %3  %68
+%352 = OpFAdd  %3  %351 %350
+OpStore %68 %352
+%353 = OpLoad  %3  %68
+%354 = OpExtInst  %3  %1 FSign %353
+%355 = OpLoad  %3  %68
+%356 = OpFAdd  %3  %355 %354
+OpStore %68 %356
+%357 = OpLoad  %3  %68
+%358 = OpFAdd  %3  %357 %25
+OpStore %68 %358
+%359 = OpAccessChain  %113  %50 %114
+%360 = OpLoad  %7  %359
+%362 = OpFConvert  %361  %360
 %363 = OpFConvert  %7  %362
-%364 = OpCompositeExtract  %338  %361 1
-%365 = OpFConvert  %7  %364
-%366 = OpCompositeConstruct  %10  %363 %365
-%367 = OpAccessChain  %137  %54 %139
-OpStore %367 %366
-%368 = OpAccessChain  %153  %49 %154
-%369 = OpLoad  %11  %368
-%371 = OpCompositeExtract  %8  %369 0
-%372 = OpFConvert  %344  %371
-%373 = OpCompositeExtract  %8  %369 1
-%374 = OpFConvert  %344  %373
-%375 = OpCompositeConstruct  %370  %372 %374
-%376 = OpCompositeExtract  %344  %375 0
-%377 = OpFConvert  %8  %376
-%378 = OpCompositeExtract  %344  %375 1
-%379 = OpFConvert  %8  %378
-%380 = OpCompositeConstruct  %11  %377 %379
-%381 = OpAccessChain  %152  %54 %154
-OpStore %381 %380
-%382 = OpAccessChain  %168  %49 %169
-%383 = OpLoad  %12  %382
-%385 = OpCompositeExtract  %9  %383 0
-%386 = OpFConvert  %350  %385
-%387 = OpCompositeExtract  %9  %383 1
-%388 = OpFConvert  %350  %387
-%389 = OpCompositeConstruct  %384  %386 %388
-%390 = OpCompositeExtract  %350  %389 0
-%391 = OpFConvert  %9  %390
-%392 = OpCompositeExtract  %350  %389 1
-%393 = OpFConvert  %9  %392
-%394 = OpCompositeConstruct  %12  %391 %393
-%395 = OpAccessChain  %167  %54 %169
-OpStore %395 %394
-%396 = OpAccessChain  %183  %49 %184
-%397 = OpLoad  %13  %396
-%399 = OpCompositeExtract  %7  %397 0
-%400 = OpFConvert  %338  %399
-%401 = OpCompositeExtract  %7  %397 1
-%402 = OpFConvert  %338  %401
-%403 = OpCompositeExtract  %7  %397 2
-%404 = OpFConvert  %338  %403
-%405 = OpCompositeConstruct  %398  %400 %402 %404
-%406 = OpCompositeExtract  %338  %405 0
-%407 = OpFConvert  %7  %406
-%408 = OpCompositeExtract  %338  %405 1
-%409 = OpFConvert  %7  %408
-%410 = OpCompositeExtract  %338  %405 2
-%411 = OpFConvert  %7  %410
-%412 = OpCompositeConstruct  %13  %407 %409 %411
-%413 = OpAccessChain  %182  %54 %184
-OpStore %413 %412
-%414 = OpAccessChain  %201  %49 %202
-%415 = OpLoad  %14  %414
-%417 = OpCompositeExtract  %8  %415 0
-%418 = OpFConvert  %344  %417
-%419 = OpCompositeExtract  %8  %415 1
-%420 = OpFConvert  %344  %419
-%421 = OpCompositeExtract  %8  %415 2
-%422 = OpFConvert  %344  %421
-%423 = OpCompositeConstruct  %416  %418 %420 %422
-%424 = OpCompositeExtract  %344  %423 0
-%425 = OpFConvert  %8  %424
-%426 = OpCompositeExtract  %344  %423 1
-%427 = OpFConvert  %8  %426
-%428 = OpCompositeExtract  %344  %423 2
-%429 = OpFConvert  %8  %428
-%430 = OpCompositeConstruct  %14  %425 %427 %429
-%431 = OpAccessChain  %200  %54 %202
-OpStore %431 %430
-%432 = OpAccessChain  %219  %49 %220
-%433 = OpLoad  %15  %432
-%435 = OpCompositeExtract  %9  %433 0
-%436 = OpFConvert  %350  %435
-%437 = OpCompositeExtract  %9  %433 1
-%438 = OpFConvert  %350  %437
-%439 = OpCompositeExtract  %9  %433 2
-%440 = OpFConvert  %350  %439
-%441 = OpCompositeConstruct  %434  %436 %438 %440
-%442 = OpCompositeExtract  %350  %441 0
-%443 = OpFConvert  %9  %442
-%444 = OpCompositeExtract  %350  %441 1
-%445 = OpFConvert  %9  %444
-%446 = OpCompositeExtract  %350  %441 2
-%447 = OpFConvert  %9  %446
-%448 = OpCompositeConstruct  %15  %443 %445 %447
-%449 = OpAccessChain  %218  %54 %220
-OpStore %449 %448
-%450 = OpAccessChain  %237  %49 %238
-%451 = OpLoad  %16  %450
-%453 = OpCompositeExtract  %7  %451 0
-%454 = OpFConvert  %338  %453
-%455 = OpCompositeExtract  %7  %451 1
-%456 = OpFConvert  %338  %455
-%457 = OpCompositeExtract  %7  %451 2
-%458 = OpFConvert  %338  %457
-%459 = OpCompositeExtract  %7  %451 3
-%460 = OpFConvert  %338  %459
-%461 = OpCompositeConstruct  %452  %454 %456 %458 %460
-%462 = OpCompositeExtract  %338  %461 0
-%463 = OpFConvert  %7  %462
-%464 = OpCompositeExtract  %338  %461 1
-%465 = OpFConvert  %7  %464
-%466 = OpCompositeExtract  %338  %461 2
-%467 = OpFConvert  %7  %466
-%468 = OpCompositeExtract  %338  %461 3
-%469 = OpFConvert  %7  %468
-%470 = OpCompositeConstruct  %16  %463 %465 %467 %469
-%471 = OpAccessChain  %236  %54 %238
-OpStore %471 %470
-%472 = OpAccessChain  %258  %49 %259
-%473 = OpLoad  %17  %472
-%475 = OpCompositeExtract  %8  %473 0
-%476 = OpFConvert  %344  %475
-%477 = OpCompositeExtract  %8  %473 1
-%478 = OpFConvert  %344  %477
-%479 = OpCompositeExtract  %8  %473 2
-%480 = OpFConvert  %344  %479
-%481 = OpCompositeExtract  %8  %473 3
-%482 = OpFConvert  %344  %481
-%483 = OpCompositeConstruct  %474  %476 %478 %480 %482
-%484 = OpCompositeExtract  %344  %483 0
-%485 = OpFConvert  %8  %484
-%486 = OpCompositeExtract  %344  %483 1
-%487 = OpFConvert  %8  %486
-%488 = OpCompositeExtract  %344  %483 2
-%489 = OpFConvert  %8  %488
-%490 = OpCompositeExtract  %344  %483 3
-%491 = OpFConvert  %8  %490
-%492 = OpCompositeConstruct  %17  %485 %487 %489 %491
-%493 = OpAccessChain  %257  %54 %259
-OpStore %493 %492
-%494 = OpAccessChain  %279  %49 %280
-%495 = OpLoad  %18  %494
-%497 = OpCompositeExtract  %9  %495 0
-%498 = OpFConvert  %350  %497
-%499 = OpCompositeExtract  %9  %495 1
-%500 = OpFConvert  %350  %499
-%501 = OpCompositeExtract  %9  %495 2
-%502 = OpFConvert  %350  %501
-%503 = OpCompositeExtract  %9  %495 3
-%504 = OpFConvert  %350  %503
-%505 = OpCompositeConstruct  %496  %498 %500 %502 %504
-%506 = OpCompositeExtract  %350  %505 0
-%507 = OpFConvert  %9  %506
-%508 = OpCompositeExtract  %350  %505 1
-%509 = OpFConvert  %9  %508
-%510 = OpCompositeExtract  %350  %505 2
-%511 = OpFConvert  %9  %510
-%512 = OpCompositeExtract  %350  %505 3
-%513 = OpFConvert  %9  %512
-%514 = OpCompositeConstruct  %18  %507 %509 %511 %513
-%515 = OpAccessChain  %278  %54 %280
-OpStore %515 %514
-%516 = OpLoad  %3  %66
-OpReturnValue %516
+%364 = OpAccessChain  %112  %55 %114
+OpStore %364 %363
+%365 = OpAccessChain  %122  %50 %123
+%366 = OpLoad  %8  %365
+%368 = OpFConvert  %367  %366
+%369 = OpFConvert  %8  %368
+%370 = OpAccessChain  %121  %55 %123
+OpStore %370 %369
+%371 = OpAccessChain  %131  %50 %132
+%372 = OpLoad  %9  %371
+%374 = OpFConvert  %373  %372
+%375 = OpFConvert  %9  %374
+%376 = OpAccessChain  %130  %55 %132
+OpStore %376 %375
+%377 = OpAccessChain  %113  %50 %141
+%378 = OpLoad  %7  %377
+%379 = OpAccessChain  %113  %50 %142
+%380 = OpLoad  %7  %379
+%381 = OpCompositeConstruct  %10  %378 %380
+%383 = OpCompositeExtract  %7  %381 0
+%384 = OpFConvert  %361  %383
+%385 = OpCompositeExtract  %7  %381 1
+%386 = OpFConvert  %361  %385
+%387 = OpCompositeConstruct  %382  %384 %386
+%388 = OpCompositeExtract  %361  %387 0
+%389 = OpFConvert  %7  %388
+%390 = OpCompositeExtract  %361  %387 1
+%391 = OpFConvert  %7  %390
+%392 = OpCompositeConstruct  %10  %389 %391
+%393 = OpAccessChain  %139  %55 %141
+OpStore %393 %392
+%394 = OpAccessChain  %159  %50 %160
+%395 = OpLoad  %11  %394
+%397 = OpCompositeExtract  %8  %395 0
+%398 = OpFConvert  %367  %397
+%399 = OpCompositeExtract  %8  %395 1
+%400 = OpFConvert  %367  %399
+%401 = OpCompositeConstruct  %396  %398 %400
+%402 = OpCompositeExtract  %367  %401 0
+%403 = OpFConvert  %8  %402
+%404 = OpCompositeExtract  %367  %401 1
+%405 = OpFConvert  %8  %404
+%406 = OpCompositeConstruct  %11  %403 %405
+%407 = OpAccessChain  %158  %55 %142
+OpStore %407 %406
+%408 = OpAccessChain  %174  %50 %175
+%409 = OpLoad  %12  %408
+%411 = OpCompositeExtract  %9  %409 0
+%412 = OpFConvert  %373  %411
+%413 = OpCompositeExtract  %9  %409 1
+%414 = OpFConvert  %373  %413
+%415 = OpCompositeConstruct  %410  %412 %414
+%416 = OpCompositeExtract  %373  %415 0
+%417 = OpFConvert  %9  %416
+%418 = OpCompositeExtract  %373  %415 1
+%419 = OpFConvert  %9  %418
+%420 = OpCompositeConstruct  %12  %417 %419
+%421 = OpAccessChain  %173  %55 %160
+OpStore %421 %420
+%422 = OpAccessChain  %113  %50 %190
+%423 = OpLoad  %7  %422
+%424 = OpAccessChain  %113  %50 %191
+%425 = OpLoad  %7  %424
+%426 = OpAccessChain  %113  %50 %192
+%427 = OpLoad  %7  %426
+%428 = OpCompositeConstruct  %13  %423 %425 %427
+%430 = OpCompositeExtract  %7  %428 0
+%431 = OpFConvert  %361  %430
+%432 = OpCompositeExtract  %7  %428 1
+%433 = OpFConvert  %361  %432
+%434 = OpCompositeExtract  %7  %428 2
+%435 = OpFConvert  %361  %434
+%436 = OpCompositeConstruct  %429  %431 %433 %435
+%437 = OpCompositeExtract  %361  %436 0
+%438 = OpFConvert  %7  %437
+%439 = OpCompositeExtract  %361  %436 1
+%440 = OpFConvert  %7  %439
+%441 = OpCompositeExtract  %361  %436 2
+%442 = OpFConvert  %7  %441
+%443 = OpCompositeConstruct  %13  %438 %440 %442
+%444 = OpAccessChain  %188  %55 %175
+OpStore %444 %443
+%445 = OpAccessChain  %214  %50 %215
+%446 = OpLoad  %14  %445
+%448 = OpCompositeExtract  %8  %446 0
+%449 = OpFConvert  %367  %448
+%450 = OpCompositeExtract  %8  %446 1
+%451 = OpFConvert  %367  %450
+%452 = OpCompositeExtract  %8  %446 2
+%453 = OpFConvert  %367  %452
+%454 = OpCompositeConstruct  %447  %449 %451 %453
+%455 = OpCompositeExtract  %367  %454 0
+%456 = OpFConvert  %8  %455
+%457 = OpCompositeExtract  %367  %454 1
+%458 = OpFConvert  %8  %457
+%459 = OpCompositeExtract  %367  %454 2
+%460 = OpFConvert  %8  %459
+%461 = OpCompositeConstruct  %14  %456 %458 %460
+%462 = OpAccessChain  %213  %55 %190
+OpStore %462 %461
+%463 = OpAccessChain  %232  %50 %233
+%464 = OpLoad  %15  %463
+%466 = OpCompositeExtract  %9  %464 0
+%467 = OpFConvert  %373  %466
+%468 = OpCompositeExtract  %9  %464 1
+%469 = OpFConvert  %373  %468
+%470 = OpCompositeExtract  %9  %464 2
+%471 = OpFConvert  %373  %470
+%472 = OpCompositeConstruct  %465  %467 %469 %471
+%473 = OpCompositeExtract  %373  %472 0
+%474 = OpFConvert  %9  %473
+%475 = OpCompositeExtract  %373  %472 1
+%476 = OpFConvert  %9  %475
+%477 = OpCompositeExtract  %373  %472 2
+%478 = OpFConvert  %9  %477
+%479 = OpCompositeConstruct  %15  %474 %476 %478
+%480 = OpAccessChain  %231  %55 %191
+OpStore %480 %479
+%481 = OpAccessChain  %113  %50 %251
+%482 = OpLoad  %7  %481
+%483 = OpAccessChain  %113  %50 %252
+%484 = OpLoad  %7  %483
+%485 = OpAccessChain  %113  %50 %253
+%486 = OpLoad  %7  %485
+%487 = OpAccessChain  %113  %50 %254
+%488 = OpLoad  %7  %487
+%489 = OpCompositeConstruct  %16  %482 %484 %486 %488
+%491 = OpCompositeExtract  %7  %489 0
+%492 = OpFConvert  %361  %491
+%493 = OpCompositeExtract  %7  %489 1
+%494 = OpFConvert  %361  %493
+%495 = OpCompositeExtract  %7  %489 2
+%496 = OpFConvert  %361  %495
+%497 = OpCompositeExtract  %7  %489 3
+%498 = OpFConvert  %361  %497
+%499 = OpCompositeConstruct  %490  %492 %494 %496 %498
+%500 = OpCompositeExtract  %361  %499 0
+%501 = OpFConvert  %7  %500
+%502 = OpCompositeExtract  %361  %499 1
+%503 = OpFConvert  %7  %502
+%504 = OpCompositeExtract  %361  %499 2
+%505 = OpFConvert  %7  %504
+%506 = OpCompositeExtract  %361  %499 3
+%507 = OpFConvert  %7  %506
+%508 = OpCompositeConstruct  %16  %501 %503 %505 %507
+%509 = OpAccessChain  %249  %55 %192
+OpStore %509 %508
+%510 = OpAccessChain  %281  %50 %282
+%511 = OpLoad  %17  %510
+%513 = OpCompositeExtract  %8  %511 0
+%514 = OpFConvert  %367  %513
+%515 = OpCompositeExtract  %8  %511 1
+%516 = OpFConvert  %367  %515
+%517 = OpCompositeExtract  %8  %511 2
+%518 = OpFConvert  %367  %517
+%519 = OpCompositeExtract  %8  %511 3
+%520 = OpFConvert  %367  %519
+%521 = OpCompositeConstruct  %512  %514 %516 %518 %520
+%522 = OpCompositeExtract  %367  %521 0
+%523 = OpFConvert  %8  %522
+%524 = OpCompositeExtract  %367  %521 1
+%525 = OpFConvert  %8  %524
+%526 = OpCompositeExtract  %367  %521 2
+%527 = OpFConvert  %8  %526
+%528 = OpCompositeExtract  %367  %521 3
+%529 = OpFConvert  %8  %528
+%530 = OpCompositeConstruct  %17  %523 %525 %527 %529
+%531 = OpAccessChain  %280  %55 %215
+OpStore %531 %530
+%532 = OpAccessChain  %302  %50 %303
+%533 = OpLoad  %18  %532
+%535 = OpCompositeExtract  %9  %533 0
+%536 = OpFConvert  %373  %535
+%537 = OpCompositeExtract  %9  %533 1
+%538 = OpFConvert  %373  %537
+%539 = OpCompositeExtract  %9  %533 2
+%540 = OpFConvert  %373  %539
+%541 = OpCompositeExtract  %9  %533 3
+%542 = OpFConvert  %373  %541
+%543 = OpCompositeConstruct  %534  %536 %538 %540 %542
+%544 = OpCompositeExtract  %373  %543 0
+%545 = OpFConvert  %9  %544
+%546 = OpCompositeExtract  %373  %543 1
+%547 = OpFConvert  %9  %546
+%548 = OpCompositeExtract  %373  %543 2
+%549 = OpFConvert  %9  %548
+%550 = OpCompositeExtract  %373  %543 3
+%551 = OpFConvert  %9  %550
+%552 = OpCompositeConstruct  %18  %545 %547 %549 %551
+%553 = OpAccessChain  %301  %55 %233
+OpStore %553 %552
+%554 = OpLoad  %3  %68
+OpReturnValue %554
 OpFunctionEnd
-%518 = OpFunction  %2  None %519
-%517 = OpLabel
-%520 = OpAccessChain  %47  %28 %48
-%521 = OpAccessChain  %50  %31 %48
-%522 = OpAccessChain  %52  %34 %48
-%523 = OpAccessChain  %50  %37 %48
-%524 = OpAccessChain  %52  %40 %48
-OpBranch %526
-%526 = OpLabel
-%527 = OpFunctionCall  %3  %45 %525
-%529 = OpAccessChain  %103  %523 %528
-OpStore %529 %527
+%556 = OpFunction  %2  None %557
+%555 = OpLabel
+%558 = OpAccessChain  %48  %29 %49
+%559 = OpAccessChain  %51  %32 %49
+%560 = OpAccessChain  %53  %35 %49
+%561 = OpAccessChain  %51  %38 %49
+%562 = OpAccessChain  %53  %41 %49
+OpBranch %564
+%564 = OpLabel
+%565 = OpFunctionCall  %3  %46 %563
+%567 = OpAccessChain  %105  %561 %566
+OpStore %567 %565
 OpReturn
 OpFunctionEnd
\ No newline at end of file
diff --git a/naga/tests/out/spv/wgsl-globals.spvasm b/naga/tests/out/spv/wgsl-globals.spvasm
index 8feae9610ce..bc952f4f2d8 100644
--- a/naga/tests/out/spv/wgsl-globals.spvasm
+++ b/naga/tests/out/spv/wgsl-globals.spvasm
@@ -1,13 +1,13 @@
 ; SPIR-V
 ; Version: 1.1
 ; Generator: rspirv
-; Bound: 172
+; Bound: 198
 OpCapability Shader
 OpExtension "SPV_KHR_storage_buffer_storage_class"
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
-OpEntryPoint GLCompute %93 "main" %116
-OpExecutionMode %93 LocalSize 1 1 1
+OpEntryPoint GLCompute %114 "main" %139
+OpExecutionMode %114 LocalSize 1 1 1
 OpDecorate %5 ArrayStride 4
 OpMemberDecorate %9 0 Offset 0
 OpMemberDecorate %9 1 Offset 12
@@ -17,42 +17,47 @@ OpDecorate %17 ArrayStride 32
 OpDecorate %19 ArrayStride 64
 OpDecorate %21 ArrayStride 32
 OpDecorate %22 ArrayStride 64
-OpDecorate %30 DescriptorSet 0
-OpDecorate %30 Binding 1
-OpDecorate %31 Block
-OpMemberDecorate %31 0 Offset 0
-OpDecorate %33 NonWritable
-OpDecorate %33 DescriptorSet 0
-OpDecorate %33 Binding 2
-OpDecorate %34 Block
-OpMemberDecorate %34 0 Offset 0
-OpDecorate %36 DescriptorSet 0
-OpDecorate %36 Binding 3
-OpDecorate %37 Block
-OpMemberDecorate %37 0 Offset 0
-OpDecorate %39 DescriptorSet 0
-OpDecorate %39 Binding 4
-OpDecorate %40 Block
-OpMemberDecorate %40 0 Offset 0
-OpDecorate %42 DescriptorSet 0
-OpDecorate %42 Binding 5
-OpDecorate %43 Block
-OpMemberDecorate %43 0 Offset 0
-OpMemberDecorate %43 0 ColMajor
-OpMemberDecorate %43 0 MatrixStride 8
-OpDecorate %45 DescriptorSet 0
-OpDecorate %45 Binding 6
-OpDecorate %46 Block
-OpMemberDecorate %46 0 Offset 0
-OpMemberDecorate %46 0 ColMajor
-OpMemberDecorate %46 0 MatrixStride 16
-OpDecorate %48 DescriptorSet 0
-OpDecorate %48 Binding 7
-OpDecorate %49 Block
-OpMemberDecorate %49 0 Offset 0
-OpMemberDecorate %49 0 ColMajor
-OpMemberDecorate %49 0 MatrixStride 8
-OpDecorate %116 BuiltIn LocalInvocationId
+OpMemberDecorate %25 0 Offset 0
+OpMemberDecorate %25 1 Offset 8
+OpMemberDecorate %25 2 Offset 16
+OpMemberDecorate %26 0 Offset 0
+OpMemberDecorate %26 1 Offset 8
+OpMemberDecorate %26 2 Offset 16
+OpMemberDecorate %26 3 Offset 24
+OpDecorate %27 ArrayStride 32
+OpDecorate %28 ArrayStride 64
+OpDecorate %34 DescriptorSet 0
+OpDecorate %34 Binding 1
+OpDecorate %35 Block
+OpMemberDecorate %35 0 Offset 0
+OpDecorate %37 NonWritable
+OpDecorate %37 DescriptorSet 0
+OpDecorate %37 Binding 2
+OpDecorate %38 Block
+OpMemberDecorate %38 0 Offset 0
+OpDecorate %40 DescriptorSet 0
+OpDecorate %40 Binding 3
+OpDecorate %41 Block
+OpMemberDecorate %41 0 Offset 0
+OpDecorate %43 DescriptorSet 0
+OpDecorate %43 Binding 4
+OpDecorate %44 Block
+OpMemberDecorate %44 0 Offset 0
+OpDecorate %46 DescriptorSet 0
+OpDecorate %46 Binding 5
+OpDecorate %47 Block
+OpMemberDecorate %47 0 Offset 0
+OpDecorate %49 DescriptorSet 0
+OpDecorate %49 Binding 6
+OpDecorate %50 Block
+OpMemberDecorate %50 0 Offset 0
+OpMemberDecorate %50 0 ColMajor
+OpMemberDecorate %50 0 MatrixStride 16
+OpDecorate %52 DescriptorSet 0
+OpDecorate %52 Binding 7
+OpDecorate %53 Block
+OpMemberDecorate %53 0 Offset 0
+OpDecorate %139 BuiltIn LocalInvocationId
 %2 = OpTypeVoid
 %3 = OpTypeBool
 %4 = OpTypeFloat 32
@@ -76,179 +81,209 @@ OpDecorate %116 BuiltIn LocalInvocationId
 %22 = OpTypeArray %21 %18
 %23 = OpTypeInt 32 1
 %24 = OpTypeMatrix %8 3
-%25 = OpConstantTrue  %3
-%27 = OpTypePointer Workgroup %5
-%26 = OpVariable  %27  Workgroup
-%29 = OpTypePointer Workgroup %7
-%28 = OpVariable  %29  Workgroup
-%31 = OpTypeStruct %9
-%32 = OpTypePointer StorageBuffer %31
-%30 = OpVariable  %32  StorageBuffer
-%34 = OpTypeStruct %11
-%35 = OpTypePointer StorageBuffer %34
-%33 = OpVariable  %35  StorageBuffer
-%37 = OpTypeStruct %13
-%38 = OpTypePointer Uniform %37
-%36 = OpVariable  %38  Uniform
-%40 = OpTypeStruct %8
-%41 = OpTypePointer Uniform %40
-%39 = OpVariable  %41  Uniform
-%43 = OpTypeStruct %15
-%44 = OpTypePointer Uniform %43
-%42 = OpVariable  %44  Uniform
-%46 = OpTypeStruct %19
-%47 = OpTypePointer Uniform %46
-%45 = OpVariable  %47  Uniform
-%49 = OpTypeStruct %22
-%50 = OpTypePointer Uniform %49
-%48 = OpVariable  %50  Uniform
-%54 = OpTypeFunction %2 %8
-%58 = OpTypeFunction %2
-%59 = OpTypePointer StorageBuffer %9
-%60 = OpConstant  %7  0
-%62 = OpConstant  %4  1
-%63 = OpConstantComposite  %8  %62 %62 %62
-%64 = OpConstant  %23  1
-%65 = OpConstant  %4  2
-%66 = OpConstant  %4  3
-%67 = OpConstantNull  %24
-%69 = OpTypePointer Function %23
-%71 = OpTypePointer StorageBuffer %8
-%73 = OpTypePointer StorageBuffer %4
-%95 = OpTypePointer StorageBuffer %11
-%97 = OpTypePointer Uniform %13
-%99 = OpTypePointer Uniform %8
-%101 = OpTypePointer Uniform %15
-%103 = OpTypePointer Uniform %19
-%105 = OpTypePointer Uniform %22
-%107 = OpConstant  %4  4
-%109 = OpTypePointer Function %4
-%111 = OpTypePointer Function %3
-%113 = OpConstantNull  %5
-%114 = OpConstantNull  %7
-%115 = OpTypeVector %7 3
-%117 = OpTypePointer Input %115
-%116 = OpVariable  %117  Input
-%119 = OpConstantNull  %115
-%120 = OpTypeVector %3 3
-%125 = OpConstant  %7  264
-%128 = OpTypePointer Workgroup %4
-%129 = OpTypePointer Uniform %21
-%130 = OpTypePointer Uniform %20
-%133 = OpTypePointer Uniform %17
-%134 = OpTypePointer Uniform %16
-%135 = OpTypePointer Uniform %12
-%140 = OpConstant  %7  7
-%146 = OpConstant  %7  6
-%148 = OpTypePointer StorageBuffer %10
-%149 = OpConstant  %7  1
-%152 = OpConstant  %7  5
-%154 = OpTypePointer Uniform %4
-%155 = OpConstant  %7  3
-%158 = OpConstant  %7  4
-%170 = OpConstant  %23  2
-%171 = OpConstant  %7  256
-%53 = OpFunction  %2  None %54
-%52 = OpFunctionParameter  %8
-%51 = OpLabel
-OpBranch %55
+%25 = OpTypeStruct %10 %10 %10
+%26 = OpTypeStruct %10 %10 %10 %10
+%27 = OpTypeArray %26 %18
+%28 = OpTypeArray %27 %18
+%29 = OpConstantTrue  %3
+%31 = OpTypePointer Workgroup %5
+%30 = OpVariable  %31  Workgroup
+%33 = OpTypePointer Workgroup %7
+%32 = OpVariable  %33  Workgroup
+%35 = OpTypeStruct %9
+%36 = OpTypePointer StorageBuffer %35
+%34 = OpVariable  %36  StorageBuffer
+%38 = OpTypeStruct %11
+%39 = OpTypePointer StorageBuffer %38
+%37 = OpVariable  %39  StorageBuffer
+%41 = OpTypeStruct %13
+%42 = OpTypePointer Uniform %41
+%40 = OpVariable  %42  Uniform
+%44 = OpTypeStruct %8
+%45 = OpTypePointer Uniform %44
+%43 = OpVariable  %45  Uniform
+%47 = OpTypeStruct %25
+%48 = OpTypePointer Uniform %47
+%46 = OpVariable  %48  Uniform
+%50 = OpTypeStruct %19
+%51 = OpTypePointer Uniform %50
+%49 = OpVariable  %51  Uniform
+%53 = OpTypeStruct %28
+%54 = OpTypePointer Uniform %53
+%52 = OpVariable  %54  Uniform
+%58 = OpTypeFunction %2 %8
+%62 = OpTypeFunction %2
+%63 = OpTypePointer StorageBuffer %9
+%64 = OpConstant  %7  0
+%66 = OpConstant  %4  1
+%67 = OpConstantComposite  %8  %66 %66 %66
+%68 = OpConstant  %23  1
+%69 = OpConstant  %4  2
+%70 = OpConstant  %4  3
+%71 = OpConstantNull  %24
+%73 = OpTypePointer Function %23
+%75 = OpTypePointer StorageBuffer %8
+%77 = OpTypePointer StorageBuffer %4
+%97 = OpTypeFunction %20 %26
+%106 = OpTypeFunction %15 %25
+%116 = OpTypePointer StorageBuffer %11
+%118 = OpTypePointer Uniform %13
+%120 = OpTypePointer Uniform %8
+%122 = OpTypePointer Uniform %25
+%124 = OpTypePointer Uniform %19
+%126 = OpTypePointer Uniform %28
+%128 = OpTypePointer Uniform %22
+%129 = OpTypePointer Uniform %15
+%130 = OpConstant  %4  4
+%132 = OpTypePointer Function %4
+%134 = OpTypePointer Function %3
+%136 = OpConstantNull  %5
+%137 = OpConstantNull  %7
+%138 = OpTypeVector %7 3
+%140 = OpTypePointer Input %138
+%139 = OpVariable  %140  Input
+%142 = OpConstantNull  %138
+%143 = OpTypeVector %3 3
+%148 = OpConstant  %7  264
+%151 = OpTypePointer Workgroup %4
+%152 = OpTypePointer Uniform %21
+%153 = OpTypePointer Uniform %20
+%154 = OpTypePointer Uniform %26
+%158 = OpTypePointer Uniform %17
+%159 = OpTypePointer Uniform %16
+%160 = OpTypePointer Uniform %12
+%165 = OpConstant  %7  7
+%172 = OpConstant  %7  6
+%174 = OpTypePointer StorageBuffer %10
+%175 = OpConstant  %7  1
+%178 = OpConstant  %7  5
+%180 = OpTypePointer Uniform %4
+%181 = OpConstant  %7  3
+%184 = OpConstant  %7  4
+%196 = OpConstant  %23  2
+%197 = OpConstant  %7  256
+%57 = OpFunction  %2  None %58
+%56 = OpFunctionParameter  %8
 %55 = OpLabel
+OpBranch %59
+%59 = OpLabel
 OpReturn
 OpFunctionEnd
-%57 = OpFunction  %2  None %58
-%56 = OpLabel
-%68 = OpVariable  %69  Function %64
-%61 = OpAccessChain  %59  %30 %60
-OpBranch %70
-%70 = OpLabel
-%72 = OpAccessChain  %71  %61 %60
-OpStore %72 %63
-%74 = OpAccessChain  %73  %61 %60 %60
-OpStore %74 %62
-%75 = OpAccessChain  %73  %61 %60 %60
-OpStore %75 %65
-%76 = OpLoad  %23  %68
-%77 = OpAccessChain  %73  %61 %60 %76
-OpStore %77 %66
-%78 = OpLoad  %9  %61
-%79 = OpCompositeExtract  %8  %78 0
-%80 = OpCompositeExtract  %8  %78 0
-%81 = OpVectorShuffle  %10  %80 %80 2 0
-%82 = OpCompositeExtract  %8  %78 0
-%83 = OpFunctionCall  %2  %53 %82
-%84 = OpCompositeExtract  %8  %78 0
-%85 = OpVectorTimesMatrix  %8  %84 %67
-%86 = OpCompositeExtract  %8  %78 0
-%87 = OpMatrixTimesVector  %8  %67 %86
-%88 = OpCompositeExtract  %8  %78 0
-%89 = OpVectorTimesScalar  %8  %88 %65
-%90 = OpCompositeExtract  %8  %78 0
-%91 = OpVectorTimesScalar  %8  %90 %65
+%61 = OpFunction  %2  None %62
+%60 = OpLabel
+%72 = OpVariable  %73  Function %68
+%65 = OpAccessChain  %63  %34 %64
+OpBranch %74
+%74 = OpLabel
+%76 = OpAccessChain  %75  %65 %64
+OpStore %76 %67
+%78 = OpAccessChain  %77  %65 %64 %64
+OpStore %78 %66
+%79 = OpAccessChain  %77  %65 %64 %64
+OpStore %79 %69
+%80 = OpLoad  %23  %72
+%81 = OpAccessChain  %77  %65 %64 %80
+OpStore %81 %70
+%82 = OpLoad  %9  %65
+%83 = OpCompositeExtract  %8  %82 0
+%84 = OpCompositeExtract  %8  %82 0
+%85 = OpVectorShuffle  %10  %84 %84 2 0
+%86 = OpCompositeExtract  %8  %82 0
+%87 = OpFunctionCall  %2  %57 %86
+%88 = OpCompositeExtract  %8  %82 0
+%89 = OpVectorTimesMatrix  %8  %88 %71
+%90 = OpCompositeExtract  %8  %82 0
+%91 = OpMatrixTimesVector  %8  %71 %90
+%92 = OpCompositeExtract  %8  %82 0
+%93 = OpVectorTimesScalar  %8  %92 %69
+%94 = OpCompositeExtract  %8  %82 0
+%95 = OpVectorTimesScalar  %8  %94 %69
 OpReturn
 OpFunctionEnd
-%93 = OpFunction  %2  None %58
-%92 = OpLabel
-%108 = OpVariable  %109  Function %62
-%110 = OpVariable  %111  Function %25
-%94 = OpAccessChain  %59  %30 %60
-%96 = OpAccessChain  %95  %33 %60
-%98 = OpAccessChain  %97  %36 %60
-%100 = OpAccessChain  %99  %39 %60
-%102 = OpAccessChain  %101  %42 %60
-%104 = OpAccessChain  %103  %45 %60
-%106 = OpAccessChain  %105  %48 %60
-OpBranch %112
-%112 = OpLabel
-%118 = OpLoad  %115  %116
-%121 = OpIEqual  %120  %118 %119
-%122 = OpAll  %3  %121
-OpSelectionMerge %123 None
-OpBranchConditional %122 %124 %123
-%124 = OpLabel
-OpStore %26 %113
-OpStore %28 %114
-OpBranch %123
-%123 = OpLabel
-OpControlBarrier %18 %18 %125
-OpBranch %126
-%126 = OpLabel
-%127 = OpFunctionCall  %2  %57
-%131 = OpAccessChain  %130  %106 %60 %60
-%132 = OpLoad  %20  %131
-%136 = OpAccessChain  %135  %104 %60 %60 %60
-%137 = OpLoad  %12  %136
-%138 = OpMatrixTimesVector  %10  %132 %137
-%139 = OpCompositeExtract  %4  %138 0
-%141 = OpAccessChain  %128  %26 %140
-OpStore %141 %139
-%142 = OpLoad  %15  %102
-%143 = OpLoad  %8  %100
-%144 = OpMatrixTimesVector  %10  %142 %143
-%145 = OpCompositeExtract  %4  %144 0
-%147 = OpAccessChain  %128  %26 %146
-OpStore %147 %145
-%150 = OpAccessChain  %73  %96 %149 %149
-%151 = OpLoad  %4  %150
-%153 = OpAccessChain  %128  %26 %152
-OpStore %153 %151
-%156 = OpAccessChain  %154  %98 %60 %155
-%157 = OpLoad  %4  %156
-%159 = OpAccessChain  %128  %26 %158
-OpStore %159 %157
-%160 = OpAccessChain  %73  %94 %149
-%161 = OpLoad  %4  %160
-%162 = OpAccessChain  %128  %26 %155
-OpStore %162 %161
-%163 = OpAccessChain  %73  %94 %60 %60
-%164 = OpLoad  %4  %163
-%165 = OpAccessChain  %128  %26 %18
-OpStore %165 %164
-%166 = OpAccessChain  %73  %94 %149
-OpStore %166 %107
-%167 = OpArrayLength  %7  %33 0
-%168 = OpConvertUToF  %4  %167
-%169 = OpAccessChain  %128  %26 %149
-OpStore %169 %168
-OpAtomicStore %28 %170 %171 %18
+%96 = OpFunction  %20  None %97
+%98 = OpFunctionParameter  %26
+%99 = OpLabel
+%100 = OpCompositeExtract  %10  %98 0
+%101 = OpCompositeExtract  %10  %98 1
+%102 = OpCompositeExtract  %10  %98 2
+%103 = OpCompositeExtract  %10  %98 3
+%104 = OpCompositeConstruct  %20  %100 %101 %102 %103
+OpReturnValue %104
+OpFunctionEnd
+%105 = OpFunction  %15  None %106
+%107 = OpFunctionParameter  %25
+%108 = OpLabel
+%109 = OpCompositeExtract  %10  %107 0
+%110 = OpCompositeExtract  %10  %107 1
+%111 = OpCompositeExtract  %10  %107 2
+%112 = OpCompositeConstruct  %15  %109 %110 %111
+OpReturnValue %112
+OpFunctionEnd
+%114 = OpFunction  %2  None %62
+%113 = OpLabel
+%131 = OpVariable  %132  Function %66
+%133 = OpVariable  %134  Function %29
+%115 = OpAccessChain  %63  %34 %64
+%117 = OpAccessChain  %116  %37 %64
+%119 = OpAccessChain  %118  %40 %64
+%121 = OpAccessChain  %120  %43 %64
+%123 = OpAccessChain  %122  %46 %64
+%125 = OpAccessChain  %124  %49 %64
+%127 = OpAccessChain  %126  %52 %64
+OpBranch %135
+%135 = OpLabel
+%141 = OpLoad  %138  %139
+%144 = OpIEqual  %143  %141 %142
+%145 = OpAll  %3  %144
+OpSelectionMerge %146 None
+OpBranchConditional %145 %147 %146
+%147 = OpLabel
+OpStore %30 %136
+OpStore %32 %137
+OpBranch %146
+%146 = OpLabel
+OpControlBarrier %18 %18 %148
+OpBranch %149
+%149 = OpLabel
+%150 = OpFunctionCall  %2  %61
+%155 = OpAccessChain  %154  %127 %64 %64
+%156 = OpLoad  %26  %155
+%157 = OpFunctionCall  %20  %96 %156
+%161 = OpAccessChain  %160  %125 %64 %64 %64
+%162 = OpLoad  %12  %161
+%163 = OpMatrixTimesVector  %10  %157 %162
+%164 = OpCompositeExtract  %4  %163 0
+%166 = OpAccessChain  %151  %30 %165
+OpStore %166 %164
+%167 = OpLoad  %25  %123
+%168 = OpFunctionCall  %15  %105 %167
+%169 = OpLoad  %8  %121
+%170 = OpMatrixTimesVector  %10  %168 %169
+%171 = OpCompositeExtract  %4  %170 0
+%173 = OpAccessChain  %151  %30 %172
+OpStore %173 %171
+%176 = OpAccessChain  %77  %117 %175 %175
+%177 = OpLoad  %4  %176
+%179 = OpAccessChain  %151  %30 %178
+OpStore %179 %177
+%182 = OpAccessChain  %180  %119 %64 %181
+%183 = OpLoad  %4  %182
+%185 = OpAccessChain  %151  %30 %184
+OpStore %185 %183
+%186 = OpAccessChain  %77  %115 %175
+%187 = OpLoad  %4  %186
+%188 = OpAccessChain  %151  %30 %181
+OpStore %188 %187
+%189 = OpAccessChain  %77  %115 %64 %64
+%190 = OpLoad  %4  %189
+%191 = OpAccessChain  %151  %30 %18
+OpStore %191 %190
+%192 = OpAccessChain  %77  %115 %175
+OpStore %192 %130
+%193 = OpArrayLength  %7  %37 0
+%194 = OpConvertUToF  %4  %193
+%195 = OpAccessChain  %151  %30 %175
+OpStore %195 %194
+OpAtomicStore %32 %196 %197 %18
 OpReturn
 OpFunctionEnd
\ No newline at end of file
diff --git a/naga/tests/out/spv/wgsl-mat_cx2.spvasm b/naga/tests/out/spv/wgsl-mat_cx2.spvasm
new file mode 100644
index 00000000000..9236d159564
--- /dev/null
+++ b/naga/tests/out/spv/wgsl-mat_cx2.spvasm
@@ -0,0 +1,964 @@
+; SPIR-V
+; Version: 1.1
+; Generator: rspirv
+; Bound: 446
+OpCapability Shader
+OpExtension "SPV_KHR_storage_buffer_storage_class"
+%1 = OpExtInstImport "GLSL.std.450"
+OpMemoryModel Logical GLSL450
+OpEntryPoint GLCompute %435 "main"
+OpExecutionMode %435 LocalSize 1 1 1
+%3 = OpString "mat_cx2.wgsl"
+OpSource Unknown 0 %3 "// Test handling of N-by-2 matrices.
+// See the doc comments on `naga::back::hlsl` and `naga::back::spv` for details.
+//
+// There are additional tests in `access.wgsl`.
+//
+// Tests that we don't apply this handling to other sizes are in mat_cx3.wgsl.
+
+// Access type (3rd item in variable names)
+// S = Struct
+// M = Matrix
+// C = Column
+// E = Element
+
+// Index type (4th item in variable names)
+// C = Constant
+// V = Variable
+
+alias Mat = mat2x2<f32>;
+
+@group(0) @binding(0)
+var<storage, read_write> s_m: Mat;
+
+@group(0) @binding(1)
+var<uniform> u_m: Mat;
+
+fn access_m() {
+    var idx = 1;
+    idx--;
+
+    // loads from storage
+    let l_s_m = s_m;
+    let l_s_c_c = s_m[0];
+    let l_s_c_v = s_m[idx];
+    let l_s_e_cc = s_m[0][0];
+    let l_s_e_cv = s_m[0][idx];
+    let l_s_e_vc = s_m[idx][0];
+    let l_s_e_vv = s_m[idx][idx];
+
+    // loads from uniform
+    let l_u_m = u_m;
+    let l_u_c_c = u_m[0];
+    let l_u_c_v = u_m[idx];
+    let l_u_e_cc = u_m[0][0];
+    let l_u_e_cv = u_m[0][idx];
+    let l_u_e_vc = u_m[idx][0];
+    let l_u_e_vv = u_m[idx][idx];
+
+    // stores to storage
+    s_m = l_u_m;
+    s_m[0] = l_u_c_c;
+    s_m[idx] = l_u_c_v;
+    s_m[0][0] = l_u_e_cc;
+    s_m[0][idx] = l_u_e_cv;
+    s_m[idx][0] = l_u_e_vc;
+    s_m[idx][idx] = l_u_e_vv;
+}
+
+struct StructWithMat {
+    m: Mat,
+}
+
+@group(1) @binding(0)
+var<storage, read_write> s_sm: StructWithMat;
+
+@group(1) @binding(1)
+var<uniform> u_sm: StructWithMat;
+
+fn access_sm() {
+    var idx = 1;
+    idx--;
+
+    // loads from storage
+    let l_s_s = s_sm;
+    let l_s_m = s_sm.m;
+    let l_s_c_c = s_sm.m[0];
+    let l_s_c_v = s_sm.m[idx];
+    let l_s_e_cc = s_sm.m[0][0];
+    let l_s_e_cv = s_sm.m[0][idx];
+    let l_s_e_vc = s_sm.m[idx][0];
+    let l_s_e_vv = s_sm.m[idx][idx];
+
+    // loads from uniform
+    let l_u_s = u_sm;
+    let l_u_m = u_sm.m;
+    let l_u_c_c = u_sm.m[0];
+    let l_u_c_v = u_sm.m[idx];
+    let l_u_e_cc = u_sm.m[0][0];
+    let l_u_e_cv = u_sm.m[0][idx];
+    let l_u_e_vc = u_sm.m[idx][0];
+    let l_u_e_vv = u_sm.m[idx][idx];
+
+    // stores to storage
+    s_sm = l_u_s;
+    s_sm.m = l_u_m;
+    s_sm.m[0] = l_u_c_c;
+    s_sm.m[idx] = l_u_c_v;
+    s_sm.m[0][0] = l_u_e_cc;
+    s_sm.m[0][idx] = l_u_e_cv;
+    s_sm.m[idx][0] = l_u_e_vc;
+    s_sm.m[idx][idx] = l_u_e_vv;
+}
+
+struct StructWithArrayOfStructOfMat {
+    a: array<StructWithMat, 4>,
+}
+
+@group(2) @binding(0)
+var<storage, read_write> s_sasm: StructWithArrayOfStructOfMat;
+
+@group(2) @binding(1)
+var<uniform> u_sasm: StructWithArrayOfStructOfMat;
+
+fn access_sasm() {
+    var idx = 1;
+    idx--;
+
+    // loads from storage
+    let l_s_s = s_sasm;
+    let l_s_a = s_sasm.a;
+    let l_s_m_c = s_sasm.a[0].m;
+    let l_s_m_v = s_sasm.a[idx].m;
+    let l_s_c_cc = s_sasm.a[0].m[0];
+    let l_s_c_cv = s_sasm.a[0].m[idx];
+    let l_s_c_vc = s_sasm.a[idx].m[0];
+    let l_s_c_vv = s_sasm.a[idx].m[idx];
+    let l_s_e_ccc = s_sasm.a[0].m[0][0];
+    let l_s_e_ccv = s_sasm.a[0].m[0][idx];
+    let l_s_e_cvc = s_sasm.a[0].m[idx][0];
+    let l_s_e_cvv = s_sasm.a[0].m[idx][idx];
+    let l_s_e_vcc = s_sasm.a[idx].m[0][0];
+    let l_s_e_vcv = s_sasm.a[idx].m[0][idx];
+    let l_s_e_vvc = s_sasm.a[idx].m[idx][0];
+    let l_s_e_vvv = s_sasm.a[idx].m[idx][idx];
+
+    // loads from uniform
+    let l_u_s = u_sasm;
+    let l_u_a = u_sasm.a;
+    let l_u_m_c = u_sasm.a[0].m;
+    let l_u_m_v = u_sasm.a[idx].m;
+    let l_u_c_cc = u_sasm.a[0].m[0];
+    let l_u_c_cv = u_sasm.a[0].m[idx];
+    let l_u_c_vc = u_sasm.a[idx].m[0];
+    let l_u_c_vv = u_sasm.a[idx].m[idx];
+    let l_u_e_ccc = u_sasm.a[0].m[0][0];
+    let l_u_e_ccv = u_sasm.a[0].m[0][idx];
+    let l_u_e_cvc = u_sasm.a[0].m[idx][0];
+    let l_u_e_cvv = u_sasm.a[0].m[idx][idx];
+    let l_u_e_vcc = u_sasm.a[idx].m[0][0];
+    let l_u_e_vcv = u_sasm.a[idx].m[0][idx];
+    let l_u_e_vvc = u_sasm.a[idx].m[idx][0];
+    let l_u_e_vvv = u_sasm.a[idx].m[idx][idx];
+
+    // stores to storage
+    s_sasm = l_u_s;
+    s_sasm.a = l_u_a;
+    s_sasm.a[0].m = l_u_m_c;
+    s_sasm.a[idx].m = l_u_m_v;
+    s_sasm.a[0].m[0] = l_u_c_cc;
+    s_sasm.a[0].m[idx] = l_u_c_cv;
+    s_sasm.a[idx].m[0] = l_u_c_vc;
+    s_sasm.a[idx].m[idx] = l_u_c_vv;
+    s_sasm.a[0].m[0][0] = l_u_e_ccc;
+    s_sasm.a[0].m[0][idx] = l_u_e_ccv;
+    s_sasm.a[0].m[idx][0] = l_u_e_cvc;
+    s_sasm.a[0].m[idx][idx] = l_u_e_cvv;
+    s_sasm.a[idx].m[0][0] = l_u_e_vcc;
+    s_sasm.a[idx].m[0][idx] = l_u_e_vcv;
+    s_sasm.a[idx].m[idx][0] = l_u_e_vvc;
+    s_sasm.a[idx].m[idx][idx] = l_u_e_vvv;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+    access_m();
+    access_sm();
+    access_sasm();
+}
+"
+OpName %4 "Mat"
+OpMemberName %8 0 "m"
+OpName %8 "StructWithMat"
+OpMemberName %12 0 "a"
+OpName %12 "StructWithArrayOfStructOfMat"
+OpMemberName %13 0 "col0"
+OpMemberName %13 1 "col1"
+OpName %13 "std140_mat2x2<f32>"
+OpMemberName %14 0 "m_col0"
+OpMemberName %14 1 "m_col1"
+OpName %14 "std140_StructWithMat"
+OpName %15 "std140_array<StructWithMat, 4>"
+OpMemberName %16 0 "a"
+OpName %16 "std140_StructWithArrayOfStructOfMat"
+OpName %17 "s_m"
+OpName %20 "u_m"
+OpName %23 "s_sm"
+OpName %26 "u_sm"
+OpName %29 "s_sasm"
+OpName %32 "u_sasm"
+OpName %35 "mat2x2<f32>_from_std140"
+OpName %42 "mat2x2<f32>_get_column"
+OpName %55 "access_m"
+OpName %64 "idx"
+OpName %129 "StructWithMat_from_std140"
+OpName %138 "access_sm"
+OpName %144 "idx"
+OpName %222 "StructWithArrayOfStructOfMat_from_std140"
+OpName %227 "array<StructWithMat, 4>_from_std140"
+OpName %243 "access_sasm"
+OpName %249 "idx"
+OpName %435 "main"
+OpMemberDecorate %8 0 Offset 0
+OpMemberDecorate %8 0 ColMajor
+OpMemberDecorate %8 0 MatrixStride 8
+OpDecorate %9 ArrayStride 16
+OpMemberDecorate %12 0 Offset 0
+OpMemberDecorate %13 0 Offset 0
+OpMemberDecorate %13 1 Offset 8
+OpMemberDecorate %14 0 Offset 0
+OpMemberDecorate %14 1 Offset 8
+OpDecorate %15 ArrayStride 16
+OpMemberDecorate %16 0 Offset 0
+OpDecorate %17 DescriptorSet 0
+OpDecorate %17 Binding 0
+OpDecorate %18 Block
+OpMemberDecorate %18 0 Offset 0
+OpMemberDecorate %18 0 ColMajor
+OpMemberDecorate %18 0 MatrixStride 8
+OpDecorate %20 DescriptorSet 0
+OpDecorate %20 Binding 1
+OpDecorate %21 Block
+OpMemberDecorate %21 0 Offset 0
+OpDecorate %23 DescriptorSet 1
+OpDecorate %23 Binding 0
+OpDecorate %24 Block
+OpMemberDecorate %24 0 Offset 0
+OpDecorate %26 DescriptorSet 1
+OpDecorate %26 Binding 1
+OpDecorate %27 Block
+OpMemberDecorate %27 0 Offset 0
+OpDecorate %29 DescriptorSet 2
+OpDecorate %29 Binding 0
+OpDecorate %30 Block
+OpMemberDecorate %30 0 Offset 0
+OpDecorate %32 DescriptorSet 2
+OpDecorate %32 Binding 1
+OpDecorate %33 Block
+OpMemberDecorate %33 0 Offset 0
+%2 = OpTypeVoid
+%6 = OpTypeFloat 32
+%5 = OpTypeVector %6 2
+%4 = OpTypeMatrix %5 2
+%7 = OpTypeInt 32 1
+%8 = OpTypeStruct %4
+%11 = OpTypeInt 32 0
+%10 = OpConstant  %11  4
+%9 = OpTypeArray %8 %10
+%12 = OpTypeStruct %9
+%13 = OpTypeStruct %5 %5
+%14 = OpTypeStruct %5 %5
+%15 = OpTypeArray %14 %10
+%16 = OpTypeStruct %15
+%18 = OpTypeStruct %4
+%19 = OpTypePointer StorageBuffer %18
+%17 = OpVariable  %19  StorageBuffer
+%21 = OpTypeStruct %13
+%22 = OpTypePointer Uniform %21
+%20 = OpVariable  %22  Uniform
+%24 = OpTypeStruct %8
+%25 = OpTypePointer StorageBuffer %24
+%23 = OpVariable  %25  StorageBuffer
+%27 = OpTypeStruct %14
+%28 = OpTypePointer Uniform %27
+%26 = OpVariable  %28  Uniform
+%30 = OpTypeStruct %12
+%31 = OpTypePointer StorageBuffer %30
+%29 = OpVariable  %31  StorageBuffer
+%33 = OpTypeStruct %16
+%34 = OpTypePointer Uniform %33
+%32 = OpVariable  %34  Uniform
+%36 = OpTypeFunction %4 %13
+%45 = OpTypeFunction %5 %4 %11
+%56 = OpTypeFunction %2
+%57 = OpTypePointer StorageBuffer %4
+%58 = OpConstant  %11  0
+%60 = OpTypePointer Uniform %13
+%62 = OpConstant  %7  1
+%63 = OpTypePointer Uniform %4
+%65 = OpTypePointer Function %7
+%70 = OpTypePointer StorageBuffer %5
+%76 = OpTypePointer StorageBuffer %6
+%91 = OpTypePointer Uniform %5
+%99 = OpTypePointer Uniform %6
+%130 = OpTypeFunction %8 %14
+%139 = OpTypePointer StorageBuffer %8
+%141 = OpTypePointer Uniform %14
+%143 = OpTypePointer Uniform %8
+%170 = OpConstant  %11  1
+%223 = OpTypeFunction %12 %16
+%228 = OpTypeFunction %9 %15
+%244 = OpTypePointer StorageBuffer %12
+%246 = OpTypePointer Uniform %16
+%248 = OpTypePointer Uniform %12
+%254 = OpTypePointer StorageBuffer %9
+%304 = OpTypePointer Uniform %9
+%305 = OpTypePointer Uniform %15
+%35 = OpFunction  %4  None %36
+%37 = OpFunctionParameter  %13
+%38 = OpLabel
+%39 = OpCompositeExtract  %5  %37 0
+%40 = OpCompositeExtract  %5  %37 1
+%41 = OpCompositeConstruct  %4  %39 %40
+OpReturnValue %41
+OpFunctionEnd
+%42 = OpFunction  %5  None %45
+%43 = OpFunctionParameter  %4
+%44 = OpFunctionParameter  %11
+%46 = OpLabel
+OpSelectionMerge %47 None
+OpSwitch %44 %50 0 %48 1 %49
+%48 = OpLabel
+%51 = OpCompositeExtract  %5  %43 0
+OpBranch %47
+%49 = OpLabel
+%52 = OpCompositeExtract  %5  %43 1
+OpBranch %47
+%50 = OpLabel
+OpUnreachable
+%47 = OpLabel
+%53 = OpPhi  %5  %51 %48 %52 %49
+OpReturnValue %53
+OpFunctionEnd
+%55 = OpFunction  %2  None %56
+%54 = OpLabel
+%64 = OpVariable  %65  Function %62
+%59 = OpAccessChain  %57  %17 %58
+%61 = OpAccessChain  %60  %20 %58
+OpBranch %66
+%66 = OpLabel
+OpLine %3 28 5
+%67 = OpLoad  %7  %64
+%68 = OpISub  %7  %67 %62
+OpLine %3 28 5
+OpStore %64 %68
+OpLine %3 31 17
+%69 = OpLoad  %4  %59
+OpLine %3 32 19
+%71 = OpAccessChain  %70  %59 %58
+%72 = OpLoad  %5  %71
+OpLine %3 33 19
+%73 = OpLoad  %7  %64
+%74 = OpAccessChain  %70  %59 %73
+%75 = OpLoad  %5  %74
+OpLine %3 34 20
+OpLine %3 34 20
+%77 = OpAccessChain  %76  %59 %58 %58
+%78 = OpLoad  %6  %77
+OpLine %3 35 20
+%79 = OpLoad  %7  %64
+%80 = OpAccessChain  %76  %59 %58 %79
+%81 = OpLoad  %6  %80
+OpLine %3 36 20
+%82 = OpLoad  %7  %64
+OpLine %3 36 20
+%83 = OpAccessChain  %76  %59 %82 %58
+%84 = OpLoad  %6  %83
+OpLine %3 37 20
+%85 = OpLoad  %7  %64
+%86 = OpLoad  %7  %64
+%87 = OpAccessChain  %76  %59 %85 %86
+%88 = OpLoad  %6  %87
+OpLine %3 40 17
+%89 = OpLoad  %13  %61
+%90 = OpFunctionCall  %4  %35 %89
+OpLine %3 41 19
+%92 = OpAccessChain  %91  %61 %58
+%93 = OpLoad  %5  %92
+OpLine %3 42 19
+%94 = OpLoad  %7  %64
+%95 = OpLoad  %13  %61
+%96 = OpFunctionCall  %4  %35 %95
+%97 = OpBitcast  %11  %94
+%98 = OpFunctionCall  %5  %42 %96 %97
+OpLine %3 43 20
+OpLine %3 43 20
+%100 = OpAccessChain  %99  %61 %58 %58
+%101 = OpLoad  %6  %100
+OpLine %3 44 20
+%102 = OpLoad  %7  %64
+%103 = OpAccessChain  %99  %61 %58 %102
+%104 = OpLoad  %6  %103
+OpLine %3 45 20
+%105 = OpLoad  %7  %64
+OpLine %3 45 20
+%106 = OpLoad  %13  %61
+%107 = OpFunctionCall  %4  %35 %106
+%108 = OpBitcast  %11  %105
+%109 = OpFunctionCall  %5  %42 %107 %108
+%110 = OpCompositeExtract  %6  %109 0
+OpLine %3 46 20
+%111 = OpLoad  %7  %64
+%112 = OpLoad  %7  %64
+%113 = OpLoad  %13  %61
+%114 = OpFunctionCall  %4  %35 %113
+%115 = OpBitcast  %11  %111
+%116 = OpFunctionCall  %5  %42 %114 %115
+%117 = OpVectorExtractDynamic  %6  %116 %112
+OpLine %3 49 5
+OpStore %59 %90
+OpLine %3 50 5
+OpLine %3 50 5
+%118 = OpAccessChain  %70  %59 %58
+OpStore %118 %93
+OpLine %3 51 5
+%119 = OpLoad  %7  %64
+OpLine %3 51 5
+%120 = OpAccessChain  %70  %59 %119
+OpStore %120 %98
+OpLine %3 52 5
+OpLine %3 52 5
+OpLine %3 52 5
+%121 = OpAccessChain  %76  %59 %58 %58
+OpStore %121 %101
+OpLine %3 53 5
+%122 = OpLoad  %7  %64
+OpLine %3 53 5
+%123 = OpAccessChain  %76  %59 %58 %122
+OpStore %123 %104
+OpLine %3 54 5
+%124 = OpLoad  %7  %64
+OpLine %3 54 5
+OpLine %3 54 5
+%125 = OpAccessChain  %76  %59 %124 %58
+OpStore %125 %110
+OpLine %3 55 5
+%126 = OpLoad  %7  %64
+%127 = OpLoad  %7  %64
+OpLine %3 55 5
+%128 = OpAccessChain  %76  %59 %126 %127
+OpStore %128 %117
+OpReturn
+OpFunctionEnd
+%129 = OpFunction  %8  None %130
+%131 = OpFunctionParameter  %14
+%132 = OpLabel
+%134 = OpCompositeExtract  %5  %131 0
+%135 = OpCompositeExtract  %5  %131 1
+%133 = OpCompositeConstruct  %4  %134 %135
+%136 = OpCompositeConstruct  %8  %133
+OpReturnValue %136
+OpFunctionEnd
+%138 = OpFunction  %2  None %56
+%137 = OpLabel
+%144 = OpVariable  %65  Function %62
+%140 = OpAccessChain  %139  %23 %58
+%142 = OpAccessChain  %141  %26 %58
+OpBranch %145
+%145 = OpLabel
+OpLine %3 70 5
+%146 = OpLoad  %7  %144
+%147 = OpISub  %7  %146 %62
+OpLine %3 70 5
+OpStore %144 %147
+OpLine %3 73 17
+%148 = OpLoad  %8  %140
+OpLine %3 74 17
+%149 = OpAccessChain  %57  %140 %58
+%150 = OpLoad  %4  %149
+OpLine %3 75 19
+OpLine %3 75 19
+%151 = OpAccessChain  %70  %140 %58 %58
+%152 = OpLoad  %5  %151
+OpLine %3 76 19
+%153 = OpLoad  %7  %144
+%154 = OpAccessChain  %70  %140 %58 %153
+%155 = OpLoad  %5  %154
+OpLine %3 77 20
+OpLine %3 77 20
+OpLine %3 77 20
+%156 = OpAccessChain  %76  %140 %58 %58 %58
+%157 = OpLoad  %6  %156
+OpLine %3 78 20
+OpLine %3 78 20
+%158 = OpLoad  %7  %144
+%159 = OpAccessChain  %76  %140 %58 %58 %158
+%160 = OpLoad  %6  %159
+OpLine %3 79 20
+%161 = OpLoad  %7  %144
+OpLine %3 79 20
+%162 = OpAccessChain  %76  %140 %58 %161 %58
+%163 = OpLoad  %6  %162
+OpLine %3 80 20
+%164 = OpLoad  %7  %144
+%165 = OpLoad  %7  %144
+%166 = OpAccessChain  %76  %140 %58 %164 %165
+%167 = OpLoad  %6  %166
+OpLine %3 83 17
+%168 = OpLoad  %14  %142
+%169 = OpFunctionCall  %8  %129 %168
+OpLine %3 84 17
+%171 = OpAccessChain  %91  %142 %58
+%172 = OpLoad  %5  %171
+%173 = OpAccessChain  %91  %142 %170
+%174 = OpLoad  %5  %173
+%175 = OpCompositeConstruct  %4  %172 %174
+OpLine %3 85 19
+OpLine %3 85 19
+%176 = OpAccessChain  %91  %142 %58
+%177 = OpLoad  %5  %176
+OpLine %3 86 19
+%178 = OpLoad  %7  %144
+%179 = OpAccessChain  %91  %142 %58
+%180 = OpLoad  %5  %179
+%181 = OpAccessChain  %91  %142 %170
+%182 = OpLoad  %5  %181
+%183 = OpCompositeConstruct  %4  %180 %182
+%184 = OpBitcast  %11  %178
+%185 = OpFunctionCall  %5  %42 %183 %184
+OpLine %3 87 20
+OpLine %3 87 20
+OpLine %3 87 20
+%186 = OpAccessChain  %99  %142 %58 %58
+%187 = OpLoad  %6  %186
+OpLine %3 88 20
+OpLine %3 88 20
+%188 = OpLoad  %7  %144
+%189 = OpAccessChain  %99  %142 %58 %188
+%190 = OpLoad  %6  %189
+OpLine %3 89 20
+%191 = OpLoad  %7  %144
+OpLine %3 89 20
+%192 = OpAccessChain  %91  %142 %58
+%193 = OpLoad  %5  %192
+%194 = OpAccessChain  %91  %142 %170
+%195 = OpLoad  %5  %194
+%196 = OpCompositeConstruct  %4  %193 %195
+%197 = OpBitcast  %11  %191
+%198 = OpFunctionCall  %5  %42 %196 %197
+%199 = OpCompositeExtract  %6  %198 0
+OpLine %3 90 20
+%200 = OpLoad  %7  %144
+%201 = OpLoad  %7  %144
+%202 = OpAccessChain  %91  %142 %58
+%203 = OpLoad  %5  %202
+%204 = OpAccessChain  %91  %142 %170
+%205 = OpLoad  %5  %204
+%206 = OpCompositeConstruct  %4  %203 %205
+%207 = OpBitcast  %11  %200
+%208 = OpFunctionCall  %5  %42 %206 %207
+%209 = OpVectorExtractDynamic  %6  %208 %201
+OpLine %3 93 5
+OpStore %140 %169
+OpLine %3 94 5
+OpLine %3 94 5
+%210 = OpAccessChain  %57  %140 %58
+OpStore %210 %175
+OpLine %3 95 5
+OpLine %3 95 5
+OpLine %3 95 5
+%211 = OpAccessChain  %70  %140 %58 %58
+OpStore %211 %177
+OpLine %3 96 5
+%212 = OpLoad  %7  %144
+OpLine %3 96 5
+%213 = OpAccessChain  %70  %140 %58 %212
+OpStore %213 %185
+OpLine %3 97 5
+OpLine %3 97 5
+OpLine %3 97 5
+OpLine %3 97 5
+%214 = OpAccessChain  %76  %140 %58 %58 %58
+OpStore %214 %187
+OpLine %3 98 5
+OpLine %3 98 5
+%215 = OpLoad  %7  %144
+OpLine %3 98 5
+%216 = OpAccessChain  %76  %140 %58 %58 %215
+OpStore %216 %190
+OpLine %3 99 5
+%217 = OpLoad  %7  %144
+OpLine %3 99 5
+OpLine %3 99 5
+%218 = OpAccessChain  %76  %140 %58 %217 %58
+OpStore %218 %199
+OpLine %3 100 5
+%219 = OpLoad  %7  %144
+%220 = OpLoad  %7  %144
+OpLine %3 100 5
+%221 = OpAccessChain  %76  %140 %58 %219 %220
+OpStore %221 %209
+OpReturn
+OpFunctionEnd
+%227 = OpFunction  %9  None %228
+%229 = OpFunctionParameter  %15
+%230 = OpLabel
+%231 = OpCompositeExtract  %14  %229 0
+%232 = OpFunctionCall  %8  %129 %231
+%233 = OpCompositeExtract  %14  %229 1
+%234 = OpFunctionCall  %8  %129 %233
+%235 = OpCompositeExtract  %14  %229 2
+%236 = OpFunctionCall  %8  %129 %235
+%237 = OpCompositeExtract  %14  %229 3
+%238 = OpFunctionCall  %8  %129 %237
+%239 = OpCompositeConstruct  %9  %232 %234 %236 %238
+OpReturnValue %239
+OpFunctionEnd
+%222 = OpFunction  %12  None %223
+%224 = OpFunctionParameter  %16
+%225 = OpLabel
+%240 = OpCompositeExtract  %15  %224 0
+%226 = OpFunctionCall  %9  %227 %240
+%241 = OpCompositeConstruct  %12  %226
+OpReturnValue %241
+OpFunctionEnd
+%243 = OpFunction  %2  None %56
+%242 = OpLabel
+%249 = OpVariable  %65  Function %62
+%245 = OpAccessChain  %244  %29 %58
+%247 = OpAccessChain  %246  %32 %58
+OpBranch %250
+%250 = OpLabel
+OpLine %3 115 5
+%251 = OpLoad  %7  %249
+%252 = OpISub  %7  %251 %62
+OpLine %3 115 5
+OpStore %249 %252
+OpLine %3 118 17
+%253 = OpLoad  %12  %245
+OpLine %3 119 17
+%255 = OpAccessChain  %254  %245 %58
+%256 = OpLoad  %9  %255
+OpLine %3 120 19
+OpLine %3 120 19
+%257 = OpAccessChain  %57  %245 %58 %58 %58
+%258 = OpLoad  %4  %257
+OpLine %3 121 19
+%259 = OpLoad  %7  %249
+%260 = OpAccessChain  %57  %245 %58 %259 %58
+%261 = OpLoad  %4  %260
+OpLine %3 122 20
+OpLine %3 122 20
+OpLine %3 122 20
+%262 = OpAccessChain  %70  %245 %58 %58 %58 %58
+%263 = OpLoad  %5  %262
+OpLine %3 123 20
+OpLine %3 123 20
+%264 = OpLoad  %7  %249
+%265 = OpAccessChain  %70  %245 %58 %58 %58 %264
+%266 = OpLoad  %5  %265
+OpLine %3 124 20
+%267 = OpLoad  %7  %249
+OpLine %3 124 20
+%268 = OpAccessChain  %70  %245 %58 %267 %58 %58
+%269 = OpLoad  %5  %268
+OpLine %3 125 20
+%270 = OpLoad  %7  %249
+%271 = OpLoad  %7  %249
+%272 = OpAccessChain  %70  %245 %58 %270 %58 %271
+%273 = OpLoad  %5  %272
+OpLine %3 126 21
+OpLine %3 126 21
+OpLine %3 126 21
+OpLine %3 126 21
+%274 = OpAccessChain  %76  %245 %58 %58 %58 %58 %58
+%275 = OpLoad  %6  %274
+OpLine %3 127 21
+OpLine %3 127 21
+OpLine %3 127 21
+%276 = OpLoad  %7  %249
+%277 = OpAccessChain  %76  %245 %58 %58 %58 %58 %276
+%278 = OpLoad  %6  %277
+OpLine %3 128 21
+OpLine %3 128 21
+%279 = OpLoad  %7  %249
+OpLine %3 128 21
+%280 = OpAccessChain  %76  %245 %58 %58 %58 %279 %58
+%281 = OpLoad  %6  %280
+OpLine %3 129 21
+OpLine %3 129 21
+%282 = OpLoad  %7  %249
+%283 = OpLoad  %7  %249
+%284 = OpAccessChain  %76  %245 %58 %58 %58 %282 %283
+%285 = OpLoad  %6  %284
+OpLine %3 130 21
+%286 = OpLoad  %7  %249
+OpLine %3 130 21
+OpLine %3 130 21
+%287 = OpAccessChain  %76  %245 %58 %286 %58 %58 %58
+%288 = OpLoad  %6  %287
+OpLine %3 131 21
+%289 = OpLoad  %7  %249
+OpLine %3 131 21
+%290 = OpLoad  %7  %249
+%291 = OpAccessChain  %76  %245 %58 %289 %58 %58 %290
+%292 = OpLoad  %6  %291
+OpLine %3 132 21
+%293 = OpLoad  %7  %249
+%294 = OpLoad  %7  %249
+OpLine %3 132 21
+%295 = OpAccessChain  %76  %245 %58 %293 %58 %294 %58
+%296 = OpLoad  %6  %295
+OpLine %3 133 21
+%297 = OpLoad  %7  %249
+%298 = OpLoad  %7  %249
+%299 = OpLoad  %7  %249
+%300 = OpAccessChain  %76  %245 %58 %297 %58 %298 %299
+%301 = OpLoad  %6  %300
+OpLine %3 136 17
+%302 = OpLoad  %16  %247
+%303 = OpFunctionCall  %12  %222 %302
+OpLine %3 137 17
+%306 = OpAccessChain  %305  %247 %58
+%307 = OpLoad  %15  %306
+%308 = OpFunctionCall  %9  %227 %307
+OpLine %3 138 19
+OpLine %3 138 19
+%309 = OpAccessChain  %141  %247 %58 %58
+%310 = OpAccessChain  %91  %309 %58
+%311 = OpLoad  %5  %310
+%312 = OpAccessChain  %91  %309 %170
+%313 = OpLoad  %5  %312
+%314 = OpCompositeConstruct  %4  %311 %313
+OpLine %3 139 19
+%315 = OpLoad  %7  %249
+%316 = OpAccessChain  %141  %247 %58 %315
+%317 = OpAccessChain  %91  %316 %58
+%318 = OpLoad  %5  %317
+%319 = OpAccessChain  %91  %316 %170
+%320 = OpLoad  %5  %319
+%321 = OpCompositeConstruct  %4  %318 %320
+OpLine %3 140 20
+OpLine %3 140 20
+OpLine %3 140 20
+%322 = OpAccessChain  %91  %247 %58 %58 %58
+%323 = OpLoad  %5  %322
+OpLine %3 141 20
+OpLine %3 141 20
+%324 = OpLoad  %7  %249
+%325 = OpAccessChain  %141  %247 %58 %58
+%326 = OpAccessChain  %91  %325 %58
+%327 = OpLoad  %5  %326
+%328 = OpAccessChain  %91  %325 %170
+%329 = OpLoad  %5  %328
+%330 = OpCompositeConstruct  %4  %327 %329
+%331 = OpBitcast  %11  %324
+%332 = OpFunctionCall  %5  %42 %330 %331
+OpLine %3 142 20
+%333 = OpLoad  %7  %249
+OpLine %3 142 20
+%334 = OpAccessChain  %91  %247 %58 %333 %58
+%335 = OpLoad  %5  %334
+OpLine %3 143 20
+%336 = OpLoad  %7  %249
+%337 = OpLoad  %7  %249
+%338 = OpAccessChain  %141  %247 %58 %336
+%339 = OpAccessChain  %91  %338 %58
+%340 = OpLoad  %5  %339
+%341 = OpAccessChain  %91  %338 %170
+%342 = OpLoad  %5  %341
+%343 = OpCompositeConstruct  %4  %340 %342
+%344 = OpBitcast  %11  %337
+%345 = OpFunctionCall  %5  %42 %343 %344
+OpLine %3 144 21
+OpLine %3 144 21
+OpLine %3 144 21
+OpLine %3 144 21
+%346 = OpAccessChain  %99  %247 %58 %58 %58 %58
+%347 = OpLoad  %6  %346
+OpLine %3 145 21
+OpLine %3 145 21
+OpLine %3 145 21
+%348 = OpLoad  %7  %249
+%349 = OpAccessChain  %99  %247 %58 %58 %58 %348
+%350 = OpLoad  %6  %349
+OpLine %3 146 21
+OpLine %3 146 21
+%351 = OpLoad  %7  %249
+OpLine %3 146 21
+%352 = OpAccessChain  %141  %247 %58 %58
+%353 = OpAccessChain  %91  %352 %58
+%354 = OpLoad  %5  %353
+%355 = OpAccessChain  %91  %352 %170
+%356 = OpLoad  %5  %355
+%357 = OpCompositeConstruct  %4  %354 %356
+%358 = OpBitcast  %11  %351
+%359 = OpFunctionCall  %5  %42 %357 %358
+%360 = OpCompositeExtract  %6  %359 0
+OpLine %3 147 21
+OpLine %3 147 21
+%361 = OpLoad  %7  %249
+%362 = OpLoad  %7  %249
+%363 = OpAccessChain  %141  %247 %58 %58
+%364 = OpAccessChain  %91  %363 %58
+%365 = OpLoad  %5  %364
+%366 = OpAccessChain  %91  %363 %170
+%367 = OpLoad  %5  %366
+%368 = OpCompositeConstruct  %4  %365 %367
+%369 = OpBitcast  %11  %361
+%370 = OpFunctionCall  %5  %42 %368 %369
+%371 = OpVectorExtractDynamic  %6  %370 %362
+OpLine %3 148 21
+%372 = OpLoad  %7  %249
+OpLine %3 148 21
+OpLine %3 148 21
+%373 = OpAccessChain  %99  %247 %58 %372 %58 %58
+%374 = OpLoad  %6  %373
+OpLine %3 149 21
+%375 = OpLoad  %7  %249
+OpLine %3 149 21
+%376 = OpLoad  %7  %249
+%377 = OpAccessChain  %99  %247 %58 %375 %58 %376
+%378 = OpLoad  %6  %377
+OpLine %3 150 21
+%379 = OpLoad  %7  %249
+%380 = OpLoad  %7  %249
+OpLine %3 150 21
+%381 = OpAccessChain  %141  %247 %58 %379
+%382 = OpAccessChain  %91  %381 %58
+%383 = OpLoad  %5  %382
+%384 = OpAccessChain  %91  %381 %170
+%385 = OpLoad  %5  %384
+%386 = OpCompositeConstruct  %4  %383 %385
+%387 = OpBitcast  %11  %380
+%388 = OpFunctionCall  %5  %42 %386 %387
+%389 = OpCompositeExtract  %6  %388 0
+OpLine %3 151 21
+%390 = OpLoad  %7  %249
+%391 = OpLoad  %7  %249
+%392 = OpLoad  %7  %249
+%393 = OpAccessChain  %141  %247 %58 %390
+%394 = OpAccessChain  %91  %393 %58
+%395 = OpLoad  %5  %394
+%396 = OpAccessChain  %91  %393 %170
+%397 = OpLoad  %5  %396
+%398 = OpCompositeConstruct  %4  %395 %397
+%399 = OpBitcast  %11  %391
+%400 = OpFunctionCall  %5  %42 %398 %399
+%401 = OpVectorExtractDynamic  %6  %400 %392
+OpLine %3 154 5
+OpStore %245 %303
+OpLine %3 155 5
+OpLine %3 155 5
+%402 = OpAccessChain  %254  %245 %58
+OpStore %402 %308
+OpLine %3 156 5
+OpLine %3 156 5
+OpLine %3 156 5
+%403 = OpAccessChain  %57  %245 %58 %58 %58
+OpStore %403 %314
+OpLine %3 157 5
+%404 = OpLoad  %7  %249
+OpLine %3 157 5
+%405 = OpAccessChain  %57  %245 %58 %404 %58
+OpStore %405 %321
+OpLine %3 158 5
+OpLine %3 158 5
+OpLine %3 158 5
+OpLine %3 158 5
+%406 = OpAccessChain  %70  %245 %58 %58 %58 %58
+OpStore %406 %323
+OpLine %3 159 5
+OpLine %3 159 5
+%407 = OpLoad  %7  %249
+OpLine %3 159 5
+%408 = OpAccessChain  %70  %245 %58 %58 %58 %407
+OpStore %408 %332
+OpLine %3 160 5
+%409 = OpLoad  %7  %249
+OpLine %3 160 5
+OpLine %3 160 5
+%410 = OpAccessChain  %70  %245 %58 %409 %58 %58
+OpStore %410 %335
+OpLine %3 161 5
+%411 = OpLoad  %7  %249
+%412 = OpLoad  %7  %249
+OpLine %3 161 5
+%413 = OpAccessChain  %70  %245 %58 %411 %58 %412
+OpStore %413 %345
+OpLine %3 162 5
+OpLine %3 162 5
+OpLine %3 162 5
+OpLine %3 162 5
+OpLine %3 162 5
+%414 = OpAccessChain  %76  %245 %58 %58 %58 %58 %58
+OpStore %414 %347
+OpLine %3 163 5
+OpLine %3 163 5
+OpLine %3 163 5
+%415 = OpLoad  %7  %249
+OpLine %3 163 5
+%416 = OpAccessChain  %76  %245 %58 %58 %58 %58 %415
+OpStore %416 %350
+OpLine %3 164 5
+OpLine %3 164 5
+%417 = OpLoad  %7  %249
+OpLine %3 164 5
+OpLine %3 164 5
+%418 = OpAccessChain  %76  %245 %58 %58 %58 %417 %58
+OpStore %418 %360
+OpLine %3 165 5
+OpLine %3 165 5
+%419 = OpLoad  %7  %249
+%420 = OpLoad  %7  %249
+OpLine %3 165 5
+%421 = OpAccessChain  %76  %245 %58 %58 %58 %419 %420
+OpStore %421 %371
+OpLine %3 166 5
+%422 = OpLoad  %7  %249
+OpLine %3 166 5
+OpLine %3 166 5
+OpLine %3 166 5
+%423 = OpAccessChain  %76  %245 %58 %422 %58 %58 %58
+OpStore %423 %374
+OpLine %3 167 5
+%424 = OpLoad  %7  %249
+OpLine %3 167 5
+%425 = OpLoad  %7  %249
+OpLine %3 167 5
+%426 = OpAccessChain  %76  %245 %58 %424 %58 %58 %425
+OpStore %426 %378
+OpLine %3 168 5
+%427 = OpLoad  %7  %249
+%428 = OpLoad  %7  %249
+OpLine %3 168 5
+OpLine %3 168 5
+%429 = OpAccessChain  %76  %245 %58 %427 %58 %428 %58
+OpStore %429 %389
+OpLine %3 169 5
+%430 = OpLoad  %7  %249
+%431 = OpLoad  %7  %249
+%432 = OpLoad  %7  %249
+OpLine %3 169 5
+%433 = OpAccessChain  %76  %245 %58 %430 %58 %431 %432
+OpStore %433 %401
+OpReturn
+OpFunctionEnd
+%435 = OpFunction  %2  None %56
+%434 = OpLabel
+%436 = OpAccessChain  %57  %17 %58
+%437 = OpAccessChain  %60  %20 %58
+%438 = OpAccessChain  %139  %23 %58
+%439 = OpAccessChain  %141  %26 %58
+%440 = OpAccessChain  %244  %29 %58
+%441 = OpAccessChain  %246  %32 %58
+OpBranch %442
+%442 = OpLabel
+OpLine %3 174 5
+%443 = OpFunctionCall  %2  %55
+OpLine %3 175 5
+%444 = OpFunctionCall  %2  %138
+OpLine %3 176 5
+%445 = OpFunctionCall  %2  %243
+OpReturn
+OpFunctionEnd
\ No newline at end of file
diff --git a/naga/tests/out/spv/wgsl-mat_cx3.spvasm b/naga/tests/out/spv/wgsl-mat_cx3.spvasm
new file mode 100644
index 00000000000..3fe4bee6c37
--- /dev/null
+++ b/naga/tests/out/spv/wgsl-mat_cx3.spvasm
@@ -0,0 +1,396 @@
+; SPIR-V
+; Version: 1.1
+; Generator: rspirv
+; Bound: 309
+OpCapability Shader
+OpExtension "SPV_KHR_storage_buffer_storage_class"
+%1 = OpExtInstImport "GLSL.std.450"
+OpMemoryModel Logical GLSL450
+OpEntryPoint GLCompute %298 "main"
+OpExecutionMode %298 LocalSize 1 1 1
+OpMemberDecorate %7 0 Offset 0
+OpMemberDecorate %7 0 ColMajor
+OpMemberDecorate %7 0 MatrixStride 16
+OpDecorate %8 ArrayStride 48
+OpMemberDecorate %11 0 Offset 0
+OpDecorate %12 DescriptorSet 0
+OpDecorate %12 Binding 0
+OpDecorate %13 Block
+OpMemberDecorate %13 0 Offset 0
+OpMemberDecorate %13 0 ColMajor
+OpMemberDecorate %13 0 MatrixStride 16
+OpDecorate %15 DescriptorSet 0
+OpDecorate %15 Binding 1
+OpDecorate %16 Block
+OpMemberDecorate %16 0 Offset 0
+OpMemberDecorate %16 0 ColMajor
+OpMemberDecorate %16 0 MatrixStride 16
+OpDecorate %18 DescriptorSet 1
+OpDecorate %18 Binding 0
+OpDecorate %19 Block
+OpMemberDecorate %19 0 Offset 0
+OpDecorate %21 DescriptorSet 1
+OpDecorate %21 Binding 1
+OpDecorate %22 Block
+OpMemberDecorate %22 0 Offset 0
+OpDecorate %24 DescriptorSet 2
+OpDecorate %24 Binding 0
+OpDecorate %25 Block
+OpMemberDecorate %25 0 Offset 0
+OpDecorate %27 DescriptorSet 2
+OpDecorate %27 Binding 1
+OpDecorate %28 Block
+OpMemberDecorate %28 0 Offset 0
+%2 = OpTypeVoid
+%5 = OpTypeFloat 32
+%4 = OpTypeVector %5 3
+%3 = OpTypeMatrix %4 3
+%6 = OpTypeInt 32 1
+%7 = OpTypeStruct %3
+%10 = OpTypeInt 32 0
+%9 = OpConstant  %10  4
+%8 = OpTypeArray %7 %9
+%11 = OpTypeStruct %8
+%13 = OpTypeStruct %3
+%14 = OpTypePointer StorageBuffer %13
+%12 = OpVariable  %14  StorageBuffer
+%16 = OpTypeStruct %3
+%17 = OpTypePointer Uniform %16
+%15 = OpVariable  %17  Uniform
+%19 = OpTypeStruct %7
+%20 = OpTypePointer StorageBuffer %19
+%18 = OpVariable  %20  StorageBuffer
+%22 = OpTypeStruct %7
+%23 = OpTypePointer Uniform %22
+%21 = OpVariable  %23  Uniform
+%25 = OpTypeStruct %11
+%26 = OpTypePointer StorageBuffer %25
+%24 = OpVariable  %26  StorageBuffer
+%28 = OpTypeStruct %11
+%29 = OpTypePointer Uniform %28
+%27 = OpVariable  %29  Uniform
+%32 = OpTypeFunction %2
+%33 = OpTypePointer StorageBuffer %3
+%34 = OpConstant  %10  0
+%36 = OpTypePointer Uniform %3
+%38 = OpConstant  %6  1
+%40 = OpTypePointer Function %6
+%45 = OpTypePointer StorageBuffer %4
+%51 = OpTypePointer StorageBuffer %5
+%65 = OpTypePointer Uniform %4
+%71 = OpTypePointer Uniform %5
+%97 = OpTypePointer StorageBuffer %7
+%99 = OpTypePointer Uniform %7
+%159 = OpTypePointer StorageBuffer %11
+%161 = OpTypePointer Uniform %11
+%168 = OpTypePointer StorageBuffer %8
+%217 = OpTypePointer Uniform %8
+%31 = OpFunction  %2  None %32
+%30 = OpLabel
+%39 = OpVariable  %40  Function %38
+%35 = OpAccessChain  %33  %12 %34
+%37 = OpAccessChain  %36  %15 %34
+OpBranch %41
+%41 = OpLabel
+%42 = OpLoad  %6  %39
+%43 = OpISub  %6  %42 %38
+OpStore %39 %43
+%44 = OpLoad  %3  %35
+%46 = OpAccessChain  %45  %35 %34
+%47 = OpLoad  %4  %46
+%48 = OpLoad  %6  %39
+%49 = OpAccessChain  %45  %35 %48
+%50 = OpLoad  %4  %49
+%52 = OpAccessChain  %51  %35 %34 %34
+%53 = OpLoad  %5  %52
+%54 = OpLoad  %6  %39
+%55 = OpAccessChain  %51  %35 %34 %54
+%56 = OpLoad  %5  %55
+%57 = OpLoad  %6  %39
+%58 = OpAccessChain  %51  %35 %57 %34
+%59 = OpLoad  %5  %58
+%60 = OpLoad  %6  %39
+%61 = OpLoad  %6  %39
+%62 = OpAccessChain  %51  %35 %60 %61
+%63 = OpLoad  %5  %62
+%64 = OpLoad  %3  %37
+%66 = OpAccessChain  %65  %37 %34
+%67 = OpLoad  %4  %66
+%68 = OpLoad  %6  %39
+%69 = OpAccessChain  %65  %37 %68
+%70 = OpLoad  %4  %69
+%72 = OpAccessChain  %71  %37 %34 %34
+%73 = OpLoad  %5  %72
+%74 = OpLoad  %6  %39
+%75 = OpAccessChain  %71  %37 %34 %74
+%76 = OpLoad  %5  %75
+%77 = OpLoad  %6  %39
+%78 = OpAccessChain  %71  %37 %77 %34
+%79 = OpLoad  %5  %78
+%80 = OpLoad  %6  %39
+%81 = OpLoad  %6  %39
+%82 = OpAccessChain  %71  %37 %80 %81
+%83 = OpLoad  %5  %82
+OpStore %35 %64
+%84 = OpAccessChain  %45  %35 %34
+OpStore %84 %67
+%85 = OpLoad  %6  %39
+%86 = OpAccessChain  %45  %35 %85
+OpStore %86 %70
+%87 = OpAccessChain  %51  %35 %34 %34
+OpStore %87 %73
+%88 = OpLoad  %6  %39
+%89 = OpAccessChain  %51  %35 %34 %88
+OpStore %89 %76
+%90 = OpLoad  %6  %39
+%91 = OpAccessChain  %51  %35 %90 %34
+OpStore %91 %79
+%92 = OpLoad  %6  %39
+%93 = OpLoad  %6  %39
+%94 = OpAccessChain  %51  %35 %92 %93
+OpStore %94 %83
+OpReturn
+OpFunctionEnd
+%96 = OpFunction  %2  None %32
+%95 = OpLabel
+%101 = OpVariable  %40  Function %38
+%98 = OpAccessChain  %97  %18 %34
+%100 = OpAccessChain  %99  %21 %34
+OpBranch %102
+%102 = OpLabel
+%103 = OpLoad  %6  %101
+%104 = OpISub  %6  %103 %38
+OpStore %101 %104
+%105 = OpLoad  %7  %98
+%106 = OpAccessChain  %33  %98 %34
+%107 = OpLoad  %3  %106
+%108 = OpAccessChain  %45  %98 %34 %34
+%109 = OpLoad  %4  %108
+%110 = OpLoad  %6  %101
+%111 = OpAccessChain  %45  %98 %34 %110
+%112 = OpLoad  %4  %111
+%113 = OpAccessChain  %51  %98 %34 %34 %34
+%114 = OpLoad  %5  %113
+%115 = OpLoad  %6  %101
+%116 = OpAccessChain  %51  %98 %34 %34 %115
+%117 = OpLoad  %5  %116
+%118 = OpLoad  %6  %101
+%119 = OpAccessChain  %51  %98 %34 %118 %34
+%120 = OpLoad  %5  %119
+%121 = OpLoad  %6  %101
+%122 = OpLoad  %6  %101
+%123 = OpAccessChain  %51  %98 %34 %121 %122
+%124 = OpLoad  %5  %123
+%125 = OpLoad  %7  %100
+%126 = OpAccessChain  %36  %100 %34
+%127 = OpLoad  %3  %126
+%128 = OpAccessChain  %65  %100 %34 %34
+%129 = OpLoad  %4  %128
+%130 = OpLoad  %6  %101
+%131 = OpAccessChain  %65  %100 %34 %130
+%132 = OpLoad  %4  %131
+%133 = OpAccessChain  %71  %100 %34 %34 %34
+%134 = OpLoad  %5  %133
+%135 = OpLoad  %6  %101
+%136 = OpAccessChain  %71  %100 %34 %34 %135
+%137 = OpLoad  %5  %136
+%138 = OpLoad  %6  %101
+%139 = OpAccessChain  %71  %100 %34 %138 %34
+%140 = OpLoad  %5  %139
+%141 = OpLoad  %6  %101
+%142 = OpLoad  %6  %101
+%143 = OpAccessChain  %71  %100 %34 %141 %142
+%144 = OpLoad  %5  %143
+OpStore %98 %125
+%145 = OpAccessChain  %33  %98 %34
+OpStore %145 %127
+%146 = OpAccessChain  %45  %98 %34 %34
+OpStore %146 %129
+%147 = OpLoad  %6  %101
+%148 = OpAccessChain  %45  %98 %34 %147
+OpStore %148 %132
+%149 = OpAccessChain  %51  %98 %34 %34 %34
+OpStore %149 %134
+%150 = OpLoad  %6  %101
+%151 = OpAccessChain  %51  %98 %34 %34 %150
+OpStore %151 %137
+%152 = OpLoad  %6  %101
+%153 = OpAccessChain  %51  %98 %34 %152 %34
+OpStore %153 %140
+%154 = OpLoad  %6  %101
+%155 = OpLoad  %6  %101
+%156 = OpAccessChain  %51  %98 %34 %154 %155
+OpStore %156 %144
+OpReturn
+OpFunctionEnd
+%158 = OpFunction  %2  None %32
+%157 = OpLabel
+%163 = OpVariable  %40  Function %38
+%160 = OpAccessChain  %159  %24 %34
+%162 = OpAccessChain  %161  %27 %34
+OpBranch %164
+%164 = OpLabel
+%165 = OpLoad  %6  %163
+%166 = OpISub  %6  %165 %38
+OpStore %163 %166
+%167 = OpLoad  %11  %160
+%169 = OpAccessChain  %168  %160 %34
+%170 = OpLoad  %8  %169
+%171 = OpAccessChain  %33  %160 %34 %34 %34
+%172 = OpLoad  %3  %171
+%173 = OpLoad  %6  %163
+%174 = OpAccessChain  %33  %160 %34 %173 %34
+%175 = OpLoad  %3  %174
+%176 = OpAccessChain  %45  %160 %34 %34 %34 %34
+%177 = OpLoad  %4  %176
+%178 = OpLoad  %6  %163
+%179 = OpAccessChain  %45  %160 %34 %34 %34 %178
+%180 = OpLoad  %4  %179
+%181 = OpLoad  %6  %163
+%182 = OpAccessChain  %45  %160 %34 %181 %34 %34
+%183 = OpLoad  %4  %182
+%184 = OpLoad  %6  %163
+%185 = OpLoad  %6  %163
+%186 = OpAccessChain  %45  %160 %34 %184 %34 %185
+%187 = OpLoad  %4  %186
+%188 = OpAccessChain  %51  %160 %34 %34 %34 %34 %34
+%189 = OpLoad  %5  %188
+%190 = OpLoad  %6  %163
+%191 = OpAccessChain  %51  %160 %34 %34 %34 %34 %190
+%192 = OpLoad  %5  %191
+%193 = OpLoad  %6  %163
+%194 = OpAccessChain  %51  %160 %34 %34 %34 %193 %34
+%195 = OpLoad  %5  %194
+%196 = OpLoad  %6  %163
+%197 = OpLoad  %6  %163
+%198 = OpAccessChain  %51  %160 %34 %34 %34 %196 %197
+%199 = OpLoad  %5  %198
+%200 = OpLoad  %6  %163
+%201 = OpAccessChain  %51  %160 %34 %200 %34 %34 %34
+%202 = OpLoad  %5  %201
+%203 = OpLoad  %6  %163
+%204 = OpLoad  %6  %163
+%205 = OpAccessChain  %51  %160 %34 %203 %34 %34 %204
+%206 = OpLoad  %5  %205
+%207 = OpLoad  %6  %163
+%208 = OpLoad  %6  %163
+%209 = OpAccessChain  %51  %160 %34 %207 %34 %208 %34
+%210 = OpLoad  %5  %209
+%211 = OpLoad  %6  %163
+%212 = OpLoad  %6  %163
+%213 = OpLoad  %6  %163
+%214 = OpAccessChain  %51  %160 %34 %211 %34 %212 %213
+%215 = OpLoad  %5  %214
+%216 = OpLoad  %11  %162
+%218 = OpAccessChain  %217  %162 %34
+%219 = OpLoad  %8  %218
+%220 = OpAccessChain  %36  %162 %34 %34 %34
+%221 = OpLoad  %3  %220
+%222 = OpLoad  %6  %163
+%223 = OpAccessChain  %36  %162 %34 %222 %34
+%224 = OpLoad  %3  %223
+%225 = OpAccessChain  %65  %162 %34 %34 %34 %34
+%226 = OpLoad  %4  %225
+%227 = OpLoad  %6  %163
+%228 = OpAccessChain  %65  %162 %34 %34 %34 %227
+%229 = OpLoad  %4  %228
+%230 = OpLoad  %6  %163
+%231 = OpAccessChain  %65  %162 %34 %230 %34 %34
+%232 = OpLoad  %4  %231
+%233 = OpLoad  %6  %163
+%234 = OpLoad  %6  %163
+%235 = OpAccessChain  %65  %162 %34 %233 %34 %234
+%236 = OpLoad  %4  %235
+%237 = OpAccessChain  %71  %162 %34 %34 %34 %34 %34
+%238 = OpLoad  %5  %237
+%239 = OpLoad  %6  %163
+%240 = OpAccessChain  %71  %162 %34 %34 %34 %34 %239
+%241 = OpLoad  %5  %240
+%242 = OpLoad  %6  %163
+%243 = OpAccessChain  %71  %162 %34 %34 %34 %242 %34
+%244 = OpLoad  %5  %243
+%245 = OpLoad  %6  %163
+%246 = OpLoad  %6  %163
+%247 = OpAccessChain  %71  %162 %34 %34 %34 %245 %246
+%248 = OpLoad  %5  %247
+%249 = OpLoad  %6  %163
+%250 = OpAccessChain  %71  %162 %34 %249 %34 %34 %34
+%251 = OpLoad  %5  %250
+%252 = OpLoad  %6  %163
+%253 = OpLoad  %6  %163
+%254 = OpAccessChain  %71  %162 %34 %252 %34 %34 %253
+%255 = OpLoad  %5  %254
+%256 = OpLoad  %6  %163
+%257 = OpLoad  %6  %163
+%258 = OpAccessChain  %71  %162 %34 %256 %34 %257 %34
+%259 = OpLoad  %5  %258
+%260 = OpLoad  %6  %163
+%261 = OpLoad  %6  %163
+%262 = OpLoad  %6  %163
+%263 = OpAccessChain  %71  %162 %34 %260 %34 %261 %262
+%264 = OpLoad  %5  %263
+OpStore %160 %216
+%265 = OpAccessChain  %168  %160 %34
+OpStore %265 %219
+%266 = OpAccessChain  %33  %160 %34 %34 %34
+OpStore %266 %221
+%267 = OpLoad  %6  %163
+%268 = OpAccessChain  %33  %160 %34 %267 %34
+OpStore %268 %224
+%269 = OpAccessChain  %45  %160 %34 %34 %34 %34
+OpStore %269 %226
+%270 = OpLoad  %6  %163
+%271 = OpAccessChain  %45  %160 %34 %34 %34 %270
+OpStore %271 %229
+%272 = OpLoad  %6  %163
+%273 = OpAccessChain  %45  %160 %34 %272 %34 %34
+OpStore %273 %232
+%274 = OpLoad  %6  %163
+%275 = OpLoad  %6  %163
+%276 = OpAccessChain  %45  %160 %34 %274 %34 %275
+OpStore %276 %236
+%277 = OpAccessChain  %51  %160 %34 %34 %34 %34 %34
+OpStore %277 %238
+%278 = OpLoad  %6  %163
+%279 = OpAccessChain  %51  %160 %34 %34 %34 %34 %278
+OpStore %279 %241
+%280 = OpLoad  %6  %163
+%281 = OpAccessChain  %51  %160 %34 %34 %34 %280 %34
+OpStore %281 %244
+%282 = OpLoad  %6  %163
+%283 = OpLoad  %6  %163
+%284 = OpAccessChain  %51  %160 %34 %34 %34 %282 %283
+OpStore %284 %248
+%285 = OpLoad  %6  %163
+%286 = OpAccessChain  %51  %160 %34 %285 %34 %34 %34
+OpStore %286 %251
+%287 = OpLoad  %6  %163
+%288 = OpLoad  %6  %163
+%289 = OpAccessChain  %51  %160 %34 %287 %34 %34 %288
+OpStore %289 %255
+%290 = OpLoad  %6  %163
+%291 = OpLoad  %6  %163
+%292 = OpAccessChain  %51  %160 %34 %290 %34 %291 %34
+OpStore %292 %259
+%293 = OpLoad  %6  %163
+%294 = OpLoad  %6  %163
+%295 = OpLoad  %6  %163
+%296 = OpAccessChain  %51  %160 %34 %293 %34 %294 %295
+OpStore %296 %264
+OpReturn
+OpFunctionEnd
+%298 = OpFunction  %2  None %32
+%297 = OpLabel
+%299 = OpAccessChain  %33  %12 %34
+%300 = OpAccessChain  %36  %15 %34
+%301 = OpAccessChain  %97  %18 %34
+%302 = OpAccessChain  %99  %21 %34
+%303 = OpAccessChain  %159  %24 %34
+%304 = OpAccessChain  %161  %27 %34
+OpBranch %305
+%305 = OpLabel
+%306 = OpFunctionCall  %2  %31
+%307 = OpFunctionCall  %2  %96
+%308 = OpFunctionCall  %2  %158
+OpReturn
+OpFunctionEnd
\ No newline at end of file
diff --git a/tests/tests/wgpu-gpu/shader/struct_layout.rs b/tests/tests/wgpu-gpu/shader/struct_layout.rs
index 8cb4501af5d..0b769a58988 100644
--- a/tests/tests/wgpu-gpu/shader/struct_layout.rs
+++ b/tests/tests/wgpu-gpu/shader/struct_layout.rs
@@ -3,7 +3,7 @@ use std::fmt::Write;
 use wgpu::{Backends, DownlevelFlags, Features, Limits};
 
 use crate::shader::{shader_input_output_test, InputStorageType, ShaderTest, MAX_BUFFER_SIZE};
-use wgpu_test::{gpu_test, FailureCase, GpuTestConfiguration, GpuTestInitializer, TestParameters};
+use wgpu_test::{gpu_test, GpuTestConfiguration, GpuTestInitializer, TestParameters};
 
 pub fn all_tests(vec: &mut Vec<GpuTestInitializer>) {
     vec.extend([
@@ -18,16 +18,13 @@ pub fn all_tests(vec: &mut Vec<GpuTestInitializer>) {
     ]);
 }
 
+// Note that some specific subtests are marked as failing on GL due to
+// https://github.com/gfx-rs/wgpu/issues/4371.
 #[gpu_test]
 static UNIFORM_INPUT: GpuTestConfiguration = GpuTestConfiguration::new()
     .parameters(
         TestParameters::default()
             .downlevel_flags(DownlevelFlags::COMPUTE_SHADERS)
-            // Validation errors thrown by the SPIR-V validator https://github.com/gfx-rs/wgpu/issues/4371
-            .expect_fail(
-                FailureCase::backend(wgpu::Backends::VULKAN)
-                    .validation_error("a matrix with stride 8 not satisfying alignment to 16"),
-            )
             .limits(Limits::downlevel_defaults()),
     )
     .run_async(|ctx| {
@@ -126,35 +123,68 @@ fn create_struct_layout_tests(storage_type: InputStorageType) -> Vec<ShaderTest>
             // - Do `input.member[0].x` (direct)
             // - Store `input.member[0]` in a variable; do `var.x` (vector_loaded)
             // - Store `input.member` in a variable; do `var[0].x` (fully_loaded)
-            let mut direct = String::new();
-            let mut vector_loaded = String::new();
-            let mut fully_loaded = String::from("let loaded = input.member;");
-            for column in 0..columns {
-                writeln!(vector_loaded, "let vec_{column} = input.member[{column}];").unwrap();
+            // For each of these, we can either use a static or dynamic index.
+            let mut direct_static = String::new();
+            let mut direct_dynamic = String::new();
+            let mut vector_loaded_static = String::new();
+            let mut vector_loaded_dynamic = String::new();
+            let mut fully_loaded_static = String::from("let loaded = input.member;");
+            let mut fully_loaded_dynamic = String::from("let loaded = input.member;");
+            let column_index_names = ["zero", "one", "two", "three"];
+            for (column, column_str) in column_index_names.iter().enumerate().take(columns) {
+                writeln!(direct_dynamic, "var {column_str} = {column};").unwrap();
+                writeln!(vector_loaded_dynamic, "var {column_str} = {column};").unwrap();
+                writeln!(fully_loaded_dynamic, "var {column_str} = {column};").unwrap();
+
+                writeln!(
+                    vector_loaded_static,
+                    "let vec_{column} = input.member[{column}];"
+                )
+                .unwrap();
+                writeln!(
+                    vector_loaded_dynamic,
+                    "let vec_{column} = input.member[{column_str}];",
+                )
+                .unwrap();
             }
 
             let mut output_values = Vec::new();
 
             let mut current_output_idx = 0;
             let mut current_input_idx = 0;
-            for column in 0..columns {
+            for (column, column_str) in column_index_names.iter().enumerate().take(columns) {
                 let component_accessors = ["x", "y", "z", "w"].into_iter().take(rows);
                 for component in component_accessors {
                     writeln!(
-                        direct,
+                        direct_static,
                         "output[{current_output_idx}] = bitcast<u32>(input.member[{column}].{component});"
                     )
                     .unwrap();
                     writeln!(
-                        vector_loaded,
+                        direct_dynamic,
+                        "output[{current_output_idx}] = bitcast<u32>(input.member[{column_str}].{component});"
+                    )
+                    .unwrap();
+                    writeln!(
+                        vector_loaded_static,
                         "output[{current_output_idx}] = bitcast<u32>(vec_{column}.{component});"
                     )
                     .unwrap();
                     writeln!(
-                        fully_loaded,
+                        vector_loaded_dynamic,
+                        "output[{current_output_idx}] = bitcast<u32>(vec_{column}.{component});"
+                    )
+                    .unwrap();
+                    writeln!(
+                        fully_loaded_static,
                         "output[{current_output_idx}] = bitcast<u32>(loaded[{column}].{component});"
                     )
                     .unwrap();
+                    writeln!(
+                        fully_loaded_dynamic,
+                        "output[{current_output_idx}] = bitcast<u32>(loaded[{column_str}].{component});"
+                    )
+                    .unwrap();
 
                     output_values.push(current_input_idx);
                     current_input_idx += 1;
@@ -175,9 +205,202 @@ fn create_struct_layout_tests(storage_type: InputStorageType) -> Vec<ShaderTest>
 
             tests.push(
                 ShaderTest::new(
-                    format!("{ty} - direct"),
+                    format!("{ty} - direct, static index"),
                     input_members.clone(),
-                    direct,
+                    direct_static,
+                    &input_values,
+                    &output_values,
+                )
+                .failures(failures),
+            );
+            tests.push(
+                ShaderTest::new(
+                    format!("{ty} - direct, dynamic index"),
+                    input_members.clone(),
+                    direct_dynamic,
+                    &input_values,
+                    &output_values,
+                )
+                .failures(failures),
+            );
+
+            tests.push(
+                ShaderTest::new(
+                    format!("{ty} - vector loaded, static index"),
+                    input_members.clone(),
+                    vector_loaded_static,
+                    &input_values,
+                    &output_values,
+                )
+                .failures(failures),
+            );
+            tests.push(
+                ShaderTest::new(
+                    format!("{ty} - vector loaded, dynamic index"),
+                    input_members.clone(),
+                    vector_loaded_dynamic,
+                    &input_values,
+                    &output_values,
+                )
+                .failures(failures),
+            );
+
+            tests.push(
+                ShaderTest::new(
+                    format!("{ty} - fully loaded, static index"),
+                    input_members.clone(),
+                    fully_loaded_static,
+                    &input_values,
+                    &output_values,
+                )
+                .failures(failures),
+            );
+            tests.push(
+                ShaderTest::new(
+                    format!("{ty} - fully loaded, dynamic index"),
+                    input_members.clone(),
+                    fully_loaded_dynamic,
+                    &input_values,
+                    &output_values,
+                )
+                .failures(failures),
+            );
+        }
+    }
+
+    // Array of matrix tests
+    for columns in [2, 4] {
+        for rows in [2, 3, 4] {
+            let array_size = 2;
+            let ty = format!("mat{columns}x{rows}<f32>");
+            let input_members = format!("members: array<{ty}, {array_size}>");
+            // There's 4 possible ways to load a component of a matrix in an array:
+            // - Do `input.members[0][0].x` (direct)
+            // - Store `input.members[0][0]` in a variable; do `var.x` (vector_loaded)
+            // - Store `input.members[0]` in a variable; do `var[0].x` (matrix_loaded)
+            // - Store `input.members` in a variable; do `var[0][0].x` (fully_loaded)
+            // For each of these, we can either use a static or dynamic index.
+            let mut direct_static = String::new();
+            let mut direct_dynamic = String::new();
+            let mut vector_loaded_static = String::new();
+            let mut vector_loaded_dynamic = String::new();
+            let mut matrix_loaded_static = String::new();
+            let mut matrix_loaded_dynamic = String::new();
+            let mut fully_loaded_static = String::from("let loaded = input.members;");
+            let mut fully_loaded_dynamic = String::from("let loaded = input.members;");
+            let column_index_names = ["zero", "one", "two", "three"];
+            for (column, column_str) in column_index_names.iter().enumerate().take(columns) {
+                writeln!(direct_dynamic, "var {column_str} = {column};").unwrap();
+                writeln!(vector_loaded_dynamic, "var {column_str} = {column};").unwrap();
+                writeln!(matrix_loaded_dynamic, "var {column_str} = {column};").unwrap();
+            }
+            for element in 0..array_size {
+                writeln!(
+                    matrix_loaded_static,
+                    "let mat_{element} = input.members[{element}];"
+                )
+                .unwrap();
+                writeln!(
+                    matrix_loaded_dynamic,
+                    "let mat_{element} = input.members[{element}];"
+                )
+                .unwrap();
+                for (column, column_str) in column_index_names.iter().enumerate().take(columns) {
+                    writeln!(
+                        vector_loaded_static,
+                        "let mat_{element}_vec_{column} = input.members[{element}][{column}];"
+                    )
+                    .unwrap();
+                    writeln!(
+                        vector_loaded_dynamic,
+                        "let mat_{element}_vec_{column} = input.members[{element}][{column_str}];",
+                    )
+                    .unwrap();
+                }
+            }
+
+            let mut output_values = Vec::new();
+
+            let mut current_output_idx = 0;
+            let mut current_input_idx = 0;
+            for element in 0..array_size {
+                for (column, column_str) in column_index_names.iter().enumerate().take(columns) {
+                    let component_accessors = ["x", "y", "z", "w"].into_iter().take(rows);
+                    for component in component_accessors {
+                        writeln!(
+                            direct_static,
+                            "output[{current_output_idx}] = bitcast<u32>(input.members[{element}][{column}].{component});"
+                        )
+                        .unwrap();
+                        writeln!(
+                            direct_dynamic,
+                            "output[{current_output_idx}] = bitcast<u32>(input.members[{element}][{column_str}].{component});"
+                        )
+                        .unwrap();
+                        writeln!(
+                            vector_loaded_static,
+                            "output[{current_output_idx}] = bitcast<u32>(mat_{element}_vec_{column}.{component});"
+                        )
+                        .unwrap();
+                        writeln!(
+                            vector_loaded_dynamic,
+                            "output[{current_output_idx}] = bitcast<u32>(mat_{element}_vec_{column}.{component});"
+                        )
+                        .unwrap();
+                        writeln!(
+                            matrix_loaded_static,
+                            "output[{current_output_idx}] = bitcast<u32>(mat_{element}[{column}].{component});"
+                        )
+                        .unwrap();
+                        writeln!(
+                            matrix_loaded_dynamic,
+                            "output[{current_output_idx}] = bitcast<u32>(mat_{element}[{column_str}].{component});"
+                        )
+                        .unwrap();
+                        writeln!(
+                            fully_loaded_static,
+                            "output[{current_output_idx}] = bitcast<u32>(loaded[{column}].{component});"
+                        )
+                        .unwrap();
+                        writeln!(
+                            fully_loaded_dynamic,
+                            "output[{current_output_idx}] = bitcast<u32>(loaded[{column_str}].{component});"
+                        )
+                        .unwrap();
+
+                        output_values.push(current_input_idx);
+                        current_input_idx += 1;
+                        current_output_idx += 1;
+                    }
+                    // Round to next vec4 if we're matrices with vec3 columns
+                    if rows == 3 {
+                        current_input_idx += 1;
+                    }
+                }
+            }
+
+            // https://github.com/gfx-rs/wgpu/issues/4371
+            let failures = if storage_type == InputStorageType::Uniform && rows == 2 {
+                Backends::GL
+            } else {
+                Backends::empty()
+            };
+
+            tests.push(
+                ShaderTest::new(
+                    format!("{ty} - direct, static index"),
+                    input_members.clone(),
+                    direct_static,
+                    &input_values,
+                    &output_values,
+                )
+                .failures(failures),
+            );
+            tests.push(
+                ShaderTest::new(
+                    format!("{ty} - direct, dynamic index"),
+                    input_members.clone(),
+                    direct_dynamic,
                     &input_values,
                     &output_values,
                 )
@@ -186,9 +409,19 @@ fn create_struct_layout_tests(storage_type: InputStorageType) -> Vec<ShaderTest>
 
             tests.push(
                 ShaderTest::new(
-                    format!("{ty} - vector loaded"),
+                    format!("{ty} - vector loaded, static index"),
+                    input_members.clone(),
+                    vector_loaded_static,
+                    &input_values,
+                    &output_values,
+                )
+                .failures(failures),
+            );
+            tests.push(
+                ShaderTest::new(
+                    format!("{ty} - vector loaded, dynamic index"),
                     input_members.clone(),
-                    vector_loaded,
+                    vector_loaded_dynamic,
                     &input_values,
                     &output_values,
                 )
@@ -197,9 +430,19 @@ fn create_struct_layout_tests(storage_type: InputStorageType) -> Vec<ShaderTest>
 
             tests.push(
                 ShaderTest::new(
-                    format!("{ty} - fully loaded"),
+                    format!("{ty} - matrix loaded, static index"),
                     input_members.clone(),
-                    fully_loaded,
+                    matrix_loaded_static,
+                    &input_values,
+                    &output_values,
+                )
+                .failures(failures),
+            );
+            tests.push(
+                ShaderTest::new(
+                    format!("{ty} - matrix loaded, dynamic index"),
+                    input_members.clone(),
+                    matrix_loaded_dynamic,
                     &input_values,
                     &output_values,
                 )
@@ -208,6 +451,41 @@ fn create_struct_layout_tests(storage_type: InputStorageType) -> Vec<ShaderTest>
         }
     }
 
+    // MatCx2 followed by other members in same struct. Since on some backends
+    // the matrix will be decomposed into separate column members in the struct,
+    // this tests that the other members can still be accessed correctly. This
+    // is especially important on SPIR-V where members are accessed by index
+    // rather than name.
+    {
+        let members = String::from("m: mat3x2<f32>,\nf: f32,");
+        let direct = String::from(
+            "\
+            output[0] = bitcast<u32>(input.m[0].x);
+            output[1] = bitcast<u32>(input.m[0].y);
+            output[2] = bitcast<u32>(input.m[1].x);
+            output[3] = bitcast<u32>(input.m[1].y);
+            output[4] = bitcast<u32>(input.m[2].x);
+            output[5] = bitcast<u32>(input.m[2].y);
+            output[6] = bitcast<u32>(input.f);
+        ",
+        );
+        tests.push(
+            ShaderTest::new(
+                String::from("MatCx2 followed by other members"),
+                members,
+                direct,
+                &input_values,
+                &[0, 1, 2, 3, 4, 5, 6],
+            )
+            // https://github.com/gfx-rs/wgpu/issues/4371
+            .failures(if storage_type == InputStorageType::Uniform {
+                Backends::GL
+            } else {
+                Backends::empty()
+            }),
+        );
+    }
+
     // Vec3 alignment tests
     for ty in ["f32", "u32", "i32"] {
         let members = format!("_vec: vec3<{ty}>,\nscalar: {ty},");