@group(0) @binding(0) var<storage, read_write> s: i32;

var<workgroup> g1 : atomic<i32>;

struct S {
  a: i32,
  b: i32,
}

fn accept_ptr_deref_pass_through(val: ptr<function, i32>) -> i32 {
  return *val + accept_ptr_deref_call_func(val);
}

fn accept_ptr_to_struct_and_access(val: ptr<function, S>) -> i32 {
  return (*val).a + (*val).b;
}

fn accept_ptr_to_struct_access_pass_ptr(val: ptr<function, S>) -> i32 {
  let b = &((*val).a);
  *b = 2;
  return *b;
}

fn accept_ptr_deref_call_func(val: ptr<function, i32>) -> i32 {
  return *val + accept_value(*val);
}

fn accept_value(val: i32) -> i32 {
  return val;
}

fn accept_ptr_vec_access_elements(v1: ptr<function, vec3f>) -> i32 {
  (*v1).x = cross(*v1, *v1).x;
  return i32((*v1).x);
}

fn call_builtin_with_mod_scope_ptr() -> i32 {
  return atomicLoad(&g1);
}

@compute @workgroup_size(1)
fn main() {
  var v1 = 0;
  var v2 = S();
  let v3 = &v2;
  var v4 = vec3f();
  let t1 = atomicLoad(&g1);

  s = accept_ptr_deref_pass_through(&v1) +
      accept_ptr_to_struct_and_access(&v2) +
      accept_ptr_to_struct_and_access(v3) +
      accept_ptr_vec_access_elements(&v4) +
      accept_ptr_to_struct_access_pass_ptr(&v2) +
      call_builtin_with_mod_scope_ptr() +
      t1;
}