From f9075fc4b84f0bf73322e582f93e8b0d07e1e919 Mon Sep 17 00:00:00 2001 From: Jim Blandy Date: Tue, 8 Oct 2024 16:32:36 -0700 Subject: [PATCH] [naga] Test access to a member/element through a pointer. --- naga/tests/in/access.wgsl | 32 + naga/tests/out/analysis/access.info.ron | 277 ++++++ .../access.assign_through_ptr.Compute.glsl | 23 + ...cess.assign_to_ptr_components.Compute.glsl | 78 ++ naga/tests/out/glsl/access.foo.Vertex.glsl | 23 + .../out/glsl/access.foo_frag.Fragment.glsl | 23 + .../out/glsl/access.foo_vert.Vertex.glsl | 23 + naga/tests/out/hlsl/access.hlsl | 41 + naga/tests/out/hlsl/access.ron | 4 + naga/tests/out/ir/access.compact.ron | 245 +++++ naga/tests/out/ir/access.ron | 245 +++++ naga/tests/out/msl/access.msl | 46 + naga/tests/out/spv/access.spvasm | 847 ++++++++++-------- naga/tests/out/wgsl/access.wgsl | 36 + 14 files changed, 1557 insertions(+), 386 deletions(-) create mode 100644 naga/tests/out/glsl/access.assign_to_ptr_components.Compute.glsl diff --git a/naga/tests/in/access.wgsl b/naga/tests/in/access.wgsl index 3336522fd..6ff68a13e 100644 --- a/naga/tests/in/access.wgsl +++ b/naga/tests/in/access.wgsl @@ -178,3 +178,35 @@ fn foo(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { fn array_by_value(a: array, i: i32) -> i32 { return a[i]; } + + +struct AssignToMember { + x: u32, +} + +fn fetch_arg_ptr_member(p: ptr) -> u32 { + return (*p).x; +} + +fn assign_to_arg_ptr_member(p: ptr) { + (*p).x = 10u; +} + +fn fetch_arg_ptr_array_element(p: ptr>) -> u32 { + return (*p)[1]; +} + +fn assign_to_arg_ptr_array_element(p: ptr>) { + (*p)[1] = 10u; +} + +@compute @workgroup_size(1) +fn assign_to_ptr_components() { + var s1: AssignToMember; + assign_to_arg_ptr_member(&s1); + fetch_arg_ptr_member(&s1); + + var a1: array; + assign_to_arg_ptr_array_element(&a1); + fetch_arg_ptr_array_element(&a1); +} diff --git a/naga/tests/out/analysis/access.info.ron b/naga/tests/out/analysis/access.info.ron index 830831cb1..8c87ce601 100644 --- a/naga/tests/out/analysis/access.info.ron +++ b/naga/tests/out/analysis/access.info.ron @@ -29,6 +29,10 @@ ("SIZED | COPY | ARGUMENT"), ("DATA | SIZED | COPY | HOST_SHAREABLE | ARGUMENT | CONSTRUCTIBLE"), ("SIZED | COPY | ARGUMENT"), + ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | ARGUMENT | CONSTRUCTIBLE"), + ("SIZED | COPY | ARGUMENT"), + ("DATA | SIZED | COPY | HOST_SHAREABLE | ARGUMENT | CONSTRUCTIBLE"), + ("SIZED | COPY | ARGUMENT"), ], functions: [ ( @@ -2783,6 +2787,216 @@ 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(30), + ), + ( + uniformity: ( + non_uniform_result: Some(0), + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Pointer( + base: 0, + space: Function, + )), + ), + ( + uniformity: ( + non_uniform_result: Some(0), + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Handle(0), + ), + ], + sampling: [], + dual_source_blending: false, + ), + ( + flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE"), + uniformity: ( + non_uniform_result: None, + requirements: (""), + ), + may_kill: false, + sampling_set: [], + global_uses: [ + (""), + (""), + (""), + (""), + (""), + ], + expressions: [ + ( + uniformity: ( + non_uniform_result: Some(0), + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Handle(30), + ), + ( + uniformity: ( + non_uniform_result: Some(0), + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Pointer( + base: 0, + space: Function, + )), + ), + ( + uniformity: ( + non_uniform_result: None, + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Scalar(( + kind: Uint, + width: 4, + ))), + ), + ], + 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(32), + ), + ( + uniformity: ( + non_uniform_result: Some(0), + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Pointer( + base: 0, + space: Function, + )), + ), + ( + uniformity: ( + non_uniform_result: Some(0), + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Handle(0), + ), + ], + sampling: [], + dual_source_blending: false, + ), + ( + flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE"), + uniformity: ( + non_uniform_result: None, + requirements: (""), + ), + may_kill: false, + sampling_set: [], + global_uses: [ + (""), + (""), + (""), + (""), + (""), + ], + expressions: [ + ( + uniformity: ( + non_uniform_result: Some(0), + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Handle(32), + ), + ( + uniformity: ( + non_uniform_result: Some(0), + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Pointer( + base: 0, + space: Function, + )), + ), + ( + uniformity: ( + non_uniform_result: None, + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Scalar(( + kind: Uint, + width: 4, + ))), + ), + ], + sampling: [], + dual_source_blending: false, + ), ], entry_points: [ ( @@ -4167,6 +4381,69 @@ 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: 2, + assignable_global: None, + ty: Value(Pointer( + base: 29, + space: Function, + )), + ), + ( + uniformity: ( + non_uniform_result: Some(0), + requirements: (""), + ), + ref_count: 0, + assignable_global: None, + ty: Handle(0), + ), + ( + uniformity: ( + non_uniform_result: Some(2), + requirements: (""), + ), + ref_count: 2, + assignable_global: None, + ty: Value(Pointer( + base: 31, + space: Function, + )), + ), + ( + uniformity: ( + non_uniform_result: Some(0), + requirements: (""), + ), + ref_count: 0, + assignable_global: None, + ty: Handle(0), + ), + ], + 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 4a4791c99..07f69f749 100644 --- a/naga/tests/out/glsl/access.assign_through_ptr.Compute.glsl +++ b/naga/tests/out/glsl/access.assign_through_ptr.Compute.glsl @@ -19,6 +19,9 @@ struct Baz { struct MatCx2InArray { mat4x2 am[2]; }; +struct AssignToMember { + uint x; +}; float read_from_private(inout float foo_2) { float _e1 = foo_2; @@ -43,6 +46,26 @@ int array_by_value(int a_1[5], int i) { return a_1[i]; } +uint fetch_arg_ptr_member(inout AssignToMember p_1) { + uint _e2 = p_1.x; + return _e2; +} + +void assign_to_arg_ptr_member(inout AssignToMember p_2) { + p_2.x = 10u; + return; +} + +uint fetch_arg_ptr_array_element(inout uint p_3[4]) { + uint _e2 = p_3[1]; + return _e2; +} + +void assign_to_arg_ptr_array_element(inout uint p_4[4]) { + p_4[1] = 10u; + return; +} + void main() { uint val = 33u; vec4 arr[2] = vec4[2](vec4(6.0), vec4(7.0)); diff --git a/naga/tests/out/glsl/access.assign_to_ptr_components.Compute.glsl b/naga/tests/out/glsl/access.assign_to_ptr_components.Compute.glsl new file mode 100644 index 000000000..1f072735e --- /dev/null +++ b/naga/tests/out/glsl/access.assign_to_ptr_components.Compute.glsl @@ -0,0 +1,78 @@ +#version 310 es + +precision highp float; +precision highp int; + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +struct GlobalConst { + uint a; + uvec3 b; + int c; +}; +struct AlignedWrapper { + int value; +}; +struct Baz { + mat3x2 m; +}; +struct MatCx2InArray { + mat4x2 am[2]; +}; +struct AssignToMember { + uint x; +}; + +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]; +} + +uint fetch_arg_ptr_member(inout AssignToMember p_1) { + uint _e2 = p_1.x; + return _e2; +} + +void assign_to_arg_ptr_member(inout AssignToMember p_2) { + p_2.x = 10u; + return; +} + +uint fetch_arg_ptr_array_element(inout uint p_3[4]) { + uint _e2 = p_3[1]; + return _e2; +} + +void assign_to_arg_ptr_array_element(inout uint p_4[4]) { + p_4[1] = 10u; + return; +} + +void main() { + AssignToMember s1_ = AssignToMember(0u); + uint a1_[4] = uint[4](0u, 0u, 0u, 0u); + assign_to_arg_ptr_member(s1_); + uint _e1 = fetch_arg_ptr_member(s1_); + assign_to_arg_ptr_array_element(a1_); + uint _e3 = fetch_arg_ptr_array_element(a1_); + return; +} + diff --git a/naga/tests/out/glsl/access.foo.Vertex.glsl b/naga/tests/out/glsl/access.foo.Vertex.glsl index e1f313840..aa2ce8b83 100644 --- a/naga/tests/out/glsl/access.foo.Vertex.glsl +++ b/naga/tests/out/glsl/access.foo.Vertex.glsl @@ -17,6 +17,9 @@ struct Baz { struct MatCx2InArray { mat4x2 am[2]; }; +struct AssignToMember { + uint x; +}; float read_from_private(inout float foo_2) { float _e1 = foo_2; @@ -41,6 +44,26 @@ int array_by_value(int a_1[5], int i) { return a_1[i]; } +uint fetch_arg_ptr_member(inout AssignToMember p_1) { + uint _e2 = p_1.x; + return _e2; +} + +void assign_to_arg_ptr_member(inout AssignToMember p_2) { + p_2.x = 10u; + return; +} + +uint fetch_arg_ptr_array_element(inout uint p_3[4]) { + uint _e2 = p_3[1]; + return _e2; +} + +void assign_to_arg_ptr_array_element(inout uint p_4[4]) { + p_4[1] = 10u; + return; +} + void main() { uint vi_1 = uint(gl_VertexID); int arr_1[5] = int[5](1, 2, 3, 4, 5); diff --git a/naga/tests/out/glsl/access.foo_frag.Fragment.glsl b/naga/tests/out/glsl/access.foo_frag.Fragment.glsl index eca6bc54c..9775b903c 100644 --- a/naga/tests/out/glsl/access.foo_frag.Fragment.glsl +++ b/naga/tests/out/glsl/access.foo_frag.Fragment.glsl @@ -17,6 +17,9 @@ struct Baz { struct MatCx2InArray { mat4x2 am[2]; }; +struct AssignToMember { + uint x; +}; layout(std430) buffer Bar_block_0Fragment { mat4x3 _matrix; mat2x2 matrix_array[2]; @@ -53,6 +56,26 @@ int array_by_value(int a_1[5], int i) { return a_1[i]; } +uint fetch_arg_ptr_member(inout AssignToMember p_1) { + uint _e2 = p_1.x; + return _e2; +} + +void assign_to_arg_ptr_member(inout AssignToMember p_2) { + p_2.x = 10u; + return; +} + +uint fetch_arg_ptr_array_element(inout uint p_3[4]) { + uint _e2 = p_3[1]; + return _e2; +} + +void assign_to_arg_ptr_array_element(inout uint p_4[4]) { + p_4[1] = 10u; + return; +} + 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 a926eadf7..d5dc00194 100644 --- a/naga/tests/out/glsl/access.foo_vert.Vertex.glsl +++ b/naga/tests/out/glsl/access.foo_vert.Vertex.glsl @@ -17,6 +17,9 @@ struct Baz { struct MatCx2InArray { mat4x2 am[2]; }; +struct AssignToMember { + uint x; +}; layout(std430) buffer Bar_block_0Vertex { mat4x3 _matrix; mat2x2 matrix_array[2]; @@ -126,6 +129,26 @@ int array_by_value(int a_1[5], int i) { return a_1[i]; } +uint fetch_arg_ptr_member(inout AssignToMember p_1) { + uint _e2 = p_1.x; + return _e2; +} + +void assign_to_arg_ptr_member(inout AssignToMember p_2) { + p_2.x = 10u; + return; +} + +uint fetch_arg_ptr_array_element(inout uint p_3[4]) { + uint _e2 = p_3[1]; + return _e2; +} + +void assign_to_arg_ptr_array_element(inout uint p_4[4]) { + p_4[1] = 10u; + return; +} + void main() { uint vi = uint(gl_VertexID); float foo = 0.0; diff --git a/naga/tests/out/hlsl/access.hlsl b/naga/tests/out/hlsl/access.hlsl index 543b84196..240b72078 100644 --- a/naga/tests/out/hlsl/access.hlsl +++ b/naga/tests/out/hlsl/access.hlsl @@ -68,6 +68,10 @@ struct MatCx2InArray { __mat4x2 am[2]; }; +struct AssignToMember { + uint x; +}; + GlobalConst ConstructGlobalConst(uint arg0, uint3 arg1, int arg2) { GlobalConst ret = (GlobalConst)0; ret.a = arg0; @@ -235,6 +239,30 @@ int array_by_value(int a_1[5], int i) return a_1[i]; } +uint fetch_arg_ptr_member(inout AssignToMember p_1) +{ + uint _e2 = p_1.x; + return _e2; +} + +void assign_to_arg_ptr_member(inout AssignToMember p_2) +{ + p_2.x = 10u; + return; +} + +uint fetch_arg_ptr_array_element(inout uint p_3[4]) +{ + uint _e2 = p_3[1]; + return _e2; +} + +void assign_to_arg_ptr_array_element(inout uint p_4[4]) +{ + p_4[1] = 10u; + return; +} + 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 }; @@ -322,3 +350,16 @@ float4 foo_1(uint vi_1 : SV_VertexID) : SV_Position int value_1 = arr_2[vi_1]; return float4((value_1).xxxx); } + +[numthreads(1, 1, 1)] +void assign_to_ptr_components() +{ + AssignToMember s1_ = (AssignToMember)0; + uint a1_[4] = (uint[4])0; + + assign_to_arg_ptr_member(s1_); + const uint _e1 = fetch_arg_ptr_member(s1_); + assign_to_arg_ptr_array_element(a1_); + const uint _e3 = fetch_arg_ptr_array_element(a1_); + return; +} diff --git a/naga/tests/out/hlsl/access.ron b/naga/tests/out/hlsl/access.ron index 8960a612e..1514e55c5 100644 --- a/naga/tests/out/hlsl/access.ron +++ b/naga/tests/out/hlsl/access.ron @@ -20,5 +20,9 @@ entry_point:"assign_through_ptr", target_profile:"cs_5_1", ), + ( + entry_point:"assign_to_ptr_components", + target_profile:"cs_5_1", + ), ], ) diff --git a/naga/tests/out/ir/access.compact.ron b/naga/tests/out/ir/access.compact.ron index 2d066b8ff..db108801c 100644 --- a/naga/tests/out/ir/access.compact.ron +++ b/naga/tests/out/ir/access.compact.ron @@ -317,6 +317,42 @@ space: Function, ), ), + ( + name: Some("AssignToMember"), + inner: Struct( + members: [ + ( + name: Some("x"), + ty: 0, + binding: None, + offset: 0, + ), + ], + span: 4, + ), + ), + ( + name: None, + inner: Pointer( + base: 29, + space: Function, + ), + ), + ( + name: None, + inner: Array( + base: 0, + size: Constant(4), + stride: 4, + ), + ), + ( + name: None, + inner: Pointer( + base: 31, + space: Function, + ), + ), ], special_types: ( ray_desc: None, @@ -1696,6 +1732,152 @@ ), ], ), + ( + name: Some("fetch_arg_ptr_member"), + arguments: [ + ( + name: Some("p"), + ty: 30, + binding: None, + ), + ], + result: Some(( + ty: 0, + binding: None, + )), + local_variables: [], + expressions: [ + FunctionArgument(0), + AccessIndex( + base: 0, + index: 0, + ), + Load( + pointer: 1, + ), + ], + named_expressions: { + 0: "p", + }, + body: [ + Emit(( + start: 1, + end: 3, + )), + Return( + value: Some(2), + ), + ], + ), + ( + name: Some("assign_to_arg_ptr_member"), + arguments: [ + ( + name: Some("p"), + ty: 30, + binding: None, + ), + ], + result: None, + local_variables: [], + expressions: [ + FunctionArgument(0), + AccessIndex( + base: 0, + index: 0, + ), + Literal(U32(10)), + ], + named_expressions: { + 0: "p", + }, + body: [ + Emit(( + start: 1, + end: 2, + )), + Store( + pointer: 1, + value: 2, + ), + Return( + value: None, + ), + ], + ), + ( + name: Some("fetch_arg_ptr_array_element"), + arguments: [ + ( + name: Some("p"), + ty: 32, + binding: None, + ), + ], + result: Some(( + ty: 0, + binding: None, + )), + local_variables: [], + expressions: [ + FunctionArgument(0), + AccessIndex( + base: 0, + index: 1, + ), + Load( + pointer: 1, + ), + ], + named_expressions: { + 0: "p", + }, + body: [ + Emit(( + start: 1, + end: 3, + )), + Return( + value: Some(2), + ), + ], + ), + ( + name: Some("assign_to_arg_ptr_array_element"), + arguments: [ + ( + name: Some("p"), + ty: 32, + binding: None, + ), + ], + result: None, + local_variables: [], + expressions: [ + FunctionArgument(0), + AccessIndex( + base: 0, + index: 1, + ), + Literal(U32(10)), + ], + named_expressions: { + 0: "p", + }, + body: [ + Emit(( + start: 1, + end: 2, + )), + Store( + pointer: 1, + value: 2, + ), + Return( + value: None, + ), + ], + ), ], entry_points: [ ( @@ -2347,5 +2529,68 @@ ], ), ), + ( + name: "assign_to_ptr_components", + stage: Compute, + early_depth_test: None, + workgroup_size: (1, 1, 1), + function: ( + name: Some("assign_to_ptr_components"), + arguments: [], + result: None, + local_variables: [ + ( + name: Some("s1"), + ty: 29, + init: None, + ), + ( + name: Some("a1"), + ty: 31, + init: None, + ), + ], + expressions: [ + LocalVariable(0), + CallResult(7), + LocalVariable(1), + CallResult(9), + ], + named_expressions: {}, + body: [ + Call( + function: 8, + arguments: [ + 0, + ], + result: None, + ), + Call( + function: 7, + arguments: [ + 0, + ], + result: Some(1), + ), + Call( + function: 10, + arguments: [ + 2, + ], + result: None, + ), + Call( + function: 9, + arguments: [ + 2, + ], + result: Some(3), + ), + Return( + value: None, + ), + ], + ), + ), ], ) \ No newline at end of file diff --git a/naga/tests/out/ir/access.ron b/naga/tests/out/ir/access.ron index 2d066b8ff..db108801c 100644 --- a/naga/tests/out/ir/access.ron +++ b/naga/tests/out/ir/access.ron @@ -317,6 +317,42 @@ space: Function, ), ), + ( + name: Some("AssignToMember"), + inner: Struct( + members: [ + ( + name: Some("x"), + ty: 0, + binding: None, + offset: 0, + ), + ], + span: 4, + ), + ), + ( + name: None, + inner: Pointer( + base: 29, + space: Function, + ), + ), + ( + name: None, + inner: Array( + base: 0, + size: Constant(4), + stride: 4, + ), + ), + ( + name: None, + inner: Pointer( + base: 31, + space: Function, + ), + ), ], special_types: ( ray_desc: None, @@ -1696,6 +1732,152 @@ ), ], ), + ( + name: Some("fetch_arg_ptr_member"), + arguments: [ + ( + name: Some("p"), + ty: 30, + binding: None, + ), + ], + result: Some(( + ty: 0, + binding: None, + )), + local_variables: [], + expressions: [ + FunctionArgument(0), + AccessIndex( + base: 0, + index: 0, + ), + Load( + pointer: 1, + ), + ], + named_expressions: { + 0: "p", + }, + body: [ + Emit(( + start: 1, + end: 3, + )), + Return( + value: Some(2), + ), + ], + ), + ( + name: Some("assign_to_arg_ptr_member"), + arguments: [ + ( + name: Some("p"), + ty: 30, + binding: None, + ), + ], + result: None, + local_variables: [], + expressions: [ + FunctionArgument(0), + AccessIndex( + base: 0, + index: 0, + ), + Literal(U32(10)), + ], + named_expressions: { + 0: "p", + }, + body: [ + Emit(( + start: 1, + end: 2, + )), + Store( + pointer: 1, + value: 2, + ), + Return( + value: None, + ), + ], + ), + ( + name: Some("fetch_arg_ptr_array_element"), + arguments: [ + ( + name: Some("p"), + ty: 32, + binding: None, + ), + ], + result: Some(( + ty: 0, + binding: None, + )), + local_variables: [], + expressions: [ + FunctionArgument(0), + AccessIndex( + base: 0, + index: 1, + ), + Load( + pointer: 1, + ), + ], + named_expressions: { + 0: "p", + }, + body: [ + Emit(( + start: 1, + end: 3, + )), + Return( + value: Some(2), + ), + ], + ), + ( + name: Some("assign_to_arg_ptr_array_element"), + arguments: [ + ( + name: Some("p"), + ty: 32, + binding: None, + ), + ], + result: None, + local_variables: [], + expressions: [ + FunctionArgument(0), + AccessIndex( + base: 0, + index: 1, + ), + Literal(U32(10)), + ], + named_expressions: { + 0: "p", + }, + body: [ + Emit(( + start: 1, + end: 2, + )), + Store( + pointer: 1, + value: 2, + ), + Return( + value: None, + ), + ], + ), ], entry_points: [ ( @@ -2347,5 +2529,68 @@ ], ), ), + ( + name: "assign_to_ptr_components", + stage: Compute, + early_depth_test: None, + workgroup_size: (1, 1, 1), + function: ( + name: Some("assign_to_ptr_components"), + arguments: [], + result: None, + local_variables: [ + ( + name: Some("s1"), + ty: 29, + init: None, + ), + ( + name: Some("a1"), + ty: 31, + init: None, + ), + ], + expressions: [ + LocalVariable(0), + CallResult(7), + LocalVariable(1), + CallResult(9), + ], + named_expressions: {}, + body: [ + Call( + function: 8, + arguments: [ + 0, + ], + result: None, + ), + Call( + function: 7, + arguments: [ + 0, + ], + result: Some(1), + ), + Call( + function: 10, + arguments: [ + 2, + ], + result: None, + ), + Call( + function: 9, + arguments: [ + 2, + ], + result: Some(3), + ), + Return( + value: None, + ), + ], + ), + ), ], ) \ No newline at end of file diff --git a/naga/tests/out/msl/access.msl b/naga/tests/out/msl/access.msl index 924b604e4..1c0fb50f4 100644 --- a/naga/tests/out/msl/access.msl +++ b/naga/tests/out/msl/access.msl @@ -57,6 +57,12 @@ struct type_20 { struct type_22 { metal::float4 inner[2]; }; +struct AssignToMember { + uint x; +}; +struct type_25 { + uint inner[4]; +}; void test_matrix_within_struct_accesses( constant Baz& baz @@ -166,6 +172,34 @@ int array_by_value( return a_1.inner[i]; } +uint fetch_arg_ptr_member( + thread AssignToMember& p_1 +) { + uint _e2 = p_1.x; + return _e2; +} + +void assign_to_arg_ptr_member( + thread AssignToMember& p_2 +) { + p_2.x = 10u; + return; +} + +uint fetch_arg_ptr_array_element( + thread type_25& p_3 +) { + uint _e2 = p_3.inner[1]; + return _e2; +} + +void assign_to_arg_ptr_array_element( + thread type_25& p_4 +) { + p_4.inner[1] = 10u; + return; +} + struct foo_vertInput { }; struct foo_vertOutput { @@ -238,3 +272,15 @@ vertex foo_1Output foo_1( int value_1 = arr_2.inner[vi_1]; return foo_1Output { static_cast(metal::int4(value_1)) }; } + + +kernel void assign_to_ptr_components( +) { + AssignToMember s1_ = {}; + type_25 a1_ = {}; + assign_to_arg_ptr_member(s1_); + uint _e1 = fetch_arg_ptr_member(s1_); + assign_to_arg_ptr_array_element(a1_); + uint _e3 = fetch_arg_ptr_array_element(a1_); + return; +} diff --git a/naga/tests/out/spv/access.spvasm b/naga/tests/out/spv/access.spvasm index 6cb87a3b1..d572baa80 100644 --- a/naga/tests/out/spv/access.spvasm +++ b/naga/tests/out/spv/access.spvasm @@ -1,17 +1,19 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 323 +; Bound: 364 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -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 +OpEntryPoint Vertex %260 "foo_vert" %255 %258 +OpEntryPoint Fragment %312 "foo_frag" %311 +OpEntryPoint GLCompute %330 "assign_through_ptr" +OpEntryPoint Vertex %344 "foo" %341 %343 +OpEntryPoint GLCompute %354 "assign_to_ptr_components" +OpExecutionMode %312 OriginUpperLeft +OpExecutionMode %330 LocalSize 1 1 1 +OpExecutionMode %354 LocalSize 1 1 1 OpMemberName %6 0 "a" OpMemberName %6 1 "b" OpMemberName %6 2 "c" @@ -29,38 +31,51 @@ OpMemberName %22 0 "m" OpName %22 "Baz" OpMemberName %26 0 "am" OpName %26 "MatCx2InArray" -OpName %40 "global_const" -OpName %42 "bar" -OpName %44 "baz" -OpName %47 "qux" -OpName %50 "nested_mat_cx2" -OpName %54 "test_matrix_within_struct_accesses" -OpName %82 "idx" -OpName %84 "t" -OpName %130 "test_matrix_within_array_within_struct_accesses" -OpName %140 "idx" -OpName %141 "t" -OpName %187 "foo" -OpName %188 "read_from_private" -OpName %193 "a" -OpName %194 "test_arr_as_arg" -OpName %200 "p" -OpName %201 "assign_through_ptr_fn" -OpName %206 "foo" -OpName %207 "assign_array_through_ptr_fn" -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" +OpMemberName %36 0 "x" +OpName %36 "AssignToMember" +OpName %45 "global_const" +OpName %47 "bar" +OpName %49 "baz" +OpName %52 "qux" +OpName %55 "nested_mat_cx2" +OpName %59 "test_matrix_within_struct_accesses" +OpName %87 "idx" +OpName %89 "t" +OpName %135 "test_matrix_within_array_within_struct_accesses" +OpName %145 "idx" +OpName %146 "t" +OpName %192 "foo" +OpName %193 "read_from_private" +OpName %198 "a" +OpName %199 "test_arr_as_arg" +OpName %205 "p" +OpName %206 "assign_through_ptr_fn" +OpName %211 "foo" +OpName %212 "assign_array_through_ptr_fn" +OpName %219 "a" +OpName %220 "i" +OpName %221 "array_by_value" +OpName %229 "p" +OpName %230 "fetch_arg_ptr_member" +OpName %236 "p" +OpName %237 "assign_to_arg_ptr_member" +OpName %242 "p" +OpName %243 "fetch_arg_ptr_array_element" +OpName %249 "p" +OpName %250 "assign_to_arg_ptr_array_element" +OpName %255 "vi" +OpName %260 "foo_vert" +OpName %272 "foo" +OpName %273 "c2" +OpName %312 "foo_frag" +OpName %330 "assign_through_ptr" +OpName %335 "val" +OpName %336 "arr" +OpName %341 "vi" +OpName %344 "foo" +OpName %354 "assign_to_ptr_components" +OpName %355 "s1" +OpName %357 "a1" OpMemberDecorate %6 0 Offset 0 OpMemberDecorate %6 1 Offset 16 OpMemberDecorate %6 2 Offset 28 @@ -91,25 +106,27 @@ OpDecorate %28 ArrayStride 4 OpDecorate %29 ArrayStride 40 OpDecorate %32 ArrayStride 4 OpDecorate %34 ArrayStride 16 -OpDecorate %42 DescriptorSet 0 -OpDecorate %42 Binding 0 -OpDecorate %44 DescriptorSet 0 -OpDecorate %44 Binding 1 -OpDecorate %45 Block -OpMemberDecorate %45 0 Offset 0 +OpMemberDecorate %36 0 Offset 0 +OpDecorate %38 ArrayStride 4 OpDecorate %47 DescriptorSet 0 -OpDecorate %47 Binding 2 -OpDecorate %48 Block -OpMemberDecorate %48 0 Offset 0 -OpDecorate %50 DescriptorSet 0 -OpDecorate %50 Binding 3 -OpDecorate %51 Block -OpMemberDecorate %51 0 Offset 0 -OpDecorate %224 BuiltIn VertexIndex -OpDecorate %227 BuiltIn Position -OpDecorate %281 Location 0 -OpDecorate %311 BuiltIn VertexIndex -OpDecorate %313 BuiltIn Position +OpDecorate %47 Binding 0 +OpDecorate %49 DescriptorSet 0 +OpDecorate %49 Binding 1 +OpDecorate %50 Block +OpMemberDecorate %50 0 Offset 0 +OpDecorate %52 DescriptorSet 0 +OpDecorate %52 Binding 2 +OpDecorate %53 Block +OpMemberDecorate %53 0 Offset 0 +OpDecorate %55 DescriptorSet 0 +OpDecorate %55 Binding 3 +OpDecorate %56 Block +OpMemberDecorate %56 0 Offset 0 +OpDecorate %255 BuiltIn VertexIndex +OpDecorate %258 BuiltIn Position +OpDecorate %311 Location 0 +OpDecorate %341 BuiltIn VertexIndex +OpDecorate %343 BuiltIn Position %2 = OpTypeVoid %3 = OpTypeInt 32 0 %4 = OpTypeVector %3 3 @@ -144,357 +161,415 @@ OpDecorate %313 BuiltIn Position %33 = OpTypePointer Function %3 %34 = OpTypeArray %31 %14 %35 = OpTypePointer Function %34 -%36 = OpConstant %3 0 -%37 = OpConstantComposite %4 %36 %36 %36 -%38 = OpConstant %5 0 -%39 = OpConstantComposite %6 %36 %37 %38 -%41 = OpTypePointer Private %6 -%40 = OpVariable %41 Private %39 -%43 = OpTypePointer StorageBuffer %20 -%42 = OpVariable %43 StorageBuffer -%45 = OpTypeStruct %22 -%46 = OpTypePointer Uniform %45 -%44 = OpVariable %46 Uniform -%48 = OpTypeStruct %23 -%49 = OpTypePointer StorageBuffer %48 -%47 = OpVariable %49 StorageBuffer -%51 = OpTypeStruct %26 -%52 = OpTypePointer Uniform %51 -%50 = OpVariable %52 Uniform -%55 = OpTypeFunction %2 -%56 = OpTypePointer Uniform %22 -%58 = OpConstant %5 1 -%59 = OpConstant %8 1.0 -%60 = OpConstantComposite %12 %59 %59 -%61 = OpConstant %8 2.0 -%62 = OpConstantComposite %12 %61 %61 -%63 = OpConstant %8 3.0 -%64 = OpConstantComposite %12 %63 %63 -%65 = OpConstantComposite %21 %60 %62 %64 -%66 = OpConstantComposite %22 %65 -%67 = OpConstant %8 6.0 -%68 = OpConstantComposite %12 %67 %67 -%69 = OpConstant %8 5.0 -%70 = OpConstantComposite %12 %69 %69 -%71 = OpConstant %8 4.0 -%72 = OpConstantComposite %12 %71 %71 -%73 = OpConstantComposite %21 %68 %70 %72 -%74 = OpConstant %8 9.0 +%36 = OpTypeStruct %3 +%37 = OpTypePointer Function %36 +%39 = OpConstant %3 4 +%38 = OpTypeArray %3 %39 +%40 = OpTypePointer Function %38 +%41 = OpConstant %3 0 +%42 = OpConstantComposite %4 %41 %41 %41 +%43 = OpConstant %5 0 +%44 = OpConstantComposite %6 %41 %42 %43 +%46 = OpTypePointer Private %6 +%45 = OpVariable %46 Private %44 +%48 = OpTypePointer StorageBuffer %20 +%47 = OpVariable %48 StorageBuffer +%50 = OpTypeStruct %22 +%51 = OpTypePointer Uniform %50 +%49 = OpVariable %51 Uniform +%53 = OpTypeStruct %23 +%54 = OpTypePointer StorageBuffer %53 +%52 = OpVariable %54 StorageBuffer +%56 = OpTypeStruct %26 +%57 = OpTypePointer Uniform %56 +%55 = OpVariable %57 Uniform +%60 = OpTypeFunction %2 +%61 = OpTypePointer Uniform %22 +%63 = OpConstant %5 1 +%64 = OpConstant %8 1.0 +%65 = OpConstantComposite %12 %64 %64 +%66 = OpConstant %8 2.0 +%67 = OpConstantComposite %12 %66 %66 +%68 = OpConstant %8 3.0 +%69 = OpConstantComposite %12 %68 %68 +%70 = OpConstantComposite %21 %65 %67 %69 +%71 = OpConstantComposite %22 %70 +%72 = OpConstant %8 6.0 +%73 = OpConstantComposite %12 %72 %72 +%74 = OpConstant %8 5.0 %75 = OpConstantComposite %12 %74 %74 -%76 = OpConstant %8 90.0 +%76 = OpConstant %8 4.0 %77 = OpConstantComposite %12 %76 %76 -%78 = OpConstant %8 10.0 -%79 = OpConstant %8 20.0 -%80 = OpConstant %8 30.0 -%81 = OpConstant %8 40.0 -%83 = OpTypePointer Function %5 -%85 = OpTypePointer Function %22 -%89 = OpTypePointer Uniform %21 -%92 = OpTypePointer Uniform %12 -%98 = OpTypePointer Uniform %8 -%99 = OpConstant %3 1 -%114 = OpTypePointer Function %21 -%116 = OpTypePointer Function %12 -%120 = OpTypePointer Function %8 -%131 = OpTypePointer Uniform %26 -%133 = OpConstantNull %25 -%134 = OpConstantComposite %26 %133 -%135 = OpConstant %8 8.0 -%136 = OpConstantComposite %12 %135 %135 -%137 = OpConstant %8 7.0 -%138 = OpConstantComposite %12 %137 %137 -%139 = OpConstantComposite %24 %136 %138 %68 %70 -%142 = OpTypePointer Function %26 -%146 = OpTypePointer Uniform %25 -%149 = OpTypePointer Uniform %24 -%171 = OpTypePointer Function %25 -%173 = OpTypePointer Function %24 -%189 = OpTypeFunction %8 %27 -%195 = OpTypeFunction %8 %29 -%202 = OpTypeFunction %2 %33 -%203 = OpConstant %3 42 -%208 = OpTypeFunction %2 %35 -%209 = OpConstantComposite %31 %59 %59 %59 %59 -%210 = OpConstantComposite %31 %61 %61 %61 %61 -%211 = OpConstantComposite %34 %209 %210 -%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 -%84 = OpVariable %85 Function %66 -%57 = OpAccessChain %56 %44 %36 -OpBranch %86 -%86 = OpLabel -%87 = OpLoad %5 %82 -%88 = OpISub %5 %87 %58 -OpStore %82 %88 -%90 = OpAccessChain %89 %57 %36 -%91 = OpLoad %21 %90 -%93 = OpAccessChain %92 %57 %36 %36 -%94 = OpLoad %12 %93 -%95 = OpLoad %5 %82 -%96 = OpAccessChain %92 %57 %36 %95 -%97 = OpLoad %12 %96 -%100 = OpAccessChain %98 %57 %36 %36 %99 -%101 = OpLoad %8 %100 -%102 = OpLoad %5 %82 -%103 = OpAccessChain %98 %57 %36 %36 %102 -%104 = OpLoad %8 %103 -%105 = OpLoad %5 %82 -%106 = OpAccessChain %98 %57 %36 %105 %99 -%107 = OpLoad %8 %106 -%108 = OpLoad %5 %82 -%109 = OpLoad %5 %82 -%110 = OpAccessChain %98 %57 %36 %108 %109 -%111 = OpLoad %8 %110 -%112 = OpLoad %5 %82 -%113 = OpIAdd %5 %112 %58 -OpStore %82 %113 -%115 = OpAccessChain %114 %84 %36 -OpStore %115 %73 -%117 = OpAccessChain %116 %84 %36 %36 -OpStore %117 %75 -%118 = OpLoad %5 %82 -%119 = OpAccessChain %116 %84 %36 %118 -OpStore %119 %77 -%121 = OpAccessChain %120 %84 %36 %36 %99 -OpStore %121 %78 -%122 = OpLoad %5 %82 -%123 = OpAccessChain %120 %84 %36 %36 %122 -OpStore %123 %79 -%124 = OpLoad %5 %82 -%125 = OpAccessChain %120 %84 %36 %124 %99 -OpStore %125 %80 -%126 = OpLoad %5 %82 -%127 = OpLoad %5 %82 -%128 = OpAccessChain %120 %84 %36 %126 %127 -OpStore %128 %81 +%78 = OpConstantComposite %21 %73 %75 %77 +%79 = OpConstant %8 9.0 +%80 = OpConstantComposite %12 %79 %79 +%81 = OpConstant %8 90.0 +%82 = OpConstantComposite %12 %81 %81 +%83 = OpConstant %8 10.0 +%84 = OpConstant %8 20.0 +%85 = OpConstant %8 30.0 +%86 = OpConstant %8 40.0 +%88 = OpTypePointer Function %5 +%90 = OpTypePointer Function %22 +%94 = OpTypePointer Uniform %21 +%97 = OpTypePointer Uniform %12 +%103 = OpTypePointer Uniform %8 +%104 = OpConstant %3 1 +%119 = OpTypePointer Function %21 +%121 = OpTypePointer Function %12 +%125 = OpTypePointer Function %8 +%136 = OpTypePointer Uniform %26 +%138 = OpConstantNull %25 +%139 = OpConstantComposite %26 %138 +%140 = OpConstant %8 8.0 +%141 = OpConstantComposite %12 %140 %140 +%142 = OpConstant %8 7.0 +%143 = OpConstantComposite %12 %142 %142 +%144 = OpConstantComposite %24 %141 %143 %73 %75 +%147 = OpTypePointer Function %26 +%151 = OpTypePointer Uniform %25 +%154 = OpTypePointer Uniform %24 +%176 = OpTypePointer Function %25 +%178 = OpTypePointer Function %24 +%194 = OpTypeFunction %8 %27 +%200 = OpTypeFunction %8 %29 +%207 = OpTypeFunction %2 %33 +%208 = OpConstant %3 42 +%213 = OpTypeFunction %2 %35 +%214 = OpConstantComposite %31 %64 %64 %64 %64 +%215 = OpConstantComposite %31 %66 %66 %66 %66 +%216 = OpConstantComposite %34 %214 %215 +%222 = OpTypeFunction %5 %32 %5 +%224 = OpTypePointer Function %32 +%231 = OpTypeFunction %3 %37 +%238 = OpTypeFunction %2 %37 +%244 = OpTypeFunction %3 %40 +%251 = OpTypeFunction %2 %40 +%256 = OpTypePointer Input %3 +%255 = OpVariable %256 Input +%259 = OpTypePointer Output %31 +%258 = OpVariable %259 Output +%262 = OpTypePointer StorageBuffer %23 +%265 = OpConstant %8 0.0 +%266 = OpConstant %3 3 +%267 = OpConstant %5 3 +%268 = OpConstant %5 4 +%269 = OpConstant %5 5 +%270 = OpConstant %5 42 +%271 = OpConstantNull %29 +%274 = OpConstantNull %32 +%279 = OpTypePointer StorageBuffer %9 +%282 = OpTypePointer StorageBuffer %18 +%285 = OpTypePointer StorageBuffer %10 +%286 = OpTypePointer StorageBuffer %8 +%289 = OpTypePointer StorageBuffer %19 +%292 = OpTypePointer StorageBuffer %7 +%293 = OpTypePointer StorageBuffer %5 +%305 = OpTypeVector %5 4 +%311 = OpVariable %259 Output +%314 = OpConstantComposite %10 %265 %265 %265 +%315 = OpConstantComposite %10 %64 %64 %64 +%316 = OpConstantComposite %10 %66 %66 %66 +%317 = OpConstantComposite %10 %68 %68 %68 +%318 = OpConstantComposite %9 %314 %315 %316 %317 +%319 = OpConstantComposite %17 %41 %41 +%320 = OpConstantComposite %17 %104 %104 +%321 = OpConstantComposite %18 %319 %320 +%322 = OpConstantNull %23 +%323 = OpConstantComposite %31 %265 %265 %265 %265 +%331 = OpConstant %3 33 +%332 = OpConstantComposite %31 %72 %72 %72 %72 +%333 = OpConstantComposite %31 %142 %142 %142 %142 +%334 = OpConstantComposite %34 %332 %333 +%341 = OpVariable %256 Input +%343 = OpVariable %259 Output +%345 = OpConstant %5 2 +%346 = OpConstantComposite %32 %63 %345 %267 %268 %269 +%356 = OpConstantNull %36 +%358 = OpConstantNull %38 +%59 = OpFunction %2 None %60 +%58 = OpLabel +%87 = OpVariable %88 Function %63 +%89 = OpVariable %90 Function %71 +%62 = OpAccessChain %61 %49 %41 +OpBranch %91 +%91 = OpLabel +%92 = OpLoad %5 %87 +%93 = OpISub %5 %92 %63 +OpStore %87 %93 +%95 = OpAccessChain %94 %62 %41 +%96 = OpLoad %21 %95 +%98 = OpAccessChain %97 %62 %41 %41 +%99 = OpLoad %12 %98 +%100 = OpLoad %5 %87 +%101 = OpAccessChain %97 %62 %41 %100 +%102 = OpLoad %12 %101 +%105 = OpAccessChain %103 %62 %41 %41 %104 +%106 = OpLoad %8 %105 +%107 = OpLoad %5 %87 +%108 = OpAccessChain %103 %62 %41 %41 %107 +%109 = OpLoad %8 %108 +%110 = OpLoad %5 %87 +%111 = OpAccessChain %103 %62 %41 %110 %104 +%112 = OpLoad %8 %111 +%113 = OpLoad %5 %87 +%114 = OpLoad %5 %87 +%115 = OpAccessChain %103 %62 %41 %113 %114 +%116 = OpLoad %8 %115 +%117 = OpLoad %5 %87 +%118 = OpIAdd %5 %117 %63 +OpStore %87 %118 +%120 = OpAccessChain %119 %89 %41 +OpStore %120 %78 +%122 = OpAccessChain %121 %89 %41 %41 +OpStore %122 %80 +%123 = OpLoad %5 %87 +%124 = OpAccessChain %121 %89 %41 %123 +OpStore %124 %82 +%126 = OpAccessChain %125 %89 %41 %41 %104 +OpStore %126 %83 +%127 = OpLoad %5 %87 +%128 = OpAccessChain %125 %89 %41 %41 %127 +OpStore %128 %84 +%129 = OpLoad %5 %87 +%130 = OpAccessChain %125 %89 %41 %129 %104 +OpStore %130 %85 +%131 = OpLoad %5 %87 +%132 = OpLoad %5 %87 +%133 = OpAccessChain %125 %89 %41 %131 %132 +OpStore %133 %86 OpReturn OpFunctionEnd -%130 = OpFunction %2 None %55 -%129 = OpLabel -%140 = OpVariable %83 Function %58 -%141 = OpVariable %142 Function %134 -%132 = OpAccessChain %131 %50 %36 -OpBranch %143 -%143 = OpLabel -%144 = OpLoad %5 %140 -%145 = OpISub %5 %144 %58 -OpStore %140 %145 -%147 = OpAccessChain %146 %132 %36 -%148 = OpLoad %25 %147 -%150 = OpAccessChain %149 %132 %36 %36 -%151 = OpLoad %24 %150 -%152 = OpAccessChain %92 %132 %36 %36 %36 -%153 = OpLoad %12 %152 -%154 = OpLoad %5 %140 -%155 = OpAccessChain %92 %132 %36 %36 %154 -%156 = OpLoad %12 %155 -%157 = OpAccessChain %98 %132 %36 %36 %36 %99 -%158 = OpLoad %8 %157 -%159 = OpLoad %5 %140 -%160 = OpAccessChain %98 %132 %36 %36 %36 %159 -%161 = OpLoad %8 %160 -%162 = OpLoad %5 %140 -%163 = OpAccessChain %98 %132 %36 %36 %162 %99 -%164 = OpLoad %8 %163 -%165 = OpLoad %5 %140 -%166 = OpLoad %5 %140 -%167 = OpAccessChain %98 %132 %36 %36 %165 %166 -%168 = OpLoad %8 %167 -%169 = OpLoad %5 %140 -%170 = OpIAdd %5 %169 %58 -OpStore %140 %170 -%172 = OpAccessChain %171 %141 %36 -OpStore %172 %133 -%174 = OpAccessChain %173 %141 %36 %36 -OpStore %174 %139 -%175 = OpAccessChain %116 %141 %36 %36 %36 -OpStore %175 %75 -%176 = OpLoad %5 %140 -%177 = OpAccessChain %116 %141 %36 %36 %176 -OpStore %177 %77 -%178 = OpAccessChain %120 %141 %36 %36 %36 %99 -OpStore %178 %78 -%179 = OpLoad %5 %140 -%180 = OpAccessChain %120 %141 %36 %36 %36 %179 -OpStore %180 %79 -%181 = OpLoad %5 %140 -%182 = OpAccessChain %120 %141 %36 %36 %181 %99 -OpStore %182 %80 -%183 = OpLoad %5 %140 -%184 = OpLoad %5 %140 -%185 = OpAccessChain %120 %141 %36 %36 %183 %184 -OpStore %185 %81 +%135 = OpFunction %2 None %60 +%134 = OpLabel +%145 = OpVariable %88 Function %63 +%146 = OpVariable %147 Function %139 +%137 = OpAccessChain %136 %55 %41 +OpBranch %148 +%148 = OpLabel +%149 = OpLoad %5 %145 +%150 = OpISub %5 %149 %63 +OpStore %145 %150 +%152 = OpAccessChain %151 %137 %41 +%153 = OpLoad %25 %152 +%155 = OpAccessChain %154 %137 %41 %41 +%156 = OpLoad %24 %155 +%157 = OpAccessChain %97 %137 %41 %41 %41 +%158 = OpLoad %12 %157 +%159 = OpLoad %5 %145 +%160 = OpAccessChain %97 %137 %41 %41 %159 +%161 = OpLoad %12 %160 +%162 = OpAccessChain %103 %137 %41 %41 %41 %104 +%163 = OpLoad %8 %162 +%164 = OpLoad %5 %145 +%165 = OpAccessChain %103 %137 %41 %41 %41 %164 +%166 = OpLoad %8 %165 +%167 = OpLoad %5 %145 +%168 = OpAccessChain %103 %137 %41 %41 %167 %104 +%169 = OpLoad %8 %168 +%170 = OpLoad %5 %145 +%171 = OpLoad %5 %145 +%172 = OpAccessChain %103 %137 %41 %41 %170 %171 +%173 = OpLoad %8 %172 +%174 = OpLoad %5 %145 +%175 = OpIAdd %5 %174 %63 +OpStore %145 %175 +%177 = OpAccessChain %176 %146 %41 +OpStore %177 %138 +%179 = OpAccessChain %178 %146 %41 %41 +OpStore %179 %144 +%180 = OpAccessChain %121 %146 %41 %41 %41 +OpStore %180 %80 +%181 = OpLoad %5 %145 +%182 = OpAccessChain %121 %146 %41 %41 %181 +OpStore %182 %82 +%183 = OpAccessChain %125 %146 %41 %41 %41 %104 +OpStore %183 %83 +%184 = OpLoad %5 %145 +%185 = OpAccessChain %125 %146 %41 %41 %41 %184 +OpStore %185 %84 +%186 = OpLoad %5 %145 +%187 = OpAccessChain %125 %146 %41 %41 %186 %104 +OpStore %187 %85 +%188 = OpLoad %5 %145 +%189 = OpLoad %5 %145 +%190 = OpAccessChain %125 %146 %41 %41 %188 %189 +OpStore %190 %86 OpReturn OpFunctionEnd -%188 = OpFunction %8 None %189 -%187 = OpFunctionParameter %27 -%186 = OpLabel -OpBranch %190 -%190 = OpLabel -%191 = OpLoad %8 %187 -OpReturnValue %191 +%193 = OpFunction %8 None %194 +%192 = OpFunctionParameter %27 +%191 = OpLabel +OpBranch %195 +%195 = OpLabel +%196 = OpLoad %8 %192 +OpReturnValue %196 OpFunctionEnd -%194 = OpFunction %8 None %195 -%193 = OpFunctionParameter %29 -%192 = OpLabel -OpBranch %196 -%196 = OpLabel -%197 = OpCompositeExtract %28 %193 4 -%198 = OpCompositeExtract %8 %197 9 -OpReturnValue %198 +%199 = OpFunction %8 None %200 +%198 = OpFunctionParameter %29 +%197 = OpLabel +OpBranch %201 +%201 = OpLabel +%202 = OpCompositeExtract %28 %198 4 +%203 = OpCompositeExtract %8 %202 9 +OpReturnValue %203 OpFunctionEnd -%201 = OpFunction %2 None %202 -%200 = OpFunctionParameter %33 -%199 = OpLabel -OpBranch %204 +%206 = OpFunction %2 None %207 +%205 = OpFunctionParameter %33 %204 = OpLabel -OpStore %200 %203 +OpBranch %209 +%209 = OpLabel +OpStore %205 %208 OpReturn OpFunctionEnd -%207 = OpFunction %2 None %208 -%206 = OpFunctionParameter %35 -%205 = OpLabel -OpBranch %212 -%212 = OpLabel -OpStore %206 %211 +%212 = OpFunction %2 None %213 +%211 = OpFunctionParameter %35 +%210 = OpLabel +OpBranch %217 +%217 = OpLabel +OpStore %211 %216 OpReturn OpFunctionEnd -%216 = OpFunction %5 None %217 -%214 = OpFunctionParameter %32 -%215 = OpFunctionParameter %5 -%213 = OpLabel -%220 = OpVariable %219 Function -OpBranch %218 +%221 = OpFunction %5 None %222 +%219 = OpFunctionParameter %32 +%220 = OpFunctionParameter %5 %218 = OpLabel -OpStore %220 %214 -%221 = OpAccessChain %83 %220 %215 -%222 = OpLoad %5 %221 -OpReturnValue %222 -OpFunctionEnd -%229 = OpFunction %2 None %55 +%225 = OpVariable %224 Function +OpBranch %223 %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 +OpStore %225 %219 +%226 = OpAccessChain %88 %225 %220 +%227 = OpLoad %5 %226 +OpReturnValue %227 +OpFunctionEnd +%230 = OpFunction %3 None %231 +%229 = OpFunctionParameter %37 +%228 = OpLabel +OpBranch %232 +%232 = OpLabel +%233 = OpAccessChain %33 %229 %41 +%234 = OpLoad %3 %233 +OpReturnValue %234 +OpFunctionEnd +%237 = OpFunction %2 None %238 +%236 = OpFunctionParameter %37 +%235 = OpLabel +OpBranch %239 +%239 = OpLabel +%240 = OpAccessChain %33 %236 %41 +OpStore %240 %16 OpReturn OpFunctionEnd -%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 +%243 = OpFunction %3 None %244 +%242 = OpFunctionParameter %40 +%241 = OpLabel +OpBranch %245 +%245 = OpLabel +%246 = OpAccessChain %33 %242 %104 +%247 = OpLoad %3 %246 +OpReturnValue %247 +OpFunctionEnd +%250 = OpFunction %2 None %251 +%249 = OpFunctionParameter %40 +%248 = OpLabel +OpBranch %252 +%252 = OpLabel +%253 = OpAccessChain %33 %249 %104 +OpStore %253 %16 OpReturn OpFunctionEnd -%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 +%260 = OpFunction %2 None %60 +%254 = OpLabel +%272 = OpVariable %27 Function %265 +%273 = OpVariable %224 Function %274 +%257 = OpLoad %3 %255 +%261 = OpAccessChain %61 %49 %41 +%263 = OpAccessChain %262 %52 %41 +%264 = OpAccessChain %136 %55 %41 +OpBranch %275 +%275 = OpLabel +%276 = OpLoad %8 %272 +OpStore %272 %64 +%277 = OpFunctionCall %2 %59 +%278 = OpFunctionCall %2 %135 +%280 = OpAccessChain %279 %47 %41 +%281 = OpLoad %9 %280 +%283 = OpAccessChain %282 %47 %39 +%284 = OpLoad %18 %283 +%287 = OpAccessChain %286 %47 %41 %266 %41 +%288 = OpLoad %8 %287 +%290 = OpArrayLength %3 %47 5 +%291 = OpISub %3 %290 %14 +%294 = OpAccessChain %293 %47 %30 %291 %41 +%295 = OpLoad %5 %294 +%296 = OpLoad %23 %263 +%297 = OpFunctionCall %8 %193 %272 +%298 = OpConvertFToS %5 %288 +%299 = OpCompositeConstruct %32 %295 %298 %267 %268 %269 +OpStore %273 %299 +%300 = OpIAdd %3 %257 %104 +%301 = OpAccessChain %88 %273 %300 +OpStore %301 %270 +%302 = OpAccessChain %88 %273 %257 +%303 = OpLoad %5 %302 +%304 = OpFunctionCall %8 %199 %271 +%306 = OpCompositeConstruct %305 %303 %303 %303 %303 +%307 = OpConvertSToF %31 %306 +%308 = OpMatrixTimesVector %10 %281 %307 +%309 = OpCompositeConstruct %31 %308 %66 +OpStore %258 %309 OpReturn OpFunctionEnd -%314 = OpFunction %2 None %55 +%312 = OpFunction %2 None %60 %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 +%313 = OpAccessChain %262 %52 %41 +OpBranch %324 +%324 = OpLabel +%325 = OpAccessChain %286 %47 %41 %104 %14 +OpStore %325 %64 +%326 = OpAccessChain %279 %47 %41 +OpStore %326 %318 +%327 = OpAccessChain %282 %47 %39 +OpStore %327 %321 +%328 = OpAccessChain %293 %47 %30 %104 %41 +OpStore %328 %63 OpStore %313 %322 +OpStore %311 %323 +OpReturn +OpFunctionEnd +%330 = OpFunction %2 None %60 +%329 = OpLabel +%335 = OpVariable %33 Function %331 +%336 = OpVariable %35 Function %334 +OpBranch %337 +%337 = OpLabel +%338 = OpFunctionCall %2 %206 %335 +%339 = OpFunctionCall %2 %212 %336 +OpReturn +OpFunctionEnd +%344 = OpFunction %2 None %60 +%340 = OpLabel +%348 = OpVariable %224 Function +%342 = OpLoad %3 %341 +OpBranch %347 +%347 = OpLabel +OpStore %348 %346 +%349 = OpAccessChain %88 %348 %342 +%350 = OpLoad %5 %349 +%351 = OpCompositeConstruct %305 %350 %350 %350 %350 +%352 = OpConvertSToF %31 %351 +OpStore %343 %352 +OpReturn +OpFunctionEnd +%354 = OpFunction %2 None %60 +%353 = OpLabel +%355 = OpVariable %37 Function %356 +%357 = OpVariable %40 Function %358 +OpBranch %359 +%359 = OpLabel +%360 = OpFunctionCall %2 %237 %355 +%361 = OpFunctionCall %3 %230 %355 +%362 = OpFunctionCall %2 %250 %357 +%363 = OpFunctionCall %3 %243 %357 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 0c29e5d60..c659b2126 100644 --- a/naga/tests/out/wgsl/access.wgsl +++ b/naga/tests/out/wgsl/access.wgsl @@ -25,6 +25,10 @@ struct MatCx2InArray { am: array, 2>, } +struct AssignToMember { + x: u32, +} + var global_const: GlobalConst = GlobalConst(0u, vec3(0u, 0u, 0u), 0i); @group(0) @binding(0) var bar: Bar; @@ -130,6 +134,26 @@ fn array_by_value(a_1: array, i: i32) -> i32 { return a_1[i]; } +fn fetch_arg_ptr_member(p_1: ptr) -> u32 { + let _e2 = (*p_1).x; + return _e2; +} + +fn assign_to_arg_ptr_member(p_2: ptr) { + (*p_2).x = 10u; + return; +} + +fn fetch_arg_ptr_array_element(p_3: ptr>) -> u32 { + let _e2 = (*p_3)[1]; + return _e2; +} + +fn assign_to_arg_ptr_array_element(p_4: ptr>) { + (*p_4)[1] = 10u; + return; +} + @vertex fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { var foo: f32 = 0f; @@ -179,3 +203,15 @@ fn foo_1(@builtin(vertex_index) vi_1: u32) -> @builtin(position) vec4 { let value_1 = arr_2[vi_1]; return vec4(vec4(value_1)); } + +@compute @workgroup_size(1, 1, 1) +fn assign_to_ptr_components() { + var s1_: AssignToMember; + var a1_: array; + + assign_to_arg_ptr_member((&s1_)); + let _e1 = fetch_arg_ptr_member((&s1_)); + assign_to_arg_ptr_array_element((&a1_)); + let _e3 = fetch_arg_ptr_array_element((&a1_)); + return; +}