dawn-cmake/test/var/uses/workgroup.wgsl.expected.msl

#include <metal_stdlib>

using namespace metal;
void uses_a(threadgroup int* const tint_symbol_3) {
  *(tint_symbol_3) = as_type<int>((as_type<uint>(*(tint_symbol_3)) + as_type<uint>(1)));
}

void uses_b(threadgroup int* const tint_symbol_4) {
  *(tint_symbol_4) = as_type<int>((as_type<uint>(*(tint_symbol_4)) * as_type<uint>(2)));
}

void uses_a_and_b(threadgroup int* const tint_symbol_5, threadgroup int* const tint_symbol_6) {
  *(tint_symbol_5) = *(tint_symbol_6);
}

void no_uses() {
}

void outer(threadgroup int* const tint_symbol_7, threadgroup int* const tint_symbol_8) {
  *(tint_symbol_7) = 0;
  uses_a(tint_symbol_7);
  uses_a_and_b(tint_symbol_8, tint_symbol_7);
  uses_b(tint_symbol_8);
  no_uses();
}

kernel void main1(uint local_invocation_index [[thread_index_in_threadgroup]]) {
  threadgroup int tint_symbol_9;
  if ((local_invocation_index == 0u)) {
    tint_symbol_9 = int();
  }
  threadgroup_barrier(mem_flags::mem_threadgroup);
  tint_symbol_9 = 42;
  uses_a(&(tint_symbol_9));
  return;
}

kernel void main2(uint local_invocation_index_1 [[thread_index_in_threadgroup]]) {
  threadgroup int tint_symbol_10;
  if ((local_invocation_index_1 == 0u)) {
    tint_symbol_10 = int();
  }
  threadgroup_barrier(mem_flags::mem_threadgroup);
  tint_symbol_10 = 7;
  uses_b(&(tint_symbol_10));
  return;
}

kernel void main3(uint local_invocation_index_2 [[thread_index_in_threadgroup]]) {
  threadgroup int tint_symbol_11;
  threadgroup int tint_symbol_12;
  if ((local_invocation_index_2 == 0u)) {
    tint_symbol_11 = int();
    tint_symbol_12 = int();
  }
  threadgroup_barrier(mem_flags::mem_threadgroup);
  outer(&(tint_symbol_11), &(tint_symbol_12));
  no_uses();
  return;
}

kernel void main4() {
  no_uses();
  return;
}