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