diff --git a/test/bug/chromium/1273451.wgsl.expected.hlsl b/test/bug/chromium/1273451.wgsl.expected.hlsl new file mode 100644 index 0000000000..dddd92f31b --- /dev/null +++ b/test/bug/chromium/1273451.wgsl.expected.hlsl @@ -0,0 +1,16 @@ +[numthreads(1, 1, 1)] +void unused_entry_point() { + return; +} + +struct A { + int a; +}; +struct B { + int b; +}; + +B f(A a) { + const B tint_symbol = (B)0; + return tint_symbol; +} diff --git a/test/bug/chromium/1273451.wgsl.expected.spvasm b/test/bug/chromium/1273451.wgsl.expected.spvasm new file mode 100644 index 0000000000..8377d686ef --- /dev/null +++ b/test/bug/chromium/1273451.wgsl.expected.spvasm @@ -0,0 +1,34 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 13 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %unused_entry_point "unused_entry_point" + OpExecutionMode %unused_entry_point LocalSize 1 1 1 + OpName %unused_entry_point "unused_entry_point" + OpName %B "B" + OpMemberName %B 0 "b" + OpName %A "A" + OpMemberName %A 0 "a" + OpName %f "f" + OpName %a "a" + OpMemberDecorate %B 0 Offset 0 + OpMemberDecorate %A 0 Offset 0 + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %B = OpTypeStruct %int + %A = OpTypeStruct %int + %5 = OpTypeFunction %B %A + %12 = OpConstantNull %B +%unused_entry_point = OpFunction %void None %1 + %4 = OpLabel + OpReturn + OpFunctionEnd + %f = OpFunction %B None %5 + %a = OpFunctionParameter %A + %11 = OpLabel + OpReturnValue %12 + OpFunctionEnd diff --git a/test/bug/chromium/1273451.wgsl.expected.wgsl b/test/bug/chromium/1273451.wgsl.expected.wgsl new file mode 100644 index 0000000000..2289c10d94 --- /dev/null +++ b/test/bug/chromium/1273451.wgsl.expected.wgsl @@ -0,0 +1,11 @@ +struct A { + a : i32; +}; + +struct B { + b : i32; +}; + +fn f(a : A) -> B { + return B(); +} diff --git a/test/bug/fxc/dyn_array_idx/write/function.wgsl.expected.hlsl b/test/bug/fxc/dyn_array_idx/write/function.wgsl.expected.hlsl new file mode 100644 index 0000000000..60c1a89bb1 --- /dev/null +++ b/test/bug/fxc/dyn_array_idx/write/function.wgsl.expected.hlsl @@ -0,0 +1,21 @@ +cbuffer cbuffer_ubo : register(b0, space0) { + uint4 ubo[1]; +}; + +struct S { + int data[64]; +}; + +RWByteAddressBuffer result : register(u1, space0); + +[numthreads(1, 1, 1)] +void f() { + S s = (S)0; + { + int tint_symbol_2[64] = s.data; + tint_symbol_2[asint(ubo[0].x)] = 1; + s.data = tint_symbol_2; + } + result.Store(0u, asuint(s.data[3])); + return; +} diff --git a/test/bug/fxc/dyn_array_idx/write/function_via_param.wgsl.expected.hlsl b/test/bug/fxc/dyn_array_idx/write/function_via_param.wgsl.expected.hlsl new file mode 100644 index 0000000000..6d73fa8f2d --- /dev/null +++ b/test/bug/fxc/dyn_array_idx/write/function_via_param.wgsl.expected.hlsl @@ -0,0 +1,25 @@ +cbuffer cbuffer_ubo : register(b0, space0) { + uint4 ubo[1]; +}; + +struct S { + int data[64]; +}; + +RWByteAddressBuffer result : register(u1, space0); + +void x(inout S p) { + { + int tint_symbol_2[64] = p.data; + tint_symbol_2[asint(ubo[0].x)] = 1; + p.data = tint_symbol_2; + } +} + +[numthreads(1, 1, 1)] +void f() { + S s = (S)0; + x(s); + result.Store(0u, asuint(s.data[3])); + return; +} diff --git a/test/bug/fxc/dyn_array_idx/write/private.wgsl.expected.hlsl b/test/bug/fxc/dyn_array_idx/write/private.wgsl.expected.hlsl new file mode 100644 index 0000000000..755fa553f4 --- /dev/null +++ b/test/bug/fxc/dyn_array_idx/write/private.wgsl.expected.hlsl @@ -0,0 +1,21 @@ +cbuffer cbuffer_ubo : register(b0, space0) { + uint4 ubo[1]; +}; + +struct S { + int data[64]; +}; + +RWByteAddressBuffer result : register(u1, space0); +static S s = (S)0; + +[numthreads(1, 1, 1)] +void f() { + { + int tint_symbol_2[64] = s.data; + tint_symbol_2[asint(ubo[0].x)] = 1; + s.data = tint_symbol_2; + } + result.Store(0u, asuint(s.data[3])); + return; +} diff --git a/test/bug/fxc/dyn_array_idx/write/private_via_param.wgsl.expected.hlsl b/test/bug/fxc/dyn_array_idx/write/private_via_param.wgsl.expected.hlsl new file mode 100644 index 0000000000..e8d8ff4e1c --- /dev/null +++ b/test/bug/fxc/dyn_array_idx/write/private_via_param.wgsl.expected.hlsl @@ -0,0 +1,25 @@ +cbuffer cbuffer_ubo : register(b0, space0) { + uint4 ubo[1]; +}; + +struct S { + int data[64]; +}; + +RWByteAddressBuffer result : register(u1, space0); +static S s = (S)0; + +void x(inout S p) { + { + int tint_symbol_2[64] = p.data; + tint_symbol_2[asint(ubo[0].x)] = 1; + p.data = tint_symbol_2; + } +} + +[numthreads(1, 1, 1)] +void f() { + x(s); + result.Store(0u, asuint(s.data[3])); + return; +} diff --git a/test/bug/fxc/indexed_assign_to_array_in_struct/1206.wgsl.expected.msl b/test/bug/fxc/indexed_assign_to_array_in_struct/1206.wgsl.expected.msl new file mode 100644 index 0000000000..96ff937ff4 --- /dev/null +++ b/test/bug/fxc/indexed_assign_to_array_in_struct/1206.wgsl.expected.msl @@ -0,0 +1,38 @@ +#include + +using namespace metal; + +template +inline vec operator*(matrix lhs, packed_vec rhs) { + return lhs * vec(rhs); +} + +template +inline vec operator*(packed_vec lhs, matrix rhs) { + return vec(lhs) * rhs; +} + +struct Simulation { + /* 0x0000 */ uint i; +}; +struct tint_array_wrapper { + /* 0x0000 */ float3 arr[8]; +}; +struct Particle { + /* 0x0000 */ tint_array_wrapper position; + /* 0x0080 */ float lifetime; + /* 0x0084 */ int8_t tint_pad[12]; + /* 0x0090 */ float4 color; + /* 0x00a0 */ packed_float3 velocity; + /* 0x00ac */ int8_t tint_pad_1[4]; +}; +struct Particles { + /* 0x0000 */ Particle p[1]; +}; + +kernel void tint_symbol(const device Particles* tint_symbol_1 [[buffer(1)]], const constant Simulation* tint_symbol_2 [[buffer(0)]]) { + Particle particle = (*(tint_symbol_1)).p[0]; + particle.position.arr[(*(tint_symbol_2)).i] = particle.position.arr[(*(tint_symbol_2)).i]; + return; +} + diff --git a/test/bug/fxc/indexed_assign_to_array_in_struct/1206.wgsl.expected.spvasm b/test/bug/fxc/indexed_assign_to_array_in_struct/1206.wgsl.expected.spvasm new file mode 100644 index 0000000000..68c8ccd8a1 --- /dev/null +++ b/test/bug/fxc/indexed_assign_to_array_in_struct/1206.wgsl.expected.spvasm @@ -0,0 +1,78 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 37 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Particles "Particles" + OpMemberName %Particles 0 "p" + OpName %Particle "Particle" + OpMemberName %Particle 0 "position" + OpMemberName %Particle 1 "lifetime" + OpMemberName %Particle 2 "color" + OpMemberName %Particle 3 "velocity" + OpName %particles "particles" + OpName %Simulation "Simulation" + OpMemberName %Simulation 0 "i" + OpName %sim "sim" + OpName %main "main" + OpName %particle "particle" + OpDecorate %Particles Block + OpMemberDecorate %Particles 0 Offset 0 + OpMemberDecorate %Particle 0 Offset 0 + OpDecorate %_arr_v3float_uint_8 ArrayStride 16 + OpMemberDecorate %Particle 1 Offset 128 + OpMemberDecorate %Particle 2 Offset 144 + OpMemberDecorate %Particle 3 Offset 160 + OpDecorate %_runtimearr_Particle ArrayStride 176 + OpDecorate %particles NonWritable + OpDecorate %particles DescriptorSet 1 + OpDecorate %particles Binding 3 + OpDecorate %Simulation Block + OpMemberDecorate %Simulation 0 Offset 0 + OpDecorate %sim NonWritable + OpDecorate %sim DescriptorSet 1 + OpDecorate %sim Binding 4 + %float = OpTypeFloat 32 + %v3float = OpTypeVector %float 3 + %uint = OpTypeInt 32 0 + %uint_8 = OpConstant %uint 8 +%_arr_v3float_uint_8 = OpTypeArray %v3float %uint_8 + %v4float = OpTypeVector %float 4 + %Particle = OpTypeStruct %_arr_v3float_uint_8 %float %v4float %v3float +%_runtimearr_Particle = OpTypeRuntimeArray %Particle + %Particles = OpTypeStruct %_runtimearr_Particle +%_ptr_StorageBuffer_Particles = OpTypePointer StorageBuffer %Particles + %particles = OpVariable %_ptr_StorageBuffer_Particles StorageBuffer + %Simulation = OpTypeStruct %uint +%_ptr_Uniform_Simulation = OpTypePointer Uniform %Simulation + %sim = OpVariable %_ptr_Uniform_Simulation Uniform + %void = OpTypeVoid + %15 = OpTypeFunction %void + %uint_0 = OpConstant %uint 0 + %int = OpTypeInt 32 1 + %int_0 = OpConstant %int 0 +%_ptr_StorageBuffer_Particle = OpTypePointer StorageBuffer %Particle +%_ptr_Function_Particle = OpTypePointer Function %Particle + %27 = OpConstantNull %Particle +%_ptr_Uniform_uint = OpTypePointer Uniform %uint +%_ptr_Function_v3float = OpTypePointer Function %v3float + %main = OpFunction %void None %15 + %18 = OpLabel + %particle = OpVariable %_ptr_Function_Particle Function %27 + %23 = OpAccessChain %_ptr_StorageBuffer_Particle %particles %uint_0 %int_0 + %24 = OpLoad %Particle %23 + OpStore %particle %24 + %29 = OpAccessChain %_ptr_Uniform_uint %sim %uint_0 + %30 = OpLoad %uint %29 + %32 = OpAccessChain %_ptr_Function_v3float %particle %uint_0 %30 + %33 = OpAccessChain %_ptr_Uniform_uint %sim %uint_0 + %34 = OpLoad %uint %33 + %35 = OpAccessChain %_ptr_Function_v3float %particle %uint_0 %34 + %36 = OpLoad %v3float %35 + OpStore %32 %36 + OpReturn + OpFunctionEnd diff --git a/test/bug/fxc/indexed_assign_to_array_in_struct/1206.wgsl.expected.wgsl b/test/bug/fxc/indexed_assign_to_array_in_struct/1206.wgsl.expected.wgsl new file mode 100644 index 0000000000..239e359145 --- /dev/null +++ b/test/bug/fxc/indexed_assign_to_array_in_struct/1206.wgsl.expected.wgsl @@ -0,0 +1,24 @@ +struct Simulation { + i : u32; +}; + +struct Particle { + position : array, 8>; + lifetime : f32; + color : vec4; + velocity : vec3; +}; + +struct Particles { + p : array; +}; + +[[group(1), binding(3)]] var particles : Particles; + +[[group(1), binding(4)]] var sim : Simulation; + +[[stage(compute), workgroup_size(1)]] +fn main() { + var particle = particles.p[0]; + particle.position[sim.i] = particle.position[sim.i]; +} diff --git a/test/bug/tint/749.spvasm.expected.hlsl b/test/bug/tint/749.spvasm.expected.hlsl new file mode 100644 index 0000000000..d6f801a9cf --- /dev/null +++ b/test/bug/tint/749.spvasm.expected.hlsl @@ -0,0 +1,1658 @@ +struct QuicksortObject { + int numbers[10]; +}; + +static QuicksortObject obj = (QuicksortObject)0; +static float4 gl_FragCoord = float4(0.0f, 0.0f, 0.0f, 0.0f); +cbuffer cbuffer_x_188 : register(b0, space0) { + uint4 x_188[1]; +}; +static float4 x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f); + +void swap_i1_i1_(inout int i, inout int j) { + int temp = 0; + const int x_932 = temp; + temp = 0; + temp = x_932; + const float3 x_523 = float3(float3(1.0f, 2.0f, 3.0f).z, float3(1.0f, 2.0f, 3.0f).y, float3(1.0f, 2.0f, 3.0f).z); + const int x_933 = i; + i = 0; + i = x_933; + const int x_28 = i; + const int x_934 = j; + j = 0; + j = x_934; + const float3 x_524 = float3(x_523.y, x_523.x, x_523.y); + const int x_935 = temp; + temp = 0; + temp = x_935; + const int x_30_save = x_28; + const int x_936 = obj.numbers[x_30_save]; + { + int tint_symbol_1[10] = obj.numbers; + tint_symbol_1[x_30_save] = 0; + obj.numbers = tint_symbol_1; + } + { + int tint_symbol_3[10] = obj.numbers; + tint_symbol_3[x_30_save] = x_936; + obj.numbers = tint_symbol_3; + } + const int x_31 = obj.numbers[x_30_save]; + const int x_937 = temp; + temp = 0; + temp = x_937; + temp = x_31; + const int x_938 = j; + j = 0; + j = x_938; + const float3 x_525 = float3(x_523.z, float3(1.0f, 2.0f, 3.0f).x, x_523.y); + const int x_939 = i; + i = 0; + i = x_939; + const int x_32 = i; + const int x_940 = obj.numbers[x_30_save]; + { + int tint_symbol_5[10] = obj.numbers; + tint_symbol_5[x_30_save] = 0; + obj.numbers = tint_symbol_5; + } + { + int tint_symbol_7[10] = obj.numbers; + tint_symbol_7[x_30_save] = x_940; + obj.numbers = tint_symbol_7; + } + const int x_33 = j; + const int x_941 = i; + i = 0; + i = x_941; + const float3 x_526 = float3(x_525.x, x_525.z, x_525.z); + const int x_942 = obj.numbers[x_30_save]; + { + int tint_symbol_9[10] = obj.numbers; + tint_symbol_9[x_30_save] = 0; + obj.numbers = tint_symbol_9; + } + { + int tint_symbol_11[10] = obj.numbers; + tint_symbol_11[x_30_save] = x_942; + obj.numbers = tint_symbol_11; + } + const int x_34_save = x_33; + const int x_35 = obj.numbers[x_34_save]; + const QuicksortObject x_943 = obj; + const int tint_symbol_52[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_53 = {tint_symbol_52}; + obj = tint_symbol_53; + obj = x_943; + const float2 x_527 = float2(x_526.x, x_526.x); + const int x_36_save = x_32; + const float3 x_528 = float3(x_524.x, x_524.z, x_524.x); + { + int tint_symbol_13[10] = obj.numbers; + tint_symbol_13[x_36_save] = x_35; + obj.numbers = tint_symbol_13; + } + const QuicksortObject x_944 = obj; + const int tint_symbol_54[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_55 = {tint_symbol_54}; + obj = tint_symbol_55; + obj = x_944; + const float3 x_529 = float3(x_526.y, x_526.z, x_526.x); + const int x_945 = i; + i = 0; + i = x_945; + const int x_37 = j; + const int x_946 = temp; + temp = 0; + temp = x_946; + const float2 x_530 = float2(x_529.z, x_529.y); + const int x_947 = obj.numbers[x_34_save]; + { + int tint_symbol_15[10] = obj.numbers; + tint_symbol_15[x_34_save] = 0; + obj.numbers = tint_symbol_15; + } + { + int tint_symbol_17[10] = obj.numbers; + tint_symbol_17[x_34_save] = x_947; + obj.numbers = tint_symbol_17; + } + const int x_38 = temp; + const int x_948 = j; + j = 0; + j = x_948; + const float3 x_531 = float3(x_527.x, x_526.y, x_526.x); + const int x_949 = obj.numbers[x_36_save]; + { + int tint_symbol_19[10] = obj.numbers; + tint_symbol_19[x_36_save] = 0; + obj.numbers = tint_symbol_19; + } + { + int tint_symbol_21[10] = obj.numbers; + tint_symbol_21[x_36_save] = x_949; + obj.numbers = tint_symbol_21; + } + const QuicksortObject x_950 = obj; + const int tint_symbol_56[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_57 = {tint_symbol_56}; + obj = tint_symbol_57; + obj = x_950; + const float3 x_532 = float3(x_528.x, x_528.y, x_528.x); + const int x_951 = obj.numbers[x_34_save]; + { + int tint_symbol_23[10] = obj.numbers; + tint_symbol_23[x_34_save] = 0; + obj.numbers = tint_symbol_23; + } + { + int tint_symbol_25[10] = obj.numbers; + tint_symbol_25[x_34_save] = x_951; + obj.numbers = tint_symbol_25; + } + { + int tint_symbol_27[10] = obj.numbers; + tint_symbol_27[x_37] = x_38; + obj.numbers = tint_symbol_27; + } + return; +} + +int performPartition_i1_i1_(inout int l, inout int h) { + int param_3 = 0; + int i_1 = 0; + int j_1 = 0; + int param_2 = 0; + int param_1 = 0; + int param = 0; + int pivot = 0; + float2 x_537 = float2(0.0f, 0.0f); + float3 x_538 = float3(0.0f, 0.0f, 0.0f); + const int x_952 = h; + h = 0; + h = x_952; + const int x_41 = h; + const int x_953 = l; + l = 0; + l = x_953; + const int x_42_save = x_41; + const int x_954 = obj.numbers[x_42_save]; + { + int tint_symbol_29[10] = obj.numbers; + tint_symbol_29[x_42_save] = 0; + obj.numbers = tint_symbol_29; + } + { + int tint_symbol_31[10] = obj.numbers; + tint_symbol_31[x_42_save] = x_954; + obj.numbers = tint_symbol_31; + } + const int x_43 = obj.numbers[x_42_save]; + const int x_955 = param_3; + param_3 = 0; + param_3 = x_955; + const float3 x_534 = float3(float3(1.0f, 2.0f, 3.0f).z, float3(1.0f, 2.0f, 3.0f).x, float3(1.0f, 2.0f, 3.0f).z); + const int x_956 = param_1; + param_1 = 0; + param_1 = x_956; + pivot = x_43; + const int x_45 = l; + const int x_957 = h; + h = 0; + h = x_957; + const int x_958 = j_1; + j_1 = 0; + j_1 = x_958; + const float3 x_535 = float3(x_534.y, x_534.z, x_534.y); + const int x_959 = l; + l = 0; + l = x_959; + i_1 = (x_45 - asint(1u)); + const int x_49 = l; + const float3 x_536 = float3(x_534.x, x_534.z, x_535.x); + j_1 = 10; + const QuicksortObject x_960 = obj; + const int tint_symbol_58[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_59 = {tint_symbol_58}; + obj = tint_symbol_59; + obj = x_960; + [loop] while (true) { + const int x_961 = pivot; + pivot = 0; + pivot = x_961; + const int x_962 = param_1; + param_1 = 0; + param_1 = x_962; + const int x_55 = j_1; + const int x_963 = pivot; + pivot = 0; + pivot = x_963; + x_537 = float2(float3(1.0f, 2.0f, 3.0f).y, float3(1.0f, 2.0f, 3.0f).z); + const QuicksortObject x_964 = obj; + const int tint_symbol_60[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_61 = {tint_symbol_60}; + obj = tint_symbol_61; + obj = x_964; + const int x_56 = h; + const int x_965 = h; + h = 0; + h = x_965; + const int x_966 = param; + param = 0; + param = x_966; + const int x_967 = j_1; + j_1 = 0; + j_1 = x_967; + x_538 = float3(x_534.x, x_537.y, x_534.z); + const int x_968 = param; + param = 0; + param = x_968; + if ((x_55 <= (x_56 - asint(1u)))) { + } else { + break; + } + const int x_60 = j_1; + const int x_969 = obj.numbers[x_42_save]; + { + int tint_symbol_33[10] = obj.numbers; + tint_symbol_33[x_42_save] = 0; + obj.numbers = tint_symbol_33; + } + { + int tint_symbol_35[10] = obj.numbers; + tint_symbol_35[x_42_save] = x_969; + obj.numbers = tint_symbol_35; + } + const int x_61_save = x_60; + const int x_970 = h; + h = 0; + h = x_970; + const float3 x_539 = float3(x_537.x, x_535.z, x_537.x); + const int x_971 = param_1; + param_1 = 0; + param_1 = x_971; + const int x_62 = obj.numbers[x_61_save]; + const QuicksortObject x_972 = obj; + const int tint_symbol_62[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_63 = {tint_symbol_62}; + obj = tint_symbol_63; + obj = x_972; + const int x_63 = pivot; + const float2 x_540 = float2(float3(1.0f, 2.0f, 3.0f).y, x_534.z); + const int x_973 = i_1; + i_1 = 0; + i_1 = x_973; + const int x_974 = l; + l = 0; + l = x_974; + const float3 x_541 = float3(x_534.y, x_534.x, x_534.y); + const int x_975 = pivot; + pivot = 0; + pivot = x_975; + if ((x_62 <= x_63)) { + const float3 x_542 = float3(x_541.z, x_541.x, x_541.x); + const int x_976 = param_3; + param_3 = 0; + param_3 = x_976; + const int x_67 = i_1; + const int x_977 = pivot; + pivot = 0; + pivot = x_977; + const float2 x_543 = float2(x_539.x, x_541.y); + const int x_978 = i_1; + i_1 = 0; + i_1 = x_978; + const int x_979 = param; + param = 0; + param = x_979; + i_1 = (x_67 + asint(1u)); + const int x_980 = l; + l = 0; + l = x_980; + const float3 x_544 = float3(float3(1.0f, 2.0f, 3.0f).z, float3(1.0f, 2.0f, 3.0f).y, x_540.x); + const int x_70 = i_1; + const float2 x_545 = float2(x_537.y, x_538.x); + const int x_981 = param; + param = 0; + param = x_981; + param = x_70; + const int x_982 = param; + param = 0; + param = x_982; + const float2 x_546 = float2(x_545.x, x_545.x); + const int x_983 = i_1; + i_1 = 0; + i_1 = x_983; + param_1 = j_1; + const int x_984 = param_3; + param_3 = 0; + param_3 = x_984; + swap_i1_i1_(param, param_1); + const int x_985 = param_1; + param_1 = 0; + param_1 = x_985; + } + const QuicksortObject x_986 = obj; + const int tint_symbol_64[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_65 = {tint_symbol_64}; + obj = tint_symbol_65; + obj = x_986; + { + const int x_987 = h; + h = 0; + h = x_987; + const int x_74 = j_1; + const int x_988 = h; + h = 0; + h = x_988; + const float3 x_547 = float3(x_539.x, x_541.z, x_541.z); + const int x_989 = obj.numbers[x_61_save]; + { + int tint_symbol_37[10] = obj.numbers; + tint_symbol_37[x_61_save] = 0; + obj.numbers = tint_symbol_37; + } + { + int tint_symbol_39[10] = obj.numbers; + tint_symbol_39[x_61_save] = x_989; + obj.numbers = tint_symbol_39; + } + const int x_990 = param; + param = 0; + param = x_990; + j_1 = (1 + x_74); + const int x_991 = param_1; + param_1 = 0; + param_1 = x_991; + const float3 x_548 = float3(x_541.y, x_541.z, x_541.x); + const int x_992 = obj.numbers[x_61_save]; + { + int tint_symbol_41[10] = obj.numbers; + tint_symbol_41[x_61_save] = 0; + obj.numbers = tint_symbol_41; + } + { + int tint_symbol_43[10] = obj.numbers; + tint_symbol_43[x_61_save] = x_992; + obj.numbers = tint_symbol_43; + } + } + } + const int x_76 = i_1; + const int x_993 = obj.numbers[x_42_save]; + { + int tint_symbol_45[10] = obj.numbers; + tint_symbol_45[x_42_save] = 0; + obj.numbers = tint_symbol_45; + } + { + int tint_symbol_47[10] = obj.numbers; + tint_symbol_47[x_42_save] = x_993; + obj.numbers = tint_symbol_47; + } + const float2 x_549 = float2(x_534.x, x_534.y); + const QuicksortObject x_994 = obj; + const int tint_symbol_66[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_67 = {tint_symbol_66}; + obj = tint_symbol_67; + obj = x_994; + const int x_995 = h; + h = 0; + h = x_995; + i_1 = (1 + x_76); + const int x_996 = param_1; + param_1 = 0; + param_1 = x_996; + const int x_79 = i_1; + const int x_997 = j_1; + j_1 = 0; + j_1 = x_997; + const float2 x_550 = float2(x_534.x, x_534.x); + const int x_998 = param_1; + param_1 = 0; + param_1 = x_998; + param_2 = x_79; + const float2 x_551 = float2(x_534.y, x_536.x); + const int x_999 = pivot; + pivot = 0; + pivot = x_999; + const int x_81 = h; + const float2 x_552 = float2(x_550.x, x_549.y); + const int x_1000 = h; + h = 0; + h = x_1000; + param_3 = x_81; + const int x_1001 = i_1; + i_1 = 0; + i_1 = x_1001; + const float2 x_553 = float2(x_549.y, x_552.x); + const int x_1002 = h; + h = 0; + h = x_1002; + swap_i1_i1_(param_2, param_3); + const int x_1003 = l; + l = 0; + l = x_1003; + const float2 x_554 = float2(x_536.z, float3(1.0f, 2.0f, 3.0f).y); + const int x_1004 = param_1; + param_1 = 0; + param_1 = x_1004; + const int x_83 = i_1; + const int x_1005 = param; + param = 0; + param = x_1005; + const float2 x_555 = float2(x_534.y, x_534.x); + const int x_1006 = j_1; + j_1 = 0; + j_1 = x_1006; + return x_83; +} + +void quicksort_() { + int param_4 = 0; + int h_1 = 0; + int p = 0; + int l_1 = 0; + int top = 0; + int stack[10] = (int[10])0; + int param_5 = 0; + l_1 = 0; + const int x_1007 = param_5; + param_5 = 0; + param_5 = x_1007; + h_1 = 9; + const int x_1008[10] = stack; + const int tint_symbol_68[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + stack = tint_symbol_68; + stack = x_1008; + const float2 x_556 = float2(float3(1.0f, 2.0f, 3.0f).y, float3(1.0f, 2.0f, 3.0f).y); + const int x_1009 = param_5; + param_5 = 0; + param_5 = x_1009; + top = -1; + const int x_1010 = p; + p = 0; + p = x_1010; + const int x_93 = top; + const float2 x_557 = float2(float3(1.0f, 2.0f, 3.0f).x, float3(1.0f, 2.0f, 3.0f).x); + const int x_1011 = p; + p = 0; + p = x_1011; + const int x_94 = (x_93 + asint(1u)); + const int x_1012 = top; + top = 0; + top = x_1012; + const float2 x_558 = float2(x_556.y, x_557.y); + const int x_1013 = param_4; + param_4 = 0; + param_4 = x_1013; + top = x_94; + const int x_1014 = h_1; + h_1 = 0; + h_1 = x_1014; + const float3 x_559 = float3(x_557.y, x_557.x, x_557.x); + const int x_1015 = param_4; + param_4 = 0; + param_4 = x_1015; + const int x_95 = l_1; + const QuicksortObject x_1016 = obj; + const int tint_symbol_69[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_70 = {tint_symbol_69}; + obj = tint_symbol_70; + obj = x_1016; + const float3 x_560 = float3(x_559.y, x_559.x, x_557.x); + const int x_96_save = x_94; + const int x_1017[10] = stack; + const int tint_symbol_71[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + stack = tint_symbol_71; + stack = x_1017; + const float3 x_561 = float3(x_556.y, x_556.y, x_556.y); + const int x_1018 = l_1; + l_1 = 0; + l_1 = 0; + stack[x_96_save] = x_95; + const int x_1019 = param_5; + param_5 = 0; + param_5 = x_1019; + const int x_97 = top; + const int x_1020 = param_4; + param_4 = 0; + param_4 = x_1020; + const float3 x_562 = float3(float3(1.0f, 2.0f, 3.0f).z, x_558.y, float3(1.0f, 2.0f, 3.0f).y); + const int x_1021 = stack[x_96_save]; + stack[x_96_save] = 0; + stack[x_96_save] = x_1021; + const int x_98 = (x_97 + 1); + const int x_1022 = stack[x_96_save]; + stack[x_96_save] = 0; + stack[x_96_save] = x_1022; + const float3 x_563 = float3(x_559.x, x_559.z, x_556.y); + top = x_98; + const int x_1023 = param_4; + param_4 = 0; + param_4 = x_1023; + const int x_99 = h_1; + const int x_1024 = param_4; + param_4 = 0; + param_4 = x_1024; + const float3 x_564 = float3(x_558.x, x_561.x, x_558.y); + const int x_1025 = l_1; + l_1 = 0; + l_1 = x_1025; + const int x_100_save = x_98; + const int x_1026 = param_5; + param_5 = 0; + param_5 = x_1026; + const float2 x_565 = float2(x_564.z, x_564.z); + const int x_1027 = p; + p = 0; + p = x_1027; + stack[x_100_save] = x_99; + [loop] while (true) { + const float3 x_566 = float3(x_563.x, x_563.x, x_563.x); + const int x_1028 = h_1; + h_1 = 0; + h_1 = x_1028; + const int x_1029[10] = stack; + const int tint_symbol_72[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + stack = tint_symbol_72; + stack = x_1029; + const int x_106 = top; + const int x_1030[10] = stack; + const int tint_symbol_73[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + stack = tint_symbol_73; + stack = x_1030; + const float2 x_567 = float2(x_558.x, x_564.z); + const int x_1031 = param_4; + param_4 = 0; + param_4 = x_1031; + if ((x_106 >= asint(0u))) { + } else { + break; + } + const QuicksortObject x_1032 = obj; + const int tint_symbol_74[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_75 = {tint_symbol_74}; + obj = tint_symbol_75; + obj = x_1032; + const float3 x_568 = float3(x_559.y, x_559.x, x_563.y); + const int x_1033 = param_4; + param_4 = 0; + param_4 = x_1033; + const int x_108 = top; + const float3 x_569 = float3(x_565.x, x_567.y, x_565.x); + const int x_1034 = h_1; + h_1 = 0; + h_1 = x_1034; + const float2 x_570 = float2(x_556.x, x_556.x); + const int x_1035 = p; + p = 0; + p = x_1035; + top = (x_108 - asint(1u)); + const int x_1036 = p; + p = 0; + p = x_1036; + const int x_110_save = x_108; + const int x_1037 = stack[x_96_save]; + stack[x_96_save] = 0; + stack[x_96_save] = x_1037; + const int x_111 = stack[x_110_save]; + const int x_1038[10] = stack; + const int tint_symbol_76[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + stack = tint_symbol_76; + stack = x_1038; + const float3 x_571 = float3(x_559.y, x_559.x, x_564.y); + const int x_1039 = l_1; + l_1 = 0; + l_1 = x_1039; + h_1 = x_111; + const int x_1040[10] = stack; + const int tint_symbol_77[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + stack = tint_symbol_77; + stack = x_1040; + const float2 x_572 = float2(x_562.y, x_561.y); + const int x_1041 = p; + p = 0; + p = x_1041; + const int x_112 = top; + const int x_1042 = param_4; + param_4 = 0; + param_4 = x_1042; + const int x_1043 = stack[x_100_save]; + stack[x_100_save] = 0; + stack[x_100_save] = x_1043; + const float2 x_573 = float2(float3(1.0f, 2.0f, 3.0f).y, float3(1.0f, 2.0f, 3.0f).z); + top = (x_112 - 1); + const int x_1044 = param_5; + param_5 = 0; + param_5 = x_1044; + const float3 x_574 = float3(x_570.y, x_565.x, x_570.y); + const int x_1045 = h_1; + h_1 = 0; + h_1 = x_1045; + const int x_114_save = x_112; + const float2 x_575 = float2(x_564.y, x_564.z); + const int x_1046 = stack[x_100_save]; + stack[x_100_save] = 0; + stack[x_100_save] = x_1046; + const int x_115 = stack[x_114_save]; + const int x_1047 = p; + p = 0; + p = x_1047; + const float3 x_576 = float3(x_573.y, x_573.y, x_565.x); + const int x_1048 = param_5; + param_5 = 0; + param_5 = x_1048; + l_1 = x_115; + const int x_1049 = top; + top = 0; + top = x_1049; + param_4 = l_1; + const int x_1050 = stack[x_110_save]; + stack[x_110_save] = 0; + stack[x_110_save] = x_1050; + const float2 x_577 = float2(x_569.y, x_569.z); + const int x_120 = h_1; + const float2 x_578 = float2(x_558.x, float3(1.0f, 2.0f, 3.0f).y); + param_5 = x_120; + const int x_1051 = stack[x_100_save]; + stack[x_100_save] = 0; + stack[x_100_save] = x_1051; + const int x_121 = performPartition_i1_i1_(param_4, param_5); + const float2 x_579 = float2(x_567.x, x_568.x); + const int x_1052 = param_5; + param_5 = 0; + param_5 = x_1052; + p = x_121; + const int x_1053 = param_4; + param_4 = 0; + param_4 = x_1053; + const int x_122 = p; + const int x_1054 = h_1; + h_1 = 0; + h_1 = x_1054; + const float2 x_580 = float2(x_568.y, x_568.y); + const int x_1055 = l_1; + l_1 = 0; + l_1 = x_1055; + const int x_1056 = h_1; + h_1 = 0; + h_1 = x_1056; + const int x_124 = l_1; + const int x_1057 = stack[x_110_save]; + stack[x_110_save] = 0; + stack[x_110_save] = x_1057; + const int x_1058 = h_1; + h_1 = 0; + h_1 = x_1058; + const float2 x_582 = float2(x_567.y, x_573.x); + const int x_1059 = stack[x_100_save]; + stack[x_100_save] = 0; + stack[x_100_save] = x_1059; + if (((x_122 - asint(1u)) > x_124)) { + const int x_1060 = param_4; + param_4 = 0; + param_4 = x_1060; + const int x_128 = top; + const float2 x_583 = float2(x_571.y, x_556.y); + const int x_1061 = stack[x_100_save]; + stack[x_100_save] = 0; + stack[x_100_save] = x_1061; + const int x_1062[10] = stack; + const int tint_symbol_78[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + stack = tint_symbol_78; + stack = x_1062; + const float2 x_584 = float2(x_569.z, x_569.y); + const float3 x_585 = float3(x_580.y, x_577.x, x_577.x); + const int x_130 = l_1; + const int x_1063 = stack[x_114_save]; + stack[x_114_save] = 0; + stack[x_114_save] = x_1063; + const float2 x_586 = float2(x_564.x, x_585.x); + const int x_1064 = param_5; + param_5 = 0; + param_5 = x_1064; + const int x_131_save = (1 + x_128); + const int x_1065 = stack[x_110_save]; + stack[x_110_save] = 0; + stack[x_110_save] = x_1065; + const float3 x_587 = float3(x_566.y, x_566.y, x_563.x); + const int x_1066 = param_5; + param_5 = 0; + param_5 = x_1066; + stack[x_131_save] = x_130; + const int x_132 = top; + const int x_1067 = stack[x_100_save]; + stack[x_100_save] = 0; + stack[x_100_save] = x_1067; + const float2 x_588 = float2(x_575.y, x_575.x); + const int x_1068 = stack[x_131_save]; + stack[x_131_save] = 0; + stack[x_131_save] = x_1068; + const int x_133 = asint((1u + asuint(x_132))); + const int x_1069 = stack[x_100_save]; + stack[x_100_save] = 0; + stack[x_100_save] = x_1069; + const float3 x_589 = float3(x_576.z, x_588.y, x_576.z); + const int x_1070 = h_1; + h_1 = 0; + h_1 = x_1070; + top = x_133; + const int x_1071[10] = stack; + const int tint_symbol_79[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + stack = tint_symbol_79; + stack = x_1071; + const int x_134 = p; + const float2 x_590 = float2(x_576.x, x_573.y); + const int x_1072 = stack[x_114_save]; + stack[x_114_save] = 0; + stack[x_114_save] = x_1072; + const int x_136_save = x_133; + const int x_1073 = stack[x_114_save]; + stack[x_114_save] = 0; + stack[x_114_save] = x_1073; + stack[x_136_save] = (x_134 - asint(1u)); + const int x_1074 = stack[x_96_save]; + stack[x_96_save] = 0; + stack[x_96_save] = x_1074; + const float2 x_591 = float2(x_569.z, x_569.y); + const int x_1075 = stack[x_136_save]; + stack[x_136_save] = 0; + stack[x_136_save] = x_1075; + } + const int x_1076 = stack[x_96_save]; + stack[x_96_save] = 0; + stack[x_96_save] = x_1076; + const float2 x_592 = float2(float3(1.0f, 2.0f, 3.0f).x, float3(1.0f, 2.0f, 3.0f).y); + const QuicksortObject x_1077 = obj; + const int tint_symbol_80[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_81 = {tint_symbol_80}; + obj = tint_symbol_81; + obj = x_1077; + const int x_137 = p; + const int x_1078 = stack[x_114_save]; + stack[x_114_save] = 0; + stack[x_114_save] = x_1078; + const float3 x_593 = float3(x_571.z, x_556.x, x_556.y); + const int x_1079 = p; + p = 0; + p = x_1079; + const float3 x_594 = float3(x_563.z, x_563.x, x_575.x); + const int x_1080 = stack[x_114_save]; + stack[x_114_save] = 0; + stack[x_114_save] = x_1080; + const int x_139 = h_1; + const int x_1081 = top; + top = 0; + top = x_1081; + const float3 x_595 = float3(x_560.z, x_568.x, x_560.x); + const int x_1082 = stack[x_100_save]; + stack[x_100_save] = 0; + stack[x_100_save] = x_1082; + const int x_1083 = p; + p = 0; + p = x_1083; + if ((asint((1u + asuint(x_137))) < x_139)) { + const int x_1084 = stack[x_114_save]; + stack[x_114_save] = 0; + stack[x_114_save] = x_1084; + const float2 x_596 = float2(x_592.y, x_582.x); + const int x_1085 = l_1; + l_1 = 0; + l_1 = x_1085; + const int x_143 = top; + const int x_1086 = stack[x_114_save]; + stack[x_114_save] = 0; + stack[x_114_save] = x_1086; + const float3 x_597 = float3(x_562.y, x_560.y, x_560.y); + const int x_144 = (x_143 + 1); + const int x_1087 = param_5; + param_5 = 0; + param_5 = x_1087; + top = x_144; + const int x_1088 = stack[x_114_save]; + stack[x_114_save] = 0; + stack[x_114_save] = x_1088; + const int x_145 = p; + const int x_1089 = param_5; + param_5 = 0; + param_5 = x_1089; + const float3 x_599 = float3(x_560.z, x_560.x, x_568.x); + const int x_1090 = p; + p = 0; + p = x_1090; + const float3 x_600 = float3(x_556.x, x_580.x, x_580.x); + const int x_1091 = stack[x_100_save]; + stack[x_100_save] = 0; + stack[x_100_save] = x_1091; + const int x_147_save = x_144; + const int x_1092 = stack[x_110_save]; + stack[x_110_save] = 0; + stack[x_110_save] = x_1092; + const float2 x_601 = float2(x_563.x, x_563.y); + stack[x_147_save] = asint((1u + asuint(x_145))); + const int x_1093[10] = stack; + const int tint_symbol_82[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + stack = tint_symbol_82; + stack = x_1093; + const int x_148 = top; + const int x_1094 = stack[x_114_save]; + stack[x_114_save] = 0; + stack[x_114_save] = x_1094; + const float2 x_602 = float2(x_565.y, x_599.y); + const int x_1095[10] = stack; + const int tint_symbol_83[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + stack = tint_symbol_83; + stack = x_1095; + const int x_149 = (x_148 + asint(1u)); + const int x_1096 = stack[x_147_save]; + stack[x_147_save] = 0; + stack[x_147_save] = x_1096; + top = x_149; + const int x_1097 = param_4; + param_4 = 0; + param_4 = x_1097; + const int x_150 = h_1; + const int x_1098 = stack[x_100_save]; + stack[x_100_save] = 0; + stack[x_100_save] = x_1098; + const int x_1099 = stack[x_96_save]; + stack[x_96_save] = 0; + stack[x_96_save] = x_1099; + stack[x_149] = x_150; + const int x_1100 = stack[x_114_save]; + stack[x_114_save] = 0; + stack[x_114_save] = x_1100; + const float3 x_603 = float3(x_568.y, x_564.x, x_564.x); + const int x_1101 = l_1; + l_1 = 0; + l_1 = x_1101; + } + const int x_1102 = stack[x_100_save]; + stack[x_100_save] = 0; + stack[x_100_save] = x_1102; + { + const int x_1103 = l_1; + l_1 = 0; + l_1 = x_1103; + const float2 x_604 = float2(x_563.z, x_564.x); + const QuicksortObject x_1104 = obj; + const int tint_symbol_84[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_85 = {tint_symbol_84}; + obj = tint_symbol_85; + obj = x_1104; + } + } + const int x_1105 = h_1; + h_1 = 0; + h_1 = x_1105; + return; +} + +void main_1() { + float3 color = float3(0.0f, 0.0f, 0.0f); + int i_2 = 0; + float2 uv = float2(0.0f, 0.0f); + const float2 x_717 = uv; + uv = float2(0.0f, 0.0f); + uv = x_717; + i_2 = 0; + const QuicksortObject x_721 = obj; + const int tint_symbol_86[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_87 = {tint_symbol_86}; + obj = tint_symbol_87; + obj = x_721; + if (true) { + const QuicksortObject x_722 = obj; + const int tint_symbol_88[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_89 = {tint_symbol_88}; + obj = tint_symbol_89; + obj = x_722; + const float2 x_431 = float2(float3(1.0f, 2.0f, 3.0f).x, float3(1.0f, 2.0f, 3.0f).x); + const int x_158 = i_2; + const float2 x_723 = uv; + uv = float2(0.0f, 0.0f); + uv = x_723; + const float3 x_725 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_725; + const float2 x_432 = float2(x_431.y, x_431.y); + const QuicksortObject x_726 = obj; + const int tint_symbol_90[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_91 = {tint_symbol_90}; + obj = tint_symbol_91; + obj = x_726; + } + const QuicksortObject x_756 = obj; + const int tint_symbol_92[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_93 = {tint_symbol_92}; + obj = tint_symbol_93; + obj = x_756; + const float2 x_446 = float2(float2(0.0f, 0.0f).x, float2(0.0f, 0.0f).x); + const int x_757 = i_2; + i_2 = 0; + i_2 = x_757; + quicksort_(); + const QuicksortObject x_758 = obj; + const int tint_symbol_94[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_95 = {tint_symbol_94}; + obj = tint_symbol_95; + obj = x_758; + const float4 x_184 = gl_FragCoord; + const float2 x_759 = uv; + uv = float2(0.0f, 0.0f); + uv = x_759; + const float2 x_447 = float2(float2(0.0f, 0.0f).y, float2(0.0f, 0.0f).y); + const float2 x_760 = uv; + uv = float2(0.0f, 0.0f); + uv = x_760; + const float2 x_185 = float2(x_184.x, x_184.y); + const float3 x_448 = float3(x_185.y, x_446.y, x_446.y); + const QuicksortObject x_761 = obj; + const int tint_symbol_96[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_97 = {tint_symbol_96}; + obj = tint_symbol_97; + obj = x_761; + const float2 x_762 = uv; + uv = float2(0.0f, 0.0f); + uv = x_762; + const float2 x_191 = asfloat(x_188[0].xy); + const QuicksortObject x_763 = obj; + const int tint_symbol_98[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_99 = {tint_symbol_98}; + obj = tint_symbol_99; + obj = x_763; + const float3 x_449 = float3(x_184.y, float3(1.0f, 2.0f, 3.0f).z, x_184.w); + const float3 x_764 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_764; + const float2 x_192 = (x_185 / x_191); + const QuicksortObject x_765 = obj; + const int tint_symbol_100[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_101 = {tint_symbol_100}; + obj = tint_symbol_101; + obj = x_765; + const float2 x_450 = float2(x_447.x, x_185.y); + const float3 x_766 = color; + color = float3(0.0f, 0.0f, 0.0f); + const float3 x_767 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_767; + color = x_766; + uv = x_192; + color = float3(1.0f, 2.0f, 3.0f); + const float3 x_768 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_768; + const float3 x_451 = float3(x_185.x, x_185.y, x_446.y); + const QuicksortObject x_769 = obj; + const int tint_symbol_102[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_103 = {tint_symbol_102}; + obj = tint_symbol_103; + obj = x_769; + const int x_770 = obj.numbers[0u]; + obj.numbers[0u] = 0; + obj.numbers[0u] = x_770; + const int x_201 = obj.numbers[0u]; + const QuicksortObject x_771 = obj; + const int tint_symbol_104[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_105 = {tint_symbol_104}; + obj = tint_symbol_105; + obj = x_771; + const int x_772 = obj.numbers[0u]; + obj.numbers[0u] = 0; + obj.numbers[0u] = x_772; + const float x_206 = color.x; + const float x_773 = color.x; + color.x = 0.0f; + color.x = x_773; + const float2 x_452 = float2(float3(1.0f, 2.0f, 3.0f).z, float3(1.0f, 2.0f, 3.0f).y); + const int x_774 = i_2; + i_2 = 0; + i_2 = x_774; + const QuicksortObject x_775 = obj; + const int tint_symbol_106[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_107 = {tint_symbol_106}; + obj = tint_symbol_107; + obj = x_775; + const float3 x_453 = float3(x_451.x, x_450.x, x_450.y); + color.x = (x_206 + float(x_201)); + const float2 x_776 = uv; + uv = float2(0.0f, 0.0f); + uv = x_776; + const float2 x_777 = uv; + uv = float2(0.0f, 0.0f); + uv = x_777; + const float2 x_454 = float2(x_184.y, x_184.y); + const float x_210 = uv.x; + const float2 x_455 = float2(x_192.y, x_192.x); + const float x_778 = uv.x; + uv.x = 0.0f; + uv.x = x_778; + const QuicksortObject x_779 = obj; + const int tint_symbol_108[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_109 = {tint_symbol_108}; + obj = tint_symbol_109; + obj = x_779; + if ((x_210 > 0.25f)) { + const int x_780 = i_2; + i_2 = 0; + i_2 = x_780; + const int x_781 = obj.numbers[0u]; + obj.numbers[0u] = 0; + obj.numbers[0u] = x_781; + const float3 x_456 = float3(float2(0.0f, 0.0f).y, x_448.y, x_448.y); + const float x_782 = uv.x; + uv.x = 0.0f; + uv.x = x_782; + const int x_216 = obj.numbers[1]; + const QuicksortObject x_783 = obj; + const int tint_symbol_110[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_111 = {tint_symbol_110}; + obj = tint_symbol_111; + obj = x_783; + const float2 x_457 = float2(x_454.x, x_454.x); + const float2 x_784 = uv; + uv = float2(0.0f, 0.0f); + uv = x_784; + const QuicksortObject x_785 = obj; + const int tint_symbol_112[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_113 = {tint_symbol_112}; + obj = tint_symbol_113; + obj = x_785; + const float2 x_458 = float2(float3(1.0f, 2.0f, 3.0f).z, float2(0.0f, 0.0f).y); + const int x_786 = i_2; + i_2 = 0; + i_2 = x_786; + const float x_219 = color[0]; + const float x_787 = color[0]; + color[0] = 0.0f; + color[0] = x_787; + const float3 x_788 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_788; + const float3 x_789 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_789; + const float3 x_459 = float3(x_454.y, x_454.y, x_447.y); + const float x_790 = color[0]; + color[0] = 0.0f; + color[0] = x_790; + color.x = (float(x_216) + x_219); + const int x_791 = obj.numbers[0u]; + obj.numbers[0u] = 0; + obj.numbers[0u] = x_791; + } + const float x_792 = uv.x; + uv.x = 0.0f; + uv.x = x_792; + const float x_793 = uv.x; + uv.x = 0.0f; + uv.x = x_793; + const float x_223 = uv.x; + const float x_794 = uv.x; + uv.x = 0.0f; + uv.x = x_794; + const float3 x_460 = float3(x_453.z, x_453.y, x_453.y); + const float2 x_795 = uv; + uv = float2(0.0f, 0.0f); + uv = x_795; + const float x_796 = uv.x; + uv.x = 0.0f; + uv.x = x_796; + const float2 x_461 = float2(float2(0.0f, 0.0f).y, float2(0.0f, 0.0f).y); + const float x_797 = uv.x; + uv.x = 0.0f; + uv.x = x_797; + if ((x_223 > 0.5f)) { + const float x_798 = uv.x; + uv.x = 0.0f; + uv.x = x_798; + const float2 x_462 = float2(x_446.x, x_446.x); + const float x_799 = color.x; + color.x = 0.0f; + color.x = x_799; + const float x_800 = color.x; + color.x = 0.0f; + color.x = x_800; + const float3 x_463 = float3(x_453.x, x_453.z, x_461.y); + const float x_801 = color.x; + color.x = 0.0f; + color.x = x_801; + const int x_230 = obj.numbers[2u]; + const float x_802 = uv.x; + uv.x = 0.0f; + uv.x = x_802; + const float x_803 = color.x; + color.x = 0.0f; + color.x = x_803; + const int x_804 = obj.numbers[2u]; + obj.numbers[2u] = 0; + obj.numbers[2u] = x_804; + const float2 x_464 = float2(x_450.y, x_191.x); + const float x_805 = color.y; + color.y = 0.0f; + color.y = x_805; + const float x_234 = color.y; + const int x_806 = obj.numbers[2u]; + obj.numbers[2u] = 0; + obj.numbers[2u] = x_806; + const float2 x_465 = float2(x_463.x, x_185.x); + const float x_807 = color.x; + color.x = 0.0f; + color.x = x_807; + const int x_808 = i_2; + i_2 = 0; + i_2 = x_808; + const float2 x_466 = float2(x_455.y, float2(0.0f, 0.0f).y); + const int x_809 = i_2; + i_2 = 0; + i_2 = x_809; + color.y = (float(x_230) + x_234); + const float x_810 = uv.x; + uv.x = 0.0f; + uv.x = x_810; + } + const int x_811 = i_2; + i_2 = 0; + i_2 = x_811; + const float2 x_467 = float2(x_191.x, x_191.x); + const float x_812 = uv.x; + uv.x = 0.0f; + uv.x = x_812; + const float x_238 = uv[0]; + const float3 x_813 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_813; + const float x_814 = color.x; + color.x = 0.0f; + color.x = x_814; + if ((x_238 > 0.75f)) { + const float x_815 = color.x; + color.x = 0.0f; + color.x = x_815; + const int x_245 = obj.numbers[3]; + const float x_816 = color.x; + color.x = 0.0f; + color.x = x_816; + const QuicksortObject x_817 = obj; + const int tint_symbol_114[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_115 = {tint_symbol_114}; + obj = tint_symbol_115; + obj = x_817; + const float3 x_468 = float3(x_467.x, x_467.x, x_467.x); + const float x_818 = uv[0]; + uv[0] = 0.0f; + uv[0] = x_818; + const float x_819 = uv.x; + uv.x = 0.0f; + uv.x = x_819; + const float x_249 = color.z; + const float3 x_820 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_820; + const float3 x_469 = float3(x_467.x, x_191.y, x_467.y); + const float x_821 = color.z; + color.z = 0.0f; + color.z = x_821; + const int x_822 = obj.numbers[0u]; + obj.numbers[0u] = 0; + obj.numbers[0u] = x_822; + const float2 x_470 = float2(float2(0.0f, 0.0f).x, float2(0.0f, 0.0f).y); + const float x_823 = color.z; + color.z = 0.0f; + color.z = x_823; + color.z = (x_249 + float(x_245)); + const float2 x_824 = uv; + uv = float2(0.0f, 0.0f); + uv = x_824; + const float2 x_471 = float2(x_470.y, x_470.y); + } + const float x_825 = uv[0]; + uv[0] = 0.0f; + uv[0] = x_825; + const float3 x_472 = float3(x_454.x, x_454.y, x_454.y); + const int x_254 = obj.numbers[4]; + const float x_826 = uv[0]; + uv[0] = 0.0f; + uv[0] = x_826; + const float3 x_827 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_827; + const float3 x_473 = float3(x_446.y, x_453.x, x_453.x); + const int x_828 = obj.numbers[4]; + obj.numbers[4] = 0; + obj.numbers[4] = x_828; + const float2 x_474 = float2(x_191.x, x_184.z); + const float x_829 = uv.x; + uv.x = 0.0f; + uv.x = x_829; + const float x_257 = color.y; + const float x_830 = color.y; + color.y = 0.0f; + color.y = x_830; + const float2 x_475 = float2(x_467.x, x_450.x); + const float x_831 = uv.x; + uv.x = 0.0f; + uv.x = x_831; + const float x_832 = color.x; + color.x = 0.0f; + color.x = x_832; + const float2 x_476 = float2(x_451.z, x_460.y); + color.y = (x_257 + float(x_254)); + const float3 x_477 = float3(float2(0.0f, 0.0f).x, x_472.x, float2(0.0f, 0.0f).y); + const float x_833 = uv.x; + uv.x = 0.0f; + uv.x = x_833; + const float x_834 = color.x; + color.x = 0.0f; + color.x = x_834; + const float2 x_478 = float2(x_472.x, x_472.y); + const float x_835 = uv.y; + uv.y = 0.0f; + uv.y = x_835; + const float x_261 = uv.y; + const int x_836 = i_2; + i_2 = 0; + i_2 = x_836; + const float3 x_479 = float3(float2(0.0f, 0.0f).y, x_454.y, float2(0.0f, 0.0f).x); + const int x_837 = obj.numbers[0u]; + obj.numbers[0u] = 0; + obj.numbers[0u] = x_837; + const float x_838 = color.y; + color.y = 0.0f; + color.y = x_838; + const float3 x_480 = float3(x_446.x, x_446.x, float2(0.0f, 0.0f).y); + const float x_839 = uv.x; + uv.x = 0.0f; + uv.x = x_839; + if ((x_261 > 0.25f)) { + const float2 x_481 = float2(x_447.x, x_480.z); + const float3 x_840 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_840; + const int x_267 = obj.numbers[5u]; + const float x_841 = color.x; + color.x = 0.0f; + color.x = x_841; + const int x_842 = i_2; + i_2 = 0; + i_2 = x_842; + const int x_843 = i_2; + i_2 = 0; + i_2 = x_843; + const float x_270 = color.x; + const float x_844 = uv[0]; + uv[0] = 0.0f; + uv[0] = x_844; + const float3 x_482 = float3(x_455.x, x_475.y, x_455.y); + const QuicksortObject x_845 = obj; + const int tint_symbol_116[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_117 = {tint_symbol_116}; + obj = tint_symbol_117; + obj = x_845; + const float x_846 = uv.y; + uv.y = 0.0f; + uv.y = x_846; + const int x_847 = i_2; + i_2 = 0; + i_2 = x_847; + const float3 x_483 = float3(x_184.w, x_184.w, x_192.x); + const float x_848 = uv.x; + uv.x = 0.0f; + uv.x = x_848; + color.x = (float(x_267) + x_270); + const float3 x_484 = float3(x_454.y, x_450.x, x_454.y); + const float x_849 = uv.x; + uv.x = 0.0f; + uv.x = x_849; + } + const float x_850 = color.x; + color.x = 0.0f; + color.x = x_850; + const float3 x_485 = float3(x_467.x, x_450.y, x_450.x); + const float x_851 = uv.y; + uv.y = 0.0f; + uv.y = x_851; + const int x_852 = obj.numbers[4]; + obj.numbers[4] = 0; + obj.numbers[4] = x_852; + const float x_274 = uv.y; + const int x_853 = obj.numbers[0u]; + obj.numbers[0u] = 0; + obj.numbers[0u] = x_853; + if ((x_274 > 0.5f)) { + const float x_854 = uv.x; + uv.x = 0.0f; + uv.x = x_854; + const float2 x_486 = float2(x_480.y, x_455.y); + const float x_855 = color.y; + color.y = 0.0f; + color.y = x_855; + const float2 x_487 = float2(x_449.z, x_449.y); + const float x_856 = uv.y; + uv.y = 0.0f; + uv.y = x_856; + const int x_280 = obj.numbers[6u]; + const float x_857 = uv.y; + uv.y = 0.0f; + uv.y = x_857; + const int x_858 = i_2; + i_2 = 0; + i_2 = x_858; + const int x_859 = obj.numbers[4]; + obj.numbers[4] = 0; + obj.numbers[4] = x_859; + const float2 x_488 = float2(x_473.z, x_473.y); + const float x_283 = color.y; + const float2 x_860 = uv; + uv = float2(0.0f, 0.0f); + uv = x_860; + const float x_861 = color.x; + color.x = 0.0f; + color.x = x_861; + const float2 x_489 = float2(x_475.y, x_475.x); + const int x_862 = obj.numbers[6u]; + obj.numbers[6u] = 0; + obj.numbers[6u] = x_862; + const int x_863 = obj.numbers[6u]; + obj.numbers[6u] = 0; + obj.numbers[6u] = x_863; + const float2 x_490 = float2(x_480.z, x_480.z); + const QuicksortObject x_864 = obj; + const int tint_symbol_118[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_119 = {tint_symbol_118}; + obj = tint_symbol_119; + obj = x_864; + color.y = (float(x_280) + x_283); + const float x_865 = color.x; + color.x = 0.0f; + color.x = x_865; + const float2 x_491 = float2(float3(1.0f, 2.0f, 3.0f).y, x_454.x); + const float x_866 = color.y; + color.y = 0.0f; + color.y = x_866; + } + const float2 x_492 = float2(x_455.y, x_455.y); + const float x_867 = color.x; + color.x = 0.0f; + color.x = x_867; + const float x_287 = uv.y; + const QuicksortObject x_868 = obj; + const int tint_symbol_120[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_121 = {tint_symbol_120}; + obj = tint_symbol_121; + obj = x_868; + const float2 x_493 = float2(x_475.x, x_475.y); + const float x_869 = uv[0]; + uv[0] = 0.0f; + uv[0] = x_869; + const float x_870 = color.y; + color.y = 0.0f; + color.y = x_870; + const float3 x_494 = float3(x_191.x, x_191.y, x_191.y); + const int x_871 = obj.numbers[4]; + obj.numbers[4] = 0; + obj.numbers[4] = x_871; + if ((x_287 > 0.75f)) { + const float3 x_872 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_872; + const float x_873 = color.x; + color.x = 0.0f; + color.x = x_873; + const float3 x_495 = float3(x_192.y, x_192.x, x_192.y); + const float3 x_874 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_874; + const int x_293 = obj.numbers[7]; + const float x_875 = uv.x; + uv.x = 0.0f; + uv.x = x_875; + const float3 x_496 = float3(x_475.x, x_467.y, x_467.x); + const float x_876 = color.y; + color.y = 0.0f; + color.y = x_876; + const float2 x_497 = float2(x_477.x, x_461.y); + const int x_877 = obj.numbers[0u]; + obj.numbers[0u] = 0; + obj.numbers[0u] = x_877; + const float x_878 = color.y; + color.y = 0.0f; + color.y = x_878; + const float3 x_498 = float3(x_478.x, x_478.y, x_478.x); + const float x_879 = color.x; + color.x = 0.0f; + color.x = x_879; + const float x_296 = color.z; + const float x_880 = uv.y; + uv.y = 0.0f; + uv.y = x_880; + const float2 x_499 = float2(x_184.x, x_184.y); + const float x_881 = uv.x; + uv.x = 0.0f; + uv.x = x_881; + const float x_882 = uv.y; + uv.y = 0.0f; + uv.y = x_882; + const float x_883 = uv.y; + uv.y = 0.0f; + uv.y = x_883; + const float3 x_500 = float3(x_499.y, x_499.y, x_494.z); + const float x_884 = color.z; + color.z = 0.0f; + color.z = x_884; + color.z = (float(x_293) + x_296); + const float x_885 = color.y; + color.y = 0.0f; + color.y = x_885; + const float2 x_501 = float2(x_453.x, x_453.z); + const float x_886 = color.x; + color.x = 0.0f; + color.x = x_886; + } + const int x_887 = i_2; + i_2 = 0; + i_2 = x_887; + const float2 x_502 = float2(x_451.y, x_192.y); + const float2 x_888 = uv; + uv = float2(0.0f, 0.0f); + uv = x_888; + const int x_301 = obj.numbers[8]; + const int x_889 = i_2; + i_2 = 0; + i_2 = x_889; + const float2 x_503 = float2(x_185.x, x_451.z); + const int x_890 = obj.numbers[8]; + obj.numbers[8] = 0; + obj.numbers[8] = x_890; + const float x_891 = color.y; + color.y = 0.0f; + color.y = x_891; + const float2 x_504 = float2(x_453.y, float2(0.0f, 0.0f).x); + const float x_892 = color.x; + color.x = 0.0f; + color.x = x_892; + const float3 x_505 = float3(x_504.x, x_504.y, x_504.x); + const float x_893 = color.z; + color.z = 0.0f; + color.z = x_893; + const float x_304 = color.z; + const float x_894 = color.x; + color.x = 0.0f; + color.x = x_894; + const float2 x_506 = float2(x_493.x, x_492.x); + const int x_895 = obj.numbers[4]; + obj.numbers[4] = 0; + obj.numbers[4] = x_895; + const float x_896 = uv.y; + uv.y = 0.0f; + uv.y = x_896; + const float2 x_507 = float2(x_461.x, x_447.x); + const float x_897 = color.y; + color.y = 0.0f; + color.y = x_897; + color.z = (x_304 + float(x_301)); + const float2 x_898 = uv; + uv = float2(0.0f, 0.0f); + uv = x_898; + const float x_899 = uv.x; + uv.x = 0.0f; + uv.x = x_899; + const float3 x_508 = float3(x_461.y, x_461.x, x_506.y); + const float x_900 = uv.x; + uv.x = 0.0f; + uv.x = x_900; + const float x_308 = uv.x; + const float x_901 = color.y; + color.y = 0.0f; + color.y = x_901; + const float3 x_509 = float3(x_503.y, x_503.x, x_448.z); + const float x_902 = uv.y; + uv.y = 0.0f; + uv.y = x_902; + const float x_310 = uv.y; + const float x_903 = uv.y; + uv.y = 0.0f; + uv.y = x_903; + const float x_904 = color.z; + color.z = 0.0f; + color.z = x_904; + const float3 x_510 = float3(float3(1.0f, 2.0f, 3.0f).y, x_485.y, x_485.z); + const float x_905 = color.z; + color.z = 0.0f; + color.z = x_905; + const int x_906 = i_2; + i_2 = 0; + i_2 = x_906; + const float2 x_511 = float2(x_485.z, x_485.y); + const float3 x_907 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_907; + const float x_908 = uv.y; + uv.y = 0.0f; + uv.y = x_908; + const float3 x_512 = float3(x_455.y, x_455.y, x_455.y); + const int x_909 = obj.numbers[4]; + obj.numbers[4] = 0; + obj.numbers[4] = x_909; + if ((abs((x_308 - x_310)) < 0.25f)) { + const float x_910 = uv.x; + uv.x = 0.0f; + uv.x = x_910; + const QuicksortObject x_911 = obj; + const int tint_symbol_122[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_123 = {tint_symbol_122}; + obj = tint_symbol_123; + obj = x_911; + const float3 x_513 = float3(x_505.z, x_505.x, x_448.x); + const int x_912 = obj.numbers[8]; + obj.numbers[8] = 0; + obj.numbers[8] = x_912; + const int x_317 = obj.numbers[9u]; + const float3 x_514 = float3(x_474.y, x_474.y, x_474.y); + const float x_913 = uv.y; + uv.y = 0.0f; + uv.y = x_913; + const float x_320 = color.x; + const float x_914 = uv.y; + uv.y = 0.0f; + uv.y = x_914; + const float2 x_515 = float2(x_502.x, x_502.y); + const float x_915 = color.x; + color.x = 0.0f; + color.x = x_915; + const float3 x_916 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_916; + const float2 x_516 = float2(x_452.x, x_452.x); + const float2 x_917 = uv; + uv = float2(0.0f, 0.0f); + uv = x_917; + const float x_918 = uv.x; + uv.x = 0.0f; + uv.x = x_918; + const float3 x_517 = float3(float2(0.0f, 0.0f).x, float2(0.0f, 0.0f).x, float2(0.0f, 0.0f).y); + color.x = (float(x_317) + x_320); + const float x_919 = color.x; + color.x = 0.0f; + color.x = x_919; + const float3 x_518 = float3(x_480.y, x_508.x, x_480.x); + const float x_920 = color.x; + color.x = 0.0f; + color.x = x_920; + } + const float x_921 = uv.y; + uv.y = 0.0f; + uv.y = x_921; + const float3 x_325 = color; + const float x_922 = uv[0]; + uv[0] = 0.0f; + uv[0] = x_922; + const float3 x_519 = float3(x_447.x, x_446.x, x_446.y); + const float3 x_326 = normalize(x_325); + const float x_923 = uv.x; + uv.x = 0.0f; + uv.x = x_923; + const QuicksortObject x_924 = obj; + const int tint_symbol_124[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_125 = {tint_symbol_124}; + obj = tint_symbol_125; + obj = x_924; + const QuicksortObject x_925 = obj; + const int tint_symbol_126[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_127 = {tint_symbol_126}; + obj = tint_symbol_127; + obj = x_925; + const float x_926 = color.y; + color.y = 0.0f; + color.y = x_926; + const float2 x_520 = float2(x_506.y, x_519.y); + const float x_927 = color.y; + color.y = 0.0f; + color.y = x_927; + const float4 x_330 = float4(x_326.x, x_326.y, x_326.z, 1.0f); + const float x_928 = uv.y; + uv.y = 0.0f; + uv.y = x_928; + const float3 x_521 = float3(float3(1.0f, 2.0f, 3.0f).y, float3(1.0f, 2.0f, 3.0f).y, x_520.y); + const float x_929 = uv.x; + uv.x = 0.0f; + uv.x = x_929; + x_GLF_color = x_330; + const QuicksortObject x_930 = obj; + const int tint_symbol_128[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_129 = {tint_symbol_128}; + obj = tint_symbol_129; + obj = x_930; + const float3 x_522 = float3(x_330.w, x_330.y, x_493.x); + const float x_931 = color.x; + color.x = 0.0f; + color.x = x_931; + return; +} + +struct main_out { + float4 x_GLF_color_1; +}; +struct tint_symbol_49 { + float4 gl_FragCoord_param : SV_Position; +}; +struct tint_symbol_50 { + float4 x_GLF_color_1 : SV_Target0; +}; + +main_out main_inner(float4 gl_FragCoord_param) { + gl_FragCoord = gl_FragCoord_param; + main_1(); + const main_out tint_symbol_130 = {x_GLF_color}; + return tint_symbol_130; +} + +tint_symbol_50 main(tint_symbol_49 tint_symbol_48) { + const main_out inner_result = main_inner(tint_symbol_48.gl_FragCoord_param); + tint_symbol_50 wrapper_result = (tint_symbol_50)0; + wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1; + return wrapper_result; +} diff --git a/test/bug/tint/998.wgsl.expected.hlsl b/test/bug/tint/998.wgsl.expected.hlsl new file mode 100644 index 0000000000..a197d251a0 --- /dev/null +++ b/test/bug/tint/998.wgsl.expected.hlsl @@ -0,0 +1,21 @@ +cbuffer cbuffer_constants : register(b0, space1) { + uint4 constants[1]; +}; + +RWByteAddressBuffer result : register(u1, space1); + +struct S { + uint data[3]; +}; + +static S s = (S)0; + +[numthreads(1, 1, 1)] +void main() { + { + uint tint_symbol_1[3] = s.data; + tint_symbol_1[constants[0].x] = 0u; + s.data = tint_symbol_1; + } + return; +} diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.msl b/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.msl new file mode 100644 index 0000000000..dad374d3f0 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.msl @@ -0,0 +1,25 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint i; +}; +struct InnerS { + int v; +}; +struct tint_array_wrapper { + InnerS arr[8]; +}; +struct OuterS { + tint_array_wrapper a1; +}; + +kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) { + InnerS v = {}; + OuterS s1 = {}; + for(int i = 0; (i < 4); i = as_type((as_type(i) + as_type(1)))) { + s1.a1.arr[(*(tint_symbol_1)).i] = v; + } + return; +} + diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.spvasm b/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.spvasm new file mode 100644 index 0000000000..60f11df715 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.spvasm @@ -0,0 +1,84 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 44 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "i" + OpName %uniforms "uniforms" + OpName %main "main" + OpName %InnerS "InnerS" + OpMemberName %InnerS 0 "v" + OpName %v "v" + OpName %OuterS "OuterS" + OpMemberName %OuterS 0 "a1" + OpName %s1 "s1" + OpName %i "i" + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpDecorate %uniforms NonWritable + OpDecorate %uniforms DescriptorSet 1 + OpDecorate %uniforms Binding 4 + OpMemberDecorate %InnerS 0 Offset 0 + OpMemberDecorate %OuterS 0 Offset 0 + OpDecorate %_arr_InnerS_uint_8 ArrayStride 4 + %uint = OpTypeInt 32 0 + %Uniforms = OpTypeStruct %uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %void = OpTypeVoid + %5 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %InnerS = OpTypeStruct %int +%_ptr_Function_InnerS = OpTypePointer Function %InnerS + %13 = OpConstantNull %InnerS + %uint_8 = OpConstant %uint 8 +%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8 + %OuterS = OpTypeStruct %_arr_InnerS_uint_8 +%_ptr_Function_OuterS = OpTypePointer Function %OuterS + %19 = OpConstantNull %OuterS + %int_0 = OpConstant %int 0 +%_ptr_Function_int = OpTypePointer Function %int + %23 = OpConstantNull %int + %int_4 = OpConstant %int 4 + %bool = OpTypeBool + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint + %int_1 = OpConstant %int 1 + %main = OpFunction %void None %5 + %8 = OpLabel + %v = OpVariable %_ptr_Function_InnerS Function %13 + %s1 = OpVariable %_ptr_Function_OuterS Function %19 + %i = OpVariable %_ptr_Function_int Function %23 + OpStore %i %int_0 + OpBranch %24 + %24 = OpLabel + OpLoopMerge %25 %26 None + OpBranch %27 + %27 = OpLabel + %29 = OpLoad %int %i + %31 = OpSLessThan %bool %29 %int_4 + %28 = OpLogicalNot %bool %31 + OpSelectionMerge %33 None + OpBranchConditional %28 %34 %33 + %34 = OpLabel + OpBranch %25 + %33 = OpLabel + %37 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %38 = OpLoad %uint %37 + %39 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %38 + %40 = OpLoad %InnerS %v + OpStore %39 %40 + OpBranch %26 + %26 = OpLabel + %41 = OpLoad %int %i + %43 = OpIAdd %int %41 %int_1 + OpStore %i %43 + OpBranch %24 + %25 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.wgsl b/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.wgsl new file mode 100644 index 0000000000..a45995105d --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.wgsl @@ -0,0 +1,22 @@ +struct Uniforms { + i : u32; +}; + +struct InnerS { + v : i32; +}; + +struct OuterS { + a1 : array; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +[[stage(compute), workgroup_size(1)]] +fn main() { + var v : InnerS; + var s1 : OuterS; + for(var i : i32 = 0; (i < 4); i = (i + 1)) { + s1.a1[uniforms.i] = v; + } +} diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.msl b/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.msl new file mode 100644 index 0000000000..6ed5562b3e --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.msl @@ -0,0 +1,25 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint i; +}; +struct InnerS { + int v; +}; +struct tint_array_wrapper { + InnerS arr[8]; +}; +struct OuterS { + tint_array_wrapper a1; +}; + +kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) { + InnerS v = {}; + OuterS s1 = {}; + for(int i = 0; (i < 4); s1.a1.arr[(*(tint_symbol_1)).i] = v) { + i = as_type((as_type(i) + as_type(1))); + } + return; +} + diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.spvasm b/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.spvasm new file mode 100644 index 0000000000..223d6261b9 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.spvasm @@ -0,0 +1,84 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 44 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "i" + OpName %uniforms "uniforms" + OpName %main "main" + OpName %InnerS "InnerS" + OpMemberName %InnerS 0 "v" + OpName %v "v" + OpName %OuterS "OuterS" + OpMemberName %OuterS 0 "a1" + OpName %s1 "s1" + OpName %i "i" + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpDecorate %uniforms NonWritable + OpDecorate %uniforms DescriptorSet 1 + OpDecorate %uniforms Binding 4 + OpMemberDecorate %InnerS 0 Offset 0 + OpMemberDecorate %OuterS 0 Offset 0 + OpDecorate %_arr_InnerS_uint_8 ArrayStride 4 + %uint = OpTypeInt 32 0 + %Uniforms = OpTypeStruct %uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %void = OpTypeVoid + %5 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %InnerS = OpTypeStruct %int +%_ptr_Function_InnerS = OpTypePointer Function %InnerS + %13 = OpConstantNull %InnerS + %uint_8 = OpConstant %uint 8 +%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8 + %OuterS = OpTypeStruct %_arr_InnerS_uint_8 +%_ptr_Function_OuterS = OpTypePointer Function %OuterS + %19 = OpConstantNull %OuterS + %int_0 = OpConstant %int 0 +%_ptr_Function_int = OpTypePointer Function %int + %23 = OpConstantNull %int + %int_4 = OpConstant %int 4 + %bool = OpTypeBool + %int_1 = OpConstant %int 1 + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint + %main = OpFunction %void None %5 + %8 = OpLabel + %v = OpVariable %_ptr_Function_InnerS Function %13 + %s1 = OpVariable %_ptr_Function_OuterS Function %19 + %i = OpVariable %_ptr_Function_int Function %23 + OpStore %i %int_0 + OpBranch %24 + %24 = OpLabel + OpLoopMerge %25 %26 None + OpBranch %27 + %27 = OpLabel + %29 = OpLoad %int %i + %31 = OpSLessThan %bool %29 %int_4 + %28 = OpLogicalNot %bool %31 + OpSelectionMerge %33 None + OpBranchConditional %28 %34 %33 + %34 = OpLabel + OpBranch %25 + %33 = OpLabel + %35 = OpLoad %int %i + %37 = OpIAdd %int %35 %int_1 + OpStore %i %37 + OpBranch %26 + %26 = OpLabel + %40 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %41 = OpLoad %uint %40 + %42 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %41 + %43 = OpLoad %InnerS %v + OpStore %42 %43 + OpBranch %24 + %25 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.wgsl b/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.wgsl new file mode 100644 index 0000000000..ac653ea519 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.wgsl @@ -0,0 +1,22 @@ +struct Uniforms { + i : u32; +}; + +struct InnerS { + v : i32; +}; + +struct OuterS { + a1 : array; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +[[stage(compute), workgroup_size(1)]] +fn main() { + var v : InnerS; + var s1 : OuterS; + for(var i : i32 = 0; (i < 4); s1.a1[uniforms.i] = v) { + i = (i + 1); + } +} diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.msl b/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.msl new file mode 100644 index 0000000000..7e27cf01fa --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.msl @@ -0,0 +1,25 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint i; +}; +struct InnerS { + int v; +}; +struct tint_array_wrapper { + InnerS arr[8]; +}; +struct OuterS { + tint_array_wrapper a1; +}; + +kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) { + InnerS v = {}; + OuterS s1 = {}; + int i = 0; + for(s1.a1.arr[(*(tint_symbol_1)).i] = v; (i < 4); i = as_type((as_type(i) + as_type(1)))) { + } + return; +} + diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.spvasm b/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.spvasm new file mode 100644 index 0000000000..ccdfdf38b6 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.spvasm @@ -0,0 +1,84 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 44 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "i" + OpName %uniforms "uniforms" + OpName %main "main" + OpName %InnerS "InnerS" + OpMemberName %InnerS 0 "v" + OpName %v "v" + OpName %OuterS "OuterS" + OpMemberName %OuterS 0 "a1" + OpName %s1 "s1" + OpName %i "i" + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpDecorate %uniforms NonWritable + OpDecorate %uniforms DescriptorSet 1 + OpDecorate %uniforms Binding 4 + OpMemberDecorate %InnerS 0 Offset 0 + OpMemberDecorate %OuterS 0 Offset 0 + OpDecorate %_arr_InnerS_uint_8 ArrayStride 4 + %uint = OpTypeInt 32 0 + %Uniforms = OpTypeStruct %uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %void = OpTypeVoid + %5 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %InnerS = OpTypeStruct %int +%_ptr_Function_InnerS = OpTypePointer Function %InnerS + %13 = OpConstantNull %InnerS + %uint_8 = OpConstant %uint 8 +%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8 + %OuterS = OpTypeStruct %_arr_InnerS_uint_8 +%_ptr_Function_OuterS = OpTypePointer Function %OuterS + %19 = OpConstantNull %OuterS + %int_0 = OpConstant %int 0 +%_ptr_Function_int = OpTypePointer Function %int + %23 = OpConstantNull %int + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint + %int_4 = OpConstant %int 4 + %bool = OpTypeBool + %int_1 = OpConstant %int 1 + %main = OpFunction %void None %5 + %8 = OpLabel + %v = OpVariable %_ptr_Function_InnerS Function %13 + %s1 = OpVariable %_ptr_Function_OuterS Function %19 + %i = OpVariable %_ptr_Function_int Function %23 + OpStore %i %int_0 + %26 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %27 = OpLoad %uint %26 + %28 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %27 + %29 = OpLoad %InnerS %v + OpStore %28 %29 + OpBranch %30 + %30 = OpLabel + OpLoopMerge %31 %32 None + OpBranch %33 + %33 = OpLabel + %35 = OpLoad %int %i + %37 = OpSLessThan %bool %35 %int_4 + %34 = OpLogicalNot %bool %37 + OpSelectionMerge %39 None + OpBranchConditional %34 %40 %39 + %40 = OpLabel + OpBranch %31 + %39 = OpLabel + OpBranch %32 + %32 = OpLabel + %41 = OpLoad %int %i + %43 = OpIAdd %int %41 %int_1 + OpStore %i %43 + OpBranch %30 + %31 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.wgsl b/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.wgsl new file mode 100644 index 0000000000..5c08a50136 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.wgsl @@ -0,0 +1,22 @@ +struct Uniforms { + i : u32; +}; + +struct InnerS { + v : i32; +}; + +struct OuterS { + a1 : array; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +[[stage(compute), workgroup_size(1)]] +fn main() { + var v : InnerS; + var s1 : OuterS; + var i : i32 = 0; + for(s1.a1[uniforms.i] = v; (i < 4); i = (i + 1)) { + } +} diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.msl b/test/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.msl new file mode 100644 index 0000000000..b01712d4fe --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.msl @@ -0,0 +1,36 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint i; + /* 0x0004 */ uint j; +}; +struct InnerS { + int v; +}; +struct tint_array_wrapper { + InnerS arr[8]; +}; +struct S1 { + tint_array_wrapper a2; +}; +struct tint_array_wrapper_1 { + S1 arr[8]; +}; +struct OuterS { + tint_array_wrapper_1 a1; +}; + +uint getNextIndex(thread uint* const tint_symbol_1) { + *(tint_symbol_1) = (*(tint_symbol_1) + 1u); + return *(tint_symbol_1); +} + +kernel void tint_symbol(const constant Uniforms* tint_symbol_3 [[buffer(0)]]) { + thread uint tint_symbol_2 = 0u; + InnerS v = {}; + OuterS s = {}; + s.a1.arr[getNextIndex(&(tint_symbol_2))].a2.arr[(*(tint_symbol_3)).j] = v; + return; +} + diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.spvasm b/test/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.spvasm new file mode 100644 index 0000000000..3fc8cbf354 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.spvasm @@ -0,0 +1,79 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 39 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %nextIndex "nextIndex" + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "i" + OpMemberName %Uniforms 1 "j" + OpName %uniforms "uniforms" + OpName %getNextIndex "getNextIndex" + OpName %main "main" + OpName %InnerS "InnerS" + OpMemberName %InnerS 0 "v" + OpName %v "v" + OpName %OuterS "OuterS" + OpMemberName %OuterS 0 "a1" + OpName %S1 "S1" + OpMemberName %S1 0 "a2" + OpName %s "s" + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpMemberDecorate %Uniforms 1 Offset 4 + OpDecorate %uniforms NonWritable + OpDecorate %uniforms DescriptorSet 1 + OpDecorate %uniforms Binding 4 + OpMemberDecorate %InnerS 0 Offset 0 + OpMemberDecorate %OuterS 0 Offset 0 + OpMemberDecorate %S1 0 Offset 0 + OpDecorate %_arr_InnerS_uint_8 ArrayStride 4 + OpDecorate %_arr_S1_uint_8 ArrayStride 32 + %uint = OpTypeInt 32 0 +%_ptr_Private_uint = OpTypePointer Private %uint + %4 = OpConstantNull %uint + %nextIndex = OpVariable %_ptr_Private_uint Private %4 + %Uniforms = OpTypeStruct %uint %uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %8 = OpTypeFunction %uint + %uint_1 = OpConstant %uint 1 + %void = OpTypeVoid + %15 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %InnerS = OpTypeStruct %int +%_ptr_Function_InnerS = OpTypePointer Function %InnerS + %23 = OpConstantNull %InnerS + %uint_8 = OpConstant %uint 8 +%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8 + %S1 = OpTypeStruct %_arr_InnerS_uint_8 +%_arr_S1_uint_8 = OpTypeArray %S1 %uint_8 + %OuterS = OpTypeStruct %_arr_S1_uint_8 +%_ptr_Function_OuterS = OpTypePointer Function %OuterS + %31 = OpConstantNull %OuterS + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint +%getNextIndex = OpFunction %uint None %8 + %10 = OpLabel + %11 = OpLoad %uint %nextIndex + %13 = OpIAdd %uint %11 %uint_1 + OpStore %nextIndex %13 + %14 = OpLoad %uint %nextIndex + OpReturnValue %14 + OpFunctionEnd + %main = OpFunction %void None %15 + %18 = OpLabel + %v = OpVariable %_ptr_Function_InnerS Function %23 + %s = OpVariable %_ptr_Function_OuterS Function %31 + %33 = OpFunctionCall %uint %getNextIndex + %35 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1 + %36 = OpLoad %uint %35 + %37 = OpAccessChain %_ptr_Function_InnerS %s %uint_0 %33 %uint_0 %36 + %38 = OpLoad %InnerS %v + OpStore %37 %38 + OpReturn + OpFunctionEnd diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.wgsl b/test/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.wgsl new file mode 100644 index 0000000000..06620e8b39 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.wgsl @@ -0,0 +1,32 @@ +struct Uniforms { + i : u32; + j : u32; +}; + +struct InnerS { + v : i32; +}; + +struct S1 { + a2 : array; +}; + +struct OuterS { + a1 : array; +}; + +var nextIndex : u32; + +fn getNextIndex() -> u32 { + nextIndex = (nextIndex + 1u); + return nextIndex; +} + +[[group(1), binding(4)]] var uniforms : Uniforms; + +[[stage(compute), workgroup_size(1)]] +fn main() { + var v : InnerS; + var s : OuterS; + s.a1[getNextIndex()].a2[uniforms.j] = v; +} diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_array.wgsl.expected.msl b/test/statements/assign/indexed_assign_to_array_in_struct/struct_array.wgsl.expected.msl new file mode 100644 index 0000000000..9430b2c971 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_array.wgsl.expected.msl @@ -0,0 +1,23 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint i; +}; +struct InnerS { + int v; +}; +struct tint_array_wrapper { + InnerS arr[8]; +}; +struct OuterS { + tint_array_wrapper a1; +}; + +kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) { + InnerS v = {}; + OuterS s1 = {}; + s1.a1.arr[(*(tint_symbol_1)).i] = v; + return; +} + diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_array.wgsl.expected.spvasm b/test/statements/assign/indexed_assign_to_array_in_struct/struct_array.wgsl.expected.spvasm new file mode 100644 index 0000000000..8507c70c66 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_array.wgsl.expected.spvasm @@ -0,0 +1,55 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 26 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "i" + OpName %uniforms "uniforms" + OpName %main "main" + OpName %InnerS "InnerS" + OpMemberName %InnerS 0 "v" + OpName %v "v" + OpName %OuterS "OuterS" + OpMemberName %OuterS 0 "a1" + OpName %s1 "s1" + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpDecorate %uniforms NonWritable + OpDecorate %uniforms DescriptorSet 1 + OpDecorate %uniforms Binding 4 + OpMemberDecorate %InnerS 0 Offset 0 + OpMemberDecorate %OuterS 0 Offset 0 + OpDecorate %_arr_InnerS_uint_8 ArrayStride 4 + %uint = OpTypeInt 32 0 + %Uniforms = OpTypeStruct %uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %void = OpTypeVoid + %5 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %InnerS = OpTypeStruct %int +%_ptr_Function_InnerS = OpTypePointer Function %InnerS + %13 = OpConstantNull %InnerS + %uint_8 = OpConstant %uint 8 +%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8 + %OuterS = OpTypeStruct %_arr_InnerS_uint_8 +%_ptr_Function_OuterS = OpTypePointer Function %OuterS + %19 = OpConstantNull %OuterS + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint + %main = OpFunction %void None %5 + %8 = OpLabel + %v = OpVariable %_ptr_Function_InnerS Function %13 + %s1 = OpVariable %_ptr_Function_OuterS Function %19 + %22 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %23 = OpLoad %uint %22 + %24 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %23 + %25 = OpLoad %InnerS %v + OpStore %24 %25 + OpReturn + OpFunctionEnd diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_array.wgsl.expected.wgsl b/test/statements/assign/indexed_assign_to_array_in_struct/struct_array.wgsl.expected.wgsl new file mode 100644 index 0000000000..9a967aa2db --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_array.wgsl.expected.wgsl @@ -0,0 +1,20 @@ +struct Uniforms { + i : u32; +}; + +struct InnerS { + v : i32; +}; + +struct OuterS { + a1 : array; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +[[stage(compute), workgroup_size(1)]] +fn main() { + var v : InnerS; + var s1 : OuterS; + s1.a1[uniforms.i] = v; +} diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_array.wgsl.expected.msl b/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_array.wgsl.expected.msl new file mode 100644 index 0000000000..020719b55f --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_array.wgsl.expected.msl @@ -0,0 +1,27 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint i; + /* 0x0004 */ uint j; +}; +struct InnerS { + int v; +}; +struct tint_array_wrapper_1 { + InnerS arr[8]; +}; +struct tint_array_wrapper { + tint_array_wrapper_1 arr[8]; +}; +struct OuterS { + tint_array_wrapper a1; +}; + +kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) { + InnerS v = {}; + OuterS s1 = {}; + s1.a1.arr[(*(tint_symbol_1)).i].arr[(*(tint_symbol_1)).j] = v; + return; +} + diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_array.wgsl.expected.spvasm b/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_array.wgsl.expected.spvasm new file mode 100644 index 0000000000..18a14e5320 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_array.wgsl.expected.spvasm @@ -0,0 +1,62 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 30 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "i" + OpMemberName %Uniforms 1 "j" + OpName %uniforms "uniforms" + OpName %main "main" + OpName %InnerS "InnerS" + OpMemberName %InnerS 0 "v" + OpName %v "v" + OpName %OuterS "OuterS" + OpMemberName %OuterS 0 "a1" + OpName %s1 "s1" + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpMemberDecorate %Uniforms 1 Offset 4 + OpDecorate %uniforms NonWritable + OpDecorate %uniforms DescriptorSet 1 + OpDecorate %uniforms Binding 4 + OpMemberDecorate %InnerS 0 Offset 0 + OpMemberDecorate %OuterS 0 Offset 0 + OpDecorate %_arr_InnerS_uint_8 ArrayStride 4 + OpDecorate %_arr__arr_InnerS_uint_8_uint_8 ArrayStride 32 + %uint = OpTypeInt 32 0 + %Uniforms = OpTypeStruct %uint %uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %void = OpTypeVoid + %5 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %InnerS = OpTypeStruct %int +%_ptr_Function_InnerS = OpTypePointer Function %InnerS + %13 = OpConstantNull %InnerS + %uint_8 = OpConstant %uint 8 +%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8 +%_arr__arr_InnerS_uint_8_uint_8 = OpTypeArray %_arr_InnerS_uint_8 %uint_8 + %OuterS = OpTypeStruct %_arr__arr_InnerS_uint_8_uint_8 +%_ptr_Function_OuterS = OpTypePointer Function %OuterS + %20 = OpConstantNull %OuterS + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint + %uint_1 = OpConstant %uint 1 + %main = OpFunction %void None %5 + %8 = OpLabel + %v = OpVariable %_ptr_Function_InnerS Function %13 + %s1 = OpVariable %_ptr_Function_OuterS Function %20 + %23 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %24 = OpLoad %uint %23 + %26 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1 + %27 = OpLoad %uint %26 + %28 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %24 %27 + %29 = OpLoad %InnerS %v + OpStore %28 %29 + OpReturn + OpFunctionEnd diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_array.wgsl.expected.wgsl b/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_array.wgsl.expected.wgsl new file mode 100644 index 0000000000..db81a61ee8 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_array.wgsl.expected.wgsl @@ -0,0 +1,21 @@ +struct Uniforms { + i : u32; + j : u32; +}; + +struct InnerS { + v : i32; +}; + +struct OuterS { + a1 : array, 8>; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +[[stage(compute), workgroup_size(1)]] +fn main() { + var v : InnerS; + var s1 : OuterS; + s1.a1[uniforms.i][uniforms.j] = v; +} diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_struct.wgsl.expected.msl b/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_struct.wgsl.expected.msl new file mode 100644 index 0000000000..c2ee63f2d0 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_struct.wgsl.expected.msl @@ -0,0 +1,26 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint i; +}; +struct InnerS { + int v; +}; +struct S1 { + InnerS s2; +}; +struct tint_array_wrapper { + S1 arr[8]; +}; +struct OuterS { + tint_array_wrapper a1; +}; + +kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) { + InnerS v = {}; + OuterS s1 = {}; + s1.a1.arr[(*(tint_symbol_1)).i].s2 = v; + return; +} + diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_struct.wgsl.expected.spvasm b/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_struct.wgsl.expected.spvasm new file mode 100644 index 0000000000..2b16c9be57 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_struct.wgsl.expected.spvasm @@ -0,0 +1,59 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 27 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "i" + OpName %uniforms "uniforms" + OpName %main "main" + OpName %InnerS "InnerS" + OpMemberName %InnerS 0 "v" + OpName %v "v" + OpName %OuterS "OuterS" + OpMemberName %OuterS 0 "a1" + OpName %S1 "S1" + OpMemberName %S1 0 "s2" + OpName %s1 "s1" + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpDecorate %uniforms NonWritable + OpDecorate %uniforms DescriptorSet 1 + OpDecorate %uniforms Binding 4 + OpMemberDecorate %InnerS 0 Offset 0 + OpMemberDecorate %OuterS 0 Offset 0 + OpMemberDecorate %S1 0 Offset 0 + OpDecorate %_arr_S1_uint_8 ArrayStride 4 + %uint = OpTypeInt 32 0 + %Uniforms = OpTypeStruct %uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %void = OpTypeVoid + %5 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %InnerS = OpTypeStruct %int +%_ptr_Function_InnerS = OpTypePointer Function %InnerS + %13 = OpConstantNull %InnerS + %S1 = OpTypeStruct %InnerS + %uint_8 = OpConstant %uint 8 +%_arr_S1_uint_8 = OpTypeArray %S1 %uint_8 + %OuterS = OpTypeStruct %_arr_S1_uint_8 +%_ptr_Function_OuterS = OpTypePointer Function %OuterS + %20 = OpConstantNull %OuterS + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint + %main = OpFunction %void None %5 + %8 = OpLabel + %v = OpVariable %_ptr_Function_InnerS Function %13 + %s1 = OpVariable %_ptr_Function_OuterS Function %20 + %23 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %24 = OpLoad %uint %23 + %25 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %24 %uint_0 + %26 = OpLoad %InnerS %v + OpStore %25 %26 + OpReturn + OpFunctionEnd diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_struct.wgsl.expected.wgsl b/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_struct.wgsl.expected.wgsl new file mode 100644 index 0000000000..d8a6c91823 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_struct.wgsl.expected.wgsl @@ -0,0 +1,24 @@ +struct Uniforms { + i : u32; +}; + +struct InnerS { + v : i32; +}; + +struct S1 { + s2 : InnerS; +}; + +struct OuterS { + a1 : array; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +[[stage(compute), workgroup_size(1)]] +fn main() { + var v : InnerS; + var s1 : OuterS; + s1.a1[uniforms.i].s2 = v; +} diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_struct_array.wgsl.expected.msl b/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_struct_array.wgsl.expected.msl new file mode 100644 index 0000000000..54092b5100 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_struct_array.wgsl.expected.msl @@ -0,0 +1,30 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint i; + /* 0x0004 */ uint j; +}; +struct InnerS { + int v; +}; +struct tint_array_wrapper { + InnerS arr[8]; +}; +struct S1 { + tint_array_wrapper a2; +}; +struct tint_array_wrapper_1 { + S1 arr[8]; +}; +struct OuterS { + tint_array_wrapper_1 a1; +}; + +kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) { + InnerS v = {}; + OuterS s = {}; + s.a1.arr[(*(tint_symbol_1)).i].a2.arr[(*(tint_symbol_1)).j] = v; + return; +} + diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_struct_array.wgsl.expected.spvasm b/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_struct_array.wgsl.expected.spvasm new file mode 100644 index 0000000000..85b4009803 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_struct_array.wgsl.expected.spvasm @@ -0,0 +1,66 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 31 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "i" + OpMemberName %Uniforms 1 "j" + OpName %uniforms "uniforms" + OpName %main "main" + OpName %InnerS "InnerS" + OpMemberName %InnerS 0 "v" + OpName %v "v" + OpName %OuterS "OuterS" + OpMemberName %OuterS 0 "a1" + OpName %S1 "S1" + OpMemberName %S1 0 "a2" + OpName %s "s" + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpMemberDecorate %Uniforms 1 Offset 4 + OpDecorate %uniforms NonWritable + OpDecorate %uniforms DescriptorSet 1 + OpDecorate %uniforms Binding 4 + OpMemberDecorate %InnerS 0 Offset 0 + OpMemberDecorate %OuterS 0 Offset 0 + OpMemberDecorate %S1 0 Offset 0 + OpDecorate %_arr_InnerS_uint_8 ArrayStride 4 + OpDecorate %_arr_S1_uint_8 ArrayStride 32 + %uint = OpTypeInt 32 0 + %Uniforms = OpTypeStruct %uint %uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %void = OpTypeVoid + %5 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %InnerS = OpTypeStruct %int +%_ptr_Function_InnerS = OpTypePointer Function %InnerS + %13 = OpConstantNull %InnerS + %uint_8 = OpConstant %uint 8 +%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8 + %S1 = OpTypeStruct %_arr_InnerS_uint_8 +%_arr_S1_uint_8 = OpTypeArray %S1 %uint_8 + %OuterS = OpTypeStruct %_arr_S1_uint_8 +%_ptr_Function_OuterS = OpTypePointer Function %OuterS + %21 = OpConstantNull %OuterS + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint + %uint_1 = OpConstant %uint 1 + %main = OpFunction %void None %5 + %8 = OpLabel + %v = OpVariable %_ptr_Function_InnerS Function %13 + %s = OpVariable %_ptr_Function_OuterS Function %21 + %24 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %25 = OpLoad %uint %24 + %27 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1 + %28 = OpLoad %uint %27 + %29 = OpAccessChain %_ptr_Function_InnerS %s %uint_0 %25 %uint_0 %28 + %30 = OpLoad %InnerS %v + OpStore %29 %30 + OpReturn + OpFunctionEnd diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_struct_array.wgsl.expected.wgsl b/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_struct_array.wgsl.expected.wgsl new file mode 100644 index 0000000000..9efd02bf1f --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_array_struct_array.wgsl.expected.wgsl @@ -0,0 +1,25 @@ +struct Uniforms { + i : u32; + j : u32; +}; + +struct InnerS { + v : i32; +}; + +struct S1 { + a2 : array; +}; + +struct OuterS { + a1 : array; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +[[stage(compute), workgroup_size(1)]] +fn main() { + var v : InnerS; + var s : OuterS; + s.a1[uniforms.i].a2[uniforms.j] = v; +} diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array.wgsl.expected.msl b/test/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array.wgsl.expected.msl new file mode 100644 index 0000000000..ffa277ba34 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array.wgsl.expected.msl @@ -0,0 +1,19 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint i; +}; +struct InnerS { + /* 0x0000 */ int v; +}; +struct OuterS { + /* 0x0000 */ InnerS a1[1]; +}; + +kernel void tint_symbol(device OuterS* tint_symbol_1 [[buffer(1)]], const constant Uniforms* tint_symbol_2 [[buffer(0)]]) { + InnerS v = {}; + (*(tint_symbol_1)).a1[(*(tint_symbol_2)).i] = v; + return; +} + diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array.wgsl.expected.spvasm b/test/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array.wgsl.expected.spvasm new file mode 100644 index 0000000000..f974ffac47 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array.wgsl.expected.spvasm @@ -0,0 +1,57 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 25 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "i" + OpName %uniforms "uniforms" + OpName %OuterS "OuterS" + OpMemberName %OuterS 0 "a1" + OpName %InnerS "InnerS" + OpMemberName %InnerS 0 "v" + OpName %s1 "s1" + OpName %main "main" + OpName %v "v" + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpDecorate %uniforms NonWritable + OpDecorate %uniforms DescriptorSet 1 + OpDecorate %uniforms Binding 4 + OpDecorate %OuterS Block + OpMemberDecorate %OuterS 0 Offset 0 + OpMemberDecorate %InnerS 0 Offset 0 + OpDecorate %_runtimearr_InnerS ArrayStride 4 + OpDecorate %s1 Binding 0 + OpDecorate %s1 DescriptorSet 0 + %uint = OpTypeInt 32 0 + %Uniforms = OpTypeStruct %uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %int = OpTypeInt 32 1 + %InnerS = OpTypeStruct %int +%_runtimearr_InnerS = OpTypeRuntimeArray %InnerS + %OuterS = OpTypeStruct %_runtimearr_InnerS +%_ptr_StorageBuffer_OuterS = OpTypePointer StorageBuffer %OuterS + %s1 = OpVariable %_ptr_StorageBuffer_OuterS StorageBuffer + %void = OpTypeVoid + %11 = OpTypeFunction %void +%_ptr_Function_InnerS = OpTypePointer Function %InnerS + %17 = OpConstantNull %InnerS + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint +%_ptr_StorageBuffer_InnerS = OpTypePointer StorageBuffer %InnerS + %main = OpFunction %void None %11 + %14 = OpLabel + %v = OpVariable %_ptr_Function_InnerS Function %17 + %20 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %21 = OpLoad %uint %20 + %23 = OpAccessChain %_ptr_StorageBuffer_InnerS %s1 %uint_0 %21 + %24 = OpLoad %InnerS %v + OpStore %23 %24 + OpReturn + OpFunctionEnd diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array.wgsl.expected.wgsl b/test/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array.wgsl.expected.wgsl new file mode 100644 index 0000000000..4eae4791e9 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array.wgsl.expected.wgsl @@ -0,0 +1,21 @@ +struct Uniforms { + i : u32; +}; + +struct InnerS { + v : i32; +}; + +struct OuterS { + a1 : array; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +[[binding(0), group(0)]] var s1 : OuterS; + +[[stage(compute), workgroup_size(1)]] +fn main() { + var v : InnerS; + s1.a1[uniforms.i] = v; +} diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array_struct_array.wgsl.expected.msl b/test/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array_struct_array.wgsl.expected.msl new file mode 100644 index 0000000000..1f40a6227a --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array_struct_array.wgsl.expected.msl @@ -0,0 +1,26 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint i; + /* 0x0004 */ uint j; +}; +struct InnerS { + /* 0x0000 */ int v; +}; +struct tint_array_wrapper { + /* 0x0000 */ InnerS arr[8]; +}; +struct S1 { + /* 0x0000 */ tint_array_wrapper a2; +}; +struct OuterS { + /* 0x0000 */ S1 a1[1]; +}; + +kernel void tint_symbol(device OuterS* tint_symbol_1 [[buffer(1)]], const constant Uniforms* tint_symbol_2 [[buffer(0)]]) { + InnerS v = {}; + (*(tint_symbol_1)).a1[(*(tint_symbol_2)).i].a2.arr[(*(tint_symbol_2)).j] = v; + return; +} + diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array_struct_array.wgsl.expected.spvasm b/test/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array_struct_array.wgsl.expected.spvasm new file mode 100644 index 0000000000..dcaad5fc74 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array_struct_array.wgsl.expected.spvasm @@ -0,0 +1,69 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 31 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "i" + OpMemberName %Uniforms 1 "j" + OpName %uniforms "uniforms" + OpName %OuterS "OuterS" + OpMemberName %OuterS 0 "a1" + OpName %S1 "S1" + OpMemberName %S1 0 "a2" + OpName %InnerS "InnerS" + OpMemberName %InnerS 0 "v" + OpName %s "s" + OpName %main "main" + OpName %v "v" + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpMemberDecorate %Uniforms 1 Offset 4 + OpDecorate %uniforms NonWritable + OpDecorate %uniforms DescriptorSet 1 + OpDecorate %uniforms Binding 4 + OpDecorate %OuterS Block + OpMemberDecorate %OuterS 0 Offset 0 + OpMemberDecorate %S1 0 Offset 0 + OpMemberDecorate %InnerS 0 Offset 0 + OpDecorate %_arr_InnerS_uint_8 ArrayStride 4 + OpDecorate %_runtimearr_S1 ArrayStride 32 + OpDecorate %s Binding 0 + OpDecorate %s DescriptorSet 0 + %uint = OpTypeInt 32 0 + %Uniforms = OpTypeStruct %uint %uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %int = OpTypeInt 32 1 + %InnerS = OpTypeStruct %int + %uint_8 = OpConstant %uint 8 +%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8 + %S1 = OpTypeStruct %_arr_InnerS_uint_8 +%_runtimearr_S1 = OpTypeRuntimeArray %S1 + %OuterS = OpTypeStruct %_runtimearr_S1 +%_ptr_StorageBuffer_OuterS = OpTypePointer StorageBuffer %OuterS + %s = OpVariable %_ptr_StorageBuffer_OuterS StorageBuffer + %void = OpTypeVoid + %14 = OpTypeFunction %void +%_ptr_Function_InnerS = OpTypePointer Function %InnerS + %20 = OpConstantNull %InnerS + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint + %uint_1 = OpConstant %uint 1 +%_ptr_StorageBuffer_InnerS = OpTypePointer StorageBuffer %InnerS + %main = OpFunction %void None %14 + %17 = OpLabel + %v = OpVariable %_ptr_Function_InnerS Function %20 + %23 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %24 = OpLoad %uint %23 + %26 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1 + %27 = OpLoad %uint %26 + %29 = OpAccessChain %_ptr_StorageBuffer_InnerS %s %uint_0 %24 %uint_0 %27 + %30 = OpLoad %InnerS %v + OpStore %29 %30 + OpReturn + OpFunctionEnd diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array_struct_array.wgsl.expected.wgsl b/test/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array_struct_array.wgsl.expected.wgsl new file mode 100644 index 0000000000..a6925a9e3c --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array_struct_array.wgsl.expected.wgsl @@ -0,0 +1,26 @@ +struct Uniforms { + i : u32; + j : u32; +}; + +struct InnerS { + v : i32; +}; + +struct S1 { + a2 : array; +}; + +struct OuterS { + a1 : array; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +[[binding(0), group(0)]] var s : OuterS; + +[[stage(compute), workgroup_size(1)]] +fn main() { + var v : InnerS; + s.a1[uniforms.i].a2[uniforms.j] = v; +} diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_matrix.wgsl.expected.msl b/test/statements/assign/indexed_assign_to_array_in_struct/struct_matrix.wgsl.expected.msl new file mode 100644 index 0000000000..5563dcff99 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_matrix.wgsl.expected.msl @@ -0,0 +1,17 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint i; +}; +struct OuterS { + float2x4 m1; +}; + +kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) { + OuterS s1 = {}; + s1.m1[(*(tint_symbol_1)).i] = float4(1.0f); + s1.m1[(*(tint_symbol_1)).i][(*(tint_symbol_1)).i] = 1.0f; + return; +} + diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_matrix.wgsl.expected.spvasm b/test/statements/assign/indexed_assign_to_array_in_struct/struct_matrix.wgsl.expected.spvasm new file mode 100644 index 0000000000..1333aeb6c1 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_matrix.wgsl.expected.spvasm @@ -0,0 +1,57 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 30 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "i" + OpName %uniforms "uniforms" + OpName %main "main" + OpName %OuterS "OuterS" + OpMemberName %OuterS 0 "m1" + OpName %s1 "s1" + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpDecorate %uniforms NonWritable + OpDecorate %uniforms DescriptorSet 1 + OpDecorate %uniforms Binding 4 + OpMemberDecorate %OuterS 0 Offset 0 + OpMemberDecorate %OuterS 0 ColMajor + OpMemberDecorate %OuterS 0 MatrixStride 16 + %uint = OpTypeInt 32 0 + %Uniforms = OpTypeStruct %uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %void = OpTypeVoid + %5 = OpTypeFunction %void + %float = OpTypeFloat 32 + %v4float = OpTypeVector %float 4 +%mat2v4float = OpTypeMatrix %v4float 2 + %OuterS = OpTypeStruct %mat2v4float +%_ptr_Function_OuterS = OpTypePointer Function %OuterS + %15 = OpConstantNull %OuterS + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint +%_ptr_Function_v4float = OpTypePointer Function %v4float + %float_1 = OpConstant %float 1 + %23 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1 +%_ptr_Function_float = OpTypePointer Function %float + %main = OpFunction %void None %5 + %8 = OpLabel + %s1 = OpVariable %_ptr_Function_OuterS Function %15 + %18 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %19 = OpLoad %uint %18 + %21 = OpAccessChain %_ptr_Function_v4float %s1 %uint_0 %19 + OpStore %21 %23 + %24 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %25 = OpLoad %uint %24 + %26 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %27 = OpLoad %uint %26 + %29 = OpAccessChain %_ptr_Function_float %s1 %uint_0 %25 %27 + OpStore %29 %float_1 + OpReturn + OpFunctionEnd diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_matrix.wgsl.expected.wgsl b/test/statements/assign/indexed_assign_to_array_in_struct/struct_matrix.wgsl.expected.wgsl new file mode 100644 index 0000000000..62a9bbbe27 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_matrix.wgsl.expected.wgsl @@ -0,0 +1,16 @@ +struct Uniforms { + i : u32; +}; + +struct OuterS { + m1 : mat2x4; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +[[stage(compute), workgroup_size(1)]] +fn main() { + var s1 : OuterS; + s1.m1[uniforms.i] = vec4(1.0); + s1.m1[uniforms.i][uniforms.i] = 1.0; +} diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_multiple_arrays.wgsl.expected.msl b/test/statements/assign/indexed_assign_to_array_in_struct/struct_multiple_arrays.wgsl.expected.msl new file mode 100644 index 0000000000..ec021bcbf3 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_multiple_arrays.wgsl.expected.msl @@ -0,0 +1,25 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint i; +}; +struct InnerS { + int v; +}; +struct tint_array_wrapper { + InnerS arr[8]; +}; +struct OuterS { + tint_array_wrapper a1; + tint_array_wrapper a2; +}; + +kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) { + InnerS v = {}; + OuterS s1 = {}; + s1.a1.arr[(*(tint_symbol_1)).i] = v; + s1.a2.arr[(*(tint_symbol_1)).i] = v; + return; +} + diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_multiple_arrays.wgsl.expected.spvasm b/test/statements/assign/indexed_assign_to_array_in_struct/struct_multiple_arrays.wgsl.expected.spvasm new file mode 100644 index 0000000000..7f0b5d5469 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_multiple_arrays.wgsl.expected.spvasm @@ -0,0 +1,63 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 31 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "i" + OpName %uniforms "uniforms" + OpName %main "main" + OpName %InnerS "InnerS" + OpMemberName %InnerS 0 "v" + OpName %v "v" + OpName %OuterS "OuterS" + OpMemberName %OuterS 0 "a1" + OpMemberName %OuterS 1 "a2" + OpName %s1 "s1" + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpDecorate %uniforms NonWritable + OpDecorate %uniforms DescriptorSet 1 + OpDecorate %uniforms Binding 4 + OpMemberDecorate %InnerS 0 Offset 0 + OpMemberDecorate %OuterS 0 Offset 0 + OpDecorate %_arr_InnerS_uint_8 ArrayStride 4 + OpMemberDecorate %OuterS 1 Offset 32 + %uint = OpTypeInt 32 0 + %Uniforms = OpTypeStruct %uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %void = OpTypeVoid + %5 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %InnerS = OpTypeStruct %int +%_ptr_Function_InnerS = OpTypePointer Function %InnerS + %13 = OpConstantNull %InnerS + %uint_8 = OpConstant %uint 8 +%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8 + %OuterS = OpTypeStruct %_arr_InnerS_uint_8 %_arr_InnerS_uint_8 +%_ptr_Function_OuterS = OpTypePointer Function %OuterS + %19 = OpConstantNull %OuterS + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint + %uint_1 = OpConstant %uint 1 + %main = OpFunction %void None %5 + %8 = OpLabel + %v = OpVariable %_ptr_Function_InnerS Function %13 + %s1 = OpVariable %_ptr_Function_OuterS Function %19 + %22 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %23 = OpLoad %uint %22 + %24 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %23 + %25 = OpLoad %InnerS %v + OpStore %24 %25 + %27 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %28 = OpLoad %uint %27 + %29 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_1 %28 + %30 = OpLoad %InnerS %v + OpStore %29 %30 + OpReturn + OpFunctionEnd diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_multiple_arrays.wgsl.expected.wgsl b/test/statements/assign/indexed_assign_to_array_in_struct/struct_multiple_arrays.wgsl.expected.wgsl new file mode 100644 index 0000000000..dcf15424f0 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_multiple_arrays.wgsl.expected.wgsl @@ -0,0 +1,22 @@ +struct Uniforms { + i : u32; +}; + +struct InnerS { + v : i32; +}; + +struct OuterS { + a1 : array; + a2 : array; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +[[stage(compute), workgroup_size(1)]] +fn main() { + var v : InnerS; + var s1 : OuterS; + s1.a1[uniforms.i] = v; + s1.a2[uniforms.i] = v; +} diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_struct_array.wgsl.expected.msl b/test/statements/assign/indexed_assign_to_array_in_struct/struct_struct_array.wgsl.expected.msl new file mode 100644 index 0000000000..9e85490823 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_struct_array.wgsl.expected.msl @@ -0,0 +1,26 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint i; +}; +struct InnerS { + int v; +}; +struct tint_array_wrapper { + InnerS arr[8]; +}; +struct S1 { + tint_array_wrapper a; +}; +struct OuterS { + S1 s2; +}; + +kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) { + InnerS v = {}; + OuterS s1 = {}; + s1.s2.a.arr[(*(tint_symbol_1)).i] = v; + return; +} + diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_struct_array.wgsl.expected.spvasm b/test/statements/assign/indexed_assign_to_array_in_struct/struct_struct_array.wgsl.expected.spvasm new file mode 100644 index 0000000000..9c0f206172 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_struct_array.wgsl.expected.spvasm @@ -0,0 +1,59 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 27 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "i" + OpName %uniforms "uniforms" + OpName %main "main" + OpName %InnerS "InnerS" + OpMemberName %InnerS 0 "v" + OpName %v "v" + OpName %OuterS "OuterS" + OpMemberName %OuterS 0 "s2" + OpName %S1 "S1" + OpMemberName %S1 0 "a" + OpName %s1 "s1" + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpDecorate %uniforms NonWritable + OpDecorate %uniforms DescriptorSet 1 + OpDecorate %uniforms Binding 4 + OpMemberDecorate %InnerS 0 Offset 0 + OpMemberDecorate %OuterS 0 Offset 0 + OpMemberDecorate %S1 0 Offset 0 + OpDecorate %_arr_InnerS_uint_8 ArrayStride 4 + %uint = OpTypeInt 32 0 + %Uniforms = OpTypeStruct %uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %void = OpTypeVoid + %5 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %InnerS = OpTypeStruct %int +%_ptr_Function_InnerS = OpTypePointer Function %InnerS + %13 = OpConstantNull %InnerS + %uint_8 = OpConstant %uint 8 +%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8 + %S1 = OpTypeStruct %_arr_InnerS_uint_8 + %OuterS = OpTypeStruct %S1 +%_ptr_Function_OuterS = OpTypePointer Function %OuterS + %20 = OpConstantNull %OuterS + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint + %main = OpFunction %void None %5 + %8 = OpLabel + %v = OpVariable %_ptr_Function_InnerS Function %13 + %s1 = OpVariable %_ptr_Function_OuterS Function %20 + %23 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %24 = OpLoad %uint %23 + %25 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %uint_0 %24 + %26 = OpLoad %InnerS %v + OpStore %25 %26 + OpReturn + OpFunctionEnd diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_struct_array.wgsl.expected.wgsl b/test/statements/assign/indexed_assign_to_array_in_struct/struct_struct_array.wgsl.expected.wgsl new file mode 100644 index 0000000000..daa4a124c8 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_struct_array.wgsl.expected.wgsl @@ -0,0 +1,24 @@ +struct Uniforms { + i : u32; +}; + +struct InnerS { + v : i32; +}; + +struct S1 { + a : array; +}; + +struct OuterS { + s2 : S1; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +[[stage(compute), workgroup_size(1)]] +fn main() { + var v : InnerS; + var s1 : OuterS; + s1.s2.a[uniforms.i] = v; +} diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_vector.wgsl.expected.msl b/test/statements/assign/indexed_assign_to_array_in_struct/struct_vector.wgsl.expected.msl new file mode 100644 index 0000000000..4e61683afc --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_vector.wgsl.expected.msl @@ -0,0 +1,16 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint i; +}; +struct OuterS { + float3 v1; +}; + +kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) { + OuterS s1 = {}; + s1.v1[(*(tint_symbol_1)).i] = 1.0f; + return; +} + diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_vector.wgsl.expected.spvasm b/test/statements/assign/indexed_assign_to_array_in_struct/struct_vector.wgsl.expected.spvasm new file mode 100644 index 0000000000..677b550a19 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_vector.wgsl.expected.spvasm @@ -0,0 +1,46 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 22 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "i" + OpName %uniforms "uniforms" + OpName %main "main" + OpName %OuterS "OuterS" + OpMemberName %OuterS 0 "v1" + OpName %s1 "s1" + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpDecorate %uniforms NonWritable + OpDecorate %uniforms DescriptorSet 1 + OpDecorate %uniforms Binding 4 + OpMemberDecorate %OuterS 0 Offset 0 + %uint = OpTypeInt 32 0 + %Uniforms = OpTypeStruct %uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %void = OpTypeVoid + %5 = OpTypeFunction %void + %float = OpTypeFloat 32 + %v3float = OpTypeVector %float 3 + %OuterS = OpTypeStruct %v3float +%_ptr_Function_OuterS = OpTypePointer Function %OuterS + %14 = OpConstantNull %OuterS + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint +%_ptr_Function_float = OpTypePointer Function %float + %float_1 = OpConstant %float 1 + %main = OpFunction %void None %5 + %8 = OpLabel + %s1 = OpVariable %_ptr_Function_OuterS Function %14 + %17 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %18 = OpLoad %uint %17 + %20 = OpAccessChain %_ptr_Function_float %s1 %uint_0 %18 + OpStore %20 %float_1 + OpReturn + OpFunctionEnd diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/struct_vector.wgsl.expected.wgsl b/test/statements/assign/indexed_assign_to_array_in_struct/struct_vector.wgsl.expected.wgsl new file mode 100644 index 0000000000..f478089ccd --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/struct_vector.wgsl.expected.wgsl @@ -0,0 +1,15 @@ +struct Uniforms { + i : u32; +}; + +struct OuterS { + v1 : vec3; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +[[stage(compute), workgroup_size(1)]] +fn main() { + var s1 : OuterS; + s1.v1[uniforms.i] = 1.0; +} diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/vector_assign.wgsl.expected.msl b/test/statements/assign/indexed_assign_to_array_in_struct/vector_assign.wgsl.expected.msl new file mode 100644 index 0000000000..584841e04b --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/vector_assign.wgsl.expected.msl @@ -0,0 +1,25 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint i; +}; +struct tint_array_wrapper { + uint arr[8]; +}; +struct OuterS { + tint_array_wrapper a1; +}; + +uint f(uint i) { + return (i + 1u); +} + +kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) { + OuterS s1 = {}; + float3 v = 0.0f; + v[s1.a1.arr[(*(tint_symbol_1)).i]] = 1.0f; + v[f(s1.a1.arr[(*(tint_symbol_1)).i])] = 1.0f; + return; +} + diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/vector_assign.wgsl.expected.spvasm b/test/statements/assign/indexed_assign_to_array_in_struct/vector_assign.wgsl.expected.spvasm new file mode 100644 index 0000000000..329d21a6b5 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/vector_assign.wgsl.expected.spvasm @@ -0,0 +1,73 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 42 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "i" + OpName %uniforms "uniforms" + OpName %f "f" + OpName %i "i" + OpName %main "main" + OpName %OuterS "OuterS" + OpMemberName %OuterS 0 "a1" + OpName %s1 "s1" + OpName %v "v" + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpDecorate %uniforms NonWritable + OpDecorate %uniforms DescriptorSet 1 + OpDecorate %uniforms Binding 4 + OpMemberDecorate %OuterS 0 Offset 0 + OpDecorate %_arr_uint_uint_8 ArrayStride 4 + %uint = OpTypeInt 32 0 + %Uniforms = OpTypeStruct %uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %5 = OpTypeFunction %uint %uint + %uint_1 = OpConstant %uint 1 + %void = OpTypeVoid + %11 = OpTypeFunction %void + %uint_8 = OpConstant %uint 8 +%_arr_uint_uint_8 = OpTypeArray %uint %uint_8 + %OuterS = OpTypeStruct %_arr_uint_uint_8 +%_ptr_Function_OuterS = OpTypePointer Function %OuterS + %20 = OpConstantNull %OuterS + %float = OpTypeFloat 32 + %v3float = OpTypeVector %float 3 +%_ptr_Function_v3float = OpTypePointer Function %v3float + %25 = OpConstantNull %v3float + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_Function_float = OpTypePointer Function %float + %float_1 = OpConstant %float 1 + %f = OpFunction %uint None %5 + %i = OpFunctionParameter %uint + %8 = OpLabel + %10 = OpIAdd %uint %i %uint_1 + OpReturnValue %10 + OpFunctionEnd + %main = OpFunction %void None %11 + %14 = OpLabel + %s1 = OpVariable %_ptr_Function_OuterS Function %20 + %v = OpVariable %_ptr_Function_v3float Function %25 + %28 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %29 = OpLoad %uint %28 + %31 = OpAccessChain %_ptr_Function_uint %s1 %uint_0 %29 + %32 = OpLoad %uint %31 + %34 = OpAccessChain %_ptr_Function_float %v %32 + OpStore %34 %float_1 + %37 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %38 = OpLoad %uint %37 + %39 = OpAccessChain %_ptr_Function_uint %s1 %uint_0 %38 + %40 = OpLoad %uint %39 + %36 = OpFunctionCall %uint %f %40 + %41 = OpAccessChain %_ptr_Function_float %v %36 + OpStore %41 %float_1 + OpReturn + OpFunctionEnd diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/vector_assign.wgsl.expected.wgsl b/test/statements/assign/indexed_assign_to_array_in_struct/vector_assign.wgsl.expected.wgsl new file mode 100644 index 0000000000..75951e57bb --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/vector_assign.wgsl.expected.wgsl @@ -0,0 +1,21 @@ +struct Uniforms { + i : u32; +}; + +struct OuterS { + a1 : array; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +fn f(i : u32) -> u32 { + return (i + 1u); +} + +[[stage(compute), workgroup_size(1)]] +fn main() { + var s1 : OuterS; + var v : vec3; + v[s1.a1[uniforms.i]] = 1.0; + v[f(s1.a1[uniforms.i])] = 1.0; +} diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/via_pointer.wgsl.expected.msl b/test/statements/assign/indexed_assign_to_array_in_struct/via_pointer.wgsl.expected.msl new file mode 100644 index 0000000000..d42fec8a95 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/via_pointer.wgsl.expected.msl @@ -0,0 +1,24 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint i; +}; +struct InnerS { + int v; +}; +struct tint_array_wrapper { + InnerS arr[8]; +}; +struct OuterS { + tint_array_wrapper a1; +}; + +kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) { + InnerS v = {}; + OuterS s1 = {}; + uint const p_save = (*(tint_symbol_1)).i; + s1.a1.arr[p_save] = v; + return; +} + diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/via_pointer.wgsl.expected.spvasm b/test/statements/assign/indexed_assign_to_array_in_struct/via_pointer.wgsl.expected.spvasm new file mode 100644 index 0000000000..8507c70c66 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/via_pointer.wgsl.expected.spvasm @@ -0,0 +1,55 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 26 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "i" + OpName %uniforms "uniforms" + OpName %main "main" + OpName %InnerS "InnerS" + OpMemberName %InnerS 0 "v" + OpName %v "v" + OpName %OuterS "OuterS" + OpMemberName %OuterS 0 "a1" + OpName %s1 "s1" + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpDecorate %uniforms NonWritable + OpDecorate %uniforms DescriptorSet 1 + OpDecorate %uniforms Binding 4 + OpMemberDecorate %InnerS 0 Offset 0 + OpMemberDecorate %OuterS 0 Offset 0 + OpDecorate %_arr_InnerS_uint_8 ArrayStride 4 + %uint = OpTypeInt 32 0 + %Uniforms = OpTypeStruct %uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %void = OpTypeVoid + %5 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %InnerS = OpTypeStruct %int +%_ptr_Function_InnerS = OpTypePointer Function %InnerS + %13 = OpConstantNull %InnerS + %uint_8 = OpConstant %uint 8 +%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8 + %OuterS = OpTypeStruct %_arr_InnerS_uint_8 +%_ptr_Function_OuterS = OpTypePointer Function %OuterS + %19 = OpConstantNull %OuterS + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint + %main = OpFunction %void None %5 + %8 = OpLabel + %v = OpVariable %_ptr_Function_InnerS Function %13 + %s1 = OpVariable %_ptr_Function_OuterS Function %19 + %22 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %23 = OpLoad %uint %22 + %24 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %23 + %25 = OpLoad %InnerS %v + OpStore %24 %25 + OpReturn + OpFunctionEnd diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/via_pointer.wgsl.expected.wgsl b/test/statements/assign/indexed_assign_to_array_in_struct/via_pointer.wgsl.expected.wgsl new file mode 100644 index 0000000000..3752e4a2d3 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/via_pointer.wgsl.expected.wgsl @@ -0,0 +1,21 @@ +struct Uniforms { + i : u32; +}; + +struct InnerS { + v : i32; +}; + +struct OuterS { + a1 : array; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +[[stage(compute), workgroup_size(1)]] +fn main() { + var v : InnerS; + var s1 : OuterS; + let p = &(s1.a1[uniforms.i]); + *(p) = v; +} diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/via_pointer_arg.wgsl.expected.msl b/test/statements/assign/indexed_assign_to_array_in_struct/via_pointer_arg.wgsl.expected.msl new file mode 100644 index 0000000000..d6c3a35c6a --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/via_pointer_arg.wgsl.expected.msl @@ -0,0 +1,27 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ uint i; +}; +struct InnerS { + int v; +}; +struct tint_array_wrapper { + InnerS arr[8]; +}; +struct OuterS { + tint_array_wrapper a1; +}; + +void f(thread OuterS* const p, const constant Uniforms* const tint_symbol_1) { + InnerS v = {}; + (*(p)).a1.arr[(*(tint_symbol_1)).i] = v; +} + +kernel void tint_symbol(const constant Uniforms* tint_symbol_2 [[buffer(0)]]) { + OuterS s1 = {}; + f(&(s1), tint_symbol_2); + return; +} + diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/via_pointer_arg.wgsl.expected.spvasm b/test/statements/assign/indexed_assign_to_array_in_struct/via_pointer_arg.wgsl.expected.spvasm new file mode 100644 index 0000000000..bfda721ac7 --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/via_pointer_arg.wgsl.expected.spvasm @@ -0,0 +1,64 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 33 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "i" + OpName %uniforms "uniforms" + OpName %OuterS "OuterS" + OpMemberName %OuterS 0 "a1" + OpName %InnerS "InnerS" + OpMemberName %InnerS 0 "v" + OpName %f "f" + OpName %p "p" + OpName %v "v" + OpName %main "main" + OpName %s1 "s1" + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpDecorate %uniforms NonWritable + OpDecorate %uniforms DescriptorSet 1 + OpDecorate %uniforms Binding 4 + OpMemberDecorate %OuterS 0 Offset 0 + OpMemberDecorate %InnerS 0 Offset 0 + OpDecorate %_arr_InnerS_uint_8 ArrayStride 4 + %uint = OpTypeInt 32 0 + %Uniforms = OpTypeStruct %uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %void = OpTypeVoid + %int = OpTypeInt 32 1 + %InnerS = OpTypeStruct %int + %uint_8 = OpConstant %uint 8 +%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8 + %OuterS = OpTypeStruct %_arr_InnerS_uint_8 +%_ptr_Function_OuterS = OpTypePointer Function %OuterS + %5 = OpTypeFunction %void %_ptr_Function_OuterS +%_ptr_Function_InnerS = OpTypePointer Function %InnerS + %18 = OpConstantNull %InnerS + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint + %26 = OpTypeFunction %void + %30 = OpConstantNull %OuterS + %f = OpFunction %void None %5 + %p = OpFunctionParameter %_ptr_Function_OuterS + %15 = OpLabel + %v = OpVariable %_ptr_Function_InnerS Function %18 + %22 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 + %23 = OpLoad %uint %22 + %24 = OpAccessChain %_ptr_Function_InnerS %p %uint_0 %23 + %25 = OpLoad %InnerS %v + OpStore %24 %25 + OpReturn + OpFunctionEnd + %main = OpFunction %void None %26 + %28 = OpLabel + %s1 = OpVariable %_ptr_Function_OuterS Function %30 + %31 = OpFunctionCall %void %f %s1 + OpReturn + OpFunctionEnd diff --git a/test/statements/assign/indexed_assign_to_array_in_struct/via_pointer_arg.wgsl.expected.wgsl b/test/statements/assign/indexed_assign_to_array_in_struct/via_pointer_arg.wgsl.expected.wgsl new file mode 100644 index 0000000000..9a80febe0d --- /dev/null +++ b/test/statements/assign/indexed_assign_to_array_in_struct/via_pointer_arg.wgsl.expected.wgsl @@ -0,0 +1,24 @@ +struct Uniforms { + i : u32; +}; + +struct InnerS { + v : i32; +}; + +struct OuterS { + a1 : array; +}; + +[[group(1), binding(4)]] var uniforms : Uniforms; + +fn f(p : ptr) { + var v : InnerS; + (*(p)).a1[uniforms.i] = v; +} + +[[stage(compute), workgroup_size(1)]] +fn main() { + var s1 : OuterS; + f(&(s1)); +}