github/wgpu
1
1
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2026-04-22 03:02:01 -04:00
Code Issues Packages Projects Releases Wiki Activity

Permit dynamic indexing of arrays and matrices only behind a pointer. (#949)

Browse Source 
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.
This commit is contained in:
Jim Blandy
2021-06-04 09:57:20 -07:00
committed by GitHub
parent 1c3baf4557
commit a2a35f2eb3
10 changed files with 187 additions and 124 deletions
Download Patch File Download Diff File Expand all files Collapse all files

78
src/back/spv/writer.rs
View File

@@ -72,7 +72,6 @@ struct Function {
signature: Option<Instruction>,
parameters: Vec<Instruction>,
variables: crate::FastHashMap<Handle<crate::LocalVariable>, LocalVariable>,
internal_variables: Vec<LocalVariable>,
blocks: Vec<Block>,
entry_point_context: Option<EntryPointContext>,
}
@@ -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<crate::Type>,
result_type_id: Word,
container_id: Word,
container_resolution: &TypeResolution,
index_id: Word,
element_ty: Handle<crate::Type>,
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<crate::Expression>,
@@ -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!(

51
src/lib.rs
View File

@@ -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<Expression>,
index: Handle<Expression>, //int
index: Handle<Expression>,
},
/// Array access with a known index.
AccessIndex {

22
src/proc/mod.rs
View File

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

8
src/proc/typifier.rs
View File

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

20
src/valid/expression.rs
View File

@@ -19,6 +19,8 @@ pub enum ExpressionError {
InvalidIndexType(Handle<crate::Expression>),
#[error("Accessing index {1} is out of {0:?} bounds")]
IndexOutOfBounds(Handle<crate::Expression>, u32),
#[error("The expression {0:?} may only be indexed by a constant")]
IndexMustBeConstant(Handle<crate::Expression>),
#[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 } => {

6
tests/in/access.wgsl
View File

@@ -15,9 +15,9 @@ fn foo([[builtin(vertex_index)]] vi: u32) -> [[builtin(position)]] vec4<f32> {
let b = bar.matrix[index].x;
let a = bar.data[arrayLength(&bar.data) - 1u];
let array = array<i32, 5>(a, i32(b), 3, 4, 5);
let value = array[vi];
var c: array<i32, 5> = array<i32, 5>(a, i32(b), 3, 4, 5);
let value = c[vi];
return vec4<f32>(vec4<i32>(value));
}

4
tests/out/access.msl
View File

@@ -24,5 +24,7 @@ vertex fooOutput foo(
, device Bar& bar [[buffer(0)]]
, constant _mslBufferSizes& _buffer_sizes [[buffer(24)]]
) {
return fooOutput { static_cast<float4>(metal::int4(type5 {bar.data[(1 + (_buffer_sizes.size0 - 64 - 4) / 4) - 1u], static_cast<int>(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<int>(bar.matrix[3u].x), 3, 4, 5}.inner[_i];
return fooOutput { static_cast<float4>(metal::int4(c.inner[vi])) };
}

73
tests/out/access.spvasm
View File

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

6
tests/out/access.wgsl
View File

@@ -9,9 +9,11 @@ var<storage> bar: [[access(read_write)]] Bar;
[[stage(vertex)]]
fn foo([[builtin(vertex_index)]] vi: u32) -> [[builtin(position)]] vec4<f32> {
var c: array<i32,5>;
let b: f32 = bar.matrix[3u].x;
let a: i32 = bar.data[(arrayLength(&bar.data) - 1u)];
let array: array<i32,5> = array<i32,5>(a, i32(b), 3, 4, 5);
let value: i32 = array[vi];
c = array<i32,5>(a, i32(b), 3, 4, 5);
let value: i32 = c[vi];
return vec4<f32>(vec4<i32>(value));
}

43
tests/wgsl-errors.rs
View File

@@ -314,3 +314,46 @@ fn missing_bindings() {
})
}
}
#[test]
fn invalid_access() {
check_validation_error! {
"
fn array_by_value(a: array<i32, 5>, i: i32) -> i32 {
return a[i];
}
",
"
fn matrix_by_value(m: mat4x4<f32>, i: i32) -> vec4<f32> {
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<i32>, i: i32) -> i32 {
return v[i];
}
",
"
fn matrix_dynamic(m: mat4x4<f32>, i: i32, j: i32) -> f32 {
var temp: mat4x4<f32> = m;
// Dynamically indexing the column vector applies
// `Access` to a `ValuePointer`.
return temp[i][j];
}
":
Ok(_)
}
}