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

#include <metal_stdlib>
using namespace metal;

struct S {
  int a;
  int b;
};

struct tint_module_vars_struct {
  device int* s;
  threadgroup atomic_int* g1;
};

struct tint_symbol_2 {
  atomic_int tint_symbol_1;
};

int accept_value(int val) {
  return val;
}

int accept_ptr_deref_call_func(thread int* const val) {
  int const v = (*val);
  return (v + accept_value((*val)));
}

int accept_ptr_deref_pass_through(thread int* const val) {
  int const v_1 = (*val);
  return (v_1 + accept_ptr_deref_call_func(val));
}

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

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

int tint_f32_to_i32(float value) {
  return select(2147483647, select((-2147483647 - 1), int(value), (value >= -2147483648.0f)), (value <= 2147483520.0f));
}

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

int call_builtin_with_mod_scope_ptr(tint_module_vars_struct tint_module_vars) {
  return atomic_load_explicit(tint_module_vars.g1, memory_order_relaxed);
}

void tint_symbol_inner(uint tint_local_index, tint_module_vars_struct tint_module_vars) {
  if ((tint_local_index == 0u)) {
    atomic_store_explicit(tint_module_vars.g1, 0, memory_order_relaxed);
  }
  threadgroup_barrier(mem_flags::mem_threadgroup);
  int v1 = 0;
  S v2 = S{};
  thread S* const v3 = (&v2);
  float3 v4 = float3(0.0f);
  int const t1 = atomic_load_explicit(tint_module_vars.g1, memory_order_relaxed);
  int const v_2 = accept_ptr_deref_pass_through((&v1));
  int const v_3 = (v_2 + accept_ptr_to_struct_and_access((&v2)));
  int const v_4 = (v_3 + accept_ptr_to_struct_and_access(v3));
  int const v_5 = (v_4 + accept_ptr_vec_access_elements((&v4)));
  int const v_6 = (v_5 + accept_ptr_to_struct_access_pass_ptr((&v2)));
  (*tint_module_vars.s) = ((v_6 + call_builtin_with_mod_scope_ptr(tint_module_vars)) + t1);
}

kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], device int* s [[buffer(0)]], threadgroup tint_symbol_2* v_7 [[threadgroup(0)]]) {
  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.s=s, .g1=(&(*v_7).tint_symbol_1)};
  tint_symbol_inner(tint_local_index, tint_module_vars);
}