mirror of
https://github.com/gfx-rs/wgpu.git
synced 2024-11-26 00:33:51 +00:00
ea77d5674d
Co-authored-by: Jacob Hughes <j@distanthills.org> Co-authored-by: Connor Fitzgerald <connorwadefitzgerald@gmail.com> Co-authored-by: atlas dostal <rodol@rivalrebels.com>
45 lines
2.1 KiB
Plaintext
45 lines
2.1 KiB
Plaintext
// language: metal2.4
|
|
#include <metal_stdlib>
|
|
#include <simd/simd.h>
|
|
|
|
using metal::uint;
|
|
|
|
struct Structure {
|
|
uint num_subgroups;
|
|
uint subgroup_size;
|
|
};
|
|
|
|
struct main_Input {
|
|
};
|
|
kernel void main_(
|
|
uint num_subgroups [[simdgroups_per_threadgroup]]
|
|
, uint subgroup_size [[threads_per_simdgroup]]
|
|
, uint subgroup_id [[simdgroup_index_in_threadgroup]]
|
|
, uint subgroup_invocation_id [[thread_index_in_simdgroup]]
|
|
) {
|
|
const Structure sizes = { num_subgroups, subgroup_size };
|
|
metal::simdgroup_barrier(metal::mem_flags::mem_threadgroup);
|
|
metal::uint4 unnamed = uint4((uint64_t)metal::simd_ballot((subgroup_invocation_id & 1u) == 1u), 0, 0, 0);
|
|
metal::uint4 unnamed_1 = uint4((uint64_t)metal::simd_ballot(true), 0, 0, 0);
|
|
bool unnamed_2 = metal::simd_all(subgroup_invocation_id != 0u);
|
|
bool unnamed_3 = metal::simd_any(subgroup_invocation_id == 0u);
|
|
uint unnamed_4 = metal::simd_sum(subgroup_invocation_id);
|
|
uint unnamed_5 = metal::simd_product(subgroup_invocation_id);
|
|
uint unnamed_6 = metal::simd_min(subgroup_invocation_id);
|
|
uint unnamed_7 = metal::simd_max(subgroup_invocation_id);
|
|
uint unnamed_8 = metal::simd_and(subgroup_invocation_id);
|
|
uint unnamed_9 = metal::simd_or(subgroup_invocation_id);
|
|
uint unnamed_10 = metal::simd_xor(subgroup_invocation_id);
|
|
uint unnamed_11 = metal::simd_prefix_exclusive_sum(subgroup_invocation_id);
|
|
uint unnamed_12 = metal::simd_prefix_exclusive_product(subgroup_invocation_id);
|
|
uint unnamed_13 = metal::simd_prefix_inclusive_sum(subgroup_invocation_id);
|
|
uint unnamed_14 = metal::simd_prefix_inclusive_product(subgroup_invocation_id);
|
|
uint unnamed_15 = metal::simd_broadcast_first(subgroup_invocation_id);
|
|
uint unnamed_16 = metal::simd_broadcast(subgroup_invocation_id, 4u);
|
|
uint unnamed_17 = metal::simd_shuffle(subgroup_invocation_id, (sizes.subgroup_size - 1u) - subgroup_invocation_id);
|
|
uint unnamed_18 = metal::simd_shuffle_down(subgroup_invocation_id, 1u);
|
|
uint unnamed_19 = metal::simd_shuffle_up(subgroup_invocation_id, 1u);
|
|
uint unnamed_20 = metal::simd_shuffle_xor(subgroup_invocation_id, sizes.subgroup_size - 1u);
|
|
return;
|
|
}
|