diff --git a/CHANGELOG.md b/CHANGELOG.md
index 0162ce356a..adc4d9b228 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -130,6 +130,7 @@ By @bradwerth [#6216](https://github.com/gfx-rs/wgpu/pull/6216).
 #### GLES / OpenGL
 
 - Fix GL debug message callbacks not being properly cleaned up (causing UB). By @Imberflur in [#6114](https://github.com/gfx-rs/wgpu/pull/6114)
+- Fix calling `slice::from_raw_parts` with unaligned pointers in push constant handling. By @Imberflur in [#6341](https://github.com/gfx-rs/wgpu/pull/6341)
 
 #### WebGPU
 
diff --git a/Cargo.lock b/Cargo.lock
index 18d0772071..482b56fdac 100644
--- a/Cargo.lock
+++ b/Cargo.lock
@@ -2121,12 +2121,9 @@ dependencies = [
 
 [[package]]
 name = "once_cell"
-version = "1.20.1"
+version = "1.19.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "82881c4be219ab5faaf2ad5e5e5ecdff8c66bd7402ca3160975c93b24961afd1"
-dependencies = [
- "portable-atomic",
-]
+checksum = "3fdb12b2476b595f9358c5161aa467c2438859caa136dec86c26fdd2efe17b92"
 
 [[package]]
 name = "oorandom"
@@ -2344,12 +2341,6 @@ version = "0.3.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "22686f4785f02a4fcc856d3b3bb19bf6c8160d103f7a99cc258bddd0251dc7f2"
 
-[[package]]
-name = "portable-atomic"
-version = "1.9.0"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "cc9c68a3f6da06753e9335d63e27f6b9754dd1920d941135b7ea8224f141adb2"
-
 [[package]]
 name = "pp-rs"
 version = "0.2.1"
@@ -3684,6 +3675,7 @@ dependencies = [
  "bit-set",
  "bitflags 2.6.0",
  "block",
+ "bytemuck",
  "cfg-if",
  "cfg_aliases",
  "core-graphics-types",
diff --git a/Cargo.toml b/Cargo.toml
index fbf06524c1..45c6fa6505 100644
--- a/Cargo.toml
+++ b/Cargo.toml
@@ -76,7 +76,7 @@ arrayvec = "0.7"
 bincode = "1"
 bit-vec = "0.8"
 bitflags = "2.6"
-bytemuck = { version = "1.18", features = ["derive"] }
+bytemuck = { version = "1.18" }
 cfg_aliases = "0.1"
 cfg-if = "1"
 criterion = "0.5"
@@ -104,7 +104,7 @@ nanorand = { version = "0.7", default-features = false, features = ["wyrand"] }
 noise = { version = "0.8", git = "https://github.com/Razaekel/noise-rs.git", rev = "c6942d4fb70af26db4441edcf41f90fa115333f2" }
 nv-flip = "0.1"
 obj = "0.10"
-once_cell = "1.20.1"
+once_cell = "1.19.0"
 parking_lot = "0.12.1"
 pico-args = { version = "0.5.0", features = [
     "eq-separator",
diff --git a/examples/Cargo.toml b/examples/Cargo.toml
index 73a575dce5..2d4fd5331b 100644
--- a/examples/Cargo.toml
+++ b/examples/Cargo.toml
@@ -29,7 +29,7 @@ webgl = ["wgpu/webgl"]
 webgpu = ["wgpu/webgpu"]
 
 [dependencies]
-bytemuck.workspace = true
+bytemuck = { workspace = true, features = ["derive"] }
 cfg-if.workspace = true
 encase = { workspace = true, features = ["glam"] }
 flume.workspace = true
diff --git a/naga/src/back/pipeline_constants.rs b/naga/src/back/pipeline_constants.rs
index 5f82862f72..0005cbcb0e 100644
--- a/naga/src/back/pipeline_constants.rs
+++ b/naga/src/back/pipeline_constants.rs
@@ -14,7 +14,10 @@ use thiserror::Error;
 pub enum PipelineConstantError {
     #[error("Missing value for pipeline-overridable constant with identifier string: '{0}'")]
     MissingValue(String),
-    #[error("Source f64 value needs to be finite (NaNs and Inifinites are not allowed) for number destinations")]
+    #[error(
+        "Source f64 value needs to be finite ({}) for number destinations",
+        "NaNs and Inifinites are not allowed"
+    )]
     SrcNeedsToBeFinite,
     #[error("Source f64 value doesn't fit in destination")]
     DstRangeTooSmall,
diff --git a/naga/src/back/spv/block.rs b/naga/src/back/spv/block.rs
index f0c3bfa848..c6adba85b1 100644
--- a/naga/src/back/spv/block.rs
+++ b/naga/src/back/spv/block.rs
@@ -4,8 +4,7 @@ Implementations for `BlockContext` methods.
 
 use super::{
     helpers, index::BoundsCheckResult, make_local, selection::Selection, Block, BlockContext,
-    Dimension, Error, Instruction, LocalType, LookupType, LoopContext, ResultMember, Writer,
-    WriterFlags,
+    Dimension, Error, Instruction, LocalType, LookupType, ResultMember, Writer, WriterFlags,
 };
 use crate::{arena::Handle, proc::TypeResolution, Statement};
 use spirv::Word;
@@ -39,7 +38,7 @@ enum ExpressionPointer {
 }
 
 /// The termination statement to be added to the end of the block
-pub enum BlockExit {
+enum BlockExit {
     /// Generates an OpReturn (void return)
     Return,
     /// Generates an OpBranch to the specified block
@@ -60,6 +59,36 @@ pub enum BlockExit {
     },
 }
 
+/// What code generation did with a provided [`BlockExit`] value.
+///
+/// A function that accepts a [`BlockExit`] argument should return a value of
+/// this type, to indicate whether the code it generated ended up using the
+/// provided exit, or ignored it and did a non-local exit of some other kind
+/// (say, [`Break`] or [`Continue`]). Some callers must use this information to
+/// decide whether to generate the target block at all.
+///
+/// [`Break`]: Statement::Break
+/// [`Continue`]: Statement::Continue
+#[must_use]
+enum BlockExitDisposition {
+    /// The generated code used the provided `BlockExit` value. If it included a
+    /// block label, the caller should be sure to actually emit the block it
+    /// refers to.
+    Used,
+
+    /// The generated code did not use the provided `BlockExit` value. If it
+    /// included a block label, the caller should not bother to actually emit
+    /// the block it refers to, unless it knows the block is needed for
+    /// something else.
+    Discarded,
+}
+
+#[derive(Clone, Copy, Default)]
+struct LoopContext {
+    continuing_id: Option<Word>,
+    break_id: Option<Word>,
+}
+
 #[derive(Debug)]
 pub(crate) struct DebugInfoInner<'a> {
     pub source_code: &'a str,
@@ -343,6 +372,32 @@ impl<'w> BlockContext<'w> {
 
                         load_id
                     }
+                    crate::TypeInner::Array {
+                        base: ty_element, ..
+                    } => {
+                        let index_id = self.cached[index];
+                        let base_id = self.cached[base];
+                        let base_ty = match self.fun_info[base].ty {
+                            TypeResolution::Handle(handle) => handle,
+                            TypeResolution::Value(_) => {
+                                return Err(Error::Validation(
+                                    "Array types should always be in the arena",
+                                ))
+                            }
+                        };
+                        let (id, variable) = self.writer.promote_access_expression_to_variable(
+                            &self.ir_module.types,
+                            result_type_id,
+                            base_id,
+                            base_ty,
+                            index_id,
+                            ty_element,
+                            block,
+                        )?;
+                        self.function.internal_variables.push(variable);
+                        id
+                    }
+                    // wgpu#4337: Support `crate::TypeInner::Matrix`
                     ref other => {
                         log::error!(
                             "Unable to access base {:?} of type {:?}",
@@ -350,7 +405,7 @@ impl<'w> BlockContext<'w> {
                             other
                         );
                         return Err(Error::Validation(
-                            "only vectors may be dynamically indexed by value",
+                            "only vectors and arrays may be dynamically indexed by value",
                         ));
                     }
                 }
@@ -2034,14 +2089,30 @@ impl<'w> BlockContext<'w> {
         }
     }
 
-    pub(super) fn write_block(
+    /// Generate one or more SPIR-V blocks for `naga_block`.
+    ///
+    /// Use `label_id` as the label for the SPIR-V entry point block.
+    ///
+    /// If control reaches the end of the SPIR-V block, terminate it according
+    /// to `exit`. This function's return value indicates whether it acted on
+    /// this parameter or not; see [`BlockExitDisposition`].
+    ///
+    /// If the block contains [`Break`] or [`Continue`] statements,
+    /// `loop_context` supplies the labels of the SPIR-V blocks to jump to. If
+    /// either of these labels are `None`, then it should have been a Naga
+    /// validation error for the corresponding statement to occur in this
+    /// context.
+    ///
+    /// [`Break`]: Statement::Break
+    /// [`Continue`]: Statement::Continue
+    fn write_block(
         &mut self,
         label_id: Word,
         naga_block: &crate::Block,
         exit: BlockExit,
         loop_context: LoopContext,
         debug_info: Option<&DebugInfoInner>,
-    ) -> Result<(), Error> {
+    ) -> Result<BlockExitDisposition, Error> {
         let mut block = Block::new(label_id);
         for (statement, span) in naga_block.span_iter() {
             if let (Some(debug_info), false) = (
@@ -2077,7 +2148,7 @@ impl<'w> BlockContext<'w> {
                     self.function.consume(block, Instruction::branch(scope_id));
 
                     let merge_id = self.gen_id();
-                    self.write_block(
+                    let merge_used = self.write_block(
                         scope_id,
                         block_statements,
                         BlockExit::Branch { target: merge_id },
@@ -2085,7 +2156,14 @@ impl<'w> BlockContext<'w> {
                         debug_info,
                     )?;
 
-                    block = Block::new(merge_id);
+                    match merge_used {
+                        BlockExitDisposition::Used => {
+                            block = Block::new(merge_id);
+                        }
+                        BlockExitDisposition::Discarded => {
+                            return Ok(BlockExitDisposition::Discarded);
+                        }
+                    }
                 }
                 Statement::If {
                     condition,
@@ -2121,7 +2199,11 @@ impl<'w> BlockContext<'w> {
                     );
 
                     if let Some(block_id) = accept_id {
-                        self.write_block(
+                        // We can ignore the `BlockExitDisposition` returned here because,
+                        // even if `merge_id` is not actually reachable, it is always
+                        // referred to by the `OpSelectionMerge` instruction we emitted
+                        // earlier.
+                        let _ = self.write_block(
                             block_id,
                             accept,
                             BlockExit::Branch { target: merge_id },
@@ -2130,7 +2212,11 @@ impl<'w> BlockContext<'w> {
                         )?;
                     }
                     if let Some(block_id) = reject_id {
-                        self.write_block(
+                        // We can ignore the `BlockExitDisposition` returned here because,
+                        // even if `merge_id` is not actually reachable, it is always
+                        // referred to by the `OpSelectionMerge` instruction we emitted
+                        // earlier.
+                        let _ = self.write_block(
                             block_id,
                             reject,
                             BlockExit::Branch { target: merge_id },
@@ -2208,7 +2294,15 @@ impl<'w> BlockContext<'w> {
                         } else {
                             merge_id
                         };
-                        self.write_block(
+                        // We can ignore the `BlockExitDisposition` returned here because
+                        // `case_finish_id` is always referred to by either:
+                        //
+                        // - the `OpSwitch`, if it's the next case's label for a
+                        //   fall-through, or
+                        //
+                        // - the `OpSelectionMerge`, if it's the switch's overall merge
+                        //   block because there's no fall-through.
+                        let _ = self.write_block(
                             *label_id,
                             &case.body,
                             BlockExit::Branch {
@@ -2254,7 +2348,10 @@ impl<'w> BlockContext<'w> {
                     ));
                     self.function.consume(block, Instruction::branch(body_id));
 
-                    self.write_block(
+                    // We can ignore the `BlockExitDisposition` returned here because,
+                    // even if `continuing_id` is not actually reachable, it is always
+                    // referred to by the `OpLoopMerge` instruction we emitted earlier.
+                    let _ = self.write_block(
                         body_id,
                         body,
                         BlockExit::Branch {
@@ -2277,7 +2374,10 @@ impl<'w> BlockContext<'w> {
                         },
                     };
 
-                    self.write_block(
+                    // We can ignore the `BlockExitDisposition` returned here because,
+                    // even if `merge_id` is not actually reachable, it is always referred
+                    // to by the `OpLoopMerge` instruction we emitted earlier.
+                    let _ = self.write_block(
                         continuing_id,
                         continuing,
                         exit,
@@ -2293,14 +2393,14 @@ impl<'w> BlockContext<'w> {
                 Statement::Break => {
                     self.function
                         .consume(block, Instruction::branch(loop_context.break_id.unwrap()));
-                    return Ok(());
+                    return Ok(BlockExitDisposition::Discarded);
                 }
                 Statement::Continue => {
                     self.function.consume(
                         block,
                         Instruction::branch(loop_context.continuing_id.unwrap()),
                     );
-                    return Ok(());
+                    return Ok(BlockExitDisposition::Discarded);
                 }
                 Statement::Return { value: Some(value) } => {
                     let value_id = self.cached[value];
@@ -2319,15 +2419,15 @@ impl<'w> BlockContext<'w> {
                         None => Instruction::return_value(value_id),
                     };
                     self.function.consume(block, instruction);
-                    return Ok(());
+                    return Ok(BlockExitDisposition::Discarded);
                 }
                 Statement::Return { value: None } => {
                     self.function.consume(block, Instruction::return_void());
-                    return Ok(());
+                    return Ok(BlockExitDisposition::Discarded);
                 }
                 Statement::Kill => {
                     self.function.consume(block, Instruction::kill());
-                    return Ok(());
+                    return Ok(BlockExitDisposition::Discarded);
                 }
                 Statement::Barrier(flags) => {
                     self.writer.write_barrier(flags, &mut block);
@@ -2693,6 +2793,24 @@ impl<'w> BlockContext<'w> {
         };
 
         self.function.consume(block, termination);
+        Ok(BlockExitDisposition::Used)
+    }
+
+    pub(super) fn write_function_body(
+        &mut self,
+        entry_id: Word,
+        debug_info: Option<&DebugInfoInner>,
+    ) -> Result<(), Error> {
+        // We can ignore the `BlockExitDisposition` returned here because
+        // `BlockExit::Return` doesn't refer to a block.
+        let _ = self.write_block(
+            entry_id,
+            &self.ir_function.body,
+            super::block::BlockExit::Return,
+            LoopContext::default(),
+            debug_info,
+        )?;
+
         Ok(())
     }
 }
diff --git a/naga/src/back/spv/helpers.rs b/naga/src/back/spv/helpers.rs
index 15c241d44e..63144abc02 100644
--- a/naga/src/back/spv/helpers.rs
+++ b/naga/src/back/spv/helpers.rs
@@ -85,22 +85,9 @@ impl crate::AddressSpace {
 
 /// Return true if the global requires a type decorated with `Block`.
 ///
-/// In the Vulkan spec 1.3.296, the section [Descriptor Set Interface][dsi] says:
+/// See [`back::spv::GlobalVariable`] for details.
 ///
-/// > Variables identified with the `Uniform` storage class are used to
-/// > access transparent buffer backed resources. Such variables must
-/// > be:
-/// >
-/// > -   typed as `OpTypeStruct`, or an array of this type,
-/// >
-/// > -   identified with a `Block` or `BufferBlock` decoration, and
-/// >
-/// > -   laid out explicitly using the `Offset`, `ArrayStride`, and
-/// >     `MatrixStride` decorations as specified in §15.6.4, "Offset
-/// >     and Stride Assignment."
-///
-/// [dsi]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#interfaces-resources-descset
-// See `back::spv::GlobalVariable::access_id` for details.
+/// [`back::spv::GlobalVariable`]: super::GlobalVariable
 pub fn global_needs_wrapper(ir_module: &crate::Module, var: &crate::GlobalVariable) -> bool {
     match var.space {
         crate::AddressSpace::Uniform
diff --git a/naga/src/back/spv/index.rs b/naga/src/back/spv/index.rs
index 0295d895b2..1b7cce9978 100644
--- a/naga/src/back/spv/index.rs
+++ b/naga/src/back/spv/index.rs
@@ -11,16 +11,31 @@ use crate::{arena::Handle, proc::BoundsCheckPolicy};
 
 /// The results of performing a bounds check.
 ///
-/// On success, `write_bounds_check` returns a value of this type.
+/// On success, [`write_bounds_check`](BlockContext::write_bounds_check)
+/// returns a value of this type. The caller can assume that the right
+/// policy has been applied, and simply do what the variant says.
 pub(super) enum BoundsCheckResult {
     /// The index is statically known and in bounds, with the given value.
     KnownInBounds(u32),
 
     /// The given instruction computes the index to be used.
+    ///
+    /// When [`BoundsCheckPolicy::Restrict`] is in force, this is a
+    /// clamped version of the index the user supplied.
+    ///
+    /// When [`BoundsCheckPolicy::Unchecked`] is in force, this is
+    /// simply the index the user supplied. This variant indicates
+    /// that we couldn't prove statically that the index was in
+    /// bounds; otherwise we would have returned [`KnownInBounds`].
+    ///
+    /// [`KnownInBounds`]: BoundsCheckResult::KnownInBounds
     Computed(Word),
 
     /// The given instruction computes a boolean condition which is true
     /// if the index is in bounds.
+    ///
+    /// This is returned when [`BoundsCheckPolicy::ReadZeroSkipWrite`]
+    /// is in force.
     Conditional(Word),
 }
 
@@ -52,13 +67,14 @@ impl<'w> BlockContext<'w> {
     /// - An optional [`Access`] or [`AccessIndex`], for case 3, applied to...
     /// - A [`GlobalVariable`].
     ///
-    /// The SPIR-V generated takes into account wrapped globals; see
-    /// [`global_needs_wrapper`].
+    /// The generated SPIR-V takes into account wrapped globals; see
+    /// [`back::spv::GlobalVariable`] for details.
     ///
     /// [`GlobalVariable`]: crate::Expression::GlobalVariable
     /// [`AccessIndex`]: crate::Expression::AccessIndex
     /// [`Access`]: crate::Expression::Access
     /// [`base`]: crate::Expression::Access::base
+    /// [`back::spv::GlobalVariable`]: super::GlobalVariable
     pub(super) fn write_runtime_array_length(
         &mut self,
         array: Handle<crate::Expression>,
@@ -357,6 +373,8 @@ impl<'w> BlockContext<'w> {
 
     /// Write an index bounds comparison to `block`, if needed.
     ///
+    /// This is used to implement [`BoundsCheckPolicy::ReadZeroSkipWrite`].
+    ///
     /// If we're able to determine statically that `index` is in bounds for
     /// `sequence`, return `KnownInBounds(value)`, where `value` is the actual
     /// value of the index. (In principle, one could know that the index is in
@@ -477,11 +495,23 @@ impl<'w> BlockContext<'w> {
 
     /// Emit code for bounds checks for an array, vector, or matrix access.
     ///
-    /// This implements either `index_bounds_check_policy` or
-    /// `buffer_bounds_check_policy`, depending on the address space of the
-    /// pointer being accessed.
+    /// This tries to handle all the critical steps for bounds checks:
+    ///
+    /// - First, select the appropriate bounds check policy for `base`,
+    ///   depending on its address space.
+    ///
+    /// - Next, analyze `index` to see if its value is known at
+    ///   compile time, in which case we can decide statically whether
+    ///   the index is in bounds.
+    ///
+    /// - If the index's value is not known at compile time, emit code to:
+    ///
+    ///     - restrict its value (for [`BoundsCheckPolicy::Restrict`]), or
+    ///
+    ///     - check whether it's in bounds (for
+    ///       [`BoundsCheckPolicy::ReadZeroSkipWrite`]).
     ///
-    /// Return a `BoundsCheckResult` indicating how the index should be
+    /// Return a [`BoundsCheckResult`] indicating how the index should be
     /// consumed. See that type's documentation for details.
     pub(super) fn write_bounds_check(
         &mut self,
diff --git a/naga/src/back/spv/mod.rs b/naga/src/back/spv/mod.rs
index 32bd1fcecf..e6397017c5 100644
--- a/naga/src/back/spv/mod.rs
+++ b/naga/src/back/spv/mod.rs
@@ -144,6 +144,7 @@ struct Function {
     signature: Option<Instruction>,
     parameters: Vec<FunctionArgument>,
     variables: crate::FastHashMap<Handle<crate::LocalVariable>, LocalVariable>,
+    internal_variables: Vec<LocalVariable>,
     blocks: Vec<TerminatedBlock>,
     entry_point_context: Option<EntryPointContext>,
 }
@@ -466,38 +467,75 @@ enum CachedConstant {
     ZeroValue(Word),
 }
 
+/// The SPIR-V representation of a [`crate::GlobalVariable`].
+///
+/// In the Vulkan spec 1.3.296, the section [Descriptor Set Interface][dsi] says:
+///
+/// > Variables identified with the `Uniform` storage class are used to access
+/// > transparent buffer backed resources. Such variables *must* be:
+/// >
+/// > -   typed as `OpTypeStruct`, or an array of this type,
+/// >
+/// > -   identified with a `Block` or `BufferBlock` decoration, and
+/// >
+/// > -   laid out explicitly using the `Offset`, `ArrayStride`, and `MatrixStride`
+/// >     decorations as specified in "Offset and Stride Assignment".
+///
+/// This is followed by identical language for the `StorageBuffer`,
+/// except that a `BufferBlock` decoration is not allowed.
+///
+/// When we encounter a global variable in the [`Storage`] or [`Uniform`]
+/// address spaces whose type is not already [`Struct`], this backend implicitly
+/// wraps the global variable in a struct: we generate a SPIR-V global variable
+/// holding an `OpTypeStruct` with a single member, whose type is what the Naga
+/// global's type would suggest, decorated as required above.
+///
+/// The [`helpers::global_needs_wrapper`] function determines whether a given
+/// [`crate::GlobalVariable`] needs to be wrapped.
+///
+/// [dsi]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#interfaces-resources-descset
+/// [`Storage`]: crate::AddressSpace::Storage
+/// [`Uniform`]: crate::AddressSpace::Uniform
+/// [`Struct`]: crate::TypeInner::Struct
 #[derive(Clone)]
 struct GlobalVariable {
-    /// ID of the OpVariable that declares the global.
+    /// The SPIR-V id of the `OpVariable` that declares the global.
     ///
-    /// If you need the variable's value, use [`access_id`] instead of this
-    /// field. If we wrapped the Naga IR `GlobalVariable`'s type in a struct to
-    /// comply with Vulkan's requirements, then this points to the `OpVariable`
-    /// with the synthesized struct type, whereas `access_id` points to the
-    /// field of said struct that holds the variable's actual value.
+    /// If this global has been implicitly wrapped in an `OpTypeStruct`, this id
+    /// refers to the wrapper, not the original Naga value it contains. If you
+    /// need the Naga value, use [`access_id`] instead of this field.
+    ///
+    /// If this global is not implicitly wrapped, this is the same as
+    /// [`access_id`].
     ///
     /// This is used to compute the `access_id` pointer in function prologues,
-    /// and used for `ArrayLength` expressions, which do need the struct.
+    /// and used for `ArrayLength` expressions, which need to pass the wrapper
+    /// struct.
     ///
     /// [`access_id`]: GlobalVariable::access_id
     var_id: Word,
 
-    /// For `AddressSpace::Handle` variables, this ID is recorded in the function
-    /// prelude block (and reset before every function) as `OpLoad` of the variable.
-    /// It is then used for all the global ops, such as `OpImageSample`.
+    /// The loaded value of a `AddressSpace::Handle` global variable.
+    ///
+    /// If the current function uses this global variable, this is the id of an
+    /// `OpLoad` instruction in the function's prologue that loads its value.
+    /// (This value is assigned as we write the prologue code of each function.)
+    /// It is then used for all operations on the global, such as `OpImageSample`.
     handle_id: Word,
 
-    /// Actual ID used to access this variable.
-    /// For wrapped buffer variables, this ID is `OpAccessChain` into the
-    /// wrapper. Otherwise, the same as `var_id`.
+    /// The SPIR-V id of a pointer to this variable's Naga IR value.
+    ///
+    /// If the current function uses this global variable, and it has been
+    /// implicitly wrapped in an `OpTypeStruct`, this is the id of an
+    /// `OpAccessChain` instruction in the function's prologue that refers to
+    /// the wrapped value inside the struct. (This value is assigned as we write
+    /// the prologue code of each function.) If you need the wrapper struct
+    /// itself, use [`var_id`] instead of this field.
     ///
-    /// Vulkan requires that globals in the `StorageBuffer` and `Uniform` storage
-    /// classes must be structs with the `Block` decoration, but WGSL and Naga IR
-    /// make no such requirement. So for such variables, we generate a wrapper struct
-    /// type with a single element of the type given by Naga, generate an
-    /// `OpAccessChain` for that member in the function prelude, and use that pointer
-    /// to refer to the global in the function body. This is the id of that access,
-    /// updated for each function in `write_function`.
+    /// If this global is not implicitly wrapped, this is the same as
+    /// [`var_id`].
+    ///
+    /// [`var_id`]: GlobalVariable::var_id
     access_id: Word,
 }
 
@@ -627,12 +665,6 @@ impl BlockContext<'_> {
     }
 }
 
-#[derive(Clone, Copy, Default)]
-struct LoopContext {
-    continuing_id: Option<Word>,
-    break_id: Option<Word>,
-}
-
 pub struct Writer {
     physical_layout: PhysicalLayout,
     logical_layout: LogicalLayout,
diff --git a/naga/src/back/spv/writer.rs b/naga/src/back/spv/writer.rs
index 678dcb4246..14f1fc0027 100644
--- a/naga/src/back/spv/writer.rs
+++ b/naga/src/back/spv/writer.rs
@@ -3,7 +3,7 @@ use super::{
     helpers::{contains_builtin, global_needs_wrapper, map_storage_class},
     make_local, Block, BlockContext, CachedConstant, CachedExpressions, DebugInfo,
     EntryPointContext, Error, Function, FunctionArgument, GlobalVariable, IdGenerator, Instruction,
-    LocalType, LocalVariable, LogicalLayout, LookupFunctionType, LookupType, LoopContext, Options,
+    LocalType, LocalVariable, LogicalLayout, LookupFunctionType, LookupType, Options,
     PhysicalLayout, PipelineOptions, ResultMember, Writer, WriterFlags, BITS_PER_BYTE,
 };
 use crate::{
@@ -32,6 +32,9 @@ impl Function {
                 for local_var in self.variables.values() {
                     local_var.instruction.to_words(sink);
                 }
+                for internal_var in self.internal_variables.iter() {
+                    internal_var.instruction.to_words(sink);
+                }
             }
             for instruction in block.body.iter() {
                 instruction.to_words(sink);
@@ -135,6 +138,56 @@ impl Writer {
         self.capabilities_used.insert(spirv::Capability::Shader);
     }
 
+    #[allow(clippy::too_many_arguments)]
+    pub(super) fn promote_access_expression_to_variable(
+        &mut self,
+        ir_types: &UniqueArena<crate::Type>,
+        result_type_id: Word,
+        container_id: Word,
+        container_ty: Handle<crate::Type>,
+        index_id: Word,
+        element_ty: Handle<crate::Type>,
+        block: &mut Block,
+    ) -> Result<(Word, LocalVariable), Error> {
+        let pointer_type_id =
+            self.get_pointer_id(ir_types, container_ty, spirv::StorageClass::Function)?;
+
+        let variable = {
+            let id = self.id_gen.next();
+            LocalVariable {
+                id,
+                instruction: Instruction::variable(
+                    pointer_type_id,
+                    id,
+                    spirv::StorageClass::Function,
+                    None,
+                ),
+            }
+        };
+        block
+            .body
+            .push(Instruction::store(variable.id, container_id, None));
+
+        let element_pointer_id = self.id_gen.next();
+        let element_pointer_type_id =
+            self.get_pointer_id(ir_types, element_ty, spirv::StorageClass::Function)?;
+        block.body.push(Instruction::access_chain(
+            element_pointer_type_id,
+            element_pointer_id,
+            variable.id,
+            &[index_id],
+        ));
+        let id = self.id_gen.next();
+        block.body.push(Instruction::load(
+            result_type_id,
+            id,
+            element_pointer_id,
+            None,
+        ));
+
+        Ok((id, variable))
+    }
+
     /// Indicate that the code requires any one of the listed capabilities.
     ///
     /// If nothing in `capabilities` appears in the available capabilities
@@ -703,13 +756,7 @@ impl Writer {
             next_id
         };
 
-        context.write_block(
-            main_id,
-            &ir_function.body,
-            super::block::BlockExit::Return,
-            LoopContext::default(),
-            debug_info.as_ref(),
-        )?;
+        context.write_function_body(main_id, debug_info.as_ref())?;
 
         // Consume the `BlockContext`, ending its borrows and letting the
         // `Writer` steal back its cached expression table and temp_list.
diff --git a/naga/src/front/glsl/context.rs b/naga/src/front/glsl/context.rs
index ee1fcc04ba..b4cb1c874e 100644
--- a/naga/src/front/glsl/context.rs
+++ b/naga/src/front/glsl/context.rs
@@ -630,7 +630,8 @@ impl<'a> Context<'a> {
                             frontend.errors.push(Error {
                                 kind: ErrorKind::SemanticError(
                                     format!(
-                                        "Cannot apply operation to {left_inner:?} and {right_inner:?}"
+                                        "Cannot apply operation to {:?} and {:?}",
+                                        left_inner, right_inner
                                     )
                                     .into(),
                                 ),
@@ -828,7 +829,8 @@ impl<'a> Context<'a> {
                             frontend.errors.push(Error {
                                 kind: ErrorKind::SemanticError(
                                     format!(
-                                        "Cannot apply operation to {left_inner:?} and {right_inner:?}"
+                                        "Cannot apply operation to {:?} and {:?}",
+                                        left_inner, right_inner
                                     )
                                     .into(),
                                 ),
@@ -908,7 +910,8 @@ impl<'a> Context<'a> {
                             frontend.errors.push(Error {
                                 kind: ErrorKind::SemanticError(
                                     format!(
-                                        "Cannot apply operation to {left_inner:?} and {right_inner:?}"
+                                        "Cannot apply operation to {:?} and {:?}",
+                                        left_inner, right_inner
                                     )
                                     .into(),
                                 ),
diff --git a/naga/src/front/glsl/functions.rs b/naga/src/front/glsl/functions.rs
index a1a6038263..2a63c7a028 100644
--- a/naga/src/front/glsl/functions.rs
+++ b/naga/src/front/glsl/functions.rs
@@ -634,7 +634,8 @@ impl Frontend {
                         self.errors.push(Error {
                             kind: ErrorKind::SemanticError(
                                 format!(
-                                    "'{name}': image needs {overload_access:?} access but only {call_access:?} was provided"
+                                    "'{}': image needs {:?} access but only {:?} was provided",
+                                    name, overload_access, call_access
                                 )
                                 .into(),
                             ),
diff --git a/naga/src/front/glsl/parser/expressions.rs b/naga/src/front/glsl/parser/expressions.rs
index 594ad6a6cd..c218e7b115 100644
--- a/naga/src/front/glsl/parser/expressions.rs
+++ b/naga/src/front/glsl/parser/expressions.rs
@@ -38,7 +38,13 @@ impl<'source> ParsingContext<'source> {
             TokenValue::FloatConstant(float) => {
                 if float.width != 32 {
                     frontend.errors.push(Error {
-                        kind: ErrorKind::SemanticError("Unsupported floating-point value (expected single-precision floating-point number)".into()),
+                        kind: ErrorKind::SemanticError(
+                            concat!(
+                                "Unsupported floating-point value ",
+                                "(expected single-precision floating-point number)"
+                            )
+                            .into(),
+                        ),
                         meta: token.meta,
                     });
                 }
diff --git a/naga/src/front/glsl/variables.rs b/naga/src/front/glsl/variables.rs
index 6b74b254bd..16c5bb65d6 100644
--- a/naga/src/front/glsl/variables.rs
+++ b/naga/src/front/glsl/variables.rs
@@ -294,14 +294,17 @@ impl Frontend {
                             .any(|i| components[i..].contains(&components[i - 1]));
                         if not_unique {
                             self.errors.push(Error {
-                                kind:
-                                ErrorKind::SemanticError(
-                                format!(
-                                    "swizzle cannot have duplicate components in left-hand-side expression for \"{name:?}\""
-                                )
-                                .into(),
-                            ),
-                                meta ,
+                                kind: ErrorKind::SemanticError(
+                                    format!(
+                                        concat!(
+                                            "swizzle cannot have duplicate components in ",
+                                            "left-hand-side expression for \"{:?}\""
+                                        ),
+                                        name
+                                    )
+                                    .into(),
+                                ),
+                                meta,
                             })
                         }
                     }
diff --git a/naga/src/front/spv/error.rs b/naga/src/front/spv/error.rs
index 219048e102..898113d446 100644
--- a/naga/src/front/spv/error.rs
+++ b/naga/src/front/spv/error.rs
@@ -47,7 +47,13 @@ pub enum Error {
     UnsupportedBinaryOperator(spirv::Word),
     #[error("Naga supports OpTypeRuntimeArray in the StorageBuffer storage class only")]
     UnsupportedRuntimeArrayStorageClass,
-    #[error("unsupported matrix stride {stride} for a {columns}x{rows} matrix with scalar width={width}")]
+    #[error(
+        "unsupported matrix stride {} for a {}x{} matrix with scalar width={}",
+        stride,
+        columns,
+        rows,
+        width
+    )]
     UnsupportedMatrixStride {
         stride: u32,
         columns: u8,
diff --git a/naga/src/front/wgsl/error.rs b/naga/src/front/wgsl/error.rs
index a7986ec898..7c65d93de3 100644
--- a/naga/src/front/wgsl/error.rs
+++ b/naga/src/front/wgsl/error.rs
@@ -298,32 +298,42 @@ impl<'a> Error<'a> {
         match *self {
             Error::Unexpected(unexpected_span, expected) => {
                 let expected_str = match expected {
-                    ExpectedToken::Token(token) => {
-                        match token {
-                            Token::Separator(c) => format!("'{c}'"),
-                            Token::Paren(c) => format!("'{c}'"),
-                            Token::Attribute => "@".to_string(),
-                            Token::Number(_) => "number".to_string(),
-                            Token::Word(s) => s.to_string(),
-                            Token::Operation(c) => format!("operation ('{c}')"),
-                            Token::LogicalOperation(c) => format!("logical operation ('{c}')"),
-                            Token::ShiftOperation(c) => format!("bitshift ('{c}{c}')"),
-                            Token::AssignmentOperation(c) if c=='<' || c=='>' => format!("bitshift ('{c}{c}=')"),
-                            Token::AssignmentOperation(c) => format!("operation ('{c}=')"),
-                            Token::IncrementOperation => "increment operation".to_string(),
-                            Token::DecrementOperation => "decrement operation".to_string(),
-                            Token::Arrow => "->".to_string(),
-                            Token::Unknown(c) => format!("unknown ('{c}')"),
-                            Token::Trivia => "trivia".to_string(),
-                            Token::End => "end".to_string(),
+                    ExpectedToken::Token(token) => match token {
+                        Token::Separator(c) => format!("'{c}'"),
+                        Token::Paren(c) => format!("'{c}'"),
+                        Token::Attribute => "@".to_string(),
+                        Token::Number(_) => "number".to_string(),
+                        Token::Word(s) => s.to_string(),
+                        Token::Operation(c) => format!("operation ('{c}')"),
+                        Token::LogicalOperation(c) => format!("logical operation ('{c}')"),
+                        Token::ShiftOperation(c) => format!("bitshift ('{c}{c}')"),
+                        Token::AssignmentOperation(c) if c == '<' || c == '>' => {
+                            format!("bitshift ('{c}{c}=')")
                         }
-                    }
+                        Token::AssignmentOperation(c) => format!("operation ('{c}=')"),
+                        Token::IncrementOperation => "increment operation".to_string(),
+                        Token::DecrementOperation => "decrement operation".to_string(),
+                        Token::Arrow => "->".to_string(),
+                        Token::Unknown(c) => format!("unknown ('{c}')"),
+                        Token::Trivia => "trivia".to_string(),
+                        Token::End => "end".to_string(),
+                    },
                     ExpectedToken::Identifier => "identifier".to_string(),
                     ExpectedToken::PrimaryExpression => "expression".to_string(),
                     ExpectedToken::Assignment => "assignment or increment/decrement".to_string(),
-                    ExpectedToken::SwitchItem => "switch item ('case' or 'default') or a closing curly bracket to signify the end of the switch statement ('}')".to_string(),
-                    ExpectedToken::WorkgroupSizeSeparator => "workgroup size separator (',') or a closing parenthesis".to_string(),
-                    ExpectedToken::GlobalItem => "global item ('struct', 'const', 'var', 'alias', ';', 'fn') or the end of the file".to_string(),
+                    ExpectedToken::SwitchItem => concat!(
+                        "switch item ('case' or 'default') or a closing curly bracket ",
+                        "to signify the end of the switch statement ('}')"
+                    )
+                    .to_string(),
+                    ExpectedToken::WorkgroupSizeSeparator => {
+                        "workgroup size separator (',') or a closing parenthesis".to_string()
+                    }
+                    ExpectedToken::GlobalItem => concat!(
+                        "global item ('struct', 'const', 'var', 'alias', ';', 'fn') ",
+                        "or the end of the file"
+                    )
+                    .to_string(),
                     ExpectedToken::Type => "type".to_string(),
                     ExpectedToken::Variable => "variable access".to_string(),
                     ExpectedToken::Function => "function name".to_string(),
@@ -384,9 +394,11 @@ impl<'a> Error<'a> {
                 notes: vec![],
             },
             Error::BadIncrDecrReferenceType(span) => ParseError {
-                message:
-                    "increment/decrement operation requires reference type to be one of i32 or u32"
-                        .to_string(),
+                message: concat!(
+                    "increment/decrement operation requires ",
+                    "reference type to be one of i32 or u32"
+                )
+                .to_string(),
                 labels: vec![(span, "must be a reference type of i32 or u32".into())],
                 notes: vec![],
             },
@@ -527,25 +539,24 @@ impl<'a> Error<'a> {
                 labels: vec![(span, "type can't be inferred".into())],
                 notes: vec![],
             },
-            Error::InitializationTypeMismatch { name, ref expected, ref got } => {
-                ParseError {
-                    message: format!(
-                        "the type of `{}` is expected to be `{}`, but got `{}`",
-                        &source[name], expected, got,
-                    ),
-                    labels: vec![(
-                        name,
-                        format!("definition of `{}`", &source[name]).into(),
-                    )],
-                    notes: vec![],
-                }
-            }
+            Error::InitializationTypeMismatch {
+                name,
+                ref expected,
+                ref got,
+            } => ParseError {
+                message: format!(
+                    "the type of `{}` is expected to be `{}`, but got `{}`",
+                    &source[name], expected, got,
+                ),
+                labels: vec![(name, format!("definition of `{}`", &source[name]).into())],
+                notes: vec![],
+            },
             Error::DeclMissingTypeAndInit(name_span) => ParseError {
-                message: format!("declaration of `{}` needs a type specifier or initializer", &source[name_span]),
-                labels: vec![(
-                    name_span,
-                    "needs a type specifier or initializer".into(),
-                )],
+                message: format!(
+                    "declaration of `{}` needs a type specifier or initializer",
+                    &source[name_span]
+                ),
+                labels: vec![(name_span, "needs a type specifier or initializer".into())],
                 notes: vec![],
             },
             Error::MissingAttribute(name, name_span) => ParseError {
@@ -725,7 +736,11 @@ impl<'a> Error<'a> {
                 notes: vec![message.into()],
             },
             Error::ExpectedConstExprConcreteIntegerScalar(span) => ParseError {
-                message: "must be a const-expression that resolves to a concrete integer scalar (u32 or i32)".to_string(),
+                message: concat!(
+                    "must be a const-expression that ",
+                    "resolves to a concrete integer scalar (u32 or i32)"
+                )
+                .to_string(),
                 labels: vec![(span, "must resolve to u32 or i32".into())],
                 notes: vec![],
             },
@@ -754,9 +769,17 @@ impl<'a> Error<'a> {
             },
             Error::AutoConversion(ref error) => {
                 // destructuring ensures all fields are handled
-                let AutoConversionError { dest_span, ref dest_type, source_span, ref source_type } = **error;
+                let AutoConversionError {
+                    dest_span,
+                    ref dest_type,
+                    source_span,
+                    ref source_type,
+                } = **error;
                 ParseError {
-                    message: format!("automatic conversions cannot convert `{source_type}` to `{dest_type}`"),
+                    message: format!(
+                        "automatic conversions cannot convert `{}` to `{}`",
+                        source_type, dest_type
+                    ),
                     labels: vec![
                         (
                             dest_span,
@@ -765,72 +788,77 @@ impl<'a> Error<'a> {
                         (
                             source_span,
                             format!("this expression has type {source_type}").into(),
-                        )
+                        ),
                     ],
                     notes: vec![],
                 }
-            },
+            }
             Error::AutoConversionLeafScalar(ref error) => {
-                let AutoConversionLeafScalarError { dest_span, ref dest_scalar, source_span, ref source_type } = **error;
+                let AutoConversionLeafScalarError {
+                    dest_span,
+                    ref dest_scalar,
+                    source_span,
+                    ref source_type,
+                } = **error;
                 ParseError {
-                    message: format!("automatic conversions cannot convert elements of `{source_type}` to `{dest_scalar}`"),
+                    message: format!(
+                        "automatic conversions cannot convert elements of `{}` to `{}`",
+                        source_type, dest_scalar
+                    ),
                     labels: vec![
                         (
                             dest_span,
-                            format!("a value with elements of type {dest_scalar} is required here").into(),
+                            format!(
+                                "a value with elements of type {} is required here",
+                                dest_scalar
+                            )
+                            .into(),
                         ),
                         (
                             source_span,
                             format!("this expression has type {source_type}").into(),
-                        )
+                        ),
                     ],
                     notes: vec![],
                 }
-            },
+            }
             Error::ConcretizationFailed(ref error) => {
-                let ConcretizationFailedError { expr_span, ref expr_type, ref scalar, ref inner } = **error;
+                let ConcretizationFailedError {
+                    expr_span,
+                    ref expr_type,
+                    ref scalar,
+                    ref inner,
+                } = **error;
                 ParseError {
                     message: format!("failed to convert expression to a concrete type: {inner}"),
-                    labels: vec![
-                        (
-                            expr_span,
-                            format!("this expression has type {expr_type}").into(),
-                        )
-                    ],
-                    notes: vec![
-                        format!("the expression should have been converted to have {} scalar type", scalar),
-                    ]
+                    labels: vec![(
+                        expr_span,
+                        format!("this expression has type {expr_type}").into(),
+                    )],
+                    notes: vec![format!(
+                        "the expression should have been converted to have {} scalar type",
+                        scalar
+                    )],
                 }
-            },
+            }
             Error::ExceededLimitForNestedBraces { span, limit } => ParseError {
                 message: "brace nesting limit reached".into(),
                 labels: vec![(span, "limit reached at this brace".into())],
-                notes: vec![
-                    format!("nesting limit is currently set to {limit}"),
-                ],
+                notes: vec![format!("nesting limit is currently set to {limit}")],
             },
             Error::PipelineConstantIDValue(span) => ParseError {
                 message: "pipeline constant ID must be between 0 and 65535 inclusive".to_string(),
-                labels: vec![(
-                    span,
-                    "must be between 0 and 65535 inclusive".into(),
-                )],
+                labels: vec![(span, "must be between 0 and 65535 inclusive".into())],
                 notes: vec![],
             },
             Error::NotBool(span) => ParseError {
                 message: "must be a const-expression that resolves to a bool".to_string(),
-                labels: vec![(
-                    span,
-                    "must resolve to bool".into(),
-                )],
+                labels: vec![(span, "must resolve to bool".into())],
                 notes: vec![],
             },
             Error::ConstAssertFailed(span) => ParseError {
                 message: "const_assert failure".to_string(),
-                labels: vec![(
-                    span,
-                    "evaluates to false".into(),
-                )],
+                labels: vec![(span, "evaluates to false".into())],
                 notes: vec![],
             },
         }
diff --git a/naga/src/lib.rs b/naga/src/lib.rs
index 85fd7a4508..038e215a6a 100644
--- a/naga/src/lib.rs
+++ b/naga/src/lib.rs
@@ -1402,21 +1402,20 @@ pub enum Expression {
     /// ## Dynamic indexing restrictions
     ///
     /// To accommodate restrictions in some of the shader languages that Naga
-    /// targets, it is not permitted to subscript a matrix or array with a
-    /// dynamically computed index unless that matrix or array appears behind a
-    /// pointer. In other words, if the inner type of `base` is [`Array`] or
-    /// [`Matrix`], then `index` must be a constant. But if the type of `base`
-    /// is a [`Pointer`] to an array or matrix or a [`ValuePointer`] with a
-    /// `size`, then the index may be any expression of integer type.
+    /// targets, it is not permitted to subscript a matrix with a dynamically
+    /// computed index unless that matrix appears behind a pointer. In other
+    /// words, if the inner type of `base` is [`Matrix`], then `index` must be a
+    /// constant. But if the type of `base` is a [`Pointer`] to an matrix, then
+    /// the index may be any expression of integer type.
     ///
     /// You can use the [`Expression::is_dynamic_index`] method to determine
-    /// whether a given index expression requires matrix or array base operands
-    /// to be behind a pointer.
+    /// whether a given index expression requires matrix base operands to be
+    /// behind a pointer.
     ///
     /// (It would be simpler to always require the use of `AccessIndex` when
-    /// subscripting arrays and matrices that are not behind pointers, but to
-    /// accommodate existing front ends, Naga also permits `Access`, with a
-    /// restricted `index`.)
+    /// subscripting matrices that are not behind pointers, but to accommodate
+    /// existing front ends, Naga also permits `Access`, with a restricted
+    /// `index`.)
     ///
     /// [`Vector`]: TypeInner::Vector
     /// [`Matrix`]: TypeInner::Matrix
diff --git a/naga/src/proc/mod.rs b/naga/src/proc/mod.rs
index a5b3ea4e38..abbe0c7e46 100644
--- a/naga/src/proc/mod.rs
+++ b/naga/src/proc/mod.rs
@@ -521,12 +521,12 @@ impl crate::Expression {
         }
     }
 
-    /// Return true if this expression is a dynamic array index, for [`Access`].
+    /// Return true if this expression is a dynamic array/vector/matrix index,
+    /// for [`Access`].
     ///
     /// This method returns true if this expression is a dynamically computed
-    /// index, and as such can only be used to index matrices and arrays when
-    /// they appear behind a pointer. See the documentation for [`Access`] for
-    /// details.
+    /// index, and as such can only be used to index matrices when they appear
+    /// behind a pointer. See the documentation for [`Access`] for details.
     ///
     /// Note, this does not check the _type_ of the given expression. It's up to
     /// the caller to establish that the `Access` expression is well-typed
diff --git a/naga/src/proc/typifier.rs b/naga/src/proc/typifier.rs
index d8af0cd236..04c776365f 100644
--- a/naga/src/proc/typifier.rs
+++ b/naga/src/proc/typifier.rs
@@ -92,6 +92,13 @@ pub enum TypeResolution {
     /// available in the associated arena. However, the `TypeInner` itself may
     /// contain `Handle<Type>` values referring to types from the arena.
     ///
+    /// The inner type must only be one of the following variants:
+    /// - TypeInner::Pointer
+    /// - TypeInner::ValuePointer
+    /// - TypeInner::Matrix (generated by matrix multiplication)
+    /// - TypeInner::Vector
+    /// - TypeInner::Scalar
+    ///
     /// [`TypeInner`]: crate::TypeInner
     Value(crate::TypeInner),
 }
@@ -631,41 +638,37 @@ impl<'a> ResolveContext<'a> {
                 use crate::MathFunction as Mf;
                 let res_arg = past(arg)?;
                 match fun {
-                    // comparison
-                    Mf::Abs |
-                    Mf::Min |
-                    Mf::Max |
-                    Mf::Clamp |
-                    Mf::Saturate |
-                    // trigonometry
-                    Mf::Cos |
-                    Mf::Cosh |
-                    Mf::Sin |
-                    Mf::Sinh |
-                    Mf::Tan |
-                    Mf::Tanh |
-                    Mf::Acos |
-                    Mf::Asin |
-                    Mf::Atan |
-                    Mf::Atan2 |
-                    Mf::Asinh |
-                    Mf::Acosh |
-                    Mf::Atanh |
-                    Mf::Radians |
-                    Mf::Degrees |
-                    // decomposition
-                    Mf::Ceil |
-                    Mf::Floor |
-                    Mf::Round |
-                    Mf::Fract |
-                    Mf::Trunc |
-                    Mf::Ldexp |
-                    // exponent
-                    Mf::Exp |
-                    Mf::Exp2 |
-                    Mf::Log |
-                    Mf::Log2 |
-                    Mf::Pow => res_arg.clone(),
+                    Mf::Abs
+                    | Mf::Min
+                    | Mf::Max
+                    | Mf::Clamp
+                    | Mf::Saturate
+                    | Mf::Cos
+                    | Mf::Cosh
+                    | Mf::Sin
+                    | Mf::Sinh
+                    | Mf::Tan
+                    | Mf::Tanh
+                    | Mf::Acos
+                    | Mf::Asin
+                    | Mf::Atan
+                    | Mf::Atan2
+                    | Mf::Asinh
+                    | Mf::Acosh
+                    | Mf::Atanh
+                    | Mf::Radians
+                    | Mf::Degrees
+                    | Mf::Ceil
+                    | Mf::Floor
+                    | Mf::Round
+                    | Mf::Fract
+                    | Mf::Trunc
+                    | Mf::Ldexp
+                    | Mf::Exp
+                    | Mf::Exp2
+                    | Mf::Log
+                    | Mf::Log2
+                    | Mf::Pow => res_arg.clone(),
                     Mf::Modf | Mf::Frexp => {
                         let (size, width) = match res_arg.inner_with(types) {
                             &Ti::Scalar(crate::Scalar {
@@ -673,77 +676,81 @@ impl<'a> ResolveContext<'a> {
                                 width,
                             }) => (None, width),
                             &Ti::Vector {
-                                scalar: crate::Scalar {
-                                    kind: crate::ScalarKind::Float,
-                                    width,
-                                },
+                                scalar:
+                                    crate::Scalar {
+                                        kind: crate::ScalarKind::Float,
+                                        width,
+                                    },
                                 size,
                             } => (Some(size), width),
-                            ref other =>
-                                return Err(ResolveError::IncompatibleOperands(format!("{fun:?}({other:?}, _)")))
+                            ref other => {
+                                return Err(ResolveError::IncompatibleOperands(format!(
+                                    "{fun:?}({other:?}, _)"
+                                )))
+                            }
                         };
                         let result = self
-                        .special_types
-                        .predeclared_types
-                        .get(&if fun == Mf::Modf {
-                            crate::PredeclaredType::ModfResult { size, width }
-                    } else {
-                            crate::PredeclaredType::FrexpResult { size, width }
-                    })
-                        .ok_or(ResolveError::MissingSpecialType)?;
+                            .special_types
+                            .predeclared_types
+                            .get(&if fun == Mf::Modf {
+                                crate::PredeclaredType::ModfResult { size, width }
+                            } else {
+                                crate::PredeclaredType::FrexpResult { size, width }
+                            })
+                            .ok_or(ResolveError::MissingSpecialType)?;
                         TypeResolution::Handle(*result)
-                    },
-                    // geometry
+                    }
                     Mf::Dot => match *res_arg.inner_with(types) {
-                        Ti::Vector {
-                            size: _,
-                            scalar,
-                        } => TypeResolution::Value(Ti::Scalar(scalar)),
-                        ref other =>
-                            return Err(ResolveError::IncompatibleOperands(
-                                format!("{fun:?}({other:?}, _)")
-                            )),
+                        Ti::Vector { size: _, scalar } => TypeResolution::Value(Ti::Scalar(scalar)),
+                        ref other => {
+                            return Err(ResolveError::IncompatibleOperands(format!(
+                                "{fun:?}({other:?}, _)"
+                            )))
+                        }
                     },
                     Mf::Outer => {
-                        let arg1 = arg1.ok_or_else(|| ResolveError::IncompatibleOperands(
-                            format!("{fun:?}(_, None)")
-                        ))?;
+                        let arg1 = arg1.ok_or_else(|| {
+                            ResolveError::IncompatibleOperands(format!("{fun:?}(_, None)"))
+                        })?;
                         match (res_arg.inner_with(types), past(arg1)?.inner_with(types)) {
                             (
-                                &Ti::Vector { size: columns, scalar },
-                                &Ti::Vector{ size: rows, .. }
+                                &Ti::Vector {
+                                    size: columns,
+                                    scalar,
+                                },
+                                &Ti::Vector { size: rows, .. },
                             ) => TypeResolution::Value(Ti::Matrix {
                                 columns,
                                 rows,
                                 scalar,
                             }),
-                            (left, right) =>
-                                return Err(ResolveError::IncompatibleOperands(
-                                    format!("{fun:?}({left:?}, {right:?})")
-                                )),
+                            (left, right) => {
+                                return Err(ResolveError::IncompatibleOperands(format!(
+                                    "{fun:?}({left:?}, {right:?})"
+                                )))
+                            }
                         }
-                    },
+                    }
                     Mf::Cross => res_arg.clone(),
-                    Mf::Distance |
-                    Mf::Length => match *res_arg.inner_with(types) {
-                        Ti::Scalar(scalar) |
-                        Ti::Vector {scalar,size:_} => TypeResolution::Value(Ti::Scalar(scalar)),
-                        ref other => return Err(ResolveError::IncompatibleOperands(
-                                format!("{fun:?}({other:?})")
-                            )),
+                    Mf::Distance | Mf::Length => match *res_arg.inner_with(types) {
+                        Ti::Scalar(scalar) | Ti::Vector { scalar, size: _ } => {
+                            TypeResolution::Value(Ti::Scalar(scalar))
+                        }
+                        ref other => {
+                            return Err(ResolveError::IncompatibleOperands(format!(
+                                "{fun:?}({other:?})"
+                            )))
+                        }
                     },
-                    Mf::Normalize |
-                    Mf::FaceForward |
-                    Mf::Reflect |
-                    Mf::Refract => res_arg.clone(),
+                    Mf::Normalize | Mf::FaceForward | Mf::Reflect | Mf::Refract => res_arg.clone(),
                     // computational
-                    Mf::Sign |
-                    Mf::Fma |
-                    Mf::Mix |
-                    Mf::Step |
-                    Mf::SmoothStep |
-                    Mf::Sqrt |
-                    Mf::InverseSqrt => res_arg.clone(),
+                    Mf::Sign
+                    | Mf::Fma
+                    | Mf::Mix
+                    | Mf::Step
+                    | Mf::SmoothStep
+                    | Mf::Sqrt
+                    | Mf::InverseSqrt => res_arg.clone(),
                     Mf::Transpose => match *res_arg.inner_with(types) {
                         Ti::Matrix {
                             columns,
@@ -754,9 +761,11 @@ impl<'a> ResolveContext<'a> {
                             rows: columns,
                             scalar,
                         }),
-                        ref other => return Err(ResolveError::IncompatibleOperands(
-                            format!("{fun:?}({other:?})")
-                        )),
+                        ref other => {
+                            return Err(ResolveError::IncompatibleOperands(format!(
+                                "{fun:?}({other:?})"
+                            )))
+                        }
                     },
                     Mf::Inverse => match *res_arg.inner_with(types) {
                         Ti::Matrix {
@@ -768,70 +777,75 @@ impl<'a> ResolveContext<'a> {
                             rows,
                             scalar,
                         }),
-                        ref other => return Err(ResolveError::IncompatibleOperands(
-                            format!("{fun:?}({other:?})")
-                        )),
+                        ref other => {
+                            return Err(ResolveError::IncompatibleOperands(format!(
+                                "{fun:?}({other:?})"
+                            )))
+                        }
                     },
                     Mf::Determinant => match *res_arg.inner_with(types) {
-                        Ti::Matrix {
-                            scalar,
-                            ..
-                        } => TypeResolution::Value(Ti::Scalar(scalar)),
-                        ref other => return Err(ResolveError::IncompatibleOperands(
-                            format!("{fun:?}({other:?})")
-                        )),
+                        Ti::Matrix { scalar, .. } => TypeResolution::Value(Ti::Scalar(scalar)),
+                        ref other => {
+                            return Err(ResolveError::IncompatibleOperands(format!(
+                                "{fun:?}({other:?})"
+                            )))
+                        }
                     },
                     // bits
-                    Mf::CountTrailingZeros |
-                    Mf::CountLeadingZeros |
-                    Mf::CountOneBits |
-                    Mf::ReverseBits |
-                    Mf::ExtractBits |
-                    Mf::InsertBits |
-                    Mf::FirstTrailingBit |
-                    Mf::FirstLeadingBit => match *res_arg.inner_with(types)  {
-                        Ti::Scalar(scalar @ crate::Scalar {
-                            kind: crate::ScalarKind::Sint | crate::ScalarKind::Uint,
-                            ..
-                        }) => TypeResolution::Value(Ti::Scalar(scalar)),
-                        Ti::Vector {
-                            size,
-                            scalar: scalar @ crate::Scalar {
+                    Mf::CountTrailingZeros
+                    | Mf::CountLeadingZeros
+                    | Mf::CountOneBits
+                    | Mf::ReverseBits
+                    | Mf::ExtractBits
+                    | Mf::InsertBits
+                    | Mf::FirstTrailingBit
+                    | Mf::FirstLeadingBit => match *res_arg.inner_with(types) {
+                        Ti::Scalar(
+                            scalar @ crate::Scalar {
                                 kind: crate::ScalarKind::Sint | crate::ScalarKind::Uint,
                                 ..
-                            }
+                            },
+                        ) => TypeResolution::Value(Ti::Scalar(scalar)),
+                        Ti::Vector {
+                            size,
+                            scalar:
+                                scalar @ crate::Scalar {
+                                    kind: crate::ScalarKind::Sint | crate::ScalarKind::Uint,
+                                    ..
+                                },
                         } => TypeResolution::Value(Ti::Vector { size, scalar }),
-                        ref other => return Err(ResolveError::IncompatibleOperands(
-                                format!("{fun:?}({other:?})")
-                            )),
+                        ref other => {
+                            return Err(ResolveError::IncompatibleOperands(format!(
+                                "{fun:?}({other:?})"
+                            )))
+                        }
                     },
                     // data packing
-                    Mf::Pack4x8snorm |
-                    Mf::Pack4x8unorm |
-                    Mf::Pack2x16snorm |
-                    Mf::Pack2x16unorm |
-                    Mf::Pack2x16float |
-                    Mf::Pack4xI8 |
-                    Mf::Pack4xU8 => TypeResolution::Value(Ti::Scalar(crate::Scalar::U32)),
+                    Mf::Pack4x8snorm
+                    | Mf::Pack4x8unorm
+                    | Mf::Pack2x16snorm
+                    | Mf::Pack2x16unorm
+                    | Mf::Pack2x16float
+                    | Mf::Pack4xI8
+                    | Mf::Pack4xU8 => TypeResolution::Value(Ti::Scalar(crate::Scalar::U32)),
                     // data unpacking
-                    Mf::Unpack4x8snorm |
-                    Mf::Unpack4x8unorm => TypeResolution::Value(Ti::Vector {
+                    Mf::Unpack4x8snorm | Mf::Unpack4x8unorm => TypeResolution::Value(Ti::Vector {
                         size: crate::VectorSize::Quad,
-                        scalar: crate::Scalar::F32
-                    }),
-                    Mf::Unpack2x16snorm |
-                    Mf::Unpack2x16unorm |
-                    Mf::Unpack2x16float => TypeResolution::Value(Ti::Vector {
-                        size: crate::VectorSize::Bi,
-                        scalar: crate::Scalar::F32
+                        scalar: crate::Scalar::F32,
                     }),
+                    Mf::Unpack2x16snorm | Mf::Unpack2x16unorm | Mf::Unpack2x16float => {
+                        TypeResolution::Value(Ti::Vector {
+                            size: crate::VectorSize::Bi,
+                            scalar: crate::Scalar::F32,
+                        })
+                    }
                     Mf::Unpack4xI8 => TypeResolution::Value(Ti::Vector {
                         size: crate::VectorSize::Quad,
-                        scalar: crate::Scalar::I32
+                        scalar: crate::Scalar::I32,
                     }),
                     Mf::Unpack4xU8 => TypeResolution::Value(Ti::Vector {
                         size: crate::VectorSize::Quad,
-                        scalar: crate::Scalar::U32
+                        scalar: crate::Scalar::U32,
                     }),
                 }
             }
diff --git a/naga/src/span.rs b/naga/src/span.rs
index 82cfbe5a4b..f8a0f67fbe 100644
--- a/naga/src/span.rs
+++ b/naga/src/span.rs
@@ -11,6 +11,7 @@ pub struct Span {
 
 impl Span {
     pub const UNDEFINED: Self = Self { start: 0, end: 0 };
+
     /// Creates a new `Span` from a range of byte indices
     ///
     /// Note: end is exclusive, it doesn't belong to the `Span`
diff --git a/naga/src/valid/expression.rs b/naga/src/valid/expression.rs
index 0b0d115c57..2b479d3a73 100644
--- a/naga/src/valid/expression.rs
+++ b/naga/src/valid/expression.rs
@@ -240,9 +240,10 @@ impl super::Validator {
                 let base_type = &resolver[base];
                 // See the documentation for `Expression::Access`.
                 let dynamic_indexing_restricted = match *base_type {
-                    Ti::Vector { .. } => false,
-                    Ti::Matrix { .. } | Ti::Array { .. } => true,
-                    Ti::Pointer { .. }
+                    Ti::Matrix { .. } => true,
+                    Ti::Vector { .. }
+                    | Ti::Array { .. }
+                    | Ti::Pointer { .. }
                     | Ti::ValuePointer { size: Some(_), .. }
                     | Ti::BindingArray { .. } => false,
                     ref other => {
diff --git a/naga/tests/in/6220-break-from-loop.param.ron b/naga/tests/in/6220-break-from-loop.param.ron
new file mode 100644
index 0000000000..72873dd667
--- /dev/null
+++ b/naga/tests/in/6220-break-from-loop.param.ron
@@ -0,0 +1,2 @@
+(
+)
diff --git a/naga/tests/in/6220-break-from-loop.wgsl b/naga/tests/in/6220-break-from-loop.wgsl
new file mode 100644
index 0000000000..424886a757
--- /dev/null
+++ b/naga/tests/in/6220-break-from-loop.wgsl
@@ -0,0 +1,43 @@
+// #6220: Don't generate unreachable SPIR-V blocks that branch into
+// structured control flow constructs.
+//
+// Suppose we have Naga code like this:
+//
+//     Block {
+//       ... prelude
+//       Block { ... nested }
+//       ... postlude
+//     }
+//
+// The SPIR-V back end used to always generate three separate SPIR-V
+// blocks for the sections labeled "prelude", "nested", and
+// "postlude", each block ending with a branch to the next, even if
+// they were empty.
+//
+// However, the function below generates code that includes the
+// following structure:
+//
+//     Loop {
+//       body: Block {
+//         ... prelude
+//         Block { Break }
+//         ... postlude
+//       }
+//       continuing: ...
+//     }
+//
+// In this case, even though the `Break` renders the "postlude"
+// unreachable, we used to generate a SPIR-V block for it anyway,
+// ending with a branch to the `Loop`'s "continuing" block. However,
+// SPIR-V's structured control flow rules forbid branches to a loop
+// construct's continue target from outside the loop, so the SPIR-V
+// module containing the unreachable block didn't pass validation.
+//
+// One might assume that unreachable blocks shouldn't affect
+// validation, but the spec doesn't clearly agree, and this doesn't
+// seem to be the way validation has been implemented.
+fn break_from_loop() {
+    for (var i = 0; i < 4; i += 1) {
+      break;
+    }
+}
diff --git a/naga/tests/in/access.wgsl b/naga/tests/in/access.wgsl
index 956a694aaa..3336522fd9 100644
--- a/naga/tests/in/access.wgsl
+++ b/naga/tests/in/access.wgsl
@@ -167,3 +167,14 @@ fn assign_through_ptr() {
 	var arr = array<vec4<f32>, 2>(vec4(6.0), vec4(7.0));
     assign_array_through_ptr_fn(&arr);
 }
+
+@vertex
+fn foo(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
+	let arr = array<i32, 5>(1, 2, 3, 4, 5);
+	let value = arr[vi];
+	return vec4<f32>(vec4<i32>(value));
+}
+
+fn array_by_value(a: array<i32, 5>, i: i32) -> i32 {
+    return a[i];
+}
diff --git a/naga/tests/out/analysis/access.info.ron b/naga/tests/out/analysis/access.info.ron
index 308bb1a8b6..830831cb1f 100644
--- a/naga/tests/out/analysis/access.info.ron
+++ b/naga/tests/out/analysis/access.info.ron
@@ -2735,6 +2735,54 @@
             sampling: [],
             dual_source_blending: false,
         ),
+        (
+            flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            uniformity: (
+                non_uniform_result: Some(0),
+                requirements: (""),
+            ),
+            may_kill: false,
+            sampling_set: [],
+            global_uses: [
+                (""),
+                (""),
+                (""),
+                (""),
+                (""),
+            ],
+            expressions: [
+                (
+                    uniformity: (
+                        non_uniform_result: Some(0),
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(25),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: Some(1),
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(2),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: Some(0),
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(2),
+                ),
+            ],
+            sampling: [],
+            dual_source_blending: false,
+        ),
     ],
     entry_points: [
         (
@@ -3981,6 +4029,144 @@
             sampling: [],
             dual_source_blending: false,
         ),
+        (
+            flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            uniformity: (
+                non_uniform_result: Some(0),
+                requirements: (""),
+            ),
+            may_kill: false,
+            sampling_set: [],
+            global_uses: [
+                (""),
+                (""),
+                (""),
+                (""),
+                (""),
+            ],
+            expressions: [
+                (
+                    uniformity: (
+                        non_uniform_result: Some(0),
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(0),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Sint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Sint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Sint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Sint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Sint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(25),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: Some(0),
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(2),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: Some(0),
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Vector(
+                        size: Quad,
+                        scalar: (
+                            kind: Sint,
+                            width: 4,
+                        ),
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: Some(0),
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Vector(
+                        size: Quad,
+                        scalar: (
+                            kind: Float,
+                            width: 4,
+                        ),
+                    )),
+                ),
+            ],
+            sampling: [],
+            dual_source_blending: false,
+        ),
     ],
     const_expression_types: [
         Value(Scalar((
diff --git a/naga/tests/out/glsl/access.assign_through_ptr.Compute.glsl b/naga/tests/out/glsl/access.assign_through_ptr.Compute.glsl
index 2e51bbde63..4a4791c994 100644
--- a/naga/tests/out/glsl/access.assign_through_ptr.Compute.glsl
+++ b/naga/tests/out/glsl/access.assign_through_ptr.Compute.glsl
@@ -20,8 +20,8 @@ struct MatCx2InArray {
     mat4x2 am[2];
 };
 
-float read_from_private(inout float foo_1) {
-    float _e1 = foo_1;
+float read_from_private(inout float foo_2) {
+    float _e1 = foo_2;
     return _e1;
 }
 
@@ -34,11 +34,15 @@ void assign_through_ptr_fn(inout uint p) {
     return;
 }
 
-void assign_array_through_ptr_fn(inout vec4 foo_2[2]) {
-    foo_2 = vec4[2](vec4(1.0), vec4(2.0));
+void assign_array_through_ptr_fn(inout vec4 foo_3[2]) {
+    foo_3 = vec4[2](vec4(1.0), vec4(2.0));
     return;
 }
 
+int array_by_value(int a_1[5], int i) {
+    return a_1[i];
+}
+
 void main() {
     uint val = 33u;
     vec4 arr[2] = vec4[2](vec4(6.0), vec4(7.0));
diff --git a/naga/tests/out/glsl/access.foo.Vertex.glsl b/naga/tests/out/glsl/access.foo.Vertex.glsl
new file mode 100644
index 0000000000..e1f3138403
--- /dev/null
+++ b/naga/tests/out/glsl/access.foo.Vertex.glsl
@@ -0,0 +1,52 @@
+#version 310 es
+
+precision highp float;
+precision highp int;
+
+struct GlobalConst {
+    uint a;
+    uvec3 b;
+    int c;
+};
+struct AlignedWrapper {
+    int value;
+};
+struct Baz {
+    mat3x2 m;
+};
+struct MatCx2InArray {
+    mat4x2 am[2];
+};
+
+float read_from_private(inout float foo_2) {
+    float _e1 = foo_2;
+    return _e1;
+}
+
+float test_arr_as_arg(float a[5][10]) {
+    return a[4][9];
+}
+
+void assign_through_ptr_fn(inout uint p) {
+    p = 42u;
+    return;
+}
+
+void assign_array_through_ptr_fn(inout vec4 foo_3[2]) {
+    foo_3 = vec4[2](vec4(1.0), vec4(2.0));
+    return;
+}
+
+int array_by_value(int a_1[5], int i) {
+    return a_1[i];
+}
+
+void main() {
+    uint vi_1 = uint(gl_VertexID);
+    int arr_1[5] = int[5](1, 2, 3, 4, 5);
+    int value = arr_1[vi_1];
+    gl_Position = vec4(ivec4(value));
+    gl_Position.yz = vec2(-gl_Position.y, gl_Position.z * 2.0 - gl_Position.w);
+    return;
+}
+
diff --git a/naga/tests/out/glsl/access.foo_frag.Fragment.glsl b/naga/tests/out/glsl/access.foo_frag.Fragment.glsl
index aacdda0130..eca6bc54c5 100644
--- a/naga/tests/out/glsl/access.foo_frag.Fragment.glsl
+++ b/naga/tests/out/glsl/access.foo_frag.Fragment.glsl
@@ -30,8 +30,8 @@ layout(std430) buffer type_13_block_1Fragment { ivec2 _group_0_binding_2_fs; };
 
 layout(location = 0) out vec4 _fs2p_location0;
 
-float read_from_private(inout float foo_1) {
-    float _e1 = foo_1;
+float read_from_private(inout float foo_2) {
+    float _e1 = foo_2;
     return _e1;
 }
 
@@ -44,11 +44,15 @@ void assign_through_ptr_fn(inout uint p) {
     return;
 }
 
-void assign_array_through_ptr_fn(inout vec4 foo_2[2]) {
-    foo_2 = vec4[2](vec4(1.0), vec4(2.0));
+void assign_array_through_ptr_fn(inout vec4 foo_3[2]) {
+    foo_3 = vec4[2](vec4(1.0), vec4(2.0));
     return;
 }
 
+int array_by_value(int a_1[5], int i) {
+    return a_1[i];
+}
+
 void main() {
     _group_0_binding_0_fs._matrix[1][2] = 1.0;
     _group_0_binding_0_fs._matrix = mat4x3(vec3(0.0), vec3(1.0), vec3(2.0), vec3(3.0));
diff --git a/naga/tests/out/glsl/access.foo_vert.Vertex.glsl b/naga/tests/out/glsl/access.foo_vert.Vertex.glsl
index d4a9b92945..a926eadf78 100644
--- a/naga/tests/out/glsl/access.foo_vert.Vertex.glsl
+++ b/naga/tests/out/glsl/access.foo_vert.Vertex.glsl
@@ -103,8 +103,8 @@ void test_matrix_within_array_within_struct_accesses() {
     return;
 }
 
-float read_from_private(inout float foo_1) {
-    float _e1 = foo_1;
+float read_from_private(inout float foo_2) {
+    float _e1 = foo_2;
     return _e1;
 }
 
@@ -117,11 +117,15 @@ void assign_through_ptr_fn(inout uint p) {
     return;
 }
 
-void assign_array_through_ptr_fn(inout vec4 foo_2[2]) {
-    foo_2 = vec4[2](vec4(1.0), vec4(2.0));
+void assign_array_through_ptr_fn(inout vec4 foo_3[2]) {
+    foo_3 = vec4[2](vec4(1.0), vec4(2.0));
     return;
 }
 
+int array_by_value(int a_1[5], int i) {
+    return a_1[i];
+}
+
 void main() {
     uint vi = uint(gl_VertexID);
     float foo = 0.0;
@@ -133,10 +137,10 @@ void main() {
     mat4x3 _matrix = _group_0_binding_0_vs._matrix;
     uvec2 arr_1[2] = _group_0_binding_0_vs.arr;
     float b = _group_0_binding_0_vs._matrix[3u][0];
-    int a_1 = _group_0_binding_0_vs.data[(uint(_group_0_binding_0_vs.data.length()) - 2u)].value;
+    int a_2 = _group_0_binding_0_vs.data[(uint(_group_0_binding_0_vs.data.length()) - 2u)].value;
     ivec2 c = _group_0_binding_2_vs;
     float _e33 = read_from_private(foo);
-    c2_ = int[5](a_1, int(b), 3, 4, 5);
+    c2_ = int[5](a_2, int(b), 3, 4, 5);
     c2_[(vi + 1u)] = 42;
     int value = c2_[vi];
     float _e47 = test_arr_as_arg(float[5][10](float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0)));
diff --git a/naga/tests/out/hlsl/access.hlsl b/naga/tests/out/hlsl/access.hlsl
index 142083be68..543b841967 100644
--- a/naga/tests/out/hlsl/access.hlsl
+++ b/naga/tests/out/hlsl/access.hlsl
@@ -201,9 +201,9 @@ void test_matrix_within_array_within_struct_accesses()
     return;
 }
 
-float read_from_private(inout float foo_1)
+float read_from_private(inout float foo_2)
 {
-    float _e1 = foo_1;
+    float _e1 = foo_2;
     return _e1;
 }
 
@@ -224,12 +224,17 @@ ret_Constructarray2_float4_ Constructarray2_float4_(float4 arg0, float4 arg1) {
     return ret;
 }
 
-void assign_array_through_ptr_fn(inout float4 foo_2[2])
+void assign_array_through_ptr_fn(inout float4 foo_3[2])
 {
-    foo_2 = Constructarray2_float4_((1.0).xxxx, (2.0).xxxx);
+    foo_3 = Constructarray2_float4_((1.0).xxxx, (2.0).xxxx);
     return;
 }
 
+int array_by_value(int a_1[5], int i)
+{
+    return a_1[i];
+}
+
 typedef int ret_Constructarray5_int_[5];
 ret_Constructarray5_int_ Constructarray5_int_(int arg0, int arg1, int arg2, int arg3, int arg4) {
     int ret[5] = { arg0, arg1, arg2, arg3, arg4 };
@@ -266,10 +271,10 @@ float4 foo_vert(uint vi : SV_VertexID) : SV_Position
     float4x3 _matrix = float4x3(asfloat(bar.Load3(0+0)), asfloat(bar.Load3(0+16)), asfloat(bar.Load3(0+32)), asfloat(bar.Load3(0+48)));
     uint2 arr_1[2] = Constructarray2_uint2_(asuint(bar.Load2(144+0)), asuint(bar.Load2(144+8)));
     float b = asfloat(bar.Load(0+3u*16+0));
-    int a_1 = asint(bar.Load(0+(((NagaBufferLengthRW(bar) - 160) / 8) - 2u)*8+160));
+    int a_2 = asint(bar.Load(0+(((NagaBufferLengthRW(bar) - 160) / 8) - 2u)*8+160));
     int2 c = asint(qux.Load2(0));
     const float _e33 = read_from_private(foo);
-    c2_ = Constructarray5_int_(a_1, int(b), 3, 4, 5);
+    c2_ = Constructarray5_int_(a_2, int(b), 3, 4, 5);
     c2_[(vi + 1u)] = 42;
     int value = c2_[vi];
     const float _e47 = test_arr_as_arg(ZeroValuearray5_array10_float__());
@@ -310,3 +315,10 @@ void assign_through_ptr()
     assign_array_through_ptr_fn(arr);
     return;
 }
+
+float4 foo_1(uint vi_1 : SV_VertexID) : SV_Position
+{
+    int arr_2[5] = Constructarray5_int_(1, 2, 3, 4, 5);
+    int value_1 = arr_2[vi_1];
+    return float4((value_1).xxxx);
+}
diff --git a/naga/tests/out/hlsl/access.ron b/naga/tests/out/hlsl/access.ron
index 73c9e44448..8960a612ed 100644
--- a/naga/tests/out/hlsl/access.ron
+++ b/naga/tests/out/hlsl/access.ron
@@ -4,6 +4,10 @@
             entry_point:"foo_vert",
             target_profile:"vs_5_1",
         ),
+        (
+            entry_point:"foo_1",
+            target_profile:"vs_5_1",
+        ),
     ],
     fragment:[
         (
diff --git a/naga/tests/out/ir/access.compact.ron b/naga/tests/out/ir/access.compact.ron
index 1b95742ff2..2d066b8ffa 100644
--- a/naga/tests/out/ir/access.compact.ron
+++ b/naga/tests/out/ir/access.compact.ron
@@ -1655,6 +1655,47 @@
                 ),
             ],
         ),
+        (
+            name: Some("array_by_value"),
+            arguments: [
+                (
+                    name: Some("a"),
+                    ty: 25,
+                    binding: None,
+                ),
+                (
+                    name: Some("i"),
+                    ty: 2,
+                    binding: None,
+                ),
+            ],
+            result: Some((
+                ty: 2,
+                binding: None,
+            )),
+            local_variables: [],
+            expressions: [
+                FunctionArgument(0),
+                FunctionArgument(1),
+                Access(
+                    base: 0,
+                    index: 1,
+                ),
+            ],
+            named_expressions: {
+                0: "a",
+                1: "i",
+            },
+            body: [
+                Emit((
+                    start: 2,
+                    end: 3,
+                )),
+                Return(
+                    value: Some(2),
+                ),
+            ],
+        ),
     ],
     entry_points: [
         (
@@ -2230,5 +2271,81 @@
                 ],
             ),
         ),
+        (
+            name: "foo",
+            stage: Vertex,
+            early_depth_test: None,
+            workgroup_size: (0, 0, 0),
+            function: (
+                name: Some("foo"),
+                arguments: [
+                    (
+                        name: Some("vi"),
+                        ty: 0,
+                        binding: Some(BuiltIn(VertexIndex)),
+                    ),
+                ],
+                result: Some((
+                    ty: 24,
+                    binding: Some(BuiltIn(Position(
+                        invariant: false,
+                    ))),
+                )),
+                local_variables: [],
+                expressions: [
+                    FunctionArgument(0),
+                    Literal(I32(1)),
+                    Literal(I32(2)),
+                    Literal(I32(3)),
+                    Literal(I32(4)),
+                    Literal(I32(5)),
+                    Compose(
+                        ty: 25,
+                        components: [
+                            1,
+                            2,
+                            3,
+                            4,
+                            5,
+                        ],
+                    ),
+                    Access(
+                        base: 6,
+                        index: 0,
+                    ),
+                    Splat(
+                        size: Quad,
+                        value: 7,
+                    ),
+                    As(
+                        expr: 8,
+                        kind: Float,
+                        convert: Some(4),
+                    ),
+                ],
+                named_expressions: {
+                    0: "vi",
+                    6: "arr",
+                    7: "value",
+                },
+                body: [
+                    Emit((
+                        start: 6,
+                        end: 7,
+                    )),
+                    Emit((
+                        start: 7,
+                        end: 8,
+                    )),
+                    Emit((
+                        start: 8,
+                        end: 10,
+                    )),
+                    Return(
+                        value: Some(9),
+                    ),
+                ],
+            ),
+        ),
     ],
 )
\ No newline at end of file
diff --git a/naga/tests/out/ir/access.ron b/naga/tests/out/ir/access.ron
index 1b95742ff2..2d066b8ffa 100644
--- a/naga/tests/out/ir/access.ron
+++ b/naga/tests/out/ir/access.ron
@@ -1655,6 +1655,47 @@
                 ),
             ],
         ),
+        (
+            name: Some("array_by_value"),
+            arguments: [
+                (
+                    name: Some("a"),
+                    ty: 25,
+                    binding: None,
+                ),
+                (
+                    name: Some("i"),
+                    ty: 2,
+                    binding: None,
+                ),
+            ],
+            result: Some((
+                ty: 2,
+                binding: None,
+            )),
+            local_variables: [],
+            expressions: [
+                FunctionArgument(0),
+                FunctionArgument(1),
+                Access(
+                    base: 0,
+                    index: 1,
+                ),
+            ],
+            named_expressions: {
+                0: "a",
+                1: "i",
+            },
+            body: [
+                Emit((
+                    start: 2,
+                    end: 3,
+                )),
+                Return(
+                    value: Some(2),
+                ),
+            ],
+        ),
     ],
     entry_points: [
         (
@@ -2230,5 +2271,81 @@
                 ],
             ),
         ),
+        (
+            name: "foo",
+            stage: Vertex,
+            early_depth_test: None,
+            workgroup_size: (0, 0, 0),
+            function: (
+                name: Some("foo"),
+                arguments: [
+                    (
+                        name: Some("vi"),
+                        ty: 0,
+                        binding: Some(BuiltIn(VertexIndex)),
+                    ),
+                ],
+                result: Some((
+                    ty: 24,
+                    binding: Some(BuiltIn(Position(
+                        invariant: false,
+                    ))),
+                )),
+                local_variables: [],
+                expressions: [
+                    FunctionArgument(0),
+                    Literal(I32(1)),
+                    Literal(I32(2)),
+                    Literal(I32(3)),
+                    Literal(I32(4)),
+                    Literal(I32(5)),
+                    Compose(
+                        ty: 25,
+                        components: [
+                            1,
+                            2,
+                            3,
+                            4,
+                            5,
+                        ],
+                    ),
+                    Access(
+                        base: 6,
+                        index: 0,
+                    ),
+                    Splat(
+                        size: Quad,
+                        value: 7,
+                    ),
+                    As(
+                        expr: 8,
+                        kind: Float,
+                        convert: Some(4),
+                    ),
+                ],
+                named_expressions: {
+                    0: "vi",
+                    6: "arr",
+                    7: "value",
+                },
+                body: [
+                    Emit((
+                        start: 6,
+                        end: 7,
+                    )),
+                    Emit((
+                        start: 7,
+                        end: 8,
+                    )),
+                    Emit((
+                        start: 8,
+                        end: 10,
+                    )),
+                    Return(
+                        value: Some(9),
+                    ),
+                ],
+            ),
+        ),
     ],
 )
\ No newline at end of file
diff --git a/naga/tests/out/msl/access.msl b/naga/tests/out/msl/access.msl
index 65dba4910e..924b604e4c 100644
--- a/naga/tests/out/msl/access.msl
+++ b/naga/tests/out/msl/access.msl
@@ -133,9 +133,9 @@ void test_matrix_within_array_within_struct_accesses(
 }
 
 float read_from_private(
-    thread float& foo_1
+    thread float& foo_2
 ) {
-    float _e1 = foo_1;
+    float _e1 = foo_2;
     return _e1;
 }
 
@@ -153,12 +153,19 @@ void assign_through_ptr_fn(
 }
 
 void assign_array_through_ptr_fn(
-    thread type_22& foo_2
+    thread type_22& foo_3
 ) {
-    foo_2 = type_22 {metal::float4(1.0), metal::float4(2.0)};
+    foo_3 = type_22 {metal::float4(1.0), metal::float4(2.0)};
     return;
 }
 
+int array_by_value(
+    type_20 a_1,
+    int i
+) {
+    return a_1.inner[i];
+}
+
 struct foo_vertInput {
 };
 struct foo_vertOutput {
@@ -181,10 +188,10 @@ vertex foo_vertOutput foo_vert(
     metal::float4x3 _matrix = bar._matrix;
     type_10 arr_1 = bar.arr;
     float b = bar._matrix[3u].x;
-    int a_1 = bar.data[(1 + (_buffer_sizes.size1 - 160 - 8) / 8) - 2u].value;
+    int a_2 = bar.data[(1 + (_buffer_sizes.size1 - 160 - 8) / 8) - 2u].value;
     metal::int2 c = qux;
     float _e33 = read_from_private(foo);
-    c2_ = type_20 {a_1, static_cast<int>(b), 3, 4, 5};
+    c2_ = type_20 {a_2, static_cast<int>(b), 3, 4, 5};
     c2_.inner[vi + 1u] = 42;
     int value = c2_.inner[vi];
     float _e47 = test_arr_as_arg(type_18 {});
@@ -217,3 +224,17 @@ kernel void assign_through_ptr(
     assign_array_through_ptr_fn(arr);
     return;
 }
+
+
+struct foo_1Input {
+};
+struct foo_1Output {
+    metal::float4 member_3 [[position]];
+};
+vertex foo_1Output foo_1(
+  uint vi_1 [[vertex_id]]
+) {
+    type_20 arr_2 = type_20 {1, 2, 3, 4, 5};
+    int value_1 = arr_2.inner[vi_1];
+    return foo_1Output { static_cast<metal::float4>(metal::int4(value_1)) };
+}
diff --git a/naga/tests/out/spv/6220-break-from-loop.spvasm b/naga/tests/out/spv/6220-break-from-loop.spvasm
new file mode 100644
index 0000000000..9dabbc2079
--- /dev/null
+++ b/naga/tests/out/spv/6220-break-from-loop.spvasm
@@ -0,0 +1,44 @@
+; SPIR-V
+; Version: 1.1
+; Generator: rspirv
+; Bound: 26
+OpCapability Shader
+OpCapability Linkage
+%1 = OpExtInstImport "GLSL.std.450"
+OpMemoryModel Logical GLSL450
+%2 = OpTypeVoid
+%3 = OpTypeInt 32 1
+%6 = OpTypeFunction %2
+%7 = OpConstant  %3  0
+%8 = OpConstant  %3  4
+%9 = OpConstant  %3  1
+%11 = OpTypePointer Function %3
+%18 = OpTypeBool
+%5 = OpFunction  %2  None %6
+%4 = OpLabel
+%10 = OpVariable  %11  Function %7
+OpBranch %12
+%12 = OpLabel
+OpBranch %13
+%13 = OpLabel
+OpLoopMerge %14 %16 None
+OpBranch %15
+%15 = OpLabel
+%17 = OpLoad  %3  %10
+%19 = OpSLessThan  %18  %17 %8
+OpSelectionMerge %20 None
+OpBranchConditional %19 %20 %21
+%21 = OpLabel
+OpBranch %14
+%20 = OpLabel
+OpBranch %22
+%22 = OpLabel
+OpBranch %14
+%16 = OpLabel
+%24 = OpLoad  %3  %10
+%25 = OpIAdd  %3  %24 %9
+OpStore %10 %25
+OpBranch %13
+%14 = OpLabel
+OpReturn
+OpFunctionEnd
\ No newline at end of file
diff --git a/naga/tests/out/spv/access.spvasm b/naga/tests/out/spv/access.spvasm
index ab0112870f..6cb87a3b10 100644
--- a/naga/tests/out/spv/access.spvasm
+++ b/naga/tests/out/spv/access.spvasm
@@ -1,16 +1,17 @@
 ; SPIR-V
 ; Version: 1.1
 ; Generator: rspirv
-; Bound: 301
+; Bound: 323
 OpCapability Shader
 OpExtension "SPV_KHR_storage_buffer_storage_class"
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
-OpEntryPoint Vertex %219 "foo_vert" %214 %217
-OpEntryPoint Fragment %273 "foo_frag" %272
-OpEntryPoint GLCompute %291 "assign_through_ptr"
-OpExecutionMode %273 OriginUpperLeft
-OpExecutionMode %291 LocalSize 1 1 1
+OpEntryPoint Vertex %229 "foo_vert" %224 %227
+OpEntryPoint Fragment %282 "foo_frag" %281
+OpEntryPoint GLCompute %300 "assign_through_ptr"
+OpEntryPoint Vertex %314 "foo" %311 %313
+OpExecutionMode %282 OriginUpperLeft
+OpExecutionMode %300 LocalSize 1 1 1
 OpMemberName %6 0 "a"
 OpMemberName %6 1 "b"
 OpMemberName %6 2 "c"
@@ -47,14 +48,19 @@ OpName %200 "p"
 OpName %201 "assign_through_ptr_fn"
 OpName %206 "foo"
 OpName %207 "assign_array_through_ptr_fn"
-OpName %214 "vi"
-OpName %219 "foo_vert"
-OpName %231 "foo"
-OpName %232 "c2"
-OpName %273 "foo_frag"
-OpName %291 "assign_through_ptr"
-OpName %296 "val"
-OpName %297 "arr"
+OpName %214 "a"
+OpName %215 "i"
+OpName %216 "array_by_value"
+OpName %224 "vi"
+OpName %229 "foo_vert"
+OpName %241 "foo"
+OpName %242 "c2"
+OpName %282 "foo_frag"
+OpName %300 "assign_through_ptr"
+OpName %305 "val"
+OpName %306 "arr"
+OpName %311 "vi"
+OpName %314 "foo"
 OpMemberDecorate %6 0 Offset 0
 OpMemberDecorate %6 1 Offset 16
 OpMemberDecorate %6 2 Offset 28
@@ -99,9 +105,11 @@ OpDecorate %50 DescriptorSet 0
 OpDecorate %50 Binding 3
 OpDecorate %51 Block
 OpMemberDecorate %51 0 Offset 0
-OpDecorate %214 BuiltIn VertexIndex
-OpDecorate %217 BuiltIn Position
-OpDecorate %272 Location 0
+OpDecorate %224 BuiltIn VertexIndex
+OpDecorate %227 BuiltIn Position
+OpDecorate %281 Location 0
+OpDecorate %311 BuiltIn VertexIndex
+OpDecorate %313 BuiltIn Position
 %2 = OpTypeVoid
 %3 = OpTypeInt 32 0
 %4 = OpTypeVector %3 3
@@ -209,44 +217,49 @@ OpDecorate %272 Location 0
 %209 = OpConstantComposite  %31  %59 %59 %59 %59
 %210 = OpConstantComposite  %31  %61 %61 %61 %61
 %211 = OpConstantComposite  %34  %209 %210
-%215 = OpTypePointer Input %3
-%214 = OpVariable  %215  Input
-%218 = OpTypePointer Output %31
-%217 = OpVariable  %218  Output
-%221 = OpTypePointer StorageBuffer %23
-%224 = OpConstant  %8  0.0
-%225 = OpConstant  %3  3
-%226 = OpConstant  %5  3
-%227 = OpConstant  %5  4
-%228 = OpConstant  %5  5
-%229 = OpConstant  %5  42
-%230 = OpConstantNull  %29
-%233 = OpTypePointer Function %32
-%234 = OpConstantNull  %32
-%239 = OpTypePointer StorageBuffer %9
-%242 = OpTypePointer StorageBuffer %18
-%243 = OpConstant  %3  4
-%246 = OpTypePointer StorageBuffer %10
-%247 = OpTypePointer StorageBuffer %8
-%250 = OpTypePointer StorageBuffer %19
-%253 = OpTypePointer StorageBuffer %7
-%254 = OpTypePointer StorageBuffer %5
-%266 = OpTypeVector %5 4
-%272 = OpVariable  %218  Output
-%275 = OpConstantComposite  %10  %224 %224 %224
-%276 = OpConstantComposite  %10  %59 %59 %59
-%277 = OpConstantComposite  %10  %61 %61 %61
-%278 = OpConstantComposite  %10  %63 %63 %63
-%279 = OpConstantComposite  %9  %275 %276 %277 %278
-%280 = OpConstantComposite  %17  %36 %36
-%281 = OpConstantComposite  %17  %99 %99
-%282 = OpConstantComposite  %18  %280 %281
-%283 = OpConstantNull  %23
-%284 = OpConstantComposite  %31  %224 %224 %224 %224
-%292 = OpConstant  %3  33
-%293 = OpConstantComposite  %31  %67 %67 %67 %67
-%294 = OpConstantComposite  %31  %137 %137 %137 %137
-%295 = OpConstantComposite  %34  %293 %294
+%217 = OpTypeFunction %5 %32 %5
+%219 = OpTypePointer Function %32
+%225 = OpTypePointer Input %3
+%224 = OpVariable  %225  Input
+%228 = OpTypePointer Output %31
+%227 = OpVariable  %228  Output
+%231 = OpTypePointer StorageBuffer %23
+%234 = OpConstant  %8  0.0
+%235 = OpConstant  %3  3
+%236 = OpConstant  %5  3
+%237 = OpConstant  %5  4
+%238 = OpConstant  %5  5
+%239 = OpConstant  %5  42
+%240 = OpConstantNull  %29
+%243 = OpConstantNull  %32
+%248 = OpTypePointer StorageBuffer %9
+%251 = OpTypePointer StorageBuffer %18
+%252 = OpConstant  %3  4
+%255 = OpTypePointer StorageBuffer %10
+%256 = OpTypePointer StorageBuffer %8
+%259 = OpTypePointer StorageBuffer %19
+%262 = OpTypePointer StorageBuffer %7
+%263 = OpTypePointer StorageBuffer %5
+%275 = OpTypeVector %5 4
+%281 = OpVariable  %228  Output
+%284 = OpConstantComposite  %10  %234 %234 %234
+%285 = OpConstantComposite  %10  %59 %59 %59
+%286 = OpConstantComposite  %10  %61 %61 %61
+%287 = OpConstantComposite  %10  %63 %63 %63
+%288 = OpConstantComposite  %9  %284 %285 %286 %287
+%289 = OpConstantComposite  %17  %36 %36
+%290 = OpConstantComposite  %17  %99 %99
+%291 = OpConstantComposite  %18  %289 %290
+%292 = OpConstantNull  %23
+%293 = OpConstantComposite  %31  %234 %234 %234 %234
+%301 = OpConstant  %3  33
+%302 = OpConstantComposite  %31  %67 %67 %67 %67
+%303 = OpConstantComposite  %31  %137 %137 %137 %137
+%304 = OpConstantComposite  %34  %302 %303
+%311 = OpVariable  %225  Input
+%313 = OpVariable  %228  Output
+%315 = OpConstant  %5  2
+%316 = OpConstantComposite  %32  %58 %315 %236 %237 %238
 %54 = OpFunction  %2  None %55
 %53 = OpLabel
 %82 = OpVariable  %83  Function %58
@@ -390,72 +403,98 @@ OpBranch %212
 OpStore %206 %211
 OpReturn
 OpFunctionEnd
-%219 = OpFunction  %2  None %55
+%216 = OpFunction  %5  None %217
+%214 = OpFunctionParameter  %32
+%215 = OpFunctionParameter  %5
 %213 = OpLabel
-%231 = OpVariable  %27  Function %224
-%232 = OpVariable  %233  Function %234
-%216 = OpLoad  %3  %214
-%220 = OpAccessChain  %56  %44 %36
-%222 = OpAccessChain  %221  %47 %36
-%223 = OpAccessChain  %131  %50 %36
-OpBranch %235
-%235 = OpLabel
-%236 = OpLoad  %8  %231
-OpStore %231 %59
-%237 = OpFunctionCall  %2  %54
-%238 = OpFunctionCall  %2  %130
-%240 = OpAccessChain  %239  %42 %36
-%241 = OpLoad  %9  %240
-%244 = OpAccessChain  %242  %42 %243
-%245 = OpLoad  %18  %244
-%248 = OpAccessChain  %247  %42 %36 %225 %36
-%249 = OpLoad  %8  %248
-%251 = OpArrayLength  %3  %42 5
-%252 = OpISub  %3  %251 %14
-%255 = OpAccessChain  %254  %42 %30 %252 %36
-%256 = OpLoad  %5  %255
-%257 = OpLoad  %23  %222
-%258 = OpFunctionCall  %8  %188 %231
-%259 = OpConvertFToS  %5  %249
-%260 = OpCompositeConstruct  %32  %256 %259 %226 %227 %228
-OpStore %232 %260
-%261 = OpIAdd  %3  %216 %99
-%262 = OpAccessChain  %83  %232 %261
-OpStore %262 %229
-%263 = OpAccessChain  %83  %232 %216
-%264 = OpLoad  %5  %263
-%265 = OpFunctionCall  %8  %194 %230
-%267 = OpCompositeConstruct  %266  %264 %264 %264 %264
-%268 = OpConvertSToF  %31  %267
-%269 = OpMatrixTimesVector  %10  %241 %268
-%270 = OpCompositeConstruct  %31  %269 %61
-OpStore %217 %270
+%220 = OpVariable  %219  Function
+OpBranch %218
+%218 = OpLabel
+OpStore %220 %214
+%221 = OpAccessChain  %83  %220 %215
+%222 = OpLoad  %5  %221
+OpReturnValue %222
+OpFunctionEnd
+%229 = OpFunction  %2  None %55
+%223 = OpLabel
+%241 = OpVariable  %27  Function %234
+%242 = OpVariable  %219  Function %243
+%226 = OpLoad  %3  %224
+%230 = OpAccessChain  %56  %44 %36
+%232 = OpAccessChain  %231  %47 %36
+%233 = OpAccessChain  %131  %50 %36
+OpBranch %244
+%244 = OpLabel
+%245 = OpLoad  %8  %241
+OpStore %241 %59
+%246 = OpFunctionCall  %2  %54
+%247 = OpFunctionCall  %2  %130
+%249 = OpAccessChain  %248  %42 %36
+%250 = OpLoad  %9  %249
+%253 = OpAccessChain  %251  %42 %252
+%254 = OpLoad  %18  %253
+%257 = OpAccessChain  %256  %42 %36 %235 %36
+%258 = OpLoad  %8  %257
+%260 = OpArrayLength  %3  %42 5
+%261 = OpISub  %3  %260 %14
+%264 = OpAccessChain  %263  %42 %30 %261 %36
+%265 = OpLoad  %5  %264
+%266 = OpLoad  %23  %232
+%267 = OpFunctionCall  %8  %188 %241
+%268 = OpConvertFToS  %5  %258
+%269 = OpCompositeConstruct  %32  %265 %268 %236 %237 %238
+OpStore %242 %269
+%270 = OpIAdd  %3  %226 %99
+%271 = OpAccessChain  %83  %242 %270
+OpStore %271 %239
+%272 = OpAccessChain  %83  %242 %226
+%273 = OpLoad  %5  %272
+%274 = OpFunctionCall  %8  %194 %240
+%276 = OpCompositeConstruct  %275  %273 %273 %273 %273
+%277 = OpConvertSToF  %31  %276
+%278 = OpMatrixTimesVector  %10  %250 %277
+%279 = OpCompositeConstruct  %31  %278 %61
+OpStore %227 %279
+OpReturn
+OpFunctionEnd
+%282 = OpFunction  %2  None %55
+%280 = OpLabel
+%283 = OpAccessChain  %231  %47 %36
+OpBranch %294
+%294 = OpLabel
+%295 = OpAccessChain  %256  %42 %36 %99 %14
+OpStore %295 %59
+%296 = OpAccessChain  %248  %42 %36
+OpStore %296 %288
+%297 = OpAccessChain  %251  %42 %252
+OpStore %297 %291
+%298 = OpAccessChain  %263  %42 %30 %99 %36
+OpStore %298 %58
+OpStore %283 %292
+OpStore %281 %293
 OpReturn
 OpFunctionEnd
-%273 = OpFunction  %2  None %55
-%271 = OpLabel
-%274 = OpAccessChain  %221  %47 %36
-OpBranch %285
-%285 = OpLabel
-%286 = OpAccessChain  %247  %42 %36 %99 %14
-OpStore %286 %59
-%287 = OpAccessChain  %239  %42 %36
-OpStore %287 %279
-%288 = OpAccessChain  %242  %42 %243
-OpStore %288 %282
-%289 = OpAccessChain  %254  %42 %30 %99 %36
-OpStore %289 %58
-OpStore %274 %283
-OpStore %272 %284
+%300 = OpFunction  %2  None %55
+%299 = OpLabel
+%305 = OpVariable  %33  Function %301
+%306 = OpVariable  %35  Function %304
+OpBranch %307
+%307 = OpLabel
+%308 = OpFunctionCall  %2  %201 %305
+%309 = OpFunctionCall  %2  %207 %306
 OpReturn
 OpFunctionEnd
-%291 = OpFunction  %2  None %55
-%290 = OpLabel
-%296 = OpVariable  %33  Function %292
-%297 = OpVariable  %35  Function %295
-OpBranch %298
-%298 = OpLabel
-%299 = OpFunctionCall  %2  %201 %296
-%300 = OpFunctionCall  %2  %207 %297
+%314 = OpFunction  %2  None %55
+%310 = OpLabel
+%318 = OpVariable  %219  Function
+%312 = OpLoad  %3  %311
+OpBranch %317
+%317 = OpLabel
+OpStore %318 %316
+%319 = OpAccessChain  %83  %318 %312
+%320 = OpLoad  %5  %319
+%321 = OpCompositeConstruct  %275  %320 %320 %320 %320
+%322 = OpConvertSToF  %31  %321
+OpStore %313 %322
 OpReturn
 OpFunctionEnd
\ No newline at end of file
diff --git a/naga/tests/out/wgsl/access.wgsl b/naga/tests/out/wgsl/access.wgsl
index 1409e80b11..0c29e5d605 100644
--- a/naga/tests/out/wgsl/access.wgsl
+++ b/naga/tests/out/wgsl/access.wgsl
@@ -107,8 +107,8 @@ fn test_matrix_within_array_within_struct_accesses() {
     return;
 }
 
-fn read_from_private(foo_1: ptr<function, f32>) -> f32 {
-    let _e1 = (*foo_1);
+fn read_from_private(foo_2: ptr<function, f32>) -> f32 {
+    let _e1 = (*foo_2);
     return _e1;
 }
 
@@ -121,11 +121,15 @@ fn assign_through_ptr_fn(p: ptr<function, u32>) {
     return;
 }
 
-fn assign_array_through_ptr_fn(foo_2: ptr<function, array<vec4<f32>, 2>>) {
-    (*foo_2) = array<vec4<f32>, 2>(vec4(1f), vec4(2f));
+fn assign_array_through_ptr_fn(foo_3: ptr<function, array<vec4<f32>, 2>>) {
+    (*foo_3) = array<vec4<f32>, 2>(vec4(1f), vec4(2f));
     return;
 }
 
+fn array_by_value(a_1: array<i32, 5>, i: i32) -> i32 {
+    return a_1[i];
+}
+
 @vertex 
 fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
     var foo: f32 = 0f;
@@ -138,11 +142,11 @@ fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
     let _matrix = bar._matrix;
     let arr_1 = bar.arr;
     let b = bar._matrix[3u][0];
-    let a_1 = bar.data[(arrayLength((&bar.data)) - 2u)].value;
+    let a_2 = bar.data[(arrayLength((&bar.data)) - 2u)].value;
     let c = qux;
     let data_pointer = (&bar.data[0].value);
     let _e33 = read_from_private((&foo));
-    c2_ = array<i32, 5>(a_1, i32(b), 3i, 4i, 5i);
+    c2_ = array<i32, 5>(a_2, i32(b), 3i, 4i, 5i);
     c2_[(vi + 1u)] = 42i;
     let value = c2_[vi];
     let _e47 = test_arr_as_arg(array<array<f32, 10>, 5>());
@@ -168,3 +172,10 @@ fn assign_through_ptr() {
     assign_array_through_ptr_fn((&arr));
     return;
 }
+
+@vertex 
+fn foo_1(@builtin(vertex_index) vi_1: u32) -> @builtin(position) vec4<f32> {
+    const arr_2 = array<i32, 5>(1i, 2i, 3i, 4i, 5i);
+    let value_1 = arr_2[vi_1];
+    return vec4<f32>(vec4(value_1));
+}
diff --git a/naga/tests/snapshots.rs b/naga/tests/snapshots.rs
index 596e4cea14..adf67f8333 100644
--- a/naga/tests/snapshots.rs
+++ b/naga/tests/snapshots.rs
@@ -932,6 +932,7 @@ fn convert_wgsl() {
             "phony_assignment",
             Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL,
         ),
+        ("6220-break-from-loop", Targets::SPIRV),
     ];
 
     for &(name, targets) in inputs.iter() {
diff --git a/naga/tests/wgsl_errors.rs b/naga/tests/wgsl_errors.rs
index 2d91ba01cf..e5fb77644d 100644
--- a/naga/tests/wgsl_errors.rs
+++ b/naga/tests/wgsl_errors.rs
@@ -1359,11 +1359,6 @@ fn missing_bindings2() {
 #[test]
 fn invalid_access() {
     check_validation! {
-        "
-        fn array_by_value(a: array<i32, 5>, i: i32) -> i32 {
-            return a[i];
-        }
-        ",
         "
         fn matrix_by_value(m: mat4x4<f32>, i: i32) -> vec4<f32> {
             return m[i];
diff --git a/wgpu-core/src/binding_model.rs b/wgpu-core/src/binding_model.rs
index 7e5e58e973..c58022e73f 100644
--- a/wgpu-core/src/binding_model.rs
+++ b/wgpu-core/src/binding_model.rs
@@ -377,10 +377,6 @@ impl BindingTypeMaxCountValidator {
             limits.max_sampled_textures_per_shader_stage,
             BindingTypeMaxCountErrorKind::SampledTextures,
         )?;
-        self.storage_buffers.validate(
-            limits.max_storage_buffers_per_shader_stage,
-            BindingTypeMaxCountErrorKind::StorageBuffers,
-        )?;
         self.samplers.validate(
             limits.max_samplers_per_shader_stage,
             BindingTypeMaxCountErrorKind::Samplers,
diff --git a/wgpu-core/src/command/compute.rs b/wgpu-core/src/command/compute.rs
index fe5f68c12f..e147ee0d5b 100644
--- a/wgpu-core/src/command/compute.rs
+++ b/wgpu-core/src/command/compute.rs
@@ -136,6 +136,8 @@ pub enum ComputePassErrorInner {
     BindGroupIndexOutOfRange { index: u32, max: u32 },
     #[error(transparent)]
     DestroyedResource(#[from] DestroyedResourceError),
+    #[error("Indirect buffer offset {0:?} is not a multiple of 4")]
+    UnalignedIndirectBufferOffset(BufferAddress),
     #[error("Indirect buffer uses bytes {offset}..{end_offset} which overruns indirect buffer of size {buffer_size}")]
     IndirectBufferOverrun {
         offset: u64,
@@ -497,8 +499,6 @@ impl Global {
             state.raw_encoder.begin_compute_pass(&hal_desc);
         }
 
-        // TODO: We should be draining the commands here, avoiding extra copies in the process.
-        //       (A command encoder can't be executed twice!)
         for command in base.commands {
             match command {
                 ArcComputeCommand::SetBindGroup {
@@ -793,7 +793,7 @@ fn set_push_constant(
         .binder
         .pipeline_layout
         .as_ref()
-        //TODO: don't error here, lazily update the push constants
+        // TODO: don't error here, lazily update the push constants using `state.push_constants`
         .ok_or(ComputePassErrorInner::Dispatch(
             DispatchError::MissingPipeline,
         ))?;
@@ -860,6 +860,10 @@ fn dispatch_indirect(
         .merge_single(&buffer, hal::BufferUses::INDIRECT)?;
     buffer.check_usage(wgt::BufferUsages::INDIRECT)?;
 
+    if offset % 4 != 0 {
+        return Err(ComputePassErrorInner::UnalignedIndirectBufferOffset(offset));
+    }
+
     let end_offset = offset + size_of::<wgt::DispatchIndirectArgs>() as u64;
     if end_offset > buffer.size {
         return Err(ComputePassErrorInner::IndirectBufferOverrun {
diff --git a/wgpu-core/src/command/render.rs b/wgpu-core/src/command/render.rs
index b6680333c2..d22eb5f0d6 100644
--- a/wgpu-core/src/command/render.rs
+++ b/wgpu-core/src/command/render.rs
@@ -638,6 +638,8 @@ pub enum RenderPassErrorInner {
     MissingFeatures(#[from] MissingFeatures),
     #[error(transparent)]
     MissingDownlevelFlags(#[from] MissingDownlevelFlags),
+    #[error("Indirect buffer offset {0:?} is not a multiple of 4")]
+    UnalignedIndirectBufferOffset(BufferAddress),
     #[error("Indirect draw uses bytes {offset}..{end_offset} {} which overruns indirect buffer of size {buffer_size}",
         count.map_or_else(String::new, |v| format!("(using count {v})")))]
     IndirectBufferOverrun {
@@ -2450,6 +2452,10 @@ fn multi_draw_indirect(
 
     let actual_count = count.map_or(1, |c| c.get());
 
+    if offset % 4 != 0 {
+        return Err(RenderPassErrorInner::UnalignedIndirectBufferOffset(offset));
+    }
+
     let end_offset = offset + stride as u64 * actual_count as u64;
     if end_offset > indirect_buffer.size {
         return Err(RenderPassErrorInner::IndirectBufferOverrun {
@@ -2534,6 +2540,10 @@ fn multi_draw_indirect_count(
     count_buffer.check_usage(BufferUsages::INDIRECT)?;
     let count_raw = count_buffer.try_raw(state.snatch_guard)?;
 
+    if offset % 4 != 0 {
+        return Err(RenderPassErrorInner::UnalignedIndirectBufferOffset(offset));
+    }
+
     let end_offset = offset + stride * max_count as u64;
     if end_offset > indirect_buffer.size {
         return Err(RenderPassErrorInner::IndirectBufferOverrun {
diff --git a/wgpu-core/src/device/resource.rs b/wgpu-core/src/device/resource.rs
index bc01e5aa55..7ce67a693e 100644
--- a/wgpu-core/src/device/resource.rs
+++ b/wgpu-core/src/device/resource.rs
@@ -38,7 +38,6 @@ use arrayvec::ArrayVec;
 use once_cell::sync::OnceCell;
 
 use smallvec::SmallVec;
-use thiserror::Error;
 use wgt::{
     math::align_to, DeviceLostReason, TextureFormat, TextureSampleType, TextureViewDimension,
 };
@@ -189,14 +188,6 @@ impl Drop for Device {
     }
 }
 
-#[derive(Clone, Debug, Error)]
-pub enum CreateDeviceError {
-    #[error("Not enough memory left to create device")]
-    OutOfMemory,
-    #[error("Failed to create internal buffer for initializing textures")]
-    FailedToCreateZeroBuffer(#[from] DeviceError),
-}
-
 impl Device {
     pub(crate) fn raw(&self) -> &dyn hal::DynDevice {
         self.raw.as_ref()
@@ -376,7 +367,7 @@ impl Device {
                     let Some(view) = view.upgrade() else {
                         continue;
                     };
-                    let Some(raw_view) = view.raw.snatch(self.snatchable_lock.write()) else {
+                    let Some(raw_view) = view.raw.snatch(&mut self.snatchable_lock.write()) else {
                         continue;
                     };
 
@@ -390,7 +381,8 @@ impl Device {
                     let Some(bind_group) = bind_group.upgrade() else {
                         continue;
                     };
-                    let Some(raw_bind_group) = bind_group.raw.snatch(self.snatchable_lock.write())
+                    let Some(raw_bind_group) =
+                        bind_group.raw.snatch(&mut self.snatchable_lock.write())
                     else {
                         continue;
                     };
@@ -441,13 +433,11 @@ impl Device {
                     .last_successful_submission_index
                     .load(Ordering::Acquire);
 
-                if let wgt::Maintain::WaitForSubmissionIndex(submission_index) = maintain {
-                    if submission_index > last_successful_submission_index {
-                        return Err(WaitIdleError::WrongSubmissionIndex(
-                            submission_index,
-                            last_successful_submission_index,
-                        ));
-                    }
+                if submission_index > last_successful_submission_index {
+                    return Err(WaitIdleError::WrongSubmissionIndex(
+                        submission_index,
+                        last_successful_submission_index,
+                    ));
                 }
 
                 submission_index
diff --git a/wgpu-core/src/present.rs b/wgpu-core/src/present.rs
index c9d0124bf4..b12e35328a 100644
--- a/wgpu-core/src/present.rs
+++ b/wgpu-core/src/present.rs
@@ -288,8 +288,11 @@ impl Global {
                     .textures
                     .remove(texture.tracker_index());
                 let suf = surface.raw(device.backend()).unwrap();
-                let exclusive_snatch_guard = device.snatchable_lock.write();
-                match texture.inner.snatch(exclusive_snatch_guard).unwrap() {
+                match texture
+                    .inner
+                    .snatch(&mut device.snatchable_lock.write())
+                    .unwrap()
+                {
                     resource::TextureInner::Surface { raw, parent_id } => {
                         if surface_id != parent_id {
                             log::error!("Presented frame is from a different surface");
@@ -359,8 +362,11 @@ impl Global {
                     .textures
                     .remove(texture.tracker_index());
                 let suf = surface.raw(device.backend());
-                let exclusive_snatch_guard = device.snatchable_lock.write();
-                match texture.inner.snatch(exclusive_snatch_guard).unwrap() {
+                match texture
+                    .inner
+                    .snatch(&mut device.snatchable_lock.write())
+                    .unwrap()
+                {
                     resource::TextureInner::Surface { raw, parent_id } => {
                         if surface_id == parent_id {
                             unsafe { suf.unwrap().discard_texture(raw) };
diff --git a/wgpu-core/src/resource.rs b/wgpu-core/src/resource.rs
index 31ba03ae12..0bf8beec28 100644
--- a/wgpu-core/src/resource.rs
+++ b/wgpu-core/src/resource.rs
@@ -738,8 +738,7 @@ impl Buffer {
         let device = &self.device;
 
         let temp = {
-            let snatch_guard = device.snatchable_lock.write();
-            let raw = match self.raw.snatch(snatch_guard) {
+            let raw = match self.raw.snatch(&mut device.snatchable_lock.write()) {
                 Some(raw) => raw,
                 None => {
                     return Err(DestroyError::AlreadyDestroyed);
@@ -1186,8 +1185,7 @@ impl Texture {
         let device = &self.device;
 
         let temp = {
-            let snatch_guard = device.snatchable_lock.write();
-            let raw = match self.inner.snatch(snatch_guard) {
+            let raw = match self.inner.snatch(&mut device.snatchable_lock.write()) {
                 Some(TextureInner::Native { raw }) => raw,
                 Some(TextureInner::Surface { .. }) => {
                     return Ok(());
diff --git a/wgpu-core/src/snatch.rs b/wgpu-core/src/snatch.rs
index 9866b77723..a817e2068c 100644
--- a/wgpu-core/src/snatch.rs
+++ b/wgpu-core/src/snatch.rs
@@ -38,7 +38,7 @@ impl<T> Snatchable<T> {
     }
 
     /// Take the value. Requires a the snatchable lock's write guard.
-    pub fn snatch(&self, _guard: ExclusiveSnatchGuard) -> Option<T> {
+    pub fn snatch(&self, _guard: &mut ExclusiveSnatchGuard) -> Option<T> {
         unsafe { (*self.value.get()).take() }
     }
 
diff --git a/wgpu-hal/Cargo.toml b/wgpu-hal/Cargo.toml
index d3a05f0790..03af10b96e 100644
--- a/wgpu-hal/Cargo.toml
+++ b/wgpu-hal/Cargo.toml
@@ -56,6 +56,7 @@ vulkan = [
 ]
 gles = [
     "naga/glsl-out",
+    "dep:bytemuck",
     "dep:glow",
     "dep:glutin_wgl_sys",
     "dep:khronos-egl",
@@ -126,6 +127,7 @@ rustc-hash.workspace = true
 log.workspace = true
 
 # backend: Gles
+bytemuck = { workspace = true, optional = true }
 glow = { workspace = true, optional = true }
 
 [dependencies.wgt]
diff --git a/wgpu-hal/examples/ray-traced-triangle/main.rs b/wgpu-hal/examples/ray-traced-triangle/main.rs
index dd91843734..4eedfe7817 100644
--- a/wgpu-hal/examples/ray-traced-triangle/main.rs
+++ b/wgpu-hal/examples/ray-traced-triangle/main.rs
@@ -124,7 +124,12 @@ impl AccelerationStructureInstance {
         &mut self,
         shader_binding_table_record_offset: u32,
     ) {
-        debug_assert!(shader_binding_table_record_offset <= Self::MAX_U24, "shader_binding_table_record_offset uses more than 24 bits! {shader_binding_table_record_offset} > {}", Self::MAX_U24);
+        debug_assert!(
+            shader_binding_table_record_offset <= Self::MAX_U24,
+            "shader_binding_table_record_offset uses more than 24 bits! {} > {}",
+            shader_binding_table_record_offset,
+            Self::MAX_U24
+        );
         self.shader_binding_table_record_offset_and_flags = (shader_binding_table_record_offset
             & Self::LOW_24_MASK)
             | (self.shader_binding_table_record_offset_and_flags & !Self::LOW_24_MASK)
@@ -151,7 +156,9 @@ impl AccelerationStructureInstance {
         );
         debug_assert!(
             shader_binding_table_record_offset <= Self::MAX_U24,
-            "shader_binding_table_record_offset uses more than 24 bits! {shader_binding_table_record_offset} > {}", Self::MAX_U24
+            "shader_binding_table_record_offset uses more than 24 bits! {} > {}",
+            shader_binding_table_record_offset,
+            Self::MAX_U24
         );
         AccelerationStructureInstance {
             transform: Self::affine_to_rows(transform),
diff --git a/wgpu-hal/src/auxil/renderdoc.rs b/wgpu-hal/src/auxil/renderdoc.rs
index 3b08955fad..3879bb9545 100644
--- a/wgpu-hal/src/auxil/renderdoc.rs
+++ b/wgpu-hal/src/auxil/renderdoc.rs
@@ -74,7 +74,8 @@ impl RenderDoc {
                 Err(e) => {
                     return RenderDoc::NotAvailable {
                         reason: format!(
-                            "Unable to get RENDERDOC_GetAPI from renderdoc library '{renderdoc_filename}': {e:?}"
+                            "Unable to get RENDERDOC_GetAPI from renderdoc library '{}': {e:?}",
+                            renderdoc_filename
                         ),
                     }
                 }
@@ -89,7 +90,8 @@ impl RenderDoc {
             },
             return_value => RenderDoc::NotAvailable {
                 reason: format!(
-                    "Unable to get API from renderdoc library '{renderdoc_filename}': {return_value}"
+                    "Unable to get API from renderdoc library '{}': {}",
+                    renderdoc_filename, return_value
                 ),
             },
         }
diff --git a/wgpu-hal/src/gles/mod.rs b/wgpu-hal/src/gles/mod.rs
index df59778065..8eb2800895 100644
--- a/wgpu-hal/src/gles/mod.rs
+++ b/wgpu-hal/src/gles/mod.rs
@@ -457,11 +457,24 @@ impl Texture {
         };
 
         log::error!(
-            "wgpu-hal heuristics assumed that the view dimension will be equal to `{got}` rather than `{view_dimension:?}`.\n{}\n{}\n{}\n{}",
-            "`D2` textures with `depth_or_array_layers == 1` are assumed to have view dimension `D2`",
-            "`D2` textures with `depth_or_array_layers > 1` are assumed to have view dimension `D2Array`",
-            "`D2` textures with `depth_or_array_layers == 6` are assumed to have view dimension `Cube`",
-            "`D2` textures with `depth_or_array_layers > 6 && depth_or_array_layers % 6 == 0` are assumed to have view dimension `CubeArray`",
+            concat!(
+                "wgpu-hal heuristics assumed that ",
+                "the view dimension will be equal to `{}` rather than `{:?}`.\n",
+                "`D2` textures with ",
+                "`depth_or_array_layers == 1` ",
+                "are assumed to have view dimension `D2`\n",
+                "`D2` textures with ",
+                "`depth_or_array_layers > 1` ",
+                "are assumed to have view dimension `D2Array`\n",
+                "`D2` textures with ",
+                "`depth_or_array_layers == 6` ",
+                "are assumed to have view dimension `Cube`\n",
+                "`D2` textures with ",
+                "`depth_or_array_layers > 6 && depth_or_array_layers % 6 == 0` ",
+                "are assumed to have view dimension `CubeArray`\n",
+            ),
+            got,
+            view_dimension,
         );
     }
 }
diff --git a/wgpu-hal/src/gles/queue.rs b/wgpu-hal/src/gles/queue.rs
index 39315f72b7..fc106eb23d 100644
--- a/wgpu-hal/src/gles/queue.rs
+++ b/wgpu-hal/src/gles/queue.rs
@@ -1603,19 +1603,13 @@ impl super::Queue {
                 ref uniform,
                 offset,
             } => {
-                // T must be POD
-                //
-                // This function is absolutely sketchy and we really should be using bytemuck.
-                unsafe fn get_data<T, const COUNT: usize>(data: &[u8], offset: u32) -> &[T; COUNT] {
+                fn get_data<T, const COUNT: usize>(data: &[u8], offset: u32) -> [T; COUNT]
+                where
+                    [T; COUNT]: bytemuck::AnyBitPattern,
+                {
                     let data_required = size_of::<T>() * COUNT;
-
                     let raw = &data[(offset as usize)..][..data_required];
-
-                    debug_assert_eq!(data_required, raw.len());
-
-                    let slice: &[T] = unsafe { slice::from_raw_parts(raw.as_ptr().cast(), COUNT) };
-
-                    slice.try_into().unwrap()
+                    bytemuck::pod_read_unaligned(raw)
                 }
 
                 let location = Some(&uniform.location);
@@ -1625,28 +1619,28 @@ impl super::Queue {
                     // --- Float 1-4 Component ---
                     //
                     naga::TypeInner::Scalar(naga::Scalar::F32) => {
-                        let data = unsafe { get_data::<f32, 1>(data_bytes, offset)[0] };
+                        let data = get_data::<f32, 1>(data_bytes, offset)[0];
                         unsafe { gl.uniform_1_f32(location, data) };
                     }
                     naga::TypeInner::Vector {
                         size: naga::VectorSize::Bi,
                         scalar: naga::Scalar::F32,
                     } => {
-                        let data = unsafe { get_data::<f32, 2>(data_bytes, offset) };
+                        let data = &get_data::<f32, 2>(data_bytes, offset);
                         unsafe { gl.uniform_2_f32_slice(location, data) };
                     }
                     naga::TypeInner::Vector {
                         size: naga::VectorSize::Tri,
                         scalar: naga::Scalar::F32,
                     } => {
-                        let data = unsafe { get_data::<f32, 3>(data_bytes, offset) };
+                        let data = &get_data::<f32, 3>(data_bytes, offset);
                         unsafe { gl.uniform_3_f32_slice(location, data) };
                     }
                     naga::TypeInner::Vector {
                         size: naga::VectorSize::Quad,
                         scalar: naga::Scalar::F32,
                     } => {
-                        let data = unsafe { get_data::<f32, 4>(data_bytes, offset) };
+                        let data = &get_data::<f32, 4>(data_bytes, offset);
                         unsafe { gl.uniform_4_f32_slice(location, data) };
                     }
 
@@ -1654,28 +1648,28 @@ impl super::Queue {
                     // --- Int 1-4 Component ---
                     //
                     naga::TypeInner::Scalar(naga::Scalar::I32) => {
-                        let data = unsafe { get_data::<i32, 1>(data_bytes, offset)[0] };
+                        let data = get_data::<i32, 1>(data_bytes, offset)[0];
                         unsafe { gl.uniform_1_i32(location, data) };
                     }
                     naga::TypeInner::Vector {
                         size: naga::VectorSize::Bi,
                         scalar: naga::Scalar::I32,
                     } => {
-                        let data = unsafe { get_data::<i32, 2>(data_bytes, offset) };
+                        let data = &get_data::<i32, 2>(data_bytes, offset);
                         unsafe { gl.uniform_2_i32_slice(location, data) };
                     }
                     naga::TypeInner::Vector {
                         size: naga::VectorSize::Tri,
                         scalar: naga::Scalar::I32,
                     } => {
-                        let data = unsafe { get_data::<i32, 3>(data_bytes, offset) };
+                        let data = &get_data::<i32, 3>(data_bytes, offset);
                         unsafe { gl.uniform_3_i32_slice(location, data) };
                     }
                     naga::TypeInner::Vector {
                         size: naga::VectorSize::Quad,
                         scalar: naga::Scalar::I32,
                     } => {
-                        let data = unsafe { get_data::<i32, 4>(data_bytes, offset) };
+                        let data = &get_data::<i32, 4>(data_bytes, offset);
                         unsafe { gl.uniform_4_i32_slice(location, data) };
                     }
 
@@ -1683,28 +1677,28 @@ impl super::Queue {
                     // --- Uint 1-4 Component ---
                     //
                     naga::TypeInner::Scalar(naga::Scalar::U32) => {
-                        let data = unsafe { get_data::<u32, 1>(data_bytes, offset)[0] };
+                        let data = get_data::<u32, 1>(data_bytes, offset)[0];
                         unsafe { gl.uniform_1_u32(location, data) };
                     }
                     naga::TypeInner::Vector {
                         size: naga::VectorSize::Bi,
                         scalar: naga::Scalar::U32,
                     } => {
-                        let data = unsafe { get_data::<u32, 2>(data_bytes, offset) };
+                        let data = &get_data::<u32, 2>(data_bytes, offset);
                         unsafe { gl.uniform_2_u32_slice(location, data) };
                     }
                     naga::TypeInner::Vector {
                         size: naga::VectorSize::Tri,
                         scalar: naga::Scalar::U32,
                     } => {
-                        let data = unsafe { get_data::<u32, 3>(data_bytes, offset) };
+                        let data = &get_data::<u32, 3>(data_bytes, offset);
                         unsafe { gl.uniform_3_u32_slice(location, data) };
                     }
                     naga::TypeInner::Vector {
                         size: naga::VectorSize::Quad,
                         scalar: naga::Scalar::U32,
                     } => {
-                        let data = unsafe { get_data::<u32, 4>(data_bytes, offset) };
+                        let data = &get_data::<u32, 4>(data_bytes, offset);
                         unsafe { gl.uniform_4_u32_slice(location, data) };
                     }
 
@@ -1716,7 +1710,7 @@ impl super::Queue {
                         rows: naga::VectorSize::Bi,
                         scalar: naga::Scalar::F32,
                     } => {
-                        let data = unsafe { get_data::<f32, 4>(data_bytes, offset) };
+                        let data = &get_data::<f32, 4>(data_bytes, offset);
                         unsafe { gl.uniform_matrix_2_f32_slice(location, false, data) };
                     }
                     naga::TypeInner::Matrix {
@@ -1725,7 +1719,7 @@ impl super::Queue {
                         scalar: naga::Scalar::F32,
                     } => {
                         // repack 2 vec3s into 6 values.
-                        let unpacked_data = unsafe { get_data::<f32, 8>(data_bytes, offset) };
+                        let unpacked_data = &get_data::<f32, 8>(data_bytes, offset);
                         #[rustfmt::skip]
                         let packed_data = [
                             unpacked_data[0], unpacked_data[1], unpacked_data[2],
@@ -1738,7 +1732,7 @@ impl super::Queue {
                         rows: naga::VectorSize::Quad,
                         scalar: naga::Scalar::F32,
                     } => {
-                        let data = unsafe { get_data::<f32, 8>(data_bytes, offset) };
+                        let data = &get_data::<f32, 8>(data_bytes, offset);
                         unsafe { gl.uniform_matrix_2x4_f32_slice(location, false, data) };
                     }
 
@@ -1750,7 +1744,7 @@ impl super::Queue {
                         rows: naga::VectorSize::Bi,
                         scalar: naga::Scalar::F32,
                     } => {
-                        let data = unsafe { get_data::<f32, 6>(data_bytes, offset) };
+                        let data = &get_data::<f32, 6>(data_bytes, offset);
                         unsafe { gl.uniform_matrix_3x2_f32_slice(location, false, data) };
                     }
                     naga::TypeInner::Matrix {
@@ -1759,7 +1753,7 @@ impl super::Queue {
                         scalar: naga::Scalar::F32,
                     } => {
                         // repack 3 vec3s into 9 values.
-                        let unpacked_data = unsafe { get_data::<f32, 12>(data_bytes, offset) };
+                        let unpacked_data = &get_data::<f32, 12>(data_bytes, offset);
                         #[rustfmt::skip]
                         let packed_data = [
                             unpacked_data[0], unpacked_data[1], unpacked_data[2],
@@ -1773,7 +1767,7 @@ impl super::Queue {
                         rows: naga::VectorSize::Quad,
                         scalar: naga::Scalar::F32,
                     } => {
-                        let data = unsafe { get_data::<f32, 12>(data_bytes, offset) };
+                        let data = &get_data::<f32, 12>(data_bytes, offset);
                         unsafe { gl.uniform_matrix_3x4_f32_slice(location, false, data) };
                     }
 
@@ -1785,7 +1779,7 @@ impl super::Queue {
                         rows: naga::VectorSize::Bi,
                         scalar: naga::Scalar::F32,
                     } => {
-                        let data = unsafe { get_data::<f32, 8>(data_bytes, offset) };
+                        let data = &get_data::<f32, 8>(data_bytes, offset);
                         unsafe { gl.uniform_matrix_4x2_f32_slice(location, false, data) };
                     }
                     naga::TypeInner::Matrix {
@@ -1794,7 +1788,7 @@ impl super::Queue {
                         scalar: naga::Scalar::F32,
                     } => {
                         // repack 4 vec3s into 12 values.
-                        let unpacked_data = unsafe { get_data::<f32, 16>(data_bytes, offset) };
+                        let unpacked_data = &get_data::<f32, 16>(data_bytes, offset);
                         #[rustfmt::skip]
                         let packed_data = [
                             unpacked_data[0], unpacked_data[1], unpacked_data[2],
@@ -1809,7 +1803,7 @@ impl super::Queue {
                         rows: naga::VectorSize::Quad,
                         scalar: naga::Scalar::F32,
                     } => {
-                        let data = unsafe { get_data::<f32, 16>(data_bytes, offset) };
+                        let data = &get_data::<f32, 16>(data_bytes, offset);
                         unsafe { gl.uniform_matrix_4_f32_slice(location, false, data) };
                     }
                     _ => panic!("Unsupported uniform datatype: {:?}!", uniform.ty),
diff --git a/wgpu-hal/src/gles/web.rs b/wgpu-hal/src/gles/web.rs
index ae7f836223..06bd871247 100644
--- a/wgpu-hal/src/gles/web.rs
+++ b/wgpu-hal/src/gles/web.rs
@@ -64,9 +64,10 @@ impl Instance {
                 // “not supported” could include “insufficient GPU resources” or “the GPU process
                 // previously crashed”. So, we must return it as an `Err` since it could occur
                 // for circumstances outside the application author's control.
-                return Err(crate::InstanceError::new(String::from(
-                    "canvas.getContext() returned null; webgl2 not available or canvas already in use"
-                )));
+                return Err(crate::InstanceError::new(String::from(concat!(
+                    "canvas.getContext() returned null; ",
+                    "webgl2 not available or canvas already in use"
+                ))));
             }
             Err(js_error) => {
                 // <https://html.spec.whatwg.org/multipage/canvas.html#dom-canvas-getcontext>
diff --git a/wgpu-hal/src/vulkan/instance.rs b/wgpu-hal/src/vulkan/instance.rs
index 5673859e45..6d56ecf964 100644
--- a/wgpu-hal/src/vulkan/instance.rs
+++ b/wgpu-hal/src/vulkan/instance.rs
@@ -745,7 +745,12 @@ impl crate::Instance for super::Instance {
                     Ok(sdk_ver) => sdk_ver,
                     Err(err) => {
                         log::error!(
-                            "Couldn't parse Android's ro.build.version.sdk system property ({val}): {err}"
+                            concat!(
+                                "Couldn't parse Android's ",
+                                "ro.build.version.sdk system property ({}): {}",
+                            ),
+                            val,
+                            err,
                         );
                         0
                     }
@@ -931,7 +936,10 @@ impl crate::Instance for super::Instance {
                         if version < (21, 2) {
                             // See https://gitlab.freedesktop.org/mesa/mesa/-/issues/4688
                             log::warn!(
-                                "Disabling presentation on '{}' (id {:?}) due to NV Optimus and Intel Mesa < v21.2",
+                                concat!(
+                                    "Disabling presentation on '{}' (id {:?}) ",
+                                    "due to NV Optimus and Intel Mesa < v21.2"
+                                ),
                                 exposed.info.name,
                                 exposed.adapter.raw
                             );
diff --git a/wgpu-hal/src/vulkan/mod.rs b/wgpu-hal/src/vulkan/mod.rs
index 843b836f46..3b0f026fde 100644
--- a/wgpu-hal/src/vulkan/mod.rs
+++ b/wgpu-hal/src/vulkan/mod.rs
@@ -418,7 +418,13 @@ impl Surface {
             swapchain.next_present_time = Some(present_timing);
         } else {
             // Ideally we'd use something like `device.required_features` here, but that's in `wgpu-core`, which we are a dependency of
-            panic!("Tried to set display timing properties without the corresponding feature ({features:?}) enabled.");
+            panic!(
+                concat!(
+                    "Tried to set display timing properties ",
+                    "without the corresponding feature ({:?}) enabled."
+                ),
+                features
+            );
         }
     }
 }