From d0e92aa6b34baee20d663cf638fa6aa981b6413a Mon Sep 17 00:00:00 2001 From: teoxoy <28601907+teoxoy@users.noreply.github.com> Date: Thu, 19 Dec 2024 12:09:40 +0100 Subject: [PATCH] resolve the size of override-sized arrays in backends --- naga/src/back/glsl/mod.rs | 21 +++---- naga/src/back/hlsl/conv.rs | 19 +++---- naga/src/back/hlsl/mod.rs | 2 + naga/src/back/hlsl/writer.rs | 10 ++-- naga/src/back/msl/mod.rs | 2 + naga/src/back/msl/writer.rs | 24 +++----- naga/src/back/pipeline_constants.rs | 87 +++++++++++++++-------------- naga/src/back/spv/index.rs | 7 +-- naga/src/back/spv/mod.rs | 2 + naga/src/back/spv/writer.rs | 32 ++++++----- naga/src/compact/mod.rs | 27 +++++---- naga/src/proc/index.rs | 34 ++++++++--- naga/src/proc/mod.rs | 47 ++++++++++++++++ naga/src/valid/expression.rs | 8 ++- naga/src/valid/mod.rs | 4 +- 15 files changed, 196 insertions(+), 130 deletions(-) diff --git a/naga/src/back/glsl/mod.rs b/naga/src/back/glsl/mod.rs index e34127b3a1..0b8c780164 100644 --- a/naga/src/back/glsl/mod.rs +++ b/naga/src/back/glsl/mod.rs @@ -502,6 +502,8 @@ pub enum Error { /// [`crate::Sampling::First`] is unsupported. #[error("`{:?}` sampling is unsupported", crate::Sampling::First)] FirstSamplingNotSupported, + #[error(transparent)] + ResolveArraySizeError(#[from] proc::ResolveArraySizeError), } /// Binary operation with a different logic on the GLSL side. @@ -973,13 +975,12 @@ impl<'a, W: Write> Writer<'a, W> { write!(self.out, "[")?; // Write the array size - // Writes nothing if `ArraySize::Dynamic` - match size { - crate::ArraySize::Constant(size) => { + // Writes nothing if `ResolvedSize::Runtime` + match proc::resolve_array_size(size, self.module.to_ctx())? { + proc::ResolvedSize::Constant(size) => { write!(self.out, "{size}")?; } - crate::ArraySize::Pending(_) => unreachable!(), - crate::ArraySize::Dynamic => (), + proc::ResolvedSize::Runtime => (), } write!(self.out, "]")?; @@ -4455,13 +4456,9 @@ impl<'a, W: Write> Writer<'a, W> { write!(self.out, ")")?; } TypeInner::Array { base, size, .. } => { - let count = match size - .to_indexable_length(self.module) - .expect("Bad array size") - { - proc::IndexableLength::Known(count) => count, - proc::IndexableLength::Pending => unreachable!(), - proc::IndexableLength::Dynamic => return Ok(()), + let count = match proc::resolve_array_size(size, self.module.to_ctx())? { + proc::ResolvedSize::Constant(size) => size, + proc::ResolvedSize::Runtime => return Ok(()), }; self.write_type(base)?; self.write_array_size(base, size)?; diff --git a/naga/src/back/hlsl/conv.rs b/naga/src/back/hlsl/conv.rs index 83c7667eab..e27dfe66e7 100644 --- a/naga/src/back/hlsl/conv.rs +++ b/naga/src/back/hlsl/conv.rs @@ -53,7 +53,7 @@ impl crate::TypeInner { } } - pub(super) fn size_hlsl(&self, gctx: crate::proc::GlobalCtx) -> u32 { + pub(super) fn size_hlsl(&self, gctx: crate::proc::GlobalCtx) -> Result { match *self { Self::Matrix { columns, @@ -62,19 +62,18 @@ impl crate::TypeInner { } => { let stride = Alignment::from(rows) * scalar.width as u32; let last_row_size = rows as u32 * scalar.width as u32; - ((columns as u32 - 1) * stride) + last_row_size + Ok(((columns as u32 - 1) * stride) + last_row_size) } Self::Array { base, size, stride } => { - let count = match size { - crate::ArraySize::Constant(size) => size.get(), - // A dynamically-sized array has to have at least one element - crate::ArraySize::Pending(_) => unreachable!(), - crate::ArraySize::Dynamic => 1, + let count = match crate::proc::resolve_array_size(size, gctx)? { + crate::proc::ResolvedSize::Constant(size) => size, + // A runtime-sized array has to have at least one element + crate::proc::ResolvedSize::Runtime => 1, }; - let last_el_size = gctx.types[base].inner.size_hlsl(gctx); - ((count - 1) * stride) + last_el_size + let last_el_size = gctx.types[base].inner.size_hlsl(gctx)?; + Ok(((count - 1) * stride) + last_el_size) } - _ => self.size(gctx), + _ => Ok(self.size(gctx)), } } diff --git a/naga/src/back/hlsl/mod.rs b/naga/src/back/hlsl/mod.rs index fe7d4f6d67..559965469b 100644 --- a/naga/src/back/hlsl/mod.rs +++ b/naga/src/back/hlsl/mod.rs @@ -267,6 +267,8 @@ pub enum Error { Custom(String), #[error("overrides should not be present at this stage")] Override, + #[error(transparent)] + ResolveArraySizeError(#[from] proc::ResolveArraySizeError), } #[derive(Default)] diff --git a/naga/src/back/hlsl/writer.rs b/naga/src/back/hlsl/writer.rs index bc6086d539..f4020b3387 100644 --- a/naga/src/back/hlsl/writer.rs +++ b/naga/src/back/hlsl/writer.rs @@ -984,12 +984,11 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { ) -> BackendResult { write!(self.out, "[")?; - match size { - crate::ArraySize::Constant(size) => { + match proc::resolve_array_size(size, module.to_ctx())? { + proc::ResolvedSize::Constant(size) => { write!(self.out, "{size}")?; } - crate::ArraySize::Pending(_) => unreachable!(), - crate::ArraySize::Dynamic => unreachable!(), + proc::ResolvedSize::Runtime => unreachable!(), } write!(self.out, "]")?; @@ -1034,7 +1033,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { } } let ty_inner = &module.types[member.ty].inner; - last_offset = member.offset + ty_inner.size_hlsl(module.to_ctx()); + last_offset = member.offset + ty_inner.size_hlsl(module.to_ctx())?; // The indentation is only for readability write!(self.out, "{}", back::INDENT)?; @@ -2635,7 +2634,6 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { index::IndexableLength::Known(limit) => { write!(self.out, "{}u", limit - 1)?; } - index::IndexableLength::Pending => unreachable!(), index::IndexableLength::Dynamic => unreachable!(), } write!(self.out, ")")?; diff --git a/naga/src/back/msl/mod.rs b/naga/src/back/msl/mod.rs index 0c85c8a9e4..8006b476d5 100644 --- a/naga/src/back/msl/mod.rs +++ b/naga/src/back/msl/mod.rs @@ -150,6 +150,8 @@ pub enum Error { UnsupportedRayTracing, #[error("overrides should not be present at this stage")] Override, + #[error(transparent)] + ResolveArraySizeError(#[from] crate::proc::ResolveArraySizeError), } #[derive(Clone, Debug, PartialEq, thiserror::Error)] diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs index f4b55bbbc5..154280dc86 100644 --- a/naga/src/back/msl/writer.rs +++ b/naga/src/back/msl/writer.rs @@ -2,8 +2,7 @@ use super::{sampler as sm, Error, LocationMode, Options, PipelineOptions, Transl use crate::{ arena::{Handle, HandleSet}, back::{self, Baked}, - proc::index, - proc::{self, NameKey, TypeResolution}, + proc::{self, index, NameKey, TypeResolution}, valid, FastHashMap, FastHashSet, }; #[cfg(test)] @@ -2409,7 +2408,6 @@ impl Writer { self.out.write_str(") < ")?; match length { index::IndexableLength::Known(value) => write!(self.out, "{value}")?, - index::IndexableLength::Pending => unreachable!(), index::IndexableLength::Dynamic => { let global = context.function.originating_global(base).ok_or_else(|| { @@ -2546,7 +2544,7 @@ impl Writer { ) -> BackendResult { let accessing_wrapped_array = match *base_ty { crate::TypeInner::Array { - size: crate::ArraySize::Constant(_), + size: crate::ArraySize::Constant(_) | crate::ArraySize::Pending(_), .. } => true, _ => false, @@ -2572,7 +2570,6 @@ impl Writer { index::IndexableLength::Known(limit) => { write!(self.out, "{}u", limit - 1)?; } - index::IndexableLength::Pending => unreachable!(), index::IndexableLength::Dynamic => { let global = context.function.originating_global(base).ok_or_else(|| { Error::GenericValidation("Could not find originating global".into()) @@ -3733,8 +3730,8 @@ impl Writer { first_time: false, }; - match size { - crate::ArraySize::Constant(size) => { + match proc::resolve_array_size(size, module.to_ctx())? { + proc::ResolvedSize::Constant(size) => { writeln!(self.out, "struct {name} {{")?; writeln!( self.out, @@ -3746,10 +3743,7 @@ impl Writer { )?; writeln!(self.out, "}};")?; } - crate::ArraySize::Pending(_) => { - unreachable!() - } - crate::ArraySize::Dynamic => { + proc::ResolvedSize::Runtime => { writeln!(self.out, "typedef {base_name} {name}[1];")?; } } @@ -6147,11 +6141,9 @@ mod workgroup_mem_init { writeln!(self.out, ", 0, {NAMESPACE}::memory_order_relaxed);")?; } crate::TypeInner::Array { base, size, .. } => { - let count = match size.to_indexable_length(module).expect("Bad array size") - { - proc::IndexableLength::Known(count) => count, - proc::IndexableLength::Pending => unreachable!(), - proc::IndexableLength::Dynamic => unreachable!(), + let count = match proc::resolve_array_size(size, module.to_ctx())? { + proc::ResolvedSize::Constant(size) => size, + proc::ResolvedSize::Runtime => unreachable!(), }; access_stack.enter_array(|access_stack, array_depth| { diff --git a/naga/src/back/pipeline_constants.rs b/naga/src/back/pipeline_constants.rs index 7d72354250..07513a1e78 100644 --- a/naga/src/back/pipeline_constants.rs +++ b/naga/src/back/pipeline_constants.rs @@ -4,7 +4,7 @@ use crate::{ proc::{ConstantEvaluator, ConstantEvaluatorError, Emitter}, valid::{Capabilities, ModuleInfo, ValidationError, ValidationFlags, Validator}, Arena, Block, Constant, Expression, Function, Handle, Literal, Module, Override, Range, Scalar, - Span, Statement, TypeInner, WithSpan, + Span, Statement, Type, TypeInner, UniqueArena, WithSpan, }; use std::{borrow::Cow, collections::HashSet, mem}; use thiserror::Error; @@ -51,6 +51,7 @@ pub fn process_overrides<'a>( return Ok((Cow::Borrowed(module), Cow::Borrowed(module_info))); } + let original_module_types = &module.types; let mut module = module.clone(); // A map from override handles to the handles of the constants @@ -196,7 +197,13 @@ pub fn process_overrides<'a>( } module.entry_points = entry_points; - process_pending(&mut module, &override_map, &adjusted_global_expressions)?; + process_pending( + &mut module.types, + original_module_types, + &module.constants, + &override_map, + &adjusted_global_expressions, + ); // Now that we've rewritten all the expressions, we need to // recompute their types and other metadata. For the time being, @@ -208,61 +215,57 @@ pub fn process_overrides<'a>( } fn process_pending( - module: &mut Module, + types: &mut UniqueArena, + original_module_types: &UniqueArena, + constants: &Arena, override_map: &HandleVec>, adjusted_global_expressions: &HandleVec>, -) -> Result<(), PipelineConstantError> { - for (handle, ty) in module.types.clone().iter() { +) { + for (handle, ty) in original_module_types.iter() { if let TypeInner::Array { base, size: crate::ArraySize::Pending(size), stride, } = ty.inner { - let expr = match size { + match size { crate::PendingArraySize::Expression(size_expr) => { - adjusted_global_expressions[size_expr] + let expr = adjusted_global_expressions[size_expr]; + if expr != size_expr { + types.replace( + handle, + Type { + name: ty.name.clone(), + inner: TypeInner::Array { + base, + size: crate::ArraySize::Pending( + crate::PendingArraySize::Expression(expr), + ), + stride, + }, + }, + ); + } } crate::PendingArraySize::Override(size_override) => { - module.constants[override_map[size_override]].init + let expr = constants[override_map[size_override]].init; + types.replace( + handle, + Type { + name: ty.name.clone(), + inner: TypeInner::Array { + base, + size: crate::ArraySize::Pending( + crate::PendingArraySize::Expression(expr), + ), + stride, + }, + }, + ); } }; - let value = module - .to_ctx() - .eval_expr_to_u32(expr) - .map(|n| { - if n == 0 { - Err(PipelineConstantError::ValidationError( - WithSpan::new(ValidationError::ArraySizeError { handle: expr }) - .with_span( - module.global_expressions.get_span(expr), - "evaluated to zero", - ), - )) - } else { - Ok(std::num::NonZeroU32::new(n).unwrap()) - } - }) - .map_err(|_| { - PipelineConstantError::ValidationError( - WithSpan::new(ValidationError::ArraySizeError { handle: expr }) - .with_span(module.global_expressions.get_span(expr), "negative"), - ) - })??; - module.types.replace( - handle, - crate::Type { - name: None, - inner: TypeInner::Array { - base, - size: crate::ArraySize::Constant(value), - stride, - }, - }, - ); } } - Ok(()) } fn process_workgroup_size_override( diff --git a/naga/src/back/spv/index.rs b/naga/src/back/spv/index.rs index 15e5df3f10..acc179b66b 100644 --- a/naga/src/back/spv/index.rs +++ b/naga/src/back/spv/index.rs @@ -267,13 +267,12 @@ impl BlockContext<'_> { block: &mut Block, ) -> Result, Error> { let sequence_ty = self.fun_info[sequence].ty.inner_with(&self.ir_module.types); - match sequence_ty.indexable_length(self.ir_module) { + match sequence_ty + .indexable_length(self.ir_module, crate::ArraySize::indexable_length_resolved) + { Ok(crate::proc::IndexableLength::Known(known_length)) => { Ok(MaybeKnown::Known(known_length)) } - Ok(crate::proc::IndexableLength::Pending) => { - unreachable!() - } Ok(crate::proc::IndexableLength::Dynamic) => { let length_id = self.write_runtime_array_length(sequence, block)?; Ok(MaybeKnown::Computed(length_id)) diff --git a/naga/src/back/spv/mod.rs b/naga/src/back/spv/mod.rs index 6385311c73..88e59941f4 100644 --- a/naga/src/back/spv/mod.rs +++ b/naga/src/back/spv/mod.rs @@ -73,6 +73,8 @@ pub enum Error { Validation(&'static str), #[error("overrides should not be present at this stage")] Override, + #[error(transparent)] + ResolveArraySizeError(#[from] crate::proc::ResolveArraySizeError), } #[derive(Default)] diff --git a/naga/src/back/spv/writer.rs b/naga/src/back/spv/writer.rs index 47f3ec513b..38d2ef2864 100644 --- a/naga/src/back/spv/writer.rs +++ b/naga/src/back/spv/writer.rs @@ -930,10 +930,10 @@ impl Writer { fn write_type_declaration_arena( &mut self, - arena: &UniqueArena, + module: &crate::Module, handle: Handle, ) -> Result { - let ty = &arena[handle]; + let ty = &module.types[handle]; // If it's a type that needs SPIR-V capabilities, request them now. // This needs to happen regardless of the LocalType lookup succeeding, // because some types which map to the same LocalType have different @@ -966,24 +966,26 @@ impl Writer { self.decorate(id, Decoration::ArrayStride, &[stride]); let type_id = self.get_type_id(LookupType::Handle(base)); - match size { - crate::ArraySize::Constant(length) => { - let length_id = self.get_index_constant(length.get()); + match crate::proc::resolve_array_size(size, module.to_ctx())? { + crate::proc::ResolvedSize::Constant(length) => { + let length_id = self.get_index_constant(length); Instruction::type_array(id, type_id, length_id) } - crate::ArraySize::Pending(_) => unreachable!(), - crate::ArraySize::Dynamic => Instruction::type_runtime_array(id, type_id), + crate::proc::ResolvedSize::Runtime => { + Instruction::type_runtime_array(id, type_id) + } } } crate::TypeInner::BindingArray { base, size } => { let type_id = self.get_type_id(LookupType::Handle(base)); - match size { - crate::ArraySize::Constant(length) => { - let length_id = self.get_index_constant(length.get()); + match crate::proc::resolve_array_size(size, module.to_ctx())? { + crate::proc::ResolvedSize::Constant(length) => { + let length_id = self.get_index_constant(length); Instruction::type_array(id, type_id, length_id) } - crate::ArraySize::Pending(_) => unreachable!(), - crate::ArraySize::Dynamic => Instruction::type_runtime_array(id, type_id), + crate::proc::ResolvedSize::Runtime => { + Instruction::type_runtime_array(id, type_id) + } } } crate::TypeInner::Struct { @@ -993,7 +995,7 @@ impl Writer { let mut has_runtime_array = false; let mut member_ids = Vec::with_capacity(members.len()); for (index, member) in members.iter().enumerate() { - let member_ty = &arena[member.ty]; + let member_ty = &module.types[member.ty]; match member_ty.inner { crate::TypeInner::Array { base: _, @@ -1004,7 +1006,7 @@ impl Writer { } _ => (), } - self.decorate_struct_member(id, index, member, arena)?; + self.decorate_struct_member(id, index, member, &module.types)?; let member_id = self.get_type_id(LookupType::Handle(member.ty)); member_ids.push(member_id); } @@ -1939,7 +1941,7 @@ impl Writer { // write all types for (handle, _) in ir_module.types.iter() { - self.write_type_declaration_arena(&ir_module.types, handle)?; + self.write_type_declaration_arena(ir_module, handle)?; } // write all const-expressions as constants diff --git a/naga/src/compact/mod.rs b/naga/src/compact/mod.rs index 9dff4a6cc2..993d82797d 100644 --- a/naga/src/compact/mod.rs +++ b/naga/src/compact/mod.rs @@ -223,20 +223,23 @@ pub fn compact(module: &mut crate::Module) { stride, } = ty.inner { + let old_size_expr = size_expr; module_map.global_expressions.adjust(&mut size_expr); - module.types.replace( - handle, - crate::Type { - name: None, - inner: crate::TypeInner::Array { - base, - size: crate::ArraySize::Pending(crate::PendingArraySize::Expression( - size_expr, - )), - stride, + if old_size_expr != size_expr { + module.types.replace( + handle, + crate::Type { + name: None, + inner: crate::TypeInner::Array { + base, + size: crate::ArraySize::Pending(crate::PendingArraySize::Expression( + size_expr, + )), + stride, + }, }, - }, - ); + ); + } } } diff --git a/naga/src/proc/index.rs b/naga/src/proc/index.rs index ac2a6589c1..eb9c59dc7c 100644 --- a/naga/src/proc/index.rs +++ b/naga/src/proc/index.rs @@ -315,7 +315,9 @@ pub fn access_needs_check( // 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(); + let length = base_inner + .indexable_length(module, crate::ArraySize::indexable_length_resolved) + .unwrap(); index.try_resolve_to_constant(expressions, module); if let (&GuardedIndex::Known(index), &IndexableLength::Known(length)) = (&index, &length) { if index < length { @@ -357,8 +359,8 @@ impl GuardedIndex { pub enum IndexableLengthError { #[error("Type is not indexable, and has no length (validation error)")] TypeNotIndexable, - #[error("Array length constant {0:?} is invalid")] - InvalidArrayLength(Handle), + #[error(transparent)] + ResolveArraySizeError(#[from] super::ResolveArraySizeError), } impl crate::TypeInner { @@ -376,13 +378,17 @@ impl crate::TypeInner { pub fn indexable_length( &self, module: &crate::Module, + array_size_indexable_length: fn( + crate::ArraySize, + &crate::Module, + ) -> Result, ) -> Result { use crate::TypeInner as Ti; let known_length = match *self { Ti::Vector { size, .. } => size as _, Ti::Matrix { columns, .. } => columns as _, Ti::Array { size, .. } | Ti::BindingArray { size, .. } => { - return size.to_indexable_length(module); + return array_size_indexable_length(size, module); } Ti::ValuePointer { size: Some(size), .. @@ -396,7 +402,7 @@ impl crate::TypeInner { Ti::Vector { size, .. } => size as _, Ti::Matrix { columns, .. } => columns as _, Ti::Array { size, .. } | Ti::BindingArray { size, .. } => { - return size.to_indexable_length(module) + return array_size_indexable_length(size, module); } _ => return Err(IndexableLengthError::TypeNotIndexable), } @@ -416,21 +422,31 @@ pub enum IndexableLength { /// Values of this type always have the given number of elements. Known(u32), - Pending, - /// The number of elements is determined at runtime. Dynamic, } impl crate::ArraySize { - pub const fn to_indexable_length( + /// This function should be used by validators that allow overrides + pub const fn indexable_length( self, _module: &crate::Module, ) -> Result { Ok(match self { Self::Constant(length) => IndexableLength::Known(length.get()), - Self::Pending(_) => IndexableLength::Pending, + Self::Pending(_) => IndexableLength::Dynamic, Self::Dynamic => IndexableLength::Dynamic, }) } + + /// This function should be used by backends and validators that reject overrides + pub fn indexable_length_resolved( + self, + module: &crate::Module, + ) -> Result { + match super::resolve_array_size(self, module.to_ctx())? { + super::ResolvedSize::Constant(length) => Ok(IndexableLength::Known(length)), + super::ResolvedSize::Runtime => Ok(IndexableLength::Dynamic), + } + } } diff --git a/naga/src/proc/mod.rs b/naga/src/proc/mod.rs index 76698fd102..1703918f78 100644 --- a/naga/src/proc/mod.rs +++ b/naga/src/proc/mod.rs @@ -18,6 +18,7 @@ pub use index::{BoundsCheckPolicies, BoundsCheckPolicy, IndexableLength, Indexab pub use layouter::{Alignment, LayoutError, LayoutErrorInner, Layouter, TypeLayout}; pub use namer::{EntryPointIndex, NameKey, Namer}; pub use terminator::ensure_block_returns; +use thiserror::Error; pub use typifier::{ResolveContext, ResolveError, TypeResolution}; impl From for super::Scalar { @@ -732,6 +733,52 @@ impl GlobalCtx<'_> { } } +pub enum ResolvedSize { + Constant(u32), + Runtime, +} + +#[derive(Error, Debug, Clone, Copy, PartialEq)] +pub enum ResolveArraySizeError { + #[error("array element count must be positive (> 0)")] + ExpectedPositiveArrayLength, +} + +/// # Panics +/// +/// - If `module` contains arrays with a pending size of [`crate::PendingArraySize::Override`] +/// - If `module` contains override expressions +pub fn resolve_array_size( + size: crate::ArraySize, + gctx: GlobalCtx, +) -> Result { + match size { + crate::ArraySize::Constant(length) => Ok(ResolvedSize::Constant(length.get())), + crate::ArraySize::Pending(pending_size) => { + let expr = match pending_size { + crate::PendingArraySize::Expression(handle) => handle, + crate::PendingArraySize::Override(_) => { + unreachable!("PendingArraySize::Override variants must have been replaced with PendingArraySize::Expression variants"); + } + }; + + let length = gctx.eval_expr_to_u32(expr).map_err(|err| match err { + U32EvalError::NonConst => { + unreachable!("non-const expression in global_expressions arena"); + } + U32EvalError::Negative => ResolveArraySizeError::ExpectedPositiveArrayLength, + })?; + + if length == 0 { + return Err(ResolveArraySizeError::ExpectedPositiveArrayLength); + } + + Ok(ResolvedSize::Constant(length)) + } + crate::ArraySize::Dynamic => Ok(ResolvedSize::Runtime), + } +} + /// Return an iterator over the individual components assembled by a /// `Compose` expression. /// diff --git a/naga/src/valid/expression.rs b/naga/src/valid/expression.rs index 28b9321cce..e7d822f761 100644 --- a/naga/src/valid/expression.rs +++ b/naga/src/valid/expression.rs @@ -271,10 +271,16 @@ impl super::Validator { } } + let array_size_indexable_length = if self.allow_overrides { + crate::ArraySize::indexable_length + } else { + crate::ArraySize::indexable_length_resolved + }; + // If we know both the length and the index, we can do the // bounds check now. if let crate::proc::IndexableLength::Known(known_length) = - base_type.indexable_length(module)? + base_type.indexable_length(module, array_size_indexable_length)? { match module .to_ctx() diff --git a/naga/src/valid/mod.rs b/naga/src/valid/mod.rs index 6a81bd7c2d..9dedebdd16 100644 --- a/naga/src/valid/mod.rs +++ b/naga/src/valid/mod.rs @@ -332,8 +332,6 @@ pub enum ValidationError { handle: Handle, source: ConstExpressionError, }, - #[error("Array size expression {handle:?} is not strictly positive")] - ArraySizeError { handle: Handle }, #[error("Constant {handle:?} '{name}' is invalid")] Constant { handle: Handle, @@ -616,7 +614,7 @@ impl Validator { })?; if !self.allow_overrides { if let crate::TypeInner::Array { - size: crate::ArraySize::Pending(_), + size: crate::ArraySize::Pending(crate::PendingArraySize::Override(_)), .. } = ty.inner {