test/tint/builtins/gen/literal/subgroupBallot/1a8251.wgsl.expected.ir.msl

#include <metal_stdlib>
using namespace metal;

struct tint_module_vars_struct {
  device uint4* prevent_dce;
  thread uint2* tint_subgroup_size_mask;
};

uint4 tint_subgroup_ballot(bool pred, tint_module_vars_struct tint_module_vars) {
  uint2 const v = as_type<uint2>((simd_vote::vote_t)simd_ballot(pred));
  return uint4((v & (*tint_module_vars.tint_subgroup_size_mask)), 0u, 0u);
}

uint4 subgroupBallot_1a8251(tint_module_vars_struct tint_module_vars) {
  uint4 res = tint_subgroup_ballot(true, tint_module_vars);
  return res;
}

fragment void fragment_main(uint tint_subgroup_size [[threads_per_simdgroup]], device uint4* prevent_dce [[buffer(0)]]) {
  thread uint2 tint_subgroup_size_mask = 0u;
  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce, .tint_subgroup_size_mask=(&tint_subgroup_size_mask)};
  (*tint_module_vars.tint_subgroup_size_mask)[0u] = select((4294967295u >> (32u - tint_subgroup_size)), 4294967295u, (tint_subgroup_size > 32u));
  (*tint_module_vars.tint_subgroup_size_mask)[1u] = select(0u, (4294967295u >> (64u - tint_subgroup_size)), (tint_subgroup_size > 32u));
  (*tint_module_vars.prevent_dce) = subgroupBallot_1a8251(tint_module_vars);
}

kernel void compute_main(uint tint_subgroup_size [[threads_per_simdgroup]], device uint4* prevent_dce [[buffer(0)]]) {
  thread uint2 tint_subgroup_size_mask = 0u;
  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce, .tint_subgroup_size_mask=(&tint_subgroup_size_mask)};
  (*tint_module_vars.tint_subgroup_size_mask)[0u] = select((4294967295u >> (32u - tint_subgroup_size)), 4294967295u, (tint_subgroup_size > 32u));
  (*tint_module_vars.tint_subgroup_size_mask)[1u] = select(0u, (4294967295u >> (64u - tint_subgroup_size)), (tint_subgroup_size > 32u));
  (*tint_module_vars.prevent_dce) = subgroupBallot_1a8251(tint_module_vars);
}