benchmarks: Add a basic set of benchmarks

Add google benchmark to the DEPs.

Implement a basic set of benchmarks for each of the writers and the WGSL parser.

Add build rules for CMake. GN build rules TODO.

Add a simple go tool (ported from Marl) to diff two benchmarks. Less
noisy than the one provided by google benchmark.

Bug: tint:1378
Change-Id: I73cf92c5d9fd2d3bfac8f264864fd774afbd5d01
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/76840
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ryan Harrison <rharrison@chromium.org>
Commit-Queue: Ben Clayton <bclayton@chromium.org>

test/benchmark/empty.wgsl.expected.hlsl

@@ -0,0 +1,4 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+  return;
+}

test/benchmark/empty.wgsl.expected.msl

@@ -0,0 +1,3 @@
+#include <metal_stdlib>
+
+using namespace metal;

test/benchmark/empty.wgsl.expected.spvasm

@@ -0,0 +1,16 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 5
+; 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"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+%unused_entry_point = OpFunction %void None %1
+          %4 = OpLabel
+               OpReturn
+               OpFunctionEnd

test/benchmark/particles.wgsl

@@ -0,0 +1,183 @@
+////////////////////////////////////////////////////////////////////////////////
+// Utilities
+////////////////////////////////////////////////////////////////////////////////
+var<private> rand_seed : vec2<f32>;
+
+fn rand() -> f32 {
+    rand_seed.x = fract(cos(dot(rand_seed, vec2<f32>(23.14077926, 232.61690225))) * 136.8168);
+    rand_seed.y = fract(cos(dot(rand_seed, vec2<f32>(54.47856553, 345.84153136))) * 534.7645);
+    return rand_seed.y;
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Vertex shader
+////////////////////////////////////////////////////////////////////////////////
+struct RenderParams {
+  modelViewProjectionMatrix : mat4x4<f32>;
+  right : vec3<f32>;
+  up    : vec3<f32>;
+};
+[[binding(0), group(0)]] var<uniform> render_params : RenderParams;
+
+struct VertexInput {
+  [[location(0)]] position : vec3<f32>;
+  [[location(1)]] color    : vec4<f32>;
+  [[location(2)]] quad_pos : vec2<f32>; // -1..+1
+};
+
+struct VertexOutput {
+  [[builtin(position)]] position : vec4<f32>;
+  [[location(0)]]       color    : vec4<f32>;
+  [[location(1)]]       quad_pos : vec2<f32>; // -1..+1
+};
+
+[[stage(vertex)]]
+fn vs_main(in : VertexInput) -> VertexOutput {
+  var quad_pos = mat2x3<f32>(render_params.right, render_params.up) * in.quad_pos;
+  var position = in.position + quad_pos * 0.01;
+  var out : VertexOutput;
+  out.position = render_params.modelViewProjectionMatrix * vec4<f32>(position, 1.0);
+  out.color = in.color;
+  out.quad_pos = in.quad_pos;
+  return out;
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Fragment shader
+////////////////////////////////////////////////////////////////////////////////
+[[stage(fragment)]]
+fn fs_main(in : VertexOutput) -> [[location(0)]] vec4<f32> {
+  var color = in.color;
+  // Apply a circular particle alpha mask
+  color.a = color.a * max(1.0 - length(in.quad_pos), 0.0);
+  return color;
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Simulation Compute shader
+////////////////////////////////////////////////////////////////////////////////
+struct SimulationParams {
+  deltaTime : f32;
+  seed : vec4<f32>;
+};
+
+struct Particle {
+  position : vec3<f32>;
+  lifetime : f32;
+  color    : vec4<f32>;
+  velocity : vec3<f32>;
+};
+
+struct Particles {
+  particles : array<Particle>;
+};
+
+[[binding(0), group(0)]] var<uniform> sim_params : SimulationParams;
+[[binding(1), group(0)]] var<storage, read_write> data : Particles;
+[[binding(2), group(0)]] var texture : texture_2d<f32>;
+
+[[stage(compute), workgroup_size(64)]]
+fn simulate([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
+  rand_seed = (sim_params.seed.xy + vec2<f32>(GlobalInvocationID.xy)) * sim_params.seed.zw;
+
+  let idx = GlobalInvocationID.x;
+  var particle = data.particles[idx];
+
+  // Apply gravity
+  particle.velocity.z = particle.velocity.z - sim_params.deltaTime * 0.5;
+
+  // Basic velocity integration
+  particle.position = particle.position + sim_params.deltaTime * particle.velocity;
+
+  // Age each particle. Fade out before vanishing.
+  particle.lifetime = particle.lifetime - sim_params.deltaTime;
+  particle.color.a = smoothStep(0.0, 0.5, particle.lifetime);
+
+  // If the lifetime has gone negative, then the particle is dead and should be
+  // respawned.
+  if (particle.lifetime < 0.0) {
+    // Use the probability map to find where the particle should be spawned.
+    // Starting with the 1x1 mip level.
+    var coord = vec2<i32>(0, 0);
+    for (var level = textureNumLevels(texture) - 1; level > 0; level = level - 1) {
+      // Load the probability value from the mip-level
+      // Generate a random number and using the probabilty values, pick the
+      // next texel in the next largest mip level:
+      //
+      // 0.0    probabilites.r    probabilites.g    probabilites.b   1.0
+      //  |              |              |              |              |
+      //  |   TOP-LEFT   |  TOP-RIGHT   | BOTTOM-LEFT  | BOTTOM_RIGHT |
+      //
+      let probabilites = textureLoad(texture, coord, level);
+      let value = vec4<f32>(rand());
+      let mask = (value >= vec4<f32>(0.0, probabilites.xyz)) & (value < probabilites);
+      coord = coord * 2;
+      coord.x = coord.x + select(0, 1, any(mask.yw)); // x  y
+      coord.y = coord.y + select(0, 1, any(mask.zw)); // z  w
+    }
+    let uv = vec2<f32>(coord) / vec2<f32>(textureDimensions(texture));
+    particle.position = vec3<f32>((uv - 0.5) * 3.0 * vec2<f32>(1.0, -1.0), 0.0);
+    particle.color = textureLoad(texture, coord, 0);
+    particle.velocity.x = (rand() - 0.5) * 0.1;
+    particle.velocity.y = (rand() - 0.5) * 0.1;
+    particle.velocity.z = rand() * 0.3;
+    particle.lifetime = 0.5 + rand() * 2.0;
+  }
+
+  // Store the new particle value
+  data.particles[idx] = particle;
+}
+
+struct UBO {
+  width : u32;
+};
+
+struct Buffer {
+  weights : array<f32>;
+};
+
+[[binding(3), group(0)]] var<uniform> ubo : UBO;
+[[binding(4), group(0)]] var<storage, read> buf_in : Buffer;
+[[binding(5), group(0)]] var<storage, read_write> buf_out : Buffer;
+[[binding(6), group(0)]] var tex_in : texture_2d<f32>;
+[[binding(7), group(0)]] var tex_out : texture_storage_2d<rgba8unorm, write>;
+
+////////////////////////////////////////////////////////////////////////////////
+// import_level
+//
+// Loads the alpha channel from a texel of the source image, and writes it to
+// the buf_out.weights.
+////////////////////////////////////////////////////////////////////////////////
+[[stage(compute), workgroup_size(64)]]
+fn import_level([[builtin(global_invocation_id)]] coord : vec3<u32>) {
+  _ = &buf_in;
+  let offset = coord.x + coord.y * ubo.width;
+  buf_out.weights[offset] = textureLoad(tex_in, vec2<i32>(coord.xy), 0).w;
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// export_level
+//
+// Loads 4 f32 weight values from buf_in.weights, and stores summed value into
+// buf_out.weights, along with the calculated 'probabilty' vec4 values into the
+// mip level of tex_out. See simulate() in particle.wgsl to understand the
+// probability logic.
+////////////////////////////////////////////////////////////////////////////////
+[[stage(compute), workgroup_size(64)]]
+fn export_level([[builtin(global_invocation_id)]] coord : vec3<u32>) {
+  if (all(coord.xy < vec2<u32>(textureDimensions(tex_out)))) {
+    let dst_offset = coord.x    + coord.y    * ubo.width;
+    let src_offset = coord.x*2u + coord.y*2u * ubo.width;
+
+    let a = buf_in.weights[src_offset + 0u];
+    let b = buf_in.weights[src_offset + 1u];
+    let c = buf_in.weights[src_offset + 0u + ubo.width];
+    let d = buf_in.weights[src_offset + 1u + ubo.width];
+    let sum = dot(vec4<f32>(a, b, c, d), vec4<f32>(1.0));
+
+    buf_out.weights[dst_offset] = sum / 4.0;
+
+    let probabilities = vec4<f32>(a, a+b, a+b+c, sum) / max(sum, 0.0001);
+    textureStore(tex_out, vec2<i32>(coord.xy), probabilities);
+  }
+}

test/benchmark/particles.wgsl.expected.hlsl

@@ -0,0 +1,204 @@
+static float2 rand_seed = float2(0.0f, 0.0f);
+
+float rand() {
+  rand_seed.x = frac((cos(dot(rand_seed, float2(23.140779495f, 232.616897583f))) * 136.816802979f));
+  rand_seed.y = frac((cos(dot(rand_seed, float2(54.478565216f, 345.841522217f))) * 534.764526367f));
+  return rand_seed.y;
+}
+
+cbuffer cbuffer_render_params : register(b0, space0) {
+  uint4 render_params[6];
+};
+
+struct VertexInput {
+  float3 position;
+  float4 color;
+  float2 quad_pos;
+};
+struct VertexOutput {
+  float4 position;
+  float4 color;
+  float2 quad_pos;
+};
+struct tint_symbol_5 {
+  float3 position : TEXCOORD0;
+  float4 color : TEXCOORD1;
+  float2 quad_pos : TEXCOORD2;
+};
+struct tint_symbol_6 {
+  float4 color : TEXCOORD0;
+  float2 quad_pos : TEXCOORD1;
+  float4 position : SV_Position;
+};
+
+float4x4 tint_symbol_17(uint4 buffer[6], uint offset) {
+  const uint scalar_offset = ((offset + 0u)) / 4;
+  const uint scalar_offset_1 = ((offset + 16u)) / 4;
+  const uint scalar_offset_2 = ((offset + 32u)) / 4;
+  const uint scalar_offset_3 = ((offset + 48u)) / 4;
+  return float4x4(asfloat(buffer[scalar_offset / 4]), asfloat(buffer[scalar_offset_1 / 4]), asfloat(buffer[scalar_offset_2 / 4]), asfloat(buffer[scalar_offset_3 / 4]));
+}
+
+VertexOutput vs_main_inner(VertexInput tint_symbol) {
+  float3 quad_pos = mul(tint_symbol.quad_pos, float2x3(asfloat(render_params[4].xyz), asfloat(render_params[5].xyz)));
+  float3 position = (tint_symbol.position + (quad_pos * 0.01f));
+  VertexOutput tint_symbol_1 = (VertexOutput)0;
+  tint_symbol_1.position = mul(float4(position, 1.0f), tint_symbol_17(render_params, 0u));
+  tint_symbol_1.color = tint_symbol.color;
+  tint_symbol_1.quad_pos = tint_symbol.quad_pos;
+  return tint_symbol_1;
+}
+
+tint_symbol_6 vs_main(tint_symbol_5 tint_symbol_4) {
+  const VertexInput tint_symbol_32 = {tint_symbol_4.position, tint_symbol_4.color, tint_symbol_4.quad_pos};
+  const VertexOutput inner_result = vs_main_inner(tint_symbol_32);
+  tint_symbol_6 wrapper_result = (tint_symbol_6)0;
+  wrapper_result.position = inner_result.position;
+  wrapper_result.color = inner_result.color;
+  wrapper_result.quad_pos = inner_result.quad_pos;
+  return wrapper_result;
+}
+
+struct tint_symbol_8 {
+  float4 color : TEXCOORD0;
+  float2 quad_pos : TEXCOORD1;
+  float4 position : SV_Position;
+};
+struct tint_symbol_9 {
+  float4 value : SV_Target0;
+};
+
+float4 fs_main_inner(VertexOutput tint_symbol) {
+  float4 color = tint_symbol.color;
+  color.a = (color.a * max((1.0f - length(tint_symbol.quad_pos)), 0.0f));
+  return color;
+}
+
+tint_symbol_9 fs_main(tint_symbol_8 tint_symbol_7) {
+  const VertexOutput tint_symbol_33 = {tint_symbol_7.position, tint_symbol_7.color, tint_symbol_7.quad_pos};
+  const float4 inner_result_1 = fs_main_inner(tint_symbol_33);
+  tint_symbol_9 wrapper_result_1 = (tint_symbol_9)0;
+  wrapper_result_1.value = inner_result_1;
+  return wrapper_result_1;
+}
+
+struct Particle {
+  float3 position;
+  float lifetime;
+  float4 color;
+  float3 velocity;
+};
+
+cbuffer cbuffer_sim_params : register(b0, space0) {
+  uint4 sim_params[2];
+};
+RWByteAddressBuffer data : register(u1, space0);
+Texture2D<float4> tint_symbol_2 : register(t2, space0);
+
+struct tint_symbol_11 {
+  uint3 GlobalInvocationID : SV_DispatchThreadID;
+};
+
+Particle tint_symbol_20(RWByteAddressBuffer buffer, uint offset) {
+  const Particle tint_symbol_34 = {asfloat(buffer.Load3((offset + 0u))), asfloat(buffer.Load((offset + 12u))), asfloat(buffer.Load4((offset + 16u))), asfloat(buffer.Load3((offset + 32u)))};
+  return tint_symbol_34;
+}
+
+void tint_symbol_25(RWByteAddressBuffer buffer, uint offset, Particle value) {
+  buffer.Store3((offset + 0u), asuint(value.position));
+  buffer.Store((offset + 12u), asuint(value.lifetime));
+  buffer.Store4((offset + 16u), asuint(value.color));
+  buffer.Store3((offset + 32u), asuint(value.velocity));
+}
+
+void simulate_inner(uint3 GlobalInvocationID) {
+  rand_seed = ((asfloat(sim_params[1]).xy + float2(GlobalInvocationID.xy)) * asfloat(sim_params[1]).zw);
+  const uint idx = GlobalInvocationID.x;
+  Particle particle = tint_symbol_20(data, (48u * idx));
+  particle.velocity.z = (particle.velocity.z - (asfloat(sim_params[0].x) * 0.5f));
+  particle.position = (particle.position + (asfloat(sim_params[0].x) * particle.velocity));
+  particle.lifetime = (particle.lifetime - asfloat(sim_params[0].x));
+  particle.color.a = smoothstep(0.0f, 0.5f, particle.lifetime);
+  if ((particle.lifetime < 0.0f)) {
+    int2 coord = int2(0, 0);
+    {
+      int3 tint_tmp;
+      tint_symbol_2.GetDimensions(0, tint_tmp.x, tint_tmp.y, tint_tmp.z);
+      int level = (tint_tmp.z - 1);
+      [loop] for(; (level > 0); level = (level - 1)) {
+        const float4 probabilites = tint_symbol_2.Load(int3(coord, level));
+        const float4 value = float4((rand()).xxxx);
+        const bool4 mask = ((value >= float4(0.0f, probabilites.xyz)) & (value < probabilites));
+        coord = (coord * 2);
+        coord.x = (coord.x + (any(mask.yw) ? 1 : 0));
+        coord.y = (coord.y + (any(mask.zw) ? 1 : 0));
+      }
+    }
+    int2 tint_tmp_1;
+    tint_symbol_2.GetDimensions(tint_tmp_1.x, tint_tmp_1.y);
+    const float2 uv = (float2(coord) / float2(tint_tmp_1));
+    particle.position = float3((((uv - 0.5f) * 3.0f) * float2(1.0f, -1.0f)), 0.0f);
+    particle.color = tint_symbol_2.Load(int3(coord, 0));
+    particle.velocity.x = ((rand() - 0.5f) * 0.100000001f);
+    particle.velocity.y = ((rand() - 0.5f) * 0.100000001f);
+    particle.velocity.z = (rand() * 0.300000012f);
+    particle.lifetime = (0.5f + (rand() * 2.0f));
+  }
+  tint_symbol_25(data, (48u * idx), particle);
+}
+
+[numthreads(64, 1, 1)]
+void simulate(tint_symbol_11 tint_symbol_10) {
+  simulate_inner(tint_symbol_10.GlobalInvocationID);
+  return;
+}
+
+cbuffer cbuffer_ubo : register(b3, space0) {
+  uint4 ubo[1];
+};
+ByteAddressBuffer buf_in : register(t4, space0);
+RWByteAddressBuffer buf_out : register(u5, space0);
+Texture2D<float4> tex_in : register(t6, space0);
+RWTexture2D<float4> tex_out : register(u7, space0);
+
+struct tint_symbol_13 {
+  uint3 coord : SV_DispatchThreadID;
+};
+
+void import_level_inner(uint3 coord) {
+  const uint offset = (coord.x + (coord.y * ubo[0].x));
+  buf_out.Store((4u * offset), asuint(tex_in.Load(int3(int2(coord.xy), 0)).w));
+}
+
+[numthreads(64, 1, 1)]
+void import_level(tint_symbol_13 tint_symbol_12) {
+  import_level_inner(tint_symbol_12.coord);
+  return;
+}
+
+struct tint_symbol_15 {
+  uint3 coord : SV_DispatchThreadID;
+};
+
+void export_level_inner(uint3 coord) {
+  int2 tint_tmp_2;
+  tex_out.GetDimensions(tint_tmp_2.x, tint_tmp_2.y);
+  if (all((coord.xy < uint2(tint_tmp_2)))) {
+    const uint dst_offset = (coord.x + (coord.y * ubo[0].x));
+    const uint src_offset = ((coord.x * 2u) + ((coord.y * 2u) * ubo[0].x));
+    const float a_1 = asfloat(buf_in.Load((4u * (src_offset + 0u))));
+    const float b = asfloat(buf_in.Load((4u * (src_offset + 1u))));
+    const float c = asfloat(buf_in.Load((4u * ((src_offset + 0u) + ubo[0].x))));
+    const float d = asfloat(buf_in.Load((4u * ((src_offset + 1u) + ubo[0].x))));
+    const float sum = dot(float4(a_1, b, c, d), float4((1.0f).xxxx));
+    buf_out.Store((4u * dst_offset), asuint((sum / 4.0f)));
+    const float4 probabilities = (float4(a_1, (a_1 + b), ((a_1 + b) + c), sum) / max(sum, 0.0001f));
+    tex_out[int2(coord.xy)] = probabilities;
+  }
+}
+
+[numthreads(64, 1, 1)]
+void export_level(tint_symbol_15 tint_symbol_14) {
+  export_level_inner(tint_symbol_14.coord);
+  return;
+}

test/benchmark/particles.wgsl.expected.msl

@@ -0,0 +1,175 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, int N, int M>
+inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
+  return lhs * vec<T, N>(rhs);
+}
+
+template<typename T, int N, int M>
+inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
+  return vec<T, M>(lhs) * rhs;
+}
+
+struct RenderParams {
+  /* 0x0000 */ float4x4 modelViewProjectionMatrix;
+  /* 0x0040 */ packed_float3 right;
+  /* 0x004c */ int8_t tint_pad[4];
+  /* 0x0050 */ packed_float3 up;
+  /* 0x005c */ int8_t tint_pad_1[4];
+};
+struct VertexInput {
+  float3 position;
+  float4 color;
+  float2 quad_pos;
+};
+struct VertexOutput {
+  float4 position;
+  float4 color;
+  float2 quad_pos;
+};
+struct tint_symbol_2 {
+  float3 position [[attribute(0)]];
+  float4 color [[attribute(1)]];
+  float2 quad_pos [[attribute(2)]];
+};
+struct tint_symbol_3 {
+  float4 color [[user(locn0)]];
+  float2 quad_pos [[user(locn1)]];
+  float4 position [[position]];
+};
+struct tint_symbol_5 {
+  float4 color [[user(locn0)]];
+  float2 quad_pos [[user(locn1)]];
+};
+struct tint_symbol_6 {
+  float4 value [[color(0)]];
+};
+struct SimulationParams {
+  /* 0x0000 */ float deltaTime;
+  /* 0x0004 */ int8_t tint_pad_2[12];
+  /* 0x0010 */ float4 seed;
+};
+struct Particle {
+  /* 0x0000 */ packed_float3 position;
+  /* 0x000c */ float lifetime;
+  /* 0x0010 */ float4 color;
+  /* 0x0020 */ packed_float3 velocity;
+  /* 0x002c */ int8_t tint_pad_3[4];
+};
+struct Particles {
+  /* 0x0000 */ Particle particles[1];
+};
+struct UBO {
+  /* 0x0000 */ uint width;
+};
+struct Buffer {
+  /* 0x0000 */ float weights[1];
+};
+
+float rand(thread float2* const tint_symbol_9) {
+  (*(tint_symbol_9))[0] = fract((cos(dot(*(tint_symbol_9), float2(23.140779495f, 232.616897583f))) * 136.816802979f));
+  (*(tint_symbol_9))[1] = fract((cos(dot(*(tint_symbol_9), float2(54.478565216f, 345.841522217f))) * 534.764526367f));
+  return (*(tint_symbol_9))[1];
+}
+
+VertexOutput vs_main_inner(VertexInput in, const constant RenderParams* const tint_symbol_10) {
+  float3 quad_pos = (float2x3((*(tint_symbol_10)).right, (*(tint_symbol_10)).up) * in.quad_pos);
+  float3 position = (in.position + (quad_pos * 0.01f));
+  VertexOutput out = {};
+  out.position = ((*(tint_symbol_10)).modelViewProjectionMatrix * float4(position, 1.0f));
+  out.color = in.color;
+  out.quad_pos = in.quad_pos;
+  return out;
+}
+
+vertex tint_symbol_3 vs_main(const constant RenderParams* tint_symbol_11 [[buffer(0)]], tint_symbol_2 tint_symbol_1 [[stage_in]]) {
+  VertexInput const tint_symbol_7 = {.position=tint_symbol_1.position, .color=tint_symbol_1.color, .quad_pos=tint_symbol_1.quad_pos};
+  VertexOutput const inner_result = vs_main_inner(tint_symbol_7, tint_symbol_11);
+  tint_symbol_3 wrapper_result = {};
+  wrapper_result.position = inner_result.position;
+  wrapper_result.color = inner_result.color;
+  wrapper_result.quad_pos = inner_result.quad_pos;
+  return wrapper_result;
+}
+
+float4 fs_main_inner(VertexOutput in) {
+  float4 color = in.color;
+  color[3] = (color[3] * fmax((1.0f - length(in.quad_pos)), 0.0f));
+  return color;
+}
+
+fragment tint_symbol_6 fs_main(float4 position [[position]], tint_symbol_5 tint_symbol_4 [[stage_in]]) {
+  VertexOutput const tint_symbol_8 = {.position=position, .color=tint_symbol_4.color, .quad_pos=tint_symbol_4.quad_pos};
+  float4 const inner_result_1 = fs_main_inner(tint_symbol_8);
+  tint_symbol_6 wrapper_result_1 = {};
+  wrapper_result_1.value = inner_result_1;
+  return wrapper_result_1;
+}
+
+void simulate_inner(uint3 GlobalInvocationID, thread float2* const tint_symbol_12, const constant SimulationParams* const tint_symbol_13, device Particles* const tint_symbol_14, texture2d<float, access::sample> tint_symbol_15) {
+  *(tint_symbol_12) = ((float4((*(tint_symbol_13)).seed).xy + float2(uint3(GlobalInvocationID).xy)) * float4((*(tint_symbol_13)).seed).zw);
+  uint const idx = GlobalInvocationID[0];
+  Particle particle = (*(tint_symbol_14)).particles[idx];
+  particle.velocity[2] = (particle.velocity[2] - ((*(tint_symbol_13)).deltaTime * 0.5f));
+  particle.position = (particle.position + ((*(tint_symbol_13)).deltaTime * particle.velocity));
+  particle.lifetime = (particle.lifetime - (*(tint_symbol_13)).deltaTime);
+  particle.color[3] = smoothstep(0.0f, 0.5f, particle.lifetime);
+  if ((particle.lifetime < 0.0f)) {
+    int2 coord = int2(0, 0);
+    for(int level = as_type<int>((as_type<uint>(int(tint_symbol_15.get_num_mip_levels())) - as_type<uint>(1))); (level > 0); level = as_type<int>((as_type<uint>(level) - as_type<uint>(1)))) {
+      float4 const probabilites = tint_symbol_15.read(uint2(coord), level);
+      float4 const value = float4(rand(tint_symbol_12));
+      bool4 const mask = ((value >= float4(0.0f, float4(probabilites).xyz)) & (value < probabilites));
+      coord = as_type<int2>((as_type<uint2>(coord) * as_type<uint>(2)));
+      coord[0] = as_type<int>((as_type<uint>(coord[0]) + as_type<uint>(select(0, 1, any(bool4(mask).yw)))));
+      coord[1] = as_type<int>((as_type<uint>(coord[1]) + as_type<uint>(select(0, 1, any(bool4(mask).zw)))));
+    }
+    float2 const uv = (float2(coord) / float2(int2(tint_symbol_15.get_width(), tint_symbol_15.get_height())));
+    particle.position = float3((((uv - 0.5f) * 3.0f) * float2(1.0f, -1.0f)), 0.0f);
+    particle.color = tint_symbol_15.read(uint2(coord), 0);
+    particle.velocity[0] = ((rand(tint_symbol_12) - 0.5f) * 0.100000001f);
+    particle.velocity[1] = ((rand(tint_symbol_12) - 0.5f) * 0.100000001f);
+    particle.velocity[2] = (rand(tint_symbol_12) * 0.300000012f);
+    particle.lifetime = (0.5f + (rand(tint_symbol_12) * 2.0f));
+  }
+  (*(tint_symbol_14)).particles[idx] = particle;
+}
+
+kernel void simulate(const constant SimulationParams* tint_symbol_17 [[buffer(0)]], device Particles* tint_symbol_18 [[buffer(1)]], texture2d<float, access::sample> tint_symbol_19 [[texture(0)]], uint3 GlobalInvocationID [[thread_position_in_grid]]) {
+  thread float2 tint_symbol_16 = 0.0f;
+  simulate_inner(GlobalInvocationID, &(tint_symbol_16), tint_symbol_17, tint_symbol_18, tint_symbol_19);
+  return;
+}
+
+void import_level_inner(uint3 coord, const constant UBO* const tint_symbol_20, device Buffer* const tint_symbol_21, texture2d<float, access::sample> tint_symbol_22) {
+  uint const offset = (coord[0] + (coord[1] * (*(tint_symbol_20)).width));
+  (*(tint_symbol_21)).weights[offset] = tint_symbol_22.read(uint2(int2(uint3(coord).xy)), 0)[3];
+}
+
+kernel void import_level(const constant UBO* tint_symbol_23 [[buffer(2)]], device Buffer* tint_symbol_24 [[buffer(3)]], texture2d<float, access::sample> tint_symbol_25 [[texture(1)]], uint3 coord [[thread_position_in_grid]]) {
+  import_level_inner(coord, tint_symbol_23, tint_symbol_24, tint_symbol_25);
+  return;
+}
+
+void export_level_inner(uint3 coord, texture2d<float, access::write> tint_symbol_26, const constant UBO* const tint_symbol_27, const device Buffer* const tint_symbol_28, device Buffer* const tint_symbol_29) {
+  if (all((uint3(coord).xy < uint2(int2(tint_symbol_26.get_width(), tint_symbol_26.get_height()))))) {
+    uint const dst_offset = (coord[0] + (coord[1] * (*(tint_symbol_27)).width));
+    uint const src_offset = ((coord[0] * 2u) + ((coord[1] * 2u) * (*(tint_symbol_27)).width));
+    float const a_1 = (*(tint_symbol_28)).weights[(src_offset + 0u)];
+    float const b = (*(tint_symbol_28)).weights[(src_offset + 1u)];
+    float const c = (*(tint_symbol_28)).weights[((src_offset + 0u) + (*(tint_symbol_27)).width)];
+    float const d = (*(tint_symbol_28)).weights[((src_offset + 1u) + (*(tint_symbol_27)).width)];
+    float const sum = dot(float4(a_1, b, c, d), float4(1.0f));
+    (*(tint_symbol_29)).weights[dst_offset] = (sum / 4.0f);
+    float4 const probabilities = (float4(a_1, (a_1 + b), ((a_1 + b) + c), sum) / fmax(sum, 0.0001f));
+    tint_symbol_26.write(probabilities, uint2(int2(uint3(coord).xy)));
+  }
+}
+
+kernel void export_level(texture2d<float, access::write> tint_symbol_30 [[texture(2)]], const constant UBO* tint_symbol_31 [[buffer(2)]], const device Buffer* tint_symbol_32 [[buffer(4)]], device Buffer* tint_symbol_33 [[buffer(3)]], uint3 coord [[thread_position_in_grid]]) {
+  export_level_inner(coord, tint_symbol_30, tint_symbol_31, tint_symbol_32, tint_symbol_33);
+  return;
+}
+

test/benchmark/particles.wgsl.expected.spvasm

@@ -0,0 +1,646 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 423
+; Schema: 0
+               OpCapability Shader
+               OpCapability ImageQuery
+         %67 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vs_main "vs_main" %position_1 %color_1 %quad_pos_1 %position_2 %color_2 %quad_pos_2 %vertex_point_size
+               OpEntryPoint Fragment %fs_main "fs_main" %position_3 %color_3 %quad_pos_3 %value_1
+               OpEntryPoint GLCompute %simulate "simulate" %GlobalInvocationID_1
+               OpEntryPoint GLCompute %import_level "import_level" %coord_1
+               OpEntryPoint GLCompute %export_level "export_level" %coord_2
+               OpExecutionMode %fs_main OriginUpperLeft
+               OpExecutionMode %simulate LocalSize 64 1 1
+               OpExecutionMode %import_level LocalSize 64 1 1
+               OpExecutionMode %export_level LocalSize 64 1 1
+               OpName %position_1 "position_1"
+               OpName %color_1 "color_1"
+               OpName %quad_pos_1 "quad_pos_1"
+               OpName %position_2 "position_2"
+               OpName %color_2 "color_2"
+               OpName %quad_pos_2 "quad_pos_2"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %position_3 "position_3"
+               OpName %color_3 "color_3"
+               OpName %quad_pos_3 "quad_pos_3"
+               OpName %value_1 "value_1"
+               OpName %GlobalInvocationID_1 "GlobalInvocationID_1"
+               OpName %coord_1 "coord_1"
+               OpName %coord_2 "coord_2"
+               OpName %rand_seed "rand_seed"
+               OpName %RenderParams "RenderParams"
+               OpMemberName %RenderParams 0 "modelViewProjectionMatrix"
+               OpMemberName %RenderParams 1 "right"
+               OpMemberName %RenderParams 2 "up"
+               OpName %render_params "render_params"
+               OpName %SimulationParams "SimulationParams"
+               OpMemberName %SimulationParams 0 "deltaTime"
+               OpMemberName %SimulationParams 1 "seed"
+               OpName %sim_params "sim_params"
+               OpName %Particles "Particles"
+               OpMemberName %Particles 0 "particles"
+               OpName %Particle "Particle"
+               OpMemberName %Particle 0 "position"
+               OpMemberName %Particle 1 "lifetime"
+               OpMemberName %Particle 2 "color"
+               OpMemberName %Particle 3 "velocity"
+               OpName %data "data"
+               OpName %texture "texture"
+               OpName %UBO "UBO"
+               OpMemberName %UBO 0 "width"
+               OpName %ubo "ubo"
+               OpName %Buffer "Buffer"
+               OpMemberName %Buffer 0 "weights"
+               OpName %buf_in "buf_in"
+               OpName %buf_out "buf_out"
+               OpName %tex_in "tex_in"
+               OpName %tex_out "tex_out"
+               OpName %rand "rand"
+               OpName %VertexOutput "VertexOutput"
+               OpMemberName %VertexOutput 0 "position"
+               OpMemberName %VertexOutput 1 "color"
+               OpMemberName %VertexOutput 2 "quad_pos"
+               OpName %VertexInput "VertexInput"
+               OpMemberName %VertexInput 0 "position"
+               OpMemberName %VertexInput 1 "color"
+               OpMemberName %VertexInput 2 "quad_pos"
+               OpName %vs_main_inner "vs_main_inner"
+               OpName %in "in"
+               OpName %quad_pos "quad_pos"
+               OpName %position "position"
+               OpName %out "out"
+               OpName %vs_main "vs_main"
+               OpName %fs_main_inner "fs_main_inner"
+               OpName %in_0 "in"
+               OpName %color "color"
+               OpName %fs_main "fs_main"
+               OpName %simulate_inner "simulate_inner"
+               OpName %GlobalInvocationID "GlobalInvocationID"
+               OpName %particle "particle"
+               OpName %coord "coord"
+               OpName %level "level"
+               OpName %simulate "simulate"
+               OpName %import_level_inner "import_level_inner"
+               OpName %coord_0 "coord"
+               OpName %import_level "import_level"
+               OpName %export_level_inner "export_level_inner"
+               OpName %coord_3 "coord"
+               OpName %export_level "export_level"
+               OpDecorate %position_1 Location 0
+               OpDecorate %color_1 Location 1
+               OpDecorate %quad_pos_1 Location 2
+               OpDecorate %position_2 BuiltIn Position
+               OpDecorate %color_2 Location 0
+               OpDecorate %quad_pos_2 Location 1
+               OpDecorate %vertex_point_size BuiltIn PointSize
+               OpDecorate %position_3 BuiltIn FragCoord
+               OpDecorate %color_3 Location 0
+               OpDecorate %quad_pos_3 Location 1
+               OpDecorate %value_1 Location 0
+               OpDecorate %GlobalInvocationID_1 BuiltIn GlobalInvocationId
+               OpDecorate %coord_1 BuiltIn GlobalInvocationId
+               OpDecorate %coord_2 BuiltIn GlobalInvocationId
+               OpDecorate %RenderParams Block
+               OpMemberDecorate %RenderParams 0 Offset 0
+               OpMemberDecorate %RenderParams 0 ColMajor
+               OpMemberDecorate %RenderParams 0 MatrixStride 16
+               OpMemberDecorate %RenderParams 1 Offset 64
+               OpMemberDecorate %RenderParams 2 Offset 80
+               OpDecorate %render_params NonWritable
+               OpDecorate %render_params Binding 0
+               OpDecorate %render_params DescriptorSet 0
+               OpDecorate %SimulationParams Block
+               OpMemberDecorate %SimulationParams 0 Offset 0
+               OpMemberDecorate %SimulationParams 1 Offset 16
+               OpDecorate %sim_params NonWritable
+               OpDecorate %sim_params Binding 0
+               OpDecorate %sim_params DescriptorSet 0
+               OpDecorate %Particles Block
+               OpMemberDecorate %Particles 0 Offset 0
+               OpMemberDecorate %Particle 0 Offset 0
+               OpMemberDecorate %Particle 1 Offset 12
+               OpMemberDecorate %Particle 2 Offset 16
+               OpMemberDecorate %Particle 3 Offset 32
+               OpDecorate %_runtimearr_Particle ArrayStride 48
+               OpDecorate %data Binding 1
+               OpDecorate %data DescriptorSet 0
+               OpDecorate %texture Binding 2
+               OpDecorate %texture DescriptorSet 0
+               OpDecorate %UBO Block
+               OpMemberDecorate %UBO 0 Offset 0
+               OpDecorate %ubo NonWritable
+               OpDecorate %ubo Binding 3
+               OpDecorate %ubo DescriptorSet 0
+               OpDecorate %Buffer Block
+               OpMemberDecorate %Buffer 0 Offset 0
+               OpDecorate %_runtimearr_float ArrayStride 4
+               OpDecorate %buf_in NonWritable
+               OpDecorate %buf_in Binding 4
+               OpDecorate %buf_in DescriptorSet 0
+               OpDecorate %buf_out Binding 5
+               OpDecorate %buf_out DescriptorSet 0
+               OpDecorate %tex_in Binding 6
+               OpDecorate %tex_in DescriptorSet 0
+               OpDecorate %tex_out NonReadable
+               OpDecorate %tex_out Binding 7
+               OpDecorate %tex_out DescriptorSet 0
+               OpMemberDecorate %VertexOutput 0 Offset 0
+               OpMemberDecorate %VertexOutput 1 Offset 16
+               OpMemberDecorate %VertexOutput 2 Offset 32
+               OpMemberDecorate %VertexInput 0 Offset 0
+               OpMemberDecorate %VertexInput 1 Offset 16
+               OpMemberDecorate %VertexInput 2 Offset 32
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+%_ptr_Input_v3float = OpTypePointer Input %v3float
+ %position_1 = OpVariable %_ptr_Input_v3float Input
+    %v4float = OpTypeVector %float 4
+%_ptr_Input_v4float = OpTypePointer Input %v4float
+    %color_1 = OpVariable %_ptr_Input_v4float Input
+    %v2float = OpTypeVector %float 2
+%_ptr_Input_v2float = OpTypePointer Input %v2float
+ %quad_pos_1 = OpVariable %_ptr_Input_v2float Input
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+         %13 = OpConstantNull %v4float
+ %position_2 = OpVariable %_ptr_Output_v4float Output %13
+    %color_2 = OpVariable %_ptr_Output_v4float Output %13
+%_ptr_Output_v2float = OpTypePointer Output %v2float
+         %17 = OpConstantNull %v2float
+ %quad_pos_2 = OpVariable %_ptr_Output_v2float Output %17
+%_ptr_Output_float = OpTypePointer Output %float
+         %20 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %20
+ %position_3 = OpVariable %_ptr_Input_v4float Input
+    %color_3 = OpVariable %_ptr_Input_v4float Input
+ %quad_pos_3 = OpVariable %_ptr_Input_v2float Input
+    %value_1 = OpVariable %_ptr_Output_v4float Output %13
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%GlobalInvocationID_1 = OpVariable %_ptr_Input_v3uint Input
+    %coord_1 = OpVariable %_ptr_Input_v3uint Input
+    %coord_2 = OpVariable %_ptr_Input_v3uint Input
+%_ptr_Private_v2float = OpTypePointer Private %v2float
+  %rand_seed = OpVariable %_ptr_Private_v2float Private %17
+%mat4v4float = OpTypeMatrix %v4float 4
+%RenderParams = OpTypeStruct %mat4v4float %v3float %v3float
+%_ptr_Uniform_RenderParams = OpTypePointer Uniform %RenderParams
+%render_params = OpVariable %_ptr_Uniform_RenderParams Uniform
+%SimulationParams = OpTypeStruct %float %v4float
+%_ptr_Uniform_SimulationParams = OpTypePointer Uniform %SimulationParams
+ %sim_params = OpVariable %_ptr_Uniform_SimulationParams Uniform
+   %Particle = OpTypeStruct %v3float %float %v4float %v3float
+%_runtimearr_Particle = OpTypeRuntimeArray %Particle
+  %Particles = OpTypeStruct %_runtimearr_Particle
+%_ptr_StorageBuffer_Particles = OpTypePointer StorageBuffer %Particles
+       %data = OpVariable %_ptr_StorageBuffer_Particles StorageBuffer
+         %47 = OpTypeImage %float 2D 0 0 0 1 Unknown
+%_ptr_UniformConstant_47 = OpTypePointer UniformConstant %47
+    %texture = OpVariable %_ptr_UniformConstant_47 UniformConstant
+        %UBO = OpTypeStruct %uint
+%_ptr_Uniform_UBO = OpTypePointer Uniform %UBO
+        %ubo = OpVariable %_ptr_Uniform_UBO Uniform
+%_runtimearr_float = OpTypeRuntimeArray %float
+     %Buffer = OpTypeStruct %_runtimearr_float
+%_ptr_StorageBuffer_Buffer = OpTypePointer StorageBuffer %Buffer
+     %buf_in = OpVariable %_ptr_StorageBuffer_Buffer StorageBuffer
+    %buf_out = OpVariable %_ptr_StorageBuffer_Buffer StorageBuffer
+     %tex_in = OpVariable %_ptr_UniformConstant_47 UniformConstant
+         %59 = OpTypeImage %float 2D 0 0 0 2 Rgba8
+%_ptr_UniformConstant_59 = OpTypePointer UniformConstant %59
+    %tex_out = OpVariable %_ptr_UniformConstant_59 UniformConstant
+         %60 = OpTypeFunction %float
+     %uint_0 = OpConstant %uint 0
+%_ptr_Private_float = OpTypePointer Private %float
+%float_23_1407795 = OpConstant %float 23.1407795
+%float_232_616898 = OpConstant %float 232.616898
+         %73 = OpConstantComposite %v2float %float_23_1407795 %float_232_616898
+%float_136_816803 = OpConstant %float 136.816803
+     %uint_1 = OpConstant %uint 1
+%float_54_4785652 = OpConstant %float 54.4785652
+%float_345_841522 = OpConstant %float 345.841522
+         %84 = OpConstantComposite %v2float %float_54_4785652 %float_345_841522
+%float_534_764526 = OpConstant %float 534.764526
+%VertexOutput = OpTypeStruct %v4float %v4float %v2float
+%VertexInput = OpTypeStruct %v3float %v4float %v2float
+         %89 = OpTypeFunction %VertexOutput %VertexInput
+%mat2v3float = OpTypeMatrix %v3float 2
+%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float
+     %uint_2 = OpConstant %uint 2
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+        %107 = OpConstantNull %v3float
+%float_0_00999999978 = OpConstant %float 0.00999999978
+%_ptr_Function_VertexOutput = OpTypePointer Function %VertexOutput
+        %116 = OpConstantNull %VertexOutput
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+%_ptr_Uniform_mat4v4float = OpTypePointer Uniform %mat4v4float
+    %float_1 = OpConstant %float 1
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+       %void = OpTypeVoid
+        %135 = OpTypeFunction %void
+        %147 = OpTypeFunction %v4float %VertexOutput
+     %uint_3 = OpConstant %uint 3
+%_ptr_Function_float = OpTypePointer Function %float
+    %float_0 = OpConstant %float 0
+        %172 = OpTypeFunction %void %v3uint
+%_ptr_Uniform_v4float = OpTypePointer Uniform %v4float
+     %v2uint = OpTypeVector %uint 2
+%_ptr_StorageBuffer_Particle = OpTypePointer StorageBuffer %Particle
+%_ptr_Function_Particle = OpTypePointer Function %Particle
+        %194 = OpConstantNull %Particle
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+  %float_0_5 = OpConstant %float 0.5
+       %bool = OpTypeBool
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+      %int_0 = OpConstant %int 0
+        %232 = OpConstantComposite %v2int %int_0 %int_0
+%_ptr_Function_v2int = OpTypePointer Function %v2int
+        %235 = OpConstantNull %v2int
+      %int_1 = OpConstant %int 1
+%_ptr_Function_int = OpTypePointer Function %int
+        %242 = OpConstantNull %int
+     %v4bool = OpTypeVector %bool 4
+      %int_2 = OpConstant %int 2
+     %v2bool = OpTypeVector %bool 2
+    %float_3 = OpConstant %float 3
+   %float_n1 = OpConstant %float -1
+        %302 = OpConstantComposite %v2float %float_1 %float_n1
+%float_0_100000001 = OpConstant %float 0.100000001
+%float_0_300000012 = OpConstant %float 0.300000012
+    %float_2 = OpConstant %float 2
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+        %402 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
+    %float_4 = OpConstant %float 4
+%float_9_99999975en05 = OpConstant %float 9.99999975e-05
+       %rand = OpFunction %float None %60
+         %62 = OpLabel
+         %65 = OpAccessChain %_ptr_Private_float %rand_seed %uint_0
+         %70 = OpLoad %v2float %rand_seed
+         %69 = OpDot %float %70 %73
+         %68 = OpExtInst %float %67 Cos %69
+         %75 = OpFMul %float %68 %float_136_816803
+         %66 = OpExtInst %float %67 Fract %75
+               OpStore %65 %66
+         %77 = OpAccessChain %_ptr_Private_float %rand_seed %uint_1
+         %81 = OpLoad %v2float %rand_seed
+         %80 = OpDot %float %81 %84
+         %79 = OpExtInst %float %67 Cos %80
+         %86 = OpFMul %float %79 %float_534_764526
+         %78 = OpExtInst %float %67 Fract %86
+               OpStore %77 %78
+         %87 = OpAccessChain %_ptr_Private_float %rand_seed %uint_1
+         %88 = OpLoad %float %87
+               OpReturnValue %88
+               OpFunctionEnd
+%vs_main_inner = OpFunction %VertexOutput None %89
+         %in = OpFunctionParameter %VertexInput
+         %94 = OpLabel
+   %quad_pos = OpVariable %_ptr_Function_v3float Function %107
+   %position = OpVariable %_ptr_Function_v3float Function %107
+        %out = OpVariable %_ptr_Function_VertexOutput Function %116
+         %97 = OpAccessChain %_ptr_Uniform_v3float %render_params %uint_1
+         %98 = OpLoad %v3float %97
+        %100 = OpAccessChain %_ptr_Uniform_v3float %render_params %uint_2
+        %101 = OpLoad %v3float %100
+        %102 = OpCompositeConstruct %mat2v3float %98 %101
+        %103 = OpCompositeExtract %v2float %in 2
+        %104 = OpMatrixTimesVector %v3float %102 %103
+               OpStore %quad_pos %104
+        %108 = OpCompositeExtract %v3float %in 0
+        %109 = OpLoad %v3float %quad_pos
+        %111 = OpVectorTimesScalar %v3float %109 %float_0_00999999978
+        %112 = OpFAdd %v3float %108 %111
+               OpStore %position %112
+        %118 = OpAccessChain %_ptr_Function_v4float %out %uint_0
+        %120 = OpAccessChain %_ptr_Uniform_mat4v4float %render_params %uint_0
+        %121 = OpLoad %mat4v4float %120
+        %122 = OpLoad %v3float %position
+        %123 = OpCompositeExtract %float %122 0
+        %124 = OpCompositeExtract %float %122 1
+        %125 = OpCompositeExtract %float %122 2
+        %127 = OpCompositeConstruct %v4float %123 %124 %125 %float_1
+        %128 = OpMatrixTimesVector %v4float %121 %127
+               OpStore %118 %128
+        %129 = OpAccessChain %_ptr_Function_v4float %out %uint_1
+        %130 = OpCompositeExtract %v4float %in 1
+               OpStore %129 %130
+        %132 = OpAccessChain %_ptr_Function_v2float %out %uint_2
+        %133 = OpCompositeExtract %v2float %in 2
+               OpStore %132 %133
+        %134 = OpLoad %VertexOutput %out
+               OpReturnValue %134
+               OpFunctionEnd
+    %vs_main = OpFunction %void None %135
+        %138 = OpLabel
+        %140 = OpLoad %v3float %position_1
+        %141 = OpLoad %v4float %color_1
+        %142 = OpLoad %v2float %quad_pos_1
+        %143 = OpCompositeConstruct %VertexInput %140 %141 %142
+        %139 = OpFunctionCall %VertexOutput %vs_main_inner %143
+        %144 = OpCompositeExtract %v4float %139 0
+               OpStore %position_2 %144
+        %145 = OpCompositeExtract %v4float %139 1
+               OpStore %color_2 %145
+        %146 = OpCompositeExtract %v2float %139 2
+               OpStore %quad_pos_2 %146
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fs_main_inner = OpFunction %v4float None %147
+       %in_0 = OpFunctionParameter %VertexOutput
+        %150 = OpLabel
+      %color = OpVariable %_ptr_Function_v4float Function %13
+        %151 = OpCompositeExtract %v4float %in_0 1
+               OpStore %color %151
+        %155 = OpAccessChain %_ptr_Function_float %color %uint_3
+        %156 = OpAccessChain %_ptr_Function_float %color %uint_3
+        %157 = OpLoad %float %156
+        %160 = OpCompositeExtract %v2float %in_0 2
+        %159 = OpExtInst %float %67 Length %160
+        %161 = OpFSub %float %float_1 %159
+        %158 = OpExtInst %float %67 NMax %161 %float_0
+        %163 = OpFMul %float %157 %158
+               OpStore %155 %163
+        %164 = OpLoad %v4float %color
+               OpReturnValue %164
+               OpFunctionEnd
+    %fs_main = OpFunction %void None %135
+        %166 = OpLabel
+        %168 = OpLoad %v4float %position_3
+        %169 = OpLoad %v4float %color_3
+        %170 = OpLoad %v2float %quad_pos_3
+        %171 = OpCompositeConstruct %VertexOutput %168 %169 %170
+        %167 = OpFunctionCall %v4float %fs_main_inner %171
+               OpStore %value_1 %167
+               OpReturn
+               OpFunctionEnd
+%simulate_inner = OpFunction %void None %172
+%GlobalInvocationID = OpFunctionParameter %v3uint
+        %175 = OpLabel
+   %particle = OpVariable %_ptr_Function_Particle Function %194
+      %coord = OpVariable %_ptr_Function_v2int Function %235
+      %level = OpVariable %_ptr_Function_int Function %242
+        %270 = OpVariable %_ptr_Function_v2int Function %235
+        %297 = OpVariable %_ptr_Function_v2float Function %17
+        %177 = OpAccessChain %_ptr_Uniform_v4float %sim_params %uint_1
+        %178 = OpLoad %v4float %177
+        %179 = OpVectorShuffle %v2float %178 %178 0 1
+        %182 = OpVectorShuffle %v2uint %GlobalInvocationID %GlobalInvocationID 0 1
+        %180 = OpConvertUToF %v2float %182
+        %183 = OpFAdd %v2float %179 %180
+        %184 = OpAccessChain %_ptr_Uniform_v4float %sim_params %uint_1
+        %185 = OpLoad %v4float %184
+        %186 = OpVectorShuffle %v2float %185 %185 2 3
+        %187 = OpFMul %v2float %183 %186
+               OpStore %rand_seed %187
+        %188 = OpCompositeExtract %uint %GlobalInvocationID 0
+        %190 = OpAccessChain %_ptr_StorageBuffer_Particle %data %uint_0 %188
+        %191 = OpLoad %Particle %190
+               OpStore %particle %191
+        %195 = OpAccessChain %_ptr_Function_float %particle %uint_3 %uint_2
+        %196 = OpAccessChain %_ptr_Function_float %particle %uint_3 %uint_2
+        %197 = OpLoad %float %196
+        %199 = OpAccessChain %_ptr_Uniform_float %sim_params %uint_0
+        %200 = OpLoad %float %199
+        %202 = OpFMul %float %200 %float_0_5
+        %203 = OpFSub %float %197 %202
+               OpStore %195 %203
+        %204 = OpAccessChain %_ptr_Function_v3float %particle %uint_0
+        %205 = OpAccessChain %_ptr_Function_v3float %particle %uint_0
+        %206 = OpLoad %v3float %205
+        %207 = OpAccessChain %_ptr_Uniform_float %sim_params %uint_0
+        %208 = OpLoad %float %207
+        %209 = OpAccessChain %_ptr_Function_v3float %particle %uint_3
+        %210 = OpLoad %v3float %209
+        %211 = OpVectorTimesScalar %v3float %210 %208
+        %212 = OpFAdd %v3float %206 %211
+               OpStore %204 %212
+        %213 = OpAccessChain %_ptr_Function_float %particle %uint_1
+        %214 = OpAccessChain %_ptr_Function_float %particle %uint_1
+        %215 = OpLoad %float %214
+        %216 = OpAccessChain %_ptr_Uniform_float %sim_params %uint_0
+        %217 = OpLoad %float %216
+        %218 = OpFSub %float %215 %217
+               OpStore %213 %218
+        %219 = OpAccessChain %_ptr_Function_float %particle %uint_2 %uint_3
+        %221 = OpAccessChain %_ptr_Function_float %particle %uint_1
+        %222 = OpLoad %float %221
+        %220 = OpExtInst %float %67 SmoothStep %float_0 %float_0_5 %222
+               OpStore %219 %220
+        %223 = OpAccessChain %_ptr_Function_float %particle %uint_1
+        %224 = OpLoad %float %223
+        %225 = OpFOrdLessThan %bool %224 %float_0
+               OpSelectionMerge %227 None
+               OpBranchConditional %225 %228 %227
+        %228 = OpLabel
+               OpStore %coord %232
+        %237 = OpLoad %47 %texture
+        %236 = OpImageQueryLevels %int %237
+        %239 = OpISub %int %236 %int_1
+               OpStore %level %239
+               OpBranch %243
+        %243 = OpLabel
+               OpLoopMerge %244 %245 None
+               OpBranch %246
+        %246 = OpLabel
+        %248 = OpLoad %int %level
+        %249 = OpSGreaterThan %bool %248 %int_0
+        %247 = OpLogicalNot %bool %249
+               OpSelectionMerge %250 None
+               OpBranchConditional %247 %251 %250
+        %251 = OpLabel
+               OpBranch %244
+        %250 = OpLabel
+        %253 = OpLoad %47 %texture
+        %254 = OpLoad %v2int %coord
+        %255 = OpLoad %int %level
+        %252 = OpImageFetch %v4float %253 %254 Lod %255
+        %256 = OpFunctionCall %float %rand
+        %257 = OpCompositeConstruct %v4float %256 %256 %256 %256
+        %258 = OpVectorShuffle %v3float %252 %252 0 1 2
+        %259 = OpCompositeExtract %float %258 0
+        %260 = OpCompositeExtract %float %258 1
+        %261 = OpCompositeExtract %float %258 2
+        %262 = OpCompositeConstruct %v4float %float_0 %259 %260 %261
+        %263 = OpFOrdGreaterThanEqual %v4bool %257 %262
+        %265 = OpFOrdLessThan %v4bool %257 %252
+        %266 = OpLogicalAnd %v4bool %263 %265
+        %267 = OpLoad %v2int %coord
+        %271 = OpCompositeConstruct %v2int %int_2 %int_2
+        %269 = OpIMul %v2int %267 %271
+               OpStore %coord %269
+        %272 = OpAccessChain %_ptr_Function_int %coord %uint_0
+        %273 = OpAccessChain %_ptr_Function_int %coord %uint_0
+        %274 = OpLoad %int %273
+        %278 = OpVectorShuffle %v2bool %266 %266 1 3
+        %276 = OpAny %bool %278
+        %275 = OpSelect %int %276 %int_1 %int_0
+        %279 = OpIAdd %int %274 %275
+               OpStore %272 %279
+        %280 = OpAccessChain %_ptr_Function_int %coord %uint_1
+        %281 = OpAccessChain %_ptr_Function_int %coord %uint_1
+        %282 = OpLoad %int %281
+        %285 = OpVectorShuffle %v2bool %266 %266 2 3
+        %284 = OpAny %bool %285
+        %283 = OpSelect %int %284 %int_1 %int_0
+        %286 = OpIAdd %int %282 %283
+               OpStore %280 %286
+               OpBranch %245
+        %245 = OpLabel
+        %287 = OpLoad %int %level
+        %288 = OpISub %int %287 %int_1
+               OpStore %level %288
+               OpBranch %243
+        %244 = OpLabel
+        %290 = OpLoad %v2int %coord
+        %289 = OpConvertSToF %v2float %290
+        %293 = OpLoad %47 %texture
+        %292 = OpImageQuerySizeLod %v2int %293 %int_0
+        %291 = OpConvertSToF %v2float %292
+        %294 = OpFDiv %v2float %289 %291
+        %295 = OpAccessChain %_ptr_Function_v3float %particle %uint_0
+        %298 = OpCompositeConstruct %v2float %float_0_5 %float_0_5
+        %296 = OpFSub %v2float %294 %298
+        %300 = OpVectorTimesScalar %v2float %296 %float_3
+        %303 = OpFMul %v2float %300 %302
+        %304 = OpCompositeExtract %float %303 0
+        %305 = OpCompositeExtract %float %303 1
+        %306 = OpCompositeConstruct %v3float %304 %305 %float_0
+               OpStore %295 %306
+        %307 = OpAccessChain %_ptr_Function_v4float %particle %uint_2
+        %309 = OpLoad %47 %texture
+        %310 = OpLoad %v2int %coord
+        %308 = OpImageFetch %v4float %309 %310 Lod %int_0
+               OpStore %307 %308
+        %311 = OpAccessChain %_ptr_Function_float %particle %uint_3 %uint_0
+        %312 = OpFunctionCall %float %rand
+        %313 = OpFSub %float %312 %float_0_5
+        %315 = OpFMul %float %313 %float_0_100000001
+               OpStore %311 %315
+        %316 = OpAccessChain %_ptr_Function_float %particle %uint_3 %uint_1
+        %317 = OpFunctionCall %float %rand
+        %318 = OpFSub %float %317 %float_0_5
+        %319 = OpFMul %float %318 %float_0_100000001
+               OpStore %316 %319
+        %320 = OpAccessChain %_ptr_Function_float %particle %uint_3 %uint_2
+        %321 = OpFunctionCall %float %rand
+        %323 = OpFMul %float %321 %float_0_300000012
+               OpStore %320 %323
+        %324 = OpAccessChain %_ptr_Function_float %particle %uint_1
+        %325 = OpFunctionCall %float %rand
+        %327 = OpFMul %float %325 %float_2
+        %328 = OpFAdd %float %float_0_5 %327
+               OpStore %324 %328
+               OpBranch %227
+        %227 = OpLabel
+        %329 = OpAccessChain %_ptr_StorageBuffer_Particle %data %uint_0 %188
+        %330 = OpLoad %Particle %particle
+               OpStore %329 %330
+               OpReturn
+               OpFunctionEnd
+   %simulate = OpFunction %void None %135
+        %332 = OpLabel
+        %334 = OpLoad %v3uint %GlobalInvocationID_1
+        %333 = OpFunctionCall %void %simulate_inner %334
+               OpReturn
+               OpFunctionEnd
+%import_level_inner = OpFunction %void None %172
+    %coord_0 = OpFunctionParameter %v3uint
+        %337 = OpLabel
+        %339 = OpCompositeExtract %uint %coord_0 0
+        %340 = OpCompositeExtract %uint %coord_0 1
+        %342 = OpAccessChain %_ptr_Uniform_uint %ubo %uint_0
+        %343 = OpLoad %uint %342
+        %344 = OpIMul %uint %340 %343
+        %345 = OpIAdd %uint %339 %344
+        %347 = OpAccessChain %_ptr_StorageBuffer_float %buf_out %uint_0 %345
+        %349 = OpLoad %47 %tex_in
+        %351 = OpVectorShuffle %v2uint %coord_0 %coord_0 0 1
+        %350 = OpBitcast %v2int %351
+        %348 = OpImageFetch %v4float %349 %350 Lod %int_0
+        %352 = OpCompositeExtract %float %348 3
+               OpStore %347 %352
+               OpReturn
+               OpFunctionEnd
+%import_level = OpFunction %void None %135
+        %354 = OpLabel
+        %356 = OpLoad %v3uint %coord_1
+        %355 = OpFunctionCall %void %import_level_inner %356
+               OpReturn
+               OpFunctionEnd
+%export_level_inner = OpFunction %void None %172
+    %coord_3 = OpFunctionParameter %v3uint
+        %359 = OpLabel
+        %413 = OpVariable %_ptr_Function_v4float Function %13
+        %361 = OpVectorShuffle %v2uint %coord_3 %coord_3 0 1
+        %364 = OpLoad %59 %tex_out
+        %363 = OpImageQuerySize %v2int %364
+        %362 = OpBitcast %v2uint %363
+        %365 = OpULessThan %v2bool %361 %362
+        %360 = OpAll %bool %365
+               OpSelectionMerge %366 None
+               OpBranchConditional %360 %367 %366
+        %367 = OpLabel
+        %368 = OpCompositeExtract %uint %coord_3 0
+        %369 = OpCompositeExtract %uint %coord_3 1
+        %370 = OpAccessChain %_ptr_Uniform_uint %ubo %uint_0
+        %371 = OpLoad %uint %370
+        %372 = OpIMul %uint %369 %371
+        %373 = OpIAdd %uint %368 %372
+        %374 = OpCompositeExtract %uint %coord_3 0
+        %375 = OpIMul %uint %374 %uint_2
+        %376 = OpCompositeExtract %uint %coord_3 1
+        %377 = OpIMul %uint %376 %uint_2
+        %378 = OpAccessChain %_ptr_Uniform_uint %ubo %uint_0
+        %379 = OpLoad %uint %378
+        %380 = OpIMul %uint %377 %379
+        %381 = OpIAdd %uint %375 %380
+        %382 = OpIAdd %uint %381 %uint_0
+        %383 = OpAccessChain %_ptr_StorageBuffer_float %buf_in %uint_0 %382
+        %384 = OpLoad %float %383
+        %385 = OpIAdd %uint %381 %uint_1
+        %386 = OpAccessChain %_ptr_StorageBuffer_float %buf_in %uint_0 %385
+        %387 = OpLoad %float %386
+        %388 = OpIAdd %uint %381 %uint_0
+        %389 = OpAccessChain %_ptr_Uniform_uint %ubo %uint_0
+        %390 = OpLoad %uint %389
+        %391 = OpIAdd %uint %388 %390
+        %392 = OpAccessChain %_ptr_StorageBuffer_float %buf_in %uint_0 %391
+        %393 = OpLoad %float %392
+        %394 = OpIAdd %uint %381 %uint_1
+        %395 = OpAccessChain %_ptr_Uniform_uint %ubo %uint_0
+        %396 = OpLoad %uint %395
+        %397 = OpIAdd %uint %394 %396
+        %398 = OpAccessChain %_ptr_StorageBuffer_float %buf_in %uint_0 %397
+        %399 = OpLoad %float %398
+        %401 = OpCompositeConstruct %v4float %384 %387 %393 %399
+        %400 = OpDot %float %401 %402
+        %403 = OpAccessChain %_ptr_StorageBuffer_float %buf_out %uint_0 %373
+        %405 = OpFDiv %float %400 %float_4
+               OpStore %403 %405
+        %406 = OpFAdd %float %384 %387
+        %407 = OpFAdd %float %384 %387
+        %408 = OpFAdd %float %407 %393
+        %409 = OpCompositeConstruct %v4float %384 %406 %408 %400
+        %410 = OpExtInst %float %67 NMax %400 %float_9_99999975en05
+        %414 = OpCompositeConstruct %v4float %410 %410 %410 %410
+        %412 = OpFDiv %v4float %409 %414
+        %416 = OpLoad %59 %tex_out
+        %418 = OpVectorShuffle %v2uint %coord_3 %coord_3 0 1
+        %417 = OpBitcast %v2int %418
+               OpImageWrite %416 %417 %412
+               OpBranch %366
+        %366 = OpLabel
+               OpReturn
+               OpFunctionEnd
+%export_level = OpFunction %void None %135
+        %420 = OpLabel
+        %422 = OpLoad %v3uint %coord_2
+        %421 = OpFunctionCall %void %export_level_inner %422
+               OpReturn
+               OpFunctionEnd

test/benchmark/particles.wgsl.expected.wgsl

@@ -0,0 +1,144 @@
+var<private> rand_seed : vec2<f32>;
+
+fn rand() -> f32 {
+  rand_seed.x = fract((cos(dot(rand_seed, vec2<f32>(23.140779495, 232.616897583))) * 136.816802979));
+  rand_seed.y = fract((cos(dot(rand_seed, vec2<f32>(54.478565216, 345.841522217))) * 534.764526367));
+  return rand_seed.y;
+}
+
+struct RenderParams {
+  modelViewProjectionMatrix : mat4x4<f32>;
+  right : vec3<f32>;
+  up : vec3<f32>;
+};
+
+[[binding(0), group(0)]] var<uniform> render_params : RenderParams;
+
+struct VertexInput {
+  [[location(0)]]
+  position : vec3<f32>;
+  [[location(1)]]
+  color : vec4<f32>;
+  [[location(2)]]
+  quad_pos : vec2<f32>;
+};
+
+struct VertexOutput {
+  [[builtin(position)]]
+  position : vec4<f32>;
+  [[location(0)]]
+  color : vec4<f32>;
+  [[location(1)]]
+  quad_pos : vec2<f32>;
+};
+
+[[stage(vertex)]]
+fn vs_main(in : VertexInput) -> VertexOutput {
+  var quad_pos = (mat2x3<f32>(render_params.right, render_params.up) * in.quad_pos);
+  var position = (in.position + (quad_pos * 0.01));
+  var out : VertexOutput;
+  out.position = (render_params.modelViewProjectionMatrix * vec4<f32>(position, 1.0));
+  out.color = in.color;
+  out.quad_pos = in.quad_pos;
+  return out;
+}
+
+[[stage(fragment)]]
+fn fs_main(in : VertexOutput) -> [[location(0)]] vec4<f32> {
+  var color = in.color;
+  color.a = (color.a * max((1.0 - length(in.quad_pos)), 0.0));
+  return color;
+}
+
+struct SimulationParams {
+  deltaTime : f32;
+  seed : vec4<f32>;
+};
+
+struct Particle {
+  position : vec3<f32>;
+  lifetime : f32;
+  color : vec4<f32>;
+  velocity : vec3<f32>;
+};
+
+struct Particles {
+  particles : array<Particle>;
+};
+
+[[binding(0), group(0)]] var<uniform> sim_params : SimulationParams;
+
+[[binding(1), group(0)]] var<storage, read_write> data : Particles;
+
+[[binding(2), group(0)]] var texture : texture_2d<f32>;
+
+[[stage(compute), workgroup_size(64)]]
+fn simulate([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
+  rand_seed = ((sim_params.seed.xy + vec2<f32>(GlobalInvocationID.xy)) * sim_params.seed.zw);
+  let idx = GlobalInvocationID.x;
+  var particle = data.particles[idx];
+  particle.velocity.z = (particle.velocity.z - (sim_params.deltaTime * 0.5));
+  particle.position = (particle.position + (sim_params.deltaTime * particle.velocity));
+  particle.lifetime = (particle.lifetime - sim_params.deltaTime);
+  particle.color.a = smoothStep(0.0, 0.5, particle.lifetime);
+  if ((particle.lifetime < 0.0)) {
+    var coord = vec2<i32>(0, 0);
+    for(var level = (textureNumLevels(texture) - 1); (level > 0); level = (level - 1)) {
+      let probabilites = textureLoad(texture, coord, level);
+      let value = vec4<f32>(rand());
+      let mask = ((value >= vec4<f32>(0.0, probabilites.xyz)) & (value < probabilites));
+      coord = (coord * 2);
+      coord.x = (coord.x + select(0, 1, any(mask.yw)));
+      coord.y = (coord.y + select(0, 1, any(mask.zw)));
+    }
+    let uv = (vec2<f32>(coord) / vec2<f32>(textureDimensions(texture)));
+    particle.position = vec3<f32>((((uv - 0.5) * 3.0) * vec2<f32>(1.0, -1.0)), 0.0);
+    particle.color = textureLoad(texture, coord, 0);
+    particle.velocity.x = ((rand() - 0.5) * 0.100000001);
+    particle.velocity.y = ((rand() - 0.5) * 0.100000001);
+    particle.velocity.z = (rand() * 0.300000012);
+    particle.lifetime = (0.5 + (rand() * 2.0));
+  }
+  data.particles[idx] = particle;
+}
+
+struct UBO {
+  width : u32;
+};
+
+struct Buffer {
+  weights : array<f32>;
+};
+
+[[binding(3), group(0)]] var<uniform> ubo : UBO;
+
+[[binding(4), group(0)]] var<storage, read> buf_in : Buffer;
+
+[[binding(5), group(0)]] var<storage, read_write> buf_out : Buffer;
+
+[[binding(6), group(0)]] var tex_in : texture_2d<f32>;
+
+[[binding(7), group(0)]] var tex_out : texture_storage_2d<rgba8unorm, write>;
+
+[[stage(compute), workgroup_size(64)]]
+fn import_level([[builtin(global_invocation_id)]] coord : vec3<u32>) {
+  _ = &(buf_in);
+  let offset = (coord.x + (coord.y * ubo.width));
+  buf_out.weights[offset] = textureLoad(tex_in, vec2<i32>(coord.xy), 0).w;
+}
+
+[[stage(compute), workgroup_size(64)]]
+fn export_level([[builtin(global_invocation_id)]] coord : vec3<u32>) {
+  if (all((coord.xy < vec2<u32>(textureDimensions(tex_out))))) {
+    let dst_offset = (coord.x + (coord.y * ubo.width));
+    let src_offset = ((coord.x * 2u) + ((coord.y * 2u) * ubo.width));
+    let a = buf_in.weights[(src_offset + 0u)];
+    let b = buf_in.weights[(src_offset + 1u)];
+    let c = buf_in.weights[((src_offset + 0u) + ubo.width)];
+    let d = buf_in.weights[((src_offset + 1u) + ubo.width)];
+    let sum = dot(vec4<f32>(a, b, c, d), vec4<f32>(1.0));
+    buf_out.weights[dst_offset] = (sum / 4.0);
+    let probabilities = (vec4<f32>(a, (a + b), ((a + b) + c), sum) / max(sum, 0.0001));
+    textureStore(tex_out, vec2<i32>(coord.xy), probabilities);
+  }
+}

test/benchmark/simple_compute.wgsl

@@ -0,0 +1,10 @@
+struct SB {
+  data : array<i32>;
+};
+
+[[group(0), binding(0)]] var<storage, read_write> buffer : SB;
+
+[[stage(compute), workgroup_size(1, 2, 3)]]
+fn main([[builtin(global_invocation_id)]] id : vec3<u32>) {
+  buffer.data[id.x] = buffer.data[id.x] + 1;
+}

test/benchmark/simple_compute.wgsl.expected.hlsl

@@ -0,0 +1,15 @@
+RWByteAddressBuffer buffer : register(u0, space0);
+
+struct tint_symbol_1 {
+  uint3 id : SV_DispatchThreadID;
+};
+
+void main_inner(uint3 id) {
+  buffer.Store((4u * id.x), asuint((asint(buffer.Load((4u * id.x))) + 1)));
+}
+
+[numthreads(1, 2, 3)]
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.id);
+  return;
+}

test/benchmark/simple_compute.wgsl.expected.msl

@@ -0,0 +1,16 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct SB {
+  /* 0x0000 */ int data[1];
+};
+
+void tint_symbol_1_inner(uint3 id, device SB* const tint_symbol_2) {
+  (*(tint_symbol_2)).data[id[0]] = as_type<int>((as_type<uint>((*(tint_symbol_2)).data[id[0]]) + as_type<uint>(1)));
+}
+
+kernel void tint_symbol_1(device SB* tint_symbol_3 [[buffer(0)]], uint3 id [[thread_position_in_grid]]) {
+  tint_symbol_1_inner(id, tint_symbol_3);
+  return;
+}
+

test/benchmark/simple_compute.wgsl.expected.spvasm

@@ -0,0 +1,55 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 29
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main" %id_1
+               OpExecutionMode %main LocalSize 1 2 3
+               OpName %id_1 "id_1"
+               OpName %SB "SB"
+               OpMemberName %SB 0 "data"
+               OpName %buffer "buffer"
+               OpName %main_inner "main_inner"
+               OpName %id "id"
+               OpName %main "main"
+               OpDecorate %id_1 BuiltIn GlobalInvocationId
+               OpDecorate %SB Block
+               OpMemberDecorate %SB 0 Offset 0
+               OpDecorate %_runtimearr_int ArrayStride 4
+               OpDecorate %buffer DescriptorSet 0
+               OpDecorate %buffer Binding 0
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+       %id_1 = OpVariable %_ptr_Input_v3uint Input
+        %int = OpTypeInt 32 1
+%_runtimearr_int = OpTypeRuntimeArray %int
+         %SB = OpTypeStruct %_runtimearr_int
+%_ptr_StorageBuffer_SB = OpTypePointer StorageBuffer %SB
+     %buffer = OpVariable %_ptr_StorageBuffer_SB StorageBuffer
+       %void = OpTypeVoid
+         %10 = OpTypeFunction %void %v3uint
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+      %int_1 = OpConstant %int 1
+         %24 = OpTypeFunction %void
+ %main_inner = OpFunction %void None %10
+         %id = OpFunctionParameter %v3uint
+         %14 = OpLabel
+         %16 = OpCompositeExtract %uint %id 0
+         %18 = OpAccessChain %_ptr_StorageBuffer_int %buffer %uint_0 %16
+         %19 = OpCompositeExtract %uint %id 0
+         %20 = OpAccessChain %_ptr_StorageBuffer_int %buffer %uint_0 %19
+         %21 = OpLoad %int %20
+         %23 = OpIAdd %int %21 %int_1
+               OpStore %18 %23
+               OpReturn
+               OpFunctionEnd
+       %main = OpFunction %void None %24
+         %26 = OpLabel
+         %28 = OpLoad %v3uint %id_1
+         %27 = OpFunctionCall %void %main_inner %28
+               OpReturn
+               OpFunctionEnd

test/benchmark/simple_compute.wgsl.expected.wgsl

@@ -0,0 +1,10 @@
+struct SB {
+  data : array<i32>;
+};
+
+[[group(0), binding(0)]] var<storage, read_write> buffer : SB;
+
+[[stage(compute), workgroup_size(1, 2, 3)]]
+fn main([[builtin(global_invocation_id)]] id : vec3<u32>) {
+  buffer.data[id.x] = (buffer.data[id.x] + 1);
+}

test/benchmark/simple_fragment.wgsl

@@ -0,0 +1,12 @@
+struct Input {
+  [[location(0)]] color: vec4<f32>;
+};
+
+struct Output {
+  [[location(0)]] color: vec4<f32>;
+};
+
+[[stage(fragment)]]
+fn main(in : Input) -> Output {
+  return Output(in.color);
+}

test/benchmark/simple_fragment.wgsl.expected.hlsl

@@ -0,0 +1,25 @@
+struct Input {
+  float4 color;
+};
+struct Output {
+  float4 color;
+};
+struct tint_symbol_2 {
+  float4 color : TEXCOORD0;
+};
+struct tint_symbol_3 {
+  float4 color : SV_Target0;
+};
+
+Output main_inner(Input tint_symbol) {
+  const Output tint_symbol_4 = {tint_symbol.color};
+  return tint_symbol_4;
+}
+
+tint_symbol_3 main(tint_symbol_2 tint_symbol_1) {
+  const Input tint_symbol_5 = {tint_symbol_1.color};
+  const Output inner_result = main_inner(tint_symbol_5);
+  tint_symbol_3 wrapper_result = (tint_symbol_3)0;
+  wrapper_result.color = inner_result.color;
+  return wrapper_result;
+}

test/benchmark/simple_fragment.wgsl.expected.msl

@@ -0,0 +1,29 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Input {
+  float4 color;
+};
+struct Output {
+  float4 color;
+};
+struct tint_symbol_2 {
+  float4 color [[user(locn0)]];
+};
+struct tint_symbol_3 {
+  float4 color [[color(0)]];
+};
+
+Output tint_symbol_inner(Input in) {
+  Output const tint_symbol_4 = {.color=in.color};
+  return tint_symbol_4;
+}
+
+fragment tint_symbol_3 tint_symbol(tint_symbol_2 tint_symbol_1 [[stage_in]]) {
+  Input const tint_symbol_5 = {.color=tint_symbol_1.color};
+  Output const inner_result = tint_symbol_inner(tint_symbol_5);
+  tint_symbol_3 wrapper_result = {};
+  wrapper_result.color = inner_result.color;
+  return wrapper_result;
+}
+

test/benchmark/simple_fragment.wgsl.expected.spvasm

@@ -0,0 +1,50 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 24
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %main "main" %color_1 %color_2
+               OpExecutionMode %main OriginUpperLeft
+               OpName %color_1 "color_1"
+               OpName %color_2 "color_2"
+               OpName %Output "Output"
+               OpMemberName %Output 0 "color"
+               OpName %Input "Input"
+               OpMemberName %Input 0 "color"
+               OpName %main_inner "main_inner"
+               OpName %in "in"
+               OpName %main "main"
+               OpDecorate %color_1 Location 0
+               OpDecorate %color_2 Location 0
+               OpMemberDecorate %Output 0 Offset 0
+               OpMemberDecorate %Input 0 Offset 0
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Input_v4float = OpTypePointer Input %v4float
+    %color_1 = OpVariable %_ptr_Input_v4float Input
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %7 = OpConstantNull %v4float
+    %color_2 = OpVariable %_ptr_Output_v4float Output %7
+     %Output = OpTypeStruct %v4float
+      %Input = OpTypeStruct %v4float
+          %8 = OpTypeFunction %Output %Input
+       %void = OpTypeVoid
+         %16 = OpTypeFunction %void
+ %main_inner = OpFunction %Output None %8
+         %in = OpFunctionParameter %Input
+         %13 = OpLabel
+         %14 = OpCompositeExtract %v4float %in 0
+         %15 = OpCompositeConstruct %Output %14
+               OpReturnValue %15
+               OpFunctionEnd
+       %main = OpFunction %void None %16
+         %19 = OpLabel
+         %21 = OpLoad %v4float %color_1
+         %22 = OpCompositeConstruct %Input %21
+         %20 = OpFunctionCall %Output %main_inner %22
+         %23 = OpCompositeExtract %v4float %20 0
+               OpStore %color_2 %23
+               OpReturn
+               OpFunctionEnd

test/benchmark/simple_fragment.wgsl.expected.wgsl

@@ -0,0 +1,14 @@
+struct Input {
+  [[location(0)]]
+  color : vec4<f32>;
+};
+
+struct Output {
+  [[location(0)]]
+  color : vec4<f32>;
+};
+
+[[stage(fragment)]]
+fn main(in : Input) -> Output {
+  return Output(in.color);
+}

test/benchmark/simple_vertex.wgsl

@@ -0,0 +1,12 @@
+struct Input {
+  [[location(0)]] position: vec4<f32>;
+};
+
+struct Output {
+  [[builtin(position)]] position : vec4<f32>;
+};
+
+[[stage(vertex)]]
+fn main(in : Input) -> Output {
+  return Output(in.position);
+}

test/benchmark/simple_vertex.wgsl.expected.hlsl

@@ -0,0 +1,25 @@
+struct Input {
+  float4 position;
+};
+struct Output {
+  float4 position;
+};
+struct tint_symbol_2 {
+  float4 position : TEXCOORD0;
+};
+struct tint_symbol_3 {
+  float4 position : SV_Position;
+};
+
+Output main_inner(Input tint_symbol) {
+  const Output tint_symbol_4 = {tint_symbol.position};
+  return tint_symbol_4;
+}
+
+tint_symbol_3 main(tint_symbol_2 tint_symbol_1) {
+  const Input tint_symbol_5 = {tint_symbol_1.position};
+  const Output inner_result = main_inner(tint_symbol_5);
+  tint_symbol_3 wrapper_result = (tint_symbol_3)0;
+  wrapper_result.position = inner_result.position;
+  return wrapper_result;
+}

test/benchmark/simple_vertex.wgsl.expected.msl

@@ -0,0 +1,29 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Input {
+  float4 position;
+};
+struct Output {
+  float4 position;
+};
+struct tint_symbol_2 {
+  float4 position [[attribute(0)]];
+};
+struct tint_symbol_3 {
+  float4 position [[position]];
+};
+
+Output tint_symbol_inner(Input in) {
+  Output const tint_symbol_4 = {.position=in.position};
+  return tint_symbol_4;
+}
+
+vertex tint_symbol_3 tint_symbol(tint_symbol_2 tint_symbol_1 [[stage_in]]) {
+  Input const tint_symbol_5 = {.position=tint_symbol_1.position};
+  Output const inner_result = tint_symbol_inner(tint_symbol_5);
+  tint_symbol_3 wrapper_result = {};
+  wrapper_result.position = inner_result.position;
+  return wrapper_result;
+}
+

test/benchmark/simple_vertex.wgsl.expected.spvasm

@@ -0,0 +1,56 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 28
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %main "main" %position_1 %position_2 %vertex_point_size
+               OpName %position_1 "position_1"
+               OpName %position_2 "position_2"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %Output "Output"
+               OpMemberName %Output 0 "position"
+               OpName %Input "Input"
+               OpMemberName %Input 0 "position"
+               OpName %main_inner "main_inner"
+               OpName %in "in"
+               OpName %main "main"
+               OpDecorate %position_1 Location 0
+               OpDecorate %position_2 BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+               OpMemberDecorate %Output 0 Offset 0
+               OpMemberDecorate %Input 0 Offset 0
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Input_v4float = OpTypePointer Input %v4float
+ %position_1 = OpVariable %_ptr_Input_v4float Input
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %7 = OpConstantNull %v4float
+ %position_2 = OpVariable %_ptr_Output_v4float Output %7
+%_ptr_Output_float = OpTypePointer Output %float
+         %10 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %10
+     %Output = OpTypeStruct %v4float
+      %Input = OpTypeStruct %v4float
+         %11 = OpTypeFunction %Output %Input
+       %void = OpTypeVoid
+         %19 = OpTypeFunction %void
+    %float_1 = OpConstant %float 1
+ %main_inner = OpFunction %Output None %11
+         %in = OpFunctionParameter %Input
+         %16 = OpLabel
+         %17 = OpCompositeExtract %v4float %in 0
+         %18 = OpCompositeConstruct %Output %17
+               OpReturnValue %18
+               OpFunctionEnd
+       %main = OpFunction %void None %19
+         %22 = OpLabel
+         %24 = OpLoad %v4float %position_1
+         %25 = OpCompositeConstruct %Input %24
+         %23 = OpFunctionCall %Output %main_inner %25
+         %26 = OpCompositeExtract %v4float %23 0
+               OpStore %position_2 %26
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd

test/benchmark/simple_vertex.wgsl.expected.wgsl

@@ -0,0 +1,14 @@
+struct Input {
+  [[location(0)]]
+  position : vec4<f32>;
+};
+
+struct Output {
+  [[builtin(position)]]
+  position : vec4<f32>;
+};
+
+[[stage(vertex)]]
+fn main(in : Input) -> Output {
+  return Output(in.position);
+}