diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index eb17c9bf89b..6b676760dca 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -304,7 +304,7 @@ jobs: # Check with all compatible features cargo clippy --target ${{ matrix.target }} ${{ matrix.extra-flags }} -p wgpu-types --no-default-features --features strict_asserts,fragile-send-sync-non-atomic-wasm,serde,counters - cargo clippy --target ${{ matrix.target }} ${{ matrix.extra-flags }} -p naga --no-default-features --features dot-out,compact + cargo clippy --target ${{ matrix.target }} ${{ matrix.extra-flags }} -p naga --no-default-features --features dot-out cargo clippy --target ${{ matrix.target }} ${{ matrix.extra-flags }} -p wgpu-hal --no-default-features --features fragile-send-sync-non-atomic-wasm cargo clippy --target ${{ matrix.target }} ${{ matrix.extra-flags }} -p wgpu --no-default-features --features serde diff --git a/CHANGELOG.md b/CHANGELOG.md index 46051e53e7b..0b5405e6143 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -67,8 +67,7 @@ Bottom level categories: #### Naga -Naga now infers the correct binding layout when a resource appears only in an assignment to `_`. By @andyleiserson in [#7540](https://github.com/gfx-rs/wgpu/pull/7540). - +- Naga now infers the correct binding layout when a resource appears only in an assignment to `_`. By @andyleiserson in [#7540](https://github.com/gfx-rs/wgpu/pull/7540). - Implement `dot4U8Packed` and `dot4I8Packed` for all backends, using specialized intrinsics on SPIR-V, HSLS, and Metal if available, and polyfills everywhere else. By @robamler in [#7494](https://github.com/gfx-rs/wgpu/pull/7494), [#7574](https://github.com/gfx-rs/wgpu/pull/7574), and [#7653](https://github.com/gfx-rs/wgpu/pull/7653). - Add polyfilled `pack4x{I,U}8Clamped` built-ins to all backends and WGSL frontend. By @ErichDonGubler in [#7546](https://github.com/gfx-rs/wgpu/pull/7546). - Allow textureLoad's sample index arg to be unsigned. By @jimblandy in [#7625](https://github.com/gfx-rs/wgpu/pull/7625). @@ -78,6 +77,7 @@ Naga now infers the correct binding layout when a resource appears only in an as - Properly evaluate `abs(most negative abstract int)`. By @jimblandy in [#7507](https://github.com/gfx-rs/wgpu/pull/7507). - Generate vectorized code for `[un]pack4x{I,U}8[Clamp]` on SPIR-V and MSL 2.1+. By @robamler in [#7664](https://github.com/gfx-rs/wgpu/pull/7664). - Fix typing for `select`, which had issues particularly with a lack of automatic type conversion. By @ErichDonGubler in [#7572](https://github.com/gfx-rs/wgpu/pull/7572). +- Allow scalars as the first argument of the `distance` built-in function. By @bernhl in [#7530](https://github.com/gfx-rs/wgpu/pull/7530). #### DX12 @@ -102,6 +102,10 @@ Naga now infers the correct binding layout when a resource appears only in an as - `naga::back::hlsl::Writer::new` has a new `pipeline_options` argument. `hlsl::PipelineOptions::default()` can be passed as a default. The `shader_stage` and `entry_point` members of `pipeline_options` can be used to write only a single entry point when using the HLSL and MSL backends (GLSL and SPIR-V already had this functionality). The Metal and DX12 HALs now write only a single entry point when loading shaders. By @andyleiserson in [#7626](https://github.com/gfx-rs/wgpu/pull/7626). - Implemented `early_depth_test` for SPIR-V backend, enabling `SHADER_EARLY_DEPTH_TEST` for Vulkan. Additionally, fixed conservative depth optimizations when using `early_depth_test`. The syntax for forcing early depth tests is now `@early_depth_test(force)` instead of `@early_depth_test`. By @dzamkov in [#7676](https://github.com/gfx-rs/wgpu/pull/7676). - `ImplementedLanguageExtension::VARIANTS` is now implemented manually rather than derived using `strum` (allowing `strum` to become a dev-only dependency) so it is no longer a member of the `strum::VARIANTS` trait. Unless you are using this trait as a bound this should have no effect. +- Compaction changes, by @andyleiserson in [#7703](https://github.com/gfx-rs/wgpu/pull/7703): + - [`process_overrides`](https://docs.rs/naga/latest/naga/back/pipeline_constants/fn.process_overrides.html) now compacts the module to remove unused items. It is no longer necessary to supply values for overrides that are not used by the active entry point. + - The `compact` Cargo feature has been removed. It is no longer possible to exclude compaction support from the build. + - [`compact`](https://docs.rs/naga/latest/naga/compact/fn.compact.html) now has an additional argument that specifies whether to remove unused functions, globals, and named types and overrides. For the previous behavior, pass `KeepUnused::Yes`. #### D3D12 @@ -119,11 +123,6 @@ Naga now infers the correct binding layout when a resource appears only in an as - Added initial `no_std` support to `wgpu-hal`. By @bushrat011899 in [#7599](https://github.com/gfx-rs/wgpu/pull/7599) -### Bug Fixes - -#### Naga - -- Allow scalars as the first argument of the `distance` built-in function. By @bernhl in [#7530](https://github.com/gfx-rs/wgpu/pull/7530). ### Documentation diff --git a/benches/benches/wgpu-benchmark/shader.rs b/benches/benches/wgpu-benchmark/shader.rs index 3c77efd44a4..d54c566faf5 100644 --- a/benches/benches/wgpu-benchmark/shader.rs +++ b/benches/benches/wgpu-benchmark/shader.rs @@ -269,6 +269,8 @@ fn validation(c: &mut Criterion) { } fn compact(c: &mut Criterion) { + use naga::compact::{compact, KeepUnused}; + let mut inputs = get_wgsl_inputs(); inputs.validate(); @@ -279,7 +281,7 @@ fn compact(c: &mut Criterion) { group.bench_function("shader: compact", |b| { b.iter(|| { for input in &mut inputs.inner { - naga::compact::compact(input.module.as_mut().unwrap()); + compact(input.module.as_mut().unwrap(), KeepUnused::No); } }); }); diff --git a/cts_runner/test.lst b/cts_runner/test.lst index e91fcd25dde..947f8e3b333 100644 --- a/cts_runner/test.lst +++ b/cts_runner/test.lst @@ -7,7 +7,9 @@ webgpu:api,operation,command_buffer,copyBufferToBuffer:state_transitions:* webgpu:api,operation,command_buffer,copyBufferToBuffer:copy_order:* webgpu:api,operation,compute,basic:memcpy:* //FAIL: webgpu:api,operation,compute,basic:large_dispatch:* +webgpu:api,operation,compute_pipeline,overrides:* webgpu:api,operation,device,lost:* +webgpu:api,operation,render_pipeline,overrides:* webgpu:api,operation,rendering,basic:clear:* webgpu:api,operation,rendering,basic:fullscreen_quad:* //FAIL: webgpu:api,operation,rendering,basic:large_draw:* diff --git a/naga-cli/Cargo.toml b/naga-cli/Cargo.toml index 53ddae80ef8..1ffadd64f8d 100644 --- a/naga-cli/Cargo.toml +++ b/naga-cli/Cargo.toml @@ -25,7 +25,6 @@ test = false [dependencies] naga = { workspace = true, features = [ - "compact", "wgsl-in", "wgsl-out", "glsl-in", diff --git a/naga-cli/src/bin/naga.rs b/naga-cli/src/bin/naga.rs index 6f95e429f68..a1e2b4a0bd1 100644 --- a/naga-cli/src/bin/naga.rs +++ b/naga-cli/src/bin/naga.rs @@ -542,7 +542,7 @@ fn run() -> anyhow::Result<()> { write_output(&module, &info, ¶ms, before_compaction)?; } - naga::compact::compact(&mut module); + naga::compact::compact(&mut module, KeepUnused::No); // Re-validate the IR after compaction. match naga::valid::Validator::new(params.validation_flags, validation_caps) @@ -717,9 +717,13 @@ fn write_output( succeed, and it failed in a previous step", ))?; - let (module, info) = - naga::back::pipeline_constants::process_overrides(module, info, ¶ms.overrides) - .unwrap_pretty(); + let (module, info) = naga::back::pipeline_constants::process_overrides( + module, + info, + None, + ¶ms.overrides, + ) + .unwrap_pretty(); let pipeline_options = msl::PipelineOptions::default(); let (msl, _) = @@ -751,9 +755,13 @@ fn write_output( succeed, and it failed in a previous step", ))?; - let (module, info) = - naga::back::pipeline_constants::process_overrides(module, info, ¶ms.overrides) - .unwrap_pretty(); + let (module, info) = naga::back::pipeline_constants::process_overrides( + module, + info, + None, + ¶ms.overrides, + ) + .unwrap_pretty(); let spv = spv::write_vec(&module, &info, ¶ms.spv_out, pipeline_options).unwrap_pretty(); @@ -788,9 +796,13 @@ fn write_output( succeed, and it failed in a previous step", ))?; - let (module, info) = - naga::back::pipeline_constants::process_overrides(module, info, ¶ms.overrides) - .unwrap_pretty(); + let (module, info) = naga::back::pipeline_constants::process_overrides( + module, + info, + None, + ¶ms.overrides, + ) + .unwrap_pretty(); let mut buffer = String::new(); let mut writer = glsl::Writer::new( @@ -819,9 +831,13 @@ fn write_output( succeed, and it failed in a previous step", ))?; - let (module, info) = - naga::back::pipeline_constants::process_overrides(module, info, ¶ms.overrides) - .unwrap_pretty(); + let (module, info) = naga::back::pipeline_constants::process_overrides( + module, + info, + None, + ¶ms.overrides, + ) + .unwrap_pretty(); let mut buffer = String::new(); let pipeline_options = Default::default(); @@ -906,4 +922,4 @@ fn bulk_validate(args: Args, params: &Parameters) -> anyhow::Result<()> { } use codespan_reporting::term::termcolor::{ColorChoice, StandardStream}; -use naga::FastHashMap; +use naga::{compact::KeepUnused, FastHashMap}; diff --git a/naga/Cargo.toml b/naga/Cargo.toml index af58c64ce1a..c03b91330c1 100644 --- a/naga/Cargo.toml +++ b/naga/Cargo.toml @@ -58,7 +58,7 @@ arbitrary = [ ] spv-in = ["dep:petgraph", "petgraph/graphmap", "dep:spirv"] spv-out = ["dep:spirv"] -wgsl-in = ["dep:hexf-parse", "dep:unicode-ident", "compact"] +wgsl-in = ["dep:hexf-parse", "dep:unicode-ident"] wgsl-out = [] ## Enables outputting to HLSL (Microsoft's High-Level Shader Language). @@ -72,8 +72,6 @@ hlsl-out = [] ## If you want to enable HLSL output it regardless of the target platform, use `naga/hlsl-out`. hlsl-out-if-target-windows = [] -compact = [] - ## Enables colored output through codespan-reporting and termcolor. termcolor = ["codespan-reporting/termcolor"] diff --git a/naga/src/arena/handle_set.rs b/naga/src/arena/handle_set.rs index 239dc429346..62691ff6cc7 100644 --- a/naga/src/arena/handle_set.rs +++ b/naga/src/arena/handle_set.rs @@ -77,6 +77,11 @@ impl HandleSet { } } + /// Add all of the handles that can be included in this set. + pub fn add_all(&mut self) { + self.members.get_mut().set_all(); + } + pub fn contains(&self, handle: Handle) -> bool { self.members.contains(handle.index()) } @@ -85,6 +90,23 @@ impl HandleSet { pub fn iter(&self) -> impl '_ + Iterator> { self.members.iter().map(Handle::from_usize) } + + /// Removes and returns the numerically largest handle in the set, or `None` + /// if the set is empty. + pub fn pop(&mut self) -> Option> { + let members = core::mem::take(&mut self.members); + let mut vec = members.into_bit_vec(); + let result = vec.iter_mut().enumerate().rev().find_map(|(i, mut bit)| { + if *bit { + *bit = false; + Some(i) + } else { + None + } + }); + self.members = bit_set::BitSet::from_bit_vec(vec); + result.map(Handle::from_usize) + } } impl Default for HandleSet { diff --git a/naga/src/arena/mod.rs b/naga/src/arena/mod.rs index 4fbfe7cc60e..84abf92a75b 100644 --- a/naga/src/arena/mod.rs +++ b/naga/src/arena/mod.rs @@ -234,7 +234,6 @@ impl Arena { Ok(()) } - #[cfg(feature = "compact")] pub(crate) fn retain_mut

(&mut self, mut predicate: P) where P: FnMut(Handle, &mut T) -> bool, diff --git a/naga/src/arena/unique_arena.rs b/naga/src/arena/unique_arena.rs index a5db7d339b5..49d481b2c0e 100644 --- a/naga/src/arena/unique_arena.rs +++ b/naga/src/arena/unique_arena.rs @@ -72,7 +72,6 @@ impl UniqueArena { .unwrap_or(&Span::default()) } - #[cfg(feature = "compact")] pub(crate) fn drain_all(&mut self) -> UniqueArenaDrain { UniqueArenaDrain { inner_elts: self.set.drain(..), @@ -82,14 +81,12 @@ impl UniqueArena { } } -#[cfg(feature = "compact")] pub struct UniqueArenaDrain<'a, T> { inner_elts: indexmap::set::Drain<'a, T>, inner_spans: alloc::vec::Drain<'a, Span>, index: Index, } -#[cfg(feature = "compact")] impl Iterator for UniqueArenaDrain<'_, T> { type Item = (Handle, T, Span); diff --git a/naga/src/back/pipeline_constants.rs b/naga/src/back/pipeline_constants.rs index 0088c6eac3f..d2c60d14872 100644 --- a/naga/src/back/pipeline_constants.rs +++ b/naga/src/back/pipeline_constants.rs @@ -10,6 +10,8 @@ use thiserror::Error; use super::PipelineConstants; use crate::{ arena::HandleVec, + compact::{compact, KeepUnused}, + ir, proc::{ConstantEvaluator, ConstantEvaluatorError, Emitter}, valid::{Capabilities, ModuleInfo, ValidationError, ValidationFlags, Validator}, Arena, Block, Constant, Expression, Function, Handle, Literal, Module, Override, Range, Scalar, @@ -55,6 +57,7 @@ pub enum PipelineConstantError { pub fn process_overrides<'a>( module: &'a Module, module_info: &'a ModuleInfo, + entry_point: Option<(ir::ShaderStage, &str)>, pipeline_constants: &PipelineConstants, ) -> Result<(Cow<'a, Module>, Cow<'a, ModuleInfo>), PipelineConstantError> { if module.overrides.is_empty() { @@ -63,6 +66,16 @@ pub fn process_overrides<'a>( let mut module = module.clone(); + // If an entry point was specified, compact the module to remove anything + // not reachable from that entry point. This is necessary because we may not + // have values for overrides that are not reachable from the entry point. + if let Some((ep_stage, ep_name)) = entry_point { + module + .entry_points + .retain(|ep| ep.stage == ep_stage && ep.name == ep_name); + } + compact(&mut module, KeepUnused::No); + // A map from override handles to the handles of the constants // we've replaced them with. let mut override_map = HandleVec::with_capacity(module.overrides.len()); diff --git a/naga/src/compact/expressions.rs b/naga/src/compact/expressions.rs index b0a858786b6..0e0f59336f8 100644 --- a/naga/src/compact/expressions.rs +++ b/naga/src/compact/expressions.rs @@ -11,6 +11,9 @@ pub struct ExpressionTracer<'tracer> { /// The used map for `types`. pub types_used: &'tracer mut HandleSet, + /// The used map for global variables. + pub global_variables_used: &'tracer mut HandleSet, + /// The used map for `constants`. pub constants_used: &'tracer mut HandleSet, @@ -76,9 +79,7 @@ impl ExpressionTracer<'_> { // Expressions that do not contain handles that need to be traced. Ex::Literal(_) | Ex::FunctionArgument(_) - | Ex::GlobalVariable(_) | Ex::LocalVariable(_) - | Ex::CallResult(_) | Ex::SubgroupBallotResult | Ex::RayQueryProceedResult => {} @@ -134,6 +135,9 @@ impl ExpressionTracer<'_> { } => { self.expressions_used.insert(vector); } + Ex::GlobalVariable(handle) => { + self.global_variables_used.insert(handle); + } Ex::Load { pointer } => { self.expressions_used.insert(pointer); } @@ -233,6 +237,10 @@ impl ExpressionTracer<'_> { Ex::ArrayLength(expr) => { self.expressions_used.insert(expr); } + // `CallResult` expressions do contain a function handle, but any used + // `CallResult` expression should have an associated `ir::Statement::Call` + // that we will trace. + Ex::CallResult(_) => {} Ex::AtomicResult { ty, comparison: _ } | Ex::WorkGroupUniformLoadResult { ty } | Ex::SubgroupOperationResult { ty } => { @@ -267,9 +275,7 @@ impl ModuleMap { // Expressions that do not contain handles that need to be adjusted. Ex::Literal(_) | Ex::FunctionArgument(_) - | Ex::GlobalVariable(_) | Ex::LocalVariable(_) - | Ex::CallResult(_) | Ex::SubgroupBallotResult | Ex::RayQueryProceedResult => {} @@ -306,6 +312,7 @@ impl ModuleMap { ref mut vector, pattern: _, } => adjust(vector), + Ex::GlobalVariable(ref mut handle) => self.globals.adjust(handle), Ex::Load { ref mut pointer } => adjust(pointer), Ex::ImageSample { ref mut image, @@ -392,6 +399,9 @@ impl ModuleMap { kind: _, convert: _, } => adjust(expr), + Ex::CallResult(ref mut function) => { + self.functions.adjust(function); + } Ex::AtomicResult { ref mut ty, comparison: _, diff --git a/naga/src/compact/functions.rs b/naga/src/compact/functions.rs index 57f6f104f21..eb829aec80e 100644 --- a/naga/src/compact/functions.rs +++ b/naga/src/compact/functions.rs @@ -6,7 +6,10 @@ pub struct FunctionTracer<'a> { pub constants: &'a crate::Arena, pub overrides: &'a crate::Arena, + pub functions_pending: &'a mut HandleSet, + pub functions_used: &'a mut HandleSet, pub types_used: &'a mut HandleSet, + pub global_variables_used: &'a mut HandleSet, pub constants_used: &'a mut HandleSet, pub overrides_used: &'a mut HandleSet, pub global_expressions_used: &'a mut HandleSet, @@ -16,6 +19,13 @@ pub struct FunctionTracer<'a> { } impl FunctionTracer<'_> { + pub fn trace_call(&mut self, function: crate::Handle) { + if !self.functions_used.contains(function) { + self.functions_used.insert(function); + self.functions_pending.insert(function); + } + } + pub fn trace(&mut self) { for argument in self.function.arguments.iter() { self.types_used.insert(argument.ty); @@ -53,6 +63,7 @@ impl FunctionTracer<'_> { expressions: &self.function.expressions, types_used: self.types_used, + global_variables_used: self.global_variables_used, constants_used: self.constants_used, overrides_used: self.overrides_used, expressions_used: &mut self.expressions_used, @@ -105,6 +116,6 @@ impl FunctionMap { assert!(reuse.is_empty()); // Adjust statements. - self.adjust_body(function); + self.adjust_body(function, &module_map.functions); } } diff --git a/naga/src/compact/handle_set_map.rs b/naga/src/compact/handle_set_map.rs index 5e80b954f1b..7a174cd9171 100644 --- a/naga/src/compact/handle_set_map.rs +++ b/naga/src/compact/handle_set_map.rs @@ -4,19 +4,43 @@ use crate::arena::{Arena, Handle, HandleSet, Range}; type Index = crate::non_max_u32::NonMaxU32; -/// A map from old handle indices to new, compressed handle indices. -pub struct HandleMap { +/// A map keyed by handles. +/// +/// In most cases, this is used to map from old handle indices to new, +/// compressed handle indices. +#[derive(Debug)] +pub struct HandleMap { /// The indices assigned to handles in the compacted module. /// /// If `new_index[i]` is `Some(n)`, then `n` is the `Index` of the /// compacted `Handle` corresponding to the pre-compacted `Handle` /// whose index is `i`. - new_index: Vec>, + new_index: Vec>, /// This type is indexed by values of type `T`. as_keys: core::marker::PhantomData, } +impl HandleMap { + pub fn with_capacity(capacity: usize) -> Self { + Self { + new_index: Vec::with_capacity(capacity), + as_keys: core::marker::PhantomData, + } + } + + pub fn get(&self, handle: Handle) -> Option<&U> { + self.new_index.get(handle.index()).unwrap_or(&None).as_ref() + } + + pub fn insert(&mut self, handle: Handle, value: U) -> Option { + if self.new_index.len() <= handle.index() { + self.new_index.resize_with(handle.index() + 1, || None); + } + core::mem::replace(&mut self.new_index[handle.index()], Some(value)) + } +} + impl HandleMap { pub fn from_set(set: HandleSet) -> Self { let mut next_index = Index::new(0).unwrap(); diff --git a/naga/src/compact/mod.rs b/naga/src/compact/mod.rs index 061ffb67c47..c7b7bfb46d8 100644 --- a/naga/src/compact/mod.rs +++ b/naga/src/compact/mod.rs @@ -6,56 +6,75 @@ mod types; use alloc::vec::Vec; -use crate::arena::HandleSet; -use crate::{arena, compact::functions::FunctionTracer}; +use crate::{ + arena::{self, HandleSet}, + compact::functions::FunctionTracer, + ir, +}; use handle_set_map::HandleMap; #[cfg(test)] use alloc::{format, string::ToString}; -/// Remove unused types, expressions, and constants from `module`. +/// Configuration option for [`compact`]. See [`compact`] for details. +#[derive(Debug, Clone, Copy, PartialEq, Eq)] +pub enum KeepUnused { + No, + Yes, +} + +impl From for bool { + fn from(keep_unused: KeepUnused) -> Self { + match keep_unused { + KeepUnused::No => false, + KeepUnused::Yes => true, + } + } +} + +/// Remove most unused objects from `module`, which must be valid. /// -/// Assume that the following are used by definition: +/// Always removes the following unused objects: +/// - anonymous types, overrides, and constants +/// - abstract-typed constants +/// - expressions +/// +/// If `keep_unused` is `Yes`, the following are never considered unused, +/// otherwise, they will also be removed if unused: +/// - functions /// - global variables -/// - named constants and overrides +/// - named types and overrides +/// +/// The following are never removed: +/// - named constants with a concrete type /// - special types -/// - functions and entry points, and within those: +/// - entry points +/// - within an entry point or a used function: /// - arguments /// - local variables /// - named expressions /// -/// Given those assumptions, determine which types, constants, -/// overrides, and expressions (both function-local and global -/// constant expressions) are actually used, and remove the rest, -/// adjusting all handles as necessary. The result should always be a -/// module functionally identical to the original. -/// -/// This may be useful to apply to modules generated in the snapshot -/// tests. Our backends often generate temporary names based on handle -/// indices, which means that adding or removing unused arena entries -/// can affect the output even though they have no semantic effect. -/// Such meaningless changes add noise to snapshot diffs, making -/// accurate patch review difficult. Compacting the modules before -/// generating snapshots makes the output independent of unused arena -/// entries. +/// After removing items according to the rules above, all handles in the +/// remaining objects are adjusted as necessary. When `KeepUnused` is `Yes`, the +/// resulting module should have all the named objects (except abstract-typed +/// constants) present in the original, and those objects should be functionally +/// identical. When `KeepUnused` is `No`, the resulting module should have the +/// entry points present in the original, and those entry points should be +/// functionally identical. /// /// # Panics /// /// If `module` would not pass validation, this may panic. -pub fn compact(module: &mut crate::Module) { +pub fn compact(module: &mut crate::Module, keep_unused: KeepUnused) { // The trickiest part of compaction is determining what is used and what is // not. Once we have computed that correctly, it's easy enough to call // `retain_mut` on each arena, drop unused elements, and fix up the handles // in what's left. // - // Some arenas' contents are considered used by definition, like - // `Module::global_variables` and `Module::functions`, so those are never - // compacted. - // - // But for every compactable arena in a `Module`, whether global to the - // `Module` or local to a function or entry point, the `ModuleTracer` type - // holds a bitmap indicating which elements of that arena are used. Our task - // is to populate those bitmaps correctly. + // For every compactable arena in a `Module`, whether global to the `Module` + // or local to a function or entry point, the `ModuleTracer` type holds a + // bitmap indicating which elements of that arena are used. Our task is to + // populate those bitmaps correctly. // // First, we mark everything that is considered used by definition, as // described in this function's documentation. @@ -99,22 +118,76 @@ pub fn compact(module: &mut crate::Module) { // anything it refers to could be compacted away. let mut module_tracer = ModuleTracer::new(module); - // We treat all globals as used by definition. - log::trace!("tracing global variables"); - { - for (_, global) in module.global_variables.iter() { - log::trace!("tracing global {:?}", global.name); - module_tracer.types_used.insert(global.ty); - if let Some(init) = global.init { - module_tracer.global_expressions_used.insert(init); + // Observe what each entry point actually uses. + log::trace!("tracing entry points"); + let entry_point_maps = module + .entry_points + .iter() + .map(|e| { + log::trace!("tracing entry point {:?}", e.function.name); + + if let Some(sizes) = e.workgroup_size_overrides { + for size in sizes.iter().filter_map(|x| *x) { + module_tracer.global_expressions_used.insert(size); + } } - } + + let mut used = module_tracer.as_function(&e.function); + used.trace(); + FunctionMap::from(used) + }) + .collect::>(); + + // Observe which types, constant expressions, constants, and expressions + // each function uses, and produce maps for each function from + // pre-compaction to post-compaction expression handles. + // + // The function tracing logic here works in conjunction with + // `FunctionTracer::trace_call`, which, when tracing a `Statement::Call` + // to a function not already identified as used, adds the called function + // to both `functions_used` and `functions_pending`. + // + // Called functions are required to appear before their callers in the + // functions arena (recursion is disallowed). We have already traced the + // entry point(s) and added any functions called directly by the entry + // point(s) to `functions_pending`. We proceed by repeatedly tracing the + // last function in `functions_pending`. By an inductive argument, any + // functions after the last function in `functions_pending` must be unused. + // + // When `KeepUnused` is active, we simply mark all functions as pending, + // and then trace all of them. + log::trace!("tracing functions"); + let mut function_maps = HandleMap::with_capacity(module.functions.len()); + if keep_unused.into() { + module_tracer.functions_used.add_all(); + module_tracer.functions_pending.add_all(); + } + while let Some(handle) = module_tracer.functions_pending.pop() { + let function = &module.functions[handle]; + log::trace!("tracing function {:?}", function); + let mut function_tracer = module_tracer.as_function(function); + function_tracer.trace(); + function_maps.insert(handle, FunctionMap::from(function_tracer)); } // We treat all special types as used by definition. log::trace!("tracing special types"); module_tracer.trace_special_types(&module.special_types); + log::trace!("tracing global variables"); + if keep_unused.into() { + module_tracer.global_variables_used.add_all(); + } + for global in module_tracer.global_variables_used.iter() { + log::trace!("tracing global {:?}", module.global_variables[global].name); + module_tracer + .types_used + .insert(module.global_variables[global].ty); + if let Some(init) = module.global_variables[global].init { + module_tracer.global_expressions_used.insert(init); + } + } + // We treat all named constants as used by definition, unless they have an // abstract type as we do not want those reaching the validator. log::trace!("tracing named constants"); @@ -129,62 +202,22 @@ pub fn compact(module: &mut crate::Module) { module_tracer.global_expressions_used.insert(constant.init); } - // We treat all named overrides as used by definition. - log::trace!("tracing named overrides"); - for (handle, r#override) in module.overrides.iter() { - if r#override.name.is_some() { - log::trace!("tracing override {:?}", r#override.name.as_ref().unwrap()); - module_tracer.overrides_used.insert(handle); - module_tracer.types_used.insert(r#override.ty); - if let Some(init) = r#override.init { - module_tracer.global_expressions_used.insert(init); + if keep_unused.into() { + // Treat all named overrides as used. + for (handle, r#override) in module.overrides.iter() { + if r#override.name.is_some() && module_tracer.overrides_used.insert(handle) { + module_tracer.types_used.insert(r#override.ty); + if let Some(init) = r#override.init { + module_tracer.global_expressions_used.insert(init); + } } } - } - // We assume that all functions are used. - // - // Observe which types, constant expressions, constants, and - // expressions each function uses, and produce maps for each - // function from pre-compaction to post-compaction expression - // handles. - log::trace!("tracing functions"); - let function_maps: Vec = module - .functions - .iter() - .map(|(_, f)| { - log::trace!("tracing function {:?}", f.name); - let mut function_tracer = module_tracer.as_function(f); - function_tracer.trace(); - FunctionMap::from(function_tracer) - }) - .collect(); - - // Similarly, observe what each entry point actually uses. - log::trace!("tracing entry points"); - let entry_point_maps: Vec = module - .entry_points - .iter() - .map(|e| { - log::trace!("tracing entry point {:?}", e.function.name); - - if let Some(sizes) = e.workgroup_size_overrides { - for size in sizes.iter().filter_map(|x| *x) { - module_tracer.global_expressions_used.insert(size); - } + // Treat all named types as used. + for (handle, ty) in module.types.iter() { + if ty.name.is_some() { + module_tracer.types_used.insert(handle); } - - let mut used = module_tracer.as_function(&e.function); - used.trace(); - FunctionMap::from(used) - }) - .collect(); - - // Treat all named types as used. - for (handle, ty) in module.types.iter() { - log::trace!("tracing type {:?}, name {:?}", handle, ty.name); - if ty.name.is_some() { - module_tracer.types_used.insert(handle); } } @@ -263,61 +296,43 @@ pub fn compact(module: &mut crate::Module) { } } - // Adjust global variables' types and initializers. + // Drop unused global variables, reusing existing storage. + // Adjust used global variables' types and initializers. log::trace!("adjusting global variables"); - for (_, global) in module.global_variables.iter_mut() { - log::trace!("adjusting global {:?}", global.name); - module_map.types.adjust(&mut global.ty); - if let Some(ref mut init) = global.init { - module_map.global_expressions.adjust(init); + module.global_variables.retain_mut(|handle, global| { + if module_map.globals.used(handle) { + log::trace!("retaining global variable {:?}", global.name); + module_map.types.adjust(&mut global.ty); + if let Some(ref mut init) = global.init { + module_map.global_expressions.adjust(init); + } + true + } else { + log::trace!("dropping global variable {:?}", global.name); + false } - } + }); + // Adjust doc comments if let Some(ref mut doc_comments) = module.doc_comments { - let crate::DocComments { - module: _, - types: ref mut doc_comments_for_types, - struct_members: ref mut doc_comments_for_struct_members, - entry_points: _, - functions: _, - constants: ref mut doc_comments_for_constants, - global_variables: _, - } = **doc_comments; - log::trace!("adjusting doc comments for types"); - for (mut ty, doc_comment) in core::mem::take(doc_comments_for_types) { - if !module_map.types.used(ty) { - continue; - } - module_map.types.adjust(&mut ty); - doc_comments_for_types.insert(ty, doc_comment); - } - log::trace!("adjusting doc comments for struct members"); - for ((mut ty, index), doc_comment) in core::mem::take(doc_comments_for_struct_members) { - if !module_map.types.used(ty) { - continue; - } - module_map.types.adjust(&mut ty); - doc_comments_for_struct_members.insert((ty, index), doc_comment); - } - log::trace!("adjusting doc comments for constants"); - for (mut constant, doc_comment) in core::mem::take(doc_comments_for_constants) { - if !module_map.constants.used(constant) { - continue; - } - module_map.constants.adjust(&mut constant); - doc_comments_for_constants.insert(constant, doc_comment); - } + module_map.adjust_doc_comments(doc_comments.as_mut()); } // Temporary storage to help us reuse allocations of existing // named expression tables. let mut reused_named_expressions = crate::NamedExpressions::default(); - // Compact each function. - for ((_, function), map) in module.functions.iter_mut().zip(function_maps.iter()) { - log::trace!("compacting function {:?}", function.name); - map.compact(function, &module_map, &mut reused_named_expressions); - } + // Drop unused functions. Compact and adjust used functions. + module.functions.retain_mut(|handle, function| { + if let Some(map) = function_maps.get(handle) { + log::trace!("retaining and compacting function {:?}", function.name); + map.compact(function, &module_map, &mut reused_named_expressions); + true + } else { + log::trace!("dropping function {:?}", function.name); + false + } + }); // Compact each entry point. for (entry, map) in module.entry_points.iter_mut().zip(entry_point_maps.iter()) { @@ -332,7 +347,14 @@ pub fn compact(module: &mut crate::Module) { struct ModuleTracer<'module> { module: &'module crate::Module, + + /// The subset of functions in `functions_used` that have not yet been + /// traced. + functions_pending: HandleSet, + + functions_used: HandleSet, types_used: HandleSet, + global_variables_used: HandleSet, constants_used: HandleSet, overrides_used: HandleSet, global_expressions_used: HandleSet, @@ -342,7 +364,10 @@ impl<'module> ModuleTracer<'module> { fn new(module: &'module crate::Module) -> Self { Self { module, + functions_pending: HandleSet::for_arena(&module.functions), + functions_used: HandleSet::for_arena(&module.functions), types_used: HandleSet::for_arena(&module.types), + global_variables_used: HandleSet::for_arena(&module.global_variables), constants_used: HandleSet::for_arena(&module.constants), overrides_used: HandleSet::for_arena(&module.overrides), global_expressions_used: HandleSet::for_arena(&module.global_expressions), @@ -450,6 +475,7 @@ impl<'module> ModuleTracer<'module> { overrides: &self.module.overrides, expressions: &self.module.global_expressions, types_used: &mut self.types_used, + global_variables_used: &mut self.global_variables_used, constants_used: &mut self.constants_used, expressions_used: &mut self.global_expressions_used, overrides_used: &mut self.overrides_used, @@ -465,7 +491,10 @@ impl<'module> ModuleTracer<'module> { function, constants: &self.module.constants, overrides: &self.module.overrides, + functions_pending: &mut self.functions_pending, + functions_used: &mut self.functions_used, types_used: &mut self.types_used, + global_variables_used: &mut self.global_variables_used, constants_used: &mut self.constants_used, overrides_used: &mut self.overrides_used, global_expressions_used: &mut self.global_expressions_used, @@ -475,7 +504,9 @@ impl<'module> ModuleTracer<'module> { } struct ModuleMap { + functions: HandleMap, types: HandleMap, + globals: HandleMap, constants: HandleMap, overrides: HandleMap, global_expressions: HandleMap, @@ -484,7 +515,9 @@ struct ModuleMap { impl From> for ModuleMap { fn from(used: ModuleTracer) -> Self { ModuleMap { + functions: HandleMap::from_set(used.functions_used), types: HandleMap::from_set(used.types_used), + globals: HandleMap::from_set(used.global_variables_used), constants: HandleMap::from_set(used.constants_used), overrides: HandleMap::from_set(used.overrides_used), global_expressions: HandleMap::from_set(used.global_expressions_used), @@ -516,6 +549,58 @@ impl ModuleMap { self.types.adjust(handle); } } + + fn adjust_doc_comments(&self, doc_comments: &mut ir::DocComments) { + let crate::DocComments { + module: _, + types: ref mut doc_types, + struct_members: ref mut doc_struct_members, + entry_points: _, + functions: ref mut doc_functions, + constants: ref mut doc_constants, + global_variables: ref mut doc_globals, + } = *doc_comments; + log::trace!("adjusting doc comments for types"); + for (mut ty, doc_comment) in core::mem::take(doc_types) { + if !self.types.used(ty) { + continue; + } + self.types.adjust(&mut ty); + doc_types.insert(ty, doc_comment); + } + log::trace!("adjusting doc comments for struct members"); + for ((mut ty, index), doc_comment) in core::mem::take(doc_struct_members) { + if !self.types.used(ty) { + continue; + } + self.types.adjust(&mut ty); + doc_struct_members.insert((ty, index), doc_comment); + } + log::trace!("adjusting doc comments for functions"); + for (mut handle, doc_comment) in core::mem::take(doc_functions) { + if !self.functions.used(handle) { + continue; + } + self.functions.adjust(&mut handle); + doc_functions.insert(handle, doc_comment); + } + log::trace!("adjusting doc comments for constants"); + for (mut constant, doc_comment) in core::mem::take(doc_constants) { + if !self.constants.used(constant) { + continue; + } + self.constants.adjust(&mut constant); + doc_constants.insert(constant, doc_comment); + } + log::trace!("adjusting doc comments for globals"); + for (mut handle, doc_comment) in core::mem::take(doc_globals) { + if !self.globals.used(handle) { + continue; + } + self.globals.adjust(&mut handle); + doc_globals.insert(handle, doc_comment); + } + } } struct FunctionMap { @@ -638,7 +723,7 @@ fn type_expression_interdependence() { let ty_init = type_needs_expression(&mut module, expr_trace); type_needed(&mut module, ty_init); let untouched = module.clone(); - compact(&mut module); + compact(&mut module, KeepUnused::Yes); assert!(cmp_modules(&module, &untouched)); let unused_expr = module.global_expressions.append( crate::Expression::Literal(crate::Literal::U32(1)), @@ -646,7 +731,7 @@ fn type_expression_interdependence() { ); type_needs_expression(&mut module, unused_expr); assert!(!cmp_modules(&module, &untouched)); - compact(&mut module); + compact(&mut module, KeepUnused::Yes); assert!(cmp_modules(&module, &untouched)); } @@ -707,7 +792,7 @@ fn array_length_override() { ); assert!(validator.validate(&module).is_ok()); - compact(&mut module); + compact(&mut module, KeepUnused::Yes); assert!(validator.validate(&module).is_ok()); } @@ -813,7 +898,7 @@ fn array_length_override_mutual() { ); assert!(validator.validate(&module).is_ok()); - compact(&mut module); + compact(&mut module, KeepUnused::Yes); assert!(validator.validate(&module).is_ok()); } @@ -862,7 +947,7 @@ fn array_length_expression() { ); assert!(validator.validate(&module).is_ok()); - compact(&mut module); + compact(&mut module, KeepUnused::Yes); assert!(validator.validate(&module).is_ok()); } @@ -918,7 +1003,7 @@ fn global_expression_override() { ); assert!(validator.validate(&module).is_ok()); - compact(&mut module); + compact(&mut module, KeepUnused::Yes); assert!(validator.validate(&module).is_ok()); } @@ -989,7 +1074,7 @@ fn local_expression_override() { ); assert!(validator.validate(&module).is_ok()); - compact(&mut module); + compact(&mut module, KeepUnused::Yes); assert!(validator.validate(&module).is_ok()); } @@ -1061,7 +1146,7 @@ fn unnamed_constant_type() { ); assert!(validator.validate(&module).is_ok()); - compact(&mut module); + compact(&mut module, KeepUnused::Yes); assert!(validator.validate(&module).is_ok()); } @@ -1133,6 +1218,6 @@ fn unnamed_override_type() { ); assert!(validator.validate(&module).is_ok()); - compact(&mut module); + compact(&mut module, KeepUnused::Yes); assert!(validator.validate(&module).is_ok()); } diff --git a/naga/src/compact/statements.rs b/naga/src/compact/statements.rs index bf8d6ec7c8f..60b6598f95d 100644 --- a/naga/src/compact/statements.rs +++ b/naga/src/compact/statements.rs @@ -3,6 +3,7 @@ use alloc::{vec, vec::Vec}; use super::functions::FunctionTracer; use super::FunctionMap; use crate::arena::Handle; +use crate::compact::handle_set_map::HandleMap; impl FunctionTracer<'_> { pub fn trace_block(&mut self, block: &[crate::Statement]) { @@ -100,10 +101,11 @@ impl FunctionTracer<'_> { self.expressions_used.insert(result); } St::Call { - function: _, + function, ref arguments, result, } => { + self.trace_call(function); for expr in arguments { self.expressions_used.insert(*expr); } @@ -204,7 +206,15 @@ impl FunctionTracer<'_> { } impl FunctionMap { - pub fn adjust_body(&self, function: &mut crate::Function) { + /// Adjust statements in the body of `function`. + /// + /// Adjusts expressions using `self.expressions`, and adjusts calls to other + /// functions using `function_map`. + pub fn adjust_body( + &self, + function: &mut crate::Function, + function_map: &HandleMap, + ) { let block = &mut function.body; let mut worklist: Vec<&mut [crate::Statement]> = vec![block]; let adjust = |handle: &mut Handle| { @@ -305,10 +315,11 @@ impl FunctionMap { adjust(result); } St::Call { - function: _, + ref mut function, ref mut arguments, ref mut result, } => { + function_map.adjust(function); for expr in arguments { adjust(expr); } diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs index eefe861ad63..eceafd08230 100644 --- a/naga/src/front/wgsl/lower/mod.rs +++ b/naga/src/front/wgsl/lower/mod.rs @@ -7,7 +7,6 @@ use alloc::{ }; use core::num::NonZeroU32; -use crate::common::wgsl::{TryToWgsl, TypeContext}; use crate::common::ForDebugWithTypes; use crate::front::wgsl::error::{Error, ExpectedToken, InvalidAssignmentType}; use crate::front::wgsl::index::Index; @@ -15,6 +14,10 @@ use crate::front::wgsl::parse::number::Number; use crate::front::wgsl::parse::{ast, conv}; use crate::front::wgsl::Result; use crate::front::Typifier; +use crate::{ + common::wgsl::{TryToWgsl, TypeContext}, + compact::KeepUnused, +}; use crate::{ir, proc}; use crate::{Arena, FastHashMap, FastIndexMap, Handle, Span}; @@ -1335,7 +1338,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { // Constant evaluation may leave abstract-typed literals and // compositions in expression arenas, so we need to compact the module // to remove unused expressions and types. - crate::compact::compact(&mut module); + crate::compact::compact(&mut module, KeepUnused::Yes); Ok(module) } diff --git a/naga/src/lib.rs b/naga/src/lib.rs index d6229742b20..e0acceec7d4 100644 --- a/naga/src/lib.rs +++ b/naga/src/lib.rs @@ -108,7 +108,6 @@ extern crate alloc; mod arena; pub mod back; pub mod common; -#[cfg(feature = "compact")] pub mod compact; pub mod diagnostic_filter; pub mod error; diff --git a/naga/tests/in/wgsl/runtime-array-in-unused-struct.toml b/naga/tests/in/wgsl/runtime-array-in-unused-struct.toml deleted file mode 100644 index 4883dbbd8db..00000000000 --- a/naga/tests/in/wgsl/runtime-array-in-unused-struct.toml +++ /dev/null @@ -1 +0,0 @@ -targets = "SPIRV" diff --git a/naga/tests/in/wgsl/runtime-array-in-unused-struct.wgsl b/naga/tests/in/wgsl/runtime-array-in-unused-struct.wgsl deleted file mode 100644 index bcee56d9b06..00000000000 --- a/naga/tests/in/wgsl/runtime-array-in-unused-struct.wgsl +++ /dev/null @@ -1,12 +0,0 @@ -struct DataStruct { - data: f32, - data_vec: vec4, -} - -struct Struct { - data: array, -}; - -struct PrimitiveStruct { - data: array, -}; diff --git a/naga/tests/in/wgsl/types_with_comments.wgsl b/naga/tests/in/wgsl/types_with_comments.wgsl index dc3cba7a68c..c873311d07c 100644 --- a/naga/tests/in/wgsl/types_with_comments.wgsl +++ b/naga/tests/in/wgsl/types_with_comments.wgsl @@ -6,22 +6,38 @@ */ @group(0) @binding(0) var mvp_matrix: mat4x4; -/// workgroup var doc comment +/// workgroup var 1 doc comment /// 2nd line of workgroup var doc comment var w_mem: mat2x2; +/// workgroup var 2 doc comment +var w_mem2: mat2x2; + /// constant doc comment const test_c: u32 = 1; -/// struct doc comment +/// struct R doc comment +struct TestR { + /// member doc comment + test_m: u32, +} + +/// struct S doc comment struct TestS { /// member doc comment test_m: u32, } -/// function doc comment +/// function f doc comment fn test_f() {} +/// function g doc comment +fn test_g() {} + /// entry point doc comment @compute @workgroup_size(1) -fn test_ep() {} \ No newline at end of file +fn test_ep() { + _ = w_mem2; + _ = TestS(); + test_g(); +} diff --git a/naga/tests/naga/snapshots.rs b/naga/tests/naga/snapshots.rs index 1abbb502d6b..08a398af30d 100644 --- a/naga/tests/naga/snapshots.rs +++ b/naga/tests/naga/snapshots.rs @@ -9,6 +9,7 @@ use std::{ path::{Path, PathBuf}, }; +use naga::compact::KeepUnused; use ron::de; const CRATE_ROOT: &str = env!("CARGO_MANIFEST_DIR"); @@ -243,6 +244,12 @@ impl Input { return None; } + if let Ok(pat) = std::env::var("NAGA_SNAPSHOT") { + if !file_name.to_string_lossy().contains(&pat) { + return None; + } + } + let input = Input::new( subdirectory, file_name.file_stem().unwrap().to_str().unwrap(), @@ -422,9 +429,14 @@ fn check_targets(input: &Input, module: &mut naga::Module, source_code: Option<& ); }); - #[cfg(feature = "compact")] let info = { - naga::compact::compact(module); + // Our backends often generate temporary names based on handle indices, + // which means that adding or removing unused arena entries can affect + // the output even though they have no semantic effect. Such + // meaningless changes add noise to snapshot diffs, making accurate + // patch review difficult. Compacting the modules before generating + // snapshots makes the output independent of unused arena entries. + naga::compact::compact(module, KeepUnused::No); #[cfg(feature = "serialize")] { @@ -609,7 +621,7 @@ fn write_output_spv( }; let (module, info) = - naga::back::pipeline_constants::process_overrides(module, info, pipeline_constants) + naga::back::pipeline_constants::process_overrides(module, info, None, pipeline_constants) .expect("override evaluation failed"); if params.separate_entry_points { @@ -673,7 +685,7 @@ fn write_output_msl( println!("generating MSL"); let (module, info) = - naga::back::pipeline_constants::process_overrides(module, info, pipeline_constants) + naga::back::pipeline_constants::process_overrides(module, info, None, pipeline_constants) .expect("override evaluation failed"); let mut options = options.clone(); @@ -715,7 +727,7 @@ fn write_output_glsl( let mut buffer = String::new(); let (module, info) = - naga::back::pipeline_constants::process_overrides(module, info, pipeline_constants) + naga::back::pipeline_constants::process_overrides(module, info, None, pipeline_constants) .expect("override evaluation failed"); let mut writer = glsl::Writer::new( &mut buffer, @@ -747,7 +759,7 @@ fn write_output_hlsl( println!("generating HLSL"); let (module, info) = - naga::back::pipeline_constants::process_overrides(module, info, pipeline_constants) + naga::back::pipeline_constants::process_overrides(module, info, None, pipeline_constants) .expect("override evaluation failed"); let mut buffer = String::new(); diff --git a/naga/tests/naga/validation.rs b/naga/tests/naga/validation.rs index 099a0f79ada..b7ba70ee028 100644 --- a/naga/tests/naga/validation.rs +++ b/naga/tests/naga/validation.rs @@ -1,4 +1,8 @@ -use naga::{valid, Expression, Function, Scalar}; +use naga::{ + ir, + valid::{self, ModuleInfo}, + Expression, Function, Module, Scalar, +}; /// Validation should fail if `AtomicResult` expressions are not /// populated by `Atomic` statements. @@ -728,7 +732,6 @@ fn bad_texture_dimensions_level() { fn arity_check() { use ir::MathFunction as Mf; use naga::Span; - use naga::{ir, valid}; let _ = env_logger::builder().is_test(true).try_init(); type Result = core::result::Result; @@ -912,3 +915,201 @@ fn main() { naga::valid::GlobalUse::READ, ); } + +/// Parse and validate the module defined in `source`. +/// +/// Panics if unsuccessful. +fn parse_validate(source: &str) -> (Module, ModuleInfo) { + let module = naga::front::wgsl::parse_str(source).expect("module should parse"); + let info = valid::Validator::new(Default::default(), valid::Capabilities::all()) + .validate(&module) + .unwrap(); + (module, info) +} + +/// Helper for `process_overrides` tests. +/// +/// The goal of these tests is to verify that `process_overrides` accepts cases +/// where all necessary overrides are specified (even if some unnecessary ones +/// are not), and does not accept cases where necessary overrides are missing. +/// "Necessary" means that the override is referenced in some way by some +/// function reachable from the specified entry point. +/// +/// Each test passes a source snippet containing a compute entry point `used` +/// that makes use of the override `ov` in some way. We augment that with (1) +/// the definition of `ov` and (2) a dummy entrypoint that does not use the +/// override, and then test the matrix of (specified or not) x (used or not). +/// +/// The optional `unused_body` can introduce additional objects to the module, +/// to verify that they are adjusted correctly by compaction. +fn override_test(test_case: &str, unused_body: Option<&str>) { + use hashbrown::HashMap; + use naga::back::pipeline_constants::PipelineConstantError; + + let source = [ + "override ov: u32;\n", + test_case, + "@compute @workgroup_size(64) +fn unused() { +", + unused_body.unwrap_or_default(), + "} +", + ] + .concat(); + let (module, info) = parse_validate(&source); + + let overrides = HashMap::from([(String::from("ov"), 1.)]); + + // Can translate `unused` with or without the override + naga::back::pipeline_constants::process_overrides( + &module, + &info, + Some((ir::ShaderStage::Compute, "unused")), + &HashMap::new(), + ) + .unwrap(); + naga::back::pipeline_constants::process_overrides( + &module, + &info, + Some((ir::ShaderStage::Compute, "unused")), + &overrides, + ) + .unwrap(); + + // Cannot translate `used` without the override + let err = naga::back::pipeline_constants::process_overrides( + &module, + &info, + Some((ir::ShaderStage::Compute, "used")), + &HashMap::new(), + ) + .unwrap_err(); + assert!(matches!(err, PipelineConstantError::MissingValue(name) if name == "ov")); + + // Can translate `used` if the override is specified + naga::back::pipeline_constants::process_overrides( + &module, + &info, + Some((ir::ShaderStage::Compute, "used")), + &overrides, + ) + .unwrap(); +} + +#[cfg(feature = "wgsl-in")] +#[test] +fn override_in_workgroup_size() { + override_test( + " +@compute @workgroup_size(ov) +fn used() { +} +", + None, + ); +} + +#[cfg(feature = "wgsl-in")] +#[test] +fn override_in_workgroup_size_nested() { + // Initializer for override used in workgroup size refers to another + // override. + override_test( + " +override ov2: u32 = ov + 5; + +@compute @workgroup_size(ov2) +fn used() { +} +", + None, + ); +} + +#[cfg(feature = "wgsl-in")] +#[test] +fn override_in_function() { + override_test( + " +fn foo() -> u32 { + return ov; +} + +@compute @workgroup_size(64) +fn used() { + foo(); +} +", + None, + ); +} + +#[cfg(feature = "wgsl-in")] +#[test] +fn override_in_entrypoint() { + override_test( + " +fn foo() -> u32 { + return ov; +} + +@compute @workgroup_size(64) +fn used() { + foo(); +} +", + None, + ); +} + +#[cfg(feature = "wgsl-in")] +#[test] +fn override_in_array_size() { + override_test( + " +var arr: array; + +@compute @workgroup_size(64) +fn used() { + _ = arr[5]; +} +", + None, + ); +} + +#[cfg(feature = "wgsl-in")] +#[test] +fn override_in_global_init() { + override_test( + " +var foo: u32 = ov; + +@compute @workgroup_size(64) +fn used() { + _ = foo; +} +", + None, + ); +} + +#[cfg(feature = "wgsl-in")] +#[test] +fn override_with_multiple_globals() { + // Test that when compaction of the `unused` entrypoint removes `arr1`, the + // handle to `arr2` is adjusted correctly. + override_test( + " +var arr1: array; +var arr2: array; + +@compute @workgroup_size(64) +fn used() { + _ = arr1[5]; +} +", + Some("_ = arr2[3];"), + ); +} diff --git a/naga/tests/naga/wgsl_errors.rs b/naga/tests/naga/wgsl_errors.rs index 1784715a384..2a330b88baf 100644 --- a/naga/tests/naga/wgsl_errors.rs +++ b/naga/tests/naga/wgsl_errors.rs @@ -3,7 +3,7 @@ Tests for the WGSL front end. */ #![cfg(feature = "wgsl-in")] -use naga::valid::Capabilities; +use naga::{compact::KeepUnused, valid::Capabilities}; #[track_caller] fn check(input: &str, snapshot: &str) { @@ -2902,7 +2902,7 @@ fn compaction_preserves_spans() { "#; // The error span should be on `x[1.0]`, which is at characters 108..114. let mut module = naga::front::wgsl::parse_str(source).expect("source ought to parse"); - naga::compact::compact(&mut module); + naga::compact::compact(&mut module, KeepUnused::No); let err = naga::valid::Validator::new( naga::valid::ValidationFlags::all(), naga::valid::Capabilities::default(), diff --git a/naga/tests/out/ir/wgsl-types_with_comments.compact.ron b/naga/tests/out/ir/wgsl-types_with_comments.compact.ron index fcb319b78cb..dde1bd367b0 100644 --- a/naga/tests/out/ir/wgsl-types_with_comments.compact.ron +++ b/naga/tests/out/ir/wgsl-types_with_comments.compact.ron @@ -1,16 +1,5 @@ ( types: [ - ( - name: None, - inner: Matrix( - columns: Quad, - rows: Quad, - scalar: ( - kind: Float, - width: 4, - ), - ), - ), ( name: None, inner: Matrix( @@ -35,7 +24,7 @@ members: [ ( name: Some("test_m"), - ty: 2, + ty: 1, binding: None, offset: 0, ), @@ -53,27 +42,17 @@ constants: [ ( name: Some("test_c"), - ty: 2, + ty: 1, init: 0, ), ], overrides: [], global_variables: [ ( - name: Some("mvp_matrix"), - space: Uniform, - binding: Some(( - group: 0, - binding: 0, - )), - ty: 0, - init: None, - ), - ( - name: Some("w_mem"), + name: Some("w_mem2"), space: WorkGroup, binding: None, - ty: 1, + ty: 0, init: None, ), ], @@ -82,7 +61,7 @@ ], functions: [ ( - name: Some("test_f"), + name: Some("test_g"), arguments: [], result: None, local_variables: [], @@ -108,9 +87,27 @@ arguments: [], result: None, local_variables: [], - expressions: [], - named_expressions: {}, + expressions: [ + GlobalVariable(0), + Load( + pointer: 0, + ), + ZeroValue(2), + ], + named_expressions: { + 1: "phony", + 2: "phony", + }, body: [ + Emit(( + start: 1, + end: 2, + )), + Call( + function: 0, + arguments: [], + result: None, + ), Return( value: None, ), @@ -123,12 +120,12 @@ diagnostic_filter_leaf: None, doc_comments: Some(( types: { - 3: [ - "/// struct doc comment", + 2: [ + "/// struct S doc comment", ], }, struct_members: { - (3, 0): [ + (2, 0): [ "/// member doc comment", ], }, @@ -139,7 +136,7 @@ }, functions: { 0: [ - "/// function doc comment", + "/// function g doc comment", ], }, constants: { @@ -149,11 +146,7 @@ }, global_variables: { 0: [ - "/**\n 🍽\u{fe0f} /* nested comment */\n */", - ], - 1: [ - "/// workgroup var doc comment", - "/// 2nd line of workgroup var doc comment", + "/// workgroup var 2 doc comment", ], }, module: [ diff --git a/naga/tests/out/ir/wgsl-types_with_comments.ron b/naga/tests/out/ir/wgsl-types_with_comments.ron index fcb319b78cb..3a23c49c4f2 100644 --- a/naga/tests/out/ir/wgsl-types_with_comments.ron +++ b/naga/tests/out/ir/wgsl-types_with_comments.ron @@ -29,6 +29,20 @@ width: 4, )), ), + ( + name: Some("TestR"), + inner: Struct( + members: [ + ( + name: Some("test_m"), + ty: 2, + binding: None, + offset: 0, + ), + ], + span: 4, + ), + ), ( name: Some("TestS"), inner: Struct( @@ -76,6 +90,13 @@ ty: 1, init: None, ), + ( + name: Some("w_mem2"), + space: WorkGroup, + binding: None, + ty: 1, + init: None, + ), ], global_expressions: [ Literal(U32(1)), @@ -95,6 +116,20 @@ ], diagnostic_filter_leaf: None, ), + ( + name: Some("test_g"), + arguments: [], + result: None, + local_variables: [], + expressions: [], + named_expressions: {}, + body: [ + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), ], entry_points: [ ( @@ -108,9 +143,27 @@ arguments: [], result: None, local_variables: [], - expressions: [], - named_expressions: {}, + expressions: [ + GlobalVariable(2), + Load( + pointer: 0, + ), + ZeroValue(4), + ], + named_expressions: { + 1: "phony", + 2: "phony", + }, body: [ + Emit(( + start: 1, + end: 2, + )), + Call( + function: 1, + arguments: [], + result: None, + ), Return( value: None, ), @@ -124,13 +177,19 @@ doc_comments: Some(( types: { 3: [ - "/// struct doc comment", + "/// struct R doc comment", + ], + 4: [ + "/// struct S doc comment", ], }, struct_members: { (3, 0): [ "/// member doc comment", ], + (4, 0): [ + "/// member doc comment", + ], }, entry_points: { 0: [ @@ -139,7 +198,10 @@ }, functions: { 0: [ - "/// function doc comment", + "/// function f doc comment", + ], + 1: [ + "/// function g doc comment", ], }, constants: { @@ -152,9 +214,12 @@ "/**\n 🍽\u{fe0f} /* nested comment */\n */", ], 1: [ - "/// workgroup var doc comment", + "/// workgroup var 1 doc comment", "/// 2nd line of workgroup var doc comment", ], + 2: [ + "/// workgroup var 2 doc comment", + ], }, module: [ "//! Module doc comment.", diff --git a/naga/tests/out/spv/wgsl-runtime-array-in-unused-struct.spvasm b/naga/tests/out/spv/wgsl-runtime-array-in-unused-struct.spvasm deleted file mode 100644 index ab3bd3d01bf..00000000000 --- a/naga/tests/out/spv/wgsl-runtime-array-in-unused-struct.spvasm +++ /dev/null @@ -1,24 +0,0 @@ -; SPIR-V -; Version: 1.1 -; Generator: rspirv -; Bound: 10 -OpCapability Shader -OpCapability Linkage -%1 = OpExtInstImport "GLSL.std.450" -OpMemoryModel Logical GLSL450 -OpMemberDecorate %5 0 Offset 0 -OpMemberDecorate %5 1 Offset 16 -OpDecorate %6 ArrayStride 32 -OpMemberDecorate %7 0 Offset 0 -OpDecorate %7 Block -OpDecorate %8 ArrayStride 4 -OpMemberDecorate %9 0 Offset 0 -OpDecorate %9 Block -%2 = OpTypeVoid -%3 = OpTypeFloat 32 -%4 = OpTypeVector %3 4 -%5 = OpTypeStruct %3 %4 -%6 = OpTypeRuntimeArray %5 -%7 = OpTypeStruct %6 -%8 = OpTypeRuntimeArray %3 -%9 = OpTypeStruct %8 \ No newline at end of file diff --git a/naga/tests/out/wgsl/glsl-896-push-constant.frag.wgsl b/naga/tests/out/wgsl/glsl-896-push-constant.frag.wgsl index 729e35a43fd..6b3c4ac973f 100644 --- a/naga/tests/out/wgsl/glsl-896-push-constant.frag.wgsl +++ b/naga/tests/out/wgsl/glsl-896-push-constant.frag.wgsl @@ -1,9 +1,3 @@ -struct PushConstants { - example: f32, -} - -var c: PushConstants; - fn main_1() { return; } diff --git a/naga/tests/out/wgsl/glsl-900-implicit-conversions.frag.wgsl b/naga/tests/out/wgsl/glsl-900-implicit-conversions.frag.wgsl index 2c468dd929a..fa537f1f789 100644 --- a/naga/tests/out/wgsl/glsl-900-implicit-conversions.frag.wgsl +++ b/naga/tests/out/wgsl/glsl-900-implicit-conversions.frag.wgsl @@ -1,63 +1,28 @@ -fn exact(a: f32) { - var a_1: f32; +fn exact(a: i32) { + var a_1: i32; a_1 = a; return; } -fn exact_1(a_2: i32) { - var a_3: i32; +fn implicit(a_2: f32) { + var a_3: f32; a_3 = a_2; return; } -fn implicit(a_4: f32) { - var a_5: f32; - - a_5 = a_4; - return; -} - -fn implicit_1(a_6: i32) { - var a_7: i32; - - a_7 = a_6; - return; -} - -fn implicit_dims(v: f32) { - var v_1: f32; +fn implicit_dims(v: vec3) { + var v_1: vec3; v_1 = v; return; } -fn implicit_dims_1(v_2: vec2) { - var v_3: vec2; - - v_3 = v_2; - return; -} - -fn implicit_dims_2(v_4: vec3) { - var v_5: vec3; - - v_5 = v_4; - return; -} - -fn implicit_dims_3(v_6: vec4) { - var v_7: vec4; - - v_7 = v_6; - return; -} - fn main_1() { - exact_1(1i); + exact(1i); implicit(1f); - implicit_dims_2(vec3(1f)); + implicit_dims(vec3(1f)); return; } diff --git a/naga/tests/out/wgsl/glsl-anonymous-entry-point-type.frag.wgsl b/naga/tests/out/wgsl/glsl-anonymous-entry-point-type.frag.wgsl index 5778897d4ce..88d33a0a055 100644 --- a/naga/tests/out/wgsl/glsl-anonymous-entry-point-type.frag.wgsl +++ b/naga/tests/out/wgsl/glsl-anonymous-entry-point-type.frag.wgsl @@ -1,8 +1,4 @@ struct FragmentOutput { - x: f32, -} - -struct FragmentOutput_1 { @location(0) o_Target: vec4, } @@ -14,8 +10,8 @@ fn main_1() { } @fragment -fn main() -> FragmentOutput_1 { +fn main() -> FragmentOutput { main_1(); let _e1 = o_Target; - return FragmentOutput_1(_e1); + return FragmentOutput(_e1); } diff --git a/naga/tests/out/wgsl/glsl-bevy-pbr.frag.wgsl b/naga/tests/out/wgsl/glsl-bevy-pbr.frag.wgsl index 0173c0f002f..fbd8d1fee90 100644 --- a/naga/tests/out/wgsl/glsl-bevy-pbr.frag.wgsl +++ b/naga/tests/out/wgsl/glsl-bevy-pbr.frag.wgsl @@ -9,10 +9,6 @@ struct DirectionalLight { color: vec4, } -struct CameraViewProj { - ViewProj: mat4x4, -} - struct CameraPosition { CameraPos: vec4, } @@ -57,28 +53,26 @@ var v_WorldNormal_1: vec3; var v_Uv_1: vec2; var v_WorldTangent_1: vec4; var o_Target: vec4; -@group(0) @binding(0) -var global: CameraViewProj; @group(0) @binding(1) -var global_1: CameraPosition; +var global: CameraPosition; @group(1) @binding(0) -var global_2: Lights; +var global_1: Lights; @group(3) @binding(0) -var global_3: StandardMaterial_base_color; +var global_2: StandardMaterial_base_color; @group(3) @binding(1) var StandardMaterial_base_color_texture: texture_2d; @group(3) @binding(2) var StandardMaterial_base_color_texture_sampler: sampler; @group(3) @binding(3) -var global_4: StandardMaterial_roughness; +var global_3: StandardMaterial_roughness; @group(3) @binding(4) -var global_5: StandardMaterial_metallic; +var global_4: StandardMaterial_metallic; @group(3) @binding(5) var StandardMaterial_metallic_roughness_texture: texture_2d; @group(3) @binding(6) var StandardMaterial_metallic_roughness_texture_sampler: sampler; @group(3) @binding(7) -var global_6: StandardMaterial_reflectance; +var global_5: StandardMaterial_reflectance; @group(3) @binding(8) var StandardMaterial_normal_map: texture_2d; @group(3) @binding(9) @@ -88,7 +82,7 @@ var StandardMaterial_occlusion_texture: texture_2d; @group(3) @binding(11) var StandardMaterial_occlusion_texture_sampler: sampler; @group(3) @binding(12) -var global_7: StandardMaterial_emissive; +var global_6: StandardMaterial_emissive; @group(3) @binding(13) var StandardMaterial_emissive_texture: texture_2d; @group(3) @binding(14) @@ -353,32 +347,6 @@ fn perceptualRoughnessToRoughness(perceptualRoughness: f32) -> f32 { return (_e7 * _e8); } -fn reinhard(color: vec3) -> vec3 { - var color_1: vec3; - - color_1 = color; - let _e2 = color_1; - let _e5 = color_1; - return (_e2 / (vec3(1f) + _e5)); -} - -fn reinhard_extended(color_2: vec3, max_white: f32) -> vec3 { - var color_3: vec3; - var max_white_1: f32; - var numerator: vec3; - - color_3 = color_2; - max_white_1 = max_white; - let _e4 = color_3; - let _e7 = color_3; - let _e8 = max_white_1; - let _e9 = max_white_1; - numerator = (_e4 * (vec3(1f) + (_e7 / vec3((_e8 * _e9))))); - let _e16 = numerator; - let _e19 = color_3; - return (_e16 / (vec3(1f) + _e19)); -} - fn luminance(v_1: vec3) -> f32 { var v_2: vec3; @@ -403,50 +371,24 @@ fn change_luminance(c_in: vec3, l_out: f32) -> vec3 { return (_e7 * (_e8 / _e9)); } -fn reinhard_luminance(color_4: vec3) -> vec3 { - var color_5: vec3; +fn reinhard_luminance(color: vec3) -> vec3 { + var color_1: vec3; var l_old: f32; var l_new: f32; - color_5 = color_4; - let _e2 = color_5; + color_1 = color; + let _e2 = color_1; let _e3 = luminance(_e2); l_old = _e3; let _e5 = l_old; let _e7 = l_old; l_new = (_e5 / (1f + _e7)); - let _e11 = color_5; + let _e11 = color_1; let _e12 = l_new; let _e13 = change_luminance(_e11, _e12); return _e13; } -fn reinhard_extended_luminance(color_6: vec3, max_white_l: f32) -> vec3 { - var color_7: vec3; - var max_white_l_1: f32; - var l_old_1: f32; - var numerator_1: f32; - var l_new_1: f32; - - color_7 = color_6; - max_white_l_1 = max_white_l; - let _e4 = color_7; - let _e5 = luminance(_e4); - l_old_1 = _e5; - let _e7 = l_old_1; - let _e9 = l_old_1; - let _e10 = max_white_l_1; - let _e11 = max_white_l_1; - numerator_1 = (_e7 * (1f + (_e9 / (_e10 * _e11)))); - let _e17 = numerator_1; - let _e19 = l_old_1; - l_new_1 = (_e17 / (1f + _e19)); - let _e23 = color_7; - let _e24 = l_new_1; - let _e25 = change_luminance(_e23, _e24); - return _e25; -} - fn point_light(light: PointLight, roughness_8: f32, NdotV: f32, N: vec3, V_1: vec3, R: vec3, F0_: vec3, diffuseColor: vec3) -> vec3 { var light_1: PointLight; var roughness_9: f32; @@ -662,7 +604,7 @@ fn main_1() { var diffuse_ambient: vec3; var specular_ambient: vec3; - let _e37 = global_3.base_color; + let _e37 = global_2.base_color; output_color = _e37; let _e39 = output_color; let _e40 = v_Uv_1; @@ -671,10 +613,10 @@ fn main_1() { let _e43 = v_Uv_1; let _e44 = textureSample(StandardMaterial_metallic_roughness_texture, StandardMaterial_metallic_roughness_texture_sampler, _e43); metallic_roughness = _e44; - let _e46 = global_5.metallic; + let _e46 = global_4.metallic; let _e47 = metallic_roughness; metallic = (_e46 * _e47.z); - let _e51 = global_4.perceptual_roughness; + let _e51 = global_3.perceptual_roughness; let _e52 = metallic_roughness; perceptual_roughness_2 = (_e51 * _e52.y); let _e56 = perceptual_roughness_2; @@ -729,7 +671,7 @@ fn main_1() { let _e120 = v_Uv_1; let _e121 = textureSample(StandardMaterial_occlusion_texture, StandardMaterial_occlusion_texture_sampler, _e120); occlusion = _e121.x; - let _e124 = global_7.emissive; + let _e124 = global_6.emissive; emissive = _e124; let _e126 = emissive; let _e128 = v_Uv_1; @@ -738,14 +680,14 @@ fn main_1() { emissive.x = _e131.x; emissive.y = _e131.y; emissive.z = _e131.z; - let _e138 = global_1.CameraPos; + let _e138 = global.CameraPos; let _e140 = v_WorldPosition_1; V_3 = normalize((_e138.xyz - _e140.xyz)); let _e145 = N_2; let _e146 = V_3; NdotV_4 = max(dot(_e145, _e146), 0.001f); - let _e152 = global_6.reflectance; - let _e154 = global_6.reflectance; + let _e152 = global_5.reflectance; + let _e154 = global_5.reflectance; let _e157 = metallic; let _e161 = output_color; let _e163 = metallic; @@ -758,7 +700,7 @@ fn main_1() { R_4 = reflect(-(_e176), _e178); loop { let _e186 = i; - let _e187 = global_2.NumLights; + let _e187 = global_1.NumLights; let _e191 = i; if !(((_e186 < i32(_e187.x)) && (_e191 < MAX_POINT_LIGHTS))) { break; @@ -766,7 +708,7 @@ fn main_1() { { let _e198 = light_accum; let _e199 = i; - let _e201 = global_2.PointLights[_e199]; + let _e201 = global_1.PointLights[_e199]; let _e202 = roughness_12; let _e203 = NdotV_4; let _e204 = N_2; @@ -784,7 +726,7 @@ fn main_1() { } loop { let _e213 = i_1; - let _e214 = global_2.NumLights; + let _e214 = global_1.NumLights; let _e218 = i_1; if !(((_e213 < i32(_e214.y)) && (_e218 < MAX_DIRECTIONAL_LIGHTS))) { break; @@ -792,7 +734,7 @@ fn main_1() { { let _e225 = light_accum; let _e226 = i_1; - let _e228 = global_2.DirectionalLights[_e226]; + let _e228 = global_1.DirectionalLights[_e226]; let _e229 = roughness_12; let _e230 = NdotV_4; let _e231 = N_2; @@ -824,7 +766,7 @@ fn main_1() { let _e255 = output_color; let _e257 = diffuse_ambient; let _e258 = specular_ambient; - let _e260 = global_2.AmbientColor; + let _e260 = global_1.AmbientColor; let _e263 = occlusion; let _e265 = (_e255.xyz + (((_e257 + _e258) * _e260.xyz) * _e263)); output_color.x = _e265.x; diff --git a/naga/tests/out/wgsl/glsl-global-constant-array.frag.wgsl b/naga/tests/out/wgsl/glsl-global-constant-array.frag.wgsl index feaca41aa54..6582dfc9ece 100644 --- a/naga/tests/out/wgsl/glsl-global-constant-array.frag.wgsl +++ b/naga/tests/out/wgsl/glsl-global-constant-array.frag.wgsl @@ -1,7 +1,5 @@ const array_: array = array(1f, 2f); -var i: u32; - fn main_1() { var local: array = array_; diff --git a/naga/tests/out/wgsl/glsl-samplers.frag.wgsl b/naga/tests/out/wgsl/glsl-samplers.frag.wgsl index 9be13832457..9ae246b03ef 100644 --- a/naga/tests/out/wgsl/glsl-samplers.frag.wgsl +++ b/naga/tests/out/wgsl/glsl-samplers.frag.wgsl @@ -26,8 +26,6 @@ var tex2DArrayShadow: texture_depth_2d_array; var texCubeShadow: texture_depth_cube; @group(1) @binding(15) var texCubeArrayShadow: texture_depth_cube_array; -@group(1) @binding(16) -var tex3DShadow: texture_3d; @group(1) @binding(17) var sampShadow: sampler_comparison; @group(0) @binding(18) diff --git a/wgpu-hal/src/dx12/device.rs b/wgpu-hal/src/dx12/device.rs index 2345803baa8..c73780bc77b 100644 --- a/wgpu-hal/src/dx12/device.rs +++ b/wgpu-hal/src/dx12/device.rs @@ -280,6 +280,7 @@ impl super::Device { let (module, info) = naga::back::pipeline_constants::process_overrides( &stage.module.naga.module, &stage.module.naga.info, + Some((naga_stage, stage.entry_point)), stage.constants, ) .map_err(|e| crate::PipelineError::PipelineConstants(stage_bit, format!("HLSL: {e:?}")))?; diff --git a/wgpu-hal/src/gles/device.rs b/wgpu-hal/src/gles/device.rs index 852042e8ef8..4e02c8775b9 100644 --- a/wgpu-hal/src/gles/device.rs +++ b/wgpu-hal/src/gles/device.rs @@ -219,6 +219,7 @@ impl super::Device { let (module, info) = naga::back::pipeline_constants::process_overrides( &stage.module.naga.module, &stage.module.naga.info, + Some((naga_stage, stage.entry_point)), stage.constants, ) .map_err(|e| { diff --git a/wgpu-hal/src/metal/device.rs b/wgpu-hal/src/metal/device.rs index 6ab22b0c3e3..b1300a31e06 100644 --- a/wgpu-hal/src/metal/device.rs +++ b/wgpu-hal/src/metal/device.rs @@ -137,6 +137,7 @@ impl super::Device { let (module, module_info) = naga::back::pipeline_constants::process_overrides( &naga_shader.module, &naga_shader.info, + Some((naga_stage, stage.entry_point)), stage.constants, ) .map_err(|e| crate::PipelineError::PipelineConstants(stage_bit, format!("MSL: {:?}", e)))?; diff --git a/wgpu-hal/src/vulkan/device.rs b/wgpu-hal/src/vulkan/device.rs index 4ae76565120..7e297421a81 100644 --- a/wgpu-hal/src/vulkan/device.rs +++ b/wgpu-hal/src/vulkan/device.rs @@ -892,6 +892,7 @@ impl super::Device { let (module, info) = naga::back::pipeline_constants::process_overrides( &naga_shader.module, &naga_shader.info, + Some((naga_stage, stage.entry_point)), stage.constants, ) .map_err(|e| {