test/tint/access/ptr.wgsl.expected.msl

#include <metal_stdlib>

using namespace metal;
int tint_ftoi(float v) {
  return select(2147483647, select(int(v), (-2147483647 - 1), (v < -2147483648.0f)), (v < 2147483520.0f));
}

struct S {
  int a;
  int b;
};

int accept_value(int val) {
  return val;
}

int accept_ptr_deref_call_func(thread int* const val) {
  int const tint_symbol_3 = *(val);
  int const tint_symbol_4 = accept_value(*(val));
  return as_type<int>((as_type<uint>(tint_symbol_3) + as_type<uint>(tint_symbol_4)));
}

int accept_ptr_deref_pass_through(thread int* const val) {
  int const tint_symbol_1 = *(val);
  int const tint_symbol_2 = accept_ptr_deref_call_func(val);
  return as_type<int>((as_type<uint>(tint_symbol_1) + as_type<uint>(tint_symbol_2)));
}

int accept_ptr_to_struct_and_access(thread S* const val) {
  return as_type<int>((as_type<uint>((*(val)).a) + as_type<uint>((*(val)).b)));
}

int accept_ptr_to_struct_access_pass_ptr(thread S* const val) {
  (*(val)).a = 2;
  return (*(val)).a;
}

int accept_ptr_vec_access_elements(thread float3* const v1) {
  (*(v1))[0] = cross(*(v1), *(v1))[0];
  return tint_ftoi((*(v1))[0]);
}

int call_builtin_with_mod_scope_ptr(threadgroup atomic_int* const tint_symbol_11) {
  return atomic_load_explicit(tint_symbol_11, memory_order_relaxed);
}

void tint_symbol_inner(uint local_invocation_index, threadgroup atomic_int* const tint_symbol_12, device int* const tint_symbol_13) {
  {
    atomic_store_explicit(tint_symbol_12, 0, memory_order_relaxed);
  }
  threadgroup_barrier(mem_flags::mem_threadgroup);
  int v1 = 0;
  S v2 = S{};
  float3 v4 = float3(0.0f);
  int const t1 = atomic_load_explicit(tint_symbol_12, memory_order_relaxed);
  int const tint_symbol_5 = accept_ptr_deref_pass_through(&(v1));
  int const tint_symbol_6 = accept_ptr_to_struct_and_access(&(v2));
  int const tint_symbol_7 = accept_ptr_to_struct_and_access(&(v2));
  int const tint_symbol_8 = accept_ptr_vec_access_elements(&(v4));
  int const tint_symbol_9 = accept_ptr_to_struct_access_pass_ptr(&(v2));
  int const tint_symbol_10 = call_builtin_with_mod_scope_ptr(tint_symbol_12);
  *(tint_symbol_13) = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(tint_symbol_5) + as_type<uint>(tint_symbol_6)))) + as_type<uint>(tint_symbol_7)))) + as_type<uint>(tint_symbol_8)))) + as_type<uint>(tint_symbol_9)))) + as_type<uint>(tint_symbol_10)))) + as_type<uint>(t1)));
}

kernel void tint_symbol(device int* tint_symbol_15 [[buffer(0)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
  threadgroup atomic_int tint_symbol_14;
  tint_symbol_inner(local_invocation_index, &(tint_symbol_14), tint_symbol_15);
  return;
}