From a2a35f2eb38a5fd413c50fd3f6aff551a1ea5ea0 Mon Sep 17 00:00:00 2001 From: Jim Blandy Date: Fri, 4 Jun 2021 09:57:20 -0700 Subject: [PATCH] Permit dynamic indexing of arrays and matrices only behind a pointer. (#949) This makes Naga IR validation impose the restrictions added to WGSL in gpuweb/gpuweb#1801. Remove code in the SPIR-V writer to spill arrays to temporary variables in order to index them dynamically. If such IR is encountered, treat it as a failure of validation. --- src/back/spv/writer.rs | 78 +++-------------------------------------- src/lib.rs | 51 ++++++++++++++++++++++++++- src/proc/mod.rs | 22 ++++++++++++ src/proc/typifier.rs | 8 +++++ src/valid/expression.rs | 20 +++++++---- tests/in/access.wgsl | 6 ++-- tests/out/access.msl | 4 ++- tests/out/access.spvasm | 73 +++++++++++++++++++------------------- tests/out/access.wgsl | 6 ++-- tests/wgsl-errors.rs | 43 +++++++++++++++++++++++ 10 files changed, 187 insertions(+), 124 deletions(-) diff --git a/src/back/spv/writer.rs b/src/back/spv/writer.rs index f3918347ad..57342be2f2 100644 --- a/src/back/spv/writer.rs +++ b/src/back/spv/writer.rs @@ -72,7 +72,6 @@ struct Function { signature: Option, parameters: Vec, variables: crate::FastHashMap, LocalVariable>, - internal_variables: Vec, blocks: Vec, entry_point_context: Option, } @@ -89,9 +88,6 @@ impl Function { for local_var in self.variables.values() { local_var.instruction.to_words(sink); } - for internal_var in self.internal_variables.iter() { - internal_var.instruction.to_words(sink); - } } for instruction in block.body.iter() { instruction.to_words(sink); @@ -1324,62 +1320,6 @@ impl Writer { }) } - #[allow(clippy::too_many_arguments)] - fn promote_access_expression_to_variable( - &mut self, - ir_types: &Arena, - result_type_id: Word, - container_id: Word, - container_resolution: &TypeResolution, - index_id: Word, - element_ty: Handle, - block: &mut Block, - ) -> Result<(Word, LocalVariable), Error> { - let container_type_id = self.get_expression_type_id(container_resolution)?; - let pointer_type_id = self.id_gen.next(); - Instruction::type_pointer( - pointer_type_id, - spirv::StorageClass::Function, - container_type_id, - ) - .to_words(&mut self.logical_layout.declarations); - - let variable = { - let id = self.id_gen.next(); - LocalVariable { - id, - instruction: Instruction::variable( - pointer_type_id, - id, - spirv::StorageClass::Function, - None, - ), - } - }; - block - .body - .push(Instruction::store(variable.id, container_id, None)); - - let element_pointer_id = self.id_gen.next(); - let element_pointer_type_id = - self.get_pointer_id(ir_types, element_ty, spirv::StorageClass::Function)?; - block.body.push(Instruction::access_chain( - element_pointer_type_id, - element_pointer_id, - variable.id, - &[index_id], - )); - let id = self.id_gen.next(); - block.body.push(Instruction::load( - result_type_id, - id, - element_pointer_id, - None, - )); - - Ok((id, variable)) - } - fn is_intermediate( &self, expr_handle: Handle, @@ -1433,20 +1373,10 @@ impl Writer { )); id } - crate::TypeInner::Array { - base: ty_element, .. - } => { - let (id, variable) = self.promote_access_expression_to_variable( - &ir_module.types, - result_type_id, - base_id, - &fun_info[base].ty, - index_id, - ty_element, - block, - )?; - function.internal_variables.push(variable); - id + crate::TypeInner::Array { .. } => { + return Err(Error::Validation( + "dynamic indexing of arrays not permitted", + )); } ref other => { log::error!( diff --git a/src/lib.rs b/src/lib.rs index ad48cd4e7a..d863ea406f 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -795,9 +795,58 @@ bitflags::bitflags! { #[cfg_attr(feature = "deserialize", derive(Deserialize))] pub enum Expression { /// Array access with a computed index. + /// + /// ## Typing rules + /// + /// The `base` operand must be some composite type: [`Vector`], [`Matrix`], + /// [`Array`], a [`Pointer`] to one of those, or a [`ValuePointer`] with a + /// `size`. + /// + /// The `index` operand must be an integer, signed or unsigned. + /// + /// Indexing a [`Vector`] or [`Array`] produces a value of its element type. + /// Indexing a [`Matrix`] produces a [`Vector`]. + /// + /// Indexing a [`Pointer`] to an [`Array`] produces a [`Pointer`] to its + /// `base` type, taking on the `Pointer`'s storage class. + /// + /// Indexing a [`Pointer`] to a [`Vector`] produces a [`ValuePointer`] whose + /// size is `None`, taking on the [`Vector`]'s scalar kind and width and the + /// [`Pointer`]'s storage class. + /// + /// Indexing a [`Pointer`] to a [`Matrix`] produces a [`ValuePointer`] for a + /// column of the matrix: its size is the matrix's height, its `kind` is + /// [`Float`], and it inherits the [`Matrix`]'s width and the [`Pointer`]'s + /// storage class. + /// + /// ## Dynamic indexing restrictions + /// + /// To accommodate restrictions in some of the shader languages that Naga + /// targets, it is not permitted to subscript a matrix or array with a + /// dynamically computed index unless that matrix or array appears behind a + /// pointer. In other words, if the inner type of `base` is [`Array`] or + /// [`Matrix`], then `index` must be a constant. But if the type of `base` + /// is a [`Pointer`] to an array or matrix or a [`ValuePointer`] with a + /// `size`, then the index may be any expression of integer type. + /// + /// You can use the [`Expression::is_dynamic_index`] method to determine + /// whether a given index expression requires matrix or array base operands + /// to be behind a pointer. + /// + /// (It would be simpler to always require the use of `AccessIndex` when + /// subscripting arrays and matrices that are not behind pointers, but to + /// accommodate existing front ends, Naga also permits `Access`, with a + /// restricted `index`.) + /// + /// [`Vector`]: TypeInner::Vector + /// [`Matrix`]: TypeInner::Matrix + /// [`Array`]: TypeInner::Array + /// [`Pointer`]: TypeInner::Pointer + /// [`ValuePointer`]: TypeInner::ValuePointer + /// [`Float`]: ScalarKind::Float Access { base: Handle, - index: Handle, //int + index: Handle, }, /// Array access with a known index. AccessIndex { diff --git a/src/proc/mod.rs b/src/proc/mod.rs index 58926c1d83..6229f9390a 100644 --- a/src/proc/mod.rs +++ b/src/proc/mod.rs @@ -183,6 +183,28 @@ impl crate::Expression { _ => false, } } + + /// Return true if this expression is a dynamic array index, for [`Access`]. + /// + /// This method returns true if this expression is a dynamically computed + /// index, and as such can only be used to index matrices and arrays when + /// they appear behind a pointer. See the documentation for [`Access`] for + /// details. + /// + /// Note, this does not check the _type_ of the given expression. It's up to + /// the caller to establish that the `Access` expression is well-typed + /// through other means, like [`ResolveContext`]. + /// + /// [`Access`]: crate::Expression::Access + /// [`ResolveContext`]: crate::proc::ResolveContext + pub fn is_dynamic_index(&self, module: &crate::Module) -> bool { + if let Self::Constant(handle) = *self { + let constant = &module.constants[handle]; + constant.specialization.is_some() + } else { + true + } + } } impl crate::SampleLevel { diff --git a/src/proc/typifier.rs b/src/proc/typifier.rs index a07e3d43ae..c8d2b02df5 100644 --- a/src/proc/typifier.rs +++ b/src/proc/typifier.rs @@ -128,7 +128,15 @@ impl<'a> ResolveContext<'a> { let types = self.types; Ok(match *expr { crate::Expression::Access { base, .. } => match *past(base).inner_with(types) { + // Arrays and matrices can only be indexed dynamically behind a + // pointer, but that's a validation error, not a type error, so + // go ahead provide a type here. Ti::Array { base, .. } => TypeResolution::Handle(base), + Ti::Matrix { rows, width, .. } => TypeResolution::Value(Ti::Vector { + size: rows, + kind: crate::ScalarKind::Float, + width, + }), Ti::Vector { size: _, kind, diff --git a/src/valid/expression.rs b/src/valid/expression.rs index 69d0d9b9dd..561a088a7e 100644 --- a/src/valid/expression.rs +++ b/src/valid/expression.rs @@ -19,6 +19,8 @@ pub enum ExpressionError { InvalidIndexType(Handle), #[error("Accessing index {1} is out of {0:?} bounds")] IndexOutOfBounds(Handle, u32), + #[error("The expression {0:?} may only be indexed by a constant")] + IndexMustBeConstant(Handle), #[error("Function argument {0:?} doesn't exist")] FunctionArgumentDoesntExist(u32), #[error("Constant {0:?} doesn't exist")] @@ -142,17 +144,16 @@ impl super::Validator { let stages = match *expression { E::Access { base, index } => { - match *resolver.resolve(base)? { - Ti::Vector { .. } - | Ti::Matrix { .. } - | Ti::Array { .. } - | Ti::Pointer { .. } - | Ti::ValuePointer { size: Some(_), .. } => {} + // See the documentation for `Expression::Access`. + let dynamic_indexing_restricted = match *resolver.resolve(base)? { + Ti::Vector { .. } => false, + Ti::Matrix { .. } | Ti::Array { .. } => true, + Ti::Pointer { .. } | Ti::ValuePointer { size: Some(_), .. } => false, ref other => { log::error!("Indexing of {:?}", other); return Err(ExpressionError::InvalidBaseType(base)); } - } + }; match *resolver.resolve(index)? { //TODO: only allow one of these Ti::Scalar { @@ -168,6 +169,11 @@ impl super::Validator { return Err(ExpressionError::InvalidIndexType(index)); } } + if dynamic_indexing_restricted + && function.expressions[index].is_dynamic_index(module) + { + return Err(ExpressionError::IndexMustBeConstant(base)); + } ShaderStages::all() } E::AccessIndex { base, index } => { diff --git a/tests/in/access.wgsl b/tests/in/access.wgsl index b6ee083557..ec963e6dff 100644 --- a/tests/in/access.wgsl +++ b/tests/in/access.wgsl @@ -15,9 +15,9 @@ fn foo([[builtin(vertex_index)]] vi: u32) -> [[builtin(position)]] vec4 { let b = bar.matrix[index].x; let a = bar.data[arrayLength(&bar.data) - 1u]; - - let array = array(a, i32(b), 3, 4, 5); - let value = array[vi]; + + var c: array = array(a, i32(b), 3, 4, 5); + let value = c[vi]; return vec4(vec4(value)); } diff --git a/tests/out/access.msl b/tests/out/access.msl index 36fd357743..1fcde25825 100644 --- a/tests/out/access.msl +++ b/tests/out/access.msl @@ -24,5 +24,7 @@ vertex fooOutput foo( , device Bar& bar [[buffer(0)]] , constant _mslBufferSizes& _buffer_sizes [[buffer(24)]] ) { - return fooOutput { static_cast(metal::int4(type5 {bar.data[(1 + (_buffer_sizes.size0 - 64 - 4) / 4) - 1u], static_cast(bar.matrix[3u].x), 3, 4, 5}.inner[vi])) }; + type5 c; + for(int _i=0; _i<5; ++_i) c.inner[_i] = type5 {bar.data[(1 + (_buffer_sizes.size0 - 64 - 4) / 4) - 1u], static_cast(bar.matrix[3u].x), 3, 4, 5}.inner[_i]; + return fooOutput { static_cast(metal::int4(c.inner[vi])) }; } diff --git a/tests/out/access.spvasm b/tests/out/access.spvasm index 2f5c21665b..843872137c 100644 --- a/tests/out/access.spvasm +++ b/tests/out/access.spvasm @@ -6,14 +6,15 @@ OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Vertex %24 "foo" %19 %22 +OpEntryPoint Vertex %26 "foo" %21 %24 OpSource GLSL 450 OpName %14 "Bar" OpMemberName %14 0 "matrix" OpMemberName %14 1 "data" OpName %16 "bar" -OpName %19 "vi" -OpName %24 "foo" +OpName %18 "c" +OpName %21 "vi" +OpName %26 "foo" OpDecorate %13 ArrayStride 4 OpDecorate %14 Block OpMemberDecorate %14 0 Offset 0 @@ -23,8 +24,8 @@ OpMemberDecorate %14 1 Offset 64 OpDecorate %15 ArrayStride 4 OpDecorate %16 DescriptorSet 0 OpDecorate %16 Binding 0 -OpDecorate %19 BuiltIn VertexIndex -OpDecorate %22 BuiltIn Position +OpDecorate %21 BuiltIn VertexIndex +OpDecorate %24 BuiltIn Position %2 = OpTypeVoid %4 = OpTypeInt 32 0 %3 = OpConstant %4 3 @@ -41,39 +42,39 @@ OpDecorate %22 BuiltIn Position %15 = OpTypeArray %7 %6 %17 = OpTypePointer StorageBuffer %14 %16 = OpVariable %17 StorageBuffer -%20 = OpTypePointer Input %4 -%19 = OpVariable %20 Input -%23 = OpTypePointer Output %11 -%22 = OpVariable %23 Output -%25 = OpTypeFunction %2 -%27 = OpTypePointer StorageBuffer %10 -%28 = OpTypePointer StorageBuffer %11 -%29 = OpConstant %4 0 -%33 = OpTypePointer StorageBuffer %13 -%36 = OpTypePointer StorageBuffer %7 -%41 = OpTypePointer Function %15 -%44 = OpTypePointer Function %7 +%19 = OpTypePointer Function %15 +%22 = OpTypePointer Input %4 +%21 = OpVariable %22 Input +%25 = OpTypePointer Output %11 +%24 = OpVariable %25 Output +%27 = OpTypeFunction %2 +%29 = OpTypePointer StorageBuffer %10 +%30 = OpTypePointer StorageBuffer %11 +%31 = OpConstant %4 0 +%35 = OpTypePointer StorageBuffer %13 +%38 = OpTypePointer StorageBuffer %7 +%43 = OpTypePointer Function %7 %46 = OpTypeVector %7 4 -%24 = OpFunction %2 None %25 -%18 = OpLabel -%42 = OpVariable %41 Function -%21 = OpLoad %4 %19 -OpBranch %26 -%26 = OpLabel -%30 = OpAccessChain %28 %16 %29 %3 -%31 = OpLoad %11 %30 -%32 = OpCompositeExtract %12 %31 0 -%34 = OpArrayLength %4 %16 1 -%35 = OpISub %4 %34 %5 -%37 = OpAccessChain %36 %16 %5 %35 -%38 = OpLoad %7 %37 -%39 = OpConvertFToS %7 %32 -%40 = OpCompositeConstruct %15 %38 %39 %8 %9 %6 -OpStore %42 %40 -%43 = OpAccessChain %44 %42 %21 -%45 = OpLoad %7 %43 +%26 = OpFunction %2 None %27 +%20 = OpLabel +%18 = OpVariable %19 Function +%23 = OpLoad %4 %21 +OpBranch %28 +%28 = OpLabel +%32 = OpAccessChain %30 %16 %31 %3 +%33 = OpLoad %11 %32 +%34 = OpCompositeExtract %12 %33 0 +%36 = OpArrayLength %4 %16 1 +%37 = OpISub %4 %36 %5 +%39 = OpAccessChain %38 %16 %5 %37 +%40 = OpLoad %7 %39 +%41 = OpConvertFToS %7 %34 +%42 = OpCompositeConstruct %15 %40 %41 %8 %9 %6 +OpStore %18 %42 +%44 = OpAccessChain %43 %18 %23 +%45 = OpLoad %7 %44 %47 = OpCompositeConstruct %46 %45 %45 %45 %45 %48 = OpConvertSToF %11 %47 -OpStore %22 %48 +OpStore %24 %48 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/access.wgsl b/tests/out/access.wgsl index 9a75615ab4..c0926ea803 100644 --- a/tests/out/access.wgsl +++ b/tests/out/access.wgsl @@ -9,9 +9,11 @@ var bar: [[access(read_write)]] Bar; [[stage(vertex)]] fn foo([[builtin(vertex_index)]] vi: u32) -> [[builtin(position)]] vec4 { + var c: array; + let b: f32 = bar.matrix[3u].x; let a: i32 = bar.data[(arrayLength(&bar.data) - 1u)]; - let array: array = array(a, i32(b), 3, 4, 5); - let value: i32 = array[vi]; + c = array(a, i32(b), 3, 4, 5); + let value: i32 = c[vi]; return vec4(vec4(value)); } diff --git a/tests/wgsl-errors.rs b/tests/wgsl-errors.rs index d11cef3f96..fb49a40564 100644 --- a/tests/wgsl-errors.rs +++ b/tests/wgsl-errors.rs @@ -314,3 +314,46 @@ fn missing_bindings() { }) } } + +#[test] +fn invalid_access() { + check_validation_error! { + " + fn array_by_value(a: array, i: i32) -> i32 { + return a[i]; + } + ", + " + fn matrix_by_value(m: mat4x4, i: i32) -> vec4 { + return m[i]; + } + ": + Err(naga::valid::ValidationError::Function { + error: naga::valid::FunctionError::Expression { + error: naga::valid::ExpressionError::IndexMustBeConstant(_), + .. + }, + .. + }) + } +} + +#[test] +fn valid_access() { + check_validation_error! { + " + fn vector_by_value(v: vec4, i: i32) -> i32 { + return v[i]; + } + ", + " + fn matrix_dynamic(m: mat4x4, i: i32, j: i32) -> f32 { + var temp: mat4x4 = m; + // Dynamically indexing the column vector applies + // `Access` to a `ValuePointer`. + return temp[i][j]; + } + ": + Ok(_) + } +}