diff --git a/naga/src/compact/expressions.rs b/naga/src/compact/expressions.rs index 78c3f0de7e..df4f173ce1 100644 --- a/naga/src/compact/expressions.rs +++ b/naga/src/compact/expressions.rs @@ -3,6 +3,7 @@ use crate::arena::{Arena, Handle}; pub struct ExpressionTracer<'tracer> { pub constants: &'tracer Arena, + pub overrides: &'tracer Arena, /// The arena in which we are currently tracing expressions. pub expressions: &'tracer Arena, @@ -13,6 +14,9 @@ pub struct ExpressionTracer<'tracer> { /// The used map for `constants`. pub constants_used: &'tracer mut HandleSet, + /// The used map for `overrides`. + pub overrides_used: &'tracer mut HandleSet, + /// The used set for `arena`. /// /// This points to whatever arena holds the expressions we are @@ -78,25 +82,32 @@ impl ExpressionTracer<'_> { | Ex::SubgroupBallotResult | Ex::RayQueryProceedResult => {} + // Expressions can refer to constants and overrides, which can refer + // in turn to expressions, which complicates our nice one-pass + // algorithm. But since constants and overrides don't refer to each + // other directly, only via expressions, we can get around this by + // looking *through* each constant/override and marking its + // initializer expression as used immediately. Since `expr` refers + // to the constant/override, which then refers to the initializer, + // the initializer must precede `expr` in the arena, so we know we + // have yet to visit the initializer, so it's not too late to mark + // it. Ex::Constant(handle) => { self.constants_used.insert(handle); - // Constants and expressions are mutually recursive, which - // complicates our nice one-pass algorithm. However, since - // constants don't refer to each other, we can get around - // this by looking *through* each constant and marking its - // initializer as used. Since `expr` refers to the constant, - // and the constant refers to the initializer, it must - // precede `expr` in the arena. let init = self.constants[handle].init; match self.global_expressions_used { Some(ref mut used) => used.insert(init), None => self.expressions_used.insert(init), }; } - Ex::Override(_) => { - // All overrides are considered used by definition. We mark - // their types and initialization expressions as used in - // `compact::compact`, so we have no more work to do here. + Ex::Override(handle) => { + self.overrides_used.insert(handle); + if let Some(init) = self.overrides[handle].init { + match self.global_expressions_used { + Some(ref mut used) => used.insert(init), + None => self.expressions_used.insert(init), + }; + } } Ex::ZeroValue(ty) => { self.types_used.insert(ty); @@ -256,11 +267,9 @@ impl ModuleMap { | Ex::SubgroupBallotResult | Ex::RayQueryProceedResult => {} - // All overrides are retained, so their handles never change. - Ex::Override(_) => {} - // Expressions that contain handles that need to be adjusted. Ex::Constant(ref mut constant) => self.constants.adjust(constant), + Ex::Override(ref mut r#override) => self.overrides.adjust(r#override), Ex::ZeroValue(ref mut ty) => self.types.adjust(ty), Ex::Compose { ref mut ty, diff --git a/naga/src/compact/functions.rs b/naga/src/compact/functions.rs index d523c1889f..b37edd7271 100644 --- a/naga/src/compact/functions.rs +++ b/naga/src/compact/functions.rs @@ -4,9 +4,11 @@ use super::{FunctionMap, ModuleMap}; pub struct FunctionTracer<'a> { pub function: &'a crate::Function, pub constants: &'a crate::Arena, + pub overrides: &'a crate::Arena, pub types_used: &'a mut HandleSet, pub constants_used: &'a mut HandleSet, + pub overrides_used: &'a mut HandleSet, pub global_expressions_used: &'a mut HandleSet, /// Function-local expressions used. @@ -47,10 +49,12 @@ impl FunctionTracer<'_> { fn as_expression(&mut self) -> super::expressions::ExpressionTracer { super::expressions::ExpressionTracer { constants: self.constants, + overrides: self.overrides, expressions: &self.function.expressions, types_used: self.types_used, constants_used: self.constants_used, + overrides_used: self.overrides_used, expressions_used: &mut self.expressions_used, global_expressions_used: Some(&mut self.global_expressions_used), } diff --git a/naga/src/compact/mod.rs b/naga/src/compact/mod.rs index cfa2e89a03..930f7a39c9 100644 --- a/naga/src/compact/mod.rs +++ b/naga/src/compact/mod.rs @@ -45,21 +45,28 @@ pub fn compact(module: &mut crate::Module) { } // 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. + log::trace!("tracing named constants"); for (handle, constant) in module.constants.iter() { if constant.name.is_some() { + log::trace!("tracing constant {:?}", constant.name.as_ref().unwrap()); module_tracer.constants_used.insert(handle); module_tracer.global_expressions_used.insert(constant.init); } } - // We treat all overrides as used by definition. - for (_, r#override) in module.overrides.iter() { - module_tracer.types_used.insert(r#override.ty); - if let Some(init) = r#override.init { - module_tracer.global_expressions_used.insert(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); + if let Some(init) = r#override.init { + module_tracer.global_expressions_used.insert(init); + } } } @@ -110,6 +117,15 @@ pub fn compact(module: &mut crate::Module) { } } + // Overrides' initializers are taken care of already, because + // expression tracing sees through overrides. But we still need to + // note type usage. + for (handle, r#override) in module.overrides.iter() { + if module_tracer.overrides_used.contains(handle) { + module_tracer.types_used.insert(r#override.ty); + } + } + // Treat all named types as used. for (handle, ty) in module.types.iter() { log::trace!("tracing type {:?}, name {:?}", handle, ty.name); @@ -167,14 +183,19 @@ pub fn compact(module: &mut crate::Module) { } }); - // Adjust override types and initializers. + // Drop unused overrides in place, reusing existing storage. log::trace!("adjusting overrides"); - for (_, r#override) in module.overrides.iter_mut() { - module_map.types.adjust(&mut r#override.ty); - if let Some(init) = r#override.init.as_mut() { - module_map.global_expressions.adjust(init); + module.overrides.retain_mut(|handle, r#override| { + if module_map.overrides.used(handle) { + module_map.types.adjust(&mut r#override.ty); + if let Some(ref mut init) = r#override.init { + module_map.global_expressions.adjust(init); + } + true + } else { + false } - } + }); // Adjust workgroup_size_overrides log::trace!("adjusting workgroup_size_overrides"); @@ -223,6 +244,7 @@ struct ModuleTracer<'module> { module: &'module crate::Module, types_used: HandleSet, constants_used: HandleSet, + overrides_used: HandleSet, global_expressions_used: HandleSet, } @@ -232,6 +254,7 @@ impl<'module> ModuleTracer<'module> { module, types_used: HandleSet::for_arena(&module.types), constants_used: HandleSet::for_arena(&module.constants), + overrides_used: HandleSet::for_arena(&module.overrides), global_expressions_used: HandleSet::for_arena(&module.global_expressions), } } @@ -275,15 +298,16 @@ impl<'module> ModuleTracer<'module> { previous = std::cmp::max( previous, match ty.inner { - crate::TypeInner::Array { - base: _, - size: crate::ArraySize::Pending(crate::PendingArraySize::Expression(expr)), - stride: _, - } - | crate::TypeInner::BindingArray { - base: _, - size: crate::ArraySize::Pending(crate::PendingArraySize::Expression(expr)), - } => Some(expr), + crate::TypeInner::Array { size, .. } + | crate::TypeInner::BindingArray { size, .. } => match size { + crate::ArraySize::Constant(_) | crate::ArraySize::Dynamic => None, + crate::ArraySize::Pending(pending) => match pending { + crate::PendingArraySize::Expression(handle) => Some(handle), + crate::PendingArraySize::Override(handle) => { + self.module.overrides[handle].init + } + }, + }, _ => None, }, ); @@ -324,18 +348,22 @@ impl<'module> ModuleTracer<'module> { fn as_type(&mut self) -> types::TypeTracer { types::TypeTracer { + overrides: &self.module.overrides, types_used: &mut self.types_used, expressions_used: &mut self.global_expressions_used, + overrides_used: &mut self.overrides_used, } } fn as_const_expression(&mut self) -> expressions::ExpressionTracer { expressions::ExpressionTracer { - expressions: &self.module.global_expressions, constants: &self.module.constants, + overrides: &self.module.overrides, + expressions: &self.module.global_expressions, types_used: &mut self.types_used, constants_used: &mut self.constants_used, expressions_used: &mut self.global_expressions_used, + overrides_used: &mut self.overrides_used, global_expressions_used: None, } } @@ -347,8 +375,10 @@ impl<'module> ModuleTracer<'module> { FunctionTracer { function, constants: &self.module.constants, + overrides: &self.module.overrides, types_used: &mut self.types_used, constants_used: &mut self.constants_used, + overrides_used: &mut self.overrides_used, global_expressions_used: &mut self.global_expressions_used, expressions_used: HandleSet::for_arena(&function.expressions), } @@ -358,6 +388,7 @@ impl<'module> ModuleTracer<'module> { struct ModuleMap { types: HandleMap, constants: HandleMap, + overrides: HandleMap, global_expressions: HandleMap, } @@ -366,6 +397,7 @@ impl From> for ModuleMap { ModuleMap { types: HandleMap::from_set(used.types_used), constants: HandleMap::from_set(used.constants_used), + overrides: HandleMap::from_set(used.overrides_used), global_expressions: HandleMap::from_set(used.global_expressions_used), } } @@ -514,3 +546,337 @@ fn type_expression_interdependence() { compact(&mut module); assert!(cmp_modules(&module, &untouched)); } + +#[test] +fn array_length_override() { + let mut module: crate::Module = Default::default(); + let ty_bool = module.types.insert( + crate::Type { + name: None, + inner: crate::TypeInner::Scalar(crate::Scalar::BOOL), + }, + crate::Span::default(), + ); + let ty_u32 = module.types.insert( + crate::Type { + name: None, + inner: crate::TypeInner::Scalar(crate::Scalar::U32), + }, + crate::Span::default(), + ); + let one = module.global_expressions.append( + crate::Expression::Literal(crate::Literal::U32(1)), + crate::Span::default(), + ); + let _unused_override = module.overrides.append( + crate::Override { + name: None, + id: Some(40), + ty: ty_u32, + init: None, + }, + crate::Span::default(), + ); + let o = module.overrides.append( + crate::Override { + name: None, + id: Some(42), + ty: ty_u32, + init: Some(one), + }, + crate::Span::default(), + ); + let _ty_array = module.types.insert( + crate::Type { + name: Some("array".to_string()), + inner: crate::TypeInner::Array { + base: ty_bool, + size: crate::ArraySize::Pending(crate::PendingArraySize::Override(o)), + stride: 4, + }, + }, + crate::Span::default(), + ); + + let mut validator = super::valid::Validator::new( + super::valid::ValidationFlags::all(), + super::valid::Capabilities::all(), + ); + + assert!(validator.validate(&module).is_ok()); + compact(&mut module); + assert!(validator.validate(&module).is_ok()); +} + +/// Test mutual references between types and expressions via override +/// lengths. +#[test] +fn array_length_override_mutual() { + use crate::Expression as Ex; + use crate::Scalar as Sc; + use crate::TypeInner as Ti; + + let nowhere = crate::Span::default(); + let mut module = crate::Module::default(); + let ty_u32 = module.types.insert( + crate::Type { + name: None, + inner: Ti::Scalar(Sc::U32), + }, + nowhere, + ); + + // This type is only referred to by the override's init + // expression, so if we visit that too early, this type will be + // removed incorrectly. + let ty_i32 = module.types.insert( + crate::Type { + name: None, + inner: Ti::Scalar(Sc::I32), + }, + nowhere, + ); + + // An override that the other override's init can refer to. + let first_override = module.overrides.append( + crate::Override { + name: None, // so it is not considered used by definition + id: Some(41), + ty: ty_i32, + init: None, + }, + nowhere, + ); + + // Initializer expression for the override: + // + // (first_override + 0) as u32 + // + // The `first_override` makes it an override expression; the `0` + // gets a use of `ty_i32` in there; and the `as` makes it match + // the type of `second_override` without actually making + // `second_override` point at `ty_i32` directly. + let first_override_expr = module + .global_expressions + .append(Ex::Override(first_override), nowhere); + let zero = module + .global_expressions + .append(Ex::ZeroValue(ty_i32), nowhere); + let sum = module.global_expressions.append( + Ex::Binary { + op: crate::BinaryOperator::Add, + left: first_override_expr, + right: zero, + }, + nowhere, + ); + let init = module.global_expressions.append( + Ex::As { + expr: sum, + kind: crate::ScalarKind::Uint, + convert: None, + }, + nowhere, + ); + + // Override that serves as the array's length. + let second_override = module.overrides.append( + crate::Override { + name: None, // so it is not considered used by definition + id: Some(42), + ty: ty_u32, + init: Some(init), + }, + nowhere, + ); + + // Array type that uses the overload as its length. + // Since this is named, it is considered used by definition. + let _ty_array = module.types.insert( + crate::Type { + name: Some("delicious_array".to_string()), + inner: Ti::Array { + base: ty_u32, + size: crate::ArraySize::Pending(crate::PendingArraySize::Override(second_override)), + stride: 4, + }, + }, + nowhere, + ); + + let mut validator = super::valid::Validator::new( + super::valid::ValidationFlags::all(), + super::valid::Capabilities::all(), + ); + + assert!(validator.validate(&module).is_ok()); + compact(&mut module); + assert!(validator.validate(&module).is_ok()); +} + +#[test] +fn array_length_expression() { + let mut module: crate::Module = Default::default(); + let ty_u32 = module.types.insert( + crate::Type { + name: None, + inner: crate::TypeInner::Scalar(crate::Scalar::U32), + }, + crate::Span::default(), + ); + let _unused_zero = module.global_expressions.append( + crate::Expression::Literal(crate::Literal::U32(0)), + crate::Span::default(), + ); + let one = module.global_expressions.append( + crate::Expression::Literal(crate::Literal::U32(1)), + crate::Span::default(), + ); + let _ty_array = module.types.insert( + crate::Type { + name: Some("array".to_string()), + inner: crate::TypeInner::Array { + base: ty_u32, + size: crate::ArraySize::Pending(crate::PendingArraySize::Expression(one)), + stride: 4, + }, + }, + crate::Span::default(), + ); + + let mut validator = super::valid::Validator::new( + super::valid::ValidationFlags::all(), + super::valid::Capabilities::all(), + ); + + assert!(validator.validate(&module).is_ok()); + compact(&mut module); + assert!(validator.validate(&module).is_ok()); +} + +#[test] +fn global_expression_override() { + let mut module: crate::Module = Default::default(); + let ty_u32 = module.types.insert( + crate::Type { + name: None, + inner: crate::TypeInner::Scalar(crate::Scalar::U32), + }, + crate::Span::default(), + ); + + // This will only be retained if we trace the initializers + // of overrides referred to by `Expression::Override` + // in global expressions. + let expr1 = module.global_expressions.append( + crate::Expression::Literal(crate::Literal::U32(1)), + crate::Span::default(), + ); + + // This will only be traced via a global `Expression::Override`. + let o = module.overrides.append( + crate::Override { + name: None, + id: Some(42), + ty: ty_u32, + init: Some(expr1), + }, + crate::Span::default(), + ); + + // This is retained by _p. + let expr2 = module + .global_expressions + .append(crate::Expression::Override(o), crate::Span::default()); + + // Since this is named, it will be retained. + let _p = module.overrides.append( + crate::Override { + name: Some("p".to_string()), + id: None, + ty: ty_u32, + init: Some(expr2), + }, + crate::Span::default(), + ); + + let mut validator = super::valid::Validator::new( + super::valid::ValidationFlags::all(), + super::valid::Capabilities::all(), + ); + + assert!(validator.validate(&module).is_ok()); + compact(&mut module); + assert!(validator.validate(&module).is_ok()); +} + +#[test] +fn local_expression_override() { + let mut module: crate::Module = Default::default(); + let ty_u32 = module.types.insert( + crate::Type { + name: None, + inner: crate::TypeInner::Scalar(crate::Scalar::U32), + }, + crate::Span::default(), + ); + + // This will only be retained if we trace the initializers + // of overrides referred to by `Expression::Override` in a function. + let expr1 = module.global_expressions.append( + crate::Expression::Literal(crate::Literal::U32(1)), + crate::Span::default(), + ); + + // This will be removed by compaction. + let _unused_override = module.overrides.append( + crate::Override { + name: None, + id: Some(41), + ty: ty_u32, + init: None, + }, + crate::Span::default(), + ); + + // This will only be traced via an `Expression::Override` in a function. + let o = module.overrides.append( + crate::Override { + name: None, + id: Some(42), + ty: ty_u32, + init: Some(expr1), + }, + crate::Span::default(), + ); + + let mut fun = crate::Function { + result: Some(crate::FunctionResult { + ty: ty_u32, + binding: None, + }), + ..crate::Function::default() + }; + + // This is used by the `Return` statement. + let o_expr = fun + .expressions + .append(crate::Expression::Override(o), crate::Span::default()); + fun.body.push( + crate::Statement::Return { + value: Some(o_expr), + }, + crate::Span::default(), + ); + + module.functions.append(fun, crate::Span::default()); + + let mut validator = super::valid::Validator::new( + super::valid::ValidationFlags::all(), + super::valid::Capabilities::all(), + ); + + assert!(validator.validate(&module).is_ok()); + compact(&mut module); + assert!(validator.validate(&module).is_ok()); +} diff --git a/naga/src/compact/types.rs b/naga/src/compact/types.rs index 068d251b87..cde1352d39 100644 --- a/naga/src/compact/types.rs +++ b/naga/src/compact/types.rs @@ -2,8 +2,10 @@ use super::{HandleSet, ModuleMap}; use crate::Handle; pub struct TypeTracer<'a> { + pub overrides: &'a crate::Arena, pub types_used: &'a mut HandleSet, pub expressions_used: &'a mut HandleSet, + pub overrides_used: &'a mut HandleSet, } impl TypeTracer<'_> { @@ -24,23 +26,29 @@ impl TypeTracer<'_> { // Types that do contain handles. Ti::Array { base, - size: crate::ArraySize::Pending(crate::PendingArraySize::Expression(expr)), + size, stride: _, } - | Ti::BindingArray { - base, - size: crate::ArraySize::Pending(crate::PendingArraySize::Expression(expr)), - } => { - self.expressions_used.insert(expr); + | Ti::BindingArray { base, size } => { self.types_used.insert(base); + match size { + crate::ArraySize::Pending(pending) => match pending { + crate::PendingArraySize::Expression(expr) => { + self.expressions_used.insert(expr); + } + crate::PendingArraySize::Override(handle) => { + self.overrides_used.insert(handle); + let r#override = &self.overrides[handle]; + self.types_used.insert(r#override.ty); + if let Some(expr) = r#override.init { + self.expressions_used.insert(expr); + } + } + }, + crate::ArraySize::Constant(_) | crate::ArraySize::Dynamic => {} + } } - Ti::Pointer { base, space: _ } - | Ti::Array { - base, - size: _, - stride: _, - } - | Ti::BindingArray { base, size: _ } => { + Ti::Pointer { base, space: _ } => { self.types_used.insert(base); } Ti::Struct { @@ -79,13 +87,24 @@ impl ModuleMap { ref mut base, ref mut size, stride: _, + } + | Ti::BindingArray { + ref mut base, + ref mut size, } => { adjust(base); - if let crate::ArraySize::Pending(crate::PendingArraySize::Expression( - ref mut size_expr, - )) = *size - { - self.global_expressions.adjust(size_expr); + match *size { + crate::ArraySize::Pending(crate::PendingArraySize::Expression( + ref mut size_expr, + )) => { + self.global_expressions.adjust(size_expr); + } + crate::ArraySize::Pending(crate::PendingArraySize::Override( + ref mut r#override, + )) => { + self.overrides.adjust(r#override); + } + crate::ArraySize::Constant(_) | crate::ArraySize::Dynamic => {} } } Ti::Struct { @@ -96,12 +115,6 @@ impl ModuleMap { self.types.adjust(&mut member.ty); } } - Ti::BindingArray { - ref mut base, - size: _, - } => { - adjust(base); - } }; } }