From 2d82054ae4af209c3cbe297039469941b2f3b4e6 Mon Sep 17 00:00:00 2001 From: sagudev <16504129+sagudev@users.noreply.github.com> Date: Sun, 1 Sep 2024 16:12:06 +0200 Subject: [PATCH] [wgsl-in, spv-out] Allow dynamic indexing of arrays by value. Bring https://github.com/gfx-rs/naga/pull/723 back from the dead. Signed-off-by: sagudev <16504129+sagudev@users.noreply.github.com> Co-authored-by: Dzmitry Malyshau Co-authored-by: Jim Blandy --- naga/src/back/spv/block.rs | 28 +- naga/src/back/spv/mod.rs | 1 + naga/src/back/spv/writer.rs | 53 ++++ naga/src/lib.rs | 21 +- naga/src/proc/mod.rs | 8 +- naga/src/proc/typifier.rs | 7 + naga/src/valid/expression.rs | 7 +- naga/tests/in/access.wgsl | 11 + naga/tests/out/analysis/access.info.ron | 186 ++++++++++++ .../access.assign_through_ptr.Compute.glsl | 12 +- naga/tests/out/glsl/access.foo.Vertex.glsl | 52 ++++ .../out/glsl/access.foo_frag.Fragment.glsl | 12 +- .../out/glsl/access.foo_vert.Vertex.glsl | 16 +- naga/tests/out/hlsl/access.hlsl | 24 +- naga/tests/out/hlsl/access.ron | 4 + naga/tests/out/ir/access.compact.ron | 117 ++++++++ naga/tests/out/ir/access.ron | 117 ++++++++ naga/tests/out/msl/access.msl | 33 ++- naga/tests/out/spv/access.spvasm | 273 ++++++++++-------- naga/tests/out/wgsl/access.wgsl | 23 +- naga/tests/wgsl_errors.rs | 5 - 21 files changed, 837 insertions(+), 173 deletions(-) create mode 100644 naga/tests/out/glsl/access.foo.Vertex.glsl diff --git a/naga/src/back/spv/block.rs b/naga/src/back/spv/block.rs index f0c3bfa84..97db1cb54 100644 --- a/naga/src/back/spv/block.rs +++ b/naga/src/back/spv/block.rs @@ -343,6 +343,32 @@ impl<'w> BlockContext<'w> { load_id } + crate::TypeInner::Array { + base: ty_element, .. + } => { + let index_id = self.cached[index]; + let base_id = self.cached[base]; + let base_ty = match self.fun_info[base].ty { + TypeResolution::Handle(handle) => handle, + TypeResolution::Value(_) => { + return Err(Error::Validation( + "Array types should always be in the arena", + )) + } + }; + let (id, variable) = self.writer.promote_access_expression_to_variable( + &self.ir_module.types, + result_type_id, + base_id, + base_ty, + index_id, + ty_element, + block, + )?; + self.function.internal_variables.push(variable); + id + } + // wgpu#4337: Support `crate::TypeInner::Matrix` ref other => { log::error!( "Unable to access base {:?} of type {:?}", @@ -350,7 +376,7 @@ impl<'w> BlockContext<'w> { other ); return Err(Error::Validation( - "only vectors may be dynamically indexed by value", + "only vectors and arrays may be dynamically indexed by value", )); } } diff --git a/naga/src/back/spv/mod.rs b/naga/src/back/spv/mod.rs index 32bd1fcec..77f915abb 100644 --- a/naga/src/back/spv/mod.rs +++ b/naga/src/back/spv/mod.rs @@ -144,6 +144,7 @@ struct Function { signature: Option, parameters: Vec, variables: crate::FastHashMap, LocalVariable>, + internal_variables: Vec, blocks: Vec, entry_point_context: Option, } diff --git a/naga/src/back/spv/writer.rs b/naga/src/back/spv/writer.rs index 678dcb424..7f41caf29 100644 --- a/naga/src/back/spv/writer.rs +++ b/naga/src/back/spv/writer.rs @@ -32,6 +32,9 @@ 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); @@ -135,6 +138,56 @@ impl Writer { self.capabilities_used.insert(spirv::Capability::Shader); } + #[allow(clippy::too_many_arguments)] + pub(super) fn promote_access_expression_to_variable( + &mut self, + ir_types: &UniqueArena, + result_type_id: Word, + container_id: Word, + container_ty: Handle, + index_id: Word, + element_ty: Handle, + block: &mut Block, + ) -> Result<(Word, LocalVariable), Error> { + let pointer_type_id = + self.get_pointer_id(ir_types, container_ty, spirv::StorageClass::Function)?; + + 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)) + } + /// Indicate that the code requires any one of the listed capabilities. /// /// If nothing in `capabilities` appears in the available capabilities diff --git a/naga/src/lib.rs b/naga/src/lib.rs index 85fd7a450..038e215a6 100644 --- a/naga/src/lib.rs +++ b/naga/src/lib.rs @@ -1402,21 +1402,20 @@ pub enum Expression { /// ## 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. + /// targets, it is not permitted to subscript a matrix with a dynamically + /// computed index unless that matrix appears behind a pointer. In other + /// words, if the inner type of `base` is [`Matrix`], then `index` must be a + /// constant. But if the type of `base` is a [`Pointer`] to an matrix, 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. + /// whether a given index expression requires matrix 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`.) + /// subscripting 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 diff --git a/naga/src/proc/mod.rs b/naga/src/proc/mod.rs index a5b3ea4e3..abbe0c7e4 100644 --- a/naga/src/proc/mod.rs +++ b/naga/src/proc/mod.rs @@ -521,12 +521,12 @@ impl crate::Expression { } } - /// Return true if this expression is a dynamic array index, for [`Access`]. + /// Return true if this expression is a dynamic array/vector/matrix 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. + /// index, and as such can only be used to index matrices 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 diff --git a/naga/src/proc/typifier.rs b/naga/src/proc/typifier.rs index d8af0cd23..f29ff40f0 100644 --- a/naga/src/proc/typifier.rs +++ b/naga/src/proc/typifier.rs @@ -92,6 +92,13 @@ pub enum TypeResolution { /// available in the associated arena. However, the `TypeInner` itself may /// contain `Handle` values referring to types from the arena. /// + /// The inner type must only be one of the following variants: + /// - TypeInner::Pointer + /// - TypeInner::ValuePointer + /// - TypeInner::Matrix (generated by matrix multiplication) + /// - TypeInner::Vector + /// - TypeInner::Scalar + /// /// [`TypeInner`]: crate::TypeInner Value(crate::TypeInner), } diff --git a/naga/src/valid/expression.rs b/naga/src/valid/expression.rs index 0b0d115c5..2b479d3a7 100644 --- a/naga/src/valid/expression.rs +++ b/naga/src/valid/expression.rs @@ -240,9 +240,10 @@ impl super::Validator { let base_type = &resolver[base]; // See the documentation for `Expression::Access`. let dynamic_indexing_restricted = match *base_type { - Ti::Vector { .. } => false, - Ti::Matrix { .. } | Ti::Array { .. } => true, - Ti::Pointer { .. } + Ti::Matrix { .. } => true, + Ti::Vector { .. } + | Ti::Array { .. } + | Ti::Pointer { .. } | Ti::ValuePointer { size: Some(_), .. } | Ti::BindingArray { .. } => false, ref other => { diff --git a/naga/tests/in/access.wgsl b/naga/tests/in/access.wgsl index 956a694aa..3336522fd 100644 --- a/naga/tests/in/access.wgsl +++ b/naga/tests/in/access.wgsl @@ -167,3 +167,14 @@ fn assign_through_ptr() { var arr = array, 2>(vec4(6.0), vec4(7.0)); assign_array_through_ptr_fn(&arr); } + +@vertex +fn foo(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { + let arr = array(1, 2, 3, 4, 5); + let value = arr[vi]; + return vec4(vec4(value)); +} + +fn array_by_value(a: array, i: i32) -> i32 { + return a[i]; +} diff --git a/naga/tests/out/analysis/access.info.ron b/naga/tests/out/analysis/access.info.ron index 308bb1a8b..830831cb1 100644 --- a/naga/tests/out/analysis/access.info.ron +++ b/naga/tests/out/analysis/access.info.ron @@ -2735,6 +2735,54 @@ sampling: [], dual_source_blending: false, ), + ( + flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE"), + uniformity: ( + non_uniform_result: Some(0), + requirements: (""), + ), + may_kill: false, + sampling_set: [], + global_uses: [ + (""), + (""), + (""), + (""), + (""), + ], + expressions: [ + ( + uniformity: ( + non_uniform_result: Some(0), + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Handle(25), + ), + ( + uniformity: ( + non_uniform_result: Some(1), + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Handle(2), + ), + ( + uniformity: ( + non_uniform_result: Some(0), + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Handle(2), + ), + ], + sampling: [], + dual_source_blending: false, + ), ], entry_points: [ ( @@ -3981,6 +4029,144 @@ sampling: [], dual_source_blending: false, ), + ( + flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE"), + uniformity: ( + non_uniform_result: Some(0), + requirements: (""), + ), + may_kill: false, + sampling_set: [], + global_uses: [ + (""), + (""), + (""), + (""), + (""), + ], + expressions: [ + ( + uniformity: ( + non_uniform_result: Some(0), + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Handle(0), + ), + ( + uniformity: ( + non_uniform_result: None, + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Scalar(( + kind: Sint, + width: 4, + ))), + ), + ( + uniformity: ( + non_uniform_result: None, + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Scalar(( + kind: Sint, + width: 4, + ))), + ), + ( + uniformity: ( + non_uniform_result: None, + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Scalar(( + kind: Sint, + width: 4, + ))), + ), + ( + uniformity: ( + non_uniform_result: None, + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Scalar(( + kind: Sint, + width: 4, + ))), + ), + ( + uniformity: ( + non_uniform_result: None, + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Scalar(( + kind: Sint, + width: 4, + ))), + ), + ( + uniformity: ( + non_uniform_result: None, + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Handle(25), + ), + ( + uniformity: ( + non_uniform_result: Some(0), + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Handle(2), + ), + ( + uniformity: ( + non_uniform_result: Some(0), + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Vector( + size: Quad, + scalar: ( + kind: Sint, + width: 4, + ), + )), + ), + ( + uniformity: ( + non_uniform_result: Some(0), + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Vector( + size: Quad, + scalar: ( + kind: Float, + width: 4, + ), + )), + ), + ], + sampling: [], + dual_source_blending: false, + ), ], const_expression_types: [ Value(Scalar(( diff --git a/naga/tests/out/glsl/access.assign_through_ptr.Compute.glsl b/naga/tests/out/glsl/access.assign_through_ptr.Compute.glsl index 2e51bbde6..4a4791c99 100644 --- a/naga/tests/out/glsl/access.assign_through_ptr.Compute.glsl +++ b/naga/tests/out/glsl/access.assign_through_ptr.Compute.glsl @@ -20,8 +20,8 @@ struct MatCx2InArray { mat4x2 am[2]; }; -float read_from_private(inout float foo_1) { - float _e1 = foo_1; +float read_from_private(inout float foo_2) { + float _e1 = foo_2; return _e1; } @@ -34,11 +34,15 @@ void assign_through_ptr_fn(inout uint p) { return; } -void assign_array_through_ptr_fn(inout vec4 foo_2[2]) { - foo_2 = vec4[2](vec4(1.0), vec4(2.0)); +void assign_array_through_ptr_fn(inout vec4 foo_3[2]) { + foo_3 = vec4[2](vec4(1.0), vec4(2.0)); return; } +int array_by_value(int a_1[5], int i) { + return a_1[i]; +} + void main() { uint val = 33u; vec4 arr[2] = vec4[2](vec4(6.0), vec4(7.0)); diff --git a/naga/tests/out/glsl/access.foo.Vertex.glsl b/naga/tests/out/glsl/access.foo.Vertex.glsl new file mode 100644 index 000000000..e1f313840 --- /dev/null +++ b/naga/tests/out/glsl/access.foo.Vertex.glsl @@ -0,0 +1,52 @@ +#version 310 es + +precision highp float; +precision highp int; + +struct GlobalConst { + uint a; + uvec3 b; + int c; +}; +struct AlignedWrapper { + int value; +}; +struct Baz { + mat3x2 m; +}; +struct MatCx2InArray { + mat4x2 am[2]; +}; + +float read_from_private(inout float foo_2) { + float _e1 = foo_2; + return _e1; +} + +float test_arr_as_arg(float a[5][10]) { + return a[4][9]; +} + +void assign_through_ptr_fn(inout uint p) { + p = 42u; + return; +} + +void assign_array_through_ptr_fn(inout vec4 foo_3[2]) { + foo_3 = vec4[2](vec4(1.0), vec4(2.0)); + return; +} + +int array_by_value(int a_1[5], int i) { + return a_1[i]; +} + +void main() { + uint vi_1 = uint(gl_VertexID); + int arr_1[5] = int[5](1, 2, 3, 4, 5); + int value = arr_1[vi_1]; + gl_Position = vec4(ivec4(value)); + gl_Position.yz = vec2(-gl_Position.y, gl_Position.z * 2.0 - gl_Position.w); + return; +} + diff --git a/naga/tests/out/glsl/access.foo_frag.Fragment.glsl b/naga/tests/out/glsl/access.foo_frag.Fragment.glsl index aacdda013..eca6bc54c 100644 --- a/naga/tests/out/glsl/access.foo_frag.Fragment.glsl +++ b/naga/tests/out/glsl/access.foo_frag.Fragment.glsl @@ -30,8 +30,8 @@ layout(std430) buffer type_13_block_1Fragment { ivec2 _group_0_binding_2_fs; }; layout(location = 0) out vec4 _fs2p_location0; -float read_from_private(inout float foo_1) { - float _e1 = foo_1; +float read_from_private(inout float foo_2) { + float _e1 = foo_2; return _e1; } @@ -44,11 +44,15 @@ void assign_through_ptr_fn(inout uint p) { return; } -void assign_array_through_ptr_fn(inout vec4 foo_2[2]) { - foo_2 = vec4[2](vec4(1.0), vec4(2.0)); +void assign_array_through_ptr_fn(inout vec4 foo_3[2]) { + foo_3 = vec4[2](vec4(1.0), vec4(2.0)); return; } +int array_by_value(int a_1[5], int i) { + return a_1[i]; +} + void main() { _group_0_binding_0_fs._matrix[1][2] = 1.0; _group_0_binding_0_fs._matrix = mat4x3(vec3(0.0), vec3(1.0), vec3(2.0), vec3(3.0)); diff --git a/naga/tests/out/glsl/access.foo_vert.Vertex.glsl b/naga/tests/out/glsl/access.foo_vert.Vertex.glsl index d4a9b9294..a926eadf7 100644 --- a/naga/tests/out/glsl/access.foo_vert.Vertex.glsl +++ b/naga/tests/out/glsl/access.foo_vert.Vertex.glsl @@ -103,8 +103,8 @@ void test_matrix_within_array_within_struct_accesses() { return; } -float read_from_private(inout float foo_1) { - float _e1 = foo_1; +float read_from_private(inout float foo_2) { + float _e1 = foo_2; return _e1; } @@ -117,11 +117,15 @@ void assign_through_ptr_fn(inout uint p) { return; } -void assign_array_through_ptr_fn(inout vec4 foo_2[2]) { - foo_2 = vec4[2](vec4(1.0), vec4(2.0)); +void assign_array_through_ptr_fn(inout vec4 foo_3[2]) { + foo_3 = vec4[2](vec4(1.0), vec4(2.0)); return; } +int array_by_value(int a_1[5], int i) { + return a_1[i]; +} + void main() { uint vi = uint(gl_VertexID); float foo = 0.0; @@ -133,10 +137,10 @@ void main() { mat4x3 _matrix = _group_0_binding_0_vs._matrix; uvec2 arr_1[2] = _group_0_binding_0_vs.arr; float b = _group_0_binding_0_vs._matrix[3u][0]; - int a_1 = _group_0_binding_0_vs.data[(uint(_group_0_binding_0_vs.data.length()) - 2u)].value; + int a_2 = _group_0_binding_0_vs.data[(uint(_group_0_binding_0_vs.data.length()) - 2u)].value; ivec2 c = _group_0_binding_2_vs; float _e33 = read_from_private(foo); - c2_ = int[5](a_1, int(b), 3, 4, 5); + c2_ = int[5](a_2, int(b), 3, 4, 5); c2_[(vi + 1u)] = 42; int value = c2_[vi]; float _e47 = test_arr_as_arg(float[5][10](float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0))); diff --git a/naga/tests/out/hlsl/access.hlsl b/naga/tests/out/hlsl/access.hlsl index 142083be6..543b84196 100644 --- a/naga/tests/out/hlsl/access.hlsl +++ b/naga/tests/out/hlsl/access.hlsl @@ -201,9 +201,9 @@ void test_matrix_within_array_within_struct_accesses() return; } -float read_from_private(inout float foo_1) +float read_from_private(inout float foo_2) { - float _e1 = foo_1; + float _e1 = foo_2; return _e1; } @@ -224,12 +224,17 @@ ret_Constructarray2_float4_ Constructarray2_float4_(float4 arg0, float4 arg1) { return ret; } -void assign_array_through_ptr_fn(inout float4 foo_2[2]) +void assign_array_through_ptr_fn(inout float4 foo_3[2]) { - foo_2 = Constructarray2_float4_((1.0).xxxx, (2.0).xxxx); + foo_3 = Constructarray2_float4_((1.0).xxxx, (2.0).xxxx); return; } +int array_by_value(int a_1[5], int i) +{ + return a_1[i]; +} + typedef int ret_Constructarray5_int_[5]; ret_Constructarray5_int_ Constructarray5_int_(int arg0, int arg1, int arg2, int arg3, int arg4) { int ret[5] = { arg0, arg1, arg2, arg3, arg4 }; @@ -266,10 +271,10 @@ float4 foo_vert(uint vi : SV_VertexID) : SV_Position float4x3 _matrix = float4x3(asfloat(bar.Load3(0+0)), asfloat(bar.Load3(0+16)), asfloat(bar.Load3(0+32)), asfloat(bar.Load3(0+48))); uint2 arr_1[2] = Constructarray2_uint2_(asuint(bar.Load2(144+0)), asuint(bar.Load2(144+8))); float b = asfloat(bar.Load(0+3u*16+0)); - int a_1 = asint(bar.Load(0+(((NagaBufferLengthRW(bar) - 160) / 8) - 2u)*8+160)); + int a_2 = asint(bar.Load(0+(((NagaBufferLengthRW(bar) - 160) / 8) - 2u)*8+160)); int2 c = asint(qux.Load2(0)); const float _e33 = read_from_private(foo); - c2_ = Constructarray5_int_(a_1, int(b), 3, 4, 5); + c2_ = Constructarray5_int_(a_2, int(b), 3, 4, 5); c2_[(vi + 1u)] = 42; int value = c2_[vi]; const float _e47 = test_arr_as_arg(ZeroValuearray5_array10_float__()); @@ -310,3 +315,10 @@ void assign_through_ptr() assign_array_through_ptr_fn(arr); return; } + +float4 foo_1(uint vi_1 : SV_VertexID) : SV_Position +{ + int arr_2[5] = Constructarray5_int_(1, 2, 3, 4, 5); + int value_1 = arr_2[vi_1]; + return float4((value_1).xxxx); +} diff --git a/naga/tests/out/hlsl/access.ron b/naga/tests/out/hlsl/access.ron index 73c9e4444..8960a612e 100644 --- a/naga/tests/out/hlsl/access.ron +++ b/naga/tests/out/hlsl/access.ron @@ -4,6 +4,10 @@ entry_point:"foo_vert", target_profile:"vs_5_1", ), + ( + entry_point:"foo_1", + target_profile:"vs_5_1", + ), ], fragment:[ ( diff --git a/naga/tests/out/ir/access.compact.ron b/naga/tests/out/ir/access.compact.ron index 1b95742ff..2d066b8ff 100644 --- a/naga/tests/out/ir/access.compact.ron +++ b/naga/tests/out/ir/access.compact.ron @@ -1655,6 +1655,47 @@ ), ], ), + ( + name: Some("array_by_value"), + arguments: [ + ( + name: Some("a"), + ty: 25, + binding: None, + ), + ( + name: Some("i"), + ty: 2, + binding: None, + ), + ], + result: Some(( + ty: 2, + binding: None, + )), + local_variables: [], + expressions: [ + FunctionArgument(0), + FunctionArgument(1), + Access( + base: 0, + index: 1, + ), + ], + named_expressions: { + 0: "a", + 1: "i", + }, + body: [ + Emit(( + start: 2, + end: 3, + )), + Return( + value: Some(2), + ), + ], + ), ], entry_points: [ ( @@ -2230,5 +2271,81 @@ ], ), ), + ( + name: "foo", + stage: Vertex, + early_depth_test: None, + workgroup_size: (0, 0, 0), + function: ( + name: Some("foo"), + arguments: [ + ( + name: Some("vi"), + ty: 0, + binding: Some(BuiltIn(VertexIndex)), + ), + ], + result: Some(( + ty: 24, + binding: Some(BuiltIn(Position( + invariant: false, + ))), + )), + local_variables: [], + expressions: [ + FunctionArgument(0), + Literal(I32(1)), + Literal(I32(2)), + Literal(I32(3)), + Literal(I32(4)), + Literal(I32(5)), + Compose( + ty: 25, + components: [ + 1, + 2, + 3, + 4, + 5, + ], + ), + Access( + base: 6, + index: 0, + ), + Splat( + size: Quad, + value: 7, + ), + As( + expr: 8, + kind: Float, + convert: Some(4), + ), + ], + named_expressions: { + 0: "vi", + 6: "arr", + 7: "value", + }, + body: [ + Emit(( + start: 6, + end: 7, + )), + Emit(( + start: 7, + end: 8, + )), + Emit(( + start: 8, + end: 10, + )), + Return( + value: Some(9), + ), + ], + ), + ), ], ) \ No newline at end of file diff --git a/naga/tests/out/ir/access.ron b/naga/tests/out/ir/access.ron index 1b95742ff..2d066b8ff 100644 --- a/naga/tests/out/ir/access.ron +++ b/naga/tests/out/ir/access.ron @@ -1655,6 +1655,47 @@ ), ], ), + ( + name: Some("array_by_value"), + arguments: [ + ( + name: Some("a"), + ty: 25, + binding: None, + ), + ( + name: Some("i"), + ty: 2, + binding: None, + ), + ], + result: Some(( + ty: 2, + binding: None, + )), + local_variables: [], + expressions: [ + FunctionArgument(0), + FunctionArgument(1), + Access( + base: 0, + index: 1, + ), + ], + named_expressions: { + 0: "a", + 1: "i", + }, + body: [ + Emit(( + start: 2, + end: 3, + )), + Return( + value: Some(2), + ), + ], + ), ], entry_points: [ ( @@ -2230,5 +2271,81 @@ ], ), ), + ( + name: "foo", + stage: Vertex, + early_depth_test: None, + workgroup_size: (0, 0, 0), + function: ( + name: Some("foo"), + arguments: [ + ( + name: Some("vi"), + ty: 0, + binding: Some(BuiltIn(VertexIndex)), + ), + ], + result: Some(( + ty: 24, + binding: Some(BuiltIn(Position( + invariant: false, + ))), + )), + local_variables: [], + expressions: [ + FunctionArgument(0), + Literal(I32(1)), + Literal(I32(2)), + Literal(I32(3)), + Literal(I32(4)), + Literal(I32(5)), + Compose( + ty: 25, + components: [ + 1, + 2, + 3, + 4, + 5, + ], + ), + Access( + base: 6, + index: 0, + ), + Splat( + size: Quad, + value: 7, + ), + As( + expr: 8, + kind: Float, + convert: Some(4), + ), + ], + named_expressions: { + 0: "vi", + 6: "arr", + 7: "value", + }, + body: [ + Emit(( + start: 6, + end: 7, + )), + Emit(( + start: 7, + end: 8, + )), + Emit(( + start: 8, + end: 10, + )), + Return( + value: Some(9), + ), + ], + ), + ), ], ) \ No newline at end of file diff --git a/naga/tests/out/msl/access.msl b/naga/tests/out/msl/access.msl index 65dba4910..924b604e4 100644 --- a/naga/tests/out/msl/access.msl +++ b/naga/tests/out/msl/access.msl @@ -133,9 +133,9 @@ void test_matrix_within_array_within_struct_accesses( } float read_from_private( - thread float& foo_1 + thread float& foo_2 ) { - float _e1 = foo_1; + float _e1 = foo_2; return _e1; } @@ -153,12 +153,19 @@ void assign_through_ptr_fn( } void assign_array_through_ptr_fn( - thread type_22& foo_2 + thread type_22& foo_3 ) { - foo_2 = type_22 {metal::float4(1.0), metal::float4(2.0)}; + foo_3 = type_22 {metal::float4(1.0), metal::float4(2.0)}; return; } +int array_by_value( + type_20 a_1, + int i +) { + return a_1.inner[i]; +} + struct foo_vertInput { }; struct foo_vertOutput { @@ -181,10 +188,10 @@ vertex foo_vertOutput foo_vert( metal::float4x3 _matrix = bar._matrix; type_10 arr_1 = bar.arr; float b = bar._matrix[3u].x; - int a_1 = bar.data[(1 + (_buffer_sizes.size1 - 160 - 8) / 8) - 2u].value; + int a_2 = bar.data[(1 + (_buffer_sizes.size1 - 160 - 8) / 8) - 2u].value; metal::int2 c = qux; float _e33 = read_from_private(foo); - c2_ = type_20 {a_1, static_cast(b), 3, 4, 5}; + c2_ = type_20 {a_2, static_cast(b), 3, 4, 5}; c2_.inner[vi + 1u] = 42; int value = c2_.inner[vi]; float _e47 = test_arr_as_arg(type_18 {}); @@ -217,3 +224,17 @@ kernel void assign_through_ptr( assign_array_through_ptr_fn(arr); return; } + + +struct foo_1Input { +}; +struct foo_1Output { + metal::float4 member_3 [[position]]; +}; +vertex foo_1Output foo_1( + uint vi_1 [[vertex_id]] +) { + type_20 arr_2 = type_20 {1, 2, 3, 4, 5}; + int value_1 = arr_2.inner[vi_1]; + return foo_1Output { static_cast(metal::int4(value_1)) }; +} diff --git a/naga/tests/out/spv/access.spvasm b/naga/tests/out/spv/access.spvasm index ab0112870..6cb87a3b1 100644 --- a/naga/tests/out/spv/access.spvasm +++ b/naga/tests/out/spv/access.spvasm @@ -1,16 +1,17 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 301 +; Bound: 323 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Vertex %219 "foo_vert" %214 %217 -OpEntryPoint Fragment %273 "foo_frag" %272 -OpEntryPoint GLCompute %291 "assign_through_ptr" -OpExecutionMode %273 OriginUpperLeft -OpExecutionMode %291 LocalSize 1 1 1 +OpEntryPoint Vertex %229 "foo_vert" %224 %227 +OpEntryPoint Fragment %282 "foo_frag" %281 +OpEntryPoint GLCompute %300 "assign_through_ptr" +OpEntryPoint Vertex %314 "foo" %311 %313 +OpExecutionMode %282 OriginUpperLeft +OpExecutionMode %300 LocalSize 1 1 1 OpMemberName %6 0 "a" OpMemberName %6 1 "b" OpMemberName %6 2 "c" @@ -47,14 +48,19 @@ OpName %200 "p" OpName %201 "assign_through_ptr_fn" OpName %206 "foo" OpName %207 "assign_array_through_ptr_fn" -OpName %214 "vi" -OpName %219 "foo_vert" -OpName %231 "foo" -OpName %232 "c2" -OpName %273 "foo_frag" -OpName %291 "assign_through_ptr" -OpName %296 "val" -OpName %297 "arr" +OpName %214 "a" +OpName %215 "i" +OpName %216 "array_by_value" +OpName %224 "vi" +OpName %229 "foo_vert" +OpName %241 "foo" +OpName %242 "c2" +OpName %282 "foo_frag" +OpName %300 "assign_through_ptr" +OpName %305 "val" +OpName %306 "arr" +OpName %311 "vi" +OpName %314 "foo" OpMemberDecorate %6 0 Offset 0 OpMemberDecorate %6 1 Offset 16 OpMemberDecorate %6 2 Offset 28 @@ -99,9 +105,11 @@ OpDecorate %50 DescriptorSet 0 OpDecorate %50 Binding 3 OpDecorate %51 Block OpMemberDecorate %51 0 Offset 0 -OpDecorate %214 BuiltIn VertexIndex -OpDecorate %217 BuiltIn Position -OpDecorate %272 Location 0 +OpDecorate %224 BuiltIn VertexIndex +OpDecorate %227 BuiltIn Position +OpDecorate %281 Location 0 +OpDecorate %311 BuiltIn VertexIndex +OpDecorate %313 BuiltIn Position %2 = OpTypeVoid %3 = OpTypeInt 32 0 %4 = OpTypeVector %3 3 @@ -209,44 +217,49 @@ OpDecorate %272 Location 0 %209 = OpConstantComposite %31 %59 %59 %59 %59 %210 = OpConstantComposite %31 %61 %61 %61 %61 %211 = OpConstantComposite %34 %209 %210 -%215 = OpTypePointer Input %3 -%214 = OpVariable %215 Input -%218 = OpTypePointer Output %31 -%217 = OpVariable %218 Output -%221 = OpTypePointer StorageBuffer %23 -%224 = OpConstant %8 0.0 -%225 = OpConstant %3 3 -%226 = OpConstant %5 3 -%227 = OpConstant %5 4 -%228 = OpConstant %5 5 -%229 = OpConstant %5 42 -%230 = OpConstantNull %29 -%233 = OpTypePointer Function %32 -%234 = OpConstantNull %32 -%239 = OpTypePointer StorageBuffer %9 -%242 = OpTypePointer StorageBuffer %18 -%243 = OpConstant %3 4 -%246 = OpTypePointer StorageBuffer %10 -%247 = OpTypePointer StorageBuffer %8 -%250 = OpTypePointer StorageBuffer %19 -%253 = OpTypePointer StorageBuffer %7 -%254 = OpTypePointer StorageBuffer %5 -%266 = OpTypeVector %5 4 -%272 = OpVariable %218 Output -%275 = OpConstantComposite %10 %224 %224 %224 -%276 = OpConstantComposite %10 %59 %59 %59 -%277 = OpConstantComposite %10 %61 %61 %61 -%278 = OpConstantComposite %10 %63 %63 %63 -%279 = OpConstantComposite %9 %275 %276 %277 %278 -%280 = OpConstantComposite %17 %36 %36 -%281 = OpConstantComposite %17 %99 %99 -%282 = OpConstantComposite %18 %280 %281 -%283 = OpConstantNull %23 -%284 = OpConstantComposite %31 %224 %224 %224 %224 -%292 = OpConstant %3 33 -%293 = OpConstantComposite %31 %67 %67 %67 %67 -%294 = OpConstantComposite %31 %137 %137 %137 %137 -%295 = OpConstantComposite %34 %293 %294 +%217 = OpTypeFunction %5 %32 %5 +%219 = OpTypePointer Function %32 +%225 = OpTypePointer Input %3 +%224 = OpVariable %225 Input +%228 = OpTypePointer Output %31 +%227 = OpVariable %228 Output +%231 = OpTypePointer StorageBuffer %23 +%234 = OpConstant %8 0.0 +%235 = OpConstant %3 3 +%236 = OpConstant %5 3 +%237 = OpConstant %5 4 +%238 = OpConstant %5 5 +%239 = OpConstant %5 42 +%240 = OpConstantNull %29 +%243 = OpConstantNull %32 +%248 = OpTypePointer StorageBuffer %9 +%251 = OpTypePointer StorageBuffer %18 +%252 = OpConstant %3 4 +%255 = OpTypePointer StorageBuffer %10 +%256 = OpTypePointer StorageBuffer %8 +%259 = OpTypePointer StorageBuffer %19 +%262 = OpTypePointer StorageBuffer %7 +%263 = OpTypePointer StorageBuffer %5 +%275 = OpTypeVector %5 4 +%281 = OpVariable %228 Output +%284 = OpConstantComposite %10 %234 %234 %234 +%285 = OpConstantComposite %10 %59 %59 %59 +%286 = OpConstantComposite %10 %61 %61 %61 +%287 = OpConstantComposite %10 %63 %63 %63 +%288 = OpConstantComposite %9 %284 %285 %286 %287 +%289 = OpConstantComposite %17 %36 %36 +%290 = OpConstantComposite %17 %99 %99 +%291 = OpConstantComposite %18 %289 %290 +%292 = OpConstantNull %23 +%293 = OpConstantComposite %31 %234 %234 %234 %234 +%301 = OpConstant %3 33 +%302 = OpConstantComposite %31 %67 %67 %67 %67 +%303 = OpConstantComposite %31 %137 %137 %137 %137 +%304 = OpConstantComposite %34 %302 %303 +%311 = OpVariable %225 Input +%313 = OpVariable %228 Output +%315 = OpConstant %5 2 +%316 = OpConstantComposite %32 %58 %315 %236 %237 %238 %54 = OpFunction %2 None %55 %53 = OpLabel %82 = OpVariable %83 Function %58 @@ -390,72 +403,98 @@ OpBranch %212 OpStore %206 %211 OpReturn OpFunctionEnd -%219 = OpFunction %2 None %55 +%216 = OpFunction %5 None %217 +%214 = OpFunctionParameter %32 +%215 = OpFunctionParameter %5 %213 = OpLabel -%231 = OpVariable %27 Function %224 -%232 = OpVariable %233 Function %234 -%216 = OpLoad %3 %214 -%220 = OpAccessChain %56 %44 %36 -%222 = OpAccessChain %221 %47 %36 -%223 = OpAccessChain %131 %50 %36 -OpBranch %235 -%235 = OpLabel -%236 = OpLoad %8 %231 -OpStore %231 %59 -%237 = OpFunctionCall %2 %54 -%238 = OpFunctionCall %2 %130 -%240 = OpAccessChain %239 %42 %36 -%241 = OpLoad %9 %240 -%244 = OpAccessChain %242 %42 %243 -%245 = OpLoad %18 %244 -%248 = OpAccessChain %247 %42 %36 %225 %36 -%249 = OpLoad %8 %248 -%251 = OpArrayLength %3 %42 5 -%252 = OpISub %3 %251 %14 -%255 = OpAccessChain %254 %42 %30 %252 %36 -%256 = OpLoad %5 %255 -%257 = OpLoad %23 %222 -%258 = OpFunctionCall %8 %188 %231 -%259 = OpConvertFToS %5 %249 -%260 = OpCompositeConstruct %32 %256 %259 %226 %227 %228 -OpStore %232 %260 -%261 = OpIAdd %3 %216 %99 -%262 = OpAccessChain %83 %232 %261 -OpStore %262 %229 -%263 = OpAccessChain %83 %232 %216 -%264 = OpLoad %5 %263 -%265 = OpFunctionCall %8 %194 %230 -%267 = OpCompositeConstruct %266 %264 %264 %264 %264 -%268 = OpConvertSToF %31 %267 -%269 = OpMatrixTimesVector %10 %241 %268 -%270 = OpCompositeConstruct %31 %269 %61 -OpStore %217 %270 +%220 = OpVariable %219 Function +OpBranch %218 +%218 = OpLabel +OpStore %220 %214 +%221 = OpAccessChain %83 %220 %215 +%222 = OpLoad %5 %221 +OpReturnValue %222 +OpFunctionEnd +%229 = OpFunction %2 None %55 +%223 = OpLabel +%241 = OpVariable %27 Function %234 +%242 = OpVariable %219 Function %243 +%226 = OpLoad %3 %224 +%230 = OpAccessChain %56 %44 %36 +%232 = OpAccessChain %231 %47 %36 +%233 = OpAccessChain %131 %50 %36 +OpBranch %244 +%244 = OpLabel +%245 = OpLoad %8 %241 +OpStore %241 %59 +%246 = OpFunctionCall %2 %54 +%247 = OpFunctionCall %2 %130 +%249 = OpAccessChain %248 %42 %36 +%250 = OpLoad %9 %249 +%253 = OpAccessChain %251 %42 %252 +%254 = OpLoad %18 %253 +%257 = OpAccessChain %256 %42 %36 %235 %36 +%258 = OpLoad %8 %257 +%260 = OpArrayLength %3 %42 5 +%261 = OpISub %3 %260 %14 +%264 = OpAccessChain %263 %42 %30 %261 %36 +%265 = OpLoad %5 %264 +%266 = OpLoad %23 %232 +%267 = OpFunctionCall %8 %188 %241 +%268 = OpConvertFToS %5 %258 +%269 = OpCompositeConstruct %32 %265 %268 %236 %237 %238 +OpStore %242 %269 +%270 = OpIAdd %3 %226 %99 +%271 = OpAccessChain %83 %242 %270 +OpStore %271 %239 +%272 = OpAccessChain %83 %242 %226 +%273 = OpLoad %5 %272 +%274 = OpFunctionCall %8 %194 %240 +%276 = OpCompositeConstruct %275 %273 %273 %273 %273 +%277 = OpConvertSToF %31 %276 +%278 = OpMatrixTimesVector %10 %250 %277 +%279 = OpCompositeConstruct %31 %278 %61 +OpStore %227 %279 OpReturn OpFunctionEnd -%273 = OpFunction %2 None %55 -%271 = OpLabel -%274 = OpAccessChain %221 %47 %36 -OpBranch %285 -%285 = OpLabel -%286 = OpAccessChain %247 %42 %36 %99 %14 -OpStore %286 %59 -%287 = OpAccessChain %239 %42 %36 -OpStore %287 %279 -%288 = OpAccessChain %242 %42 %243 -OpStore %288 %282 -%289 = OpAccessChain %254 %42 %30 %99 %36 -OpStore %289 %58 -OpStore %274 %283 -OpStore %272 %284 +%282 = OpFunction %2 None %55 +%280 = OpLabel +%283 = OpAccessChain %231 %47 %36 +OpBranch %294 +%294 = OpLabel +%295 = OpAccessChain %256 %42 %36 %99 %14 +OpStore %295 %59 +%296 = OpAccessChain %248 %42 %36 +OpStore %296 %288 +%297 = OpAccessChain %251 %42 %252 +OpStore %297 %291 +%298 = OpAccessChain %263 %42 %30 %99 %36 +OpStore %298 %58 +OpStore %283 %292 +OpStore %281 %293 OpReturn OpFunctionEnd -%291 = OpFunction %2 None %55 -%290 = OpLabel -%296 = OpVariable %33 Function %292 -%297 = OpVariable %35 Function %295 -OpBranch %298 -%298 = OpLabel -%299 = OpFunctionCall %2 %201 %296 -%300 = OpFunctionCall %2 %207 %297 +%300 = OpFunction %2 None %55 +%299 = OpLabel +%305 = OpVariable %33 Function %301 +%306 = OpVariable %35 Function %304 +OpBranch %307 +%307 = OpLabel +%308 = OpFunctionCall %2 %201 %305 +%309 = OpFunctionCall %2 %207 %306 +OpReturn +OpFunctionEnd +%314 = OpFunction %2 None %55 +%310 = OpLabel +%318 = OpVariable %219 Function +%312 = OpLoad %3 %311 +OpBranch %317 +%317 = OpLabel +OpStore %318 %316 +%319 = OpAccessChain %83 %318 %312 +%320 = OpLoad %5 %319 +%321 = OpCompositeConstruct %275 %320 %320 %320 %320 +%322 = OpConvertSToF %31 %321 +OpStore %313 %322 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/wgsl/access.wgsl b/naga/tests/out/wgsl/access.wgsl index 1409e80b1..0c29e5d60 100644 --- a/naga/tests/out/wgsl/access.wgsl +++ b/naga/tests/out/wgsl/access.wgsl @@ -107,8 +107,8 @@ fn test_matrix_within_array_within_struct_accesses() { return; } -fn read_from_private(foo_1: ptr) -> f32 { - let _e1 = (*foo_1); +fn read_from_private(foo_2: ptr) -> f32 { + let _e1 = (*foo_2); return _e1; } @@ -121,11 +121,15 @@ fn assign_through_ptr_fn(p: ptr) { return; } -fn assign_array_through_ptr_fn(foo_2: ptr, 2>>) { - (*foo_2) = array, 2>(vec4(1f), vec4(2f)); +fn assign_array_through_ptr_fn(foo_3: ptr, 2>>) { + (*foo_3) = array, 2>(vec4(1f), vec4(2f)); return; } +fn array_by_value(a_1: array, i: i32) -> i32 { + return a_1[i]; +} + @vertex fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { var foo: f32 = 0f; @@ -138,11 +142,11 @@ fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { let _matrix = bar._matrix; let arr_1 = bar.arr; let b = bar._matrix[3u][0]; - let a_1 = bar.data[(arrayLength((&bar.data)) - 2u)].value; + let a_2 = bar.data[(arrayLength((&bar.data)) - 2u)].value; let c = qux; let data_pointer = (&bar.data[0].value); let _e33 = read_from_private((&foo)); - c2_ = array(a_1, i32(b), 3i, 4i, 5i); + c2_ = array(a_2, i32(b), 3i, 4i, 5i); c2_[(vi + 1u)] = 42i; let value = c2_[vi]; let _e47 = test_arr_as_arg(array, 5>()); @@ -168,3 +172,10 @@ fn assign_through_ptr() { assign_array_through_ptr_fn((&arr)); return; } + +@vertex +fn foo_1(@builtin(vertex_index) vi_1: u32) -> @builtin(position) vec4 { + const arr_2 = array(1i, 2i, 3i, 4i, 5i); + let value_1 = arr_2[vi_1]; + return vec4(vec4(value_1)); +} diff --git a/naga/tests/wgsl_errors.rs b/naga/tests/wgsl_errors.rs index 2d91ba01c..e5fb77644 100644 --- a/naga/tests/wgsl_errors.rs +++ b/naga/tests/wgsl_errors.rs @@ -1359,11 +1359,6 @@ fn missing_bindings2() { #[test] fn invalid_access() { check_validation! { - " - fn array_by_value(a: array, i: i32) -> i32 { - return a[i]; - } - ", " fn matrix_by_value(m: mat4x4, i: i32) -> vec4 { return m[i];