Repository: gfx-rs/wgpu
Pull request #6553 (2024-11-15 19:00:20)

Improve binding error
**Description**
Error messages in case of a mismatch between resource binding as it is in the shader and as it is in the binding layout only say "there is a mismatch" without precising what binding type was found where. This improves the error message.

**Checklist**

- [x] Run `cargo fmt`.
- [x] Run `taplo format`.
- [x] Run `cargo clippy`. If applicable, add:
  - [x] `--target wasm32-unknown-unknown`
  - [x] `--target wasm32-unknown-emscripten`
- [x] Run `cargo xtask test` to run tests.
- [x] Add change to `CHANGELOG.md`. See simple instructions inside file.



Patch:
diff --git a/CHANGELOG.md b/CHANGELOG.md
index 5149ae6c4..f482cc829 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -90,7 +90,7 @@ There are some limitations to keep in mind with this new functionality:
 - Not all lints can be controlled with `diagnostic(…)` rules. In fact, only the `derivative_uniformity` triggering rule exists in the WGSL standard. That said, Naga contributors are excited to see how this level of control unlocks a new ecosystem of configurable diagnostics.
 - Finally, `diagnostic(…)` rules are not yet emitted in WGSL output. This means that `wgsl-in` → `wgsl-out` is currently a lossy process. We felt that it was important to unblock users who needed `diagnostic(…)` rules (i.e., <https://github.com/gfx-rs/wgpu/issues/3135>) before we took significant effort to fix this (tracked in <https://github.com/gfx-rs/wgpu/issues/6496>).
 
-By @ErichDonGubler in [#6456](https://github.com/gfx-rs/wgpu/pull/6456), [#6148](https://github.com/gfx-rs/wgpu/pull/6148), [#6533](https://github.com/gfx-rs/wgpu/pull/6533), [#6353](https://github.com/gfx-rs/wgpu/pull/6353).
+By @ErichDonGubler in [#6456](https://github.com/gfx-rs/wgpu/pull/6456), [#6148](https://github.com/gfx-rs/wgpu/pull/6148), [#6533](https://github.com/gfx-rs/wgpu/pull/6533), [#6353](https://github.com/gfx-rs/wgpu/pull/6353), [#6537](https://github.com/gfx-rs/wgpu/pull/6537).
 
 ### New Features
 
@@ -104,6 +104,7 @@ By @ErichDonGubler in [#6456](https://github.com/gfx-rs/wgpu/pull/6456), [#6148]
 #### General
 
 - Return submission index in `map_async` and `on_submitted_work_done` to track down completion of async callbacks. By @eliemichel in [#6360](https://github.com/gfx-rs/wgpu/pull/6360).
+- Move raytracing alignments into HAL instead of in core. By @Vecvec in [#6563](https://github.com/gfx-rs/wgpu/pull/6563).
 
 ### Changes
 
@@ -115,6 +116,7 @@ By @ErichDonGubler in [#6456](https://github.com/gfx-rs/wgpu/pull/6456), [#6148]
 
 - Make `Surface::as_hal` take an immutable reference to the surface. By @jerzywilczek in [#9999](https://github.com/gfx-rs/wgpu/pull/9999)
 - Add actual sample type to `CreateBindGroupError::InvalidTextureSampleType` error message. By @ErichDonGubler in [#6530](https://github.com/gfx-rs/wgpu/pull/6530).
+- Improve binding error to give a clearer message when there is a mismatch between resource binding as it is in the shader and as it is in the binding layout. By @eliemichel in [#6553](https://github.com/gfx-rs/wgpu/pull/6553).
 
 #### HAL
 
@@ -130,6 +132,12 @@ By @ErichDonGubler in [#6456](https://github.com/gfx-rs/wgpu/pull/6456), [#6148]
 - Lower `QUERY_SET_MAX_QUERIES` (and enforced limits) from 8192 to 4096 to match WebGPU spec. By @ErichDonGubler in [#6525](https://github.com/gfx-rs/wgpu/pull/6525).
 - Allow non-filterable float on texture bindings never used with samplers when using a derived bind group layout. By @ErichDonGubler in [#6531](https://github.com/gfx-rs/wgpu/pull/6531/).
 - Replace potentially unsound usage of `PreHashedMap` with `FastHashMap`. By @jamienicol in [#6541](https://github.com/gfx-rs/wgpu/pull/6541).
+- Add missing validation for timestamp writes in compute and render passes. By @ErichDonGubler in [#6578](https://github.com/gfx-rs/wgpu/pull/6578), [#6583](https://github.com/gfx-rs/wgpu/pull/6583).
+  - Check the status of the `TIMESTAMP_QUERY` feature before other validation.
+  - Check that indices are in-bounds for the query set.
+  - Check that begin and end indices are not equal.
+  - Check that at least one index is specified.
+- Reject destroyed buffers in query set resolution. By @ErichDonGubler in [#6579](https://github.com/gfx-rs/wgpu/pull/6579).
 
 #### Naga
 
@@ -3482,4 +3490,4 @@ DeviceDescriptor {
 - concept of the storage hub
 - basic recording of passes and command buffers
 - submission-based lifetime tracking and command buffer recycling
-- automatic resource transitions
\ No newline at end of file
+- automatic resource transitions
diff --git a/Cargo.lock b/Cargo.lock
index 8ff6df69e..14d1594e4 100644
--- a/Cargo.lock
+++ b/Cargo.lock
@@ -1311,9 +1311,9 @@ dependencies = [
 
 [[package]]
 name = "glow"
-version = "0.14.2"
+version = "0.16.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "d51fa363f025f5c111e03f13eda21162faeacb6911fe8caa0c0349f9cf0c4483"
+checksum = "c5e5ea60d70410161c8bf5da3fdfeaa1c72ed2c15f8bbb9d19fe3a4fad085f08"
 dependencies = [
  "js-sys",
  "slotmap",
@@ -1822,9 +1822,9 @@ dependencies = [
 
 [[package]]
 name = "metal"
-version = "0.29.0"
+version = "0.30.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "7ecfd3296f8c56b7c1f6fbac3c71cefa9d78ce009850c45000015f206dc7fa21"
+checksum = "9c3572083504c43e14aec05447f8a3d57cce0f66d7a3c1b9058572eca4d70ab9"
 dependencies = [
  "bitflags 2.6.0",
  "block",
diff --git a/Cargo.toml b/Cargo.toml
index 5e1142575..fb3f68a8f 100644
--- a/Cargo.toml
+++ b/Cargo.toml
@@ -136,7 +136,7 @@ winit = { version = "0.29", features = ["android-native-activity"] }
 # Metal dependencies
 block = "0.1"
 core-graphics-types = "0.1"
-metal = { version = "0.29.0" }
+metal = { version = "0.30.0" }
 objc = "0.2.5"
 
 # Vulkan dependencies
@@ -153,7 +153,7 @@ windows-core = { version = "0.58", default-features = false }
 
 # Gles dependencies
 khronos-egl = "6"
-glow = "0.14.2"
+glow = "0.16"
 glutin = { version = "0.31", default-features = false }
 glutin-winit = { version = "0.4", default-features = false }
 glutin_wgl_sys = "0.6"
diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs
index 282e28fe1..ed91aa248 100644
--- a/naga/src/back/msl/writer.rs
+++ b/naga/src/back/msl/writer.rs
@@ -383,10 +383,10 @@ pub struct Writer<W> {
     /// padding inserted **before** them (i.e. between fields at index - 1 and index)
     struct_member_pads: FastHashSet<(Handle<crate::Type>, u32)>,
 
-    /// Name of the loop reachability macro.
+    /// Name of the force-bounded-loop macro.
     ///
-    /// See `emit_loop_reachable_macro` for details.
-    loop_reachable_macro_name: String,
+    /// See `emit_force_bounded_loop_macro` for details.
+    force_bounded_loop_macro_name: String,
 }
 
 impl crate::Scalar {
@@ -682,7 +682,7 @@ impl<W: Write> Writer<W> {
             #[cfg(test)]
             put_block_stack_pointers: Default::default(),
             struct_member_pads: FastHashSet::default(),
-            loop_reachable_macro_name: String::default(),
+            force_bounded_loop_macro_name: String::default(),
         }
     }
 
@@ -693,12 +693,13 @@ impl<W: Write> Writer<W> {
         self.out
     }
 
-    /// Define a macro to invoke before loops, to defeat MSL infinite loop
-    /// reasoning.
+    /// Define a macro to invoke at the bottom of each loop body, to
+    /// defeat MSL infinite loop reasoning.
     ///
     /// If we haven't done so already, emit the definition of a preprocessor
-    /// macro to be invoked before each loop in the generated MSL, to ensure
-    /// that the MSL compiler's optimizations do not remove bounds checks.
+    /// macro to be invoked at the end of each loop body in the generated MSL,
+    /// to ensure that the MSL compiler's optimizations do not remove bounds
+    /// checks.
     ///
     /// Only the first call to this function for a given module actually causes
     /// the macro definition to be written. Subsequent loops can simply use the
@@ -764,52 +765,51 @@ impl<W: Write> Writer<W> {
     /// nicely, after having stolen data from elsewhere in the GPU address
     /// space.
     ///
-    /// Ideally, Naga would prevent UB entirely via some means that persuades
-    /// the MSL compiler that no loop Naga generates is infinite. One approach
-    /// would be to add inline assembly to each loop that is annotated as
-    /// potentially branching out of the loop, but which in fact generates no
-    /// instructions. Unfortunately, inline assembly is not handled correctly by
-    /// some Metal device drivers. Further experimentation hasn't produced a
-    /// satisfactory approach.
+    /// To avoid UB, Naga must persuade the MSL compiler that no loop Naga
+    /// generates is infinite. One approach would be to add inline assembly to
+    /// each loop that is annotated as potentially branching out of the loop,
+    /// but which in fact generates no instructions. Unfortunately, inline
+    /// assembly is not handled correctly by some Metal device drivers.
     ///
-    /// Instead, we accept that the MSL compiler may determine that some loops
-    /// are infinite, and focus instead on preventing the range analysis from
-    /// being affected. We transform *every* loop into something like this:
+    /// Instead, we add the following code to the bottom of every loop:
     ///
     /// ```ignore
-    /// if (volatile bool unpredictable = true; unpredictable)
-    ///     while (true) { }
+    /// if (volatile bool unpredictable = false; unpredictable)
+    ///     break;
     /// ```
     ///
-    /// Since the `volatile` qualifier prevents the compiler from assuming that
-    /// the `if` condition is true, it cannot be sure the infinite loop is
-    /// reached, and thus it cannot assume the entire structure is unreachable.
-    /// This prevents the range analysis impact described above.
+    /// Although the `if` condition will always be false in any real execution,
+    /// the `volatile` qualifier prevents the compiler from assuming this. Thus,
+    /// it must assume that the `break` might be reached, and hence that the
+    /// loop is not unbounded. This prevents the range analysis impact described
+    /// above.
     ///
     /// Unfortunately, what makes this a kludge, not a hack, is that this
     /// solution leaves the GPU executing a pointless conditional branch, at
-    /// runtime, before each loop. There's no part of the system that has a
-    /// global enough view to be sure that `unpredictable` is true, and remove
-    /// it from the code.
+    /// runtime, in every iteration of the loop. There's no part of the system
+    /// that has a global enough view to be sure that `unpredictable` is true,
+    /// and remove it from the code. Adding the branch also affects
+    /// optimization: for example, it's impossible to unroll this loop. This
+    /// transformation has been observed to significantly hurt performance.
     ///
     /// To make our output a bit more legible, we pull the condition out into a
     /// preprocessor macro defined at the top of the module.
     ///
-    /// This approach is also used by Chromium WebGPU's Dawn shader compiler, as of
-    /// <https://github.com/google/dawn/commit/ffd485c685040edb1e678165dcbf0e841cfa0298>.
-    fn emit_loop_reachable_macro(&mut self) -> BackendResult {
-        if !self.loop_reachable_macro_name.is_empty() {
+    /// This approach is also used by Chromium WebGPU's Dawn shader compiler:
+    /// <https://dawn.googlesource.com/dawn/+/a37557db581c2b60fb1cd2c01abdb232927dd961/src/tint/lang/msl/writer/printer/printer.cc#222>
+    fn emit_force_bounded_loop_macro(&mut self) -> BackendResult {
+        if !self.force_bounded_loop_macro_name.is_empty() {
             return Ok(());
         }
 
-        self.loop_reachable_macro_name = self.namer.call("LOOP_IS_REACHABLE");
-        let loop_reachable_volatile_name = self.namer.call("unpredictable_jump_over_loop");
+        self.force_bounded_loop_macro_name = self.namer.call("LOOP_IS_BOUNDED");
+        let loop_bounded_volatile_name = self.namer.call("unpredictable_break_from_loop");
         writeln!(
             self.out,
-            "#define {} if (volatile bool {} = true; {})",
-            self.loop_reachable_macro_name,
-            loop_reachable_volatile_name,
-            loop_reachable_volatile_name,
+            "#define {} {{ volatile bool {} = false; if ({}) break; }}",
+            self.force_bounded_loop_macro_name,
+            loop_bounded_volatile_name,
+            loop_bounded_volatile_name,
         )?;
 
         Ok(())
@@ -3045,15 +3045,10 @@ impl<W: Write> Writer<W> {
                     ref continuing,
                     break_if,
                 } => {
-                    self.emit_loop_reachable_macro()?;
                     if !continuing.is_empty() || break_if.is_some() {
                         let gate_name = self.namer.call("loop_init");
                         writeln!(self.out, "{level}bool {gate_name} = true;")?;
-                        writeln!(
-                            self.out,
-                            "{level}{} while(true) {{",
-                            self.loop_reachable_macro_name,
-                        )?;
+                        writeln!(self.out, "{level}while(true) {{",)?;
                         let lif = level.next();
                         let lcontinuing = lif.next();
                         writeln!(self.out, "{lif}if (!{gate_name}) {{")?;
@@ -3068,13 +3063,16 @@ impl<W: Write> Writer<W> {
                         writeln!(self.out, "{lif}}}")?;
                         writeln!(self.out, "{lif}{gate_name} = false;")?;
                     } else {
-                        writeln!(
-                            self.out,
-                            "{level}{} while(true) {{",
-                            self.loop_reachable_macro_name,
-                        )?;
+                        writeln!(self.out, "{level}while(true) {{",)?;
                     }
                     self.put_block(level.next(), body, context)?;
+                    self.emit_force_bounded_loop_macro()?;
+                    writeln!(
+                        self.out,
+                        "{}{}",
+                        level.next(),
+                        self.force_bounded_loop_macro_name
+                    )?;
                     writeln!(self.out, "{level}}}")?;
                 }
                 crate::Statement::Break => {
@@ -3553,7 +3551,7 @@ impl<W: Write> Writer<W> {
             &[CLAMPED_LOD_LOAD_PREFIX],
             &mut self.names,
         );
-        self.loop_reachable_macro_name.clear();
+        self.force_bounded_loop_macro_name.clear();
         self.struct_member_pads.clear();
 
         writeln!(
diff --git a/naga/src/back/spv/image.rs b/naga/src/back/spv/image.rs
index a76d015f3..46b1dd22d 100644
--- a/naga/src/back/spv/image.rs
+++ b/naga/src/back/spv/image.rs
@@ -924,7 +924,43 @@ impl<'w> BlockContext<'w> {
                     depth_id,
                 );
 
-                let lod_id = self.cached[lod_handle];
+                let mut lod_id = self.cached[lod_handle];
+                // SPIR-V expects the LOD to be a float for all image classes.
+                // lod_id, however, will be an integer for depth images,
+                // therefore we must do a conversion.
+                if matches!(
+                    self.ir_module.types[image_type].inner,
+                    crate::TypeInner::Image {
+                        class: crate::ImageClass::Depth { .. },
+                        ..
+                    }
+                ) {
+                    let lod_f32_id = self.gen_id();
+                    let f32_type_id = self.get_type_id(LookupType::Local(LocalType::Numeric(
+                        NumericType::Scalar(crate::Scalar::F32),
+                    )));
+                    let convert_op = match *self.fun_info[lod_handle]
+                        .ty
+                        .inner_with(&self.ir_module.types)
+                    {
+                        crate::TypeInner::Scalar(crate::Scalar {
+                            kind: crate::ScalarKind::Uint,
+                            width: 4,
+                        }) => spirv::Op::ConvertUToF,
+                        crate::TypeInner::Scalar(crate::Scalar {
+                            kind: crate::ScalarKind::Sint,
+                            width: 4,
+                        }) => spirv::Op::ConvertSToF,
+                        _ => unreachable!(),
+                    };
+                    block.body.push(Instruction::unary(
+                        convert_op,
+                        f32_type_id,
+                        lod_f32_id,
+                        lod_id,
+                    ));
+                    lod_id = lod_f32_id;
+                }
                 mask |= spirv::ImageOperands::LOD;
                 inst.add_operand(mask.bits());
                 inst.add_operand(lod_id);
diff --git a/naga/src/back/spv/writer.rs b/naga/src/back/spv/writer.rs
index 5cc37de78..b3dd14532 100644
--- a/naga/src/back/spv/writer.rs
+++ b/naga/src/back/spv/writer.rs
@@ -1800,12 +1800,13 @@ impl Writer {
             }
         }
 
-        // Matrices and arrays of matrices both require decorations,
-        // so "see through" an array to determine if they're needed.
-        let member_array_subty_inner = match arena[member.ty].inner {
-            crate::TypeInner::Array { base, .. } => &arena[base].inner,
-            ref other => other,
-        };
+        // Matrices and (potentially nested) arrays of matrices both require decorations,
+        // so "see through" any arrays to determine if they're needed.
+        let mut member_array_subty_inner = &arena[member.ty].inner;
+        while let crate::TypeInner::Array { base, .. } = *member_array_subty_inner {
+            member_array_subty_inner = &arena[base].inner;
+        }
+
         if let crate::TypeInner::Matrix {
             columns: _,
             rows,
diff --git a/naga/src/diagnostic_filter.rs b/naga/src/diagnostic_filter.rs
index 29499fbae..2fa5464cd 100644
--- a/naga/src/diagnostic_filter.rs
+++ b/naga/src/diagnostic_filter.rs
@@ -54,27 +54,40 @@ impl Severity {
 /// A filterable triggering rule in a [`DiagnosticFilter`].
 ///
 /// <https://www.w3.org/TR/WGSL/#filterable-triggering-rules>
-#[derive(Clone, Copy, Debug, Eq, Hash, Ord, PartialEq, PartialOrd)]
+#[derive(Clone, Debug, Eq, Hash, Ord, PartialEq, PartialOrd)]
 #[cfg_attr(feature = "serialize", derive(Serialize))]
 #[cfg_attr(feature = "deserialize", derive(Deserialize))]
 #[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
 pub enum FilterableTriggeringRule {
+    Standard(StandardFilterableTriggeringRule),
+    Unknown(Box<str>),
+    User(Box<[Box<str>; 2]>),
+}
+
+/// A filterable triggering rule in a [`DiagnosticFilter`].
+///
+/// <https://www.w3.org/TR/WGSL/#filterable-triggering-rules>
+#[derive(Clone, Copy, Debug, Eq, Hash, Ord, PartialEq, PartialOrd)]
+#[cfg_attr(feature = "serialize", derive(Serialize))]
+#[cfg_attr(feature = "deserialize", derive(Deserialize))]
+#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
+pub enum StandardFilterableTriggeringRule {
     DerivativeUniformity,
 }
 
-impl FilterableTriggeringRule {
+impl StandardFilterableTriggeringRule {
     /// The default severity associated with this triggering rule.
     ///
     /// See <https://www.w3.org/TR/WGSL/#filterable-triggering-rules> for a table of default
     /// severities.
     pub(crate) const fn default_severity(self) -> Severity {
         match self {
-            FilterableTriggeringRule::DerivativeUniformity => Severity::Error,
+            Self::DerivativeUniformity => Severity::Error,
         }
     }
 }
 
-/// A filter that modifies how diagnostics are emitted for shaders.
+/// A filtering rule that modifies how diagnostics are emitted for shaders.
 ///
 /// <https://www.w3.org/TR/WGSL/#diagnostic-filter>
 #[derive(Clone, Debug)]
@@ -140,7 +153,7 @@ impl DiagnosticFilterMap {
             triggering_rule,
         } = diagnostic_filter;
 
-        match diagnostic_filters.entry(triggering_rule) {
+        match diagnostic_filters.entry(triggering_rule.clone()) {
             Entry::Vacant(entry) => {
                 entry.insert((new_severity, span));
             }
@@ -152,7 +165,6 @@ impl DiagnosticFilterMap {
                 };
                 if first_severity != new_severity || should_conflict_on_full_duplicate {
                     return Err(ConflictingDiagnosticRuleError {
-                        triggering_rule,
                         triggering_rule_spans: [first_span, span],
                     });
                 }
@@ -190,7 +202,6 @@ impl IntoIterator for DiagnosticFilterMap {
 #[cfg(feature = "wgsl-in")]
 #[derive(Clone, Debug)]
 pub(crate) struct ConflictingDiagnosticRuleError {
-    pub triggering_rule: FilterableTriggeringRule,
     pub triggering_rule_spans: [Span; 2],
 }
 
@@ -232,24 +243,24 @@ pub struct DiagnosticFilterNode {
 impl DiagnosticFilterNode {
     /// Finds the most specific filter rule applicable to `triggering_rule` from the chain of
     /// diagnostic filter rules in `arena`, starting with `node`, and returns its severity. If none
-    /// is found, return the value of [`FilterableTriggeringRule::default_severity`].
+    /// is found, return the value of [`StandardFilterableTriggeringRule::default_severity`].
     ///
     /// When `triggering_rule` is not applicable to this node, its parent is consulted recursively.
     pub(crate) fn search(
         node: Option<Handle<Self>>,
         arena: &Arena<Self>,
-        triggering_rule: FilterableTriggeringRule,
+        triggering_rule: StandardFilterableTriggeringRule,
     ) -> Severity {
         let mut next = node;
         while let Some(handle) = next {
             let node = &arena[handle];
             let &Self { ref inner, parent } = node;
             let &DiagnosticFilter {
-                triggering_rule: rule,
+                triggering_rule: ref rule,
                 new_severity,
             } = inner;
 
-            if rule == triggering_rule {
+            if rule == &FilterableTriggeringRule::Standard(triggering_rule) {
                 return new_severity;
             }
 
diff --git a/naga/src/front/spv/image.rs b/naga/src/front/spv/image.rs
index 71ba79e1e..1d2b8dc64 100644
--- a/naga/src/front/spv/image.rs
+++ b/naga/src/front/spv/image.rs
@@ -30,6 +30,7 @@ impl<'function> super::BlockContext<'function> {
         match self.expressions[handle] {
             crate::Expression::GlobalVariable(handle) => Ok(self.global_arena[handle].ty),
             crate::Expression::FunctionArgument(i) => Ok(self.arguments[i as usize].ty),
+            crate::Expression::Access { base, .. } => Ok(self.get_image_expr_ty(base)?),
             ref other => Err(Error::InvalidImageExpression(other.clone())),
         }
     }
@@ -460,6 +461,7 @@ impl<I: Iterator<Item = u32>> super::Frontend<I> {
         } else {
             None
         };
+        let span = self.span_from_with_op(start);
 
         let mut image_ops = if words_left != 0 {
             words_left -= 1;
@@ -486,9 +488,34 @@ impl<I: Iterator<Item = u32>> super::Frontend<I> {
                     let lod_lexp = self.lookup_expression.lookup(lod_expr)?;
                     let lod_handle =
                         self.get_expr_handle(lod_expr, lod_lexp, ctx, emitter, block, body_idx);
+
+                    let is_depth_image = {
+                        let image_lexp = self.lookup_sampled_image.lookup(sampled_image_id)?;
+                        let image_ty = ctx.get_image_expr_ty(image_lexp.image)?;
+                        matches!(
+                            ctx.type_arena[image_ty].inner,
+                            crate::TypeInner::Image {
+                                class: crate::ImageClass::Depth { .. },
+                                ..
+                            }
+                        )
+                    };
+
                     level = if options.compare {
                         log::debug!("Assuming {:?} is zero", lod_handle);
                         crate::SampleLevel::Zero
+                    } else if is_depth_image {
+                        log::debug!(
+                            "Assuming level {:?} converts losslessly to an integer",
+                            lod_handle
+                        );
+                        let expr = crate::Expression::As {
+                            expr: lod_handle,
+                            kind: crate::ScalarKind::Sint,
+                            convert: Some(4),
+                        };
+                        let s32_lod_handle = ctx.expressions.append(expr, span);
+                        crate::SampleLevel::Exact(s32_lod_handle)
                     } else {
                         crate::SampleLevel::Exact(lod_handle)
                     };
diff --git a/naga/src/front/wgsl/diagnostic_filter.rs b/naga/src/front/wgsl/diagnostic_filter.rs
index f505205b6..b3cd7956b 100644
--- a/naga/src/front/wgsl/diagnostic_filter.rs
+++ b/naga/src/front/wgsl/diagnostic_filter.rs
@@ -1,4 +1,8 @@
-use crate::diagnostic_filter::{FilterableTriggeringRule, Severity};
+use std::fmt::{self, Display, Formatter};
+
+use crate::diagnostic_filter::{
+    FilterableTriggeringRule, Severity, StandardFilterableTriggeringRule,
+};
 
 impl Severity {
     const ERROR: &'static str = "error";
@@ -18,10 +22,34 @@ impl Severity {
     }
 }
 
+struct DisplayFilterableTriggeringRule<'a>(&'a FilterableTriggeringRule);
+
+impl Display for DisplayFilterableTriggeringRule<'_> {
+    fn fmt(&self, f: &mut Formatter<'_>) -> fmt::Result {
+        let &Self(inner) = self;
+        match *inner {
+            FilterableTriggeringRule::Standard(rule) => write!(f, "{}", rule.to_wgsl_ident()),
+            FilterableTriggeringRule::Unknown(ref rule) => write!(f, "{rule}"),
+            FilterableTriggeringRule::User(ref rules) => {
+                let &[ref seg1, ref seg2] = rules.as_ref();
+                write!(f, "{seg1}.{seg2}")
+            }
+        }
+    }
+}
+
 impl FilterableTriggeringRule {
+    /// [`Display`] this rule's identifiers in WGSL.
+    pub const fn display_wgsl_ident(&self) -> impl Display + '_ {
+        DisplayFilterableTriggeringRule(self)
+    }
+}
+
+impl StandardFilterableTriggeringRule {
     const DERIVATIVE_UNIFORMITY: &'static str = "derivative_uniformity";
 
-    /// Convert from a sentinel word in WGSL into its associated [`FilterableTriggeringRule`], if possible.
+    /// Convert from a sentinel word in WGSL into its associated
+    /// [`StandardFilterableTriggeringRule`], if possible.
     pub fn from_wgsl_ident(s: &str) -> Option<Self> {
         Some(match s {
             Self::DERIVATIVE_UNIFORMITY => Self::DerivativeUniformity,
@@ -29,10 +57,152 @@ impl FilterableTriggeringRule {
         })
     }
 
-    /// Maps this [`FilterableTriggeringRule`] into the sentinel word associated with it in WGSL.
+    /// Maps this [`StandardFilterableTriggeringRule`] into the sentinel word associated with it in
+    /// WGSL.
     pub const fn to_wgsl_ident(self) -> &'static str {
         match self {
             Self::DerivativeUniformity => Self::DERIVATIVE_UNIFORMITY,
         }
     }
 }
+
+#[cfg(test)]
+mod test {
+    mod parse_sites_not_yet_supported {
+        use crate::front::wgsl::assert_parse_err;
+
+        #[test]
+        fn user_rules() {
+            let shader = "
+fn myfunc() {
+    if (true) @diagnostic(off, my.lint) {
+        //    ^^^^^^^^^^^^^^^^^^^^^^^^^ not yet supported, should report an error
+    }
+}
+";
+            assert_parse_err(shader, "\
+error: `@diagnostic(…)` attribute(s) not yet implemented
+  ┌─ wgsl:3:15
+  │
+3 │     if (true) @diagnostic(off, my.lint) {
+  │               ^^^^^^^^^^^^^^^^^^^^^^^^^ can't use this on compound statements (yet)
+  │
+  = note: Let Naga maintainers know that you ran into this at <https://github.com/gfx-rs/wgpu/issues/5320>, so they can prioritize it!
+
+");
+        }
+
+        #[test]
+        fn unknown_rules() {
+            let shader = "
+fn myfunc() {
+	if (true) @diagnostic(off, wat_is_this) {
+		//    ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ should emit a warning
+	}
+}
+";
+            assert_parse_err(shader, "\
+error: `@diagnostic(…)` attribute(s) not yet implemented
+  ┌─ wgsl:3:12
+  │
+3 │     if (true) @diagnostic(off, wat_is_this) {
+  │               ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ can't use this on compound statements (yet)
+  │
+  = note: Let Naga maintainers know that you ran into this at <https://github.com/gfx-rs/wgpu/issues/5320>, so they can prioritize it!
+
+");
+        }
+    }
+
+    mod directive_conflict {
+        use crate::front::wgsl::assert_parse_err;
+
+        #[test]
+        fn user_rules() {
+            let shader = "
+diagnostic(off, my.lint);
+diagnostic(warning, my.lint);
+";
+            assert_parse_err(shader, "\
+error: found conflicting `diagnostic(…)` rule(s)
+  ┌─ wgsl:2:1
+  │
+2 │ diagnostic(off, my.lint);
+  │ ^^^^^^^^^^^^^^^^^^^^^^^^ first rule
+3 │ diagnostic(warning, my.lint);
+  │ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^ second rule
+  │
+  = note: Multiple `diagnostic(…)` rules with the same rule name conflict unless they are directives and the severity is the same.
+  = note: You should delete the rule you don't want.
+
+");
+        }
+
+        #[test]
+        fn unknown_rules() {
+            let shader = "
+diagnostic(off, wat_is_this);
+diagnostic(warning, wat_is_this);
+";
+            assert_parse_err(shader, "\
+error: found conflicting `diagnostic(…)` rule(s)
+  ┌─ wgsl:2:1
+  │
+2 │ diagnostic(off, wat_is_this);
+  │ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^ first rule
+3 │ diagnostic(warning, wat_is_this);
+  │ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ second rule
+  │
+  = note: Multiple `diagnostic(…)` rules with the same rule name conflict unless they are directives and the severity is the same.
+  = note: You should delete the rule you don't want.
+
+");
+        }
+    }
+
+    mod attribute_conflict {
+        use crate::front::wgsl::assert_parse_err;
+
+        #[test]
+        fn user_rules() {
+            let shader = "
+diagnostic(off, my.lint);
+diagnostic(warning, my.lint);
+";
+            assert_parse_err(shader, "\
+error: found conflicting `diagnostic(…)` rule(s)
+  ┌─ wgsl:2:1
+  │
+2 │ diagnostic(off, my.lint);
+  │ ^^^^^^^^^^^^^^^^^^^^^^^^ first rule
+3 │ diagnostic(warning, my.lint);
+  │ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^ second rule
+  │
+  = note: Multiple `diagnostic(…)` rules with the same rule name conflict unless they are directives and the severity is the same.
+  = note: You should delete the rule you don't want.
+
+");
+        }
+
+        #[test]
+        fn unknown_rules() {
+            let shader = "
+diagnostic(off, wat_is_this);
+diagnostic(warning, wat_is_this);
+";
+            assert_parse_err(shader, "\
+error: found conflicting `diagnostic(…)` rule(s)
+  ┌─ wgsl:2:1
+  │
+2 │ diagnostic(off, wat_is_this);
+  │ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^ first rule
+3 │ diagnostic(warning, wat_is_this);
+  │ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ second rule
+  │
+  = note: Multiple `diagnostic(…)` rules with the same rule name conflict unless they are directives and the severity is the same.
+  = note: You should delete the rule you don't want.
+
+");
+        }
+    }
+}
diff --git a/naga/src/front/wgsl/error.rs b/naga/src/front/wgsl/error.rs
index 26199b5e5..f827ca50f 100644
--- a/naga/src/front/wgsl/error.rs
+++ b/naga/src/front/wgsl/error.rs
@@ -1034,15 +1034,11 @@ impl<'a> Error<'a> {
                 .into()],
             },
             Error::DiagnosticDuplicateTriggeringRule(ConflictingDiagnosticRuleError {
-                triggering_rule,
                 triggering_rule_spans,
             }) => {
                 let [first_span, second_span] = triggering_rule_spans;
                 ParseError {
-                    message: format!(
-                        "found conflicting `diagnostic(…)` rule(s) for `{}`",
-                        triggering_rule.to_wgsl_ident()
-                    ),
+                    message: "found conflicting `diagnostic(…)` rule(s)".into(),
                     labels: vec![
                         (first_span, "first rule".into()),
                         (second_span, "second rule".into()),
@@ -1050,7 +1046,7 @@ impl<'a> Error<'a> {
                     notes: vec![
                         concat!(
                             "Multiple `diagnostic(…)` rules with the same rule name ",
-                            "conflict unless it is a directive and the severity is the same.",
+                            "conflict unless they are directives and the severity is the same.",
                         )
                         .into(),
                         "You should delete the rule you don't want.".into(),
diff --git a/naga/src/front/wgsl/parse/mod.rs b/naga/src/front/wgsl/parse/mod.rs
index ae22eb38e..93b934332 100644
--- a/naga/src/front/wgsl/parse/mod.rs
+++ b/naga/src/front/wgsl/parse/mod.rs
@@ -1,6 +1,6 @@
 use crate::diagnostic_filter::{
     self, DiagnosticFilter, DiagnosticFilterMap, DiagnosticFilterNode, FilterableTriggeringRule,
-    ShouldConflictOnFullDuplicate,
+    ShouldConflictOnFullDuplicate, StandardFilterableTriggeringRule,
 };
 use crate::front::wgsl::error::{Error, ExpectedToken};
 use crate::front::wgsl::parse::directive::enable_extension::{
@@ -2166,10 +2166,9 @@ impl Parser {
         while lexer.skip(Token::Attribute) {
             let (name, name_span) = lexer.next_ident_with_span()?;
             if let Some(DirectiveKind::Diagnostic) = DirectiveKind::from_ident(name) {
-                if let Some(filter) = self.diagnostic_filter(lexer)? {
-                    let span = self.peek_rule_span(lexer);
-                    diagnostic_filters.add(filter, span, ShouldConflictOnFullDuplicate::Yes)?;
-                }
+                let filter = self.diagnostic_filter(lexer)?;
+                let span = self.peek_rule_span(lexer);
+                diagnostic_filters.add(filter, span, ShouldConflictOnFullDuplicate::Yes)?;
             } else {
                 return Err(Error::Unexpected(
                     name_span,
@@ -2370,10 +2369,9 @@ impl Parser {
         while lexer.skip(Token::Attribute) {
             let (name, name_span) = lexer.next_ident_with_span()?;
             if let Some(DirectiveKind::Diagnostic) = DirectiveKind::from_ident(name) {
-                if let Some(filter) = self.diagnostic_filter(lexer)? {
-                    let span = self.peek_rule_span(lexer);
-                    diagnostic_filters.add(filter, span, ShouldConflictOnFullDuplicate::Yes)?;
-                }
+                let filter = self.diagnostic_filter(lexer)?;
+                let span = self.peek_rule_span(lexer);
+                diagnostic_filters.add(filter, span, ShouldConflictOnFullDuplicate::Yes)?;
                 continue;
             }
             match name {
@@ -2603,14 +2601,13 @@ impl Parser {
                 let _ = lexer.next_ident_with_span().unwrap();
                 match kind {
                     DirectiveKind::Diagnostic => {
-                        if let Some(diagnostic_filter) = self.diagnostic_filter(&mut lexer)? {
-                            let span = self.peek_rule_span(&lexer);
-                            diagnostic_filters.add(
-                                diagnostic_filter,
-                                span,
-                                ShouldConflictOnFullDuplicate::No,
-                            )?;
-                        }
+                        let diagnostic_filter = self.diagnostic_filter(&mut lexer)?;
+                        let span = self.peek_rule_span(&lexer);
+                        diagnostic_filters.add(
+                            diagnostic_filter,
+                            span,
+                            ShouldConflictOnFullDuplicate::No,
+                        )?;
                         lexer.expect(Token::Separator(';'))?;
                     }
                     DirectiveKind::Enable => {
@@ -2698,10 +2695,7 @@ impl Parser {
         Ok(brace_nesting_level + 1)
     }
 
-    fn diagnostic_filter<'a>(
-        &self,
-        lexer: &mut Lexer<'a>,
-    ) -> Result<Option<DiagnosticFilter>, Error<'a>> {
+    fn diagnostic_filter<'a>(&self, lexer: &mut Lexer<'a>) -> Result<DiagnosticFilter, Error<'a>> {
         lexer.expect(Token::Paren('('))?;
 
         let (severity_control_name, severity_control_name_span) = lexer.next_ident_with_span()?;
@@ -2713,35 +2707,28 @@ impl Parser {
         lexer.expect(Token::Separator(','))?;
 
         let (diagnostic_name_token, diagnostic_name_token_span) = lexer.next_ident_with_span()?;
-        let diagnostic_rule_name = if lexer.skip(Token::Separator('.')) {
-            // Don't try to validate these name tokens on two tokens, which is conventionally used
-            // for third-party tooling.
-            lexer.next_ident_with_span()?;
-            None
+        let triggering_rule = if lexer.skip(Token::Separator('.')) {
+            let (ident, _span) = lexer.next_ident_with_span()?;
+            FilterableTriggeringRule::User(Box::new([diagnostic_name_token.into(), ident.into()]))
         } else {
-            Some(diagnostic_name_token)
+            let diagnostic_rule_name = diagnostic_name_token;
+            let diagnostic_rule_name_span = diagnostic_name_token_span;
+            if let Some(triggering_rule) =
+                StandardFilterableTriggeringRule::from_wgsl_ident(diagnostic_rule_name)
+            {
+                FilterableTriggeringRule::Standard(triggering_rule)
+            } else {
+                diagnostic_filter::Severity::Warning.report_wgsl_parse_diag(
+                    Error::UnknownDiagnosticRuleName(diagnostic_rule_name_span),
+                    lexer.source,
+                )?;
+                FilterableTriggeringRule::Unknown(diagnostic_rule_name.into())
+            }
+        };
+        let filter = DiagnosticFilter {
+            triggering_rule,
+            new_severity,
         };
-        let diagnostic_rule_name_span = diagnostic_name_token_span;
-
-        let filter = diagnostic_rule_name
-            .and_then(|name| {
-                FilterableTriggeringRule::from_wgsl_ident(name)
-                    .map(Ok)
-                    .or_else(|| {
-                        diagnostic_filter::Severity::Warning
-                            .report_wgsl_parse_diag(
-                                Error::UnknownDiagnosticRuleName(diagnostic_rule_name_span),
-                                lexer.source,
-                            )
-                            .err()
-                            .map(Err)
-                    })
-            })
-            .transpose()?
-            .map(|triggering_rule| DiagnosticFilter {
-                new_severity,
-                triggering_rule,
-            });
         lexer.skip(Token::Separator(','));
         lexer.expect(Token::Paren(')'))?;
 
diff --git a/naga/src/valid/analyzer.rs b/naga/src/valid/analyzer.rs
index 0ca9829d6..ac9bb9cd5 100644
--- a/naga/src/valid/analyzer.rs
+++ b/naga/src/valid/analyzer.rs
@@ -6,7 +6,7 @@
 //! - expression reference counts
 
 use super::{ExpressionError, FunctionError, ModuleInfo, ShaderStages, ValidationFlags};
-use crate::diagnostic_filter::{DiagnosticFilterNode, FilterableTriggeringRule};
+use crate::diagnostic_filter::{DiagnosticFilterNode, StandardFilterableTriggeringRule};
 use crate::span::{AddSpan as _, WithSpan};
 use crate::{
     arena::{Arena, Handle},
@@ -852,7 +852,7 @@ impl FunctionInfo {
                                 let severity = DiagnosticFilterNode::search(
                                     self.diagnostic_filter_leaf,
                                     diagnostic_filter_arena,
-                                    FilterableTriggeringRule::DerivativeUniformity,
+                                    StandardFilterableTriggeringRule::DerivativeUniformity,
                                 );
                                 severity.report_diag(
                                     FunctionError::NonUniformControlFlow(req, expr, cause)
@@ -1372,7 +1372,10 @@ fn uniform_control_flow() {
                 DiagnosticFilterNode {
                     inner: crate::diagnostic_filter::DiagnosticFilter {
                         new_severity: crate::diagnostic_filter::Severity::Off,
-                        triggering_rule: FilterableTriggeringRule::DerivativeUniformity,
+                        triggering_rule:
+                            crate::diagnostic_filter::FilterableTriggeringRule::Standard(
+                                StandardFilterableTriggeringRule::DerivativeUniformity,
+                            ),
                     },
                     parent: None,
                 },
diff --git a/naga/src/valid/expression.rs b/naga/src/valid/expression.rs
index 9a1503401..9f376c65c 100644
--- a/naga/src/valid/expression.rs
+++ b/naga/src/valid/expression.rs
@@ -106,10 +106,12 @@ pub enum ExpressionError {
     InvalidGatherComponent(crate::SwizzleComponent),
     #[error("Gather can't be done for image dimension {0:?}")]
     InvalidGatherDimension(crate::ImageDimension),
-    #[error("Sample level (exact) type {0:?} is not a scalar float")]
+    #[error("Sample level (exact) type {0:?} has an invalid type")]
     InvalidSampleLevelExactType(Handle<crate::Expression>),
     #[error("Sample level (bias) type {0:?} is not a scalar float")]
     InvalidSampleLevelBiasType(Handle<crate::Expression>),
+    #[error("Bias can't be done for image dimension {0:?}")]
+    InvalidSampleLevelBiasDimension(crate::ImageDimension),
     #[error("Sample level (gradient) of {1:?} doesn't match the image dimension {0:?}")]
     InvalidSampleLevelGradientType(crate::ImageDimension, Handle<crate::Expression>),
     #[error("Unable to cast")]
@@ -530,11 +532,24 @@ impl super::Validator {
                     crate::SampleLevel::Auto => ShaderStages::FRAGMENT,
                     crate::SampleLevel::Zero => ShaderStages::all(),
                     crate::SampleLevel::Exact(expr) => {
-                        match resolver[expr] {
-                            Ti::Scalar(Sc {
-                                kind: Sk::Float, ..
-                            }) => {}
-                            _ => return Err(ExpressionError::InvalidSampleLevelExactType(expr)),
+                        match class {
+                            crate::ImageClass::Depth { .. } => match resolver[expr] {
+                                Ti::Scalar(Sc {
+                                    kind: Sk::Sint | Sk::Uint,
+                                    ..
+                                }) => {}
+                                _ => {
+                                    return Err(ExpressionError::InvalidSampleLevelExactType(expr))
+                                }
+                            },
+                            _ => match resolver[expr] {
+                                Ti::Scalar(Sc {
+                                    kind: Sk::Float, ..
+                                }) => {}
+                                _ => {
+                                    return Err(ExpressionError::InvalidSampleLevelExactType(expr))
+                                }
+                            },
                         }
                         ShaderStages::all()
                     }
@@ -545,6 +560,19 @@ impl super::Validator {
                             }) => {}
                             _ => return Err(ExpressionError::InvalidSampleLevelBiasType(expr)),
                         }
+                        match class {
+                            crate::ImageClass::Sampled {
+                                kind: Sk::Float,
+                                multi: false,
+                            } => {
+                                if dim == crate::ImageDimension::D1 {
+                                    return Err(ExpressionError::InvalidSampleLevelBiasDimension(
+                                        dim,
+                                    ));
+                                }
+                            }
+                            _ => return Err(ExpressionError::InvalidImageClass(class)),
+                        }
                         ShaderStages::FRAGMENT
                     }
                     crate::SampleLevel::Gradient { x, y } => {
diff --git a/wgpu-core/src/command/bundle.rs b/wgpu-core/src/command/bundle.rs
index b0d90976d..ac7a5280b 100644
--- a/wgpu-core/src/command/bundle.rs
+++ b/wgpu-core/src/command/bundle.rs
@@ -543,8 +543,8 @@ impl RenderBundleEncoder {
                 label: desc.label.as_ref().map(|cow| cow.to_string()),
                 commands,
                 dynamic_offsets: flat_dynamic_offsets,
-                string_data: Vec::new(),
-                push_constant_data: Vec::new(),
+                string_data: self.base.string_data,
+                push_constant_data: self.base.push_constant_data,
             },
             is_depth_read_only: self.is_depth_read_only,
             is_stencil_read_only: self.is_stencil_read_only,
diff --git a/wgpu-core/src/command/compute.rs b/wgpu-core/src/command/compute.rs
index 28e76c111..89681b8e2 100644
--- a/wgpu-core/src/command/compute.rs
+++ b/wgpu-core/src/command/compute.rs
@@ -308,30 +308,15 @@ impl Global {
             Err(e) => return make_err(e, arc_desc),
         };
 
-        arc_desc.timestamp_writes = if let Some(tw) = desc.timestamp_writes {
-            let query_set = match hub.query_sets.get(tw.query_set).get() {
-                Ok(query_set) => query_set,
-                Err(e) => return make_err(e.into(), arc_desc),
-            };
-            match query_set.same_device(&cmd_buf.device) {
-                Ok(()) => (),
-                Err(e) => return make_err(e.into(), arc_desc),
-            }
-            match cmd_buf
-                .device
-                .require_features(wgt::Features::TIMESTAMP_QUERY)
-            {
-                Ok(()) => (),
-                Err(e) => return make_err(e.into(), arc_desc),
-            }
-
-            Some(ArcPassTimestampWrites {
-                query_set,
-                beginning_of_pass_write_index: tw.beginning_of_pass_write_index,
-                end_of_pass_write_index: tw.end_of_pass_write_index,
+        arc_desc.timestamp_writes = match desc
+            .timestamp_writes
+            .map(|tw| {
+                Self::validate_pass_timestamp_writes(&cmd_buf.device, &hub.query_sets.read(), tw)
             })
-        } else {
-            None
+            .transpose()
+        {
+            Ok(ok) => ok,
+            Err(e) => return make_err(e, arc_desc),
         };
 
         (ComputePass::new(Some(cmd_buf), arc_desc), None)
diff --git a/wgpu-core/src/command/mod.rs b/wgpu-core/src/command/mod.rs
index 6a06d1f83..fab58515f 100644
--- a/wgpu-core/src/command/mod.rs
+++ b/wgpu-core/src/command/mod.rs
@@ -33,7 +33,8 @@ use crate::snatch::SnatchGuard;
 
 use crate::init_tracker::BufferInitTrackerAction;
 use crate::ray_tracing::{BlasAction, TlasAction};
-use crate::resource::{InvalidResourceError, Labeled};
+use crate::resource::{Fallible, InvalidResourceError, Labeled, ParentDevice as _, QuerySet};
+use crate::storage::Storage;
 use crate::track::{DeviceTracker, Tracker, UsageScope};
 use crate::LabelHelpers;
 use crate::{api_log, global::Global, id, resource_log, Label};
@@ -654,6 +655,14 @@ pub enum CommandEncoderError {
     InvalidResource(#[from] InvalidResourceError),
     #[error(transparent)]
     MissingFeatures(#[from] MissingFeatures),
+    #[error(
+        "begin and end indices of pass timestamp writes are both set to {idx}, which is not allowed"
+    )]
+    TimestampWriteIndicesEqual { idx: u32 },
+    #[error(transparent)]
+    TimestampWritesInvalid(#[from] QueryUseError),
+    #[error("no begin or end indices were specified for pass timestamp writes, expected at least one to be set")]
+    TimestampWriteIndicesMissing,
 }
 
 impl Global {
@@ -774,6 +783,50 @@ impl Global {
         }
         Ok(())
     }
+
+    fn validate_pass_timestamp_writes(
+        device: &Device,
+        query_sets: &Storage<Fallible<QuerySet>>,
+        timestamp_writes: &PassTimestampWrites,
+    ) -> Result<ArcPassTimestampWrites, CommandEncoderError> {
+        let &PassTimestampWrites {
+            query_set,
+            beginning_of_pass_write_index,
+            end_of_pass_write_index,
+        } = timestamp_writes;
+
+        device.require_features(wgt::Features::TIMESTAMP_QUERY)?;
+
+        let query_set = query_sets.get(query_set).get()?;
+
+        query_set.same_device(device)?;
+
+        for idx in [beginning_of_pass_write_index, end_of_pass_write_index]
+            .into_iter()
+            .flatten()
+        {
+            query_set.validate_query(SimplifiedQueryType::Timestamp, idx, None)?;
+        }
+
+        if let Some((begin, end)) = beginning_of_pass_write_index.zip(end_of_pass_write_index) {
+            if begin == end {
+                return Err(CommandEncoderError::TimestampWriteIndicesEqual { idx: begin });
+            }
+        }
+
+        if beginning_of_pass_write_index
+            .or(end_of_pass_write_index)
+            .is_none()
+        {
+            return Err(CommandEncoderError::TimestampWriteIndicesMissing);
+        }
+
+        Ok(ArcPassTimestampWrites {
+            query_set,
+            beginning_of_pass_write_index,
+            end_of_pass_write_index,
+        })
+    }
 }
 
 fn push_constant_clear<PushFn>(offset: u32, size_bytes: u32, mut push_fn: PushFn)
diff --git a/wgpu-core/src/command/query.rs b/wgpu-core/src/command/query.rs
index d783721fb..e99144663 100644
--- a/wgpu-core/src/command/query.rs
+++ b/wgpu-core/src/command/query.rs
@@ -160,7 +160,7 @@ pub enum ResolveError {
 }
 
 impl QuerySet {
-    fn validate_query(
+    pub(crate) fn validate_query(
         self: &Arc<Self>,
         query_type: SimplifiedQueryType,
         query_index: u32,
@@ -387,13 +387,14 @@ impl Global {
 
         dst_buffer.same_device_as(cmd_buf.as_ref())?;
 
+        let snatch_guard = dst_buffer.device.snatchable_lock.read();
+        dst_buffer.check_destroyed(&snatch_guard)?;
+
         let dst_pending = cmd_buf_data
             .trackers
             .buffers
             .set_single(&dst_buffer, hal::BufferUses::COPY_DST);
 
-        let snatch_guard = dst_buffer.device.snatchable_lock.read();
-
         let dst_barrier = dst_pending.map(|pending| pending.into_hal(&dst_buffer, &snatch_guard));
 
         dst_buffer
diff --git a/wgpu-core/src/command/ray_tracing.rs b/wgpu-core/src/command/ray_tracing.rs
index efc3ebbe2..ec90a5d28 100644
--- a/wgpu-core/src/command/ray_tracing.rs
+++ b/wgpu-core/src/command/ray_tracing.rs
@@ -5,7 +5,7 @@ use crate::{
     id::CommandEncoderId,
     init_tracker::MemoryInitKind,
     ray_tracing::{
-        tlas_instance_into_bytes, BlasAction, BlasBuildEntry, BlasGeometries, BlasTriangleGeometry,
+        BlasAction, BlasBuildEntry, BlasGeometries, BlasTriangleGeometry,
         BuildAccelerationStructureError, TlasAction, TlasBuildEntry, TlasInstance, TlasPackage,
         TraceBlasBuildEntry, TraceBlasGeometries, TraceBlasTriangleGeometry, TraceTlasInstance,
         TraceTlasPackage, ValidateBlasActionsError, ValidateTlasActionsError,
@@ -60,9 +60,6 @@ struct TlasBufferStore {
     entry: TlasBuildEntry,
 }
 
-// TODO: Get this from the device (e.g. VkPhysicalDeviceAccelerationStructurePropertiesKHR.minAccelerationStructureScratchOffsetAlignment) this is currently the largest possible some devices have 0, 64, 128 (lower limits) so this could create excess allocation (Note: dx12 has 256).
-const SCRATCH_BUFFER_ALIGNMENT: u32 = 256;
-
 impl Global {
     // Currently this function is very similar to its safe counterpart, however certain parts of it are very different,
     // making for the two to be implemented differently, the main difference is this function has separate buffers for each
@@ -193,6 +190,7 @@ impl Global {
             &mut scratch_buffer_blas_size,
             &mut blas_storage,
             hub,
+            device.alignments.ray_tracing_scratch_buffer_alignment,
         )?;
 
         let mut scratch_buffer_tlas_size = 0;
@@ -260,7 +258,7 @@ impl Global {
             let scratch_buffer_offset = scratch_buffer_tlas_size;
             scratch_buffer_tlas_size += align_to(
                 tlas.size_info.build_scratch_size as u32,
-                SCRATCH_BUFFER_ALIGNMENT,
+                device.alignments.ray_tracing_scratch_buffer_alignment,
             ) as u64;
 
             tlas_storage.push(UnsafeTlasStore {
@@ -508,6 +506,7 @@ impl Global {
             &mut scratch_buffer_blas_size,
             &mut blas_storage,
             hub,
+            device.alignments.ray_tracing_scratch_buffer_alignment,
         )?;
         let mut tlas_lock_store = Vec::<(Option<TlasPackage>, Arc<Tlas>)>::new();
 
@@ -535,7 +534,7 @@ impl Global {
             let scratch_buffer_offset = scratch_buffer_tlas_size;
             scratch_buffer_tlas_size += align_to(
                 tlas.size_info.build_scratch_size as u32,
-                SCRATCH_BUFFER_ALIGNMENT,
+                device.alignments.ray_tracing_scratch_buffer_alignment,
             ) as u64;
 
             let first_byte_index = instance_buffer_staging_source.len();
@@ -558,10 +557,13 @@ impl Global {
 
                 cmd_buf_data.trackers.blas_s.set_single(blas.clone());
 
-                instance_buffer_staging_source.extend(tlas_instance_into_bytes(
-                    &instance,
-                    blas.handle,
-                    device.backend(),
+                instance_buffer_staging_source.extend(device.raw().tlas_instance_to_bytes(
+                    hal::TlasInstance {
+                        transform: *instance.transform,
+                        custom_index: instance.custom_index,
+                        mask: instance.mask,
+                        blas_address: blas.handle,
+                    },
                 ));
 
                 instance_count += 1;
@@ -1013,6 +1015,7 @@ fn iter_buffers<'a, 'b>(
     scratch_buffer_blas_size: &mut u64,
     blas_storage: &mut Vec<BlasStore<'a>>,
     hub: &Hub,
+    ray_tracing_scratch_buffer_alignment: u32,
 ) -> Result<(), BuildAccelerationStructureError> {
     let mut triangle_entries =
         Vec::<hal::AccelerationStructureTriangles<dyn hal::DynBuffer>>::new();
@@ -1192,7 +1195,7 @@ fn iter_buffers<'a, 'b>(
             let scratch_buffer_offset = *scratch_buffer_blas_size;
             *scratch_buffer_blas_size += align_to(
                 blas.size_info.build_scratch_size as u32,
-                SCRATCH_BUFFER_ALIGNMENT,
+                ray_tracing_scratch_buffer_alignment,
             ) as u64;
 
             blas_storage.push(BlasStore {
diff --git a/wgpu-core/src/command/render.rs b/wgpu-core/src/command/render.rs
index ad3a5b6ca..7211edc71 100644
--- a/wgpu-core/src/command/render.rs
+++ b/wgpu-core/src/command/render.rs
@@ -1391,20 +1391,10 @@ impl Global {
                     None
                 };
 
-            arc_desc.timestamp_writes = if let Some(tw) = desc.timestamp_writes {
-                let query_set = query_sets.get(tw.query_set).get()?;
-                query_set.same_device(device)?;
-
-                device.require_features(wgt::Features::TIMESTAMP_QUERY)?;
-
-                Some(ArcPassTimestampWrites {
-                    query_set,
-                    beginning_of_pass_write_index: tw.beginning_of_pass_write_index,
-                    end_of_pass_write_index: tw.end_of_pass_write_index,
-                })
-            } else {
-                None
-            };
+            arc_desc.timestamp_writes = desc
+                .timestamp_writes
+                .map(|tw| Global::validate_pass_timestamp_writes(device, &query_sets, tw))
+                .transpose()?;
 
             arc_desc.occlusion_query_set =
                 if let Some(occlusion_query_set) = desc.occlusion_query_set {
diff --git a/wgpu-core/src/device/global.rs b/wgpu-core/src/device/global.rs
index f9b6794b4..42881f1bd 100644
--- a/wgpu-core/src/device/global.rs
+++ b/wgpu-core/src/device/global.rs
@@ -2102,6 +2102,14 @@ impl Global {
         }
     }
 
+    pub fn device_unregister_device_lost_closure(&self, device_id: DeviceId) {
+        let device = self.hub.devices.get(device_id);
+        let closure = device.device_lost_closure.lock().take();
+        if let Some(closure) = closure {
+            closure.call(DeviceLostReason::ReplacedCallback, "".to_string());
+        }
+    }
+
     pub fn device_destroy(&self, device_id: DeviceId) {
         api_log!("Device::destroy {device_id:?}");
 
diff --git a/wgpu-core/src/device/ray_tracing.rs b/wgpu-core/src/device/ray_tracing.rs
index baab42091..12afc7e6a 100644
--- a/wgpu-core/src/device/ray_tracing.rs
+++ b/wgpu-core/src/device/ray_tracing.rs
@@ -12,7 +12,7 @@ use crate::{
     global::Global,
     id::{self, BlasId, TlasId},
     lock::RwLock,
-    ray_tracing::{get_raw_tlas_instance_size, CreateBlasError, CreateTlasError},
+    ray_tracing::{CreateBlasError, CreateTlasError},
     resource, LabelHelpers,
 };
 use hal::AccelerationStructureTriangleIndices;
@@ -135,7 +135,7 @@ impl Device {
         .map_err(DeviceError::from_hal)?;
 
         let instance_buffer_size =
-            get_raw_tlas_instance_size(self.backend()) * desc.max_instances.max(1) as usize;
+            self.alignments.raw_tlas_instance_size * desc.max_instances.max(1) as usize;
         let instance_buffer = unsafe {
             self.raw().create_buffer(&hal::BufferDescriptor {
                 label: Some("(wgpu-core) instances_buffer"),
diff --git a/wgpu-core/src/instance.rs b/wgpu-core/src/instance.rs
index 5c8aef3cd..89b37c78a 100644
--- a/wgpu-core/src/instance.rs
+++ b/wgpu-core/src/instance.rs
@@ -653,7 +653,7 @@ pub enum GetSurfaceSupportError {
 
 #[derive(Clone, Debug, Error)]
 #[cfg_attr(feature = "serde", derive(serde::Serialize, serde::Deserialize))]
-/// Error when requesting a device from the adaptor
+/// Error when requesting a device from the adapter
 #[non_exhaustive]
 pub enum RequestDeviceError {
     #[error(transparent)]
diff --git a/wgpu-core/src/ray_tracing.rs b/wgpu-core/src/ray_tracing.rs
index a094b3530..9f4a11946 100644
--- a/wgpu-core/src/ray_tracing.rs
+++ b/wgpu-core/src/ray_tracing.rs
@@ -13,8 +13,8 @@ use crate::{
     id::{BlasId, BufferId, TlasId},
     resource::CreateBufferError,
 };
-use std::{mem::size_of, sync::Arc};
-use std::{num::NonZeroU64, slice};
+use std::num::NonZeroU64;
+use std::sync::Arc;
 
 use crate::resource::{Blas, ResourceErrorIdent, Tlas};
 use thiserror::Error;
@@ -276,48 +276,3 @@ pub struct TraceTlasPackage {
     pub instances: Vec<Option<TraceTlasInstance>>,
     pub lowest_unmodified: u32,
 }
-
-pub(crate) fn get_raw_tlas_instance_size(backend: wgt::Backend) -> usize {
-    // TODO: this should be provided by the backend
-    match backend {
-        wgt::Backend::Empty => 0,
-        wgt::Backend::Vulkan => 64,
-        _ => unimplemented!(),
-    }
-}
-
-#[derive(Clone)]
-#[repr(C)]
-struct RawTlasInstance {
-    transform: [f32; 12],
-    custom_index_and_mask: u32,
-    shader_binding_table_record_offset_and_flags: u32,
-    acceleration_structure_reference: u64,
-}
-
-pub(crate) fn tlas_instance_into_bytes(
-    instance: &TlasInstance,
-    blas_address: u64,
-    backend: wgt::Backend,
-) -> Vec<u8> {
-    // TODO: get the device to do this
-    match backend {
-        wgt::Backend::Empty => vec![],
-        wgt::Backend::Vulkan => {
-            const MAX_U24: u32 = (1u32 << 24u32) - 1u32;
-            let temp = RawTlasInstance {
-                transform: *instance.transform,
-                custom_index_and_mask: (instance.custom_index & MAX_U24)
-                    | (u32::from(instance.mask) << 24),
-                shader_binding_table_record_offset_and_flags: 0,
-                acceleration_structure_reference: blas_address,
-            };
-            let temp: *const _ = &temp;
-            unsafe {
-                slice::from_raw_parts::<u8>(temp.cast::<u8>(), size_of::<RawTlasInstance>())
-                    .to_vec()
-            }
-        }
-        _ => unimplemented!(),
-    }
-}
diff --git a/wgpu-core/src/validation.rs b/wgpu-core/src/validation.rs
index 52273b4f8..b1c005190 100644
--- a/wgpu-core/src/validation.rs
+++ b/wgpu-core/src/validation.rs
@@ -20,6 +20,37 @@ enum ResourceType {
     AccelerationStructure,
 }
 
+#[derive(Clone, Debug)]
+pub enum BindingTypeName {
+    Buffer,
+    Texture,
+    Sampler,
+    AccelerationStructure,
+}
+
+impl From<&ResourceType> for BindingTypeName {
+    fn from(ty: &ResourceType) -> BindingTypeName {
+        match ty {
+            ResourceType::Buffer { .. } => BindingTypeName::Buffer,
+            ResourceType::Texture { .. } => BindingTypeName::Texture,
+            ResourceType::Sampler { .. } => BindingTypeName::Sampler,
+            ResourceType::AccelerationStructure { .. } => BindingTypeName::AccelerationStructure,
+        }
+    }
+}
+
+impl From<&BindingType> for BindingTypeName {
+    fn from(ty: &BindingType) -> BindingTypeName {
+        match ty {
+            BindingType::Buffer { .. } => BindingTypeName::Buffer,
+            BindingType::Texture { .. } => BindingTypeName::Texture,
+            BindingType::StorageTexture { .. } => BindingTypeName::Texture,
+            BindingType::Sampler { .. } => BindingTypeName::Sampler,
+            BindingType::AccelerationStructure { .. } => BindingTypeName::AccelerationStructure,
+        }
+    }
+}
+
 #[derive(Debug)]
 struct Resource {
     #[allow(unused)]
@@ -140,13 +171,20 @@ pub enum BindingError {
     Missing,
     #[error("Visibility flags don't include the shader stage")]
     Invisible,
-    #[error("Type on the shader side does not match the pipeline binding")]
-    WrongType,
+    #[error(
+        "Type on the shader side ({shader:?}) does not match the pipeline binding ({binding:?})"
+    )]
+    WrongType {
+        binding: BindingTypeName,
+        shader: BindingTypeName,
+    },
     #[error("Storage class {binding:?} doesn't match the shader {shader:?}")]
     WrongAddressSpace {
         binding: naga::AddressSpace,
         shader: naga::AddressSpace,
     },
+    #[error("Address space {space:?} is not a valid Buffer address space")]
+    WrongBufferAddressSpace { space: naga::AddressSpace },
     #[error("Buffer structure size {buffer_size}, added to one element of an unbound array, if it's the last field, ended up greater than the given `min_binding_size`, which is {min_binding_size}")]
     WrongBufferSize {
         buffer_size: wgt::BufferSize,
@@ -378,7 +416,12 @@ impl Resource {
                         }
                         min_binding_size
                     }
-                    _ => return Err(BindingError::WrongType),
+                    _ => {
+                        return Err(BindingError::WrongType {
+                            binding: (&entry.ty).into(),
+                            shader: (&self.ty).into(),
+                        })
+                    }
                 };
                 match min_size {
                     Some(non_zero) if non_zero < size => {
@@ -396,7 +439,12 @@ impl Resource {
                         return Err(BindingError::WrongSamplerComparison);
                     }
                 }
-                _ => return Err(BindingError::WrongType),
+                _ => {
+                    return Err(BindingError::WrongType {
+                        binding: (&entry.ty).into(),
+                        shader: (&self.ty).into(),
+                    })
+                }
             },
             ResourceType::Texture {
                 dim,
@@ -478,7 +526,12 @@ impl Resource {
                             access: naga_access,
                         }
                     }
-                    _ => return Err(BindingError::WrongType),
+                    _ => {
+                        return Err(BindingError::WrongType {
+                            binding: (&entry.ty).into(),
+                            shader: (&self.ty).into(),
+                        })
+                    }
                 };
                 if class != expected_class {
                     return Err(BindingError::WrongTextureClass {
@@ -487,7 +540,15 @@ impl Resource {
                     });
                 }
             }
-            ResourceType::AccelerationStructure => (),
+            ResourceType::AccelerationStructure => match entry.ty {
+                BindingType::AccelerationStructure => (),
+                _ => {
+                    return Err(BindingError::WrongType {
+                        binding: (&entry.ty).into(),
+                        shader: (&self.ty).into(),
+                    })
+                }
+            },
         };
 
         Ok(())
@@ -504,7 +565,7 @@ impl Resource {
                     naga::AddressSpace::Storage { access } => wgt::BufferBindingType::Storage {
                         read_only: access == naga::StorageAccess::LOAD,
                     },
-                    _ => return Err(BindingError::WrongType),
+                    _ => return Err(BindingError::WrongBufferAddressSpace { space: self.class }),
                 },
                 has_dynamic_offset: false,
                 min_binding_size: Some(size),
diff --git a/wgpu-core/src/weak_vec.rs b/wgpu-core/src/weak_vec.rs
index b645cc714..f6034c6e9 100644
--- a/wgpu-core/src/weak_vec.rs
+++ b/wgpu-core/src/weak_vec.rs
@@ -47,7 +47,7 @@ impl<T> WeakVec<T> {
         }
         if let Some(i) = self.empty_slots.pop() {
             self.inner[i] = Some(value);
-            self.scan_slots_on_next_push = false;
+            self.scan_slots_on_next_push = self.empty_slots.is_empty();
         } else {
             self.inner.push(Some(value));
             self.scan_slots_on_next_push = self.inner.len() == self.inner.capacity();
diff --git a/wgpu-hal/src/dx12/adapter.rs b/wgpu-hal/src/dx12/adapter.rs
index 45d69f558..d037adc96 100644
--- a/wgpu-hal/src/dx12/adapter.rs
+++ b/wgpu-hal/src/dx12/adapter.rs
@@ -293,6 +293,7 @@ impl super::Adapter {
             }
         };
 
+        // these should always be available on d3d12
         let mut features = wgt::Features::empty()
             | wgt::Features::DEPTH_CLIP_CONTROL
             | wgt::Features::DEPTH32FLOAT_STENCIL8
@@ -315,7 +316,8 @@ impl super::Adapter {
             | wgt::Features::SHADER_PRIMITIVE_INDEX
             | wgt::Features::RG11B10UFLOAT_RENDERABLE
             | wgt::Features::DUAL_SOURCE_BLENDING
-            | wgt::Features::TEXTURE_FORMAT_NV12;
+            | wgt::Features::TEXTURE_FORMAT_NV12
+            | wgt::Features::FLOAT32_FILTERABLE;
 
         //TODO: in order to expose this, we need to run a compute shader
         // that extract the necessary statistics out of the D3D12 result.
@@ -403,9 +405,6 @@ impl super::Adapter {
             atomic_int64_on_typed_resource_supported,
         );
 
-        // float32-filterable should always be available on d3d12
-        features.set(wgt::Features::FLOAT32_FILTERABLE, true);
-
         // TODO: Determine if IPresentationManager is supported
         let presentation_timer = auxil::dxgi::time::PresentationTimer::new_dxgi();
 
@@ -522,6 +521,8 @@ impl super::Adapter {
                     // Direct3D correctly bounds-checks all array accesses:
                     // https://microsoft.github.io/DirectX-Specs/d3d/archive/D3D11_3_FunctionalSpec.htm#18.6.8.2%20Device%20Memory%20Reads
                     uniform_bounds_check_alignment: wgt::BufferSize::new(1).unwrap(),
+                    raw_tlas_instance_size: 0,
+                    ray_tracing_scratch_buffer_alignment: 0,
                 },
                 downlevel,
             },
diff --git a/wgpu-hal/src/dx12/device.rs b/wgpu-hal/src/dx12/device.rs
index b2b9629ed..9e089bbd1 100644
--- a/wgpu-hal/src/dx12/device.rs
+++ b/wgpu-hal/src/dx12/device.rs
@@ -21,6 +21,7 @@ use super::{conv, descriptor, D3D12Lib};
 use crate::{
     auxil::{self, dxgi::result::HResult},
     dx12::{borrow_optional_interface_temporarily, shader_compilation, Event},
+    TlasInstance,
 };
 
 // this has to match Naga's HLSL backend, and also needs to be null-terminated
@@ -1939,4 +1940,8 @@ impl crate::Device for super::Device {
             total_reserved_bytes: upstream.total_reserved_bytes,
         })
     }
+
+    fn tlas_instance_to_bytes(&self, _instance: TlasInstance) -> Vec<u8> {
+        todo!()
+    }
 }
diff --git a/wgpu-hal/src/dynamic/device.rs b/wgpu-hal/src/dynamic/device.rs
index f38d9fd32..28ef35287 100644
--- a/wgpu-hal/src/dynamic/device.rs
+++ b/wgpu-hal/src/dynamic/device.rs
@@ -5,7 +5,7 @@ use crate::{
     GetAccelerationStructureBuildSizesDescriptor, Label, MemoryRange, PipelineCacheDescriptor,
     PipelineCacheError, PipelineError, PipelineLayoutDescriptor, RenderPipelineDescriptor,
     SamplerDescriptor, ShaderError, ShaderInput, ShaderModuleDescriptor, TextureDescriptor,
-    TextureViewDescriptor,
+    TextureViewDescriptor, TlasInstance,
 };
 
 use super::{
@@ -158,6 +158,7 @@ pub trait DynDevice: DynResource {
         &self,
         acceleration_structure: Box<dyn DynAccelerationStructure>,
     );
+    fn tlas_instance_to_bytes(&self, instance: TlasInstance) -> Vec<u8>;
 
     fn get_internal_counters(&self) -> wgt::HalCounters;
     fn generate_allocator_report(&self) -> Option<wgt::AllocatorReport>;
@@ -520,6 +521,10 @@ impl<D: Device + DynResource> DynDevice for D {
         unsafe { D::destroy_acceleration_structure(self, acceleration_structure.unbox()) }
     }
 
+    fn tlas_instance_to_bytes(&self, instance: TlasInstance) -> Vec<u8> {
+        D::tlas_instance_to_bytes(self, instance)
+    }
+
     fn get_internal_counters(&self) -> wgt::HalCounters {
         D::get_internal_counters(self)
     }
diff --git a/wgpu-hal/src/empty.rs b/wgpu-hal/src/empty.rs
index e74c91c07..d4fc7c312 100644
--- a/wgpu-hal/src/empty.rs
+++ b/wgpu-hal/src/empty.rs
@@ -1,5 +1,6 @@
 #![allow(unused_variables)]
 
+use crate::TlasInstance;
 use std::ops::Range;
 
 #[derive(Clone, Debug)]
@@ -306,6 +307,10 @@ impl crate::Device for Context {
     }
     unsafe fn destroy_acceleration_structure(&self, _acceleration_structure: Resource) {}
 
+    fn tlas_instance_to_bytes(&self, instance: TlasInstance) -> Vec<u8> {
+        vec![]
+    }
+
     fn get_internal_counters(&self) -> wgt::HalCounters {
         Default::default()
     }
diff --git a/wgpu-hal/src/gles/adapter.rs b/wgpu-hal/src/gles/adapter.rs
index a654215d2..032ad1b46 100644
--- a/wgpu-hal/src/gles/adapter.rs
+++ b/wgpu-hal/src/gles/adapter.rs
@@ -851,6 +851,8 @@ impl super::Adapter {
                     // being, provide 1 as the value here, to cause as little
                     // trouble as possible.
                     uniform_bounds_check_alignment: wgt::BufferSize::new(1).unwrap(),
+                    raw_tlas_instance_size: 0,
+                    ray_tracing_scratch_buffer_alignment: 0,
                 },
             },
         })
diff --git a/wgpu-hal/src/gles/device.rs b/wgpu-hal/src/gles/device.rs
index 365ac7d72..acc0388d7 100644
--- a/wgpu-hal/src/gles/device.rs
+++ b/wgpu-hal/src/gles/device.rs
@@ -8,7 +8,7 @@ use std::{
     sync::{Arc, Mutex},
 };
 
-use crate::AtomicFenceValue;
+use crate::{AtomicFenceValue, TlasInstance};
 use arrayvec::ArrayVec;
 use std::sync::atomic::Ordering;
 
@@ -826,7 +826,7 @@ impl crate::Device for super::Device {
                                 0,
                                 format_desc.external,
                                 format_desc.data_type,
-                                None,
+                                glow::PixelUnpackData::Slice(None),
                             );
                             width = max(1, width / 2);
                             height = max(1, height / 2);
@@ -846,7 +846,7 @@ impl crate::Device for super::Device {
                                 0,
                                 format_desc.external,
                                 format_desc.data_type,
-                                None,
+                                glow::PixelUnpackData::Slice(None),
                             );
                             width = max(1, width / 2);
                             height = max(1, height / 2);
@@ -899,7 +899,7 @@ impl crate::Device for super::Device {
                                     0,
                                     format_desc.external,
                                     format_desc.data_type,
-                                    None,
+                                    glow::PixelUnpackData::Slice(None),
                                 );
                             }
                             width = max(1, width / 2);
@@ -918,7 +918,7 @@ impl crate::Device for super::Device {
                                 0,
                                 format_desc.external,
                                 format_desc.data_type,
-                                None,
+                                glow::PixelUnpackData::Slice(None),
                             );
                             width = max(1, width / 2);
                             height = max(1, height / 2);
@@ -1633,6 +1633,10 @@ impl crate::Device for super::Device {
     ) {
     }
 
+    fn tlas_instance_to_bytes(&self, _instance: TlasInstance) -> Vec<u8> {
+        unimplemented!()
+    }
+
     fn get_internal_counters(&self) -> wgt::HalCounters {
         self.counters.clone()
     }
diff --git a/wgpu-hal/src/gles/queue.rs b/wgpu-hal/src/gles/queue.rs
index fc106eb23..f1bd47920 100644
--- a/wgpu-hal/src/gles/queue.rs
+++ b/wgpu-hal/src/gles/queue.rs
@@ -744,7 +744,7 @@ impl super::Queue {
                             buffer_data = src.data.as_ref().unwrap().lock().unwrap();
                             let src_data =
                                 &buffer_data.as_slice()[copy.buffer_layout.offset as usize..];
-                            glow::PixelUnpackData::Slice(src_data)
+                            glow::PixelUnpackData::Slice(Some(src_data))
                         }
                     };
                     if is_layered_target(dst_target) {
@@ -889,7 +889,7 @@ impl super::Queue {
                         None => {
                             buffer_data = dst.data.as_ref().unwrap().lock().unwrap();
                             let dst_data = &mut buffer_data.as_mut_slice()[offset as usize..];
-                            glow::PixelPackData::Slice(dst_data)
+                            glow::PixelPackData::Slice(Some(dst_data))
                         }
                     };
                     unsafe {
@@ -1549,7 +1549,7 @@ impl super::Queue {
                 unsafe {
                     gl.bind_image_texture(
                         slot,
-                        binding.raw,
+                        Some(binding.raw),
                         binding.mip_level as i32,
                         binding.array_layer.is_none(),
                         binding.array_layer.unwrap_or_default() as i32,
diff --git a/wgpu-hal/src/lib.rs b/wgpu-hal/src/lib.rs
index ad4b458cd..15fa9a287 100644
--- a/wgpu-hal/src/lib.rs
+++ b/wgpu-hal/src/lib.rs
@@ -971,6 +971,7 @@ pub trait Device: WasmNotSendSync {
         &self,
         acceleration_structure: <Self::A as Api>::AccelerationStructure,
     );
+    fn tlas_instance_to_bytes(&self, instance: TlasInstance) -> Vec<u8>;
 
     fn get_internal_counters(&self) -> wgt::HalCounters;
 
@@ -1771,6 +1772,12 @@ pub struct Alignments {
     /// [`Uniform`]: wgt::BufferBindingType::Uniform
     /// [size]: BufferBinding::size
     pub uniform_bounds_check_alignment: wgt::BufferSize,
+
+    /// The size of the raw TLAS instance
+    pub raw_tlas_instance_size: usize,
+
+    /// What the scratch buffer for building an acceleration structure must be aligned to
+    pub ray_tracing_scratch_buffer_alignment: u32,
 }
 
 #[derive(Clone, Debug)]
@@ -2519,3 +2526,11 @@ bitflags::bitflags! {
 pub struct AccelerationStructureBarrier {
     pub usage: Range<AccelerationStructureUses>,
 }
+
+#[derive(Debug, Copy, Clone)]
+pub struct TlasInstance {
+    pub transform: [f32; 12],
+    pub custom_index: u32,
+    pub mask: u8,
+    pub blas_address: u64,
+}
diff --git a/wgpu-hal/src/metal/adapter.rs b/wgpu-hal/src/metal/adapter.rs
index e7db97a1f..fa7622a37 100644
--- a/wgpu-hal/src/metal/adapter.rs
+++ b/wgpu-hal/src/metal/adapter.rs
@@ -1001,6 +1001,8 @@ impl super::PrivateCapabilities {
                 // Metal Shading Language it generates, so from `wgpu_hal`'s
                 // users' point of view, references are tightly checked.
                 uniform_bounds_check_alignment: wgt::BufferSize::new(1).unwrap(),
+                raw_tlas_instance_size: 0,
+                ray_tracing_scratch_buffer_alignment: 0,
             },
             downlevel,
         }
diff --git a/wgpu-hal/src/metal/device.rs b/wgpu-hal/src/metal/device.rs
index 81887aac7..745ad598a 100644
--- a/wgpu-hal/src/metal/device.rs
+++ b/wgpu-hal/src/metal/device.rs
@@ -8,6 +8,7 @@ use std::{
 
 use super::conv;
 use crate::auxil::map_naga_stage;
+use crate::TlasInstance;
 
 type DeviceResult<T> = Result<T, crate::DeviceError>;
 
@@ -1426,6 +1427,10 @@ impl crate::Device for super::Device {
         unimplemented!()
     }
 
+    fn tlas_instance_to_bytes(&self, _instance: TlasInstance) -> Vec<u8> {
+        unimplemented!()
+    }
+
     fn get_internal_counters(&self) -> wgt::HalCounters {
         self.counters.clone()
     }
diff --git a/wgpu-hal/src/vulkan/adapter.rs b/wgpu-hal/src/vulkan/adapter.rs
index 2916928e4..6e88e03b8 100644
--- a/wgpu-hal/src/vulkan/adapter.rs
+++ b/wgpu-hal/src/vulkan/adapter.rs
@@ -1140,6 +1140,13 @@ impl PhysicalDeviceProperties {
                 };
                 wgt::BufferSize::new(alignment).unwrap()
             },
+            raw_tlas_instance_size: 64,
+            ray_tracing_scratch_buffer_alignment: self.acceleration_structure.map_or(
+                0,
+                |acceleration_structure| {
+                    acceleration_structure.min_acceleration_structure_scratch_offset_alignment
+                },
+            ),
         }
     }
 }
diff --git a/wgpu-hal/src/vulkan/device.rs b/wgpu-hal/src/vulkan/device.rs
index fe72ec40f..114307c9e 100644
--- a/wgpu-hal/src/vulkan/device.rs
+++ b/wgpu-hal/src/vulkan/device.rs
@@ -1,16 +1,17 @@
-use super::conv;
+use super::{conv, RawTlasInstance};
 
 use arrayvec::ArrayVec;
 use ash::{khr, vk};
 use parking_lot::Mutex;
 
+use crate::TlasInstance;
 use std::{
     borrow::Cow,
     collections::{hash_map::Entry, BTreeMap},
     ffi::{CStr, CString},
-    mem::MaybeUninit,
+    mem::{self, size_of, MaybeUninit},
     num::NonZeroU32,
-    ptr,
+    ptr, slice,
     sync::Arc,
 };
 
@@ -1662,7 +1663,7 @@ impl crate::Device for super::Device {
                     // Additional safety docs from unstable slice_assume_init_mut
                     // SAFETY: similar to safety notes for `slice_get_ref`, but we have a
                     // mutable reference which is also guaranteed to be valid for writes.
-                    unsafe { std::mem::transmute::<&mut [MaybeUninit<T>], &mut [T]>(to_init) }
+                    unsafe { mem::transmute::<&mut [MaybeUninit<T>], &mut [T]>(to_init) }
                 };
                 (Self { remainder }, init)
             }
@@ -2557,6 +2558,21 @@ impl crate::Device for super::Device {
 
         self.counters.clone()
     }
+
+    fn tlas_instance_to_bytes(&self, instance: TlasInstance) -> Vec<u8> {
+        const MAX_U24: u32 = (1u32 << 24u32) - 1u32;
+        let temp = RawTlasInstance {
+            transform: instance.transform,
+            custom_index_and_mask: (instance.custom_index & MAX_U24)
+                | (u32::from(instance.mask) << 24),
+            shader_binding_table_record_offset_and_flags: 0,
+            acceleration_structure_reference: instance.blas_address,
+        };
+        let temp: *const _ = &temp;
+        unsafe {
+            slice::from_raw_parts::<u8>(temp.cast::<u8>(), size_of::<RawTlasInstance>()).to_vec()
+        }
+    }
 }
 
 impl super::DeviceShared {
diff --git a/wgpu-hal/src/vulkan/mod.rs b/wgpu-hal/src/vulkan/mod.rs
index ff3c865fc..c4c5f70d7 100644
--- a/wgpu-hal/src/vulkan/mod.rs
+++ b/wgpu-hal/src/vulkan/mod.rs
@@ -1423,3 +1423,12 @@ fn get_lost_err() -> crate::DeviceError {
     #[allow(unreachable_code)]
     crate::DeviceError::Lost
 }
+
+#[derive(Clone)]
+#[repr(C)]
+struct RawTlasInstance {
+    transform: [f32; 12],
+    custom_index_and_mask: u32,
+    shader_binding_table_record_offset_and_flags: u32,
+    acceleration_structure_reference: u64,
+}


Test patch:
diff --git a/naga/tests/in/glsl/samplers.frag b/naga/tests/in/glsl/samplers.frag
index efb1c0ceb..53cc25d1c 100644
--- a/naga/tests/in/glsl/samplers.frag
+++ b/naga/tests/in/glsl/samplers.frag
@@ -45,17 +45,18 @@ void testTex1D(in float coord) {
     int levels = textureQueryLevels(sampler1D(tex1D, samp));
     vec4 c;
     c = texture(sampler1D(tex1D, samp), coord);
-    c = texture(sampler1D(tex1D, samp), coord, 2.0);
+    /* 1 dimensional textures are not supported in the WGSL specification with `textureSampleBias()` */
+    // c = texture(sampler1D(tex1D, samp), coord, 2.0);
     c = textureGrad(sampler1D(tex1D, samp), coord, 4.0, 4.0);
     c = textureGradOffset(sampler1D(tex1D, samp), coord, 4.0, 4.0, 5);
     c = textureLod(sampler1D(tex1D, samp), coord, 3.0);
     c = textureLodOffset(sampler1D(tex1D, samp), coord, 3.0, 5);
     c = textureOffset(sampler1D(tex1D, samp), coord, 5);
-    c = textureOffset(sampler1D(tex1D, samp), coord, 5, 2.0);
+    // c = textureOffset(sampler1D(tex1D, samp), coord, 5, 2.0);
     c = textureProj(sampler1D(tex1D, samp), vec2(coord, 6.0));
     c = textureProj(sampler1D(tex1D, samp), vec4(coord, 0.0, 0.0, 6.0));
-    c = textureProj(sampler1D(tex1D, samp), vec2(coord, 6.0), 2.0);
-    c = textureProj(sampler1D(tex1D, samp), vec4(coord, 0.0, 0.0, 6.0), 2.0);
+    // c = textureProj(sampler1D(tex1D, samp), vec2(coord, 6.0), 2.0);
+    // c = textureProj(sampler1D(tex1D, samp), vec4(coord, 0.0, 0.0, 6.0), 2.0);
     c = textureProjGrad(sampler1D(tex1D, samp), vec2(coord, 6.0), 4.0, 4.0);
     c = textureProjGrad(sampler1D(tex1D, samp), vec4(coord, 0.0, 0.0, 6.0), 4.0, 4.0);
     c = textureProjGradOffset(sampler1D(tex1D, samp), vec2(coord, 6.0), 4.0, 4.0, 5);
@@ -66,8 +67,8 @@ void testTex1D(in float coord) {
     c = textureProjLodOffset(sampler1D(tex1D, samp), vec4(coord, 0.0, 0.0, 6.0), 3.0, 5);
     c = textureProjOffset(sampler1D(tex1D, samp), vec2(coord, 6.0), 5);
     c = textureProjOffset(sampler1D(tex1D, samp), vec4(coord, 0.0, 0.0, 6.0), 5);
-    c = textureProjOffset(sampler1D(tex1D, samp), vec2(coord, 6.0), 5, 2.0);
-    c = textureProjOffset(sampler1D(tex1D, samp), vec4(coord, 0.0, 0.0, 6.0), 5, 2.0);
+    // c = textureProjOffset(sampler1D(tex1D, samp), vec2(coord, 6.0), 5, 2.0);
+    // c = textureProjOffset(sampler1D(tex1D, samp), vec4(coord, 0.0, 0.0, 6.0), 5, 2.0);
     c = texelFetch(sampler1D(tex1D, samp), int(coord), 3);
     c = texelFetchOffset(sampler1D(tex1D, samp), int(coord), 3, 5);
 }
@@ -101,13 +102,13 @@ void testTex1DArray(in vec2 coord) {
     int levels = textureQueryLevels(sampler1DArray(tex1DArray, samp));
     vec4 c;
     c = texture(sampler1DArray(tex1DArray, samp), coord);
-    c = texture(sampler1DArray(tex1DArray, samp), coord, 2.0);
+    // c = texture(sampler1DArray(tex1DArray, samp), coord, 2.0);
     c = textureGrad(sampler1DArray(tex1DArray, samp), coord, 4.0, 4.0);
     c = textureGradOffset(sampler1DArray(tex1DArray, samp), coord, 4.0, 4.0, 5);
     c = textureLod(sampler1DArray(tex1DArray, samp), coord, 3.0);
     c = textureLodOffset(sampler1DArray(tex1DArray, samp), coord, 3.0, 5);
     c = textureOffset(sampler1DArray(tex1DArray, samp), coord, 5);
-    c = textureOffset(sampler1DArray(tex1DArray, samp), coord, 5, 2.0);
+    // c = textureOffset(sampler1DArray(tex1DArray, samp), coord, 5, 2.0);
     c = texelFetch(sampler1DArray(tex1DArray, samp), ivec2(coord), 3);
     c = texelFetchOffset(sampler1DArray(tex1DArray, samp), ivec2(coord), 3, 5);
 }
diff --git a/naga/tests/in/image.wgsl b/naga/tests/in/image.wgsl
index e78480118..50eabcb6f 100644
--- a/naga/tests/in/image.wgsl
+++ b/naga/tests/in/image.wgsl
@@ -185,7 +185,9 @@ fn gather() -> @location(0) vec4<f32> {
 @fragment
 fn depth_no_comparison() -> @location(0) vec4<f32> {
     let tc = vec2<f32>(0.5);
+    let level = 1;
     let s2d = textureSample(image_2d_depth, sampler_reg, tc);
     let s2d_gather = textureGather(image_2d_depth, sampler_reg, tc);
-    return s2d + s2d_gather;
+    let s2d_level = textureSampleLevel(image_2d_depth, sampler_reg, tc, level);
+    return s2d + s2d_gather + s2d_level;
 }
diff --git a/naga/tests/out/hlsl/image.hlsl b/naga/tests/out/hlsl/image.hlsl
index a81625b05..84e9f6d70 100644
--- a/naga/tests/out/hlsl/image.hlsl
+++ b/naga/tests/out/hlsl/image.hlsl
@@ -369,5 +369,6 @@ float4 depth_no_comparison() : SV_Target0
     float2 tc_3 = (0.5).xx;
     float s2d_1 = image_2d_depth.Sample(sampler_reg, tc_3);
     float4 s2d_gather = image_2d_depth.Gather(sampler_reg, tc_3);
-    return ((s2d_1).xxxx + s2d_gather);
+    float s2d_level = image_2d_depth.SampleLevel(sampler_reg, tc_3, 1);
+    return (((s2d_1).xxxx + s2d_gather) + (s2d_level).xxxx);
 }
diff --git a/naga/tests/out/ir/diagnostic-filter.compact.ron b/naga/tests/out/ir/diagnostic-filter.compact.ron
index 4294b06e9..315c8b342 100644
--- a/naga/tests/out/ir/diagnostic-filter.compact.ron
+++ b/naga/tests/out/ir/diagnostic-filter.compact.ron
@@ -44,14 +44,14 @@
         (
             inner: (
                 new_severity: Off,
-                triggering_rule: DerivativeUniformity,
+                triggering_rule: Standard(DerivativeUniformity),
             ),
             parent: None,
         ),
         (
             inner: (
                 new_severity: Warning,
-                triggering_rule: DerivativeUniformity,
+                triggering_rule: Standard(DerivativeUniformity),
             ),
             parent: Some(0),
         ),
diff --git a/naga/tests/out/ir/diagnostic-filter.ron b/naga/tests/out/ir/diagnostic-filter.ron
index 4294b06e9..315c8b342 100644
--- a/naga/tests/out/ir/diagnostic-filter.ron
+++ b/naga/tests/out/ir/diagnostic-filter.ron
@@ -44,14 +44,14 @@
         (
             inner: (
                 new_severity: Off,
-                triggering_rule: DerivativeUniformity,
+                triggering_rule: Standard(DerivativeUniformity),
             ),
             parent: None,
         ),
         (
             inner: (
                 new_severity: Warning,
-                triggering_rule: DerivativeUniformity,
+                triggering_rule: Standard(DerivativeUniformity),
             ),
             parent: Some(0),
         ),
diff --git a/naga/tests/out/msl/atomicCompareExchange.msl b/naga/tests/out/msl/atomicCompareExchange.msl
index 800b5b201..6655fad7e 100644
--- a/naga/tests/out/msl/atomicCompareExchange.msl
+++ b/naga/tests/out/msl/atomicCompareExchange.msl
@@ -76,9 +76,8 @@ kernel void test_atomic_compare_exchange_i32_(
     uint i = 0u;
     int old = {};
     bool exchanged = {};
-#define LOOP_IS_REACHABLE if (volatile bool unpredictable_jump_over_loop = true; unpredictable_jump_over_loop)
     bool loop_init = true;
-    LOOP_IS_REACHABLE while(true) {
+    while(true) {
         if (!loop_init) {
             uint _e27 = i;
             i = _e27 + 1u;
@@ -94,7 +93,7 @@ kernel void test_atomic_compare_exchange_i32_(
             int _e8 = metal::atomic_load_explicit(&arr_i32_.inner[_e6], metal::memory_order_relaxed);
             old = _e8;
             exchanged = false;
-            LOOP_IS_REACHABLE while(true) {
+            while(true) {
                 bool _e12 = exchanged;
                 if (!(_e12)) {
                 } else {
@@ -109,8 +108,11 @@ kernel void test_atomic_compare_exchange_i32_(
                     old = _e23.old_value;
                     exchanged = _e23.exchanged;
                 }
+#define LOOP_IS_BOUNDED { volatile bool unpredictable_break_from_loop = false; if (unpredictable_break_from_loop) break; }
+                LOOP_IS_BOUNDED
             }
         }
+        LOOP_IS_BOUNDED
     }
     return;
 }
@@ -123,7 +125,7 @@ kernel void test_atomic_compare_exchange_u32_(
     uint old_1 = {};
     bool exchanged_1 = {};
     bool loop_init_1 = true;
-    LOOP_IS_REACHABLE while(true) {
+    while(true) {
         if (!loop_init_1) {
             uint _e27 = i_1;
             i_1 = _e27 + 1u;
@@ -139,7 +141,7 @@ kernel void test_atomic_compare_exchange_u32_(
             uint _e8 = metal::atomic_load_explicit(&arr_u32_.inner[_e6], metal::memory_order_relaxed);
             old_1 = _e8;
             exchanged_1 = false;
-            LOOP_IS_REACHABLE while(true) {
+            while(true) {
                 bool _e12 = exchanged_1;
                 if (!(_e12)) {
                 } else {
@@ -154,8 +156,10 @@ kernel void test_atomic_compare_exchange_u32_(
                     old_1 = _e23.old_value;
                     exchanged_1 = _e23.exchanged;
                 }
+                LOOP_IS_BOUNDED
             }
         }
+        LOOP_IS_BOUNDED
     }
     return;
 }
diff --git a/naga/tests/out/msl/boids.msl b/naga/tests/out/msl/boids.msl
index 0dd520ac7..07acd7cf6 100644
--- a/naga/tests/out/msl/boids.msl
+++ b/naga/tests/out/msl/boids.msl
@@ -55,9 +55,8 @@ kernel void main_(
     vPos = _e8;
     metal::float2 _e14 = particlesSrc.particles[index].vel;
     vVel = _e14;
-#define LOOP_IS_REACHABLE if (volatile bool unpredictable_jump_over_loop = true; unpredictable_jump_over_loop)
     bool loop_init = true;
-    LOOP_IS_REACHABLE while(true) {
+    while(true) {
         if (!loop_init) {
             uint _e91 = i;
             i = _e91 + 1u;
@@ -106,6 +105,8 @@ kernel void main_(
             int _e88 = cVelCount;
             cVelCount = _e88 + 1;
         }
+#define LOOP_IS_BOUNDED { volatile bool unpredictable_break_from_loop = false; if (unpredictable_break_from_loop) break; }
+        LOOP_IS_BOUNDED
     }
     int _e94 = cMassCount;
     if (_e94 > 0) {
diff --git a/naga/tests/out/msl/break-if.msl b/naga/tests/out/msl/break-if.msl
index 3684f7222..4d3397234 100644
--- a/naga/tests/out/msl/break-if.msl
+++ b/naga/tests/out/msl/break-if.msl
@@ -7,15 +7,16 @@ using metal::uint;
 
 void breakIfEmpty(
 ) {
-#define LOOP_IS_REACHABLE if (volatile bool unpredictable_jump_over_loop = true; unpredictable_jump_over_loop)
     bool loop_init = true;
-    LOOP_IS_REACHABLE while(true) {
+    while(true) {
         if (!loop_init) {
             if (true) {
                 break;
             }
         }
         loop_init = false;
+#define LOOP_IS_BOUNDED { volatile bool unpredictable_break_from_loop = false; if (unpredictable_break_from_loop) break; }
+        LOOP_IS_BOUNDED
     }
     return;
 }
@@ -26,7 +27,7 @@ void breakIfEmptyBody(
     bool b = {};
     bool c = {};
     bool loop_init_1 = true;
-    LOOP_IS_REACHABLE while(true) {
+    while(true) {
         if (!loop_init_1) {
             b = a;
             bool _e2 = b;
@@ -37,6 +38,7 @@ void breakIfEmptyBody(
             }
         }
         loop_init_1 = false;
+        LOOP_IS_BOUNDED
     }
     return;
 }
@@ -47,7 +49,7 @@ void breakIf(
     bool d = {};
     bool e = {};
     bool loop_init_2 = true;
-    LOOP_IS_REACHABLE while(true) {
+    while(true) {
         if (!loop_init_2) {
             bool _e5 = e;
             if (a_1 == e) {
@@ -58,6 +60,7 @@ void breakIf(
         d = a_1;
         bool _e2 = d;
         e = a_1 != _e2;
+        LOOP_IS_BOUNDED
     }
     return;
 }
@@ -66,7 +69,7 @@ void breakIfSeparateVariable(
 ) {
     uint counter = 0u;
     bool loop_init_3 = true;
-    LOOP_IS_REACHABLE while(true) {
+    while(true) {
         if (!loop_init_3) {
             uint _e5 = counter;
             if (counter == 5u) {
@@ -76,6 +79,7 @@ void breakIfSeparateVariable(
         loop_init_3 = false;
         uint _e3 = counter;
         counter = _e3 + 1u;
+        LOOP_IS_BOUNDED
     }
     return;
 }
diff --git a/naga/tests/out/msl/collatz.msl b/naga/tests/out/msl/collatz.msl
index 1ae910de6..e282d13ab 100644
--- a/naga/tests/out/msl/collatz.msl
+++ b/naga/tests/out/msl/collatz.msl
@@ -19,8 +19,7 @@ uint collatz_iterations(
     uint n = {};
     uint i = 0u;
     n = n_base;
-#define LOOP_IS_REACHABLE if (volatile bool unpredictable_jump_over_loop = true; unpredictable_jump_over_loop)
-    LOOP_IS_REACHABLE while(true) {
+    while(true) {
         uint _e4 = n;
         if (_e4 > 1u) {
         } else {
@@ -38,6 +37,8 @@ uint collatz_iterations(
             uint _e20 = i;
             i = _e20 + 1u;
         }
+#define LOOP_IS_BOUNDED { volatile bool unpredictable_break_from_loop = false; if (unpredictable_break_from_loop) break; }
+        LOOP_IS_BOUNDED
     }
     uint _e23 = i;
     return _e23;
diff --git a/naga/tests/out/msl/control-flow.msl b/naga/tests/out/msl/control-flow.msl
index dbf75163a..1b35249f3 100644
--- a/naga/tests/out/msl/control-flow.msl
+++ b/naga/tests/out/msl/control-flow.msl
@@ -31,8 +31,7 @@ void switch_case_break(
 void loop_switch_continue(
     int x
 ) {
-#define LOOP_IS_REACHABLE if (volatile bool unpredictable_jump_over_loop = true; unpredictable_jump_over_loop)
-    LOOP_IS_REACHABLE while(true) {
+    while(true) {
         switch(x) {
             case 1: {
                 continue;
@@ -41,6 +40,8 @@ void loop_switch_continue(
                 break;
             }
         }
+#define LOOP_IS_BOUNDED { volatile bool unpredictable_break_from_loop = false; if (unpredictable_break_from_loop) break; }
+        LOOP_IS_BOUNDED
     }
     return;
 }
@@ -50,7 +51,7 @@ void loop_switch_continue_nesting(
     int y,
     int z
 ) {
-    LOOP_IS_REACHABLE while(true) {
+    while(true) {
         switch(x_1) {
             case 1: {
                 continue;
@@ -61,7 +62,7 @@ void loop_switch_continue_nesting(
                         continue;
                     }
                     default: {
-                        LOOP_IS_REACHABLE while(true) {
+                        while(true) {
                             switch(z) {
                                 case 1: {
                                     continue;
@@ -70,6 +71,7 @@ void loop_switch_continue_nesting(
                                     break;
                                 }
                             }
+                            LOOP_IS_BOUNDED
                         }
                         break;
                     }
@@ -85,8 +87,9 @@ void loop_switch_continue_nesting(
                 continue;
             }
         }
+        LOOP_IS_BOUNDED
     }
-    LOOP_IS_REACHABLE while(true) {
+    while(true) {
         switch(y) {
             case 1:
             default: {
@@ -98,6 +101,7 @@ void loop_switch_continue_nesting(
                 break;
             }
         }
+        LOOP_IS_BOUNDED
     }
     return;
 }
@@ -109,7 +113,7 @@ void loop_switch_omit_continue_variable_checks(
     int w
 ) {
     int pos_1 = 0;
-    LOOP_IS_REACHABLE while(true) {
+    while(true) {
         switch(x_2) {
             case 1: {
                 pos_1 = 1;
@@ -119,8 +123,9 @@ void loop_switch_omit_continue_variable_checks(
                 break;
             }
         }
+        LOOP_IS_BOUNDED
     }
-    LOOP_IS_REACHABLE while(true) {
+    while(true) {
         switch(x_2) {
             case 1: {
                 break;
@@ -149,6 +154,7 @@ void loop_switch_omit_continue_variable_checks(
                 break;
             }
         }
+        LOOP_IS_BOUNDED
     }
     return;
 }
diff --git a/naga/tests/out/msl/do-while.msl b/naga/tests/out/msl/do-while.msl
index b093da1dc..2a883304d 100644
--- a/naga/tests/out/msl/do-while.msl
+++ b/naga/tests/out/msl/do-while.msl
@@ -8,9 +8,8 @@ using metal::uint;
 void fb1_(
     thread bool& cond
 ) {
-#define LOOP_IS_REACHABLE if (volatile bool unpredictable_jump_over_loop = true; unpredictable_jump_over_loop)
     bool loop_init = true;
-    LOOP_IS_REACHABLE while(true) {
+    while(true) {
         if (!loop_init) {
             bool _e1 = cond;
             if (!(cond)) {
@@ -19,6 +18,8 @@ void fb1_(
         }
         loop_init = false;
         continue;
+#define LOOP_IS_BOUNDED { volatile bool unpredictable_break_from_loop = false; if (unpredictable_break_from_loop) break; }
+        LOOP_IS_BOUNDED
     }
     return;
 }
diff --git a/naga/tests/out/msl/image.msl b/naga/tests/out/msl/image.msl
index 9f0a24300..114ed3655 100644
--- a/naga/tests/out/msl/image.msl
+++ b/naga/tests/out/msl/image.msl
@@ -265,5 +265,6 @@ fragment depth_no_comparisonOutput depth_no_comparison(
     metal::float2 tc_3 = metal::float2(0.5);
     float s2d_1 = image_2d_depth.sample(sampler_reg, tc_3);
     metal::float4 s2d_gather = image_2d_depth.gather(sampler_reg, tc_3);
-    return depth_no_comparisonOutput { metal::float4(s2d_1) + s2d_gather };
+    float s2d_level = image_2d_depth.sample(sampler_reg, tc_3, metal::level(1));
+    return depth_no_comparisonOutput { (metal::float4(s2d_1) + s2d_gather) + metal::float4(s2d_level) };
 }
diff --git a/naga/tests/out/msl/overrides-ray-query.msl b/naga/tests/out/msl/overrides-ray-query.msl
index f2ad45c98..3aa0ee035 100644
--- a/naga/tests/out/msl/overrides-ray-query.msl
+++ b/naga/tests/out/msl/overrides-ray-query.msl
@@ -33,14 +33,15 @@ kernel void main_(
     rq.intersector.force_opacity((desc.flags & 1) != 0 ? metal::raytracing::forced_opacity::opaque : (desc.flags & 2) != 0 ? metal::raytracing::forced_opacity::non_opaque : metal::raytracing::forced_opacity::none);
     rq.intersector.accept_any_intersection((desc.flags & 4) != 0);
     rq.intersection = rq.intersector.intersect(metal::raytracing::ray(desc.origin, desc.dir, desc.tmin, desc.tmax), acc_struct, desc.cull_mask);    rq.ready = true;
-#define LOOP_IS_REACHABLE if (volatile bool unpredictable_jump_over_loop = true; unpredictable_jump_over_loop)
-    LOOP_IS_REACHABLE while(true) {
+    while(true) {
         bool _e31 = rq.ready;
         rq.ready = false;
         if (_e31) {
         } else {
             break;
         }
+#define LOOP_IS_BOUNDED { volatile bool unpredictable_break_from_loop = false; if (unpredictable_break_from_loop) break; }
+        LOOP_IS_BOUNDED
     }
     return;
 }
diff --git a/naga/tests/out/msl/ray-query.msl b/naga/tests/out/msl/ray-query.msl
index 129ad108a..e5e44d0b5 100644
--- a/naga/tests/out/msl/ray-query.msl
+++ b/naga/tests/out/msl/ray-query.msl
@@ -53,14 +53,15 @@ RayIntersection query_loop(
     rq.intersector.force_opacity((_e8.flags & 1) != 0 ? metal::raytracing::forced_opacity::opaque : (_e8.flags & 2) != 0 ? metal::raytracing::forced_opacity::non_opaque : metal::raytracing::forced_opacity::none);
     rq.intersector.accept_any_intersection((_e8.flags & 4) != 0);
     rq.intersection = rq.intersector.intersect(metal::raytracing::ray(_e8.origin, _e8.dir, _e8.tmin, _e8.tmax), acs, _e8.cull_mask);    rq.ready = true;
-#define LOOP_IS_REACHABLE if (volatile bool unpredictable_jump_over_loop = true; unpredictable_jump_over_loop)
-    LOOP_IS_REACHABLE while(true) {
+    while(true) {
         bool _e9 = rq.ready;
         rq.ready = false;
         if (_e9) {
         } else {
             break;
         }
+#define LOOP_IS_BOUNDED { volatile bool unpredictable_break_from_loop = false; if (unpredictable_break_from_loop) break; }
+        LOOP_IS_BOUNDED
     }
     return RayIntersection {_map_intersection_type(rq.intersection.type), rq.intersection.distance, rq.intersection.user_instance_id, rq.intersection.instance_id, {}, rq.intersection.geometry_id, rq.intersection.primitive_id, rq.intersection.triangle_barycentric_coord, rq.intersection.triangle_front_facing, {}, rq.intersection.object_to_world_transform, rq.intersection.world_to_object_transform};
 }
diff --git a/naga/tests/out/msl/shadow.msl b/naga/tests/out/msl/shadow.msl
index f8aeef9d4..18cc84211 100644
--- a/naga/tests/out/msl/shadow.msl
+++ b/naga/tests/out/msl/shadow.msl
@@ -100,9 +100,8 @@ fragment fs_mainOutput fs_main(
     metal::float3 color = c_ambient;
     uint i = 0u;
     metal::float3 normal_1 = metal::normalize(in.world_normal);
-#define LOOP_IS_REACHABLE if (volatile bool unpredictable_jump_over_loop = true; unpredictable_jump_over_loop)
     bool loop_init = true;
-    LOOP_IS_REACHABLE while(true) {
+    while(true) {
         if (!loop_init) {
             uint _e40 = i;
             i = _e40 + 1u;
@@ -124,6 +123,8 @@ fragment fs_mainOutput fs_main(
             metal::float3 _e37 = color;
             color = _e37 + ((_e23 * diffuse) * light.color.xyz);
         }
+#define LOOP_IS_BOUNDED { volatile bool unpredictable_break_from_loop = false; if (unpredictable_break_from_loop) break; }
+        LOOP_IS_BOUNDED
     }
     metal::float3 _e42 = color;
     metal::float4 _e47 = u_entity.color;
@@ -152,7 +153,7 @@ fragment fs_main_without_storageOutput fs_main_without_storage(
     uint i_1 = 0u;
     metal::float3 normal_2 = metal::normalize(in_1.world_normal);
     bool loop_init_1 = true;
-    LOOP_IS_REACHABLE while(true) {
+    while(true) {
         if (!loop_init_1) {
             uint _e40 = i_1;
             i_1 = _e40 + 1u;
@@ -174,6 +175,7 @@ fragment fs_main_without_storageOutput fs_main_without_storage(
             metal::float3 _e37 = color_1;
             color_1 = _e37 + ((_e23 * diffuse_1) * light_1.color.xyz);
         }
+        LOOP_IS_BOUNDED
     }
     metal::float3 _e42 = color_1;
     metal::float4 _e47 = u_entity.color;
diff --git a/naga/tests/out/spv/globals.spvasm b/naga/tests/out/spv/globals.spvasm
index 4aa6a10ad..31e836531 100644
--- a/naga/tests/out/spv/globals.spvasm
+++ b/naga/tests/out/spv/globals.spvasm
@@ -44,10 +44,14 @@ OpDecorate %45 DescriptorSet 0
 OpDecorate %45 Binding 6
 OpDecorate %46 Block
 OpMemberDecorate %46 0 Offset 0
+OpMemberDecorate %46 0 ColMajor
+OpMemberDecorate %46 0 MatrixStride 16
 OpDecorate %48 DescriptorSet 0
 OpDecorate %48 Binding 7
 OpDecorate %49 Block
 OpMemberDecorate %49 0 Offset 0
+OpMemberDecorate %49 0 ColMajor
+OpMemberDecorate %49 0 MatrixStride 8
 OpDecorate %116 BuiltIn LocalInvocationId
 %2 = OpTypeVoid
 %3 = OpTypeBool
diff --git a/naga/tests/out/spv/image.spvasm b/naga/tests/out/spv/image.spvasm
index 974f29e16..5ec2925dc 100644
--- a/naga/tests/out/spv/image.spvasm
+++ b/naga/tests/out/spv/image.spvasm
@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.1
 ; Generator: rspirv
-; Bound: 520
+; Bound: 526
 OpCapability Shader
 OpCapability Image1D
 OpCapability Sampled1D
@@ -687,8 +687,14 @@ OpBranch %512
 %515 = OpCompositeExtract  %7  %514 0
 %516 = OpSampledImage  %428  %511 %510
 %517 = OpImageGather  %23  %516 %280 %198
-%518 = OpCompositeConstruct  %23  %515 %515 %515 %515
-%519 = OpFAdd  %23  %518 %517
-OpStore %508 %519
+%518 = OpSampledImage  %428  %511 %510
+%520 = OpConvertSToF  %7  %29
+%519 = OpImageSampleExplicitLod  %23  %518 %280 Lod %520
+%521 = OpCompositeExtract  %7  %519 0
+%522 = OpCompositeConstruct  %23  %515 %515 %515 %515
+%523 = OpFAdd  %23  %522 %517
+%524 = OpCompositeConstruct  %23  %521 %521 %521 %521
+%525 = OpFAdd  %23  %523 %524
+OpStore %508 %525
 OpReturn
 OpFunctionEnd
\ No newline at end of file
diff --git a/naga/tests/out/wgsl/image.wgsl b/naga/tests/out/wgsl/image.wgsl
index dddc79404..0c10eda85 100644
--- a/naga/tests/out/wgsl/image.wgsl
+++ b/naga/tests/out/wgsl/image.wgsl
@@ -235,5 +235,6 @@ fn depth_no_comparison() -> @location(0) vec4<f32> {
     const tc_3 = vec2(0.5f);
     let s2d_1 = textureSample(image_2d_depth, sampler_reg, tc_3);
     let s2d_gather = textureGather(image_2d_depth, sampler_reg, tc_3);
-    return (vec4(s2d_1) + s2d_gather);
+    let s2d_level = textureSampleLevel(image_2d_depth, sampler_reg, tc_3, 1i);
+    return ((vec4(s2d_1) + s2d_gather) + vec4(s2d_level));
 }
diff --git a/naga/tests/out/wgsl/samplers.frag.wgsl b/naga/tests/out/wgsl/samplers.frag.wgsl
index 1e7514804..c51f05305 100644
--- a/naga/tests/out/wgsl/samplers.frag.wgsl
+++ b/naga/tests/out/wgsl/samplers.frag.wgsl
@@ -49,115 +49,89 @@ fn testTex1D(coord: f32) {
     let _e30 = coord_1;
     let _e31 = textureSample(tex1D, samp, _e30);
     c = _e31;
-    let _e34 = coord_1;
-    let _e36 = textureSampleBias(tex1D, samp, _e34, 2f);
-    c = _e36;
-    let _e40 = coord_1;
-    let _e43 = textureSampleGrad(tex1D, samp, _e40, 4f, 4f);
-    c = _e43;
-    let _e48 = coord_1;
-    let _e52 = textureSampleGrad(tex1D, samp, _e48, 4f, 4f, 5i);
+    let _e35 = coord_1;
+    let _e38 = textureSampleGrad(tex1D, samp, _e35, 4f, 4f);
+    c = _e38;
+    let _e43 = coord_1;
+    let _e47 = textureSampleGrad(tex1D, samp, _e43, 4f, 4f, 5i);
+    c = _e47;
+    let _e50 = coord_1;
+    let _e52 = textureSampleLevel(tex1D, samp, _e50, 3f);
     c = _e52;
-    let _e55 = coord_1;
-    let _e57 = textureSampleLevel(tex1D, samp, _e55, 3f);
-    c = _e57;
-    let _e61 = coord_1;
-    let _e64 = textureSampleLevel(tex1D, samp, _e61, 3f, 5i);
+    let _e56 = coord_1;
+    let _e59 = textureSampleLevel(tex1D, samp, _e56, 3f, 5i);
+    c = _e59;
+    let _e62 = coord_1;
+    let _e64 = textureSample(tex1D, samp, _e62, 5i);
     c = _e64;
-    let _e67 = coord_1;
-    let _e69 = textureSample(tex1D, samp, _e67, 5i);
-    c = _e69;
-    let _e73 = coord_1;
-    let _e76 = textureSampleBias(tex1D, samp, _e73, 2f, 5i);
-    c = _e76;
-    let _e77 = coord_1;
+    let _e65 = coord_1;
+    let _e68 = coord_1;
+    let _e70 = vec2<f32>(_e68, 6f);
+    let _e74 = textureSample(tex1D, samp, (_e70.x / _e70.y));
+    c = _e74;
+    let _e75 = coord_1;
     let _e80 = coord_1;
-    let _e82 = vec2<f32>(_e80, 6f);
-    let _e86 = textureSample(tex1D, samp, (_e82.x / _e82.y));
-    c = _e86;
-    let _e87 = coord_1;
-    let _e92 = coord_1;
-    let _e96 = vec4<f32>(_e92, 0f, 0f, 6f);
-    let _e102 = textureSample(tex1D, samp, (_e96.xyz / vec3(_e96.w)).x);
-    c = _e102;
-    let _e103 = coord_1;
-    let _e107 = coord_1;
-    let _e109 = vec2<f32>(_e107, 6f);
-    let _e114 = textureSampleBias(tex1D, samp, (_e109.x / _e109.y), 2f);
-    c = _e114;
-    let _e115 = coord_1;
-    let _e121 = coord_1;
-    let _e125 = vec4<f32>(_e121, 0f, 0f, 6f);
-    let _e132 = textureSampleBias(tex1D, samp, (_e125.xyz / vec3(_e125.w)).x, 2f);
-    c = _e132;
-    let _e133 = coord_1;
-    let _e138 = coord_1;
-    let _e140 = vec2<f32>(_e138, 6f);
-    let _e146 = textureSampleGrad(tex1D, samp, (_e140.x / _e140.y), 4f, 4f);
-    c = _e146;
-    let _e147 = coord_1;
-    let _e154 = coord_1;
-    let _e158 = vec4<f32>(_e154, 0f, 0f, 6f);
-    let _e166 = textureSampleGrad(tex1D, samp, (_e158.xyz / vec3(_e158.w)).x, 4f, 4f);
-    c = _e166;
+    let _e84 = vec4<f32>(_e80, 0f, 0f, 6f);
+    let _e90 = textureSample(tex1D, samp, (_e84.xyz / vec3(_e84.w)).x);
+    c = _e90;
+    let _e91 = coord_1;
+    let _e96 = coord_1;
+    let _e98 = vec2<f32>(_e96, 6f);
+    let _e104 = textureSampleGrad(tex1D, samp, (_e98.x / _e98.y), 4f, 4f);
+    c = _e104;
+    let _e105 = coord_1;
+    let _e112 = coord_1;
+    let _e116 = vec4<f32>(_e112, 0f, 0f, 6f);
+    let _e124 = textureSampleGrad(tex1D, samp, (_e116.xyz / vec3(_e116.w)).x, 4f, 4f);
+    c = _e124;
+    let _e125 = coord_1;
+    let _e131 = coord_1;
+    let _e133 = vec2<f32>(_e131, 6f);
+    let _e140 = textureSampleGrad(tex1D, samp, (_e133.x / _e133.y), 4f, 4f, 5i);
+    c = _e140;
+    let _e141 = coord_1;
+    let _e149 = coord_1;
+    let _e153 = vec4<f32>(_e149, 0f, 0f, 6f);
+    let _e162 = textureSampleGrad(tex1D, samp, (_e153.xyz / vec3(_e153.w)).x, 4f, 4f, 5i);
+    c = _e162;
+    let _e163 = coord_1;
     let _e167 = coord_1;
-    let _e173 = coord_1;
-    let _e175 = vec2<f32>(_e173, 6f);
-    let _e182 = textureSampleGrad(tex1D, samp, (_e175.x / _e175.y), 4f, 4f, 5i);
-    c = _e182;
-    let _e183 = coord_1;
-    let _e191 = coord_1;
-    let _e195 = vec4<f32>(_e191, 0f, 0f, 6f);
-    let _e204 = textureSampleGrad(tex1D, samp, (_e195.xyz / vec3(_e195.w)).x, 4f, 4f, 5i);
-    c = _e204;
-    let _e205 = coord_1;
-    let _e209 = coord_1;
-    let _e211 = vec2<f32>(_e209, 6f);
-    let _e216 = textureSampleLevel(tex1D, samp, (_e211.x / _e211.y), 3f);
-    c = _e216;
-    let _e217 = coord_1;
-    let _e223 = coord_1;
-    let _e227 = vec4<f32>(_e223, 0f, 0f, 6f);
-    let _e234 = textureSampleLevel(tex1D, samp, (_e227.xyz / vec3(_e227.w)).x, 3f);
-    c = _e234;
-    let _e235 = coord_1;
-    let _e240 = coord_1;
-    let _e242 = vec2<f32>(_e240, 6f);
-    let _e248 = textureSampleLevel(tex1D, samp, (_e242.x / _e242.y), 3f, 5i);
-    c = _e248;
-    let _e249 = coord_1;
-    let _e256 = coord_1;
-    let _e260 = vec4<f32>(_e256, 0f, 0f, 6f);
-    let _e268 = textureSampleLevel(tex1D, samp, (_e260.xyz / vec3(_e260.w)).x, 3f, 5i);
-    c = _e268;
-    let _e269 = coord_1;
-    let _e273 = coord_1;
-    let _e275 = vec2<f32>(_e273, 6f);
-    let _e280 = textureSample(tex1D, samp, (_e275.x / _e275.y), 5i);
-    c = _e280;
-    let _e281 = coord_1;
-    let _e287 = coord_1;
-    let _e291 = vec4<f32>(_e287, 0f, 0f, 6f);
-    let _e298 = textureSample(tex1D, samp, (_e291.xyz / vec3(_e291.w)).x, 5i);
-    c = _e298;
-    let _e299 = coord_1;
-    let _e304 = coord_1;
-    let _e306 = vec2<f32>(_e304, 6f);
-    let _e312 = textureSampleBias(tex1D, samp, (_e306.x / _e306.y), 2f, 5i);
-    c = _e312;
-    let _e313 = coord_1;
-    let _e320 = coord_1;
-    let _e324 = vec4<f32>(_e320, 0f, 0f, 6f);
-    let _e332 = textureSampleBias(tex1D, samp, (_e324.xyz / vec3(_e324.w)).x, 2f, 5i);
-    c = _e332;
-    let _e333 = coord_1;
-    let _e336 = coord_1;
-    let _e339 = textureLoad(tex1D, i32(_e336), 3i);
-    c = _e339;
-    let _e340 = coord_1;
-    let _e344 = coord_1;
-    let _e348 = textureLoad(tex1D, i32(_e344), 3i);
-    c = _e348;
+    let _e169 = vec2<f32>(_e167, 6f);
+    let _e174 = textureSampleLevel(tex1D, samp, (_e169.x / _e169.y), 3f);
+    c = _e174;
+    let _e175 = coord_1;
+    let _e181 = coord_1;
+    let _e185 = vec4<f32>(_e181, 0f, 0f, 6f);
+    let _e192 = textureSampleLevel(tex1D, samp, (_e185.xyz / vec3(_e185.w)).x, 3f);
+    c = _e192;
+    let _e193 = coord_1;
+    let _e198 = coord_1;
+    let _e200 = vec2<f32>(_e198, 6f);
+    let _e206 = textureSampleLevel(tex1D, samp, (_e200.x / _e200.y), 3f, 5i);
+    c = _e206;
+    let _e207 = coord_1;
+    let _e214 = coord_1;
+    let _e218 = vec4<f32>(_e214, 0f, 0f, 6f);
+    let _e226 = textureSampleLevel(tex1D, samp, (_e218.xyz / vec3(_e218.w)).x, 3f, 5i);
+    c = _e226;
+    let _e227 = coord_1;
+    let _e231 = coord_1;
+    let _e233 = vec2<f32>(_e231, 6f);
+    let _e238 = textureSample(tex1D, samp, (_e233.x / _e233.y), 5i);
+    c = _e238;
+    let _e239 = coord_1;
+    let _e245 = coord_1;
+    let _e249 = vec4<f32>(_e245, 0f, 0f, 6f);
+    let _e256 = textureSample(tex1D, samp, (_e249.xyz / vec3(_e249.w)).x, 5i);
+    c = _e256;
+    let _e257 = coord_1;
+    let _e260 = coord_1;
+    let _e263 = textureLoad(tex1D, i32(_e260), 3i);
+    c = _e263;
+    let _e264 = coord_1;
+    let _e268 = coord_1;
+    let _e272 = textureLoad(tex1D, i32(_e268), 3i);
+    c = _e272;
     return;
 }
 
@@ -176,37 +150,31 @@ fn testTex1DArray(coord_2: vec2<f32>) {
     let _e32 = coord_3;
     let _e36 = textureSample(tex1DArray, samp, _e32.x, i32(_e32.y));
     c_1 = _e36;
-    let _e39 = coord_3;
-    let _e44 = textureSampleBias(tex1DArray, samp, _e39.x, i32(_e39.y), 2f);
-    c_1 = _e44;
-    let _e48 = coord_3;
-    let _e54 = textureSampleGrad(tex1DArray, samp, _e48.x, i32(_e48.y), 4f, 4f);
-    c_1 = _e54;
-    let _e59 = coord_3;
-    let _e66 = textureSampleGrad(tex1DArray, samp, _e59.x, i32(_e59.y), 4f, 4f, 5i);
+    let _e40 = coord_3;
+    let _e46 = textureSampleGrad(tex1DArray, samp, _e40.x, i32(_e40.y), 4f, 4f);
+    c_1 = _e46;
+    let _e51 = coord_3;
+    let _e58 = textureSampleGrad(tex1DArray, samp, _e51.x, i32(_e51.y), 4f, 4f, 5i);
+    c_1 = _e58;
+    let _e61 = coord_3;
+    let _e66 = textureSampleLevel(tex1DArray, samp, _e61.x, i32(_e61.y), 3f);
     c_1 = _e66;
-    let _e69 = coord_3;
-    let _e74 = textureSampleLevel(tex1DArray, samp, _e69.x, i32(_e69.y), 3f);
-    c_1 = _e74;
-    let _e78 = coord_3;
-    let _e84 = textureSampleLevel(tex1DArray, samp, _e78.x, i32(_e78.y), 3f, 5i);
+    let _e70 = coord_3;
+    let _e76 = textureSampleLevel(tex1DArray, samp, _e70.x, i32(_e70.y), 3f, 5i);
+    c_1 = _e76;
+    let _e79 = coord_3;
+    let _e84 = textureSample(tex1DArray, samp, _e79.x, i32(_e79.y), 5i);
     c_1 = _e84;
-    let _e87 = coord_3;
-    let _e92 = textureSample(tex1DArray, samp, _e87.x, i32(_e87.y), 5i);
-    c_1 = _e92;
-    let _e96 = coord_3;
-    let _e102 = textureSampleBias(tex1DArray, samp, _e96.x, i32(_e96.y), 2f, 5i);
-    c_1 = _e102;
-    let _e103 = coord_3;
-    let _e106 = coord_3;
-    let _e107 = vec2<i32>(_e106);
-    let _e111 = textureLoad(tex1DArray, _e107.x, _e107.y, 3i);
-    c_1 = _e111;
-    let _e112 = coord_3;
-    let _e116 = coord_3;
-    let _e117 = vec2<i32>(_e116);
-    let _e122 = textureLoad(tex1DArray, _e117.x, _e117.y, 3i);
-    c_1 = _e122;
+    let _e85 = coord_3;
+    let _e88 = coord_3;
+    let _e89 = vec2<i32>(_e88);
+    let _e93 = textureLoad(tex1DArray, _e89.x, _e89.y, 3i);
+    c_1 = _e93;
+    let _e94 = coord_3;
+    let _e98 = coord_3;
+    let _e99 = vec2<i32>(_e98);
+    let _e104 = textureLoad(tex1DArray, _e99.x, _e99.y, 3i);
+    c_1 = _e104;
     return;
 }
 
diff --git a/tests/tests/push_constants.rs b/tests/tests/push_constants.rs
index 905578d53..714f1b918 100644
--- a/tests/tests/push_constants.rs
+++ b/tests/tests/push_constants.rs
@@ -1,5 +1,8 @@
+use std::mem::size_of;
 use std::num::NonZeroU64;
 
+use wgpu::util::RenderEncoder;
+use wgpu::*;
 use wgpu_test::{gpu_test, GpuTestConfiguration, TestParameters, TestingContext};
 
 /// We want to test that partial updates to push constants work as expected.
@@ -153,3 +156,219 @@ async fn partial_update_test(ctx: TestingContext) {
     // second 4 floats the first update
     assert_eq!(floats, [1.0, 2.0, 3.0, 4.0, 1.0, 5.0, 3.0, 4.0]);
 }
+#[gpu_test]
+static RENDER_PASS_TEST: GpuTestConfiguration = GpuTestConfiguration::new()
+    .parameters(
+        TestParameters::default()
+            .features(Features::PUSH_CONSTANTS | Features::VERTEX_WRITABLE_STORAGE)
+            .limits(wgpu::Limits {
+                max_push_constant_size: 64,
+                ..Default::default()
+            }),
+    )
+    .run_async(move |ctx| async move {
+        for use_render_bundle in [false, true] {
+            render_pass_test(&ctx, use_render_bundle).await;
+        }
+    });
+
+// This shader simply  moves the values from vector_constants and push_constants into the
+// result buffer.  It expects to be called 4 times (with vector_index in 0..4) with its
+// topology being PointList, so that each vertex shader call leads to exactly one fragment
+// call.
+const SHADER2: &str = "
+    const POSITION: vec4f = vec4f(0, 0, 0, 1);
+
+    struct PushConstants {
+        vertex_constants: vec4i,
+        fragment_constants: vec4i,
+    }
+
+    var<push_constant> push_constants: PushConstants;
+
+    @group(0) @binding(0) var<storage, read_write> result: array<i32>;
+
+    struct VertexOutput {
+        @builtin(position) position: vec4f,
+        @location(0) index: u32,
+    }
+
+    @vertex fn vertex(
+        @builtin(vertex_index) ix: u32,
+    ) -> VertexOutput {
+        result[ix] = push_constants.vertex_constants[ix];
+        return VertexOutput(POSITION, ix);
+    }
+
+    @fragment fn fragment(
+        @location(0) ix: u32,
+     ) -> @location(0) vec4f {
+        result[ix + 4u] = push_constants.fragment_constants[ix];
+        return vec4f();
+    }
+";
+
+async fn render_pass_test(ctx: &TestingContext, use_render_bundle: bool) {
+    let output_buffer = ctx.device.create_buffer(&BufferDescriptor {
+        label: Some("output buffer"),
+        size: 8 * size_of::<u32>() as BufferAddress,
+        usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
+        mapped_at_creation: false,
+    });
+
+    let cpu_buffer = ctx.device.create_buffer(&BufferDescriptor {
+        label: Some("cpu buffer"),
+        size: output_buffer.size(),
+        usage: BufferUsages::COPY_DST | BufferUsages::MAP_READ,
+        mapped_at_creation: false,
+    });
+
+    // We need an output texture, even though we're not ever going to look at it.
+    let output_texture = ctx.device.create_texture(&TextureDescriptor {
+        size: Extent3d {
+            width: 2,
+            height: 2,
+            depth_or_array_layers: 1,
+        },
+        mip_level_count: 1,
+        sample_count: 1,
+        dimension: TextureDimension::D2,
+        format: TextureFormat::Rgba8UnormSrgb,
+        usage: TextureUsages::RENDER_ATTACHMENT,
+        label: Some("Output Texture"),
+        view_formats: &[],
+    });
+    let output_texture_view = output_texture.create_view(&Default::default());
+
+    let shader = ctx.device.create_shader_module(ShaderModuleDescriptor {
+        label: Some("Shader"),
+        source: ShaderSource::Wgsl(SHADER2.into()),
+    });
+
+    let bind_group_layout = ctx
+        .device
+        .create_bind_group_layout(&BindGroupLayoutDescriptor {
+            label: None,
+            entries: &[BindGroupLayoutEntry {
+                binding: 0,
+                visibility: ShaderStages::VERTEX_FRAGMENT,
+                ty: BindingType::Buffer {
+                    ty: BufferBindingType::Storage { read_only: false },
+                    has_dynamic_offset: false,
+                    min_binding_size: None,
+                },
+                count: None,
+            }],
+        });
+
+    let render_pipeline_layout = ctx
+        .device
+        .create_pipeline_layout(&PipelineLayoutDescriptor {
+            bind_group_layouts: &[&bind_group_layout],
+            push_constant_ranges: &[PushConstantRange {
+                stages: ShaderStages::VERTEX_FRAGMENT,
+                range: 0..8 * size_of::<u32>() as u32,
+            }],
+            ..Default::default()
+        });
+
+    let pipeline = ctx
+        .device
+        .create_render_pipeline(&RenderPipelineDescriptor {
+            label: Some("Render Pipeline"),
+            layout: Some(&render_pipeline_layout),
+            vertex: VertexState {
+                module: &shader,
+                entry_point: None,
+                buffers: &[],
+                compilation_options: Default::default(),
+            },
+            fragment: Some(FragmentState {
+                module: &shader,
+                entry_point: None,
+                targets: &[Some(output_texture.format().into())],
+                compilation_options: Default::default(),
+            }),
+            primitive: PrimitiveState {
+                topology: PrimitiveTopology::PointList,
+                ..Default::default()
+            },
+            depth_stencil: None,
+            multisample: MultisampleState::default(),
+            multiview: None,
+            cache: None,
+        });
+
+    let render_pass_desc = RenderPassDescriptor {
+        label: Some("Render Pass"),
+        color_attachments: &[Some(RenderPassColorAttachment {
+            view: &output_texture_view,
+            resolve_target: None,
+            ops: Operations {
+                load: LoadOp::Clear(Color::default()),
+                store: StoreOp::Store,
+            },
+        })],
+        ..Default::default()
+    };
+
+    let bind_group = ctx.device.create_bind_group(&BindGroupDescriptor {
+        label: Some("bind group"),
+        layout: &pipeline.get_bind_group_layout(0),
+        entries: &[BindGroupEntry {
+            binding: 0,
+            resource: output_buffer.as_entire_binding(),
+        }],
+    });
+
+    let data: Vec<i32> = (0..8).map(|i| (i * i) - 1).collect();
+
+    fn do_encoding<'a>(
+        encoder: &mut dyn RenderEncoder<'a>,
+        pipeline: &'a RenderPipeline,
+        bind_group: &'a BindGroup,
+        data: &'a Vec<i32>,
+    ) {
+        let data_as_u8: &[u8] = bytemuck::cast_slice(data.as_slice());
+        encoder.set_pipeline(pipeline);
+        encoder.set_push_constants(ShaderStages::VERTEX_FRAGMENT, 0, data_as_u8);
+        encoder.set_bind_group(0, Some(bind_group), &[]);
+        encoder.draw(0..4, 0..1);
+    }
+
+    let mut command_encoder = ctx
+        .device
+        .create_command_encoder(&CommandEncoderDescriptor::default());
+    {
+        let mut render_pass = command_encoder.begin_render_pass(&render_pass_desc);
+        if use_render_bundle {
+            // Execute the commands in a render_bundle_encoder.
+            let mut render_bundle_encoder =
+                ctx.device
+                    .create_render_bundle_encoder(&RenderBundleEncoderDescriptor {
+                        color_formats: &[Some(output_texture.format())],
+                        sample_count: 1,
+                        ..RenderBundleEncoderDescriptor::default()
+                    });
+            do_encoding(&mut render_bundle_encoder, &pipeline, &bind_group, &data);
+            let render_bundle = render_bundle_encoder.finish(&RenderBundleDescriptor::default());
+            render_pass.execute_bundles([&render_bundle]);
+        } else {
+            // Execute the commands directly.
+            do_encoding(&mut render_pass, &pipeline, &bind_group, &data);
+        }
+    }
+    // Move the result to the cpu buffer, so that we can read them.
+    command_encoder.copy_buffer_to_buffer(&output_buffer, 0, &cpu_buffer, 0, output_buffer.size());
+    let command_buffer = command_encoder.finish();
+    ctx.queue.submit([command_buffer]);
+    cpu_buffer.slice(..).map_async(MapMode::Read, |_| ());
+    ctx.async_poll(wgpu::Maintain::wait())
+        .await
+        .panic_on_timeout();
+    let mapped_data = cpu_buffer.slice(..).get_mapped_range();
+    let result = bytemuck::cast_slice::<u8, i32>(&mapped_data).to_vec();
+    drop(mapped_data);
+    cpu_buffer.unmap();
+    assert_eq!(&result, &data);
+}