nordic-dev.net/wgpu
nordic-dev.net/wgpu
2
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2024-11-25 00:03:29 +00:00
Code Issues Packages Projects Releases Wiki Activity

[wgsl-in, spv-out] Allow dynamic indexing of arrays by value.

Browse Source 
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 <kvarkus@gmail.com>
Co-authored-by: Jim Blandy <jimb@red-bean.com>
This commit is contained in:
sagudev 2024-09-01 16:12:06 +02:00 committed by Jim Blandy
parent 71b4f361e2
commit 2d82054ae4
21 changed files with 837 additions and 173 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

28
naga/src/back/spv/block.rs
View File

@ -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",
));
}
}

1
naga/src/back/spv/mod.rs
View File

@ -144,6 +144,7 @@ struct Function {
signature: Option<Instruction>,
parameters: Vec<FunctionArgument>,
variables: crate::FastHashMap<Handle<crate::LocalVariable>, LocalVariable>,
internal_variables: Vec<LocalVariable>,
blocks: Vec<TerminatedBlock>,
entry_point_context: Option<EntryPointContext>,
}

53
naga/src/back/spv/writer.rs
View File

@ -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<crate::Type>,
result_type_id: Word,
container_id: Word,
container_ty: Handle<crate::Type>,
index_id: Word,
element_ty: Handle<crate::Type>,
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

21
naga/src/lib.rs
View File

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

8
naga/src/proc/mod.rs
View File

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

7
naga/src/proc/typifier.rs
View File

@ -92,6 +92,13 @@ pub enum TypeResolution {
/// available in the associated arena. However, the `TypeInner` itself may
/// contain `Handle<Type>` 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),
}

7
naga/src/valid/expression.rs
View File

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

11
naga/tests/in/access.wgsl
View File

@ -167,3 +167,14 @@ fn assign_through_ptr() {
var arr = array<vec4<f32>, 2>(vec4(6.0), vec4(7.0));
assign_array_through_ptr_fn(&arr);
}
@vertex
fn foo(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
let arr = array<i32, 5>(1, 2, 3, 4, 5);
let value = arr[vi];
return vec4<f32>(vec4<i32>(value));
}
fn array_by_value(a: array<i32, 5>, i: i32) -> i32 {
return a[i];
}

186
naga/tests/out/analysis/access.info.ron
View File

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

12
naga/tests/out/glsl/access.assign_through_ptr.Compute.glsl
View File

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

52
naga/tests/out/glsl/access.foo.Vertex.glsl Normal file
View File

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

12
naga/tests/out/glsl/access.foo_frag.Fragment.glsl
View File

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

16
naga/tests/out/glsl/access.foo_vert.Vertex.glsl
View File

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

24
naga/tests/out/hlsl/access.hlsl
View File

@ -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);
}

4
naga/tests/out/hlsl/access.ron
View File

@ -4,6 +4,10 @@
entry_point:"foo_vert",
target_profile:"vs_5_1",
),
(
entry_point:"foo_1",
target_profile:"vs_5_1",
),
],
fragment:[
(

117
naga/tests/out/ir/access.compact.ron
View File

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

117
naga/tests/out/ir/access.ron
View File

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

33
naga/tests/out/msl/access.msl
View File

@ -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<int>(b), 3, 4, 5};
c2_ = type_20 {a_2, static_cast<int>(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::float4>(metal::int4(value_1)) };
}

273
naga/tests/out/spv/access.spvasm
View File

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

23
naga/tests/out/wgsl/access.wgsl
View File

@ -107,8 +107,8 @@ fn test_matrix_within_array_within_struct_accesses() {
return;
}
fn read_from_private(foo_1: ptr<function, f32>) -> f32 {
let _e1 = (*foo_1);
fn read_from_private(foo_2: ptr<function, f32>) -> f32 {
let _e1 = (*foo_2);
return _e1;
}
@ -121,11 +121,15 @@ fn assign_through_ptr_fn(p: ptr<function, u32>) {
return;
}
fn assign_array_through_ptr_fn(foo_2: ptr<function, array<vec4<f32>, 2>>) {
(*foo_2) = array<vec4<f32>, 2>(vec4(1f), vec4(2f));
fn assign_array_through_ptr_fn(foo_3: ptr<function, array<vec4<f32>, 2>>) {
(*foo_3) = array<vec4<f32>, 2>(vec4(1f), vec4(2f));
return;
}
fn array_by_value(a_1: array<i32, 5>, i: i32) -> i32 {
return a_1[i];
}
@vertex
fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
var foo: f32 = 0f;
@ -138,11 +142,11 @@ fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
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<i32, 5>(a_1, i32(b), 3i, 4i, 5i);
c2_ = array<i32, 5>(a_2, i32(b), 3i, 4i, 5i);
c2_[(vi + 1u)] = 42i;
let value = c2_[vi];
let _e47 = test_arr_as_arg(array<array<f32, 10>, 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<f32> {
const arr_2 = array<i32, 5>(1i, 2i, 3i, 4i, 5i);
let value_1 = arr_2[vi_1];
return vec4<f32>(vec4(value_1));
}

5
naga/tests/wgsl_errors.rs
View File

@ -1359,11 +1359,6 @@ fn missing_bindings2() {
#[test]
fn invalid_access() {
check_validation! {
"
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];