diff --git a/cli/src/main.rs b/cli/src/main.rs index 56306a5016..b1a863ca24 100644 --- a/cli/src/main.rs +++ b/cli/src/main.rs @@ -335,6 +335,9 @@ fn run() -> Result<(), Box> { "metal" => { use naga::back::msl; + let mut options = params.msl.clone(); + options.bounds_check_policies = params.bounds_check_policies; + let pipeline_options = msl::PipelineOptions::default(); let (msl, _) = msl::write_string( &module, @@ -342,7 +345,7 @@ fn run() -> Result<(), Box> { "Generating metal output requires validation to \ succeed, and it failed in a previous step", ))?, - ¶ms.msl, + &options, &pipeline_options, ) .unwrap_pretty(); diff --git a/src/back/msl/mod.rs b/src/back/msl/mod.rs index ca7b9682c0..4295ae0e95 100644 --- a/src/back/msl/mod.rs +++ b/src/back/msl/mod.rs @@ -23,7 +23,7 @@ For the result type, if it's a structure, we re-compose it with a temporary valu holding the result. !*/ -use crate::{arena::Handle, valid::ModuleInfo}; +use crate::{arena::Handle, proc::index, valid::ModuleInfo}; use std::{ fmt::{Error as FmtError, Write}, ops, @@ -177,6 +177,9 @@ pub struct Options { pub spirv_cross_compatibility: bool, /// Don't panic on missing bindings, instead generate invalid MSL. pub fake_missing_bindings: bool, + /// Bounds checking policies. + #[cfg_attr(feature = "deserialize", serde(default))] + pub bounds_check_policies: index::BoundsCheckPolicies, } impl Default for Options { @@ -187,6 +190,7 @@ impl Default for Options { inline_samplers: Vec::new(), spirv_cross_compatibility: false, fake_missing_bindings: true, + bounds_check_policies: index::BoundsCheckPolicies::default(), } } } diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index df8f388aa3..393cd8b4e5 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -2,9 +2,11 @@ use super::{sampler as sm, Error, LocationMode, Options, PipelineOptions, Transl use crate::{ arena::Handle, back, + proc::index, proc::{self, NameKey, TypeResolution}, valid, FastHashMap, FastHashSet, }; +use bit_set::BitSet; use std::{ fmt::{Display, Error as FmtError, Formatter, Write}, iter, @@ -450,12 +452,35 @@ struct ExpressionContext<'a> { info: &'a valid::FunctionInfo, module: &'a crate::Module, pipeline_options: &'a PipelineOptions, + policies: index::BoundsCheckPolicies, + + /// A bitset containing the `Expression` handle indexes of expressions used + /// as indices in `ReadZeroSkipWrite`-policy accesses. These may need to be + /// cached in temporary variables. See `index::find_checked_indexes` for + /// details. + guarded_indices: BitSet, } impl<'a> ExpressionContext<'a> { fn resolve_type(&self, handle: Handle) -> &'a crate::TypeInner { self.info[handle].ty.inner_with(&self.module.types) } + + fn choose_bounds_check_policy( + &self, + pointer: Handle, + ) -> index::BoundsCheckPolicy { + self.policies + .choose_policy(pointer, &self.module.types, self.info) + } + + fn access_needs_check( + &self, + base: Handle, + index: index::GuardedIndex, + ) -> Option { + index::access_needs_check(base, index, self.module, self.function, self.info) + } } struct StatementContext<'a> { @@ -671,11 +696,13 @@ impl Writer { Ok(()) } - /// Write the length of the dynamically sized array at the end of `handle`. + /// Write the maximum valid index of the dynamically sized array at the end of `handle`. /// - /// `expr` must be the handle of a global variable whose final member is a dynamically - /// sized array. - fn put_dynamic_array_length( + /// The 'maximum valid index' is simply one less than the array's length. + /// + /// `handle` must be the handle of a global variable whose final member is a + /// dynamically sized array. + fn put_dynamic_array_max_index( &mut self, handle: Handle, context: &ExpressionContext, @@ -715,7 +742,7 @@ impl Writer { // prevent that. write!( self.out, - "(1 + (_buffer_sizes.size{idx} - {offset} - {span}) / {stride})", + "(_buffer_sizes.size{idx} - {offset} - {span}) / {stride}", idx = handle.index(), offset = offset, span = span, @@ -743,6 +770,17 @@ impl Writer { Ok(()) } + /// Emit code for the expression `expr_handle`. + /// + /// The `is_scoped` argument is true if the surrounding operators have the + /// precedence of the comma operator, or lower. So, for example: + /// + /// - Pass `true` for `is_scoped` when writing function arguments, an + /// expression statement, an initializer expression, or anything already + /// wrapped in parenthesis. + /// + /// - Pass `false` if it is an operand of a `?:` operator, a `[]`, or really + /// almost anything else. fn put_expression( &mut self, expr_handle: Handle, @@ -763,62 +801,30 @@ impl Writer { let expression = &context.function.expressions[expr_handle]; log::trace!("expression {:?} = {:?}", expr_handle, expression); match *expression { - crate::Expression::Access { base, index } => { - let accessing_wrapped_array = - match *context.info[base].ty.inner_with(&context.module.types) { - crate::TypeInner::Array { .. } => true, - crate::TypeInner::Pointer { - base: pointer_base, .. - } => match context.module.types[pointer_base].inner { - crate::TypeInner::Array { - size: crate::ArraySize::Constant(_), - .. - } => true, - _ => false, - }, - _ => false, - }; + crate::Expression::Access { .. } | crate::Expression::AccessIndex { .. } => { + // This is an acceptable place to generate a `ReadZeroSkipWrite` check. + // Since `put_bounds_checks` and `put_access_chain` handle an entire + // access chain at a time, recursing back through `put_expression` only + // for index expressions and the base object, we will never see intermediate + // `Access` or `AccessIndex` expressions here. + let policy = context.choose_bounds_check_policy(expr_handle); + if policy == index::BoundsCheckPolicy::ReadZeroSkipWrite + && self.put_bounds_checks( + expr_handle, + context, + back::Level(0), + if is_scoped { "" } else { "(" }, + )? + { + write!(self.out, " ? ")?; + self.put_access_chain(expr_handle, policy, context)?; + write!(self.out, " : 0")?; - self.put_expression(base, context, false)?; - if accessing_wrapped_array { - write!(self.out, ".{}", WRAPPED_ARRAY_FIELD)?; - } - write!(self.out, "[")?; - self.put_expression(index, context, true)?; - write!(self.out, "]")?; - } - crate::Expression::AccessIndex { base, index } => { - self.put_expression(base, context, false)?; - let base_res = &context.info[base].ty; - let mut resolved = base_res.inner_with(&context.module.types); - let base_ty_handle = match *resolved { - crate::TypeInner::Pointer { base, class: _ } => { - resolved = &context.module.types[base].inner; - Some(base) - } - _ => base_res.handle(), - }; - match *resolved { - crate::TypeInner::Struct { .. } => { - let base_ty = base_ty_handle.unwrap(); - let name = &self.names[&NameKey::StructMember(base_ty, index)]; - write!(self.out, ".{}", name)?; - } - crate::TypeInner::ValuePointer { .. } | crate::TypeInner::Vector { .. } => { - write!(self.out, ".{}", back::COMPONENTS[index as usize])?; - } - crate::TypeInner::Array { - size: crate::ArraySize::Constant(_), - .. - } => { - write!(self.out, ".{}[{}]", WRAPPED_ARRAY_FIELD, index)?; - } - crate::TypeInner::Array { .. } | crate::TypeInner::Matrix { .. } => { - write!(self.out, "[{}]", index)?; - } - _ => { - // unexpected indexing, should fail validation + if !is_scoped { + write!(self.out, ")")?; } + } else { + self.put_access_chain(expr_handle, policy, context)?; } } crate::Expression::Constant(handle) => { @@ -1220,20 +1226,326 @@ impl Writer { _ => return Err(Error::Validation), }; - self.put_dynamic_array_length(global, context)?; + if !is_scoped { + write!(self.out, "(")?; + } + write!(self.out, "1 + ")?; + self.put_dynamic_array_max_index(global, context)?; + if !is_scoped { + write!(self.out, ")")?; + } + } + } + Ok(()) + } + + /// Write a `GuardedIndex` as a Metal expression. + fn put_index( + &mut self, + index: index::GuardedIndex, + context: &ExpressionContext, + is_scoped: bool, + ) -> BackendResult { + match index { + index::GuardedIndex::Expression(expr) => { + self.put_expression(expr, context, is_scoped)? } + index::GuardedIndex::Known(value) => write!(self.out, "{}", value)?, } Ok(()) } + /// Emit an index bounds check condition for `chain`, if required. + /// + /// `chain` is a subtree of `Access` and `AccessIndex` expressions, + /// operating either on a pointer to a value, or on a value directly. If we cannot + /// statically determine that all indexing operations in `chain` are within + /// bounds, then write a conditional expression to check them dynamically, + /// and return true. All accesses in the chain are checked by the generated + /// expression. + /// + /// This assumes that the [`BoundsCheckPolicy`] for `chain` is [`ReadZeroSkipWrite`]. + /// + /// The text written is of the form: + /// + /// ```ignore + /// {level}{prefix}metal::uint(i) < 4 && metal::uint(j) < 10 + /// ``` + /// + /// where `{level}` and `{prefix}` are the arguments to this function. For [`Store`] + /// statements, presumably these arguments start an indented `if` statement; for + /// [`Load`] expressions, the caller is probably building up a ternary `?:` + /// expression. In either case, what is written is not a complete syntactic structure + /// in its own right, and the caller will have to finish it off if we return `true`. + /// + /// If no expression is written, return false. + /// + /// [`BoundsCheckPolicy`]: index::BoundsCheckPolicy + /// [`ReadZeroSkipWrite`]: index::BoundsCheckPolicy::ReadZeroSkipWrite + /// [`Store`]: crate::Statement::Store + /// [`Load`]: crate::Expression::Load + #[allow(unused_variables)] + fn put_bounds_checks( + &mut self, + mut chain: Handle, + context: &ExpressionContext, + level: back::Level, + prefix: &'static str, + ) -> Result { + let mut check_written = false; + + // Pull this out into a closure so that we can handle both `Access` and + // `AccessIndex` expressions using the same code, using `GuardedIndex` to + // represent the index, computed or constant. + let mut consider = |base, index| -> BackendResult { + if let Some(length) = context.access_needs_check(base, index) { + if check_written { + write!(self.out, " && ")?; + } else { + write!(self.out, "{}{}", level, prefix)?; + check_written = true; + } + + // Check that the index falls within bounds. Do this with a single + // comparison, by casting the index to `uint` first, so that negative + // indices become large positive values. + write!(self.out, "{}::uint(", NAMESPACE)?; + self.put_index(index, context, true)?; + self.out.write_str(") < ")?; + match length { + index::IndexableLength::Known(value) => write!(self.out, "{}", value)?, + index::IndexableLength::Dynamic => { + let global = base + .originating_global(context.function) + .ok_or(Error::Validation)?; + write!(self.out, "1 + ")?; + self.put_dynamic_array_max_index(global, context)? + } + } + } + + Ok(()) + }; + + // Iterate over the access chain, handling each expression. + loop { + chain = match context.function.expressions[chain] { + crate::Expression::Access { base, index } => { + consider(base, index::GuardedIndex::Expression(index))?; + base + } + crate::Expression::AccessIndex { base, index } => { + // Don't try to check indices into structs. Validation already took + // care of them, and index::needs_guard doesn't handle that case. + let mut base_inner = context.info[base].ty.inner_with(&context.module.types); + if let crate::TypeInner::Pointer { base, .. } = *base_inner { + base_inner = &context.module.types[base].inner; + } + match *base_inner { + crate::TypeInner::Struct { .. } => {} + _ => consider(base, index::GuardedIndex::Known(index))?, + } + base + } + _ => break, + } + } + + Ok(check_written) + } + + /// Write the access chain `chain`. + /// + /// `chain` is a subtree of [`Access`] and [`AccessIndex`] expressions, + /// operating either on a pointer to a value, or on a value directly. + /// + /// Generate bounds checks code only if `policy` is [`Restrict`]. The + /// [`ReadZeroSkipWrite`] policy requires checks before any accesses take place, so + /// that must be handled in the caller. + /// + /// Handle the entire chain, recursing back into `put_expression` only for index + /// expressions and the base expression that originates the pointer or composite value + /// being accessed. This allows `put_expression` to assume that any `Access` or + /// `AccessIndex` expressions it sees are the top of a chain, so it can emit + /// `ReadZeroSkipWrite` checks. + /// + /// [`Access`]: crate::Expression::Access + /// [`AccessIndex`]: crate::Expression::AccessIndex + /// [`Restrict`]: crate::proc::index::BoundsCheckPolicy::Restrict + /// [`ReadZeroSkipWrite`]: crate::proc::index::BoundsCheckPolicy::ReadZeroSkipWrite + fn put_access_chain( + &mut self, + chain: Handle, + policy: index::BoundsCheckPolicy, + context: &ExpressionContext, + ) -> BackendResult { + match context.function.expressions[chain] { + crate::Expression::Access { base, index } => { + let mut base_ty = context.info[base].ty.inner_with(&context.module.types); + + // Look through any pointers to see what we're really indexing. + if let crate::TypeInner::Pointer { base, class: _ } = *base_ty { + base_ty = &context.module.types[base].inner; + } + + self.put_subscripted_access_chain( + base, + base_ty, + index::GuardedIndex::Expression(index), + policy, + context, + )?; + } + crate::Expression::AccessIndex { base, index } => { + let base_resolution = &context.info[base].ty; + let mut base_ty = base_resolution.inner_with(&context.module.types); + let mut base_ty_handle = base_resolution.handle(); + + // Look through any pointers to see what we're really indexing. + if let crate::TypeInner::Pointer { base, class: _ } = *base_ty { + base_ty = &context.module.types[base].inner; + base_ty_handle = Some(base); + } + + // Handle structs and anything else that can use `.x` syntax here, so + // `put_subscripted_access_chain` won't have to handle the absurd case of + // indexing a struct with an expression. + match *base_ty { + crate::TypeInner::Struct { .. } => { + let base_ty = base_ty_handle.unwrap(); + self.put_access_chain(base, policy, context)?; + let name = &self.names[&NameKey::StructMember(base_ty, index)]; + write!(self.out, ".{}", name)?; + } + crate::TypeInner::ValuePointer { .. } | crate::TypeInner::Vector { .. } => { + self.put_access_chain(base, policy, context)?; + write!(self.out, ".{}", back::COMPONENTS[index as usize])?; + } + _ => { + self.put_subscripted_access_chain( + base, + base_ty, + index::GuardedIndex::Known(index), + policy, + context, + )?; + } + } + } + _ => self.put_expression(chain, context, false)?, + } + + Ok(()) + } + + /// Write a `[]`-style access of `base` by `index`. + /// + /// If `policy` is [`Restrict`], then generate code as needed to force all index + /// values within bounds. + /// + /// The `base_ty` argument must be the type we are actually indexing, like [`Array`] or + /// [`Vector`]. In other words, it's `base`'s type with any surrounding [`Pointer`] + /// removed. Our callers often already have this handy. + /// + /// This only emits `[]` expressions; it doesn't handle struct member accesses or + /// referencing vector components by name. + /// + /// [`Restrict`]: crate::proc::index::BoundsCheckPolicy::Restrict + /// [`Array`]: crate::TypeInner::Array + /// [`Vector`]: crate::TypeInner::Vector + /// [`Pointer`]: crate::TypeInner::Pointer + fn put_subscripted_access_chain( + &mut self, + base: Handle, + base_ty: &crate::TypeInner, + index: index::GuardedIndex, + policy: index::BoundsCheckPolicy, + context: &ExpressionContext, + ) -> BackendResult { + let accessing_wrapped_array = match *base_ty { + crate::TypeInner::Array { + size: crate::ArraySize::Constant(_), + .. + } => true, + _ => false, + }; + + self.put_access_chain(base, policy, context)?; + if accessing_wrapped_array { + write!(self.out, ".{}", WRAPPED_ARRAY_FIELD)?; + } + write!(self.out, "[")?; + + // Decide whether this index needs to be clamped to fall within range. + let restriction_needed = if policy == index::BoundsCheckPolicy::Restrict { + context.access_needs_check(base, index) + } else { + None + }; + if let Some(limit) = restriction_needed { + write!(self.out, "{}::min(unsigned(", NAMESPACE)?; + self.put_index(index, context, true)?; + write!(self.out, "), ")?; + match limit { + index::IndexableLength::Known(limit) => { + write!(self.out, "{}u", limit - 1)?; + } + index::IndexableLength::Dynamic => { + let global = base + .originating_global(context.function) + .ok_or(Error::Validation)?; + self.put_dynamic_array_max_index(global, context)?; + } + } + write!(self.out, ")")?; + } else { + self.put_index(index, context, true)?; + } + + write!(self.out, "]")?; + + Ok(()) + } + fn put_load( &mut self, pointer: Handle, context: &ExpressionContext, is_scoped: bool, + ) -> BackendResult { + // Since access chains never cross storage classes, we can just check the index + // bounds check policy once at the top. + let policy = context.choose_bounds_check_policy(pointer); + if policy == index::BoundsCheckPolicy::ReadZeroSkipWrite + && self.put_bounds_checks( + pointer, + context, + back::Level(0), + if is_scoped { "" } else { "(" }, + )? + { + write!(self.out, " ? ")?; + self.put_unchecked_load(pointer, policy, context)?; + write!(self.out, " : 0")?; + + if !is_scoped { + write!(self.out, ")")?; + } + } else { + self.put_unchecked_load(pointer, policy, context)?; + } + + Ok(()) + } + + fn put_unchecked_load( + &mut self, + pointer: Handle, + policy: index::BoundsCheckPolicy, + context: &ExpressionContext, ) -> BackendResult { // Because packed vectors such as `packed_float3` cannot be directly multipied by - // matrices, we wrap them with `float3` on load. + // matrices, we convert them to unpacked vectors like `float3` on load. let wrap_packed_vec_scalar_kind = match context.function.expressions[pointer] { crate::Expression::AccessIndex { base, index } => { let ty = match *context.resolve_type(base) { @@ -1264,7 +1576,7 @@ impl Writer { NAMESPACE, scalar_kind_string(scalar_kind) )?; - self.put_expression(pointer, context, true)?; + self.put_access_chain(pointer, policy, context)?; write!(self.out, ")")?; } else if is_atomic { write!( @@ -1272,13 +1584,13 @@ impl Writer { "{}::atomic_load_explicit({}", NAMESPACE, ATOMIC_REFERENCE )?; - self.put_expression(pointer, context, true)?; + self.put_access_chain(pointer, policy, context)?; write!(self.out, ", {}::memory_order_relaxed)", NAMESPACE)?; } else { // We don't do any dereferencing with `*` here as pointer arguments to functions // are done by `&` references and not `*` pointers. These do not need to be // dereferenced. - self.put_expression(pointer, context, is_scoped)?; + self.put_access_chain(pointer, policy, context)?; } Ok(()) @@ -1440,15 +1752,33 @@ impl Writer { } else if let Some(name) = context.expression.function.named_expressions.get(&handle) { - // Front end provides names for all variables at the start of writing. - // But we write them to step by step. We need to recache them - // Otherwise, we could accidentally write variable name instead of full expression. - // Also, we use sanitized names! It defense backend from generating variable with name from reserved keywords. + // The `crate::Function::named_expressions` table holds + // expressions that should be saved in temporaries once they + // are `Emit`ted. We only add them to `self.named_expressions` + // when we reach the `Emit` that covers them, so that we don't + // try to use their names before we've actually initialized + // the temporary that holds them. + // + // Don't assume the names in `named_expressions` are unique, + // or even valid. Use the `Namer`. Some(self.namer.call(name)) } else { - let min_ref_count = - context.expression.function.expressions[handle].bake_ref_count(); - if min_ref_count <= info.ref_count { + // If this expression is an index that we're going to first compare + // against a limit, and then actually use as an index, then we may + // want to cache it in a temporary, to avoid evaluating it twice. + let bake = + if context.expression.guarded_indices.contains(handle.index()) { + true + } else { + // Expressions whose reference count is above the + // threshold should always be stored in temporaries. + let min_ref_count = context.expression.function.expressions + [handle] + .bake_ref_count(); + min_ref_count <= info.ref_count + }; + + if bake { Some(format!("{}{}", back::BAKE_PREFIX, handle.index())) } else { None @@ -1730,6 +2060,28 @@ impl Writer { value: Handle, level: back::Level, context: &StatementContext, + ) -> BackendResult { + let policy = context.expression.choose_bounds_check_policy(pointer); + if policy == index::BoundsCheckPolicy::ReadZeroSkipWrite + && self.put_bounds_checks(pointer, &context.expression, level, "if (")? + { + writeln!(self.out, ") {{")?; + self.put_unchecked_store(pointer, value, policy, level.next(), context)?; + writeln!(self.out, "{}}}", level)?; + } else { + self.put_unchecked_store(pointer, value, policy, level, context)?; + } + + Ok(()) + } + + fn put_unchecked_store( + &mut self, + pointer: Handle, + value: Handle, + policy: index::BoundsCheckPolicy, + level: back::Level, + context: &StatementContext, ) -> BackendResult { let pointer_info = &context.expression.info[pointer]; let (array_size, is_atomic) = @@ -1753,7 +2105,7 @@ impl Writer { .to_array_length() .unwrap(); write!(self.out, "{}for(int _i=0; _i<{}; ++_i) ", level, size)?; - self.put_expression(pointer, &context.expression, true)?; + self.put_access_chain(pointer, policy, &context.expression)?; write!(self.out, ".{}[_i] = ", WRAPPED_ARRAY_FIELD)?; self.put_expression(value, &context.expression, true)?; writeln!(self.out, ".{}[_i];", WRAPPED_ARRAY_FIELD)?; @@ -1763,13 +2115,13 @@ impl Writer { "{}{}::atomic_store_explicit({}", level, NAMESPACE, ATOMIC_REFERENCE )?; - self.put_expression(pointer, &context.expression, true)?; + self.put_access_chain(pointer, policy, &context.expression)?; write!(self.out, ", ")?; self.put_expression(value, &context.expression, true)?; writeln!(self.out, ", {}::memory_order_relaxed);", NAMESPACE)?; } else { write!(self.out, "{}", level)?; - self.put_expression(pointer, &context.expression, true)?; + self.put_access_chain(pointer, policy, &context.expression)?; write!(self.out, " = ")?; self.put_expression(value, &context.expression, true)?; writeln!(self.out, ";")?; @@ -2216,11 +2568,16 @@ impl Writer { writeln!(self.out, ";")?; } + let guarded_indices = + index::find_checked_indexes(module, fun, fun_info, options.bounds_check_policies); + let context = StatementContext { expression: ExpressionContext { function: fun, origin: FunctionOrigin::Handle(fun_handle), info: fun_info, + policies: options.bounds_check_policies, + guarded_indices, module, pipeline_options, }, @@ -2638,11 +2995,16 @@ impl Writer { writeln!(self.out, ";")?; } + let guarded_indices = + index::find_checked_indexes(module, fun, fun_info, options.bounds_check_policies); + let context = StatementContext { expression: ExpressionContext { function: fun, origin: FunctionOrigin::EntryPoint(ep_index as _), info: fun_info, + policies: options.bounds_check_policies, + guarded_indices, module, pipeline_options, }, @@ -2721,7 +3083,7 @@ fn test_stack_size() { let stack_size = addresses.end - addresses.start; // check the size (in debug only) // last observed macOS value: 20528 (CI) - if !(15000..=25000).contains(&stack_size) { + if !(14000..=25000).contains(&stack_size) { panic!("`put_expression` stack size {} has changed!", stack_size); } } diff --git a/src/proc/index.rs b/src/proc/index.rs index 27110cb5a2..15381c9d23 100644 --- a/src/proc/index.rs +++ b/src/proc/index.rs @@ -3,6 +3,7 @@ use super::ProcError; use crate::valid; use crate::{Handle, UniqueArena}; +use bit_set::BitSet; /// How should code generated by Naga do bounds checks? /// @@ -33,7 +34,8 @@ use crate::{Handle, UniqueArena}; /// If you're using native checks like this, then having Naga inject its own /// checks as well would be redundant, and the `Unchecked` policy is /// appropriate. -#[derive(Clone, Copy, Debug)] +#[derive(Clone, Copy, Debug, Eq, Hash, PartialEq)] +#[cfg_attr(feature = "serialize", derive(serde::Serialize))] #[cfg_attr(feature = "deserialize", derive(serde::Deserialize))] pub enum BoundsCheckPolicy { /// Replace out-of-bounds indexes with some arbitrary in-bounds index. @@ -44,7 +46,12 @@ pub enum BoundsCheckPolicy { /// the last element, not the first.) Restrict, - /// Out-of-bounds reads return zero, and writes have no effect. + /// Out-of-bounds reads return zero, and writes have no effect. + /// + /// When applied to a chain of accesses, like `a[i][j].b[k]`, all index + /// expressions are evaluated, regardless of whether prior or later index + /// expressions were in bounds. But all the accesses per se are skipped + /// if any index is out of bounds. ReadZeroSkipWrite, /// Naga adds no checks to indexing operations. Generate the fastest code @@ -53,9 +60,10 @@ pub enum BoundsCheckPolicy { Unchecked, } -#[derive(Clone, Copy, Debug, Default)] -#[cfg_attr(feature = "deserialize", derive(serde::Deserialize))] /// Policies for injecting bounds checks during code generation. +#[derive(Clone, Copy, Debug, Default, Eq, Hash, PartialEq)] +#[cfg_attr(feature = "serialize", derive(serde::Serialize))] +#[cfg_attr(feature = "deserialize", derive(serde::Deserialize))] pub struct BoundsCheckPolicies { /// How should the generated code handle array, vector, or matrix indices /// that are out of range? @@ -119,26 +127,172 @@ impl Default for BoundsCheckPolicy { } impl BoundsCheckPolicies { - /// Determine which policy applies to a load or store of `pointer`. + /// Determine which policy applies to `access`. + /// + /// `access` is a subtree of `Access` and `AccessIndex` expressions, + /// operating either on a pointer to a value, or on a value directly. /// /// See the documentation for [`BoundsCheckPolicy`] for details about /// when each policy applies. pub fn choose_policy( &self, - pointer: Handle, + access: Handle, types: &UniqueArena, info: &valid::FunctionInfo, ) -> BoundsCheckPolicy { - let is_buffer = match info[pointer].ty.inner_with(types).pointer_class() { + match info[access].ty.inner_with(types).pointer_class() { Some(crate::StorageClass::Storage { access: _ }) - | Some(crate::StorageClass::Uniform) => true, - _ => false, - }; + | Some(crate::StorageClass::Uniform) => self.buffer, + // This covers other storage classes, but also accessing vectors and + // matrices by value, where no pointer is involved. + _ => self.index, + } + } +} + +/// An index that may be statically known, or may need to be computed at runtime. +/// +/// This enum lets us handle both [`Access`] and [`AccessIndex`] expressions +/// with the same code. +/// +/// [`Access`]: crate::Expression::Access +/// [`AccessIndex`]: crate::Expression::AccessIndex +#[derive(Clone, Copy, Debug)] +pub enum GuardedIndex { + Known(u32), + Expression(Handle), +} + +/// Build a set of expressions used as indices, to cache in temporary variables when +/// emitted. +/// +/// Given the bounds-check policies `policies`, construct a `BitSet` containing the handle +/// indices of all the expressions in `function` that are ever used as guarded indices +/// under the [`ReadZeroSkipWrite`] policy. The `module` argument must be the module to +/// which `function` belongs, and `info` should be that function's analysis results. +/// +/// Such index expressions will be used twice in the generated code: first for the +/// comparison to see if the index is in bounds, and then for the access itself, should +/// the comparison succeed. To avoid computing the expressions twice, they should be +/// cached in temporary variables. +/// +/// Why do we need to build such a set before processing a function's statements? Whether +/// an expression needs to be cached depends on whether it appears as the [`index`] +/// operand of any [`Access`] expression, and on the index bounds check policies that +/// apply to those accesses. But [`Emit`] statements just identify a range of expressions +/// by index; there's no good way to tell what an expression is used for. The only way to +/// do it is to just iterate over all the expressions looking for relevant `Access` +/// expressions --- which is what this function does. +/// +/// Simple expressions like variable loads and constants don't make sense to cache: it's +/// no better than just re-evaluating them. But constants are not covered by `Emit` +/// statements, and `Load`s are always cached to ensure they occur at the right time, so +/// we don't bother filtering them out from this set. +/// +/// [`ReadZeroSkipWrite`]: BoundsCheckPolicy::ReadZeroSkipWrite +/// [`index`]: crate::Expression::Access::index +/// [`Access`]: crate::Expression::Access +/// [`Emit`]: crate::Statement::Emit +pub fn find_checked_indexes( + module: &crate::Module, + function: &crate::Function, + info: &crate::valid::FunctionInfo, + policies: BoundsCheckPolicies, +) -> BitSet { + use crate::Expression as Ex; + + let mut guarded_indices = BitSet::new(); - if is_buffer { - self.buffer - } else { - self.index + // Don't bother scanning if we never need `ReadZeroSkipWrite`. + if policies.index == BoundsCheckPolicy::ReadZeroSkipWrite + || policies.buffer == BoundsCheckPolicy::ReadZeroSkipWrite + { + for (_handle, expr) in function.expressions.iter() { + // There's no need to handle `AccessIndex` expressions, as their + // indices never need to be cached. + if let Ex::Access { base, index } = *expr { + if policies.choose_policy(base, &module.types, info) + == BoundsCheckPolicy::ReadZeroSkipWrite + && access_needs_check( + base, + GuardedIndex::Expression(index), + module, + function, + info, + ) + .is_some() + { + guarded_indices.insert(index.index()); + } + } + } + } + + guarded_indices +} + +/// Determine whether `index` is statically known to be in bounds for `base`. +/// +/// If we can't be sure that the index is in bounds, return the limit within +/// which valid indices must fall. +/// +/// The return value is one of the following: +/// +/// - `Some(Known(n))` indicates that `n` is the largest valid index. +/// +/// - `Some(Computed(global))` indicates that the largest valid index is one +/// less than the length of the array that is the last member of the +/// struct held in `global`. +/// +/// - `None` indicates that the index need not be checked, either because it +/// is statically known to be in bounds, or because the applicable policy +/// is `Unchecked`. +/// +/// This function only handles subscriptable types: arrays, vectors, and +/// matrices. It does not handle struct member indices; those never require +/// run-time checks, so it's best to deal with them further up the call +/// chain. +pub fn access_needs_check( + base: Handle, + mut index: GuardedIndex, + module: &crate::Module, + function: &crate::Function, + info: &crate::valid::FunctionInfo, +) -> Option { + let base_inner = info[base].ty.inner_with(&module.types); + // Unwrap safety: `Err` here indicates unindexable base types and invalid + // length constants, but `access_needs_check` is only used by back ends, so + // validation should have caught those problems. + let length = base_inner.indexable_length(module).unwrap(); + index.try_resolve_to_constant(function, module); + if let (&GuardedIndex::Known(index), &IndexableLength::Known(length)) = (&index, &length) { + if index < length { + // Index is statically known to be in bounds, no check needed. + return None; + } + }; + + Some(length) +} + +impl GuardedIndex { + /// Make A `GuardedIndex::Known` from a `GuardedIndex::Expression` if possible. + /// + /// If the expression is a [`Constant`] whose value is a non-specialized, scalar + /// integer constant that can be converted to a `u32`, do so and return a + /// `GuardedIndex::Known`. Otherwise, return the `GuardedIndex::Expression` + /// unchanged. + /// + /// Return values that are already `Known` unchanged. + /// + /// [`Constant`]: crate::Expression::Constant + fn try_resolve_to_constant(&mut self, function: &crate::Function, module: &crate::Module) { + if let GuardedIndex::Expression(expr) = *self { + if let crate::Expression::Constant(handle) = function.expressions[expr] { + if let Some(value) = module.constants[handle].to_array_length() { + *self = GuardedIndex::Known(value); + } + } } } } @@ -149,6 +303,8 @@ impl crate::TypeInner { /// The `self` parameter should be a handle to a vector, matrix, or array /// type, a pointer to one of those, or a value pointer. Arrays may be /// fixed-size, dynamically sized, or sized by a specializable constant. + /// This function does not handle struct member references, as with + /// `AccessIndex`. /// /// The value returned is appropriate for bounds checks on subscripting. /// @@ -186,6 +342,7 @@ impl crate::TypeInner { /// /// This summarizes the length of vectors, matrices, and arrays in a way that is /// convenient for indexing and bounds-checking code. +#[derive(Debug)] pub enum IndexableLength { /// Values of this type always have the given number of elements. Known(u32), diff --git a/src/proc/mod.rs b/src/proc/mod.rs index 4c40d845d1..b2fc9c7575 100644 --- a/src/proc/mod.rs +++ b/src/proc/mod.rs @@ -1,6 +1,6 @@ //! Module processing functionality. -mod index; +pub mod index; mod layouter; mod namer; mod terminator; @@ -313,6 +313,33 @@ impl crate::Expression { } } +impl crate::Handle { + /// Return the global variable being accessed by a pointer expression. + /// + /// Assuming that `pointer` is a series of `Access` and `AccessIndex` + /// expressions that ultimately access some part of a `GlobalVariable`, + /// return a handle for that global. + /// + /// If the expression does not ultimately access a global variable, return + /// `None`. + pub fn originating_global( + mut self, + func: &crate::Function, + ) -> Option> { + loop { + self = match func.expressions[self] { + crate::Expression::Access { base, .. } => base, + crate::Expression::AccessIndex { base, .. } => base, + crate::Expression::GlobalVariable(handle) => return Some(handle), + crate::Expression::LocalVariable(_) => return None, + crate::Expression::FunctionArgument(_) => return None, + // There are no other expressions that produce pointer values. + _ => unreachable!(), + } + } + } +} + impl crate::SampleLevel { pub fn implicit_derivatives(&self) -> bool { match *self { diff --git a/tests/in/bounds-check-restrict.param.ron b/tests/in/bounds-check-restrict.param.ron new file mode 100644 index 0000000000..93c734c146 --- /dev/null +++ b/tests/in/bounds-check-restrict.param.ron @@ -0,0 +1,6 @@ +( + bounds_check_policies: ( + index: Restrict, + buffer: Restrict, + ), +) diff --git a/tests/in/bounds-check-restrict.wgsl b/tests/in/bounds-check-restrict.wgsl new file mode 100644 index 0000000000..6f4ae998d8 --- /dev/null +++ b/tests/in/bounds-check-restrict.wgsl @@ -0,0 +1,73 @@ +// Tests for `naga::back::BoundsCheckPolicy::Restrict`. + +[[block]] +struct Globals { + a: array; + v: vec4; + m: mat3x4; + d: array; +}; + +[[group(0), binding(0)]] var globals: Globals; + +fn index_array(i: i32) -> f32 { + return globals.a[i]; +} + +fn index_dynamic_array(i: i32) -> f32 { + return globals.d[i]; +} + +fn index_vector(i: i32) -> f32 { + return globals.v[i]; +} + +fn index_vector_by_value(v: vec4, i: i32) -> f32 { + return v[i]; +} + +fn index_matrix(i: i32) -> vec4 { + return globals.m[i]; +} + +fn index_twice(i: i32, j: i32) -> f32 { + return globals.m[i][j]; +} + +fn index_expensive(i: i32) -> f32 { + return globals.a[i32(sin(f32(i) / 100.0) * 100.0)]; +} + +fn index_in_bounds() -> f32 { + return globals.a[9] + globals.v[3] + globals.m[2][3]; +} + +fn set_array(i: i32, v: f32) { + globals.a[i] = v; +} + +fn set_dynamic_array(i: i32, v: f32) { + globals.d[i] = v; +} + +fn set_vector(i: i32, v: f32) { + globals.v[i] = v; +} + +fn set_matrix(i: i32, v: vec4) { + globals.m[i] = v; +} + +fn set_index_twice(i: i32, j: i32, v: f32) { + globals.m[i][j] = v; +} + +fn set_expensive(i: i32, v: f32) { + globals.a[i32(sin(f32(i) / 100.0) * 100.0)] = v; +} + +fn set_in_bounds(v: f32) { + globals.a[9] = v; + globals.v[3] = v; + globals.m[2][3] = v; +} diff --git a/tests/in/bounds-check-zero.wgsl b/tests/in/bounds-check-zero.wgsl index 51602cbd65..b4fbebcd78 100644 --- a/tests/in/bounds-check-zero.wgsl +++ b/tests/in/bounds-check-zero.wgsl @@ -5,6 +5,7 @@ struct Globals { a: array; v: vec4; m: mat3x4; + d: array; }; [[group(0), binding(0)]] var globals: Globals; @@ -13,6 +14,10 @@ fn index_array(i: i32) -> f32 { return globals.a[i]; } +fn index_dynamic_array(i: i32) -> f32 { + return globals.d[i]; +} + fn index_vector(i: i32) -> f32 { return globals.v[i]; } @@ -29,10 +34,22 @@ fn index_twice(i: i32, j: i32) -> f32 { return globals.m[i][j]; } +fn index_expensive(i: i32) -> f32 { + return globals.a[i32(sin(f32(i) / 100.0) * 100.0)]; +} + +fn index_in_bounds() -> f32 { + return globals.a[9] + globals.v[3] + globals.m[2][3]; +} + fn set_array(i: i32, v: f32) { globals.a[i] = v; } +fn set_dynamic_array(i: i32, v: f32) { + globals.d[i] = v; +} + fn set_vector(i: i32, v: f32) { globals.v[i] = v; } @@ -44,3 +61,13 @@ fn set_matrix(i: i32, v: vec4) { fn set_index_twice(i: i32, j: i32, v: f32) { globals.m[i][j] = v; } + +fn set_expensive(i: i32, v: f32) { + globals.a[i32(sin(f32(i) / 100.0) * 100.0)] = v; +} + +fn set_in_bounds(v: f32) { + globals.a[9] = v; + globals.v[3] = v; + globals.m[2][3] = v; +} diff --git a/tests/out/msl/bounds-check-restrict.msl b/tests/out/msl/bounds-check-restrict.msl new file mode 100644 index 0000000000..16e4320085 --- /dev/null +++ b/tests/out/msl/bounds-check-restrict.msl @@ -0,0 +1,163 @@ +// language: metal1.1 +#include +#include + +struct _mslBufferSizes { + metal::uint size0; +}; + +struct type_1 { + float inner[10]; +}; +typedef float type_4[1]; +struct Globals { + type_1 a; + char _pad1[8]; + metal::float4 v; + metal::float3x4 m; + type_4 d; +}; + +float index_array( + int i, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + float _e4 = globals.a.inner[metal::min(unsigned(i), 9u)]; + return _e4; +} + +float index_dynamic_array( + int i_1, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + float _e4 = globals.d[metal::min(unsigned(i_1), (_buffer_sizes.size0 - 112 - 4) / 4)]; + return _e4; +} + +float index_vector( + int i_2, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + float _e4 = globals.v[metal::min(unsigned(i_2), 3u)]; + return _e4; +} + +float index_vector_by_value( + metal::float4 v, + int i_3 +) { + return v[metal::min(unsigned(i_3), 3u)]; +} + +metal::float4 index_matrix( + int i_4, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + metal::float4 _e4 = globals.m[metal::min(unsigned(i_4), 2u)]; + return _e4; +} + +float index_twice( + int i_5, + int j, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + float _e6 = globals.m[metal::min(unsigned(i_5), 2u)][metal::min(unsigned(j), 3u)]; + return _e6; +} + +float index_expensive( + int i_6, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + float _e11 = globals.a.inner[metal::min(unsigned(static_cast(metal::sin(static_cast(i_6) / 100.0) * 100.0)), 9u)]; + return _e11; +} + +float index_in_bounds( + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + float _e4 = globals.a.inner[9]; + float _e8 = globals.v.w; + float _e15 = globals.m[2].w; + return (_e4 + _e8) + _e15; +} + +void set_array( + int i_7, + float v_1, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + globals.a.inner[metal::min(unsigned(i_7), 9u)] = v_1; + return; +} + +void set_dynamic_array( + int i_8, + float v_2, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + globals.d[metal::min(unsigned(i_8), (_buffer_sizes.size0 - 112 - 4) / 4)] = v_2; + return; +} + +void set_vector( + int i_9, + float v_3, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + globals.v[metal::min(unsigned(i_9), 3u)] = v_3; + return; +} + +void set_matrix( + int i_10, + metal::float4 v_4, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + globals.m[metal::min(unsigned(i_10), 2u)] = v_4; + return; +} + +void set_index_twice( + int i_11, + int j_1, + float v_5, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + globals.m[metal::min(unsigned(i_11), 2u)][metal::min(unsigned(j_1), 3u)] = v_5; + return; +} + +void set_expensive( + int i_12, + float v_6, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + globals.a.inner[metal::min(unsigned(static_cast(metal::sin(static_cast(i_12) / 100.0) * 100.0)), 9u)] = v_6; + return; +} + +void set_in_bounds( + float v_7, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + globals.a.inner[9] = v_7; + globals.v.w = v_7; + globals.m[2].w = v_7; + return; +} diff --git a/tests/out/msl/bounds-check-zero.msl b/tests/out/msl/bounds-check-zero.msl new file mode 100644 index 0000000000..6fa6613204 --- /dev/null +++ b/tests/out/msl/bounds-check-zero.msl @@ -0,0 +1,177 @@ +// language: metal1.1 +#include +#include + +struct _mslBufferSizes { + metal::uint size0; +}; + +struct type_1 { + float inner[10]; +}; +typedef float type_4[1]; +struct Globals { + type_1 a; + char _pad1[8]; + metal::float4 v; + metal::float3x4 m; + type_4 d; +}; + +float index_array( + int i, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + float _e4 = metal::uint(i) < 10 ? globals.a.inner[i] : 0; + return _e4; +} + +float index_dynamic_array( + int i_1, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + float _e4 = metal::uint(i_1) < 1 + (_buffer_sizes.size0 - 112 - 4) / 4 ? globals.d[i_1] : 0; + return _e4; +} + +float index_vector( + int i_2, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + float _e4 = metal::uint(i_2) < 4 ? globals.v[i_2] : 0; + return _e4; +} + +float index_vector_by_value( + metal::float4 v, + int i_3 +) { + return metal::uint(i_3) < 4 ? v[i_3] : 0; +} + +metal::float4 index_matrix( + int i_4, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + metal::float4 _e4 = metal::uint(i_4) < 3 ? globals.m[i_4] : 0; + return _e4; +} + +float index_twice( + int i_5, + int j, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + float _e6 = metal::uint(j) < 4 && metal::uint(i_5) < 3 ? globals.m[i_5][j] : 0; + return _e6; +} + +float index_expensive( + int i_6, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + int _e9 = static_cast(metal::sin(static_cast(i_6) / 100.0) * 100.0); + float _e11 = metal::uint(_e9) < 10 ? globals.a.inner[_e9] : 0; + return _e11; +} + +float index_in_bounds( + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + float _e4 = globals.a.inner[9]; + float _e8 = globals.v.w; + float _e15 = globals.m[2].w; + return (_e4 + _e8) + _e15; +} + +void set_array( + int i_7, + float v_1, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + if (metal::uint(i_7) < 10) { + globals.a.inner[i_7] = v_1; + } + return; +} + +void set_dynamic_array( + int i_8, + float v_2, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + if (metal::uint(i_8) < 1 + (_buffer_sizes.size0 - 112 - 4) / 4) { + globals.d[i_8] = v_2; + } + return; +} + +void set_vector( + int i_9, + float v_3, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + if (metal::uint(i_9) < 4) { + globals.v[i_9] = v_3; + } + return; +} + +void set_matrix( + int i_10, + metal::float4 v_4, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + if (metal::uint(i_10) < 3) { + globals.m[i_10] = v_4; + } + return; +} + +void set_index_twice( + int i_11, + int j_1, + float v_5, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + if (metal::uint(j_1) < 4 && metal::uint(i_11) < 3) { + globals.m[i_11][j_1] = v_5; + } + return; +} + +void set_expensive( + int i_12, + float v_6, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + int _e10 = static_cast(metal::sin(static_cast(i_12) / 100.0) * 100.0); + if (metal::uint(_e10) < 10) { + globals.a.inner[_e10] = v_6; + } + return; +} + +void set_in_bounds( + float v_7, + device Globals& globals, + constant _mslBufferSizes& _buffer_sizes +) { + globals.a.inner[9] = v_7; + globals.v.w = v_7; + globals.m[2].w = v_7; + return; +} diff --git a/tests/out/msl/policy-mix.msl b/tests/out/msl/policy-mix.msl new file mode 100644 index 0000000000..c3b5bb1dda --- /dev/null +++ b/tests/out/msl/policy-mix.msl @@ -0,0 +1,46 @@ +// language: metal1.1 +#include +#include + +struct type_1 { + metal::float4 inner[10]; +}; +struct InStorage { + type_1 a; +}; +struct type_2 { + metal::float4 inner[20]; +}; +struct InUniform { + type_2 a; +}; +struct type_5 { + float inner[30]; +}; +struct type_6 { + float inner[40]; +}; +struct type_9 { + metal::float4 inner[2]; +}; + +metal::float4 mock_function( + metal::int2 c, + int i, + int l, + constant InStorage& in_storage, + constant InUniform& in_uniform, + metal::texture2d_array image_2d_array, + threadgroup type_5 const& in_workgroup, + thread type_6 const& in_private +) { + type_9 in_function; + for(int _i=0; _i<2; ++_i) in_function.inner[_i] = type_9 {metal::float4(0.7070000171661377, 0.0, 0.0, 1.0), metal::float4(0.0, 0.7070000171661377, 0.0, 1.0)}.inner[_i]; + metal::float4 _e22 = in_storage.a.inner[i]; + metal::float4 _e25 = in_uniform.a.inner[i]; + metal::float4 _e27 = image_2d_array.read(metal::uint2(c), i, l); + float _e30 = in_workgroup.inner[metal::min(unsigned(i), 29u)]; + float _e34 = in_private.inner[metal::min(unsigned(i), 39u)]; + metal::float4 _e38 = in_function.inner[metal::min(unsigned(i), 1u)]; + return ((((_e22 + _e25) + _e27) + metal::float4(_e30)) + metal::float4(_e34)) + _e38; +} diff --git a/tests/out/spv/bounds-check-restrict.spvasm b/tests/out/spv/bounds-check-restrict.spvasm new file mode 100644 index 0000000000..1f83ea976c --- /dev/null +++ b/tests/out/spv/bounds-check-restrict.spvasm @@ -0,0 +1,238 @@ +; SPIR-V +; Version: 1.1 +; Generator: rspirv +; Bound: 166 +OpCapability Shader +OpCapability Linkage +OpExtension "SPV_KHR_storage_buffer_storage_class" +%1 = OpExtInstImport "GLSL.std.450" +OpMemoryModel Logical GLSL450 +OpDecorate %10 ArrayStride 4 +OpDecorate %13 ArrayStride 4 +OpDecorate %14 Block +OpMemberDecorate %14 0 Offset 0 +OpMemberDecorate %14 1 Offset 48 +OpMemberDecorate %14 2 Offset 64 +OpMemberDecorate %14 2 ColMajor +OpMemberDecorate %14 2 MatrixStride 16 +OpMemberDecorate %14 3 Offset 112 +OpDecorate %15 DescriptorSet 0 +OpDecorate %15 Binding 0 +%2 = OpTypeVoid +%4 = OpTypeInt 32 1 +%3 = OpConstant %4 10 +%6 = OpTypeFloat 32 +%5 = OpConstant %6 100.0 +%7 = OpConstant %4 9 +%8 = OpConstant %4 3 +%9 = OpConstant %4 2 +%10 = OpTypeArray %6 %3 +%11 = OpTypeVector %6 4 +%12 = OpTypeMatrix %11 3 +%13 = OpTypeRuntimeArray %6 +%14 = OpTypeStruct %10 %11 %12 %13 +%16 = OpTypePointer StorageBuffer %14 +%15 = OpVariable %16 StorageBuffer +%20 = OpTypeFunction %6 %4 +%22 = OpTypePointer StorageBuffer %10 +%23 = OpTypePointer StorageBuffer %6 +%25 = OpTypeInt 32 0 +%24 = OpConstant %25 9 +%27 = OpConstant %25 0 +%34 = OpTypePointer StorageBuffer %13 +%36 = OpConstant %25 1 +%39 = OpConstant %25 3 +%46 = OpTypePointer StorageBuffer %11 +%47 = OpTypePointer StorageBuffer %6 +%55 = OpTypeFunction %6 %11 %4 +%62 = OpTypeFunction %11 %4 +%64 = OpTypePointer StorageBuffer %12 +%65 = OpTypePointer StorageBuffer %11 +%66 = OpConstant %25 2 +%74 = OpTypeFunction %6 %4 %4 +%94 = OpTypeFunction %6 +%108 = OpTypeFunction %2 %4 %6 +%132 = OpTypeFunction %2 %4 %11 +%141 = OpTypeFunction %2 %4 %4 %6 +%161 = OpTypeFunction %2 %6 +%19 = OpFunction %6 None %20 +%18 = OpFunctionParameter %4 +%17 = OpLabel +OpBranch %21 +%21 = OpLabel +%26 = OpExtInst %25 %1 UMin %18 %24 +%28 = OpAccessChain %23 %15 %27 %26 +%29 = OpLoad %6 %28 +OpReturnValue %29 +OpFunctionEnd +%32 = OpFunction %6 None %20 +%31 = OpFunctionParameter %4 +%30 = OpLabel +OpBranch %33 +%33 = OpLabel +%35 = OpArrayLength %25 %15 3 +%37 = OpISub %25 %35 %36 +%38 = OpExtInst %25 %1 UMin %31 %37 +%40 = OpAccessChain %23 %15 %39 %38 +%41 = OpLoad %6 %40 +OpReturnValue %41 +OpFunctionEnd +%44 = OpFunction %6 None %20 +%43 = OpFunctionParameter %4 +%42 = OpLabel +OpBranch %45 +%45 = OpLabel +%48 = OpExtInst %25 %1 UMin %43 %39 +%49 = OpAccessChain %47 %15 %36 %48 +%50 = OpLoad %6 %49 +OpReturnValue %50 +OpFunctionEnd +%54 = OpFunction %6 None %55 +%52 = OpFunctionParameter %11 +%53 = OpFunctionParameter %4 +%51 = OpLabel +OpBranch %56 +%56 = OpLabel +%57 = OpExtInst %25 %1 UMin %53 %39 +%58 = OpVectorExtractDynamic %6 %52 %57 +OpReturnValue %58 +OpFunctionEnd +%61 = OpFunction %11 None %62 +%60 = OpFunctionParameter %4 +%59 = OpLabel +OpBranch %63 +%63 = OpLabel +%67 = OpExtInst %25 %1 UMin %60 %66 +%68 = OpAccessChain %65 %15 %66 %67 +%69 = OpLoad %11 %68 +OpReturnValue %69 +OpFunctionEnd +%73 = OpFunction %6 None %74 +%71 = OpFunctionParameter %4 +%72 = OpFunctionParameter %4 +%70 = OpLabel +OpBranch %75 +%75 = OpLabel +%76 = OpExtInst %25 %1 UMin %72 %39 +%77 = OpExtInst %25 %1 UMin %71 %66 +%78 = OpAccessChain %47 %15 %66 %77 %76 +%79 = OpLoad %6 %78 +OpReturnValue %79 +OpFunctionEnd +%82 = OpFunction %6 None %20 +%81 = OpFunctionParameter %4 +%80 = OpLabel +OpBranch %83 +%83 = OpLabel +%84 = OpConvertSToF %6 %81 +%85 = OpFDiv %6 %84 %5 +%86 = OpExtInst %6 %1 Sin %85 +%87 = OpFMul %6 %86 %5 +%88 = OpConvertFToS %4 %87 +%89 = OpExtInst %25 %1 UMin %88 %24 +%90 = OpAccessChain %23 %15 %27 %89 +%91 = OpLoad %6 %90 +OpReturnValue %91 +OpFunctionEnd +%93 = OpFunction %6 None %94 +%92 = OpLabel +OpBranch %95 +%95 = OpLabel +%96 = OpAccessChain %23 %15 %27 %24 +%97 = OpLoad %6 %96 +%98 = OpAccessChain %47 %15 %36 %39 +%99 = OpLoad %6 %98 +%100 = OpFAdd %6 %97 %99 +%101 = OpAccessChain %47 %15 %66 %66 %39 +%102 = OpLoad %6 %101 +%103 = OpFAdd %6 %100 %102 +OpReturnValue %103 +OpFunctionEnd +%107 = OpFunction %2 None %108 +%105 = OpFunctionParameter %4 +%106 = OpFunctionParameter %6 +%104 = OpLabel +OpBranch %109 +%109 = OpLabel +%110 = OpExtInst %25 %1 UMin %105 %24 +%111 = OpAccessChain %23 %15 %27 %110 +OpStore %111 %106 +OpReturn +OpFunctionEnd +%115 = OpFunction %2 None %108 +%113 = OpFunctionParameter %4 +%114 = OpFunctionParameter %6 +%112 = OpLabel +OpBranch %116 +%116 = OpLabel +%117 = OpArrayLength %25 %15 3 +%118 = OpISub %25 %117 %36 +%119 = OpExtInst %25 %1 UMin %113 %118 +%120 = OpAccessChain %23 %15 %39 %119 +OpStore %120 %114 +OpReturn +OpFunctionEnd +%124 = OpFunction %2 None %108 +%122 = OpFunctionParameter %4 +%123 = OpFunctionParameter %6 +%121 = OpLabel +OpBranch %125 +%125 = OpLabel +%126 = OpExtInst %25 %1 UMin %122 %39 +%127 = OpAccessChain %47 %15 %36 %126 +OpStore %127 %123 +OpReturn +OpFunctionEnd +%131 = OpFunction %2 None %132 +%129 = OpFunctionParameter %4 +%130 = OpFunctionParameter %11 +%128 = OpLabel +OpBranch %133 +%133 = OpLabel +%134 = OpExtInst %25 %1 UMin %129 %66 +%135 = OpAccessChain %65 %15 %66 %134 +OpStore %135 %130 +OpReturn +OpFunctionEnd +%140 = OpFunction %2 None %141 +%137 = OpFunctionParameter %4 +%138 = OpFunctionParameter %4 +%139 = OpFunctionParameter %6 +%136 = OpLabel +OpBranch %142 +%142 = OpLabel +%143 = OpExtInst %25 %1 UMin %138 %39 +%144 = OpExtInst %25 %1 UMin %137 %66 +%145 = OpAccessChain %47 %15 %66 %144 %143 +OpStore %145 %139 +OpReturn +OpFunctionEnd +%149 = OpFunction %2 None %108 +%147 = OpFunctionParameter %4 +%148 = OpFunctionParameter %6 +%146 = OpLabel +OpBranch %150 +%150 = OpLabel +%151 = OpConvertSToF %6 %147 +%152 = OpFDiv %6 %151 %5 +%153 = OpExtInst %6 %1 Sin %152 +%154 = OpFMul %6 %153 %5 +%155 = OpConvertFToS %4 %154 +%156 = OpExtInst %25 %1 UMin %155 %24 +%157 = OpAccessChain %23 %15 %27 %156 +OpStore %157 %148 +OpReturn +OpFunctionEnd +%160 = OpFunction %2 None %161 +%159 = OpFunctionParameter %6 +%158 = OpLabel +OpBranch %162 +%162 = OpLabel +%163 = OpAccessChain %23 %15 %27 %24 +OpStore %163 %159 +%164 = OpAccessChain %47 %15 %36 %39 +OpStore %164 %159 +%165 = OpAccessChain %47 %15 %66 %66 %39 +OpStore %165 %159 +OpReturn +OpFunctionEnd \ No newline at end of file diff --git a/tests/out/spv/bounds-check-zero.spvasm b/tests/out/spv/bounds-check-zero.spvasm index eb3a2da42a..82138feeb9 100644 --- a/tests/out/spv/bounds-check-zero.spvasm +++ b/tests/out/spv/bounds-check-zero.spvasm @@ -1,204 +1,320 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 130 +; Bound: 209 OpCapability Shader OpCapability Linkage OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpDecorate %6 ArrayStride 4 -OpDecorate %9 Block -OpMemberDecorate %9 0 Offset 0 -OpMemberDecorate %9 1 Offset 48 -OpMemberDecorate %9 2 Offset 64 -OpMemberDecorate %9 2 ColMajor -OpMemberDecorate %9 2 MatrixStride 16 -OpDecorate %10 DescriptorSet 0 -OpDecorate %10 Binding 0 +OpDecorate %10 ArrayStride 4 +OpDecorate %13 ArrayStride 4 +OpDecorate %14 Block +OpMemberDecorate %14 0 Offset 0 +OpMemberDecorate %14 1 Offset 48 +OpMemberDecorate %14 2 Offset 64 +OpMemberDecorate %14 2 ColMajor +OpMemberDecorate %14 2 MatrixStride 16 +OpMemberDecorate %14 3 Offset 112 +OpDecorate %15 DescriptorSet 0 +OpDecorate %15 Binding 0 %2 = OpTypeVoid %4 = OpTypeInt 32 1 %3 = OpConstant %4 10 -%5 = OpTypeFloat 32 -%6 = OpTypeArray %5 %3 -%7 = OpTypeVector %5 4 -%8 = OpTypeMatrix %7 3 -%9 = OpTypeStruct %6 %7 %8 -%11 = OpTypePointer StorageBuffer %9 -%10 = OpVariable %11 StorageBuffer -%15 = OpTypeFunction %5 %4 -%17 = OpTypePointer StorageBuffer %6 -%18 = OpTypePointer StorageBuffer %5 -%20 = OpTypeInt 32 0 -%19 = OpConstant %20 10 -%22 = OpTypeBool -%23 = OpConstant %20 0 -%25 = OpConstantNull %5 -%34 = OpTypePointer StorageBuffer %7 -%35 = OpTypePointer StorageBuffer %5 -%36 = OpConstant %20 4 -%38 = OpConstant %20 1 -%40 = OpConstantNull %5 -%49 = OpTypeFunction %5 %7 %4 -%52 = OpConstantNull %5 -%60 = OpTypeFunction %7 %4 -%62 = OpTypePointer StorageBuffer %8 -%63 = OpTypePointer StorageBuffer %7 -%64 = OpConstant %20 3 -%66 = OpConstant %20 2 -%68 = OpConstantNull %7 -%77 = OpTypeFunction %5 %4 %4 -%83 = OpConstantNull %5 -%92 = OpTypeFunction %2 %4 %5 -%111 = OpTypeFunction %2 %4 %7 -%122 = OpTypeFunction %2 %4 %4 %5 -%14 = OpFunction %5 None %15 -%13 = OpFunctionParameter %4 -%12 = OpLabel -OpBranch %16 -%16 = OpLabel -%21 = OpULessThan %22 %13 %19 -OpSelectionMerge %26 None -OpBranchConditional %21 %27 %26 -%27 = OpLabel -%24 = OpAccessChain %18 %10 %23 %13 -%28 = OpLoad %5 %24 -OpBranch %26 -%26 = OpLabel -%29 = OpPhi %5 %25 %16 %28 %27 -OpReturnValue %29 +%6 = OpTypeFloat 32 +%5 = OpConstant %6 100.0 +%7 = OpConstant %4 9 +%8 = OpConstant %4 3 +%9 = OpConstant %4 2 +%10 = OpTypeArray %6 %3 +%11 = OpTypeVector %6 4 +%12 = OpTypeMatrix %11 3 +%13 = OpTypeRuntimeArray %6 +%14 = OpTypeStruct %10 %11 %12 %13 +%16 = OpTypePointer StorageBuffer %14 +%15 = OpVariable %16 StorageBuffer +%20 = OpTypeFunction %6 %4 +%22 = OpTypePointer StorageBuffer %10 +%23 = OpTypePointer StorageBuffer %6 +%25 = OpTypeInt 32 0 +%24 = OpConstant %25 10 +%27 = OpTypeBool +%28 = OpConstant %25 0 +%30 = OpConstantNull %6 +%39 = OpTypePointer StorageBuffer %13 +%42 = OpConstant %25 3 +%44 = OpConstantNull %6 +%53 = OpTypePointer StorageBuffer %11 +%54 = OpTypePointer StorageBuffer %6 +%55 = OpConstant %25 4 +%57 = OpConstant %25 1 +%59 = OpConstantNull %6 +%68 = OpTypeFunction %6 %11 %4 +%71 = OpConstantNull %6 +%79 = OpTypeFunction %11 %4 +%81 = OpTypePointer StorageBuffer %12 +%82 = OpTypePointer StorageBuffer %11 +%84 = OpConstant %25 2 +%86 = OpConstantNull %11 +%95 = OpTypeFunction %6 %4 %4 +%101 = OpConstantNull %6 +%117 = OpConstantNull %6 +%124 = OpTypeFunction %6 +%126 = OpConstant %25 9 +%139 = OpTypeFunction %2 %4 %6 +%168 = OpTypeFunction %2 %4 %11 +%179 = OpTypeFunction %2 %4 %4 %6 +%204 = OpTypeFunction %2 %6 +%19 = OpFunction %6 None %20 +%18 = OpFunctionParameter %4 +%17 = OpLabel +OpBranch %21 +%21 = OpLabel +%26 = OpULessThan %27 %18 %24 +OpSelectionMerge %31 None +OpBranchConditional %26 %32 %31 +%32 = OpLabel +%29 = OpAccessChain %23 %15 %28 %18 +%33 = OpLoad %6 %29 +OpBranch %31 +%31 = OpLabel +%34 = OpPhi %6 %30 %21 %33 %32 +OpReturnValue %34 OpFunctionEnd -%32 = OpFunction %5 None %15 -%31 = OpFunctionParameter %4 -%30 = OpLabel -OpBranch %33 -%33 = OpLabel -%37 = OpULessThan %22 %31 %36 -OpSelectionMerge %41 None -OpBranchConditional %37 %42 %41 -%42 = OpLabel -%39 = OpAccessChain %35 %10 %38 %31 -%43 = OpLoad %5 %39 -OpBranch %41 -%41 = OpLabel -%44 = OpPhi %5 %40 %33 %43 %42 -OpReturnValue %44 -OpFunctionEnd -%48 = OpFunction %5 None %49 -%46 = OpFunctionParameter %7 -%47 = OpFunctionParameter %4 +%37 = OpFunction %6 None %20 +%36 = OpFunctionParameter %4 +%35 = OpLabel +OpBranch %38 +%38 = OpLabel +%40 = OpArrayLength %25 %15 3 +%41 = OpULessThan %27 %36 %40 +OpSelectionMerge %45 None +OpBranchConditional %41 %46 %45 +%46 = OpLabel +%43 = OpAccessChain %23 %15 %42 %36 +%47 = OpLoad %6 %43 +OpBranch %45 %45 = OpLabel -OpBranch %50 -%50 = OpLabel -%51 = OpULessThan %22 %47 %36 -OpSelectionMerge %53 None -OpBranchConditional %51 %54 %53 -%54 = OpLabel -%55 = OpVectorExtractDynamic %5 %46 %47 -OpBranch %53 -%53 = OpLabel -%56 = OpPhi %5 %52 %50 %55 %54 -OpReturnValue %56 +%48 = OpPhi %6 %44 %38 %47 %46 +OpReturnValue %48 OpFunctionEnd -%59 = OpFunction %7 None %60 -%58 = OpFunctionParameter %4 -%57 = OpLabel -OpBranch %61 +%51 = OpFunction %6 None %20 +%50 = OpFunctionParameter %4 +%49 = OpLabel +OpBranch %52 +%52 = OpLabel +%56 = OpULessThan %27 %50 %55 +OpSelectionMerge %60 None +OpBranchConditional %56 %61 %60 %61 = OpLabel -%65 = OpULessThan %22 %58 %64 -OpSelectionMerge %69 None -OpBranchConditional %65 %70 %69 -%70 = OpLabel -%67 = OpAccessChain %63 %10 %66 %58 -%71 = OpLoad %7 %67 +%58 = OpAccessChain %54 %15 %57 %50 +%62 = OpLoad %6 %58 +OpBranch %60 +%60 = OpLabel +%63 = OpPhi %6 %59 %52 %62 %61 +OpReturnValue %63 +OpFunctionEnd +%67 = OpFunction %6 None %68 +%65 = OpFunctionParameter %11 +%66 = OpFunctionParameter %4 +%64 = OpLabel OpBranch %69 %69 = OpLabel -%72 = OpPhi %7 %68 %61 %71 %70 -OpReturnValue %72 -OpFunctionEnd -%76 = OpFunction %5 None %77 -%74 = OpFunctionParameter %4 -%75 = OpFunctionParameter %4 +%70 = OpULessThan %27 %66 %55 +OpSelectionMerge %72 None +OpBranchConditional %70 %73 %72 %73 = OpLabel -OpBranch %78 -%78 = OpLabel -%79 = OpULessThan %22 %75 %36 -%80 = OpULessThan %22 %74 %64 -%81 = OpLogicalAnd %22 %79 %80 -OpSelectionMerge %84 None -OpBranchConditional %81 %85 %84 -%85 = OpLabel -%82 = OpAccessChain %35 %10 %66 %74 %75 -%86 = OpLoad %5 %82 -OpBranch %84 -%84 = OpLabel -%87 = OpPhi %5 %83 %78 %86 %85 -OpReturnValue %87 +%74 = OpVectorExtractDynamic %6 %65 %66 +OpBranch %72 +%72 = OpLabel +%75 = OpPhi %6 %71 %69 %74 %73 +OpReturnValue %75 OpFunctionEnd -%91 = OpFunction %2 None %92 -%89 = OpFunctionParameter %4 -%90 = OpFunctionParameter %5 +%78 = OpFunction %11 None %79 +%77 = OpFunctionParameter %4 +%76 = OpLabel +OpBranch %80 +%80 = OpLabel +%83 = OpULessThan %27 %77 %42 +OpSelectionMerge %87 None +OpBranchConditional %83 %88 %87 %88 = OpLabel -OpBranch %93 -%93 = OpLabel -%94 = OpULessThan %22 %89 %19 -OpSelectionMerge %96 None -OpBranchConditional %94 %97 %96 -%97 = OpLabel -%95 = OpAccessChain %18 %10 %23 %89 -OpStore %95 %90 +%85 = OpAccessChain %82 %15 %84 %77 +%89 = OpLoad %11 %85 +OpBranch %87 +%87 = OpLabel +%90 = OpPhi %11 %86 %80 %89 %88 +OpReturnValue %90 +OpFunctionEnd +%94 = OpFunction %6 None %95 +%92 = OpFunctionParameter %4 +%93 = OpFunctionParameter %4 +%91 = OpLabel OpBranch %96 %96 = OpLabel -OpReturn -OpFunctionEnd -%101 = OpFunction %2 None %92 -%99 = OpFunctionParameter %4 -%100 = OpFunctionParameter %5 -%98 = OpLabel +%97 = OpULessThan %27 %93 %55 +%98 = OpULessThan %27 %92 %42 +%99 = OpLogicalAnd %27 %97 %98 +OpSelectionMerge %102 None +OpBranchConditional %99 %103 %102 +%103 = OpLabel +%100 = OpAccessChain %54 %15 %84 %92 %93 +%104 = OpLoad %6 %100 OpBranch %102 %102 = OpLabel -%103 = OpULessThan %22 %99 %36 -OpSelectionMerge %105 None -OpBranchConditional %103 %106 %105 +%105 = OpPhi %6 %101 %96 %104 %103 +OpReturnValue %105 +OpFunctionEnd +%108 = OpFunction %6 None %20 +%107 = OpFunctionParameter %4 %106 = OpLabel -%104 = OpAccessChain %35 %10 %38 %99 -OpStore %104 %100 -OpBranch %105 -%105 = OpLabel +OpBranch %109 +%109 = OpLabel +%110 = OpConvertSToF %6 %107 +%111 = OpFDiv %6 %110 %5 +%112 = OpExtInst %6 %1 Sin %111 +%113 = OpFMul %6 %112 %5 +%114 = OpConvertFToS %4 %113 +%115 = OpULessThan %27 %114 %24 +OpSelectionMerge %118 None +OpBranchConditional %115 %119 %118 +%119 = OpLabel +%116 = OpAccessChain %23 %15 %28 %114 +%120 = OpLoad %6 %116 +OpBranch %118 +%118 = OpLabel +%121 = OpPhi %6 %117 %109 %120 %119 +OpReturnValue %121 +OpFunctionEnd +%123 = OpFunction %6 None %124 +%122 = OpLabel +OpBranch %125 +%125 = OpLabel +%127 = OpAccessChain %23 %15 %28 %126 +%128 = OpLoad %6 %127 +%129 = OpAccessChain %54 %15 %57 %42 +%130 = OpLoad %6 %129 +%131 = OpFAdd %6 %128 %130 +%132 = OpAccessChain %54 %15 %84 %84 %42 +%133 = OpLoad %6 %132 +%134 = OpFAdd %6 %131 %133 +OpReturnValue %134 +OpFunctionEnd +%138 = OpFunction %2 None %139 +%136 = OpFunctionParameter %4 +%137 = OpFunctionParameter %6 +%135 = OpLabel +OpBranch %140 +%140 = OpLabel +%141 = OpULessThan %27 %136 %24 +OpSelectionMerge %143 None +OpBranchConditional %141 %144 %143 +%144 = OpLabel +%142 = OpAccessChain %23 %15 %28 %136 +OpStore %142 %137 +OpBranch %143 +%143 = OpLabel +OpReturn +OpFunctionEnd +%148 = OpFunction %2 None %139 +%146 = OpFunctionParameter %4 +%147 = OpFunctionParameter %6 +%145 = OpLabel +OpBranch %149 +%149 = OpLabel +%150 = OpArrayLength %25 %15 3 +%151 = OpULessThan %27 %146 %150 +OpSelectionMerge %153 None +OpBranchConditional %151 %154 %153 +%154 = OpLabel +%152 = OpAccessChain %23 %15 %42 %146 +OpStore %152 %147 +OpBranch %153 +%153 = OpLabel +OpReturn +OpFunctionEnd +%158 = OpFunction %2 None %139 +%156 = OpFunctionParameter %4 +%157 = OpFunctionParameter %6 +%155 = OpLabel +OpBranch %159 +%159 = OpLabel +%160 = OpULessThan %27 %156 %55 +OpSelectionMerge %162 None +OpBranchConditional %160 %163 %162 +%163 = OpLabel +%161 = OpAccessChain %54 %15 %57 %156 +OpStore %161 %157 +OpBranch %162 +%162 = OpLabel +OpReturn +OpFunctionEnd +%167 = OpFunction %2 None %168 +%165 = OpFunctionParameter %4 +%166 = OpFunctionParameter %11 +%164 = OpLabel +OpBranch %169 +%169 = OpLabel +%170 = OpULessThan %27 %165 %42 +OpSelectionMerge %172 None +OpBranchConditional %170 %173 %172 +%173 = OpLabel +%171 = OpAccessChain %82 %15 %84 %165 +OpStore %171 %166 +OpBranch %172 +%172 = OpLabel +OpReturn +OpFunctionEnd +%178 = OpFunction %2 None %179 +%175 = OpFunctionParameter %4 +%176 = OpFunctionParameter %4 +%177 = OpFunctionParameter %6 +%174 = OpLabel +OpBranch %180 +%180 = OpLabel +%181 = OpULessThan %27 %176 %55 +%182 = OpULessThan %27 %175 %42 +%183 = OpLogicalAnd %27 %181 %182 +OpSelectionMerge %185 None +OpBranchConditional %183 %186 %185 +%186 = OpLabel +%184 = OpAccessChain %54 %15 %84 %175 %176 +OpStore %184 %177 +OpBranch %185 +%185 = OpLabel OpReturn OpFunctionEnd -%110 = OpFunction %2 None %111 -%108 = OpFunctionParameter %4 -%109 = OpFunctionParameter %7 -%107 = OpLabel -OpBranch %112 -%112 = OpLabel -%113 = OpULessThan %22 %108 %64 -OpSelectionMerge %115 None -OpBranchConditional %113 %116 %115 -%116 = OpLabel -%114 = OpAccessChain %63 %10 %66 %108 -OpStore %114 %109 -OpBranch %115 -%115 = OpLabel +%190 = OpFunction %2 None %139 +%188 = OpFunctionParameter %4 +%189 = OpFunctionParameter %6 +%187 = OpLabel +OpBranch %191 +%191 = OpLabel +%192 = OpConvertSToF %6 %188 +%193 = OpFDiv %6 %192 %5 +%194 = OpExtInst %6 %1 Sin %193 +%195 = OpFMul %6 %194 %5 +%196 = OpConvertFToS %4 %195 +%197 = OpULessThan %27 %196 %24 +OpSelectionMerge %199 None +OpBranchConditional %197 %200 %199 +%200 = OpLabel +%198 = OpAccessChain %23 %15 %28 %196 +OpStore %198 %189 +OpBranch %199 +%199 = OpLabel OpReturn OpFunctionEnd -%121 = OpFunction %2 None %122 -%118 = OpFunctionParameter %4 -%119 = OpFunctionParameter %4 -%120 = OpFunctionParameter %5 -%117 = OpLabel -OpBranch %123 -%123 = OpLabel -%124 = OpULessThan %22 %119 %36 -%125 = OpULessThan %22 %118 %64 -%126 = OpLogicalAnd %22 %124 %125 -OpSelectionMerge %128 None -OpBranchConditional %126 %129 %128 -%129 = OpLabel -%127 = OpAccessChain %35 %10 %66 %118 %119 -OpStore %127 %120 -OpBranch %128 -%128 = OpLabel +%203 = OpFunction %2 None %204 +%202 = OpFunctionParameter %6 +%201 = OpLabel +OpBranch %205 +%205 = OpLabel +%206 = OpAccessChain %23 %15 %28 %126 +OpStore %206 %202 +%207 = OpAccessChain %54 %15 %57 %42 +OpStore %207 %202 +%208 = OpAccessChain %54 %15 %84 %84 %42 +OpStore %208 %202 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/snapshots.rs b/tests/snapshots.rs index 74a2b5cf26..565f326190 100644 --- a/tests/snapshots.rs +++ b/tests/snapshots.rs @@ -82,6 +82,7 @@ fn check_targets(module: &naga::Module, name: &str, targets: Targets) { Ok(string) => ron::de::from_str(&string).expect("Couldn't parse param file"), Err(_) => Parameters::default(), }; + let capabilities = if params.god_mode { naga::valid::Capabilities::all() } else { @@ -128,7 +129,14 @@ fn check_targets(module: &naga::Module, name: &str, targets: Targets) { #[cfg(all(feature = "deserialize", feature = "msl-out"))] { if targets.contains(Targets::METAL) { - write_output_msl(module, &info, &dest, name, ¶ms.msl); + write_output_msl( + module, + &info, + &dest, + name, + ¶ms.msl, + params.bounds_check_policies, + ); } } #[cfg(all(feature = "deserialize", feature = "glsl-out"))] @@ -227,6 +235,7 @@ fn write_output_msl( destination: &PathBuf, file_name: &str, options: &naga::back::msl::Options, + bounds_check_policies: naga::proc::BoundsCheckPolicies, ) { use naga::back::msl; @@ -236,8 +245,10 @@ fn write_output_msl( allow_point_size: true, }; + let mut options = options.clone(); + options.bounds_check_policies = bounds_check_policies; let (string, tr_info) = - msl::write_string(module, info, options, &pipeline_options).expect("Metal write failed"); + msl::write_string(module, info, &options, &pipeline_options).expect("Metal write failed"); for (ep, result) in module.entry_points.iter().zip(tr_info.entry_point_names) { if let Err(error) = result { @@ -458,10 +469,11 @@ fn convert_wgsl() { "globals", Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL, ), - ("bounds-check-zero", Targets::SPIRV), + ("bounds-check-zero", Targets::SPIRV | Targets::METAL), + ("bounds-check-restrict", Targets::SPIRV | Targets::METAL), ("bounds-check-image-restrict", Targets::SPIRV), ("bounds-check-image-rzsw", Targets::SPIRV), - ("policy-mix", Targets::SPIRV), + ("policy-mix", Targets::SPIRV | Targets::METAL), ( "texture-arg", Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL,