test: Generate expected output for all tests

The expected output is far from perfect, and the generated HLSL and MSL
isn't even validated yet, so may be incorrect.

However, by committing the generated output, we get clear examples of
the currently generated output of each backend. As we land fixes and
improvements to each backend, the presubmits will require us to update
the expected test output, and so code reviews will include diffs of
each backend's generated output.

Change-Id: I5c2a9e5b796d0ab75b3ec4c7f8ad00a0a2ab166f
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/51224
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Reviewed-by: David Neto <dneto@google.com>

test/bug/tint/782.wgsl.expected.hlsl

@@ -0,0 +1,12 @@
+
+void foo() {
+  int tint_symbol[2] = {0, 0};
+  int implict[2] = {0, 0};
+  implict = tint_symbol;
+}
+
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+  return;
+}
+

test/bug/tint/782.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+void foo() {
+  int tint_symbol[2] = {0};
+  int implict[2] = {0};
+  implict = tint_symbol;
+}
+

test/bug/tint/782.wgsl.expected.spvasm

@@ -0,0 +1,34 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; 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 %foo "foo"
+               OpName %explicit "explicit"
+               OpName %implict "implict"
+               OpDecorate %_arr_int_uint_2 ArrayStride 4
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+       %uint = OpTypeInt 32 0
+     %uint_2 = OpConstant %uint 2
+%_arr_int_uint_2 = OpTypeArray %int %uint_2
+%_ptr_Function__arr_int_uint_2 = OpTypePointer Function %_arr_int_uint_2
+         %13 = OpConstantNull %_arr_int_uint_2
+%unused_entry_point = OpFunction %void None %1
+          %4 = OpLabel
+               OpReturn
+               OpFunctionEnd
+        %foo = OpFunction %void None %1
+          %6 = OpLabel
+   %explicit = OpVariable %_ptr_Function__arr_int_uint_2 Function %13
+    %implict = OpVariable %_ptr_Function__arr_int_uint_2 Function %13
+         %15 = OpLoad %_arr_int_uint_2 %explicit
+               OpStore %implict %15
+               OpReturn
+               OpFunctionEnd

test/bug/tint/782.wgsl.expected.wgsl

@@ -0,0 +1,9 @@
+type ArrayExplicitStride = [[stride(4)]] array<i32, 2>;
+
+type ArrayImplicitStride = array<i32, 2>;
+
+fn foo() {
+  var explicit : ArrayExplicitStride;
+  var implict : ArrayImplicitStride;
+  implict = explicit;
+}

test/samples/compute_boids.wgsl.expected.hlsl

@@ -0,0 +1,118 @@
+struct tint_symbol_1 {
+  float2 a_particlePos : TEXCOORD0;
+  float2 a_particleVel : TEXCOORD1;
+  float2 a_pos : TEXCOORD2;
+};
+struct tint_symbol_2 {
+  float4 value : SV_Position;
+};
+struct tint_symbol_3 {
+  float4 value : SV_Target0;
+};
+struct SimParams {
+  float deltaT;
+  float rule1Distance;
+  float rule2Distance;
+  float rule3Distance;
+  float rule1Scale;
+  float rule2Scale;
+  float rule3Scale;
+};
+struct tint_symbol_5 {
+  uint3 gl_GlobalInvocationID : SV_DispatchThreadID;
+};
+
+ConstantBuffer<SimParams> params : register(b0, space0);
+
+RWByteAddressBuffer particlesA : register(u1, space0);
+RWByteAddressBuffer particlesB : register(u2, space0);
+
+tint_symbol_2 vert_main(tint_symbol_1 tint_symbol) {
+  const float2 a_particlePos = tint_symbol.a_particlePos;
+  const float2 a_particleVel = tint_symbol.a_particleVel;
+  const float2 a_pos = tint_symbol.a_pos;
+  float angle = -(  atan2(a_particleVel.x, a_particleVel.y));
+  float2 pos = float2(((a_pos.x *   cos(angle)) - (a_pos.y *   sin(angle))), ((a_pos.x *   sin(angle)) + (a_pos.y *   cos(angle))));
+  const tint_symbol_2 tint_symbol_8 = {float4((pos + a_particlePos), 0.0f, 1.0f)};
+  return tint_symbol_8;
+}
+
+tint_symbol_3 frag_main() {
+  const tint_symbol_3 tint_symbol_9 = {float4(1.0f, 1.0f, 1.0f, 1.0f)};
+  return tint_symbol_9;
+}
+
+[numthreads(1, 1, 1)]
+void comp_main(tint_symbol_5 tint_symbol_4) {
+  const uint3 gl_GlobalInvocationID = tint_symbol_4.gl_GlobalInvocationID;
+  uint index = gl_GlobalInvocationID.x;
+  if ((index >= 5u)) {
+    return;
+  }
+  float2 vPos = asfloat(particlesA.Load2((16u * index)));
+  float2 vVel = asfloat(particlesA.Load2(((16u * index) + 8u)));
+  float2 cMass = float2(0.0f, 0.0f);
+  float2 cVel = float2(0.0f, 0.0f);
+  float2 colVel = float2(0.0f, 0.0f);
+  int cMassCount = 0;
+  int cVelCount = 0;
+  float2 pos = float2(0.0f, 0.0f);
+  float2 vel = float2(0.0f, 0.0f);
+  {
+    uint i = 0u;
+    {
+      bool tint_hlsl_is_first_1 = true;
+      for(;;) {
+        if (!tint_hlsl_is_first_1) {
+          i = (i + 1u);
+        }
+        tint_hlsl_is_first_1 = false;
+
+        if (!((i < 5u))) {
+          break;
+        }
+        if ((i == index)) {
+          continue;
+        }
+        pos = asfloat(particlesA.Load2((16u * i))).xy;
+        vel = asfloat(particlesA.Load2(((16u * i) + 8u))).xy;
+        if ((        distance(pos, vPos) < params.rule1Distance)) {
+          cMass = (cMass + pos);
+          cMassCount = (cMassCount + 1);
+        }
+        if ((        distance(pos, vPos) < params.rule2Distance)) {
+          colVel = (colVel - (pos - vPos));
+        }
+        if ((        distance(pos, vPos) < params.rule3Distance)) {
+          cVel = (cVel + vel);
+          cVelCount = (cVelCount + 1);
+        }
+      }
+    }
+  }
+  if ((cMassCount > 0)) {
+    cMass = ((cMass / float2(float(cMassCount), float(cMassCount))) - vPos);
+  }
+  if ((cVelCount > 0)) {
+    cVel = (cVel / float2(float(cVelCount), float(cVelCount)));
+  }
+  vVel = (((vVel + (cMass * params.rule1Scale)) + (colVel * params.rule2Scale)) + (cVel * params.rule3Scale));
+  vVel = (  normalize(vVel) *   clamp(  length(vVel), 0.0f, 0.100000001f));
+  vPos = (vPos + (vVel * params.deltaT));
+  if ((vPos.x < -1.0f)) {
+    vPos.x = 1.0f;
+  }
+  if ((vPos.x > 1.0f)) {
+    vPos.x = -1.0f;
+  }
+  if ((vPos.y < -1.0f)) {
+    vPos.y = 1.0f;
+  }
+  if ((vPos.y > 1.0f)) {
+    vPos.y = -1.0f;
+  }
+  particlesB.Store2((16u * index), asuint(vPos));
+  particlesB.Store2(((16u * index) + 8u), asuint(vVel));
+  return;
+}
+

test/samples/compute_boids.wgsl.expected.msl

@@ -0,0 +1,120 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_1 {
+  float2 a_particlePos [[attribute(0)]];
+  float2 a_particleVel [[attribute(1)]];
+  float2 a_pos [[attribute(2)]];
+};
+struct tint_symbol_2 {
+  float4 value [[position]];
+};
+struct tint_symbol_3 {
+  float4 value [[color(0)]];
+};
+struct Particle {
+  /* 0x0000 */ packed_float2 pos;
+  /* 0x0008 */ packed_float2 vel;
+};
+struct SimParams {
+  /* 0x0000 */ float deltaT;
+  /* 0x0004 */ float rule1Distance;
+  /* 0x0008 */ float rule2Distance;
+  /* 0x000c */ float rule3Distance;
+  /* 0x0010 */ float rule1Scale;
+  /* 0x0014 */ float rule2Scale;
+  /* 0x0018 */ float rule3Scale;
+};
+struct Particles {
+  /* 0x0000 */ Particle particles[5];
+};
+struct tint_symbol_5 {
+  uint3 gl_GlobalInvocationID [[thread_position_in_grid]];
+};
+
+vertex tint_symbol_2 vert_main(tint_symbol_1 tint_symbol [[stage_in]]) {
+  const float2 a_particlePos = tint_symbol.a_particlePos;
+  const float2 a_particleVel = tint_symbol.a_particleVel;
+  const float2 a_pos = tint_symbol.a_pos;
+  float angle = -(  atan2(a_particleVel.x, a_particleVel.y));
+  float2 pos = float2(((a_pos.x *   cos(angle)) - (a_pos.y *   sin(angle))), ((a_pos.x *   sin(angle)) + (a_pos.y *   cos(angle))));
+  return {float4((pos + a_particlePos), 0.0f, 1.0f)};
+}
+
+fragment tint_symbol_3 frag_main() {
+  return {float4(1.0f, 1.0f, 1.0f, 1.0f)};
+}
+
+kernel void comp_main(tint_symbol_5 tint_symbol_4 [[stage_in]], constant SimParams& params [[buffer(0)]], device Particles& particlesA [[buffer(1)]], device Particles& particlesB [[buffer(2)]]) {
+  const uint3 gl_GlobalInvocationID = tint_symbol_4.gl_GlobalInvocationID;
+  uint index = gl_GlobalInvocationID.x;
+  if ((index >= 5u)) {
+    return;
+  }
+  float2 vPos = particlesA.particles[index].pos;
+  float2 vVel = particlesA.particles[index].vel;
+  float2 cMass = float2(0.0f, 0.0f);
+  float2 cVel = float2(0.0f, 0.0f);
+  float2 colVel = float2(0.0f, 0.0f);
+  int cMassCount = 0;
+  int cVelCount = 0;
+  float2 pos = 0.0f;
+  float2 vel = 0.0f;
+  {
+    uint i = 0u;
+    {
+      bool tint_msl_is_first_1 = true;
+      for(;;) {
+        if (!tint_msl_is_first_1) {
+          i = (i + 1u);
+        }
+        tint_msl_is_first_1 = false;
+
+        if (!((i < 5u))) {
+          break;
+        }
+        if ((i == index)) {
+          continue;
+        }
+        pos = particlesA.particles[i].pos.xy;
+        vel = particlesA.particles[i].vel.xy;
+        if ((        distance(pos, vPos) < params.rule1Distance)) {
+          cMass = (cMass + pos);
+          cMassCount = (cMassCount + 1);
+        }
+        if ((        distance(pos, vPos) < params.rule2Distance)) {
+          colVel = (colVel - (pos - vPos));
+        }
+        if ((        distance(pos, vPos) < params.rule3Distance)) {
+          cVel = (cVel + vel);
+          cVelCount = (cVelCount + 1);
+        }
+      }
+    }
+  }
+  if ((cMassCount > 0)) {
+    cMass = ((cMass / float2(float(cMassCount), float(cMassCount))) - vPos);
+  }
+  if ((cVelCount > 0)) {
+    cVel = (cVel / float2(float(cVelCount), float(cVelCount)));
+  }
+  vVel = (((vVel + (cMass * params.rule1Scale)) + (colVel * params.rule2Scale)) + (cVel * params.rule3Scale));
+  vVel = (  normalize(vVel) *   clamp(  length(vVel), 0.0f, 0.100000001f));
+  vPos = (vPos + (vVel * params.deltaT));
+  if ((vPos.x < -1.0f)) {
+    vPos.x = 1.0f;
+  }
+  if ((vPos.x > 1.0f)) {
+    vPos.x = -1.0f;
+  }
+  if ((vPos.y < -1.0f)) {
+    vPos.y = 1.0f;
+  }
+  if ((vPos.y > 1.0f)) {
+    vPos.y = -1.0f;
+  }
+  particlesB.particles[index].pos = vPos;
+  particlesB.particles[index].vel = vVel;
+  return;
+}
+

test/samples/compute_boids.wgsl.expected.spvasm

@@ -0,0 +1,443 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 279
+; Schema: 0
+               OpCapability Shader
+         %40 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vert_main "vert_main" %tint_pointsize %tint_symbol_1 %tint_symbol_2 %tint_symbol %tint_symbol_4
+               OpEntryPoint Fragment %frag_main "frag_main" %tint_symbol_7
+               OpEntryPoint GLCompute %comp_main "comp_main" %tint_symbol_9
+               OpExecutionMode %frag_main OriginUpperLeft
+               OpExecutionMode %comp_main LocalSize 1 1 1
+               OpName %tint_pointsize "tint_pointsize"
+               OpName %tint_symbol "tint_symbol"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %tint_symbol_2 "tint_symbol_2"
+               OpName %tint_symbol_4 "tint_symbol_4"
+               OpName %tint_symbol_7 "tint_symbol_7"
+               OpName %SimParams "SimParams"
+               OpMemberName %SimParams 0 "deltaT"
+               OpMemberName %SimParams 1 "rule1Distance"
+               OpMemberName %SimParams 2 "rule2Distance"
+               OpMemberName %SimParams 3 "rule3Distance"
+               OpMemberName %SimParams 4 "rule1Scale"
+               OpMemberName %SimParams 5 "rule2Scale"
+               OpMemberName %SimParams 6 "rule3Scale"
+               OpName %params "params"
+               OpName %Particles "Particles"
+               OpMemberName %Particles 0 "particles"
+               OpName %Particle "Particle"
+               OpMemberName %Particle 0 "pos"
+               OpMemberName %Particle 1 "vel"
+               OpName %particlesA "particlesA"
+               OpName %particlesB "particlesB"
+               OpName %tint_symbol_9 "tint_symbol_9"
+               OpName %tint_symbol_5 "tint_symbol_5"
+               OpName %tint_symbol_3 "tint_symbol_3"
+               OpName %vert_main "vert_main"
+               OpName %angle "angle"
+               OpName %pos "pos"
+               OpName %tint_symbol_8 "tint_symbol_8"
+               OpName %tint_symbol_6 "tint_symbol_6"
+               OpName %frag_main "frag_main"
+               OpName %comp_main "comp_main"
+               OpName %index "index"
+               OpName %vPos "vPos"
+               OpName %vVel "vVel"
+               OpName %cMass "cMass"
+               OpName %cVel "cVel"
+               OpName %colVel "colVel"
+               OpName %cMassCount "cMassCount"
+               OpName %cVelCount "cVelCount"
+               OpName %pos_0 "pos"
+               OpName %vel "vel"
+               OpName %i "i"
+               OpDecorate %tint_pointsize BuiltIn PointSize
+               OpDecorate %tint_symbol Location 0
+               OpDecorate %tint_symbol_1 Location 1
+               OpDecorate %tint_symbol_2 Location 2
+               OpDecorate %tint_symbol_4 BuiltIn Position
+               OpDecorate %tint_symbol_7 Location 0
+               OpDecorate %SimParams Block
+               OpMemberDecorate %SimParams 0 Offset 0
+               OpMemberDecorate %SimParams 1 Offset 4
+               OpMemberDecorate %SimParams 2 Offset 8
+               OpMemberDecorate %SimParams 3 Offset 12
+               OpMemberDecorate %SimParams 4 Offset 16
+               OpMemberDecorate %SimParams 5 Offset 20
+               OpMemberDecorate %SimParams 6 Offset 24
+               OpDecorate %params Binding 0
+               OpDecorate %params DescriptorSet 0
+               OpDecorate %Particles Block
+               OpMemberDecorate %Particles 0 Offset 0
+               OpMemberDecorate %Particle 0 Offset 0
+               OpMemberDecorate %Particle 1 Offset 8
+               OpDecorate %_arr_Particle_uint_5 ArrayStride 16
+               OpDecorate %particlesA Binding 1
+               OpDecorate %particlesA DescriptorSet 0
+               OpDecorate %particlesB Binding 2
+               OpDecorate %particlesB DescriptorSet 0
+               OpDecorate %tint_symbol_9 BuiltIn GlobalInvocationId
+      %float = OpTypeFloat 32
+%_ptr_Output_float = OpTypePointer Output %float
+          %4 = OpConstantNull %float
+%tint_pointsize = OpVariable %_ptr_Output_float Output %4
+    %v2float = OpTypeVector %float 2
+%_ptr_Input_v2float = OpTypePointer Input %v2float
+%tint_symbol = OpVariable %_ptr_Input_v2float Input
+%tint_symbol_1 = OpVariable %_ptr_Input_v2float Input
+%tint_symbol_2 = OpVariable %_ptr_Input_v2float Input
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+         %13 = OpConstantNull %v4float
+%tint_symbol_4 = OpVariable %_ptr_Output_v4float Output %13
+%tint_symbol_7 = OpVariable %_ptr_Output_v4float Output %13
+  %SimParams = OpTypeStruct %float %float %float %float %float %float %float
+%_ptr_Uniform_SimParams = OpTypePointer Uniform %SimParams
+     %params = OpVariable %_ptr_Uniform_SimParams Uniform
+   %Particle = OpTypeStruct %v2float %v2float
+       %uint = OpTypeInt 32 0
+     %uint_5 = OpConstant %uint 5
+%_arr_Particle_uint_5 = OpTypeArray %Particle %uint_5
+  %Particles = OpTypeStruct %_arr_Particle_uint_5
+%_ptr_StorageBuffer_Particles = OpTypePointer StorageBuffer %Particles
+ %particlesA = OpVariable %_ptr_StorageBuffer_Particles StorageBuffer
+ %particlesB = OpVariable %_ptr_StorageBuffer_Particles StorageBuffer
+     %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%tint_symbol_9 = OpVariable %_ptr_Input_v3uint Input
+       %void = OpTypeVoid
+         %29 = OpTypeFunction %void %v4float
+         %34 = OpTypeFunction %void
+    %float_1 = OpConstant %float 1
+     %uint_0 = OpConstant %uint 0
+%_ptr_Input_float = OpTypePointer Input %float
+     %uint_1 = OpConstant %uint 1
+%_ptr_Function_float = OpTypePointer Function %float
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+         %75 = OpConstantNull %v2float
+    %float_0 = OpConstant %float 0
+         %90 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
+%_ptr_Input_uint = OpTypePointer Input %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %98 = OpConstantNull %uint
+       %bool = OpTypeBool
+%_ptr_StorageBuffer_v2float = OpTypePointer StorageBuffer %v2float
+        %113 = OpConstantComposite %v2float %float_0 %float_0
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+%_ptr_Function_int = OpTypePointer Function %int
+        %121 = OpConstantNull %int
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+      %int_1 = OpConstant %int 1
+     %uint_2 = OpConstant %uint 2
+     %uint_3 = OpConstant %uint 3
+     %uint_4 = OpConstant %uint 4
+     %uint_6 = OpConstant %uint 6
+%float_0_100000001 = OpConstant %float 0.100000001
+   %float_n1 = OpConstant %float -1
+%tint_symbol_5 = OpFunction %void None %29
+%tint_symbol_3 = OpFunctionParameter %v4float
+         %33 = OpLabel
+               OpStore %tint_symbol_4 %tint_symbol_3
+               OpReturn
+               OpFunctionEnd
+  %vert_main = OpFunction %void None %34
+         %36 = OpLabel
+      %angle = OpVariable %_ptr_Function_float Function %4
+        %pos = OpVariable %_ptr_Function_v2float Function %75
+               OpStore %tint_pointsize %float_1
+         %43 = OpAccessChain %_ptr_Input_float %tint_symbol_1 %uint_0
+         %44 = OpLoad %float %43
+         %46 = OpAccessChain %_ptr_Input_float %tint_symbol_1 %uint_1
+         %47 = OpLoad %float %46
+         %39 = OpExtInst %float %40 Atan2 %44 %47
+         %38 = OpFNegate %float %39
+               OpStore %angle %38
+         %50 = OpAccessChain %_ptr_Input_float %tint_symbol_2 %uint_0
+         %51 = OpLoad %float %50
+         %53 = OpLoad %float %angle
+         %52 = OpExtInst %float %40 Cos %53
+         %54 = OpFMul %float %51 %52
+         %55 = OpAccessChain %_ptr_Input_float %tint_symbol_2 %uint_1
+         %56 = OpLoad %float %55
+         %58 = OpLoad %float %angle
+         %57 = OpExtInst %float %40 Sin %58
+         %59 = OpFMul %float %56 %57
+         %60 = OpFSub %float %54 %59
+         %61 = OpAccessChain %_ptr_Input_float %tint_symbol_2 %uint_0
+         %62 = OpLoad %float %61
+         %64 = OpLoad %float %angle
+         %63 = OpExtInst %float %40 Sin %64
+         %65 = OpFMul %float %62 %63
+         %66 = OpAccessChain %_ptr_Input_float %tint_symbol_2 %uint_1
+         %67 = OpLoad %float %66
+         %69 = OpLoad %float %angle
+         %68 = OpExtInst %float %40 Cos %69
+         %70 = OpFMul %float %67 %68
+         %71 = OpFAdd %float %65 %70
+         %72 = OpCompositeConstruct %v2float %60 %71
+               OpStore %pos %72
+         %77 = OpLoad %v2float %pos
+         %78 = OpLoad %v2float %tint_symbol
+         %79 = OpFAdd %v2float %77 %78
+         %80 = OpCompositeExtract %float %79 0
+         %81 = OpCompositeExtract %float %79 1
+         %83 = OpCompositeConstruct %v4float %80 %81 %float_0 %float_1
+         %76 = OpFunctionCall %void %tint_symbol_5 %83
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_8 = OpFunction %void None %29
+%tint_symbol_6 = OpFunctionParameter %v4float
+         %86 = OpLabel
+               OpStore %tint_symbol_7 %tint_symbol_6
+               OpReturn
+               OpFunctionEnd
+  %frag_main = OpFunction %void None %34
+         %88 = OpLabel
+         %89 = OpFunctionCall %void %tint_symbol_8 %90
+               OpReturn
+               OpFunctionEnd
+  %comp_main = OpFunction %void None %34
+         %92 = OpLabel
+      %index = OpVariable %_ptr_Function_uint Function %98
+       %vPos = OpVariable %_ptr_Function_v2float Function %75
+       %vVel = OpVariable %_ptr_Function_v2float Function %75
+      %cMass = OpVariable %_ptr_Function_v2float Function %75
+       %cVel = OpVariable %_ptr_Function_v2float Function %75
+     %colVel = OpVariable %_ptr_Function_v2float Function %75
+ %cMassCount = OpVariable %_ptr_Function_int Function %121
+  %cVelCount = OpVariable %_ptr_Function_int Function %121
+      %pos_0 = OpVariable %_ptr_Function_v2float Function %75
+        %vel = OpVariable %_ptr_Function_v2float Function %75
+          %i = OpVariable %_ptr_Function_uint Function %98
+         %94 = OpAccessChain %_ptr_Input_uint %tint_symbol_9 %uint_0
+         %95 = OpLoad %uint %94
+               OpStore %index %95
+         %99 = OpLoad %uint %index
+        %100 = OpUGreaterThanEqual %bool %99 %uint_5
+               OpSelectionMerge %102 None
+               OpBranchConditional %100 %103 %102
+        %103 = OpLabel
+               OpReturn
+        %102 = OpLabel
+        %104 = OpLoad %uint %index
+        %106 = OpAccessChain %_ptr_StorageBuffer_v2float %particlesA %uint_0 %104 %uint_0
+        %107 = OpLoad %v2float %106
+               OpStore %vPos %107
+        %109 = OpLoad %uint %index
+        %110 = OpAccessChain %_ptr_StorageBuffer_v2float %particlesA %uint_0 %109 %uint_1
+        %111 = OpLoad %v2float %110
+               OpStore %vVel %111
+               OpStore %cMass %113
+               OpStore %cVel %113
+               OpStore %colVel %113
+               OpStore %cMassCount %int_0
+               OpStore %cVelCount %int_0
+               OpStore %i %uint_0
+               OpBranch %126
+        %126 = OpLabel
+               OpLoopMerge %127 %128 None
+               OpBranch %129
+        %129 = OpLabel
+        %131 = OpLoad %uint %i
+        %132 = OpULessThan %bool %131 %uint_5
+        %130 = OpLogicalNot %bool %132
+               OpSelectionMerge %133 None
+               OpBranchConditional %130 %134 %133
+        %134 = OpLabel
+               OpBranch %127
+        %133 = OpLabel
+        %135 = OpLoad %uint %i
+        %136 = OpLoad %uint %index
+        %137 = OpIEqual %bool %135 %136
+               OpSelectionMerge %138 None
+               OpBranchConditional %137 %139 %138
+        %139 = OpLabel
+               OpBranch %128
+        %138 = OpLabel
+        %140 = OpLoad %uint %i
+        %141 = OpAccessChain %_ptr_StorageBuffer_v2float %particlesA %uint_0 %140 %uint_0
+        %142 = OpLoad %v2float %141
+        %143 = OpVectorShuffle %v2float %142 %142 0 1
+               OpStore %pos_0 %143
+        %144 = OpLoad %uint %i
+        %145 = OpAccessChain %_ptr_StorageBuffer_v2float %particlesA %uint_0 %144 %uint_1
+        %146 = OpLoad %v2float %145
+        %147 = OpVectorShuffle %v2float %146 %146 0 1
+               OpStore %vel %147
+        %149 = OpLoad %v2float %pos_0
+        %150 = OpLoad %v2float %vPos
+        %148 = OpExtInst %float %40 Distance %149 %150
+        %152 = OpAccessChain %_ptr_Uniform_float %params %uint_1
+        %153 = OpLoad %float %152
+        %154 = OpFOrdLessThan %bool %148 %153
+               OpSelectionMerge %155 None
+               OpBranchConditional %154 %156 %155
+        %156 = OpLabel
+        %157 = OpLoad %v2float %cMass
+        %158 = OpLoad %v2float %pos_0
+        %159 = OpFAdd %v2float %157 %158
+               OpStore %cMass %159
+        %160 = OpLoad %int %cMassCount
+        %162 = OpIAdd %int %160 %int_1
+               OpStore %cMassCount %162
+               OpBranch %155
+        %155 = OpLabel
+        %164 = OpLoad %v2float %pos_0
+        %165 = OpLoad %v2float %vPos
+        %163 = OpExtInst %float %40 Distance %164 %165
+        %167 = OpAccessChain %_ptr_Uniform_float %params %uint_2
+        %168 = OpLoad %float %167
+        %169 = OpFOrdLessThan %bool %163 %168
+               OpSelectionMerge %170 None
+               OpBranchConditional %169 %171 %170
+        %171 = OpLabel
+        %172 = OpLoad %v2float %colVel
+        %173 = OpLoad %v2float %pos_0
+        %174 = OpLoad %v2float %vPos
+        %175 = OpFSub %v2float %173 %174
+        %176 = OpFSub %v2float %172 %175
+               OpStore %colVel %176
+               OpBranch %170
+        %170 = OpLabel
+        %178 = OpLoad %v2float %pos_0
+        %179 = OpLoad %v2float %vPos
+        %177 = OpExtInst %float %40 Distance %178 %179
+        %181 = OpAccessChain %_ptr_Uniform_float %params %uint_3
+        %182 = OpLoad %float %181
+        %183 = OpFOrdLessThan %bool %177 %182
+               OpSelectionMerge %184 None
+               OpBranchConditional %183 %185 %184
+        %185 = OpLabel
+        %186 = OpLoad %v2float %cVel
+        %187 = OpLoad %v2float %vel
+        %188 = OpFAdd %v2float %186 %187
+               OpStore %cVel %188
+        %189 = OpLoad %int %cVelCount
+        %190 = OpIAdd %int %189 %int_1
+               OpStore %cVelCount %190
+               OpBranch %184
+        %184 = OpLabel
+               OpBranch %128
+        %128 = OpLabel
+        %191 = OpLoad %uint %i
+        %192 = OpIAdd %uint %191 %uint_1
+               OpStore %i %192
+               OpBranch %126
+        %127 = OpLabel
+        %193 = OpLoad %int %cMassCount
+        %194 = OpSGreaterThan %bool %193 %int_0
+               OpSelectionMerge %195 None
+               OpBranchConditional %194 %196 %195
+        %196 = OpLabel
+        %197 = OpLoad %v2float %cMass
+        %199 = OpLoad %int %cMassCount
+        %198 = OpConvertSToF %float %199
+        %201 = OpLoad %int %cMassCount
+        %200 = OpConvertSToF %float %201
+        %202 = OpCompositeConstruct %v2float %198 %200
+        %203 = OpFDiv %v2float %197 %202
+        %204 = OpLoad %v2float %vPos
+        %205 = OpFSub %v2float %203 %204
+               OpStore %cMass %205
+               OpBranch %195
+        %195 = OpLabel
+        %206 = OpLoad %int %cVelCount
+        %207 = OpSGreaterThan %bool %206 %int_0
+               OpSelectionMerge %208 None
+               OpBranchConditional %207 %209 %208
+        %209 = OpLabel
+        %210 = OpLoad %v2float %cVel
+        %212 = OpLoad %int %cVelCount
+        %211 = OpConvertSToF %float %212
+        %214 = OpLoad %int %cVelCount
+        %213 = OpConvertSToF %float %214
+        %215 = OpCompositeConstruct %v2float %211 %213
+        %216 = OpFDiv %v2float %210 %215
+               OpStore %cVel %216
+               OpBranch %208
+        %208 = OpLabel
+        %217 = OpLoad %v2float %vVel
+        %218 = OpLoad %v2float %cMass
+        %220 = OpAccessChain %_ptr_Uniform_float %params %uint_4
+        %221 = OpLoad %float %220
+        %222 = OpVectorTimesScalar %v2float %218 %221
+        %223 = OpFAdd %v2float %217 %222
+        %224 = OpLoad %v2float %colVel
+        %225 = OpAccessChain %_ptr_Uniform_float %params %uint_5
+        %226 = OpLoad %float %225
+        %227 = OpVectorTimesScalar %v2float %224 %226
+        %228 = OpFAdd %v2float %223 %227
+        %229 = OpLoad %v2float %cVel
+        %231 = OpAccessChain %_ptr_Uniform_float %params %uint_6
+        %232 = OpLoad %float %231
+        %233 = OpVectorTimesScalar %v2float %229 %232
+        %234 = OpFAdd %v2float %228 %233
+               OpStore %vVel %234
+        %236 = OpLoad %v2float %vVel
+        %235 = OpExtInst %v2float %40 Normalize %236
+        %239 = OpLoad %v2float %vVel
+        %238 = OpExtInst %float %40 Length %239
+        %237 = OpExtInst %float %40 NClamp %238 %float_0 %float_0_100000001
+        %241 = OpVectorTimesScalar %v2float %235 %237
+               OpStore %vVel %241
+        %242 = OpLoad %v2float %vPos
+        %243 = OpLoad %v2float %vVel
+        %244 = OpAccessChain %_ptr_Uniform_float %params %uint_0
+        %245 = OpLoad %float %244
+        %246 = OpVectorTimesScalar %v2float %243 %245
+        %247 = OpFAdd %v2float %242 %246
+               OpStore %vPos %247
+        %248 = OpAccessChain %_ptr_Function_float %vPos %uint_0
+        %249 = OpLoad %float %248
+        %251 = OpFOrdLessThan %bool %249 %float_n1
+               OpSelectionMerge %252 None
+               OpBranchConditional %251 %253 %252
+        %253 = OpLabel
+        %254 = OpAccessChain %_ptr_Function_float %vPos %uint_0
+               OpStore %254 %float_1
+               OpBranch %252
+        %252 = OpLabel
+        %255 = OpAccessChain %_ptr_Function_float %vPos %uint_0
+        %256 = OpLoad %float %255
+        %257 = OpFOrdGreaterThan %bool %256 %float_1
+               OpSelectionMerge %258 None
+               OpBranchConditional %257 %259 %258
+        %259 = OpLabel
+        %260 = OpAccessChain %_ptr_Function_float %vPos %uint_0
+               OpStore %260 %float_n1
+               OpBranch %258
+        %258 = OpLabel
+        %261 = OpAccessChain %_ptr_Function_float %vPos %uint_1
+        %262 = OpLoad %float %261
+        %263 = OpFOrdLessThan %bool %262 %float_n1
+               OpSelectionMerge %264 None
+               OpBranchConditional %263 %265 %264
+        %265 = OpLabel
+        %266 = OpAccessChain %_ptr_Function_float %vPos %uint_1
+               OpStore %266 %float_1
+               OpBranch %264
+        %264 = OpLabel
+        %267 = OpAccessChain %_ptr_Function_float %vPos %uint_1
+        %268 = OpLoad %float %267
+        %269 = OpFOrdGreaterThan %bool %268 %float_1
+               OpSelectionMerge %270 None
+               OpBranchConditional %269 %271 %270
+        %271 = OpLabel
+        %272 = OpAccessChain %_ptr_Function_float %vPos %uint_1
+               OpStore %272 %float_n1
+               OpBranch %270
+        %270 = OpLabel
+        %273 = OpLoad %uint %index
+        %274 = OpAccessChain %_ptr_StorageBuffer_v2float %particlesB %uint_0 %273 %uint_0
+        %275 = OpLoad %v2float %vPos
+               OpStore %274 %275
+        %276 = OpLoad %uint %index
+        %277 = OpAccessChain %_ptr_StorageBuffer_v2float %particlesB %uint_0 %276 %uint_1
+        %278 = OpLoad %v2float %vVel
+               OpStore %277 %278
+               OpReturn
+               OpFunctionEnd

test/samples/compute_boids.wgsl.expected.wgsl

@@ -0,0 +1,106 @@
+[[stage(vertex)]]
+fn vert_main([[location(0)]] a_particlePos : vec2<f32>, [[location(1)]] a_particleVel : vec2<f32>, [[location(2)]] a_pos : vec2<f32>) -> [[builtin(position)]] vec4<f32> {
+  var angle : f32 = -(atan2(a_particleVel.x, a_particleVel.y));
+  var pos : vec2<f32> = vec2<f32>(((a_pos.x * cos(angle)) - (a_pos.y * sin(angle))), ((a_pos.x * sin(angle)) + (a_pos.y * cos(angle))));
+  return vec4<f32>((pos + a_particlePos), 0.0, 1.0);
+}
+
+[[stage(fragment)]]
+fn frag_main() -> [[location(0)]] vec4<f32> {
+  return vec4<f32>(1.0, 1.0, 1.0, 1.0);
+}
+
+struct Particle {
+  pos : vec2<f32>;
+  vel : vec2<f32>;
+};
+
+[[block]]
+struct SimParams {
+  deltaT : f32;
+  rule1Distance : f32;
+  rule2Distance : f32;
+  rule3Distance : f32;
+  rule1Scale : f32;
+  rule2Scale : f32;
+  rule3Scale : f32;
+};
+
+[[block]]
+struct Particles {
+  particles : array<Particle, 5>;
+};
+
+[[binding(0), group(0)]] var<uniform> params : SimParams;
+
+[[binding(1), group(0)]] var<storage> particlesA : [[access(read_write)]] Particles;
+
+[[binding(2), group(0)]] var<storage> particlesB : [[access(read_write)]] Particles;
+
+[[stage(compute)]]
+fn comp_main([[builtin(global_invocation_id)]] gl_GlobalInvocationID : vec3<u32>) {
+  var index : u32 = gl_GlobalInvocationID.x;
+  if ((index >= 5u)) {
+    return;
+  }
+  var vPos : vec2<f32> = particlesA.particles[index].pos;
+  var vVel : vec2<f32> = particlesA.particles[index].vel;
+  var cMass : vec2<f32> = vec2<f32>(0.0, 0.0);
+  var cVel : vec2<f32> = vec2<f32>(0.0, 0.0);
+  var colVel : vec2<f32> = vec2<f32>(0.0, 0.0);
+  var cMassCount : i32 = 0;
+  var cVelCount : i32 = 0;
+  var pos : vec2<f32>;
+  var vel : vec2<f32>;
+  {
+    var i : u32 = 0u;
+    loop {
+      if (!((i < 5u))) {
+        break;
+      }
+      if ((i == index)) {
+        continue;
+      }
+      pos = particlesA.particles[i].pos.xy;
+      vel = particlesA.particles[i].vel.xy;
+      if ((distance(pos, vPos) < params.rule1Distance)) {
+        cMass = (cMass + pos);
+        cMassCount = (cMassCount + 1);
+      }
+      if ((distance(pos, vPos) < params.rule2Distance)) {
+        colVel = (colVel - (pos - vPos));
+      }
+      if ((distance(pos, vPos) < params.rule3Distance)) {
+        cVel = (cVel + vel);
+        cVelCount = (cVelCount + 1);
+      }
+
+      continuing {
+        i = (i + 1u);
+      }
+    }
+  }
+  if ((cMassCount > 0)) {
+    cMass = ((cMass / vec2<f32>(f32(cMassCount), f32(cMassCount))) - vPos);
+  }
+  if ((cVelCount > 0)) {
+    cVel = (cVel / vec2<f32>(f32(cVelCount), f32(cVelCount)));
+  }
+  vVel = (((vVel + (cMass * params.rule1Scale)) + (colVel * params.rule2Scale)) + (cVel * params.rule3Scale));
+  vVel = (normalize(vVel) * clamp(length(vVel), 0.0, 0.100000001));
+  vPos = (vPos + (vVel * params.deltaT));
+  if ((vPos.x < -1.0)) {
+    vPos.x = 1.0;
+  }
+  if ((vPos.x > 1.0)) {
+    vPos.x = -1.0;
+  }
+  if ((vPos.y < -1.0)) {
+    vPos.y = 1.0;
+  }
+  if ((vPos.y > 1.0)) {
+    vPos.y = -1.0;
+  }
+  particlesB.particles[index].pos = vPos;
+  particlesB.particles[index].vel = vVel;
+}

test/samples/cube.wgsl.expected.hlsl

@@ -0,0 +1,43 @@
+struct Uniforms {
+  float4x4 modelViewProjectionMatrix;
+};
+struct VertexInput {
+  float4 cur_position;
+  float4 color;
+};
+struct VertexOutput {
+  float4 vtxFragColor;
+  float4 Position;
+};
+struct tint_symbol_1 {
+  float4 cur_position : TEXCOORD0;
+  float4 color : TEXCOORD1;
+};
+struct tint_symbol_2 {
+  float4 vtxFragColor : TEXCOORD0;
+  float4 Position : SV_Position;
+};
+struct tint_symbol_4 {
+  float4 fragColor : TEXCOORD0;
+};
+struct tint_symbol_5 {
+  float4 value : SV_Target0;
+};
+
+ConstantBuffer<Uniforms> uniforms : register(b0, space0);
+
+tint_symbol_2 vtx_main(tint_symbol_1 tint_symbol) {
+  const VertexInput input = {tint_symbol.cur_position, tint_symbol.color};
+  VertexOutput output = {float4(0.0f, 0.0f, 0.0f, 0.0f), float4(0.0f, 0.0f, 0.0f, 0.0f)};
+  output.Position = mul(input.cur_position, uniforms.modelViewProjectionMatrix);
+  output.vtxFragColor = input.color;
+  const tint_symbol_2 tint_symbol_6 = {output.vtxFragColor, output.Position};
+  return tint_symbol_6;
+}
+
+tint_symbol_5 frag_main(tint_symbol_4 tint_symbol_3) {
+  const float4 fragColor = tint_symbol_3.fragColor;
+  const tint_symbol_5 tint_symbol_7 = {fragColor};
+  return tint_symbol_7;
+}
+

test/samples/cube.wgsl.expected.msl

@@ -0,0 +1,42 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ float4x4 modelViewProjectionMatrix;
+};
+struct VertexInput {
+  float4 cur_position;
+  float4 color;
+};
+struct VertexOutput {
+  float4 vtxFragColor;
+  float4 Position;
+};
+struct tint_symbol_1 {
+  float4 cur_position [[attribute(0)]];
+  float4 color [[attribute(1)]];
+};
+struct tint_symbol_2 {
+  float4 vtxFragColor [[user(locn0)]];
+  float4 Position [[position]];
+};
+struct tint_symbol_4 {
+  float4 fragColor [[user(locn0)]];
+};
+struct tint_symbol_5 {
+  float4 value [[color(0)]];
+};
+
+vertex tint_symbol_2 vtx_main(tint_symbol_1 tint_symbol [[stage_in]], constant Uniforms& uniforms [[buffer(0)]]) {
+  const VertexInput input = {tint_symbol.cur_position, tint_symbol.color};
+  VertexOutput output = {};
+  output.Position = (uniforms.modelViewProjectionMatrix * input.cur_position);
+  output.vtxFragColor = input.color;
+  return {output.vtxFragColor, output.Position};
+}
+
+fragment tint_symbol_5 frag_main(tint_symbol_4 tint_symbol_3 [[stage_in]]) {
+  const float4 fragColor = tint_symbol_3.fragColor;
+  return {fragColor};
+}
+

test/samples/cube.wgsl.expected.spvasm

@@ -0,0 +1,123 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 60
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vtx_main "vtx_main" %tint_pointsize %tint_symbol %tint_symbol_1 %tint_symbol_4 %tint_symbol_5
+               OpEntryPoint Fragment %frag_main "frag_main" %tint_symbol_7 %tint_symbol_9
+               OpExecutionMode %frag_main OriginUpperLeft
+               OpName %tint_pointsize "tint_pointsize"
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "modelViewProjectionMatrix"
+               OpName %uniforms "uniforms"
+               OpName %tint_symbol "tint_symbol"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %tint_symbol_4 "tint_symbol_4"
+               OpName %tint_symbol_5 "tint_symbol_5"
+               OpName %tint_symbol_7 "tint_symbol_7"
+               OpName %tint_symbol_9 "tint_symbol_9"
+               OpName %VertexOutput "VertexOutput"
+               OpMemberName %VertexOutput 0 "vtxFragColor"
+               OpMemberName %VertexOutput 1 "Position"
+               OpName %tint_symbol_6 "tint_symbol_6"
+               OpName %tint_symbol_3 "tint_symbol_3"
+               OpName %vtx_main "vtx_main"
+               OpName %VertexInput "VertexInput"
+               OpMemberName %VertexInput 0 "cur_position"
+               OpMemberName %VertexInput 1 "color"
+               OpName %output "output"
+               OpName %tint_symbol_10 "tint_symbol_10"
+               OpName %tint_symbol_8 "tint_symbol_8"
+               OpName %frag_main "frag_main"
+               OpDecorate %tint_pointsize BuiltIn PointSize
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpMemberDecorate %Uniforms 0 ColMajor
+               OpMemberDecorate %Uniforms 0 MatrixStride 16
+               OpDecorate %uniforms Binding 0
+               OpDecorate %uniforms DescriptorSet 0
+               OpDecorate %tint_symbol Location 0
+               OpDecorate %tint_symbol_1 Location 1
+               OpDecorate %tint_symbol_4 Location 0
+               OpDecorate %tint_symbol_5 BuiltIn Position
+               OpDecorate %tint_symbol_7 Location 0
+               OpDecorate %tint_symbol_9 Location 0
+               OpMemberDecorate %VertexOutput 0 Offset 0
+               OpMemberDecorate %VertexOutput 1 Offset 16
+               OpMemberDecorate %VertexInput 0 Offset 0
+               OpMemberDecorate %VertexInput 1 Offset 16
+      %float = OpTypeFloat 32
+%_ptr_Output_float = OpTypePointer Output %float
+          %4 = OpConstantNull %float
+%tint_pointsize = OpVariable %_ptr_Output_float Output %4
+    %v4float = OpTypeVector %float 4
+%mat4v4float = OpTypeMatrix %v4float 4
+   %Uniforms = OpTypeStruct %mat4v4float
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+%_ptr_Input_v4float = OpTypePointer Input %v4float
+%tint_symbol = OpVariable %_ptr_Input_v4float Input
+%tint_symbol_1 = OpVariable %_ptr_Input_v4float Input
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+         %15 = OpConstantNull %v4float
+%tint_symbol_4 = OpVariable %_ptr_Output_v4float Output %15
+%tint_symbol_5 = OpVariable %_ptr_Output_v4float Output %15
+%tint_symbol_7 = OpVariable %_ptr_Input_v4float Input
+%tint_symbol_9 = OpVariable %_ptr_Output_v4float Output %15
+       %void = OpTypeVoid
+%VertexOutput = OpTypeStruct %v4float %v4float
+         %19 = OpTypeFunction %void %VertexOutput
+         %27 = OpTypeFunction %void
+    %float_1 = OpConstant %float 1
+%VertexInput = OpTypeStruct %v4float %v4float
+%_ptr_Function_VertexOutput = OpTypePointer Function %VertexOutput
+         %37 = OpConstantNull %VertexOutput
+       %uint = OpTypeInt 32 0
+     %uint_1 = OpConstant %uint 1
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_mat4v4float = OpTypePointer Uniform %mat4v4float
+         %52 = OpTypeFunction %void %v4float
+%tint_symbol_6 = OpFunction %void None %19
+%tint_symbol_3 = OpFunctionParameter %VertexOutput
+         %24 = OpLabel
+         %25 = OpCompositeExtract %v4float %tint_symbol_3 0
+               OpStore %tint_symbol_4 %25
+         %26 = OpCompositeExtract %v4float %tint_symbol_3 1
+               OpStore %tint_symbol_5 %26
+               OpReturn
+               OpFunctionEnd
+   %vtx_main = OpFunction %void None %27
+         %29 = OpLabel
+     %output = OpVariable %_ptr_Function_VertexOutput Function %37
+               OpStore %tint_pointsize %float_1
+         %32 = OpLoad %v4float %tint_symbol
+         %33 = OpLoad %v4float %tint_symbol_1
+         %34 = OpCompositeConstruct %VertexInput %32 %33
+         %41 = OpAccessChain %_ptr_Function_v4float %output %uint_1
+         %44 = OpAccessChain %_ptr_Uniform_mat4v4float %uniforms %uint_0
+         %45 = OpLoad %mat4v4float %44
+         %46 = OpCompositeExtract %v4float %34 0
+         %47 = OpMatrixTimesVector %v4float %45 %46
+               OpStore %41 %47
+         %48 = OpAccessChain %_ptr_Function_v4float %output %uint_0
+         %49 = OpCompositeExtract %v4float %34 1
+               OpStore %48 %49
+         %51 = OpLoad %VertexOutput %output
+         %50 = OpFunctionCall %void %tint_symbol_6 %51
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_10 = OpFunction %void None %52
+%tint_symbol_8 = OpFunctionParameter %v4float
+         %55 = OpLabel
+               OpStore %tint_symbol_9 %tint_symbol_8
+               OpReturn
+               OpFunctionEnd
+  %frag_main = OpFunction %void None %27
+         %57 = OpLabel
+         %59 = OpLoad %v4float %tint_symbol_7
+         %58 = OpFunctionCall %void %tint_symbol_10 %59
+               OpReturn
+               OpFunctionEnd

test/samples/cube.wgsl.expected.wgsl

@@ -0,0 +1,33 @@
+[[block]]
+struct Uniforms {
+  modelViewProjectionMatrix : mat4x4<f32>;
+};
+
+[[binding(0), group(0)]] var<uniform> uniforms : Uniforms;
+
+struct VertexInput {
+  [[location(0)]]
+  cur_position : vec4<f32>;
+  [[location(1)]]
+  color : vec4<f32>;
+};
+
+struct VertexOutput {
+  [[location(0)]]
+  vtxFragColor : vec4<f32>;
+  [[builtin(position)]]
+  Position : vec4<f32>;
+};
+
+[[stage(vertex)]]
+fn vtx_main(input : VertexInput) -> VertexOutput {
+  var output : VertexOutput;
+  output.Position = (uniforms.modelViewProjectionMatrix * input.cur_position);
+  output.vtxFragColor = input.color;
+  return output;
+}
+
+[[stage(fragment)]]
+fn frag_main([[location(0)]] fragColor : vec4<f32>) -> [[location(0)]] vec4<f32> {
+  return fragColor;
+}

test/samples/function.wgsl.expected.hlsl

@@ -0,0 +1,9 @@
+float main() {
+  return (((2.0f * 3.0f) - 4.0f) / 5.0f);
+}
+
+[numthreads(2, 1, 1)]
+void ep() {
+  return;
+}
+

test/samples/function.wgsl.expected.msl

@@ -0,0 +1,11 @@
+#include <metal_stdlib>
+
+using namespace metal;
+float tint_symbol() {
+  return (((2.0f * 3.0f) - 4.0f) / 5.0f);
+}
+
+kernel void ep() {
+  return;
+}
+

test/samples/function.wgsl.expected.spvasm

@@ -0,0 +1,30 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %ep "ep"
+               OpExecutionMode %ep LocalSize 2 1 1
+               OpName %main "main"
+               OpName %ep "ep"
+      %float = OpTypeFloat 32
+          %1 = OpTypeFunction %float
+    %float_2 = OpConstant %float 2
+    %float_3 = OpConstant %float 3
+    %float_4 = OpConstant %float 4
+    %float_5 = OpConstant %float 5
+       %void = OpTypeVoid
+         %12 = OpTypeFunction %void
+       %main = OpFunction %float None %1
+          %4 = OpLabel
+          %7 = OpFMul %float %float_2 %float_3
+          %9 = OpFSub %float %7 %float_4
+         %11 = OpFDiv %float %9 %float_5
+               OpReturnValue %11
+               OpFunctionEnd
+         %ep = OpFunction %void None %12
+         %15 = OpLabel
+               OpReturn
+               OpFunctionEnd

test/samples/function.wgsl.expected.wgsl

@@ -0,0 +1,7 @@
+fn main() -> f32 {
+  return (((2.0 * 3.0) - 4.0) / 5.0);
+}
+
+[[stage(compute), workgroup_size(2, 1, 1)]]
+fn ep() {
+}

test/samples/simple.wgsl.expected.hlsl

@@ -0,0 +1,14 @@
+struct tint_symbol {
+  float4 value : SV_Target0;
+};
+
+void bar() {
+}
+
+tint_symbol main() {
+  float2 a = float2(0.0f, 0.0f);
+  bar();
+  const tint_symbol tint_symbol_1 = {float4(0.400000006f, 0.400000006f, 0.800000012f, 1.0f)};
+  return tint_symbol_1;
+}
+

test/samples/simple.wgsl.expected.msl

@@ -0,0 +1,16 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_1 {
+  float4 value [[color(0)]];
+};
+
+void bar() {
+}
+
+fragment tint_symbol_1 tint_symbol() {
+  float2 a = float2(0.0f);
+  bar();
+  return {float4(0.400000006f, 0.400000006f, 0.800000012f, 1.0f)};
+}
+

test/samples/simple.wgsl.expected.spvasm

@@ -0,0 +1,49 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 26
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %main "main" %tint_symbol_1
+               OpExecutionMode %main OriginUpperLeft
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %bar "bar"
+               OpName %tint_symbol_2 "tint_symbol_2"
+               OpName %tint_symbol "tint_symbol"
+               OpName %main "main"
+               OpName %a "a"
+               OpDecorate %tint_symbol_1 Location 0
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %5
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+         %10 = OpTypeFunction %void %v4float
+    %v2float = OpTypeVector %float 2
+         %17 = OpConstantNull %v2float
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+%float_0_400000006 = OpConstant %float 0.400000006
+%float_0_800000012 = OpConstant %float 0.800000012
+    %float_1 = OpConstant %float 1
+         %25 = OpConstantComposite %v4float %float_0_400000006 %float_0_400000006 %float_0_800000012 %float_1
+        %bar = OpFunction %void None %6
+          %9 = OpLabel
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %10
+%tint_symbol = OpFunctionParameter %v4float
+         %13 = OpLabel
+               OpStore %tint_symbol_1 %tint_symbol
+               OpReturn
+               OpFunctionEnd
+       %main = OpFunction %void None %6
+         %15 = OpLabel
+          %a = OpVariable %_ptr_Function_v2float Function %17
+               OpStore %a %17
+         %20 = OpFunctionCall %void %bar
+         %21 = OpFunctionCall %void %tint_symbol_2 %25
+               OpReturn
+               OpFunctionEnd

test/samples/simple.wgsl.expected.wgsl

@@ -0,0 +1,9 @@
+fn bar() {
+}
+
+[[stage(fragment)]]
+fn main() -> [[location(0)]] vec4<f32> {
+  var a : vec2<f32> = vec2<f32>();
+  bar();
+  return vec4<f32>(0.400000006, 0.400000006, 0.800000012, 1.0);
+}

test/samples/simple_vertex.spvasm.expected.hlsl

@@ -0,0 +1,10 @@
+struct main_out {
+  float4 gl_Position : SV_Position;
+};
+
+main_out main() {
+  main_out tint_out = (main_out)0;
+  tint_out.gl_Position = float4(0.0f, 0.0f, 0.0f, 0.0f);
+  return tint_out;
+}
+

test/samples/simple_vertex.spvasm.expected.msl

@@ -0,0 +1,13 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_out {
+  float4 gl_Position [[position]];
+};
+
+vertex tint_symbol_out tint_symbol() {
+  tint_symbol_out _tint_out = {};
+  _tint_out.gl_Position = float4(0.0f, 0.0f, 0.0f, 0.0f);
+  return _tint_out;
+}
+

test/samples/simple_vertex.spvasm.expected.spvasm

@@ -0,0 +1,32 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %main "main" %tint_pointsize %gl_Position
+               OpName %tint_pointsize "tint_pointsize"
+               OpName %gl_Position "gl_Position"
+               OpName %main "main"
+               OpDecorate %tint_pointsize BuiltIn PointSize
+               OpDecorate %gl_Position BuiltIn Position
+      %float = OpTypeFloat 32
+%_ptr_Output_float = OpTypePointer Output %float
+          %4 = OpConstantNull %float
+%tint_pointsize = OpVariable %_ptr_Output_float Output %4
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %8 = OpConstantNull %v4float
+%gl_Position = OpVariable %_ptr_Output_v4float Output %8
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+    %float_1 = OpConstant %float 1
+    %float_0 = OpConstant %float 0
+         %15 = OpConstantComposite %v4float %float_0 %float_0 %float_0 %float_0
+       %main = OpFunction %void None %9
+         %12 = OpLabel
+               OpStore %tint_pointsize %float_1
+               OpStore %gl_Position %15
+               OpReturn
+               OpFunctionEnd

test/samples/simple_vertex.spvasm.expected.wgsl

@@ -0,0 +1,7 @@
+[[builtin(position)]] var<out> gl_Position : vec4<f32>;
+
+[[stage(vertex)]]
+fn main() {
+  gl_Position = vec4<f32>(0.0, 0.0, 0.0, 0.0);
+  return;
+}

test/samples/triangle.wgsl.expected.hlsl

@@ -0,0 +1,22 @@
+struct tint_symbol_1 {
+  int VertexIndex : SV_VertexID;
+};
+struct tint_symbol_2 {
+  float4 value : SV_Position;
+};
+struct tint_symbol_3 {
+  float4 value : SV_Target0;
+};
+
+static const float2 pos[3] = {float2(0.0f, 0.5f), float2(-0.5f, -0.5f), float2(0.5f, -0.5f)};
+tint_symbol_2 vtx_main(tint_symbol_1 tint_symbol) {
+  const int VertexIndex = tint_symbol.VertexIndex;
+  const tint_symbol_2 tint_symbol_4 = {float4(pos[VertexIndex], 0.0f, 1.0f)};
+  return tint_symbol_4;
+}
+
+tint_symbol_3 frag_main() {
+  const tint_symbol_3 tint_symbol_5 = {float4(1.0f, 0.0f, 0.0f, 1.0f)};
+  return tint_symbol_5;
+}
+

test/samples/triangle.wgsl.expected.msl

@@ -0,0 +1,23 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_1 {
+  int VertexIndex [[vertex_id]];
+};
+struct tint_symbol_2 {
+  float4 value [[position]];
+};
+struct tint_symbol_3 {
+  float4 value [[color(0)]];
+};
+
+constant float2 pos[3] = {float2(0.0f, 0.5f), float2(-0.5f, -0.5f), float2(0.5f, -0.5f)};
+vertex tint_symbol_2 vtx_main(tint_symbol_1 tint_symbol [[stage_in]]) {
+  const int VertexIndex = tint_symbol.VertexIndex;
+  return {float4(pos[VertexIndex], 0.0f, 1.0f)};
+}
+
+fragment tint_symbol_3 frag_main() {
+  return {float4(1.0f, 0.0f, 0.0f, 1.0f)};
+}
+

test/samples/triangle.wgsl.expected.spvasm

@@ -0,0 +1,88 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 51
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vtx_main "vtx_main" %tint_pointsize %tint_symbol %tint_symbol_2
+               OpEntryPoint Fragment %frag_main "frag_main" %tint_symbol_5
+               OpExecutionMode %frag_main OriginUpperLeft
+               OpName %tint_pointsize "tint_pointsize"
+               OpName %pos "pos"
+               OpName %tint_symbol "tint_symbol"
+               OpName %tint_symbol_2 "tint_symbol_2"
+               OpName %tint_symbol_5 "tint_symbol_5"
+               OpName %tint_symbol_3 "tint_symbol_3"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %vtx_main "vtx_main"
+               OpName %tint_symbol_6 "tint_symbol_6"
+               OpName %tint_symbol_4 "tint_symbol_4"
+               OpName %frag_main "frag_main"
+               OpDecorate %tint_pointsize BuiltIn PointSize
+               OpDecorate %_arr_v2float_uint_3 ArrayStride 8
+               OpDecorate %tint_symbol BuiltIn VertexIndex
+               OpDecorate %tint_symbol_2 BuiltIn Position
+               OpDecorate %tint_symbol_5 Location 0
+      %float = OpTypeFloat 32
+%_ptr_Output_float = OpTypePointer Output %float
+          %4 = OpConstantNull %float
+%tint_pointsize = OpVariable %_ptr_Output_float Output %4
+    %v2float = OpTypeVector %float 2
+       %uint = OpTypeInt 32 0
+     %uint_3 = OpConstant %uint 3
+%_arr_v2float_uint_3 = OpTypeArray %v2float %uint_3
+    %float_0 = OpConstant %float 0
+  %float_0_5 = OpConstant %float 0.5
+         %11 = OpConstantComposite %v2float %float_0 %float_0_5
+ %float_n0_5 = OpConstant %float -0.5
+         %13 = OpConstantComposite %v2float %float_n0_5 %float_n0_5
+         %14 = OpConstantComposite %v2float %float_0_5 %float_n0_5
+        %pos = OpConstantComposite %_arr_v2float_uint_3 %11 %13 %14
+        %int = OpTypeInt 32 1
+%_ptr_Input_int = OpTypePointer Input %int
+%tint_symbol = OpVariable %_ptr_Input_int Input
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+         %22 = OpConstantNull %v4float
+%tint_symbol_2 = OpVariable %_ptr_Output_v4float Output %22
+%tint_symbol_5 = OpVariable %_ptr_Output_v4float Output %22
+       %void = OpTypeVoid
+         %24 = OpTypeFunction %void %v4float
+         %29 = OpTypeFunction %void
+    %float_1 = OpConstant %float 1
+%_ptr_Function__arr_v2float_uint_3 = OpTypePointer Function %_arr_v2float_uint_3
+         %36 = OpConstantNull %_arr_v2float_uint_3
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+         %50 = OpConstantComposite %v4float %float_1 %float_0 %float_0 %float_1
+%tint_symbol_3 = OpFunction %void None %24
+%tint_symbol_1 = OpFunctionParameter %v4float
+         %28 = OpLabel
+               OpStore %tint_symbol_2 %tint_symbol_1
+               OpReturn
+               OpFunctionEnd
+   %vtx_main = OpFunction %void None %29
+         %31 = OpLabel
+         %35 = OpVariable %_ptr_Function__arr_v2float_uint_3 Function %36
+               OpStore %tint_pointsize %float_1
+               OpStore %35 %pos
+         %37 = OpLoad %int %tint_symbol
+         %39 = OpAccessChain %_ptr_Function_v2float %35 %37
+         %40 = OpLoad %v2float %39
+         %41 = OpCompositeExtract %float %40 0
+         %42 = OpCompositeExtract %float %40 1
+         %43 = OpCompositeConstruct %v4float %41 %42 %float_0 %float_1
+         %33 = OpFunctionCall %void %tint_symbol_3 %43
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_6 = OpFunction %void None %24
+%tint_symbol_4 = OpFunctionParameter %v4float
+         %46 = OpLabel
+               OpStore %tint_symbol_5 %tint_symbol_4
+               OpReturn
+               OpFunctionEnd
+  %frag_main = OpFunction %void None %29
+         %48 = OpLabel
+         %49 = OpFunctionCall %void %tint_symbol_6 %50
+               OpReturn
+               OpFunctionEnd

test/samples/triangle.wgsl.expected.wgsl

@@ -0,0 +1,11 @@
+let pos : array<vec2<f32>, 3> = array<vec2<f32>, 3>(vec2<f32>(0.0, 0.5), vec2<f32>(-0.5, -0.5), vec2<f32>(0.5, -0.5));
+
+[[stage(vertex)]]
+fn vtx_main([[builtin(vertex_index)]] VertexIndex : i32) -> [[builtin(position)]] vec4<f32> {
+  return vec4<f32>(pos[VertexIndex], 0.0, 1.0);
+}
+
+[[stage(fragment)]]
+fn frag_main() -> [[location(0)]] vec4<f32> {
+  return vec4<f32>(1.0, 0.0, 0.0, 1.0);
+}