From 58105a06e2bd5aefeb9330984d47976e63c11dc4 Mon Sep 17 00:00:00 2001 From: Leah Date: Fri, 17 Feb 2023 08:40:21 +0800 Subject: [PATCH] [glsl/hlsl-out] Write sizes of arrays behind pointers in function arguments (#2250) arrays can be put behind pointers in inout and out parameters in GLSL and HLSL, whose dimensions must be specified to let array access compile. so, we specify their dimensions. fixes #2248 --- src/back/glsl/mod.rs | 12 +- src/back/hlsl/writer.rs | 2 +- tests/in/access.wgsl | 9 +- tests/out/analysis/access.info.ron | 208 +++++ .../access.assign_through_ptr.Compute.glsl | 8 + tests/out/glsl/access.atomics.Compute.glsl | 5 + tests/out/glsl/access.foo_frag.Fragment.glsl | 5 + tests/out/glsl/access.foo_vert.Vertex.glsl | 7 +- tests/out/hlsl/access.hlsl | 18 +- tests/out/ir/access.ron | 114 ++- tests/out/msl/access.msl | 15 +- tests/out/spv/access.spvasm | 799 +++++++++--------- tests/out/wgsl/access.wgsl | 11 +- 13 files changed, 816 insertions(+), 397 deletions(-) diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index addcbadbcd..3433c2f1ca 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -1455,8 +1455,16 @@ impl<'a, W: Write> Writer<'a, W> { write!(this.out, " {}", &this.names[&ctx.argument_key(i as u32)])?; // Write array size - if let TypeInner::Array { base, size, .. } = this.module.types[arg.ty].inner { - this.write_array_size(base, size)?; + match this.module.types[arg.ty].inner { + TypeInner::Array { base, size, .. } => { + this.write_array_size(base, size)?; + } + TypeInner::Pointer { base, .. } => { + if let TypeInner::Array { base, size, .. } = this.module.types[base].inner { + this.write_array_size(base, size)?; + } + } + _ => {} } Ok(()) diff --git a/src/back/hlsl/writer.rs b/src/back/hlsl/writer.rs index 2af7a3524b..c075fe3fe2 100644 --- a/src/back/hlsl/writer.rs +++ b/src/back/hlsl/writer.rs @@ -1143,7 +1143,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { // Write argument name. Space is important. write!(self.out, " {argument_name}")?; - if let TypeInner::Array { base, size, .. } = module.types[arg.ty].inner { + if let TypeInner::Array { base, size, .. } = module.types[arg_ty].inner { self.write_array_size(module, base, size)?; } } diff --git a/tests/in/access.wgsl b/tests/in/access.wgsl index 473f54213e..a9c27ee78c 100644 --- a/tests/in/access.wgsl +++ b/tests/in/access.wgsl @@ -174,7 +174,14 @@ fn assign_through_ptr_fn(p: ptr) { *p = 42u; } +fn assign_array_through_ptr_fn(foo: ptr, 2>>) { + *foo = array, 2>(vec4(1.0), vec4(2.0)); +} + @compute @workgroup_size(1) fn assign_through_ptr() { + var arr = array, 2>(vec4(6.0), vec4(7.0)); + assign_through_ptr_fn(&val); -} + assign_array_through_ptr_fn(&arr); +} \ No newline at end of file diff --git a/tests/out/analysis/access.info.ron b/tests/out/analysis/access.info.ron index 75f53126b8..01e08a63ba 100644 --- a/tests/out/analysis/access.info.ron +++ b/tests/out/analysis/access.info.ron @@ -93,6 +93,12 @@ ( bits: 70, ), + ( + bits: 215, + ), + ( + bits: 70, + ), ], functions: [ ( @@ -3588,6 +3594,125 @@ ], sampling: [], ), + ( + flags: ( + bits: 63, + ), + available_stages: ( + bits: 7, + ), + uniformity: ( + non_uniform_result: None, + requirements: ( + bits: 0, + ), + ), + may_kill: false, + sampling_set: [], + global_uses: [ + ( + bits: 0, + ), + ( + bits: 0, + ), + ( + bits: 0, + ), + ( + bits: 0, + ), + ( + bits: 0, + ), + ( + bits: 0, + ), + ], + expressions: [ + ( + uniformity: ( + non_uniform_result: Some(1), + requirements: ( + bits: 0, + ), + ), + ref_count: 1, + assignable_global: None, + ty: Handle(33), + ), + ( + uniformity: ( + non_uniform_result: None, + requirements: ( + bits: 0, + ), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Scalar( + kind: Float, + width: 4, + )), + ), + ( + uniformity: ( + non_uniform_result: None, + requirements: ( + bits: 0, + ), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Vector( + size: Quad, + kind: Float, + width: 4, + )), + ), + ( + uniformity: ( + non_uniform_result: None, + requirements: ( + bits: 0, + ), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Scalar( + kind: Float, + width: 4, + )), + ), + ( + uniformity: ( + non_uniform_result: None, + requirements: ( + bits: 0, + ), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Vector( + size: Quad, + kind: Float, + width: 4, + )), + ), + ( + uniformity: ( + non_uniform_result: None, + requirements: ( + bits: 0, + ), + ), + ref_count: 1, + assignable_global: None, + ty: Handle(32), + ), + ], + sampling: [], + ), ], entry_points: [ ( @@ -5661,6 +5786,89 @@ ), ], expressions: [ + ( + uniformity: ( + non_uniform_result: None, + requirements: ( + bits: 0, + ), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Scalar( + kind: Float, + width: 4, + )), + ), + ( + uniformity: ( + non_uniform_result: None, + requirements: ( + bits: 0, + ), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Vector( + size: Quad, + kind: Float, + width: 4, + )), + ), + ( + uniformity: ( + non_uniform_result: None, + requirements: ( + bits: 0, + ), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Scalar( + kind: Float, + width: 4, + )), + ), + ( + uniformity: ( + non_uniform_result: None, + requirements: ( + bits: 0, + ), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Vector( + size: Quad, + kind: Float, + width: 4, + )), + ), + ( + uniformity: ( + non_uniform_result: None, + requirements: ( + bits: 0, + ), + ), + ref_count: 1, + assignable_global: None, + ty: Handle(32), + ), + ( + uniformity: ( + non_uniform_result: Some(6), + requirements: ( + bits: 0, + ), + ), + ref_count: 2, + assignable_global: None, + ty: Value(Pointer( + base: 32, + space: Function, + )), + ), ( uniformity: ( non_uniform_result: None, diff --git a/tests/out/glsl/access.assign_through_ptr.Compute.glsl b/tests/out/glsl/access.assign_through_ptr.Compute.glsl index 1306fe514d..f79b30e3e9 100644 --- a/tests/out/glsl/access.assign_through_ptr.Compute.glsl +++ b/tests/out/glsl/access.assign_through_ptr.Compute.glsl @@ -36,13 +36,21 @@ 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)); + return; +} + void main() { if (gl_GlobalInvocationID == uvec3(0u)) { val = 0u; } memoryBarrierShared(); barrier(); + vec4 arr[2] = vec4[2](vec4(0.0), vec4(0.0)); + arr = vec4[2](vec4(6.0), vec4(7.0)); assign_through_ptr_fn(val); + assign_array_through_ptr_fn(arr); return; } diff --git a/tests/out/glsl/access.atomics.Compute.glsl b/tests/out/glsl/access.atomics.Compute.glsl index ab164cf6f5..8ccefe422b 100644 --- a/tests/out/glsl/access.atomics.Compute.glsl +++ b/tests/out/glsl/access.atomics.Compute.glsl @@ -43,6 +43,11 @@ 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)); + return; +} + void main() { int tmp = 0; int value = _group_0_binding_0_cs.atom; diff --git a/tests/out/glsl/access.foo_frag.Fragment.glsl b/tests/out/glsl/access.foo_frag.Fragment.glsl index e7656b8c1c..c71617572b 100644 --- a/tests/out/glsl/access.foo_frag.Fragment.glsl +++ b/tests/out/glsl/access.foo_frag.Fragment.glsl @@ -44,6 +44,11 @@ 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)); + 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/tests/out/glsl/access.foo_vert.Vertex.glsl b/tests/out/glsl/access.foo_vert.Vertex.glsl index 331b7cda9b..03e33ccb81 100644 --- a/tests/out/glsl/access.foo_vert.Vertex.glsl +++ b/tests/out/glsl/access.foo_vert.Vertex.glsl @@ -121,6 +121,11 @@ 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)); + return; +} + void main() { uint vi = uint(gl_VertexID); float foo = 0.0; @@ -131,7 +136,7 @@ void main() { test_matrix_within_struct_accesses(); test_matrix_within_array_within_struct_accesses(); mat4x3 _matrix = _group_0_binding_0_vs._matrix; - uvec2 arr[2] = _group_0_binding_0_vs.arr; + uvec2 arr_1[2] = _group_0_binding_0_vs.arr; float b = _group_0_binding_0_vs._matrix[3][0]; int a_1 = _group_0_binding_0_vs.data[(uint(_group_0_binding_0_vs.data.length()) - 2u)].value; ivec2 c = _group_0_binding_2_vs; diff --git a/tests/out/hlsl/access.hlsl b/tests/out/hlsl/access.hlsl index 9f898b0e9a..c87f39cf27 100644 --- a/tests/out/hlsl/access.hlsl +++ b/tests/out/hlsl/access.hlsl @@ -237,6 +237,18 @@ void assign_through_ptr_fn(inout uint p) return; } +typedef float4 ret_Constructarray2_float4_[2]; +ret_Constructarray2_float4_ Constructarray2_float4_(float4 arg0, float4 arg1) { + float4 ret[2] = { arg0, arg1 }; + return ret; +} + +void assign_array_through_ptr_fn(inout float4 foo_2[2]) +{ + foo_2 = Constructarray2_float4_((1.0).xxxx, (2.0).xxxx); + return; +} + uint NagaBufferLengthRW(RWByteAddressBuffer buffer) { uint ret; @@ -261,7 +273,7 @@ float4 foo_vert(uint vi : SV_VertexID) : SV_Position test_matrix_within_struct_accesses(); test_matrix_within_array_within_struct_accesses(); 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[2] = {asuint(bar.Load2(144+0)), asuint(bar.Load2(144+8))}; + uint2 arr_1[2] = {asuint(bar.Load2(144+0)), asuint(bar.Load2(144+8))}; float b = asfloat(bar.Load(0+48+0)); int a_1 = asint(bar.Load(0+(((NagaBufferLengthRW(bar) - 160) / 8) - 2u)*8+160)); int2 c = asint(qux.Load2(0)); @@ -332,6 +344,10 @@ void assign_through_ptr(uint3 __global_invocation_id : SV_DispatchThreadID) val = (uint)0; } GroupMemoryBarrierWithGroupSync(); + float4 arr[2] = (float4[2])0; + + arr = Constructarray2_float4_((6.0).xxxx, (7.0).xxxx); assign_through_ptr_fn(val); + assign_array_through_ptr_fn(arr); return; } diff --git a/tests/out/ir/access.ron b/tests/out/ir/access.ron index 3c37839986..e544ee1a5d 100644 --- a/tests/out/ir/access.ron +++ b/tests/out/ir/access.ron @@ -317,6 +317,21 @@ space: WorkGroup, ), ), + ( + name: None, + inner: Array( + base: 26, + size: Constant(5), + stride: 16, + ), + ), + ( + name: None, + inner: Pointer( + base: 32, + space: Function, + ), + ), ], constants: [ ( @@ -1964,6 +1979,58 @@ ), ], ), + ( + name: Some("assign_array_through_ptr_fn"), + arguments: [ + ( + name: Some("foo"), + ty: 33, + binding: None, + ), + ], + result: None, + local_variables: [], + expressions: [ + FunctionArgument(0), + Constant(8), + Splat( + size: Quad, + value: 2, + ), + Constant(9), + Splat( + size: Quad, + value: 4, + ), + Compose( + ty: 32, + components: [ + 3, + 5, + ], + ), + ], + named_expressions: { + 1: "foo", + }, + body: [ + Emit(( + start: 2, + end: 3, + )), + Emit(( + start: 4, + end: 6, + )), + Store( + pointer: 1, + value: 6, + ), + Return( + value: None, + ), + ], + ), ], entry_points: [ ( @@ -2721,16 +2788,59 @@ name: Some("assign_through_ptr"), arguments: [], result: None, - local_variables: [], + local_variables: [ + ( + name: Some("arr"), + ty: 32, + init: None, + ), + ], expressions: [ + Constant(11), + Splat( + size: Quad, + value: 1, + ), + Constant(25), + Splat( + size: Quad, + value: 3, + ), + Compose( + ty: 32, + components: [ + 2, + 4, + ], + ), + LocalVariable(1), GlobalVariable(6), ], named_expressions: {}, body: [ + Emit(( + start: 1, + end: 2, + )), + Emit(( + start: 3, + end: 5, + )), + Store( + pointer: 6, + value: 5, + ), Call( function: 5, arguments: [ - 1, + 7, + ], + result: None, + ), + Call( + function: 6, + arguments: [ + 6, ], result: None, ), diff --git a/tests/out/msl/access.msl b/tests/out/msl/access.msl index 30b2bdbc61..e565643bb6 100644 --- a/tests/out/msl/access.msl +++ b/tests/out/msl/access.msl @@ -54,6 +54,9 @@ struct type_19 { struct type_22 { int inner[5]; }; +struct type_26 { + metal::float4 inner[2]; +}; constant metal::uint3 const_type_1_ = {0u, 0u, 0u}; constant GlobalConst const_GlobalConst = {0u, {}, const_type_1_, 0}; constant metal::float2 const_type_13_ = {0.0, 0.0}; @@ -161,6 +164,13 @@ void assign_through_ptr_fn( return; } +void assign_array_through_ptr_fn( + thread type_26& foo_2 +) { + for(int _i=0; _i<2; ++_i) foo_2.inner[_i] = type_26 {metal::float4(1.0), metal::float4(2.0)}.inner[_i]; + return; +} + struct foo_vertInput { }; struct foo_vertOutput { @@ -182,7 +192,7 @@ vertex foo_vertOutput foo_vert( test_matrix_within_struct_accesses(baz); test_matrix_within_array_within_struct_accesses(nested_mat_cx2_); metal::float4x3 _matrix = bar._matrix; - type_9 arr = bar.arr; + type_9 arr_1 = bar.arr; float b = bar._matrix[3].x; int a_1 = bar.data[(1 + (_buffer_sizes.size1 - 160 - 8) / 8) - 2u].value; metal::int2 c = qux; @@ -247,6 +257,9 @@ kernel void assign_through_ptr( val = {}; } metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); + type_26 arr = {}; + for(int _i=0; _i<2; ++_i) arr.inner[_i] = type_26 {metal::float4(6.0), metal::float4(7.0)}.inner[_i]; assign_through_ptr_fn(val); + assign_array_through_ptr_fn(arr); return; } diff --git a/tests/out/spv/access.spvasm b/tests/out/spv/access.spvasm index 128e0158b0..42ab59eb50 100644 --- a/tests/out/spv/access.spvasm +++ b/tests/out/spv/access.spvasm @@ -1,18 +1,18 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 338 +; Bound: 354 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Vertex %237 "foo_vert" %232 %235 -OpEntryPoint Fragment %279 "foo_frag" %278 -OpEntryPoint GLCompute %298 "atomics" -OpEntryPoint GLCompute %322 "assign_through_ptr" %325 -OpExecutionMode %279 OriginUpperLeft -OpExecutionMode %298 LocalSize 1 1 1 -OpExecutionMode %322 LocalSize 1 1 1 +OpEntryPoint Vertex %247 "foo_vert" %242 %245 +OpEntryPoint Fragment %289 "foo_frag" %288 +OpEntryPoint GLCompute %308 "atomics" +OpEntryPoint GLCompute %334 "assign_through_ptr" %337 +OpExecutionMode %289 OriginUpperLeft +OpExecutionMode %308 LocalSize 1 1 1 +OpExecutionMode %334 LocalSize 1 1 1 OpSource GLSL 450 OpMemberName %36 0 "a" OpMemberName %36 1 "b" @@ -31,32 +31,35 @@ OpMemberName %49 0 "m" OpName %49 "Baz" OpMemberName %53 0 "am" OpName %53 "MatCx2InArray" -OpName %70 "global_const" -OpName %72 "bar" -OpName %74 "baz" -OpName %77 "qux" -OpName %80 "nested_mat_cx2" -OpName %83 "val" -OpName %84 "idx" -OpName %87 "t" -OpName %91 "test_matrix_within_struct_accesses" -OpName %148 "idx" -OpName %150 "t" -OpName %154 "test_matrix_within_array_within_struct_accesses" -OpName %209 "foo" -OpName %210 "read_from_private" -OpName %215 "a" -OpName %216 "test_arr_as_arg" -OpName %222 "p" -OpName %223 "assign_through_ptr_fn" -OpName %226 "foo" -OpName %228 "c2" -OpName %232 "vi" -OpName %237 "foo_vert" -OpName %279 "foo_frag" -OpName %295 "tmp" -OpName %298 "atomics" -OpName %322 "assign_through_ptr" +OpName %72 "global_const" +OpName %74 "bar" +OpName %76 "baz" +OpName %79 "qux" +OpName %82 "nested_mat_cx2" +OpName %85 "val" +OpName %86 "idx" +OpName %89 "t" +OpName %93 "test_matrix_within_struct_accesses" +OpName %150 "idx" +OpName %152 "t" +OpName %156 "test_matrix_within_array_within_struct_accesses" +OpName %211 "foo" +OpName %212 "read_from_private" +OpName %217 "a" +OpName %218 "test_arr_as_arg" +OpName %224 "p" +OpName %225 "assign_through_ptr_fn" +OpName %229 "foo" +OpName %230 "assign_array_through_ptr_fn" +OpName %236 "foo" +OpName %238 "c2" +OpName %242 "vi" +OpName %247 "foo_vert" +OpName %289 "foo_frag" +OpName %305 "tmp" +OpName %308 "atomics" +OpName %331 "arr" +OpName %334 "assign_through_ptr" OpMemberDecorate %36 0 Offset 0 OpMemberDecorate %36 1 Offset 16 OpMemberDecorate %36 2 Offset 28 @@ -85,25 +88,26 @@ OpMemberDecorate %53 0 MatrixStride 8 OpDecorate %55 ArrayStride 4 OpDecorate %56 ArrayStride 40 OpDecorate %59 ArrayStride 4 -OpDecorate %72 DescriptorSet 0 -OpDecorate %72 Binding 0 -OpDecorate %47 Block +OpDecorate %62 ArrayStride 16 OpDecorate %74 DescriptorSet 0 -OpDecorate %74 Binding 1 -OpDecorate %75 Block -OpMemberDecorate %75 0 Offset 0 -OpDecorate %77 DescriptorSet 0 -OpDecorate %77 Binding 2 -OpDecorate %78 Block -OpMemberDecorate %78 0 Offset 0 -OpDecorate %80 DescriptorSet 0 -OpDecorate %80 Binding 3 -OpDecorate %81 Block -OpMemberDecorate %81 0 Offset 0 -OpDecorate %232 BuiltIn VertexIndex -OpDecorate %235 BuiltIn Position -OpDecorate %278 Location 0 -OpDecorate %325 BuiltIn GlobalInvocationId +OpDecorate %74 Binding 0 +OpDecorate %47 Block +OpDecorate %76 DescriptorSet 0 +OpDecorate %76 Binding 1 +OpDecorate %77 Block +OpMemberDecorate %77 0 Offset 0 +OpDecorate %79 DescriptorSet 0 +OpDecorate %79 Binding 2 +OpDecorate %80 Block +OpMemberDecorate %80 0 Offset 0 +OpDecorate %82 DescriptorSet 0 +OpDecorate %82 Binding 3 +OpDecorate %83 Block +OpMemberDecorate %83 0 Offset 0 +OpDecorate %242 BuiltIn VertexIndex +OpDecorate %245 BuiltIn Position +OpDecorate %288 Location 0 +OpDecorate %337 BuiltIn GlobalInvocationId %2 = OpTypeVoid %4 = OpTypeInt 32 0 %3 = OpConstant %4 0 @@ -164,356 +168,377 @@ OpDecorate %325 BuiltIn GlobalInvocationId %59 = OpTypeArray %6 %26 %60 = OpTypeVector %6 4 %61 = OpTypePointer Workgroup %4 -%62 = OpConstantComposite %35 %3 %3 %3 -%63 = OpConstantComposite %36 %3 %62 %5 -%64 = OpConstantComposite %41 %23 %23 -%65 = OpConstantComposite %51 %64 %64 %64 %64 -%66 = OpConstantComposite %52 %65 %65 -%67 = OpConstantComposite %55 %23 %23 %23 %23 %23 %23 %23 %23 %23 %23 -%68 = OpConstantComposite %56 %67 %67 %67 %67 %67 -%69 = OpConstantComposite %50 %5 %5 -%71 = OpTypePointer Private %36 -%70 = OpVariable %71 Private %63 -%73 = OpTypePointer StorageBuffer %47 -%72 = OpVariable %73 StorageBuffer -%75 = OpTypeStruct %49 -%76 = OpTypePointer Uniform %75 -%74 = OpVariable %76 Uniform -%78 = OpTypeStruct %50 -%79 = OpTypePointer StorageBuffer %78 -%77 = OpVariable %79 StorageBuffer -%81 = OpTypeStruct %53 -%82 = OpTypePointer Uniform %81 -%80 = OpVariable %82 Uniform -%83 = OpVariable %61 Workgroup -%85 = OpTypePointer Function %6 -%86 = OpConstantNull %6 -%88 = OpTypePointer Function %49 -%89 = OpConstantNull %49 -%92 = OpTypeFunction %2 -%93 = OpTypePointer Uniform %49 -%98 = OpTypePointer Uniform %48 -%101 = OpTypePointer Uniform %41 -%107 = OpTypePointer Uniform %11 -%127 = OpTypePointer Function %48 -%133 = OpTypePointer Function %41 -%139 = OpTypePointer Function %11 -%149 = OpConstantNull %6 -%151 = OpTypePointer Function %53 -%152 = OpConstantNull %53 -%155 = OpTypePointer Uniform %53 -%160 = OpTypePointer Uniform %52 -%163 = OpTypePointer Uniform %51 -%186 = OpTypePointer Function %52 -%188 = OpTypePointer Function %51 -%211 = OpTypeFunction %11 %54 -%217 = OpTypeFunction %11 %56 -%224 = OpTypeFunction %2 %61 -%227 = OpConstantNull %11 -%229 = OpTypePointer Function %59 -%230 = OpConstantNull %59 -%233 = OpTypePointer Input %4 -%232 = OpVariable %233 Input -%236 = OpTypePointer Output %57 -%235 = OpVariable %236 Output -%239 = OpTypePointer StorageBuffer %50 -%246 = OpTypePointer StorageBuffer %38 -%249 = OpTypePointer StorageBuffer %45 -%250 = OpConstant %4 4 -%253 = OpTypePointer StorageBuffer %39 -%254 = OpTypePointer StorageBuffer %11 -%257 = OpTypePointer StorageBuffer %46 -%260 = OpTypePointer StorageBuffer %37 -%261 = OpConstant %4 5 -%278 = OpVariable %236 Output -%296 = OpConstantNull %6 -%300 = OpTypePointer StorageBuffer %6 -%303 = OpConstant %4 64 -%324 = OpConstantNull %4 -%326 = OpTypePointer Input %35 -%325 = OpVariable %326 Input -%328 = OpConstantNull %35 -%330 = OpTypeBool -%329 = OpTypeVector %330 3 -%335 = OpConstant %4 264 -%91 = OpFunction %2 None %92 -%90 = OpLabel -%84 = OpVariable %85 Function %86 -%87 = OpVariable %88 Function %89 -%94 = OpAccessChain %93 %74 %3 -OpBranch %95 -%95 = OpLabel -OpStore %84 %9 -%96 = OpLoad %6 %84 -%97 = OpISub %6 %96 %9 -OpStore %84 %97 -%99 = OpAccessChain %98 %94 %3 -%100 = OpLoad %48 %99 -%102 = OpAccessChain %101 %94 %3 %3 -%103 = OpLoad %41 %102 -%104 = OpLoad %6 %84 -%105 = OpAccessChain %101 %94 %3 %104 -%106 = OpLoad %41 %105 -%108 = OpAccessChain %107 %94 %3 %3 %32 -%109 = OpLoad %11 %108 -%110 = OpLoad %6 %84 -%111 = OpAccessChain %107 %94 %3 %3 %110 -%112 = OpLoad %11 %111 -%113 = OpLoad %6 %84 -%114 = OpAccessChain %107 %94 %3 %113 %32 -%115 = OpLoad %11 %114 -%116 = OpLoad %6 %84 -%117 = OpLoad %6 %84 -%118 = OpAccessChain %107 %94 %3 %116 %117 -%119 = OpLoad %11 %118 -%120 = OpCompositeConstruct %41 %10 %10 -%121 = OpCompositeConstruct %41 %12 %12 -%122 = OpCompositeConstruct %41 %13 %13 -%123 = OpCompositeConstruct %48 %120 %121 %122 -%124 = OpCompositeConstruct %49 %123 -OpStore %87 %124 -%125 = OpLoad %6 %84 -%126 = OpIAdd %6 %125 %9 -OpStore %84 %126 -%128 = OpCompositeConstruct %41 %14 %14 -%129 = OpCompositeConstruct %41 %15 %15 -%130 = OpCompositeConstruct %41 %16 %16 -%131 = OpCompositeConstruct %48 %128 %129 %130 -%132 = OpAccessChain %127 %87 %3 -OpStore %132 %131 -%134 = OpCompositeConstruct %41 %17 %17 -%135 = OpAccessChain %133 %87 %3 %3 -OpStore %135 %134 -%136 = OpLoad %6 %84 -%137 = OpCompositeConstruct %41 %18 %18 -%138 = OpAccessChain %133 %87 %3 %136 -OpStore %138 %137 -%140 = OpAccessChain %139 %87 %3 %3 %32 -OpStore %140 %19 -%141 = OpLoad %6 %84 -%142 = OpAccessChain %139 %87 %3 %3 %141 -OpStore %142 %20 -%143 = OpLoad %6 %84 -%144 = OpAccessChain %139 %87 %3 %143 %32 -OpStore %144 %21 -%145 = OpLoad %6 %84 -%146 = OpLoad %6 %84 -%147 = OpAccessChain %139 %87 %3 %145 %146 -OpStore %147 %22 +%62 = OpTypeArray %57 %7 +%63 = OpTypePointer Function %62 +%64 = OpConstantComposite %35 %3 %3 %3 +%65 = OpConstantComposite %36 %3 %64 %5 +%66 = OpConstantComposite %41 %23 %23 +%67 = OpConstantComposite %51 %66 %66 %66 %66 +%68 = OpConstantComposite %52 %67 %67 +%69 = OpConstantComposite %55 %23 %23 %23 %23 %23 %23 %23 %23 %23 %23 +%70 = OpConstantComposite %56 %69 %69 %69 %69 %69 +%71 = OpConstantComposite %50 %5 %5 +%73 = OpTypePointer Private %36 +%72 = OpVariable %73 Private %65 +%75 = OpTypePointer StorageBuffer %47 +%74 = OpVariable %75 StorageBuffer +%77 = OpTypeStruct %49 +%78 = OpTypePointer Uniform %77 +%76 = OpVariable %78 Uniform +%80 = OpTypeStruct %50 +%81 = OpTypePointer StorageBuffer %80 +%79 = OpVariable %81 StorageBuffer +%83 = OpTypeStruct %53 +%84 = OpTypePointer Uniform %83 +%82 = OpVariable %84 Uniform +%85 = OpVariable %61 Workgroup +%87 = OpTypePointer Function %6 +%88 = OpConstantNull %6 +%90 = OpTypePointer Function %49 +%91 = OpConstantNull %49 +%94 = OpTypeFunction %2 +%95 = OpTypePointer Uniform %49 +%100 = OpTypePointer Uniform %48 +%103 = OpTypePointer Uniform %41 +%109 = OpTypePointer Uniform %11 +%129 = OpTypePointer Function %48 +%135 = OpTypePointer Function %41 +%141 = OpTypePointer Function %11 +%151 = OpConstantNull %6 +%153 = OpTypePointer Function %53 +%154 = OpConstantNull %53 +%157 = OpTypePointer Uniform %53 +%162 = OpTypePointer Uniform %52 +%165 = OpTypePointer Uniform %51 +%188 = OpTypePointer Function %52 +%190 = OpTypePointer Function %51 +%213 = OpTypeFunction %11 %54 +%219 = OpTypeFunction %11 %56 +%226 = OpTypeFunction %2 %61 +%231 = OpTypeFunction %2 %63 +%237 = OpConstantNull %11 +%239 = OpTypePointer Function %59 +%240 = OpConstantNull %59 +%243 = OpTypePointer Input %4 +%242 = OpVariable %243 Input +%246 = OpTypePointer Output %57 +%245 = OpVariable %246 Output +%249 = OpTypePointer StorageBuffer %50 +%256 = OpTypePointer StorageBuffer %38 +%259 = OpTypePointer StorageBuffer %45 +%260 = OpConstant %4 4 +%263 = OpTypePointer StorageBuffer %39 +%264 = OpTypePointer StorageBuffer %11 +%267 = OpTypePointer StorageBuffer %46 +%270 = OpTypePointer StorageBuffer %37 +%271 = OpConstant %4 5 +%288 = OpVariable %246 Output +%306 = OpConstantNull %6 +%310 = OpTypePointer StorageBuffer %6 +%313 = OpConstant %4 64 +%332 = OpConstantNull %62 +%336 = OpConstantNull %4 +%338 = OpTypePointer Input %35 +%337 = OpVariable %338 Input +%340 = OpConstantNull %35 +%342 = OpTypeBool +%341 = OpTypeVector %342 3 +%347 = OpConstant %4 264 +%93 = OpFunction %2 None %94 +%92 = OpLabel +%86 = OpVariable %87 Function %88 +%89 = OpVariable %90 Function %91 +%96 = OpAccessChain %95 %76 %3 +OpBranch %97 +%97 = OpLabel +OpStore %86 %9 +%98 = OpLoad %6 %86 +%99 = OpISub %6 %98 %9 +OpStore %86 %99 +%101 = OpAccessChain %100 %96 %3 +%102 = OpLoad %48 %101 +%104 = OpAccessChain %103 %96 %3 %3 +%105 = OpLoad %41 %104 +%106 = OpLoad %6 %86 +%107 = OpAccessChain %103 %96 %3 %106 +%108 = OpLoad %41 %107 +%110 = OpAccessChain %109 %96 %3 %3 %32 +%111 = OpLoad %11 %110 +%112 = OpLoad %6 %86 +%113 = OpAccessChain %109 %96 %3 %3 %112 +%114 = OpLoad %11 %113 +%115 = OpLoad %6 %86 +%116 = OpAccessChain %109 %96 %3 %115 %32 +%117 = OpLoad %11 %116 +%118 = OpLoad %6 %86 +%119 = OpLoad %6 %86 +%120 = OpAccessChain %109 %96 %3 %118 %119 +%121 = OpLoad %11 %120 +%122 = OpCompositeConstruct %41 %10 %10 +%123 = OpCompositeConstruct %41 %12 %12 +%124 = OpCompositeConstruct %41 %13 %13 +%125 = OpCompositeConstruct %48 %122 %123 %124 +%126 = OpCompositeConstruct %49 %125 +OpStore %89 %126 +%127 = OpLoad %6 %86 +%128 = OpIAdd %6 %127 %9 +OpStore %86 %128 +%130 = OpCompositeConstruct %41 %14 %14 +%131 = OpCompositeConstruct %41 %15 %15 +%132 = OpCompositeConstruct %41 %16 %16 +%133 = OpCompositeConstruct %48 %130 %131 %132 +%134 = OpAccessChain %129 %89 %3 +OpStore %134 %133 +%136 = OpCompositeConstruct %41 %17 %17 +%137 = OpAccessChain %135 %89 %3 %3 +OpStore %137 %136 +%138 = OpLoad %6 %86 +%139 = OpCompositeConstruct %41 %18 %18 +%140 = OpAccessChain %135 %89 %3 %138 +OpStore %140 %139 +%142 = OpAccessChain %141 %89 %3 %3 %32 +OpStore %142 %19 +%143 = OpLoad %6 %86 +%144 = OpAccessChain %141 %89 %3 %3 %143 +OpStore %144 %20 +%145 = OpLoad %6 %86 +%146 = OpAccessChain %141 %89 %3 %145 %32 +OpStore %146 %21 +%147 = OpLoad %6 %86 +%148 = OpLoad %6 %86 +%149 = OpAccessChain %141 %89 %3 %147 %148 +OpStore %149 %22 OpReturn OpFunctionEnd -%154 = OpFunction %2 None %92 -%153 = OpLabel -%148 = OpVariable %85 Function %149 -%150 = OpVariable %151 Function %152 -%156 = OpAccessChain %155 %80 %3 -OpBranch %157 -%157 = OpLabel -OpStore %148 %9 -%158 = OpLoad %6 %148 -%159 = OpISub %6 %158 %9 -OpStore %148 %159 -%161 = OpAccessChain %160 %156 %3 -%162 = OpLoad %52 %161 -%164 = OpAccessChain %163 %156 %3 %3 -%165 = OpLoad %51 %164 -%166 = OpAccessChain %101 %156 %3 %3 %3 -%167 = OpLoad %41 %166 -%168 = OpLoad %6 %148 -%169 = OpAccessChain %101 %156 %3 %3 %168 -%170 = OpLoad %41 %169 -%171 = OpAccessChain %107 %156 %3 %3 %3 %32 -%172 = OpLoad %11 %171 -%173 = OpLoad %6 %148 -%174 = OpAccessChain %107 %156 %3 %3 %3 %173 -%175 = OpLoad %11 %174 -%176 = OpLoad %6 %148 -%177 = OpAccessChain %107 %156 %3 %3 %176 %32 -%178 = OpLoad %11 %177 -%179 = OpLoad %6 %148 -%180 = OpLoad %6 %148 -%181 = OpAccessChain %107 %156 %3 %3 %179 %180 -%182 = OpLoad %11 %181 -%183 = OpCompositeConstruct %53 %66 -OpStore %150 %183 -%184 = OpLoad %6 %148 -%185 = OpIAdd %6 %184 %9 -OpStore %148 %185 -%187 = OpAccessChain %186 %150 %3 -OpStore %187 %66 -%189 = OpCompositeConstruct %41 %24 %24 -%190 = OpCompositeConstruct %41 %25 %25 -%191 = OpCompositeConstruct %41 %14 %14 -%192 = OpCompositeConstruct %41 %15 %15 -%193 = OpCompositeConstruct %51 %189 %190 %191 %192 -%194 = OpAccessChain %188 %150 %3 %3 -OpStore %194 %193 -%195 = OpCompositeConstruct %41 %17 %17 -%196 = OpAccessChain %133 %150 %3 %3 %3 +%156 = OpFunction %2 None %94 +%155 = OpLabel +%150 = OpVariable %87 Function %151 +%152 = OpVariable %153 Function %154 +%158 = OpAccessChain %157 %82 %3 +OpBranch %159 +%159 = OpLabel +OpStore %150 %9 +%160 = OpLoad %6 %150 +%161 = OpISub %6 %160 %9 +OpStore %150 %161 +%163 = OpAccessChain %162 %158 %3 +%164 = OpLoad %52 %163 +%166 = OpAccessChain %165 %158 %3 %3 +%167 = OpLoad %51 %166 +%168 = OpAccessChain %103 %158 %3 %3 %3 +%169 = OpLoad %41 %168 +%170 = OpLoad %6 %150 +%171 = OpAccessChain %103 %158 %3 %3 %170 +%172 = OpLoad %41 %171 +%173 = OpAccessChain %109 %158 %3 %3 %3 %32 +%174 = OpLoad %11 %173 +%175 = OpLoad %6 %150 +%176 = OpAccessChain %109 %158 %3 %3 %3 %175 +%177 = OpLoad %11 %176 +%178 = OpLoad %6 %150 +%179 = OpAccessChain %109 %158 %3 %3 %178 %32 +%180 = OpLoad %11 %179 +%181 = OpLoad %6 %150 +%182 = OpLoad %6 %150 +%183 = OpAccessChain %109 %158 %3 %3 %181 %182 +%184 = OpLoad %11 %183 +%185 = OpCompositeConstruct %53 %68 +OpStore %152 %185 +%186 = OpLoad %6 %150 +%187 = OpIAdd %6 %186 %9 +OpStore %150 %187 +%189 = OpAccessChain %188 %152 %3 +OpStore %189 %68 +%191 = OpCompositeConstruct %41 %24 %24 +%192 = OpCompositeConstruct %41 %25 %25 +%193 = OpCompositeConstruct %41 %14 %14 +%194 = OpCompositeConstruct %41 %15 %15 +%195 = OpCompositeConstruct %51 %191 %192 %193 %194 +%196 = OpAccessChain %190 %152 %3 %3 OpStore %196 %195 -%197 = OpLoad %6 %148 -%198 = OpCompositeConstruct %41 %18 %18 -%199 = OpAccessChain %133 %150 %3 %3 %197 -OpStore %199 %198 -%200 = OpAccessChain %139 %150 %3 %3 %3 %32 -OpStore %200 %19 -%201 = OpLoad %6 %148 -%202 = OpAccessChain %139 %150 %3 %3 %3 %201 -OpStore %202 %20 -%203 = OpLoad %6 %148 -%204 = OpAccessChain %139 %150 %3 %3 %203 %32 -OpStore %204 %21 -%205 = OpLoad %6 %148 -%206 = OpLoad %6 %148 -%207 = OpAccessChain %139 %150 %3 %3 %205 %206 -OpStore %207 %22 +%197 = OpCompositeConstruct %41 %17 %17 +%198 = OpAccessChain %135 %152 %3 %3 %3 +OpStore %198 %197 +%199 = OpLoad %6 %150 +%200 = OpCompositeConstruct %41 %18 %18 +%201 = OpAccessChain %135 %152 %3 %3 %199 +OpStore %201 %200 +%202 = OpAccessChain %141 %152 %3 %3 %3 %32 +OpStore %202 %19 +%203 = OpLoad %6 %150 +%204 = OpAccessChain %141 %152 %3 %3 %3 %203 +OpStore %204 %20 +%205 = OpLoad %6 %150 +%206 = OpAccessChain %141 %152 %3 %3 %205 %32 +OpStore %206 %21 +%207 = OpLoad %6 %150 +%208 = OpLoad %6 %150 +%209 = OpAccessChain %141 %152 %3 %3 %207 %208 +OpStore %209 %22 OpReturn OpFunctionEnd -%210 = OpFunction %11 None %211 -%209 = OpFunctionParameter %54 -%208 = OpLabel -OpBranch %212 -%212 = OpLabel -%213 = OpLoad %11 %209 -OpReturnValue %213 -OpFunctionEnd -%216 = OpFunction %11 None %217 -%215 = OpFunctionParameter %56 +%212 = OpFunction %11 None %213 +%211 = OpFunctionParameter %54 +%210 = OpLabel +OpBranch %214 %214 = OpLabel -OpBranch %218 -%218 = OpLabel -%219 = OpCompositeExtract %55 %215 4 -%220 = OpCompositeExtract %11 %219 9 -OpReturnValue %220 +%215 = OpLoad %11 %211 +OpReturnValue %215 OpFunctionEnd -%223 = OpFunction %2 None %224 -%222 = OpFunctionParameter %61 -%221 = OpLabel -OpBranch %225 -%225 = OpLabel -OpStore %222 %34 +%218 = OpFunction %11 None %219 +%217 = OpFunctionParameter %56 +%216 = OpLabel +OpBranch %220 +%220 = OpLabel +%221 = OpCompositeExtract %55 %217 4 +%222 = OpCompositeExtract %11 %221 9 +OpReturnValue %222 +OpFunctionEnd +%225 = OpFunction %2 None %226 +%224 = OpFunctionParameter %61 +%223 = OpLabel +OpBranch %227 +%227 = OpLabel +OpStore %224 %34 OpReturn OpFunctionEnd -%237 = OpFunction %2 None %92 -%231 = OpLabel -%226 = OpVariable %54 Function %227 -%228 = OpVariable %229 Function %230 -%234 = OpLoad %4 %232 -%238 = OpAccessChain %93 %74 %3 -%240 = OpAccessChain %239 %77 %3 -%241 = OpAccessChain %155 %80 %3 -OpBranch %242 -%242 = OpLabel -OpStore %226 %23 -%243 = OpLoad %11 %226 -OpStore %226 %10 -%244 = OpFunctionCall %2 %91 -%245 = OpFunctionCall %2 %154 -%247 = OpAccessChain %246 %72 %3 -%248 = OpLoad %38 %247 -%251 = OpAccessChain %249 %72 %250 -%252 = OpLoad %45 %251 -%255 = OpAccessChain %254 %72 %3 %29 %3 -%256 = OpLoad %11 %255 -%258 = OpArrayLength %4 %72 5 -%259 = OpISub %4 %258 %30 -%262 = OpAccessChain %58 %72 %261 %259 %3 -%263 = OpLoad %6 %262 -%264 = OpLoad %50 %240 -%265 = OpFunctionCall %11 %210 %226 -%266 = OpConvertFToS %6 %256 -%267 = OpCompositeConstruct %59 %263 %266 %31 %27 %26 -OpStore %228 %267 -%268 = OpIAdd %4 %234 %32 -%269 = OpAccessChain %85 %228 %268 -OpStore %269 %33 -%270 = OpAccessChain %85 %228 %234 -%271 = OpLoad %6 %270 -%272 = OpFunctionCall %11 %216 %68 -%273 = OpCompositeConstruct %60 %271 %271 %271 %271 -%274 = OpConvertSToF %57 %273 -%275 = OpMatrixTimesVector %39 %248 %274 -%276 = OpCompositeConstruct %57 %275 %12 -OpStore %235 %276 +%230 = OpFunction %2 None %231 +%229 = OpFunctionParameter %63 +%228 = OpLabel +OpBranch %232 +%232 = OpLabel +%233 = OpCompositeConstruct %57 %10 %10 %10 %10 +%234 = OpCompositeConstruct %57 %12 %12 %12 %12 +%235 = OpCompositeConstruct %62 %233 %234 +OpStore %229 %235 OpReturn OpFunctionEnd -%279 = OpFunction %2 None %92 -%277 = OpLabel -%280 = OpAccessChain %239 %77 %3 -OpBranch %281 -%281 = OpLabel -%282 = OpAccessChain %254 %72 %3 %32 %30 -OpStore %282 %10 -%283 = OpCompositeConstruct %39 %23 %23 %23 -%284 = OpCompositeConstruct %39 %10 %10 %10 -%285 = OpCompositeConstruct %39 %12 %12 %12 -%286 = OpCompositeConstruct %39 %13 %13 %13 -%287 = OpCompositeConstruct %38 %283 %284 %285 %286 -%288 = OpAccessChain %246 %72 %3 -OpStore %288 %287 -%289 = OpCompositeConstruct %44 %3 %3 -%290 = OpCompositeConstruct %44 %32 %32 -%291 = OpCompositeConstruct %45 %289 %290 -%292 = OpAccessChain %249 %72 %250 -OpStore %292 %291 -%293 = OpAccessChain %58 %72 %261 %32 %3 -OpStore %293 %9 -OpStore %280 %69 -%294 = OpCompositeConstruct %57 %23 %23 %23 %23 -OpStore %278 %294 +%247 = OpFunction %2 None %94 +%241 = OpLabel +%236 = OpVariable %54 Function %237 +%238 = OpVariable %239 Function %240 +%244 = OpLoad %4 %242 +%248 = OpAccessChain %95 %76 %3 +%250 = OpAccessChain %249 %79 %3 +%251 = OpAccessChain %157 %82 %3 +OpBranch %252 +%252 = OpLabel +OpStore %236 %23 +%253 = OpLoad %11 %236 +OpStore %236 %10 +%254 = OpFunctionCall %2 %93 +%255 = OpFunctionCall %2 %156 +%257 = OpAccessChain %256 %74 %3 +%258 = OpLoad %38 %257 +%261 = OpAccessChain %259 %74 %260 +%262 = OpLoad %45 %261 +%265 = OpAccessChain %264 %74 %3 %29 %3 +%266 = OpLoad %11 %265 +%268 = OpArrayLength %4 %74 5 +%269 = OpISub %4 %268 %30 +%272 = OpAccessChain %58 %74 %271 %269 %3 +%273 = OpLoad %6 %272 +%274 = OpLoad %50 %250 +%275 = OpFunctionCall %11 %212 %236 +%276 = OpConvertFToS %6 %266 +%277 = OpCompositeConstruct %59 %273 %276 %31 %27 %26 +OpStore %238 %277 +%278 = OpIAdd %4 %244 %32 +%279 = OpAccessChain %87 %238 %278 +OpStore %279 %33 +%280 = OpAccessChain %87 %238 %244 +%281 = OpLoad %6 %280 +%282 = OpFunctionCall %11 %218 %70 +%283 = OpCompositeConstruct %60 %281 %281 %281 %281 +%284 = OpConvertSToF %57 %283 +%285 = OpMatrixTimesVector %39 %258 %284 +%286 = OpCompositeConstruct %57 %285 %12 +OpStore %245 %286 OpReturn OpFunctionEnd -%298 = OpFunction %2 None %92 -%297 = OpLabel -%295 = OpVariable %85 Function %296 -OpBranch %299 -%299 = OpLabel -%301 = OpAccessChain %300 %72 %30 -%302 = OpAtomicLoad %6 %301 %9 %303 -%305 = OpAccessChain %300 %72 %30 -%304 = OpAtomicIAdd %6 %305 %9 %303 %26 -OpStore %295 %304 -%307 = OpAccessChain %300 %72 %30 -%306 = OpAtomicISub %6 %307 %9 %303 %26 -OpStore %295 %306 -%309 = OpAccessChain %300 %72 %30 -%308 = OpAtomicAnd %6 %309 %9 %303 %26 -OpStore %295 %308 -%311 = OpAccessChain %300 %72 %30 -%310 = OpAtomicOr %6 %311 %9 %303 %26 -OpStore %295 %310 -%313 = OpAccessChain %300 %72 %30 -%312 = OpAtomicXor %6 %313 %9 %303 %26 -OpStore %295 %312 -%315 = OpAccessChain %300 %72 %30 -%314 = OpAtomicSMin %6 %315 %9 %303 %26 -OpStore %295 %314 -%317 = OpAccessChain %300 %72 %30 -%316 = OpAtomicSMax %6 %317 %9 %303 %26 -OpStore %295 %316 -%319 = OpAccessChain %300 %72 %30 -%318 = OpAtomicExchange %6 %319 %9 %303 %26 -OpStore %295 %318 -%320 = OpAccessChain %300 %72 %30 -OpAtomicStore %320 %9 %303 %302 +%289 = OpFunction %2 None %94 +%287 = OpLabel +%290 = OpAccessChain %249 %79 %3 +OpBranch %291 +%291 = OpLabel +%292 = OpAccessChain %264 %74 %3 %32 %30 +OpStore %292 %10 +%293 = OpCompositeConstruct %39 %23 %23 %23 +%294 = OpCompositeConstruct %39 %10 %10 %10 +%295 = OpCompositeConstruct %39 %12 %12 %12 +%296 = OpCompositeConstruct %39 %13 %13 %13 +%297 = OpCompositeConstruct %38 %293 %294 %295 %296 +%298 = OpAccessChain %256 %74 %3 +OpStore %298 %297 +%299 = OpCompositeConstruct %44 %3 %3 +%300 = OpCompositeConstruct %44 %32 %32 +%301 = OpCompositeConstruct %45 %299 %300 +%302 = OpAccessChain %259 %74 %260 +OpStore %302 %301 +%303 = OpAccessChain %58 %74 %271 %32 %3 +OpStore %303 %9 +OpStore %290 %71 +%304 = OpCompositeConstruct %57 %23 %23 %23 %23 +OpStore %288 %304 OpReturn OpFunctionEnd -%322 = OpFunction %2 None %92 -%321 = OpLabel -OpBranch %323 -%323 = OpLabel -%327 = OpLoad %35 %325 -%331 = OpIEqual %329 %327 %328 -%332 = OpAll %330 %331 -OpSelectionMerge %333 None -OpBranchConditional %332 %334 %333 -%334 = OpLabel -OpStore %83 %324 -OpBranch %333 +%308 = OpFunction %2 None %94 +%307 = OpLabel +%305 = OpVariable %87 Function %306 +OpBranch %309 +%309 = OpLabel +%311 = OpAccessChain %310 %74 %30 +%312 = OpAtomicLoad %6 %311 %9 %313 +%315 = OpAccessChain %310 %74 %30 +%314 = OpAtomicIAdd %6 %315 %9 %313 %26 +OpStore %305 %314 +%317 = OpAccessChain %310 %74 %30 +%316 = OpAtomicISub %6 %317 %9 %313 %26 +OpStore %305 %316 +%319 = OpAccessChain %310 %74 %30 +%318 = OpAtomicAnd %6 %319 %9 %313 %26 +OpStore %305 %318 +%321 = OpAccessChain %310 %74 %30 +%320 = OpAtomicOr %6 %321 %9 %313 %26 +OpStore %305 %320 +%323 = OpAccessChain %310 %74 %30 +%322 = OpAtomicXor %6 %323 %9 %313 %26 +OpStore %305 %322 +%325 = OpAccessChain %310 %74 %30 +%324 = OpAtomicSMin %6 %325 %9 %313 %26 +OpStore %305 %324 +%327 = OpAccessChain %310 %74 %30 +%326 = OpAtomicSMax %6 %327 %9 %313 %26 +OpStore %305 %326 +%329 = OpAccessChain %310 %74 %30 +%328 = OpAtomicExchange %6 %329 %9 %313 %26 +OpStore %305 %328 +%330 = OpAccessChain %310 %74 %30 +OpAtomicStore %330 %9 %313 %312 +OpReturn +OpFunctionEnd +%334 = OpFunction %2 None %94 %333 = OpLabel -OpControlBarrier %30 %30 %335 -OpBranch %336 -%336 = OpLabel -%337 = OpFunctionCall %2 %223 %83 +%331 = OpVariable %63 Function %332 +OpBranch %335 +%335 = OpLabel +%339 = OpLoad %35 %337 +%343 = OpIEqual %341 %339 %340 +%344 = OpAll %342 %343 +OpSelectionMerge %345 None +OpBranchConditional %344 %346 %345 +%346 = OpLabel +OpStore %85 %336 +OpBranch %345 +%345 = OpLabel +OpControlBarrier %30 %30 %347 +OpBranch %348 +%348 = OpLabel +%349 = OpCompositeConstruct %57 %14 %14 %14 %14 +%350 = OpCompositeConstruct %57 %25 %25 %25 %25 +%351 = OpCompositeConstruct %62 %349 %350 +OpStore %331 %351 +%352 = OpFunctionCall %2 %225 %85 +%353 = OpFunctionCall %2 %230 %331 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/access.wgsl b/tests/out/wgsl/access.wgsl index 5510a7988d..7133f53d69 100644 --- a/tests/out/wgsl/access.wgsl +++ b/tests/out/wgsl/access.wgsl @@ -126,6 +126,11 @@ fn assign_through_ptr_fn(p: ptr) { return; } +fn assign_array_through_ptr_fn(foo_2: ptr,2>>) { + (*foo_2) = array,2>(vec4(1.0), vec4(2.0)); + return; +} + @vertex fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { var foo: f32; @@ -137,7 +142,7 @@ fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { test_matrix_within_struct_accesses(); test_matrix_within_array_within_struct_accesses(); let _matrix = bar._matrix; - let arr = bar.arr; + let arr_1 = bar.arr; let b = bar._matrix[3][0]; let a_1 = bar.data[(arrayLength((&bar.data)) - 2u)].value; let c = qux; @@ -187,6 +192,10 @@ fn atomics() { @compute @workgroup_size(1, 1, 1) fn assign_through_ptr() { + var arr: array,2>; + + arr = array,2>(vec4(6.0), vec4(7.0)); assign_through_ptr_fn((&val)); + assign_array_through_ptr_fn((&arr)); return; }