archive-gh-me/wgpu
1
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2025-12-08 21:26:17 +00:00
Code Issues Projects Releases Wiki Activity

[naga] Remove non-essential override references via compaction

Browse Source 
Adds a mode to compaction that removes unused functions, global
variables, and named types and overrides. This mode is used
everywhere except the compaction at the end of lowering, where
it is important to preserve unused items for type checking and
other validation of the module.

Pruning all but the active entry point and then compacting makes
`process_overrides` tolerant of missing values for overrides that are
not used by the active entry point.

Fixes #5885
This commit is contained in:
Andy Leiserson 2025-05-19 11:45:41 -07:00 committed by Erich Gubler
parent 096f1f1f6d
commit 82fa8e2a94
31 changed files with 734 additions and 396 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

13
CHANGELOG.md
View File

@ -69,8 +69,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).
@ -80,6 +79,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
@ -107,6 +107,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
@ -124,11 +128,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

4
benches/benches/wgpu-benchmark/shader.rs
View File

@ -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);
}
});
});

2
cts_runner/test.lst
View File

@ -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:*

44
naga-cli/src/bin/naga.rs
View File

@ -542,7 +542,7 @@ fn run() -> anyhow::Result<()> {
write_output(&module, &info, &params, 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, &params.overrides)
.unwrap_pretty();
let (module, info) = naga::back::pipeline_constants::process_overrides(
module,
info,
None,
&params.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, &params.overrides)
.unwrap_pretty();
let (module, info) = naga::back::pipeline_constants::process_overrides(
module,
info,
None,
&params.overrides,
)
.unwrap_pretty();
let spv =
spv::write_vec(&module, &info, &params.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, &params.overrides)
.unwrap_pretty();
let (module, info) = naga::back::pipeline_constants::process_overrides(
module,
info,
None,
&params.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, &params.overrides)
.unwrap_pretty();
let (module, info) = naga::back::pipeline_constants::process_overrides(
module,
info,
None,
&params.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};

22
naga/src/arena/handle_set.rs
View File

@ -77,6 +77,11 @@ impl<T> HandleSet<T> {
}
}
/// 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<T>) -> bool {
self.members.contains(handle.index())
}
@ -85,6 +90,23 @@ impl<T> HandleSet<T> {
pub fn iter(&self) -> impl '_ + Iterator<Item = Handle<T>> {
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<Handle<T>> {
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<T> Default for HandleSet<T> {

13
naga/src/back/pipeline_constants.rs
View File

@ -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());

18
naga/src/compact/expressions.rs
View File

@ -11,6 +11,9 @@ pub struct ExpressionTracer<'tracer> {
/// The used map for `types`.
pub types_used: &'tracer mut HandleSet<crate::Type>,
/// The used map for global variables.
pub global_variables_used: &'tracer mut HandleSet<crate::GlobalVariable>,
/// The used map for `constants`.
pub constants_used: &'tracer mut HandleSet<crate::Constant>,
@ -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: _,

13
naga/src/compact/functions.rs
View File

@ -6,7 +6,10 @@ pub struct FunctionTracer<'a> {
pub constants: &'a crate::Arena<crate::Constant>,
pub overrides: &'a crate::Arena<crate::Override>,
pub functions_pending: &'a mut HandleSet<crate::Function>,
pub functions_used: &'a mut HandleSet<crate::Function>,
pub types_used: &'a mut HandleSet<crate::Type>,
pub global_variables_used: &'a mut HandleSet<crate::GlobalVariable>,
pub constants_used: &'a mut HandleSet<crate::Constant>,
pub overrides_used: &'a mut HandleSet<crate::Override>,
pub global_expressions_used: &'a mut HandleSet<crate::Expression>,
@ -16,6 +19,13 @@ pub struct FunctionTracer<'a> {
}
impl FunctionTracer<'_> {
pub fn trace_call(&mut self, function: crate::Handle<crate::Function>) {
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);
}
}

30
naga/src/compact/handle_set_map.rs
View File

@ -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<T> {
/// 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<T, U = Index> {
/// 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<Option<Index>>,
new_index: Vec<Option<U>>,
/// This type is indexed by values of type `T`.
as_keys: core::marker::PhantomData<T>,
}
impl<T, U> HandleMap<T, U> {
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<T>) -> Option<&U> {
self.new_index.get(handle.index()).unwrap_or(&None).as_ref()
}
pub fn insert(&mut self, handle: Handle<T>, value: U) -> Option<U> {
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<T: 'static> HandleMap<T> {
pub fn from_set(set: HandleSet<T>) -> Self {
let mut next_index = Index::new(0).unwrap();

333
naga/src/compact/mod.rs
View File

@ -16,49 +16,65 @@ 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<KeepUnused> 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.
@ -102,70 +118,9 @@ 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);
}
}
}
// We treat all special types as used by definition.
log::trace!("tracing special types");
module_tracer.trace_special_types(&module.special_types);
// 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");
for (handle, constant) in module.constants.iter() {
if constant.name.is_none() || module.types[constant.ty].inner.is_abstract(&module.types) {
continue;
}
log::trace!("tracing constant {:?}", constant.name.as_ref().unwrap());
module_tracer.constants_used.insert(handle);
module_tracer.types_used.insert(constant.ty);
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);
}
}
}
// 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<FunctionMap> = 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.
// Observe what each entry point actually uses.
log::trace!("tracing entry points");
let entry_point_maps: Vec<FunctionMap> = module
let entry_point_maps = module
.entry_points
.iter()
.map(|e| {
@ -181,13 +136,88 @@ pub fn compact(module: &mut crate::Module) {
used.trace();
FunctionMap::from(used)
})
.collect();
.collect::<Vec<_>>();
// 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);
// 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");
for (handle, constant) in module.constants.iter() {
if constant.name.is_none() || module.types[constant.ty].inner.is_abstract(&module.types) {
continue;
}
log::trace!("tracing constant {:?}", constant.name.as_ref().unwrap());
module_tracer.constants_used.insert(handle);
module_tracer.types_used.insert(constant.ty);
module_tracer.global_expressions_used.insert(constant.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);
}
}
}
// Treat all named types as used.
for (handle, ty) in module.types.iter() {
if ty.name.is_some() {
module_tracer.types_used.insert(handle);
}
}
}
@ -266,15 +296,22 @@ 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 {
@ -285,11 +322,17 @@ pub fn compact(module: &mut crate::Module) {
// 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()) {
@ -304,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<crate::Function>,
functions_used: HandleSet<crate::Function>,
types_used: HandleSet<crate::Type>,
global_variables_used: HandleSet<crate::GlobalVariable>,
constants_used: HandleSet<crate::Constant>,
overrides_used: HandleSet<crate::Override>,
global_expressions_used: HandleSet<crate::Expression>,
@ -314,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),
@ -422,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,
@ -437,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,
@ -447,7 +504,9 @@ impl<'module> ModuleTracer<'module> {
}
struct ModuleMap {
functions: HandleMap<crate::Function>,
types: HandleMap<crate::Type>,
globals: HandleMap<crate::GlobalVariable>,
constants: HandleMap<crate::Constant>,
overrides: HandleMap<crate::Override>,
global_expressions: HandleMap<crate::Expression>,
@ -456,7 +515,9 @@ struct ModuleMap {
impl From<ModuleTracer<'_>> 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),
@ -492,36 +553,52 @@ impl ModuleMap {
fn adjust_doc_comments(&self, doc_comments: &mut ir::DocComments) {
let crate::DocComments {
module: _,
types: ref mut doc_comments_for_types,
struct_members: ref mut doc_comments_for_struct_members,
types: ref mut doc_types,
struct_members: ref mut doc_struct_members,
entry_points: _,
functions: _,
constants: ref mut doc_comments_for_constants,
global_variables: _,
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_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_comments_for_types.insert(ty, doc_comment);
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_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_comments_for_struct_members.insert((ty, index), doc_comment);
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_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_comments_for_constants.insert(constant, doc_comment);
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);
}
}
}
@ -646,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)),
@ -654,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));
}
@ -715,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());
}
@ -821,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());
}
@ -870,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());
}
@ -926,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());
}
@ -997,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());
}
@ -1069,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());
}
@ -1141,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());
}

17
naga/src/compact/statements.rs
View File

@ -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<crate::Function>,
) {
let block = &mut function.body;
let mut worklist: Vec<&mut [crate::Statement]> = vec![block];
let adjust = |handle: &mut Handle<crate::Expression>| {
@ -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);
}

7
naga/src/front/wgsl/lower/mod.rs
View File

@ -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)
}

1
naga/tests/in/wgsl/runtime-array-in-unused-struct.toml
View File

@ -1 +0,0 @@
targets = "SPIRV"

12
naga/tests/in/wgsl/runtime-array-in-unused-struct.wgsl
View File

@ -1,12 +0,0 @@
struct DataStruct {
data: f32,
data_vec: vec4<f32>,
}
struct Struct {
data: array<DataStruct>,
};
struct PrimitiveStruct {
data: array<f32>,
};

24
naga/tests/in/wgsl/types_with_comments.wgsl
View File

@ -6,22 +6,38 @@
*/
@group(0) @binding(0) var<uniform> mvp_matrix: mat4x4<f32>;
/// workgroup var doc comment
/// workgroup var 1 doc comment
/// 2nd line of workgroup var doc comment
var<workgroup> w_mem: mat2x2<f32>;
/// workgroup var 2 doc comment
var<workgroup> w_mem2: mat2x2<f32>;
/// 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() {}
fn test_ep() {
_ = w_mem2;
_ = TestS();
test_g();
}

23
naga/tests/naga/snapshots.rs
View File

@ -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(),
@ -423,7 +430,13 @@ fn check_targets(input: &Input, module: &mut naga::Module, source_code: Option<&
});
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")]
{
@ -608,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 {
@ -672,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();
@ -714,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,
@ -746,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();

205
naga/tests/naga/validation.rs
View File

@ -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<naga::valid::ModuleInfo, naga::valid::ValidationError>;
@ -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<workgroup> arr: array<u32, ov>;
@compute @workgroup_size(64)
fn used() {
_ = arr[5];
}
",
None,
);
}
#[cfg(feature = "wgsl-in")]
#[test]
fn override_in_global_init() {
override_test(
"
var<private> 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<workgroup> arr1: array<u32, ov>;
var<workgroup> arr2: array<u32, 4>;
@compute @workgroup_size(64)
fn used() {
_ = arr1[5];
}
",
Some("_ = arr2[3];"),
);
}

4
naga/tests/naga/wgsl_errors.rs
View File

@ -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(),

67
naga/tests/out/ir/wgsl-types_with_comments.compact.ron
View File

@ -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: [

75
naga/tests/out/ir/wgsl-types_with_comments.ron
View File

@ -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.",

24
naga/tests/out/spv/wgsl-runtime-array-in-unused-struct.spvasm
View File

@ -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

6
naga/tests/out/wgsl/glsl-896-push-constant.frag.wgsl
View File

@ -1,9 +1,3 @@
struct PushConstants {
example: f32,
}
var<push_constant> c: PushConstants;
fn main_1() {
return;
}

51
naga/tests/out/wgsl/glsl-900-implicit-conversions.frag.wgsl
View File

@ -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<f32>) {
var v_1: vec3<f32>;
v_1 = v;
return;
}
fn implicit_dims_1(v_2: vec2<f32>) {
var v_3: vec2<f32>;
v_3 = v_2;
return;
}
fn implicit_dims_2(v_4: vec3<f32>) {
var v_5: vec3<f32>;
v_5 = v_4;
return;
}
fn implicit_dims_3(v_6: vec4<f32>) {
var v_7: vec4<f32>;
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;
}

8
naga/tests/out/wgsl/glsl-anonymous-entry-point-type.frag.wgsl
View File

@ -1,8 +1,4 @@
struct FragmentOutput {
x: f32,
}
struct FragmentOutput_1 {
@location(0) o_Target: vec4<f32>,
}
@ -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);
}

106
naga/tests/out/wgsl/glsl-bevy-pbr.frag.wgsl
View File

@ -9,10 +9,6 @@ struct DirectionalLight {
color: vec4<f32>,
}
struct CameraViewProj {
ViewProj: mat4x4<f32>,
}
struct CameraPosition {
CameraPos: vec4<f32>,
}
@ -57,28 +53,26 @@ var<private> v_WorldNormal_1: vec3<f32>;
var<private> v_Uv_1: vec2<f32>;
var<private> v_WorldTangent_1: vec4<f32>;
var<private> o_Target: vec4<f32>;
@group(0) @binding(0)
var<uniform> global: CameraViewProj;
@group(0) @binding(1)
var<uniform> global_1: CameraPosition;
var<uniform> global: CameraPosition;
@group(1) @binding(0)
var<uniform> global_2: Lights;
var<uniform> global_1: Lights;
@group(3) @binding(0)
var<uniform> global_3: StandardMaterial_base_color;
var<uniform> global_2: StandardMaterial_base_color;
@group(3) @binding(1)
var StandardMaterial_base_color_texture: texture_2d<f32>;
@group(3) @binding(2)
var StandardMaterial_base_color_texture_sampler: sampler;
@group(3) @binding(3)
var<uniform> global_4: StandardMaterial_roughness;
var<uniform> global_3: StandardMaterial_roughness;
@group(3) @binding(4)
var<uniform> global_5: StandardMaterial_metallic;
var<uniform> global_4: StandardMaterial_metallic;
@group(3) @binding(5)
var StandardMaterial_metallic_roughness_texture: texture_2d<f32>;
@group(3) @binding(6)
var StandardMaterial_metallic_roughness_texture_sampler: sampler;
@group(3) @binding(7)
var<uniform> global_6: StandardMaterial_reflectance;
var<uniform> global_5: StandardMaterial_reflectance;
@group(3) @binding(8)
var StandardMaterial_normal_map: texture_2d<f32>;
@group(3) @binding(9)
@ -88,7 +82,7 @@ var StandardMaterial_occlusion_texture: texture_2d<f32>;
@group(3) @binding(11)
var StandardMaterial_occlusion_texture_sampler: sampler;
@group(3) @binding(12)
var<uniform> global_7: StandardMaterial_emissive;
var<uniform> global_6: StandardMaterial_emissive;
@group(3) @binding(13)
var StandardMaterial_emissive_texture: texture_2d<f32>;
@group(3) @binding(14)
@ -353,32 +347,6 @@ fn perceptualRoughnessToRoughness(perceptualRoughness: f32) -> f32 {
return (_e7 * _e8);
}
fn reinhard(color: vec3<f32>) -> vec3<f32> {
var color_1: vec3<f32>;
color_1 = color;
let _e2 = color_1;
let _e5 = color_1;
return (_e2 / (vec3(1f) + _e5));
}
fn reinhard_extended(color_2: vec3<f32>, max_white: f32) -> vec3<f32> {
var color_3: vec3<f32>;
var max_white_1: f32;
var numerator: vec3<f32>;
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>) -> f32 {
var v_2: vec3<f32>;
@ -403,50 +371,24 @@ fn change_luminance(c_in: vec3<f32>, l_out: f32) -> vec3<f32> {
return (_e7 * (_e8 / _e9));
}
fn reinhard_luminance(color_4: vec3<f32>) -> vec3<f32> {
var color_5: vec3<f32>;
fn reinhard_luminance(color: vec3<f32>) -> vec3<f32> {
var color_1: vec3<f32>;
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<f32>, max_white_l: f32) -> vec3<f32> {
var color_7: vec3<f32>;
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<f32>, V_1: vec3<f32>, R: vec3<f32>, F0_: vec3<f32>, diffuseColor: vec3<f32>) -> vec3<f32> {
var light_1: PointLight;
var roughness_9: f32;
@ -662,7 +604,7 @@ fn main_1() {
var diffuse_ambient: vec3<f32>;
var specular_ambient: vec3<f32>;
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;

2
naga/tests/out/wgsl/glsl-global-constant-array.frag.wgsl
View File

@ -1,7 +1,5 @@
const array_: array<f32, 2> = array<f32, 2>(1f, 2f);
var<private> i: u32;
fn main_1() {
var local: array<f32, 2> = array_;

2
naga/tests/out/wgsl/glsl-samplers.frag.wgsl
View File

@ -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<f32>;
@group(1) @binding(17)
var sampShadow: sampler_comparison;
@group(0) @binding(18)

1
wgpu-hal/src/dx12/device.rs
View File

@ -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:?}")))?;

1
wgpu-hal/src/gles/device.rs
View File

@ -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| {

1
wgpu-hal/src/metal/device.rs
View File

@ -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)))?;

1
wgpu-hal/src/vulkan/device.rs
View File

@ -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| {