| #include <metal_stdlib> |
| |
| using namespace metal; |
| struct tint_private_vars_struct { |
| uint4 tint_subgroup_size_mask; |
| }; |
| |
| |
| uint4 tint_msl_subgroup_ballot(thread tint_private_vars_struct* const tint_private_vars) { |
| uint4 const tint_symbol = uint4(as_type<uint2>((ulong)simd_active_threads_mask()), 0u, 0u); |
| return (tint_symbol & (*(tint_private_vars)).tint_subgroup_size_mask); |
| } |
| |
| uint4 subgroupBallot_7e6d0e(thread tint_private_vars_struct* const tint_private_vars) { |
| uint4 res = tint_msl_subgroup_ballot(tint_private_vars); |
| return res; |
| } |
| |
| kernel void compute_main(device uint4* tint_symbol_1 [[buffer(0)]], uint tint_subgroup_size [[threads_per_simdgroup]]) { |
| thread tint_private_vars_struct tint_private_vars = {}; |
| { |
| bool const gt = (tint_subgroup_size > 32u); |
| tint_private_vars.tint_subgroup_size_mask[0u] = select((4294967295u >> (32u - tint_subgroup_size)), 4294967295u, gt); |
| tint_private_vars.tint_subgroup_size_mask[1u] = select(0u, (4294967295u >> (64u - tint_subgroup_size)), gt); |
| } |
| *(tint_symbol_1) = subgroupBallot_7e6d0e(&(tint_private_vars)); |
| return; |
| } |
| |