diff --git a/test/samples/compute_boids.wgsl.expected.hlsl b/test/samples/compute_boids.wgsl.expected.hlsl
new file mode 100644
index 0000000..e29d25b
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/samples/compute_boids.wgsl.expected.msl b/test/samples/compute_boids.wgsl.expected.msl
new file mode 100644
index 0000000..05afb51
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/samples/compute_boids.wgsl.expected.spvasm b/test/samples/compute_boids.wgsl.expected.spvasm
new file mode 100644
index 0000000..a164a17
--- /dev/null
+++ b/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
diff --git a/test/samples/compute_boids.wgsl.expected.wgsl b/test/samples/compute_boids.wgsl.expected.wgsl
new file mode 100644
index 0000000..882e263
--- /dev/null
+++ b/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;
+}
diff --git a/test/samples/cube.wgsl.expected.hlsl b/test/samples/cube.wgsl.expected.hlsl
new file mode 100644
index 0000000..c15b578
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/samples/cube.wgsl.expected.msl b/test/samples/cube.wgsl.expected.msl
new file mode 100644
index 0000000..4d86279
--- /dev/null
+++ b/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};
+}
+
diff --git a/test/samples/cube.wgsl.expected.spvasm b/test/samples/cube.wgsl.expected.spvasm
new file mode 100644
index 0000000..8323589
--- /dev/null
+++ b/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
diff --git a/test/samples/cube.wgsl.expected.wgsl b/test/samples/cube.wgsl.expected.wgsl
new file mode 100644
index 0000000..a53fdec
--- /dev/null
+++ b/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;
+}
diff --git a/test/samples/function.wgsl.expected.hlsl b/test/samples/function.wgsl.expected.hlsl
new file mode 100644
index 0000000..deace9c
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/samples/function.wgsl.expected.msl b/test/samples/function.wgsl.expected.msl
new file mode 100644
index 0000000..4ab43dd
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/samples/function.wgsl.expected.spvasm b/test/samples/function.wgsl.expected.spvasm
new file mode 100644
index 0000000..affe051
--- /dev/null
+++ b/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
diff --git a/test/samples/function.wgsl.expected.wgsl b/test/samples/function.wgsl.expected.wgsl
new file mode 100644
index 0000000..d8a4c35
--- /dev/null
+++ b/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() {
+}
diff --git a/test/samples/simple.wgsl.expected.hlsl b/test/samples/simple.wgsl.expected.hlsl
new file mode 100644
index 0000000..f48bbae
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/samples/simple.wgsl.expected.msl b/test/samples/simple.wgsl.expected.msl
new file mode 100644
index 0000000..a132f06
--- /dev/null
+++ b/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)};
+}
+
diff --git a/test/samples/simple.wgsl.expected.spvasm b/test/samples/simple.wgsl.expected.spvasm
new file mode 100644
index 0000000..ac34752
--- /dev/null
+++ b/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
diff --git a/test/samples/simple.wgsl.expected.wgsl b/test/samples/simple.wgsl.expected.wgsl
new file mode 100644
index 0000000..b5a4dc0
--- /dev/null
+++ b/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);
+}
diff --git a/test/samples/simple_vertex.spvasm.expected.hlsl b/test/samples/simple_vertex.spvasm.expected.hlsl
new file mode 100644
index 0000000..40f75d7
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/samples/simple_vertex.spvasm.expected.msl b/test/samples/simple_vertex.spvasm.expected.msl
new file mode 100644
index 0000000..58199d4
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/samples/simple_vertex.spvasm.expected.spvasm b/test/samples/simple_vertex.spvasm.expected.spvasm
new file mode 100644
index 0000000..89ff63f
--- /dev/null
+++ b/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
diff --git a/test/samples/simple_vertex.spvasm.expected.wgsl b/test/samples/simple_vertex.spvasm.expected.wgsl
new file mode 100644
index 0000000..75c87cd
--- /dev/null
+++ b/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;
+}
diff --git a/test/samples/triangle.wgsl.expected.hlsl b/test/samples/triangle.wgsl.expected.hlsl
new file mode 100644
index 0000000..803a381
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/samples/triangle.wgsl.expected.msl b/test/samples/triangle.wgsl.expected.msl
new file mode 100644
index 0000000..00638c3
--- /dev/null
+++ b/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)};
+}
+
diff --git a/test/samples/triangle.wgsl.expected.spvasm b/test/samples/triangle.wgsl.expected.spvasm
new file mode 100644
index 0000000..8a7bf08
--- /dev/null
+++ b/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
diff --git a/test/samples/triangle.wgsl.expected.wgsl b/test/samples/triangle.wgsl.expected.wgsl
new file mode 100644
index 0000000..0085989
--- /dev/null
+++ b/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);
+}