MSL: Expand subgroup support.
Add support for declaring a fixed subgroup size. Metal, like Vulkan with `VK_EXT_subgroup_size_control`, allows the thread execution width to vary depending on factors such as register usage. Unfortunately, this breaks several tests that depend on the subgroup size being what the device says it is. So we'll fix the subgroup size at the size the device declares. The extra invocations in the subgroup will appear to be inactive. Because of this, the ballot mask builtins are now ANDed with the active subgroup mask. Add support for emulating a subgroup of size 1. This is intended to be used by Vulkan Portability implementations (e.g. MoltenVK) when the hardware/software combo provides insufficient support for subgroups. Luckily for us, Vulkan 1.1 only requires that the subgroup size be at least 1. Add support for quadgroup and SIMD-group functions which were added to iOS in Metal 2.2 and 2.3. This will allow clients to take advantage of expanded quadgroup and SIMD-group support in recent Metal versions and on recent Apple GPUs (families 6 and 7). Gut emulation of subgroup builtins in fragment shaders. It turns out codegen for the SIMD-group functions in fragment wasn't implemented for AMD on Mojave; it's a safe bet that it wasn't implemented for the other drivers either. Subgroup support in fragment shaders now requires Metal 2.2.
This commit is contained in:
Chip Davis
parent
b3c59263a0
commit
68908355a9
@ -323,7 +323,7 @@ if (SPIRV_CROSS_STATIC)
|
||||
endif()
|
||||
|
||||
set(spirv-cross-abi-major 0)
|
||||
set(spirv-cross-abi-minor 42)
|
||||
set(spirv-cross-abi-minor 43)
|
||||
set(spirv-cross-abi-patch 0)
|
||||
|
||||
if (SPIRV_CROSS_SHARED)
|
||||
|
||||
19
main.cpp
19
main.cpp
@ -566,6 +566,9 @@ struct CLIArguments
|
||||
uint32_t msl_r32ui_linear_texture_alignment = 4;
|
||||
uint32_t msl_r32ui_alignment_constant_id = 65535;
|
||||
bool msl_texture_1d_as_2d = false;
|
||||
bool msl_ios_use_simdgroup_functions = false;
|
||||
bool msl_emulate_subgroups = false;
|
||||
uint32_t msl_fixed_subgroup_size = 0;
|
||||
bool glsl_emit_push_constant_as_ubo = false;
|
||||
bool glsl_emit_ubo_as_plain_uniforms = false;
|
||||
bool glsl_force_flattened_io_blocks = false;
|
||||
@ -779,7 +782,14 @@ static void print_help_msl()
|
||||
"\t[--msl-r32ui-linear-texture-align-constant-id <id>]:\n\t\tThe function constant ID to use for the linear texture alignment.\n"
|
||||
"\t\tOn MSL 1.2 or later, you can override the alignment by setting this function constant.\n"
|
||||
"\t[--msl-texture-1d-as-2d]:\n\t\tEmit Image variables of dimension Dim1D as texture2d.\n"
|
||||
"\t\tIn Metal, 1D textures do not support all features that 2D textures do. Use this option if your code relies on these features.\n");
|
||||
"\t\tIn Metal, 1D textures do not support all features that 2D textures do. Use this option if your code relies on these features.\n"
|
||||
"\t[--msl-ios-use-simdgroup-functions]:\n\t\tUse simd_*() functions for subgroup ops instead of quad_*().\n"
|
||||
"\t\tRecent Apple GPUs support SIMD-groups larger than a quad. Use this option to take advantage of this support.\n"
|
||||
"\t[--msl-emulate-subgroups]:\n\t\tAssume subgroups of size 1.\n"
|
||||
"\t\tIntended for Vulkan Portability implementations where Metal support for SIMD-groups is insufficient for true subgroups.\n"
|
||||
"\t[--msl-fixed-subgroup-size <size>]:\n\t\tAssign a constant <size> to the SubgroupSize builtin.\n"
|
||||
"\t\tIntended for Vulkan Portability implementations where VK_EXT_subgroup_size_control is not supported or disabled.\n"
|
||||
"\t\tIf 0, assume variable subgroup size as actually exposed by Metal.\n");
|
||||
// clang-format on
|
||||
}
|
||||
|
||||
@ -1021,6 +1031,9 @@ static string compile_iteration(const CLIArguments &args, std::vector<uint32_t>
|
||||
msl_opts.r32ui_linear_texture_alignment = args.msl_r32ui_linear_texture_alignment;
|
||||
msl_opts.r32ui_alignment_constant_id = args.msl_r32ui_alignment_constant_id;
|
||||
msl_opts.texture_1D_as_2D = args.msl_texture_1d_as_2d;
|
||||
msl_opts.ios_use_simdgroup_functions = args.msl_ios_use_simdgroup_functions;
|
||||
msl_opts.emulate_subgroups = args.msl_emulate_subgroups;
|
||||
msl_opts.fixed_subgroup_size = args.msl_fixed_subgroup_size;
|
||||
msl_comp->set_msl_options(msl_opts);
|
||||
for (auto &v : args.msl_discrete_descriptor_sets)
|
||||
msl_comp->add_discrete_descriptor_set(v);
|
||||
@ -1449,6 +1462,10 @@ static int main_inner(int argc, char *argv[])
|
||||
cbs.add("--msl-r32ui-linear-texture-align-constant-id",
|
||||
[&args](CLIParser &parser) { args.msl_r32ui_alignment_constant_id = parser.next_uint(); });
|
||||
cbs.add("--msl-texture-1d-as-2d", [&args](CLIParser &) { args.msl_texture_1d_as_2d = true; });
|
||||
cbs.add("--msl-ios-use-simdgroup-functions", [&args](CLIParser &) { args.msl_ios_use_simdgroup_functions = true; });
|
||||
cbs.add("--msl-emulate-subgroups", [&args](CLIParser &) { args.msl_emulate_subgroups = true; });
|
||||
cbs.add("--msl-fixed-subgroup-size",
|
||||
[&args](CLIParser &parser) { args.msl_fixed_subgroup_size = parser.next_uint(); });
|
||||
cbs.add("--extension", [&args](CLIParser &parser) { args.extensions.push_back(parser.next_string()); });
|
||||
cbs.add("--rename-entry-point", [&args](CLIParser &parser) {
|
||||
auto old_name = parser.next_string();
|
||||
|
||||
@ -0,0 +1,30 @@
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
float FragColor;
|
||||
};
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
||||
|
||||
kernel void main0(device SSBO& _9 [[buffer(0)]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]], unsupported-built-in-type gl_WorkGroupSize [[unsupported-built-in]])
|
||||
{
|
||||
uint gl_NumSubgroups = gl_WorkGroupSize.x * gl_WorkGroupSize.y * gl_WorkGroupSize.z;
|
||||
uint gl_SubgroupID = gl_LocalInvocationIndex;
|
||||
uint gl_SubgroupSize = 1;
|
||||
uint gl_SubgroupInvocationID = 0;
|
||||
_9.FragColor = float(gl_NumSubgroups);
|
||||
_9.FragColor = float(gl_SubgroupID);
|
||||
_9.FragColor = float(gl_SubgroupSize);
|
||||
_9.FragColor = float(gl_SubgroupInvocationID);
|
||||
threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture);
|
||||
threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture);
|
||||
threadgroup_barrier(mem_flags::mem_device);
|
||||
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||
threadgroup_barrier(mem_flags::mem_texture);
|
||||
bool elected = true;
|
||||
}
|
||||
|
||||
@ -242,7 +242,8 @@ kernel void main0(device SSBO& _9 [[buffer(0)]], uint gl_NumSubgroups [[simdgrou
|
||||
simdgroup_barrier(mem_flags::mem_device);
|
||||
simdgroup_barrier(mem_flags::mem_threadgroup);
|
||||
simdgroup_barrier(mem_flags::mem_texture);
|
||||
bool elected = simd_is_first();
|
||||
bool _39 = simd_is_first();
|
||||
bool elected = _39;
|
||||
_9.FragColor = float4(gl_SubgroupEqMask).x;
|
||||
_9.FragColor = float4(gl_SubgroupGeMask).x;
|
||||
_9.FragColor = float4(gl_SubgroupGtMask).x;
|
||||
|
||||
@ -0,0 +1,322 @@
|
||||
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
||||
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
float FragColor;
|
||||
};
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
||||
|
||||
template<typename T>
|
||||
inline T spvSubgroupBroadcast(T value, ushort lane)
|
||||
{
|
||||
return simd_broadcast(value, lane);
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvSubgroupBroadcast(bool value, ushort lane)
|
||||
{
|
||||
return !!simd_broadcast((ushort)value, lane);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline vec<bool, N> spvSubgroupBroadcast(vec<bool, N> value, ushort lane)
|
||||
{
|
||||
return (vec<bool, N>)simd_broadcast((vec<ushort, N>)value, lane);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline T spvSubgroupBroadcastFirst(T value)
|
||||
{
|
||||
return simd_broadcast_first(value);
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvSubgroupBroadcastFirst(bool value)
|
||||
{
|
||||
return !!simd_broadcast_first((ushort)value);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline vec<bool, N> spvSubgroupBroadcastFirst(vec<bool, N> value)
|
||||
{
|
||||
return (vec<bool, N>)simd_broadcast_first((vec<ushort, N>)value);
|
||||
}
|
||||
|
||||
inline uint4 spvSubgroupBallot(bool value)
|
||||
{
|
||||
simd_vote vote = simd_ballot(value);
|
||||
// simd_ballot() returns a 64-bit integer-like object, but
|
||||
// SPIR-V callers expect a uint4. We must convert.
|
||||
// FIXME: This won't include higher bits if Apple ever supports
|
||||
// 128 lanes in an SIMD-group.
|
||||
return uint4((uint)((simd_vote::vote_t)vote & 0xFFFFFFFF), (uint)(((simd_vote::vote_t)vote >> 32) & 0xFFFFFFFF), 0, 0);
|
||||
}
|
||||
|
||||
inline bool spvSubgroupBallotBitExtract(uint4 ballot, uint bit)
|
||||
{
|
||||
return !!extract_bits(ballot[bit / 32], bit % 32, 1);
|
||||
}
|
||||
|
||||
inline uint spvSubgroupBallotFindLSB(uint4 ballot, uint gl_SubgroupSize)
|
||||
{
|
||||
uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupSize, 32u)), extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupSize - 32, 0)), uint2(0));
|
||||
ballot &= mask;
|
||||
return select(ctz(ballot.x), select(32 + ctz(ballot.y), select(64 + ctz(ballot.z), select(96 + ctz(ballot.w), uint(-1), ballot.w == 0), ballot.z == 0), ballot.y == 0), ballot.x == 0);
|
||||
}
|
||||
|
||||
inline uint spvSubgroupBallotFindMSB(uint4 ballot, uint gl_SubgroupSize)
|
||||
{
|
||||
uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupSize, 32u)), extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupSize - 32, 0)), uint2(0));
|
||||
ballot &= mask;
|
||||
return select(128 - (clz(ballot.w) + 1), select(96 - (clz(ballot.z) + 1), select(64 - (clz(ballot.y) + 1), select(32 - (clz(ballot.x) + 1), uint(-1), ballot.x == 0), ballot.y == 0), ballot.z == 0), ballot.w == 0);
|
||||
}
|
||||
|
||||
inline uint spvPopCount4(uint4 ballot)
|
||||
{
|
||||
return popcount(ballot.x) + popcount(ballot.y) + popcount(ballot.z) + popcount(ballot.w);
|
||||
}
|
||||
|
||||
inline uint spvSubgroupBallotBitCount(uint4 ballot, uint gl_SubgroupSize)
|
||||
{
|
||||
uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupSize, 32u)), extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupSize - 32, 0)), uint2(0));
|
||||
return spvPopCount4(ballot & mask);
|
||||
}
|
||||
|
||||
inline uint spvSubgroupBallotInclusiveBitCount(uint4 ballot, uint gl_SubgroupInvocationID)
|
||||
{
|
||||
uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupInvocationID + 1, 32u)), extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupInvocationID + 1 - 32, 0)), uint2(0));
|
||||
return spvPopCount4(ballot & mask);
|
||||
}
|
||||
|
||||
inline uint spvSubgroupBallotExclusiveBitCount(uint4 ballot, uint gl_SubgroupInvocationID)
|
||||
{
|
||||
uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupInvocationID, 32u)), extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupInvocationID - 32, 0)), uint2(0));
|
||||
return spvPopCount4(ballot & mask);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline bool spvSubgroupAllEqual(T value)
|
||||
{
|
||||
return simd_all(all(value == simd_broadcast_first(value)));
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvSubgroupAllEqual(bool value)
|
||||
{
|
||||
return simd_all(value) || !simd_any(value);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline bool spvSubgroupAllEqual(vec<bool, N> value)
|
||||
{
|
||||
return simd_all(all(value == (vec<bool, N>)simd_broadcast_first((vec<ushort, N>)value)));
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline T spvSubgroupShuffle(T value, ushort lane)
|
||||
{
|
||||
return simd_shuffle(value, lane);
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvSubgroupShuffle(bool value, ushort lane)
|
||||
{
|
||||
return !!simd_shuffle((ushort)value, lane);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline vec<bool, N> spvSubgroupShuffle(vec<bool, N> value, ushort lane)
|
||||
{
|
||||
return (vec<bool, N>)simd_shuffle((vec<ushort, N>)value, lane);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline T spvSubgroupShuffleXor(T value, ushort mask)
|
||||
{
|
||||
return simd_shuffle_xor(value, mask);
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvSubgroupShuffleXor(bool value, ushort mask)
|
||||
{
|
||||
return !!simd_shuffle_xor((ushort)value, mask);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline vec<bool, N> spvSubgroupShuffleXor(vec<bool, N> value, ushort mask)
|
||||
{
|
||||
return (vec<bool, N>)simd_shuffle_xor((vec<ushort, N>)value, mask);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline T spvSubgroupShuffleUp(T value, ushort delta)
|
||||
{
|
||||
return simd_shuffle_up(value, delta);
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvSubgroupShuffleUp(bool value, ushort delta)
|
||||
{
|
||||
return !!simd_shuffle_up((ushort)value, delta);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline vec<bool, N> spvSubgroupShuffleUp(vec<bool, N> value, ushort delta)
|
||||
{
|
||||
return (vec<bool, N>)simd_shuffle_up((vec<ushort, N>)value, delta);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline T spvSubgroupShuffleDown(T value, ushort delta)
|
||||
{
|
||||
return simd_shuffle_down(value, delta);
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvSubgroupShuffleDown(bool value, ushort delta)
|
||||
{
|
||||
return !!simd_shuffle_down((ushort)value, delta);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline vec<bool, N> spvSubgroupShuffleDown(vec<bool, N> value, ushort delta)
|
||||
{
|
||||
return (vec<bool, N>)simd_shuffle_down((vec<ushort, N>)value, delta);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline T spvQuadBroadcast(T value, uint lane)
|
||||
{
|
||||
return quad_broadcast(value, lane);
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvQuadBroadcast(bool value, uint lane)
|
||||
{
|
||||
return !!quad_broadcast((ushort)value, lane);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline vec<bool, N> spvQuadBroadcast(vec<bool, N> value, uint lane)
|
||||
{
|
||||
return (vec<bool, N>)quad_broadcast((vec<ushort, N>)value, lane);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline T spvQuadSwap(T value, uint dir)
|
||||
{
|
||||
return quad_shuffle_xor(value, dir + 1);
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvQuadSwap(bool value, uint dir)
|
||||
{
|
||||
return !!quad_shuffle_xor((ushort)value, dir + 1);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline vec<bool, N> spvQuadSwap(vec<bool, N> value, uint dir)
|
||||
{
|
||||
return (vec<bool, N>)quad_shuffle_xor((vec<ushort, N>)value, dir + 1);
|
||||
}
|
||||
|
||||
kernel void main0(device SSBO& _9 [[buffer(0)]], uint gl_NumSubgroups [[simdgroups_per_threadgroup]], uint gl_SubgroupID [[simdgroup_index_in_threadgroup]], uint gl_SubgroupInvocationID [[thread_index_in_simdgroup]])
|
||||
{
|
||||
uint gl_SubgroupSize = 32;
|
||||
uint4 gl_SubgroupEqMask = gl_SubgroupInvocationID >= 32 ? uint4(0, (1 << (gl_SubgroupInvocationID - 32)), uint2(0)) : uint4(1 << gl_SubgroupInvocationID, uint3(0));
|
||||
uint4 gl_SubgroupGeMask = uint4(insert_bits(0u, 0xFFFFFFFF, gl_SubgroupInvocationID, 32 - gl_SubgroupInvocationID), uint3(0)) & spvSubgroupBallot(true);
|
||||
uint4 gl_SubgroupGtMask = uint4(insert_bits(0u, 0xFFFFFFFF, gl_SubgroupInvocationID + 1, 32 - gl_SubgroupInvocationID - 1), uint3(0)) & spvSubgroupBallot(true);
|
||||
uint4 gl_SubgroupLeMask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupInvocationID + 1, 32u)), extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupInvocationID + 1 - 32, 0)), uint2(0));
|
||||
uint4 gl_SubgroupLtMask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupInvocationID, 32u)), extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupInvocationID - 32, 0)), uint2(0));
|
||||
_9.FragColor = float(gl_NumSubgroups);
|
||||
_9.FragColor = float(gl_SubgroupID);
|
||||
_9.FragColor = float(gl_SubgroupSize);
|
||||
_9.FragColor = float(gl_SubgroupInvocationID);
|
||||
simdgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture);
|
||||
simdgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture);
|
||||
simdgroup_barrier(mem_flags::mem_device);
|
||||
simdgroup_barrier(mem_flags::mem_threadgroup);
|
||||
simdgroup_barrier(mem_flags::mem_texture);
|
||||
bool _39 = simd_is_first();
|
||||
bool elected = _39;
|
||||
_9.FragColor = float4(gl_SubgroupEqMask).x;
|
||||
_9.FragColor = float4(gl_SubgroupGeMask).x;
|
||||
_9.FragColor = float4(gl_SubgroupGtMask).x;
|
||||
_9.FragColor = float4(gl_SubgroupLeMask).x;
|
||||
_9.FragColor = float4(gl_SubgroupLtMask).x;
|
||||
float4 broadcasted = spvSubgroupBroadcast(float4(10.0), 8u);
|
||||
bool2 broadcasted_bool = spvSubgroupBroadcast(bool2(true), 8u);
|
||||
float3 first = spvSubgroupBroadcastFirst(float3(20.0));
|
||||
bool4 first_bool = spvSubgroupBroadcastFirst(bool4(false));
|
||||
uint4 ballot_value = spvSubgroupBallot(true);
|
||||
bool inverse_ballot_value = spvSubgroupBallotBitExtract(ballot_value, gl_SubgroupInvocationID);
|
||||
bool bit_extracted = spvSubgroupBallotBitExtract(uint4(10u), 8u);
|
||||
uint bit_count = spvSubgroupBallotBitCount(ballot_value, gl_SubgroupSize);
|
||||
uint inclusive_bit_count = spvSubgroupBallotInclusiveBitCount(ballot_value, gl_SubgroupInvocationID);
|
||||
uint exclusive_bit_count = spvSubgroupBallotExclusiveBitCount(ballot_value, gl_SubgroupInvocationID);
|
||||
uint lsb = spvSubgroupBallotFindLSB(ballot_value, gl_SubgroupSize);
|
||||
uint msb = spvSubgroupBallotFindMSB(ballot_value, gl_SubgroupSize);
|
||||
uint shuffled = spvSubgroupShuffle(10u, 8u);
|
||||
bool shuffled_bool = spvSubgroupShuffle(true, 9u);
|
||||
uint shuffled_xor = spvSubgroupShuffleXor(30u, 8u);
|
||||
bool shuffled_xor_bool = spvSubgroupShuffleXor(false, 9u);
|
||||
uint shuffled_up = spvSubgroupShuffleUp(20u, 4u);
|
||||
bool shuffled_up_bool = spvSubgroupShuffleUp(true, 4u);
|
||||
uint shuffled_down = spvSubgroupShuffleDown(20u, 4u);
|
||||
bool shuffled_down_bool = spvSubgroupShuffleDown(false, 4u);
|
||||
bool has_all = simd_all(true);
|
||||
bool has_any = simd_any(true);
|
||||
bool has_equal = spvSubgroupAllEqual(0);
|
||||
has_equal = spvSubgroupAllEqual(true);
|
||||
has_equal = spvSubgroupAllEqual(float3(0.0, 1.0, 2.0));
|
||||
has_equal = spvSubgroupAllEqual(bool4(true, true, false, true));
|
||||
float4 added = simd_sum(float4(20.0));
|
||||
int4 iadded = simd_sum(int4(20));
|
||||
float4 multiplied = simd_product(float4(20.0));
|
||||
int4 imultiplied = simd_product(int4(20));
|
||||
float4 lo = simd_min(float4(20.0));
|
||||
float4 hi = simd_max(float4(20.0));
|
||||
int4 slo = simd_min(int4(20));
|
||||
int4 shi = simd_max(int4(20));
|
||||
uint4 ulo = simd_min(uint4(20u));
|
||||
uint4 uhi = simd_max(uint4(20u));
|
||||
uint4 anded = simd_and(ballot_value);
|
||||
uint4 ored = simd_or(ballot_value);
|
||||
uint4 xored = simd_xor(ballot_value);
|
||||
added = simd_prefix_inclusive_sum(added);
|
||||
iadded = simd_prefix_inclusive_sum(iadded);
|
||||
multiplied = simd_prefix_inclusive_product(multiplied);
|
||||
imultiplied = simd_prefix_inclusive_product(imultiplied);
|
||||
added = simd_prefix_exclusive_sum(multiplied);
|
||||
multiplied = simd_prefix_exclusive_product(multiplied);
|
||||
iadded = simd_prefix_exclusive_sum(imultiplied);
|
||||
imultiplied = simd_prefix_exclusive_product(imultiplied);
|
||||
added = quad_sum(added);
|
||||
multiplied = quad_product(multiplied);
|
||||
iadded = quad_sum(iadded);
|
||||
imultiplied = quad_product(imultiplied);
|
||||
lo = quad_min(lo);
|
||||
hi = quad_max(hi);
|
||||
ulo = quad_min(ulo);
|
||||
uhi = quad_max(uhi);
|
||||
slo = quad_min(slo);
|
||||
shi = quad_max(shi);
|
||||
anded = quad_and(anded);
|
||||
ored = quad_or(ored);
|
||||
xored = quad_xor(xored);
|
||||
float4 swap_horiz = spvQuadSwap(float4(20.0), 0u);
|
||||
bool4 swap_horiz_bool = spvQuadSwap(bool4(true), 0u);
|
||||
float4 swap_vertical = spvQuadSwap(float4(20.0), 1u);
|
||||
bool4 swap_vertical_bool = spvQuadSwap(bool4(true), 1u);
|
||||
float4 swap_diagonal = spvQuadSwap(float4(20.0), 2u);
|
||||
bool4 swap_diagonal_bool = spvQuadSwap(bool4(true), 2u);
|
||||
float4 quad_broadcast0 = spvQuadBroadcast(float4(20.0), 3u);
|
||||
bool4 quad_broadcast_bool = spvQuadBroadcast(bool4(true), 3u);
|
||||
}
|
||||
|
||||
@ -0,0 +1,282 @@
|
||||
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
||||
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
float FragColor;
|
||||
};
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
||||
|
||||
template<typename T>
|
||||
inline T spvSubgroupBroadcast(T value, ushort lane)
|
||||
{
|
||||
return quad_broadcast(value, lane);
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvSubgroupBroadcast(bool value, ushort lane)
|
||||
{
|
||||
return !!quad_broadcast((ushort)value, lane);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline vec<bool, N> spvSubgroupBroadcast(vec<bool, N> value, ushort lane)
|
||||
{
|
||||
return (vec<bool, N>)quad_broadcast((vec<ushort, N>)value, lane);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline T spvSubgroupBroadcastFirst(T value)
|
||||
{
|
||||
return quad_broadcast_first(value);
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvSubgroupBroadcastFirst(bool value)
|
||||
{
|
||||
return !!quad_broadcast_first((ushort)value);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline vec<bool, N> spvSubgroupBroadcastFirst(vec<bool, N> value)
|
||||
{
|
||||
return (vec<bool, N>)quad_broadcast_first((vec<ushort, N>)value);
|
||||
}
|
||||
|
||||
inline uint4 spvSubgroupBallot(bool value)
|
||||
{
|
||||
return uint4((quad_vote::vote_t)quad_ballot(value), 0, 0, 0);
|
||||
}
|
||||
|
||||
inline bool spvSubgroupBallotBitExtract(uint4 ballot, uint bit)
|
||||
{
|
||||
return !!extract_bits(ballot[bit / 32], bit % 32, 1);
|
||||
}
|
||||
|
||||
inline uint spvSubgroupBallotFindLSB(uint4 ballot, uint gl_SubgroupSize)
|
||||
{
|
||||
uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, gl_SubgroupSize), uint3(0));
|
||||
ballot &= mask;
|
||||
return select(ctz(ballot.x), select(32 + ctz(ballot.y), select(64 + ctz(ballot.z), select(96 + ctz(ballot.w), uint(-1), ballot.w == 0), ballot.z == 0), ballot.y == 0), ballot.x == 0);
|
||||
}
|
||||
|
||||
inline uint spvSubgroupBallotFindMSB(uint4 ballot, uint gl_SubgroupSize)
|
||||
{
|
||||
uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, gl_SubgroupSize), uint3(0));
|
||||
ballot &= mask;
|
||||
return select(128 - (clz(ballot.w) + 1), select(96 - (clz(ballot.z) + 1), select(64 - (clz(ballot.y) + 1), select(32 - (clz(ballot.x) + 1), uint(-1), ballot.x == 0), ballot.y == 0), ballot.z == 0), ballot.w == 0);
|
||||
}
|
||||
|
||||
inline uint spvPopCount4(uint4 ballot)
|
||||
{
|
||||
return popcount(ballot.x) + popcount(ballot.y) + popcount(ballot.z) + popcount(ballot.w);
|
||||
}
|
||||
|
||||
inline uint spvSubgroupBallotBitCount(uint4 ballot, uint gl_SubgroupSize)
|
||||
{
|
||||
uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, gl_SubgroupSize), uint3(0));
|
||||
return spvPopCount4(ballot & mask);
|
||||
}
|
||||
|
||||
inline uint spvSubgroupBallotInclusiveBitCount(uint4 ballot, uint gl_SubgroupInvocationID)
|
||||
{
|
||||
uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, gl_SubgroupInvocationID + 1), uint3(0));
|
||||
return spvPopCount4(ballot & mask);
|
||||
}
|
||||
|
||||
inline uint spvSubgroupBallotExclusiveBitCount(uint4 ballot, uint gl_SubgroupInvocationID)
|
||||
{
|
||||
uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, gl_SubgroupInvocationID), uint2(0));
|
||||
return spvPopCount4(ballot & mask);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline bool spvSubgroupAllEqual(T value)
|
||||
{
|
||||
return quad_all(all(value == quad_broadcast_first(value)));
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvSubgroupAllEqual(bool value)
|
||||
{
|
||||
return quad_all(value) || !quad_any(value);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline bool spvSubgroupAllEqual(vec<bool, N> value)
|
||||
{
|
||||
return quad_all(all(value == (vec<bool, N>)quad_broadcast_first((vec<ushort, N>)value)));
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline T spvSubgroupShuffle(T value, ushort lane)
|
||||
{
|
||||
return quad_shuffle(value, lane);
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvSubgroupShuffle(bool value, ushort lane)
|
||||
{
|
||||
return !!quad_shuffle((ushort)value, lane);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline vec<bool, N> spvSubgroupShuffle(vec<bool, N> value, ushort lane)
|
||||
{
|
||||
return (vec<bool, N>)quad_shuffle((vec<ushort, N>)value, lane);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline T spvSubgroupShuffleXor(T value, ushort mask)
|
||||
{
|
||||
return quad_shuffle_xor(value, mask);
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvSubgroupShuffleXor(bool value, ushort mask)
|
||||
{
|
||||
return !!quad_shuffle_xor((ushort)value, mask);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline vec<bool, N> spvSubgroupShuffleXor(vec<bool, N> value, ushort mask)
|
||||
{
|
||||
return (vec<bool, N>)quad_shuffle_xor((vec<ushort, N>)value, mask);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline T spvSubgroupShuffleUp(T value, ushort delta)
|
||||
{
|
||||
return quad_shuffle_up(value, delta);
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvSubgroupShuffleUp(bool value, ushort delta)
|
||||
{
|
||||
return !!quad_shuffle_up((ushort)value, delta);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline vec<bool, N> spvSubgroupShuffleUp(vec<bool, N> value, ushort delta)
|
||||
{
|
||||
return (vec<bool, N>)quad_shuffle_up((vec<ushort, N>)value, delta);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline T spvSubgroupShuffleDown(T value, ushort delta)
|
||||
{
|
||||
return quad_shuffle_down(value, delta);
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvSubgroupShuffleDown(bool value, ushort delta)
|
||||
{
|
||||
return !!quad_shuffle_down((ushort)value, delta);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline vec<bool, N> spvSubgroupShuffleDown(vec<bool, N> value, ushort delta)
|
||||
{
|
||||
return (vec<bool, N>)quad_shuffle_down((vec<ushort, N>)value, delta);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline T spvQuadBroadcast(T value, uint lane)
|
||||
{
|
||||
return quad_broadcast(value, lane);
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvQuadBroadcast(bool value, uint lane)
|
||||
{
|
||||
return !!quad_broadcast((ushort)value, lane);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline vec<bool, N> spvQuadBroadcast(vec<bool, N> value, uint lane)
|
||||
{
|
||||
return (vec<bool, N>)quad_broadcast((vec<ushort, N>)value, lane);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline T spvQuadSwap(T value, uint dir)
|
||||
{
|
||||
return quad_shuffle_xor(value, dir + 1);
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvQuadSwap(bool value, uint dir)
|
||||
{
|
||||
return !!quad_shuffle_xor((ushort)value, dir + 1);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline vec<bool, N> spvQuadSwap(vec<bool, N> value, uint dir)
|
||||
{
|
||||
return (vec<bool, N>)quad_shuffle_xor((vec<ushort, N>)value, dir + 1);
|
||||
}
|
||||
|
||||
kernel void main0(device SSBO& _9 [[buffer(0)]], uint gl_NumSubgroups [[quadgroups_per_threadgroup]], uint gl_SubgroupID [[quadgroup_index_in_threadgroup]], uint gl_SubgroupSize [[thread_execution_width]], uint gl_SubgroupInvocationID [[thread_index_in_quadgroup]])
|
||||
{
|
||||
uint4 gl_SubgroupEqMask = uint4(1 << gl_SubgroupInvocationID, uint3(0));
|
||||
uint4 gl_SubgroupGeMask = uint4(insert_bits(0u, 0xFFFFFFFF, gl_SubgroupInvocationID, gl_SubgroupSize - gl_SubgroupInvocationID), uint3(0));
|
||||
uint4 gl_SubgroupGtMask = uint4(insert_bits(0u, 0xFFFFFFFF, gl_SubgroupInvocationID + 1, gl_SubgroupSize - gl_SubgroupInvocationID - 1), uint3(0));
|
||||
uint4 gl_SubgroupLeMask = uint4(extract_bits(0xFFFFFFFF, 0, gl_SubgroupInvocationID + 1), uint3(0));
|
||||
uint4 gl_SubgroupLtMask = uint4(extract_bits(0xFFFFFFFF, 0, gl_SubgroupInvocationID), uint3(0));
|
||||
_9.FragColor = float(gl_NumSubgroups);
|
||||
_9.FragColor = float(gl_SubgroupID);
|
||||
_9.FragColor = float(gl_SubgroupSize);
|
||||
_9.FragColor = float(gl_SubgroupInvocationID);
|
||||
simdgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture);
|
||||
simdgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture);
|
||||
simdgroup_barrier(mem_flags::mem_device);
|
||||
simdgroup_barrier(mem_flags::mem_threadgroup);
|
||||
simdgroup_barrier(mem_flags::mem_texture);
|
||||
bool _39 = quad_is_first();
|
||||
bool elected = _39;
|
||||
_9.FragColor = float4(gl_SubgroupEqMask).x;
|
||||
_9.FragColor = float4(gl_SubgroupGeMask).x;
|
||||
_9.FragColor = float4(gl_SubgroupGtMask).x;
|
||||
_9.FragColor = float4(gl_SubgroupLeMask).x;
|
||||
_9.FragColor = float4(gl_SubgroupLtMask).x;
|
||||
float4 broadcasted = spvSubgroupBroadcast(float4(10.0), 8u);
|
||||
bool2 broadcasted_bool = spvSubgroupBroadcast(bool2(true), 8u);
|
||||
float3 first = spvSubgroupBroadcastFirst(float3(20.0));
|
||||
bool4 first_bool = spvSubgroupBroadcastFirst(bool4(false));
|
||||
uint4 ballot_value = spvSubgroupBallot(true);
|
||||
bool inverse_ballot_value = spvSubgroupBallotBitExtract(ballot_value, gl_SubgroupInvocationID);
|
||||
bool bit_extracted = spvSubgroupBallotBitExtract(uint4(10u), 8u);
|
||||
uint bit_count = spvSubgroupBallotBitCount(ballot_value, gl_SubgroupSize);
|
||||
uint inclusive_bit_count = spvSubgroupBallotInclusiveBitCount(ballot_value, gl_SubgroupInvocationID);
|
||||
uint exclusive_bit_count = spvSubgroupBallotExclusiveBitCount(ballot_value, gl_SubgroupInvocationID);
|
||||
uint lsb = spvSubgroupBallotFindLSB(ballot_value, gl_SubgroupSize);
|
||||
uint msb = spvSubgroupBallotFindMSB(ballot_value, gl_SubgroupSize);
|
||||
uint shuffled = spvSubgroupShuffle(10u, 8u);
|
||||
bool shuffled_bool = spvSubgroupShuffle(true, 9u);
|
||||
uint shuffled_xor = spvSubgroupShuffleXor(30u, 8u);
|
||||
bool shuffled_xor_bool = spvSubgroupShuffleXor(false, 9u);
|
||||
uint shuffled_up = spvSubgroupShuffleUp(20u, 4u);
|
||||
bool shuffled_up_bool = spvSubgroupShuffleUp(true, 4u);
|
||||
uint shuffled_down = spvSubgroupShuffleDown(20u, 4u);
|
||||
bool shuffled_down_bool = spvSubgroupShuffleDown(false, 4u);
|
||||
bool has_all = quad_all(true);
|
||||
bool has_any = quad_any(true);
|
||||
bool has_equal = spvSubgroupAllEqual(0);
|
||||
has_equal = spvSubgroupAllEqual(true);
|
||||
has_equal = spvSubgroupAllEqual(float3(0.0, 1.0, 2.0));
|
||||
has_equal = spvSubgroupAllEqual(bool4(true, true, false, true));
|
||||
float4 swap_horiz = spvQuadSwap(float4(20.0), 0u);
|
||||
bool4 swap_horiz_bool = spvQuadSwap(bool4(true), 0u);
|
||||
float4 swap_vertical = spvQuadSwap(float4(20.0), 1u);
|
||||
bool4 swap_vertical_bool = spvQuadSwap(bool4(true), 1u);
|
||||
float4 swap_diagonal = spvQuadSwap(float4(20.0), 2u);
|
||||
bool4 swap_diagonal_bool = spvQuadSwap(bool4(true), 2u);
|
||||
float4 quad_broadcast0 = spvQuadBroadcast(float4(20.0), 3u);
|
||||
bool4 quad_broadcast_bool = spvQuadBroadcast(bool4(true), 3u);
|
||||
}
|
||||
|
||||
@ -0,0 +1,316 @@
|
||||
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
||||
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
float FragColor;
|
||||
};
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
||||
|
||||
template<typename T>
|
||||
inline T spvSubgroupBroadcast(T value, ushort lane)
|
||||
{
|
||||
return simd_broadcast(value, lane);
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvSubgroupBroadcast(bool value, ushort lane)
|
||||
{
|
||||
return !!simd_broadcast((ushort)value, lane);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline vec<bool, N> spvSubgroupBroadcast(vec<bool, N> value, ushort lane)
|
||||
{
|
||||
return (vec<bool, N>)simd_broadcast((vec<ushort, N>)value, lane);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline T spvSubgroupBroadcastFirst(T value)
|
||||
{
|
||||
return simd_broadcast_first(value);
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvSubgroupBroadcastFirst(bool value)
|
||||
{
|
||||
return !!simd_broadcast_first((ushort)value);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline vec<bool, N> spvSubgroupBroadcastFirst(vec<bool, N> value)
|
||||
{
|
||||
return (vec<bool, N>)simd_broadcast_first((vec<ushort, N>)value);
|
||||
}
|
||||
|
||||
inline uint4 spvSubgroupBallot(bool value)
|
||||
{
|
||||
return uint4((simd_vote::vote_t)simd_ballot(value), 0, 0, 0);
|
||||
}
|
||||
|
||||
inline bool spvSubgroupBallotBitExtract(uint4 ballot, uint bit)
|
||||
{
|
||||
return !!extract_bits(ballot[bit / 32], bit % 32, 1);
|
||||
}
|
||||
|
||||
inline uint spvSubgroupBallotFindLSB(uint4 ballot, uint gl_SubgroupSize)
|
||||
{
|
||||
uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, gl_SubgroupSize), uint3(0));
|
||||
ballot &= mask;
|
||||
return select(ctz(ballot.x), select(32 + ctz(ballot.y), select(64 + ctz(ballot.z), select(96 + ctz(ballot.w), uint(-1), ballot.w == 0), ballot.z == 0), ballot.y == 0), ballot.x == 0);
|
||||
}
|
||||
|
||||
inline uint spvSubgroupBallotFindMSB(uint4 ballot, uint gl_SubgroupSize)
|
||||
{
|
||||
uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, gl_SubgroupSize), uint3(0));
|
||||
ballot &= mask;
|
||||
return select(128 - (clz(ballot.w) + 1), select(96 - (clz(ballot.z) + 1), select(64 - (clz(ballot.y) + 1), select(32 - (clz(ballot.x) + 1), uint(-1), ballot.x == 0), ballot.y == 0), ballot.z == 0), ballot.w == 0);
|
||||
}
|
||||
|
||||
inline uint spvPopCount4(uint4 ballot)
|
||||
{
|
||||
return popcount(ballot.x) + popcount(ballot.y) + popcount(ballot.z) + popcount(ballot.w);
|
||||
}
|
||||
|
||||
inline uint spvSubgroupBallotBitCount(uint4 ballot, uint gl_SubgroupSize)
|
||||
{
|
||||
uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, gl_SubgroupSize), uint3(0));
|
||||
return spvPopCount4(ballot & mask);
|
||||
}
|
||||
|
||||
inline uint spvSubgroupBallotInclusiveBitCount(uint4 ballot, uint gl_SubgroupInvocationID)
|
||||
{
|
||||
uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, gl_SubgroupInvocationID + 1), uint3(0));
|
||||
return spvPopCount4(ballot & mask);
|
||||
}
|
||||
|
||||
inline uint spvSubgroupBallotExclusiveBitCount(uint4 ballot, uint gl_SubgroupInvocationID)
|
||||
{
|
||||
uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, gl_SubgroupInvocationID), uint2(0));
|
||||
return spvPopCount4(ballot & mask);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline bool spvSubgroupAllEqual(T value)
|
||||
{
|
||||
return simd_all(all(value == simd_broadcast_first(value)));
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvSubgroupAllEqual(bool value)
|
||||
{
|
||||
return simd_all(value) || !simd_any(value);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline bool spvSubgroupAllEqual(vec<bool, N> value)
|
||||
{
|
||||
return simd_all(all(value == (vec<bool, N>)simd_broadcast_first((vec<ushort, N>)value)));
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline T spvSubgroupShuffle(T value, ushort lane)
|
||||
{
|
||||
return simd_shuffle(value, lane);
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvSubgroupShuffle(bool value, ushort lane)
|
||||
{
|
||||
return !!simd_shuffle((ushort)value, lane);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline vec<bool, N> spvSubgroupShuffle(vec<bool, N> value, ushort lane)
|
||||
{
|
||||
return (vec<bool, N>)simd_shuffle((vec<ushort, N>)value, lane);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline T spvSubgroupShuffleXor(T value, ushort mask)
|
||||
{
|
||||
return simd_shuffle_xor(value, mask);
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvSubgroupShuffleXor(bool value, ushort mask)
|
||||
{
|
||||
return !!simd_shuffle_xor((ushort)value, mask);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline vec<bool, N> spvSubgroupShuffleXor(vec<bool, N> value, ushort mask)
|
||||
{
|
||||
return (vec<bool, N>)simd_shuffle_xor((vec<ushort, N>)value, mask);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline T spvSubgroupShuffleUp(T value, ushort delta)
|
||||
{
|
||||
return simd_shuffle_up(value, delta);
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvSubgroupShuffleUp(bool value, ushort delta)
|
||||
{
|
||||
return !!simd_shuffle_up((ushort)value, delta);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline vec<bool, N> spvSubgroupShuffleUp(vec<bool, N> value, ushort delta)
|
||||
{
|
||||
return (vec<bool, N>)simd_shuffle_up((vec<ushort, N>)value, delta);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline T spvSubgroupShuffleDown(T value, ushort delta)
|
||||
{
|
||||
return simd_shuffle_down(value, delta);
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvSubgroupShuffleDown(bool value, ushort delta)
|
||||
{
|
||||
return !!simd_shuffle_down((ushort)value, delta);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline vec<bool, N> spvSubgroupShuffleDown(vec<bool, N> value, ushort delta)
|
||||
{
|
||||
return (vec<bool, N>)simd_shuffle_down((vec<ushort, N>)value, delta);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline T spvQuadBroadcast(T value, uint lane)
|
||||
{
|
||||
return quad_broadcast(value, lane);
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvQuadBroadcast(bool value, uint lane)
|
||||
{
|
||||
return !!quad_broadcast((ushort)value, lane);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline vec<bool, N> spvQuadBroadcast(vec<bool, N> value, uint lane)
|
||||
{
|
||||
return (vec<bool, N>)quad_broadcast((vec<ushort, N>)value, lane);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline T spvQuadSwap(T value, uint dir)
|
||||
{
|
||||
return quad_shuffle_xor(value, dir + 1);
|
||||
}
|
||||
|
||||
template<>
|
||||
inline bool spvQuadSwap(bool value, uint dir)
|
||||
{
|
||||
return !!quad_shuffle_xor((ushort)value, dir + 1);
|
||||
}
|
||||
|
||||
template<uint N>
|
||||
inline vec<bool, N> spvQuadSwap(vec<bool, N> value, uint dir)
|
||||
{
|
||||
return (vec<bool, N>)quad_shuffle_xor((vec<ushort, N>)value, dir + 1);
|
||||
}
|
||||
|
||||
kernel void main0(device SSBO& _9 [[buffer(0)]], uint gl_NumSubgroups [[quadgroups_per_threadgroup]], uint gl_SubgroupID [[quadgroup_index_in_threadgroup]], uint gl_SubgroupSize [[thread_execution_width]], uint gl_SubgroupInvocationID [[thread_index_in_quadgroup]])
|
||||
{
|
||||
uint4 gl_SubgroupEqMask = uint4(1 << gl_SubgroupInvocationID, uint3(0));
|
||||
uint4 gl_SubgroupGeMask = uint4(insert_bits(0u, 0xFFFFFFFF, gl_SubgroupInvocationID, gl_SubgroupSize - gl_SubgroupInvocationID), uint3(0));
|
||||
uint4 gl_SubgroupGtMask = uint4(insert_bits(0u, 0xFFFFFFFF, gl_SubgroupInvocationID + 1, gl_SubgroupSize - gl_SubgroupInvocationID - 1), uint3(0));
|
||||
uint4 gl_SubgroupLeMask = uint4(extract_bits(0xFFFFFFFF, 0, gl_SubgroupInvocationID + 1), uint3(0));
|
||||
uint4 gl_SubgroupLtMask = uint4(extract_bits(0xFFFFFFFF, 0, gl_SubgroupInvocationID), uint3(0));
|
||||
_9.FragColor = float(gl_NumSubgroups);
|
||||
_9.FragColor = float(gl_SubgroupID);
|
||||
_9.FragColor = float(gl_SubgroupSize);
|
||||
_9.FragColor = float(gl_SubgroupInvocationID);
|
||||
simdgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture);
|
||||
simdgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture);
|
||||
simdgroup_barrier(mem_flags::mem_device);
|
||||
simdgroup_barrier(mem_flags::mem_threadgroup);
|
||||
simdgroup_barrier(mem_flags::mem_texture);
|
||||
bool _39 = simd_is_first();
|
||||
bool elected = _39;
|
||||
_9.FragColor = float4(gl_SubgroupEqMask).x;
|
||||
_9.FragColor = float4(gl_SubgroupGeMask).x;
|
||||
_9.FragColor = float4(gl_SubgroupGtMask).x;
|
||||
_9.FragColor = float4(gl_SubgroupLeMask).x;
|
||||
_9.FragColor = float4(gl_SubgroupLtMask).x;
|
||||
float4 broadcasted = spvSubgroupBroadcast(float4(10.0), 8u);
|
||||
bool2 broadcasted_bool = spvSubgroupBroadcast(bool2(true), 8u);
|
||||
float3 first = spvSubgroupBroadcastFirst(float3(20.0));
|
||||
bool4 first_bool = spvSubgroupBroadcastFirst(bool4(false));
|
||||
uint4 ballot_value = spvSubgroupBallot(true);
|
||||
bool inverse_ballot_value = spvSubgroupBallotBitExtract(ballot_value, gl_SubgroupInvocationID);
|
||||
bool bit_extracted = spvSubgroupBallotBitExtract(uint4(10u), 8u);
|
||||
uint bit_count = spvSubgroupBallotBitCount(ballot_value, gl_SubgroupSize);
|
||||
uint inclusive_bit_count = spvSubgroupBallotInclusiveBitCount(ballot_value, gl_SubgroupInvocationID);
|
||||
uint exclusive_bit_count = spvSubgroupBallotExclusiveBitCount(ballot_value, gl_SubgroupInvocationID);
|
||||
uint lsb = spvSubgroupBallotFindLSB(ballot_value, gl_SubgroupSize);
|
||||
uint msb = spvSubgroupBallotFindMSB(ballot_value, gl_SubgroupSize);
|
||||
uint shuffled = spvSubgroupShuffle(10u, 8u);
|
||||
bool shuffled_bool = spvSubgroupShuffle(true, 9u);
|
||||
uint shuffled_xor = spvSubgroupShuffleXor(30u, 8u);
|
||||
bool shuffled_xor_bool = spvSubgroupShuffleXor(false, 9u);
|
||||
uint shuffled_up = spvSubgroupShuffleUp(20u, 4u);
|
||||
bool shuffled_up_bool = spvSubgroupShuffleUp(true, 4u);
|
||||
uint shuffled_down = spvSubgroupShuffleDown(20u, 4u);
|
||||
bool shuffled_down_bool = spvSubgroupShuffleDown(false, 4u);
|
||||
bool has_all = simd_all(true);
|
||||
bool has_any = simd_any(true);
|
||||
bool has_equal = spvSubgroupAllEqual(0);
|
||||
has_equal = spvSubgroupAllEqual(true);
|
||||
has_equal = spvSubgroupAllEqual(float3(0.0, 1.0, 2.0));
|
||||
has_equal = spvSubgroupAllEqual(bool4(true, true, false, true));
|
||||
float4 added = simd_sum(float4(20.0));
|
||||
int4 iadded = simd_sum(int4(20));
|
||||
float4 multiplied = simd_product(float4(20.0));
|
||||
int4 imultiplied = simd_product(int4(20));
|
||||
float4 lo = simd_min(float4(20.0));
|
||||
float4 hi = simd_max(float4(20.0));
|
||||
int4 slo = simd_min(int4(20));
|
||||
int4 shi = simd_max(int4(20));
|
||||
uint4 ulo = simd_min(uint4(20u));
|
||||
uint4 uhi = simd_max(uint4(20u));
|
||||
uint4 anded = simd_and(ballot_value);
|
||||
uint4 ored = simd_or(ballot_value);
|
||||
uint4 xored = simd_xor(ballot_value);
|
||||
added = simd_prefix_inclusive_sum(added);
|
||||
iadded = simd_prefix_inclusive_sum(iadded);
|
||||
multiplied = simd_prefix_inclusive_product(multiplied);
|
||||
imultiplied = simd_prefix_inclusive_product(imultiplied);
|
||||
added = simd_prefix_exclusive_sum(multiplied);
|
||||
multiplied = simd_prefix_exclusive_product(multiplied);
|
||||
iadded = simd_prefix_exclusive_sum(imultiplied);
|
||||
imultiplied = simd_prefix_exclusive_product(imultiplied);
|
||||
added = quad_sum(added);
|
||||
multiplied = quad_product(multiplied);
|
||||
iadded = quad_sum(iadded);
|
||||
imultiplied = quad_product(imultiplied);
|
||||
lo = quad_min(lo);
|
||||
hi = quad_max(hi);
|
||||
ulo = quad_min(ulo);
|
||||
uhi = quad_max(uhi);
|
||||
slo = quad_min(slo);
|
||||
shi = quad_max(shi);
|
||||
anded = quad_and(anded);
|
||||
ored = quad_or(ored);
|
||||
xored = quad_xor(xored);
|
||||
float4 swap_horiz = spvQuadSwap(float4(20.0), 0u);
|
||||
bool4 swap_horiz_bool = spvQuadSwap(bool4(true), 0u);
|
||||
float4 swap_vertical = spvQuadSwap(float4(20.0), 1u);
|
||||
bool4 swap_vertical_bool = spvQuadSwap(bool4(true), 1u);
|
||||
float4 swap_diagonal = spvQuadSwap(float4(20.0), 2u);
|
||||
bool4 swap_diagonal_bool = spvQuadSwap(bool4(true), 2u);
|
||||
float4 quad_broadcast0 = spvQuadBroadcast(float4(20.0), 3u);
|
||||
bool4 quad_broadcast_bool = spvQuadBroadcast(bool4(true), 3u);
|
||||
}
|
||||
|
||||
@ -224,11 +224,9 @@ inline vec<bool, N> spvQuadSwap(vec<bool, N> value, uint dir)
|
||||
return (vec<bool, N>)quad_shuffle_xor((vec<ushort, N>)value, dir + 1);
|
||||
}
|
||||
|
||||
fragment main0_out main0()
|
||||
fragment main0_out main0(uint gl_SubgroupSize [[threads_per_simdgroup]], uint gl_SubgroupInvocationID [[thread_index_in_simdgroup]])
|
||||
{
|
||||
main0_out out = {};
|
||||
uint gl_SubgroupSize = simd_sum(1);
|
||||
uint gl_SubgroupInvocationID = simd_prefix_exclusive_sum(1);
|
||||
uint4 gl_SubgroupEqMask = gl_SubgroupInvocationID >= 32 ? uint4(0, (1 << (gl_SubgroupInvocationID - 32)), uint2(0)) : uint4(1 << gl_SubgroupInvocationID, uint3(0));
|
||||
uint4 gl_SubgroupGeMask = uint4(insert_bits(0u, 0xFFFFFFFF, min(gl_SubgroupInvocationID, 32u), (uint)max(min((int)gl_SubgroupSize, 32) - (int)gl_SubgroupInvocationID, 0)), insert_bits(0u, 0xFFFFFFFF, (uint)max((int)gl_SubgroupInvocationID - 32, 0), (uint)max((int)gl_SubgroupSize - (int)max(gl_SubgroupInvocationID, 32u), 0)), uint2(0));
|
||||
uint4 gl_SubgroupGtMask = uint4(insert_bits(0u, 0xFFFFFFFF, min(gl_SubgroupInvocationID + 1, 32u), (uint)max(min((int)gl_SubgroupSize, 32) - (int)gl_SubgroupInvocationID - 1, 0)), insert_bits(0u, 0xFFFFFFFF, (uint)max((int)gl_SubgroupInvocationID + 1 - 32, 0), (uint)max((int)gl_SubgroupSize - (int)max(gl_SubgroupInvocationID + 1, 32u), 0)), uint2(0));
|
||||
@ -236,7 +234,8 @@ fragment main0_out main0()
|
||||
uint4 gl_SubgroupLtMask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupInvocationID, 32u)), extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupInvocationID - 32, 0)), uint2(0));
|
||||
out.FragColor = float(gl_SubgroupSize);
|
||||
out.FragColor = float(gl_SubgroupInvocationID);
|
||||
bool elected = simd_is_first();
|
||||
bool _24 = simd_is_first();
|
||||
bool elected = _24;
|
||||
out.FragColor = float4(gl_SubgroupEqMask).x;
|
||||
out.FragColor = float4(gl_SubgroupGeMask).x;
|
||||
out.FragColor = float4(gl_SubgroupGtMask).x;
|
||||
@ -0,0 +1,25 @@
|
||||
#version 450
|
||||
#extension GL_KHR_shader_subgroup_basic : require
|
||||
layout(local_size_x = 1) in;
|
||||
|
||||
layout(std430, binding = 0) buffer SSBO
|
||||
{
|
||||
float FragColor;
|
||||
};
|
||||
|
||||
// Reduced test for emulated functionality.
|
||||
|
||||
void main()
|
||||
{
|
||||
// basic
|
||||
FragColor = float(gl_NumSubgroups);
|
||||
FragColor = float(gl_SubgroupID);
|
||||
FragColor = float(gl_SubgroupSize);
|
||||
FragColor = float(gl_SubgroupInvocationID);
|
||||
subgroupBarrier();
|
||||
subgroupMemoryBarrier();
|
||||
subgroupMemoryBarrierBuffer();
|
||||
subgroupMemoryBarrierShared();
|
||||
subgroupMemoryBarrierImage();
|
||||
bool elected = subgroupElect();
|
||||
}
|
||||
@ -0,0 +1,138 @@
|
||||
#version 450
|
||||
#extension GL_KHR_shader_subgroup_basic : require
|
||||
#extension GL_KHR_shader_subgroup_ballot : require
|
||||
#extension GL_KHR_shader_subgroup_vote : require
|
||||
#extension GL_KHR_shader_subgroup_shuffle : require
|
||||
#extension GL_KHR_shader_subgroup_shuffle_relative : require
|
||||
#extension GL_KHR_shader_subgroup_arithmetic : require
|
||||
#extension GL_KHR_shader_subgroup_clustered : require
|
||||
#extension GL_KHR_shader_subgroup_quad : require
|
||||
layout(local_size_x = 1) in;
|
||||
|
||||
layout(std430, binding = 0) buffer SSBO
|
||||
{
|
||||
float FragColor;
|
||||
};
|
||||
|
||||
void main()
|
||||
{
|
||||
// basic
|
||||
FragColor = float(gl_NumSubgroups);
|
||||
FragColor = float(gl_SubgroupID);
|
||||
FragColor = float(gl_SubgroupSize);
|
||||
FragColor = float(gl_SubgroupInvocationID);
|
||||
subgroupBarrier();
|
||||
subgroupMemoryBarrier();
|
||||
subgroupMemoryBarrierBuffer();
|
||||
subgroupMemoryBarrierShared();
|
||||
subgroupMemoryBarrierImage();
|
||||
bool elected = subgroupElect();
|
||||
|
||||
// ballot
|
||||
FragColor = float(gl_SubgroupEqMask);
|
||||
FragColor = float(gl_SubgroupGeMask);
|
||||
FragColor = float(gl_SubgroupGtMask);
|
||||
FragColor = float(gl_SubgroupLeMask);
|
||||
FragColor = float(gl_SubgroupLtMask);
|
||||
vec4 broadcasted = subgroupBroadcast(vec4(10.0), 8u);
|
||||
bvec2 broadcasted_bool = subgroupBroadcast(bvec2(true), 8u);
|
||||
vec3 first = subgroupBroadcastFirst(vec3(20.0));
|
||||
bvec4 first_bool = subgroupBroadcastFirst(bvec4(false));
|
||||
uvec4 ballot_value = subgroupBallot(true);
|
||||
bool inverse_ballot_value = subgroupInverseBallot(ballot_value);
|
||||
bool bit_extracted = subgroupBallotBitExtract(uvec4(10u), 8u);
|
||||
uint bit_count = subgroupBallotBitCount(ballot_value);
|
||||
uint inclusive_bit_count = subgroupBallotInclusiveBitCount(ballot_value);
|
||||
uint exclusive_bit_count = subgroupBallotExclusiveBitCount(ballot_value);
|
||||
uint lsb = subgroupBallotFindLSB(ballot_value);
|
||||
uint msb = subgroupBallotFindMSB(ballot_value);
|
||||
|
||||
// shuffle
|
||||
uint shuffled = subgroupShuffle(10u, 8u);
|
||||
bool shuffled_bool = subgroupShuffle(true, 9u);
|
||||
uint shuffled_xor = subgroupShuffleXor(30u, 8u);
|
||||
bool shuffled_xor_bool = subgroupShuffleXor(false, 9u);
|
||||
|
||||
// shuffle relative
|
||||
uint shuffled_up = subgroupShuffleUp(20u, 4u);
|
||||
bool shuffled_up_bool = subgroupShuffleUp(true, 4u);
|
||||
uint shuffled_down = subgroupShuffleDown(20u, 4u);
|
||||
bool shuffled_down_bool = subgroupShuffleDown(false, 4u);
|
||||
|
||||
// vote
|
||||
bool has_all = subgroupAll(true);
|
||||
bool has_any = subgroupAny(true);
|
||||
bool has_equal = subgroupAllEqual(0);
|
||||
has_equal = subgroupAllEqual(true);
|
||||
has_equal = subgroupAllEqual(vec3(0.0, 1.0, 2.0));
|
||||
has_equal = subgroupAllEqual(bvec4(true, true, false, true));
|
||||
|
||||
// arithmetic
|
||||
vec4 added = subgroupAdd(vec4(20.0));
|
||||
ivec4 iadded = subgroupAdd(ivec4(20));
|
||||
vec4 multiplied = subgroupMul(vec4(20.0));
|
||||
ivec4 imultiplied = subgroupMul(ivec4(20));
|
||||
vec4 lo = subgroupMin(vec4(20.0));
|
||||
vec4 hi = subgroupMax(vec4(20.0));
|
||||
ivec4 slo = subgroupMin(ivec4(20));
|
||||
ivec4 shi = subgroupMax(ivec4(20));
|
||||
uvec4 ulo = subgroupMin(uvec4(20));
|
||||
uvec4 uhi = subgroupMax(uvec4(20));
|
||||
uvec4 anded = subgroupAnd(ballot_value);
|
||||
uvec4 ored = subgroupOr(ballot_value);
|
||||
uvec4 xored = subgroupXor(ballot_value);
|
||||
|
||||
added = subgroupInclusiveAdd(added);
|
||||
iadded = subgroupInclusiveAdd(iadded);
|
||||
multiplied = subgroupInclusiveMul(multiplied);
|
||||
imultiplied = subgroupInclusiveMul(imultiplied);
|
||||
//lo = subgroupInclusiveMin(lo); // FIXME: Unsupported by Metal
|
||||
//hi = subgroupInclusiveMax(hi);
|
||||
//slo = subgroupInclusiveMin(slo);
|
||||
//shi = subgroupInclusiveMax(shi);
|
||||
//ulo = subgroupInclusiveMin(ulo);
|
||||
//uhi = subgroupInclusiveMax(uhi);
|
||||
//anded = subgroupInclusiveAnd(anded);
|
||||
//ored = subgroupInclusiveOr(ored);
|
||||
//xored = subgroupInclusiveXor(ored);
|
||||
//added = subgroupExclusiveAdd(lo);
|
||||
|
||||
added = subgroupExclusiveAdd(multiplied);
|
||||
multiplied = subgroupExclusiveMul(multiplied);
|
||||
iadded = subgroupExclusiveAdd(imultiplied);
|
||||
imultiplied = subgroupExclusiveMul(imultiplied);
|
||||
//lo = subgroupExclusiveMin(lo); // FIXME: Unsupported by Metal
|
||||
//hi = subgroupExclusiveMax(hi);
|
||||
//ulo = subgroupExclusiveMin(ulo);
|
||||
//uhi = subgroupExclusiveMax(uhi);
|
||||
//slo = subgroupExclusiveMin(slo);
|
||||
//shi = subgroupExclusiveMax(shi);
|
||||
//anded = subgroupExclusiveAnd(anded);
|
||||
//ored = subgroupExclusiveOr(ored);
|
||||
//xored = subgroupExclusiveXor(ored);
|
||||
|
||||
// clustered
|
||||
added = subgroupClusteredAdd(added, 4u);
|
||||
multiplied = subgroupClusteredMul(multiplied, 4u);
|
||||
iadded = subgroupClusteredAdd(iadded, 4u);
|
||||
imultiplied = subgroupClusteredMul(imultiplied, 4u);
|
||||
lo = subgroupClusteredMin(lo, 4u);
|
||||
hi = subgroupClusteredMax(hi, 4u);
|
||||
ulo = subgroupClusteredMin(ulo, 4u);
|
||||
uhi = subgroupClusteredMax(uhi, 4u);
|
||||
slo = subgroupClusteredMin(slo, 4u);
|
||||
shi = subgroupClusteredMax(shi, 4u);
|
||||
anded = subgroupClusteredAnd(anded, 4u);
|
||||
ored = subgroupClusteredOr(ored, 4u);
|
||||
xored = subgroupClusteredXor(xored, 4u);
|
||||
|
||||
// quad
|
||||
vec4 swap_horiz = subgroupQuadSwapHorizontal(vec4(20.0));
|
||||
bvec4 swap_horiz_bool = subgroupQuadSwapHorizontal(bvec4(true));
|
||||
vec4 swap_vertical = subgroupQuadSwapVertical(vec4(20.0));
|
||||
bvec4 swap_vertical_bool = subgroupQuadSwapVertical(bvec4(true));
|
||||
vec4 swap_diagonal = subgroupQuadSwapDiagonal(vec4(20.0));
|
||||
bvec4 swap_diagonal_bool = subgroupQuadSwapDiagonal(bvec4(true));
|
||||
vec4 quad_broadcast = subgroupQuadBroadcast(vec4(20.0), 3u);
|
||||
bvec4 quad_broadcast_bool = subgroupQuadBroadcast(bvec4(true), 3u);
|
||||
}
|
||||
@ -0,0 +1,79 @@
|
||||
#version 450
|
||||
#extension GL_KHR_shader_subgroup_basic : require
|
||||
#extension GL_KHR_shader_subgroup_ballot : require
|
||||
#extension GL_KHR_shader_subgroup_vote : require
|
||||
#extension GL_KHR_shader_subgroup_shuffle : require
|
||||
#extension GL_KHR_shader_subgroup_shuffle_relative : require
|
||||
#extension GL_KHR_shader_subgroup_quad : require
|
||||
layout(local_size_x = 1) in;
|
||||
|
||||
layout(std430, binding = 0) buffer SSBO
|
||||
{
|
||||
float FragColor;
|
||||
};
|
||||
|
||||
// Reduced test for functionality exposed on iOS.
|
||||
|
||||
void main()
|
||||
{
|
||||
// basic
|
||||
FragColor = float(gl_NumSubgroups);
|
||||
FragColor = float(gl_SubgroupID);
|
||||
FragColor = float(gl_SubgroupSize);
|
||||
FragColor = float(gl_SubgroupInvocationID);
|
||||
subgroupBarrier();
|
||||
subgroupMemoryBarrier();
|
||||
subgroupMemoryBarrierBuffer();
|
||||
subgroupMemoryBarrierShared();
|
||||
subgroupMemoryBarrierImage();
|
||||
bool elected = subgroupElect();
|
||||
|
||||
// ballot
|
||||
FragColor = float(gl_SubgroupEqMask);
|
||||
FragColor = float(gl_SubgroupGeMask);
|
||||
FragColor = float(gl_SubgroupGtMask);
|
||||
FragColor = float(gl_SubgroupLeMask);
|
||||
FragColor = float(gl_SubgroupLtMask);
|
||||
vec4 broadcasted = subgroupBroadcast(vec4(10.0), 8u);
|
||||
bvec2 broadcasted_bool = subgroupBroadcast(bvec2(true), 8u);
|
||||
vec3 first = subgroupBroadcastFirst(vec3(20.0));
|
||||
bvec4 first_bool = subgroupBroadcastFirst(bvec4(false));
|
||||
uvec4 ballot_value = subgroupBallot(true);
|
||||
bool inverse_ballot_value = subgroupInverseBallot(ballot_value);
|
||||
bool bit_extracted = subgroupBallotBitExtract(uvec4(10u), 8u);
|
||||
uint bit_count = subgroupBallotBitCount(ballot_value);
|
||||
uint inclusive_bit_count = subgroupBallotInclusiveBitCount(ballot_value);
|
||||
uint exclusive_bit_count = subgroupBallotExclusiveBitCount(ballot_value);
|
||||
uint lsb = subgroupBallotFindLSB(ballot_value);
|
||||
uint msb = subgroupBallotFindMSB(ballot_value);
|
||||
|
||||
// shuffle
|
||||
uint shuffled = subgroupShuffle(10u, 8u);
|
||||
bool shuffled_bool = subgroupShuffle(true, 9u);
|
||||
uint shuffled_xor = subgroupShuffleXor(30u, 8u);
|
||||
bool shuffled_xor_bool = subgroupShuffleXor(false, 9u);
|
||||
|
||||
// shuffle relative
|
||||
uint shuffled_up = subgroupShuffleUp(20u, 4u);
|
||||
bool shuffled_up_bool = subgroupShuffleUp(true, 4u);
|
||||
uint shuffled_down = subgroupShuffleDown(20u, 4u);
|
||||
bool shuffled_down_bool = subgroupShuffleDown(false, 4u);
|
||||
|
||||
// vote
|
||||
bool has_all = subgroupAll(true);
|
||||
bool has_any = subgroupAny(true);
|
||||
bool has_equal = subgroupAllEqual(0);
|
||||
has_equal = subgroupAllEqual(true);
|
||||
has_equal = subgroupAllEqual(vec3(0.0, 1.0, 2.0));
|
||||
has_equal = subgroupAllEqual(bvec4(true, true, false, true));
|
||||
|
||||
// quad
|
||||
vec4 swap_horiz = subgroupQuadSwapHorizontal(vec4(20.0));
|
||||
bvec4 swap_horiz_bool = subgroupQuadSwapHorizontal(bvec4(true));
|
||||
vec4 swap_vertical = subgroupQuadSwapVertical(vec4(20.0));
|
||||
bvec4 swap_vertical_bool = subgroupQuadSwapVertical(bvec4(true));
|
||||
vec4 swap_diagonal = subgroupQuadSwapDiagonal(vec4(20.0));
|
||||
bvec4 swap_diagonal_bool = subgroupQuadSwapDiagonal(bvec4(true));
|
||||
vec4 quad_broadcast = subgroupQuadBroadcast(vec4(20.0), 3u);
|
||||
bvec4 quad_broadcast_bool = subgroupQuadBroadcast(bvec4(true), 3u);
|
||||
}
|
||||
@ -0,0 +1,138 @@
|
||||
#version 450
|
||||
#extension GL_KHR_shader_subgroup_basic : require
|
||||
#extension GL_KHR_shader_subgroup_ballot : require
|
||||
#extension GL_KHR_shader_subgroup_vote : require
|
||||
#extension GL_KHR_shader_subgroup_shuffle : require
|
||||
#extension GL_KHR_shader_subgroup_shuffle_relative : require
|
||||
#extension GL_KHR_shader_subgroup_arithmetic : require
|
||||
#extension GL_KHR_shader_subgroup_clustered : require
|
||||
#extension GL_KHR_shader_subgroup_quad : require
|
||||
layout(local_size_x = 1) in;
|
||||
|
||||
layout(std430, binding = 0) buffer SSBO
|
||||
{
|
||||
float FragColor;
|
||||
};
|
||||
|
||||
void main()
|
||||
{
|
||||
// basic
|
||||
FragColor = float(gl_NumSubgroups);
|
||||
FragColor = float(gl_SubgroupID);
|
||||
FragColor = float(gl_SubgroupSize);
|
||||
FragColor = float(gl_SubgroupInvocationID);
|
||||
subgroupBarrier();
|
||||
subgroupMemoryBarrier();
|
||||
subgroupMemoryBarrierBuffer();
|
||||
subgroupMemoryBarrierShared();
|
||||
subgroupMemoryBarrierImage();
|
||||
bool elected = subgroupElect();
|
||||
|
||||
// ballot
|
||||
FragColor = float(gl_SubgroupEqMask);
|
||||
FragColor = float(gl_SubgroupGeMask);
|
||||
FragColor = float(gl_SubgroupGtMask);
|
||||
FragColor = float(gl_SubgroupLeMask);
|
||||
FragColor = float(gl_SubgroupLtMask);
|
||||
vec4 broadcasted = subgroupBroadcast(vec4(10.0), 8u);
|
||||
bvec2 broadcasted_bool = subgroupBroadcast(bvec2(true), 8u);
|
||||
vec3 first = subgroupBroadcastFirst(vec3(20.0));
|
||||
bvec4 first_bool = subgroupBroadcastFirst(bvec4(false));
|
||||
uvec4 ballot_value = subgroupBallot(true);
|
||||
bool inverse_ballot_value = subgroupInverseBallot(ballot_value);
|
||||
bool bit_extracted = subgroupBallotBitExtract(uvec4(10u), 8u);
|
||||
uint bit_count = subgroupBallotBitCount(ballot_value);
|
||||
uint inclusive_bit_count = subgroupBallotInclusiveBitCount(ballot_value);
|
||||
uint exclusive_bit_count = subgroupBallotExclusiveBitCount(ballot_value);
|
||||
uint lsb = subgroupBallotFindLSB(ballot_value);
|
||||
uint msb = subgroupBallotFindMSB(ballot_value);
|
||||
|
||||
// shuffle
|
||||
uint shuffled = subgroupShuffle(10u, 8u);
|
||||
bool shuffled_bool = subgroupShuffle(true, 9u);
|
||||
uint shuffled_xor = subgroupShuffleXor(30u, 8u);
|
||||
bool shuffled_xor_bool = subgroupShuffleXor(false, 9u);
|
||||
|
||||
// shuffle relative
|
||||
uint shuffled_up = subgroupShuffleUp(20u, 4u);
|
||||
bool shuffled_up_bool = subgroupShuffleUp(true, 4u);
|
||||
uint shuffled_down = subgroupShuffleDown(20u, 4u);
|
||||
bool shuffled_down_bool = subgroupShuffleDown(false, 4u);
|
||||
|
||||
// vote
|
||||
bool has_all = subgroupAll(true);
|
||||
bool has_any = subgroupAny(true);
|
||||
bool has_equal = subgroupAllEqual(0);
|
||||
has_equal = subgroupAllEqual(true);
|
||||
has_equal = subgroupAllEqual(vec3(0.0, 1.0, 2.0));
|
||||
has_equal = subgroupAllEqual(bvec4(true, true, false, true));
|
||||
|
||||
// arithmetic
|
||||
vec4 added = subgroupAdd(vec4(20.0));
|
||||
ivec4 iadded = subgroupAdd(ivec4(20));
|
||||
vec4 multiplied = subgroupMul(vec4(20.0));
|
||||
ivec4 imultiplied = subgroupMul(ivec4(20));
|
||||
vec4 lo = subgroupMin(vec4(20.0));
|
||||
vec4 hi = subgroupMax(vec4(20.0));
|
||||
ivec4 slo = subgroupMin(ivec4(20));
|
||||
ivec4 shi = subgroupMax(ivec4(20));
|
||||
uvec4 ulo = subgroupMin(uvec4(20));
|
||||
uvec4 uhi = subgroupMax(uvec4(20));
|
||||
uvec4 anded = subgroupAnd(ballot_value);
|
||||
uvec4 ored = subgroupOr(ballot_value);
|
||||
uvec4 xored = subgroupXor(ballot_value);
|
||||
|
||||
added = subgroupInclusiveAdd(added);
|
||||
iadded = subgroupInclusiveAdd(iadded);
|
||||
multiplied = subgroupInclusiveMul(multiplied);
|
||||
imultiplied = subgroupInclusiveMul(imultiplied);
|
||||
//lo = subgroupInclusiveMin(lo); // FIXME: Unsupported by Metal
|
||||
//hi = subgroupInclusiveMax(hi);
|
||||
//slo = subgroupInclusiveMin(slo);
|
||||
//shi = subgroupInclusiveMax(shi);
|
||||
//ulo = subgroupInclusiveMin(ulo);
|
||||
//uhi = subgroupInclusiveMax(uhi);
|
||||
//anded = subgroupInclusiveAnd(anded);
|
||||
//ored = subgroupInclusiveOr(ored);
|
||||
//xored = subgroupInclusiveXor(ored);
|
||||
//added = subgroupExclusiveAdd(lo);
|
||||
|
||||
added = subgroupExclusiveAdd(multiplied);
|
||||
multiplied = subgroupExclusiveMul(multiplied);
|
||||
iadded = subgroupExclusiveAdd(imultiplied);
|
||||
imultiplied = subgroupExclusiveMul(imultiplied);
|
||||
//lo = subgroupExclusiveMin(lo); // FIXME: Unsupported by Metal
|
||||
//hi = subgroupExclusiveMax(hi);
|
||||
//ulo = subgroupExclusiveMin(ulo);
|
||||
//uhi = subgroupExclusiveMax(uhi);
|
||||
//slo = subgroupExclusiveMin(slo);
|
||||
//shi = subgroupExclusiveMax(shi);
|
||||
//anded = subgroupExclusiveAnd(anded);
|
||||
//ored = subgroupExclusiveOr(ored);
|
||||
//xored = subgroupExclusiveXor(ored);
|
||||
|
||||
// clustered
|
||||
added = subgroupClusteredAdd(added, 4u);
|
||||
multiplied = subgroupClusteredMul(multiplied, 4u);
|
||||
iadded = subgroupClusteredAdd(iadded, 4u);
|
||||
imultiplied = subgroupClusteredMul(imultiplied, 4u);
|
||||
lo = subgroupClusteredMin(lo, 4u);
|
||||
hi = subgroupClusteredMax(hi, 4u);
|
||||
ulo = subgroupClusteredMin(ulo, 4u);
|
||||
uhi = subgroupClusteredMax(uhi, 4u);
|
||||
slo = subgroupClusteredMin(slo, 4u);
|
||||
shi = subgroupClusteredMax(shi, 4u);
|
||||
anded = subgroupClusteredAnd(anded, 4u);
|
||||
ored = subgroupClusteredOr(ored, 4u);
|
||||
xored = subgroupClusteredXor(xored, 4u);
|
||||
|
||||
// quad
|
||||
vec4 swap_horiz = subgroupQuadSwapHorizontal(vec4(20.0));
|
||||
bvec4 swap_horiz_bool = subgroupQuadSwapHorizontal(bvec4(true));
|
||||
vec4 swap_vertical = subgroupQuadSwapVertical(vec4(20.0));
|
||||
bvec4 swap_vertical_bool = subgroupQuadSwapVertical(bvec4(true));
|
||||
vec4 swap_diagonal = subgroupQuadSwapDiagonal(vec4(20.0));
|
||||
bvec4 swap_diagonal_bool = subgroupQuadSwapDiagonal(bvec4(true));
|
||||
vec4 quad_broadcast = subgroupQuadBroadcast(vec4(20.0), 3u);
|
||||
bvec4 quad_broadcast_bool = subgroupQuadBroadcast(bvec4(true), 3u);
|
||||
}
|
||||
@ -678,6 +678,18 @@ spvc_result spvc_compiler_options_set_uint(spvc_compiler_options options, spvc_c
|
||||
case SPVC_COMPILER_OPTION_MSL_R32UI_ALIGNMENT_CONSTANT_ID:
|
||||
options->msl.r32ui_alignment_constant_id = value;
|
||||
break;
|
||||
|
||||
case SPVC_COMPILER_OPTION_MSL_IOS_USE_SIMDGROUP_FUNCTIONS:
|
||||
options->msl.ios_use_simdgroup_functions = value != 0;
|
||||
break;
|
||||
|
||||
case SPVC_COMPILER_OPTION_MSL_EMULATE_SUBGROUPS:
|
||||
options->msl.emulate_subgroups = value != 0;
|
||||
break;
|
||||
|
||||
case SPVC_COMPILER_OPTION_MSL_FIXED_SUBGROUP_SIZE:
|
||||
options->msl.fixed_subgroup_size = value;
|
||||
break;
|
||||
#endif
|
||||
|
||||
default:
|
||||
|
||||
@ -33,7 +33,7 @@ extern "C" {
|
||||
/* Bumped if ABI or API breaks backwards compatibility. */
|
||||
#define SPVC_C_API_VERSION_MAJOR 0
|
||||
/* Bumped if APIs or enumerations are added in a backwards compatible way. */
|
||||
#define SPVC_C_API_VERSION_MINOR 42
|
||||
#define SPVC_C_API_VERSION_MINOR 43
|
||||
/* Bumped if internal implementation details change. */
|
||||
#define SPVC_C_API_VERSION_PATCH 0
|
||||
|
||||
@ -647,6 +647,10 @@ typedef enum spvc_compiler_option
|
||||
|
||||
SPVC_COMPILER_OPTION_HLSL_FLATTEN_MATRIX_VERTEX_INPUT_SEMANTICS = 71 | SPVC_COMPILER_OPTION_HLSL_BIT,
|
||||
|
||||
SPVC_COMPILER_OPTION_MSL_IOS_USE_SIMDGROUP_FUNCTIONS = 72 | SPVC_COMPILER_OPTION_MSL_BIT,
|
||||
SPVC_COMPILER_OPTION_MSL_EMULATE_SUBGROUPS = 73 | SPVC_COMPILER_OPTION_MSL_BIT,
|
||||
SPVC_COMPILER_OPTION_MSL_FIXED_SUBGROUP_SIZE = 74 | SPVC_COMPILER_OPTION_MSL_BIT,
|
||||
|
||||
SPVC_COMPILER_OPTION_INT_MAX = 0x7fffffff
|
||||
} spvc_compiler_option;
|
||||
|
||||
|
||||
569
spirv_msl.cpp
569
spirv_msl.cpp
@ -168,9 +168,12 @@ void CompilerMSL::build_implicit_builtins()
|
||||
active_input_builtins.get(BuiltInBaseVertex) || active_input_builtins.get(BuiltInInstanceId) ||
|
||||
active_input_builtins.get(BuiltInInstanceIndex) || active_input_builtins.get(BuiltInBaseInstance));
|
||||
bool need_sample_mask = msl_options.additional_fixed_sample_mask != 0xffffffff;
|
||||
bool need_local_invocation_index = msl_options.emulate_subgroups && active_input_builtins.get(BuiltInSubgroupId);
|
||||
bool need_workgroup_size = msl_options.emulate_subgroups && active_input_builtins.get(BuiltInNumSubgroups);
|
||||
if (need_subpass_input || need_sample_pos || need_subgroup_mask || need_vertex_params || need_tesc_params ||
|
||||
need_multiview || need_dispatch_base || need_vertex_base_params || need_grid_params || needs_sample_id ||
|
||||
needs_subgroup_invocation_id || needs_subgroup_size || need_sample_mask)
|
||||
needs_subgroup_invocation_id || needs_subgroup_size || need_sample_mask || need_local_invocation_index ||
|
||||
need_workgroup_size)
|
||||
{
|
||||
bool has_frag_coord = false;
|
||||
bool has_sample_id = false;
|
||||
@ -184,6 +187,8 @@ void CompilerMSL::build_implicit_builtins()
|
||||
bool has_subgroup_size = false;
|
||||
bool has_view_idx = false;
|
||||
bool has_layer = false;
|
||||
bool has_local_invocation_index = false;
|
||||
bool has_workgroup_size = false;
|
||||
uint32_t workgroup_id_type = 0;
|
||||
|
||||
// FIXME: Investigate the fact that there are no checks for the entry point interface variables.
|
||||
@ -191,7 +196,6 @@ void CompilerMSL::build_implicit_builtins()
|
||||
if (!ir.meta[var.self].decoration.builtin)
|
||||
return;
|
||||
|
||||
// Use Metal's native frame-buffer fetch API for subpass inputs.
|
||||
BuiltIn builtin = ir.meta[var.self].decoration.builtin_type;
|
||||
|
||||
if (var.storage == StorageClassOutput)
|
||||
@ -207,6 +211,7 @@ void CompilerMSL::build_implicit_builtins()
|
||||
if (var.storage != StorageClassInput)
|
||||
return;
|
||||
|
||||
// Use Metal's native frame-buffer fetch API for subpass inputs.
|
||||
if (need_subpass_input && (!msl_options.use_framebuffer_fetch_subpasses))
|
||||
{
|
||||
switch (builtin)
|
||||
@ -330,6 +335,20 @@ void CompilerMSL::build_implicit_builtins()
|
||||
}
|
||||
}
|
||||
|
||||
if (need_local_invocation_index && builtin == BuiltInLocalInvocationIndex)
|
||||
{
|
||||
builtin_local_invocation_index_id = var.self;
|
||||
mark_implicit_builtin(StorageClassInput, BuiltInLocalInvocationIndex, var.self);
|
||||
has_local_invocation_index = true;
|
||||
}
|
||||
|
||||
if (need_workgroup_size && builtin == BuiltInLocalInvocationId)
|
||||
{
|
||||
builtin_workgroup_size_id = var.self;
|
||||
mark_implicit_builtin(StorageClassInput, BuiltInWorkgroupSize, var.self);
|
||||
has_workgroup_size = true;
|
||||
}
|
||||
|
||||
// The base workgroup needs to have the same type and vector size
|
||||
// as the workgroup or invocation ID, so keep track of the type that
|
||||
// was used.
|
||||
@ -681,6 +700,48 @@ void CompilerMSL::build_implicit_builtins()
|
||||
builtin_sample_mask_id = var_id;
|
||||
mark_implicit_builtin(StorageClassOutput, BuiltInSampleMask, var_id);
|
||||
}
|
||||
|
||||
if (need_local_invocation_index && !has_local_invocation_index)
|
||||
{
|
||||
uint32_t offset = ir.increase_bound_by(2);
|
||||
uint32_t type_ptr_id = offset;
|
||||
uint32_t var_id = offset + 1;
|
||||
|
||||
// Create gl_LocalInvocationIndex.
|
||||
SPIRType uint_type_ptr;
|
||||
uint_type_ptr = get_uint_type();
|
||||
uint_type_ptr.pointer = true;
|
||||
uint_type_ptr.parent_type = get_uint_type_id();
|
||||
uint_type_ptr.storage = StorageClassInput;
|
||||
|
||||
auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr);
|
||||
ptr_type.self = get_uint_type_id();
|
||||
set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
|
||||
set_decoration(var_id, DecorationBuiltIn, BuiltInLocalInvocationIndex);
|
||||
builtin_local_invocation_index_id = var_id;
|
||||
mark_implicit_builtin(StorageClassInput, BuiltInLocalInvocationIndex, var_id);
|
||||
}
|
||||
|
||||
if (need_workgroup_size && !has_workgroup_size)
|
||||
{
|
||||
uint32_t offset = ir.increase_bound_by(2);
|
||||
uint32_t type_ptr_id = offset;
|
||||
uint32_t var_id = offset + 1;
|
||||
|
||||
// Create gl_WorkgroupSize.
|
||||
uint32_t type_id = build_extended_vector_type(get_uint_type_id(), 3);
|
||||
SPIRType uint_type_ptr = get<SPIRType>(type_id);
|
||||
uint_type_ptr.pointer = true;
|
||||
uint_type_ptr.parent_type = type_id;
|
||||
uint_type_ptr.storage = StorageClassInput;
|
||||
|
||||
auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr);
|
||||
ptr_type.self = type_id;
|
||||
set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
|
||||
set_decoration(var_id, DecorationBuiltIn, BuiltInWorkgroupSize);
|
||||
builtin_workgroup_size_id = var_id;
|
||||
mark_implicit_builtin(StorageClassInput, BuiltInWorkgroupSize, var_id);
|
||||
}
|
||||
}
|
||||
|
||||
if (needs_swizzle_buffer_def)
|
||||
@ -4787,7 +4848,7 @@ void CompilerMSL::emit_custom_functions()
|
||||
statement("template<typename T>");
|
||||
statement("inline T spvSubgroupBroadcast(T value, ushort lane)");
|
||||
begin_scope();
|
||||
if (msl_options.is_ios())
|
||||
if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
|
||||
statement("return quad_broadcast(value, lane);");
|
||||
else
|
||||
statement("return simd_broadcast(value, lane);");
|
||||
@ -4796,7 +4857,7 @@ void CompilerMSL::emit_custom_functions()
|
||||
statement("template<>");
|
||||
statement("inline bool spvSubgroupBroadcast(bool value, ushort lane)");
|
||||
begin_scope();
|
||||
if (msl_options.is_ios())
|
||||
if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
|
||||
statement("return !!quad_broadcast((ushort)value, lane);");
|
||||
else
|
||||
statement("return !!simd_broadcast((ushort)value, lane);");
|
||||
@ -4805,7 +4866,7 @@ void CompilerMSL::emit_custom_functions()
|
||||
statement("template<uint N>");
|
||||
statement("inline vec<bool, N> spvSubgroupBroadcast(vec<bool, N> value, ushort lane)");
|
||||
begin_scope();
|
||||
if (msl_options.is_ios())
|
||||
if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
|
||||
statement("return (vec<bool, N>)quad_broadcast((vec<ushort, N>)value, lane);");
|
||||
else
|
||||
statement("return (vec<bool, N>)simd_broadcast((vec<ushort, N>)value, lane);");
|
||||
@ -4817,19 +4878,28 @@ void CompilerMSL::emit_custom_functions()
|
||||
statement("template<typename T>");
|
||||
statement("inline T spvSubgroupBroadcastFirst(T value)");
|
||||
begin_scope();
|
||||
statement("return simd_broadcast_first(value);");
|
||||
if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
|
||||
statement("return quad_broadcast_first(value);");
|
||||
else
|
||||
statement("return simd_broadcast_first(value);");
|
||||
end_scope();
|
||||
statement("");
|
||||
statement("template<>");
|
||||
statement("inline bool spvSubgroupBroadcastFirst(bool value)");
|
||||
begin_scope();
|
||||
statement("return !!simd_broadcast_first((ushort)value);");
|
||||
if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
|
||||
statement("return !!quad_broadcast_first((ushort)value);");
|
||||
else
|
||||
statement("return !!simd_broadcast_first((ushort)value);");
|
||||
end_scope();
|
||||
statement("");
|
||||
statement("template<uint N>");
|
||||
statement("inline vec<bool, N> spvSubgroupBroadcastFirst(vec<bool, N> value)");
|
||||
begin_scope();
|
||||
statement("return (vec<bool, N>)simd_broadcast_first((vec<ushort, N>)value);");
|
||||
if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
|
||||
statement("return (vec<bool, N>)quad_broadcast_first((vec<ushort, N>)value);");
|
||||
else
|
||||
statement("return (vec<bool, N>)simd_broadcast_first((vec<ushort, N>)value);");
|
||||
end_scope();
|
||||
statement("");
|
||||
break;
|
||||
@ -4837,13 +4907,26 @@ void CompilerMSL::emit_custom_functions()
|
||||
case SPVFuncImplSubgroupBallot:
|
||||
statement("inline uint4 spvSubgroupBallot(bool value)");
|
||||
begin_scope();
|
||||
statement("simd_vote vote = simd_ballot(value);");
|
||||
statement("// simd_ballot() returns a 64-bit integer-like object, but");
|
||||
statement("// SPIR-V callers expect a uint4. We must convert.");
|
||||
statement("// FIXME: This won't include higher bits if Apple ever supports");
|
||||
statement("// 128 lanes in an SIMD-group.");
|
||||
statement("return uint4((uint)((simd_vote::vote_t)vote & 0xFFFFFFFF), (uint)(((simd_vote::vote_t)vote >> "
|
||||
"32) & 0xFFFFFFFF), 0, 0);");
|
||||
if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
|
||||
{
|
||||
statement("return uint4((quad_vote::vote_t)quad_ballot(value), 0, 0, 0);");
|
||||
}
|
||||
else if (msl_options.is_ios())
|
||||
{
|
||||
// The current simd_vote on iOS uses a 32-bit integer-like object.
|
||||
statement("return uint4((simd_vote::vote_t)simd_ballot(value), 0, 0, 0);");
|
||||
}
|
||||
else
|
||||
{
|
||||
statement("simd_vote vote = simd_ballot(value);");
|
||||
statement("// simd_ballot() returns a 64-bit integer-like object, but");
|
||||
statement("// SPIR-V callers expect a uint4. We must convert.");
|
||||
statement("// FIXME: This won't include higher bits if Apple ever supports");
|
||||
statement("// 128 lanes in an SIMD-group.");
|
||||
statement(
|
||||
"return uint4((uint)((simd_vote::vote_t)vote & 0xFFFFFFFF), (uint)(((simd_vote::vote_t)vote >> "
|
||||
"32) & 0xFFFFFFFF), 0, 0);");
|
||||
}
|
||||
end_scope();
|
||||
statement("");
|
||||
break;
|
||||
@ -4859,8 +4942,15 @@ void CompilerMSL::emit_custom_functions()
|
||||
case SPVFuncImplSubgroupBallotFindLSB:
|
||||
statement("inline uint spvSubgroupBallotFindLSB(uint4 ballot, uint gl_SubgroupSize)");
|
||||
begin_scope();
|
||||
statement("uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupSize, 32u)), "
|
||||
"extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupSize - 32, 0)), uint2(0));");
|
||||
if (msl_options.is_ios())
|
||||
{
|
||||
statement("uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, gl_SubgroupSize), uint3(0));");
|
||||
}
|
||||
else
|
||||
{
|
||||
statement("uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupSize, 32u)), "
|
||||
"extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupSize - 32, 0)), uint2(0));");
|
||||
}
|
||||
statement("ballot &= mask;");
|
||||
statement("return select(ctz(ballot.x), select(32 + ctz(ballot.y), select(64 + ctz(ballot.z), select(96 + "
|
||||
"ctz(ballot.w), uint(-1), ballot.w == 0), ballot.z == 0), ballot.y == 0), ballot.x == 0);");
|
||||
@ -4871,8 +4961,15 @@ void CompilerMSL::emit_custom_functions()
|
||||
case SPVFuncImplSubgroupBallotFindMSB:
|
||||
statement("inline uint spvSubgroupBallotFindMSB(uint4 ballot, uint gl_SubgroupSize)");
|
||||
begin_scope();
|
||||
statement("uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupSize, 32u)), "
|
||||
"extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupSize - 32, 0)), uint2(0));");
|
||||
if (msl_options.is_ios())
|
||||
{
|
||||
statement("uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, gl_SubgroupSize), uint3(0));");
|
||||
}
|
||||
else
|
||||
{
|
||||
statement("uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupSize, 32u)), "
|
||||
"extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupSize - 32, 0)), uint2(0));");
|
||||
}
|
||||
statement("ballot &= mask;");
|
||||
statement("return select(128 - (clz(ballot.w) + 1), select(96 - (clz(ballot.z) + 1), select(64 - "
|
||||
"(clz(ballot.y) + 1), select(32 - (clz(ballot.x) + 1), uint(-1), ballot.x == 0), ballot.y == 0), "
|
||||
@ -4889,23 +4986,44 @@ void CompilerMSL::emit_custom_functions()
|
||||
statement("");
|
||||
statement("inline uint spvSubgroupBallotBitCount(uint4 ballot, uint gl_SubgroupSize)");
|
||||
begin_scope();
|
||||
statement("uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupSize, 32u)), "
|
||||
"extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupSize - 32, 0)), uint2(0));");
|
||||
if (msl_options.is_ios())
|
||||
{
|
||||
statement("uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, gl_SubgroupSize), uint3(0));");
|
||||
}
|
||||
else
|
||||
{
|
||||
statement("uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupSize, 32u)), "
|
||||
"extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupSize - 32, 0)), uint2(0));");
|
||||
}
|
||||
statement("return spvPopCount4(ballot & mask);");
|
||||
end_scope();
|
||||
statement("");
|
||||
statement("inline uint spvSubgroupBallotInclusiveBitCount(uint4 ballot, uint gl_SubgroupInvocationID)");
|
||||
begin_scope();
|
||||
statement("uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupInvocationID + 1, 32u)), "
|
||||
"extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupInvocationID + 1 - 32, 0)), "
|
||||
"uint2(0));");
|
||||
if (msl_options.is_ios())
|
||||
{
|
||||
statement("uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, gl_SubgroupInvocationID + 1), uint3(0));");
|
||||
}
|
||||
else
|
||||
{
|
||||
statement("uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupInvocationID + 1, 32u)), "
|
||||
"extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupInvocationID + 1 - 32, 0)), "
|
||||
"uint2(0));");
|
||||
}
|
||||
statement("return spvPopCount4(ballot & mask);");
|
||||
end_scope();
|
||||
statement("");
|
||||
statement("inline uint spvSubgroupBallotExclusiveBitCount(uint4 ballot, uint gl_SubgroupInvocationID)");
|
||||
begin_scope();
|
||||
statement("uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupInvocationID, 32u)), "
|
||||
"extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupInvocationID - 32, 0)), uint2(0));");
|
||||
if (msl_options.is_ios())
|
||||
{
|
||||
statement("uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, gl_SubgroupInvocationID), uint2(0));");
|
||||
}
|
||||
else
|
||||
{
|
||||
statement("uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupInvocationID, 32u)), "
|
||||
"extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupInvocationID - 32, 0)), uint2(0));");
|
||||
}
|
||||
statement("return spvPopCount4(ballot & mask);");
|
||||
end_scope();
|
||||
statement("");
|
||||
@ -4919,19 +5037,28 @@ void CompilerMSL::emit_custom_functions()
|
||||
statement("template<typename T>");
|
||||
statement("inline bool spvSubgroupAllEqual(T value)");
|
||||
begin_scope();
|
||||
statement("return simd_all(all(value == simd_broadcast_first(value)));");
|
||||
if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
|
||||
statement("return quad_all(all(value == quad_broadcast_first(value)));");
|
||||
else
|
||||
statement("return simd_all(all(value == simd_broadcast_first(value)));");
|
||||
end_scope();
|
||||
statement("");
|
||||
statement("template<>");
|
||||
statement("inline bool spvSubgroupAllEqual(bool value)");
|
||||
begin_scope();
|
||||
statement("return simd_all(value) || !simd_any(value);");
|
||||
if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
|
||||
statement("return quad_all(value) || !quad_any(value);");
|
||||
else
|
||||
statement("return simd_all(value) || !simd_any(value);");
|
||||
end_scope();
|
||||
statement("");
|
||||
statement("template<uint N>");
|
||||
statement("inline bool spvSubgroupAllEqual(vec<bool, N> value)");
|
||||
begin_scope();
|
||||
statement("return simd_all(all(value == (vec<bool, N>)simd_broadcast_first((vec<ushort, N>)value)));");
|
||||
if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
|
||||
statement("return quad_all(all(value == (vec<bool, N>)quad_broadcast_first((vec<ushort, N>)value)));");
|
||||
else
|
||||
statement("return simd_all(all(value == (vec<bool, N>)simd_broadcast_first((vec<ushort, N>)value)));");
|
||||
end_scope();
|
||||
statement("");
|
||||
break;
|
||||
@ -4940,7 +5067,7 @@ void CompilerMSL::emit_custom_functions()
|
||||
statement("template<typename T>");
|
||||
statement("inline T spvSubgroupShuffle(T value, ushort lane)");
|
||||
begin_scope();
|
||||
if (msl_options.is_ios())
|
||||
if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
|
||||
statement("return quad_shuffle(value, lane);");
|
||||
else
|
||||
statement("return simd_shuffle(value, lane);");
|
||||
@ -4949,7 +5076,7 @@ void CompilerMSL::emit_custom_functions()
|
||||
statement("template<>");
|
||||
statement("inline bool spvSubgroupShuffle(bool value, ushort lane)");
|
||||
begin_scope();
|
||||
if (msl_options.is_ios())
|
||||
if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
|
||||
statement("return !!quad_shuffle((ushort)value, lane);");
|
||||
else
|
||||
statement("return !!simd_shuffle((ushort)value, lane);");
|
||||
@ -4958,7 +5085,7 @@ void CompilerMSL::emit_custom_functions()
|
||||
statement("template<uint N>");
|
||||
statement("inline vec<bool, N> spvSubgroupShuffle(vec<bool, N> value, ushort lane)");
|
||||
begin_scope();
|
||||
if (msl_options.is_ios())
|
||||
if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
|
||||
statement("return (vec<bool, N>)quad_shuffle((vec<ushort, N>)value, lane);");
|
||||
else
|
||||
statement("return (vec<bool, N>)simd_shuffle((vec<ushort, N>)value, lane);");
|
||||
@ -4970,7 +5097,7 @@ void CompilerMSL::emit_custom_functions()
|
||||
statement("template<typename T>");
|
||||
statement("inline T spvSubgroupShuffleXor(T value, ushort mask)");
|
||||
begin_scope();
|
||||
if (msl_options.is_ios())
|
||||
if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
|
||||
statement("return quad_shuffle_xor(value, mask);");
|
||||
else
|
||||
statement("return simd_shuffle_xor(value, mask);");
|
||||
@ -4979,7 +5106,7 @@ void CompilerMSL::emit_custom_functions()
|
||||
statement("template<>");
|
||||
statement("inline bool spvSubgroupShuffleXor(bool value, ushort mask)");
|
||||
begin_scope();
|
||||
if (msl_options.is_ios())
|
||||
if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
|
||||
statement("return !!quad_shuffle_xor((ushort)value, mask);");
|
||||
else
|
||||
statement("return !!simd_shuffle_xor((ushort)value, mask);");
|
||||
@ -4988,7 +5115,7 @@ void CompilerMSL::emit_custom_functions()
|
||||
statement("template<uint N>");
|
||||
statement("inline vec<bool, N> spvSubgroupShuffleXor(vec<bool, N> value, ushort mask)");
|
||||
begin_scope();
|
||||
if (msl_options.is_ios())
|
||||
if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
|
||||
statement("return (vec<bool, N>)quad_shuffle_xor((vec<ushort, N>)value, mask);");
|
||||
else
|
||||
statement("return (vec<bool, N>)simd_shuffle_xor((vec<ushort, N>)value, mask);");
|
||||
@ -5000,7 +5127,7 @@ void CompilerMSL::emit_custom_functions()
|
||||
statement("template<typename T>");
|
||||
statement("inline T spvSubgroupShuffleUp(T value, ushort delta)");
|
||||
begin_scope();
|
||||
if (msl_options.is_ios())
|
||||
if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
|
||||
statement("return quad_shuffle_up(value, delta);");
|
||||
else
|
||||
statement("return simd_shuffle_up(value, delta);");
|
||||
@ -5009,7 +5136,7 @@ void CompilerMSL::emit_custom_functions()
|
||||
statement("template<>");
|
||||
statement("inline bool spvSubgroupShuffleUp(bool value, ushort delta)");
|
||||
begin_scope();
|
||||
if (msl_options.is_ios())
|
||||
if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
|
||||
statement("return !!quad_shuffle_up((ushort)value, delta);");
|
||||
else
|
||||
statement("return !!simd_shuffle_up((ushort)value, delta);");
|
||||
@ -5018,7 +5145,7 @@ void CompilerMSL::emit_custom_functions()
|
||||
statement("template<uint N>");
|
||||
statement("inline vec<bool, N> spvSubgroupShuffleUp(vec<bool, N> value, ushort delta)");
|
||||
begin_scope();
|
||||
if (msl_options.is_ios())
|
||||
if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
|
||||
statement("return (vec<bool, N>)quad_shuffle_up((vec<ushort, N>)value, delta);");
|
||||
else
|
||||
statement("return (vec<bool, N>)simd_shuffle_up((vec<ushort, N>)value, delta);");
|
||||
@ -5030,7 +5157,7 @@ void CompilerMSL::emit_custom_functions()
|
||||
statement("template<typename T>");
|
||||
statement("inline T spvSubgroupShuffleDown(T value, ushort delta)");
|
||||
begin_scope();
|
||||
if (msl_options.is_ios())
|
||||
if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
|
||||
statement("return quad_shuffle_down(value, delta);");
|
||||
else
|
||||
statement("return simd_shuffle_down(value, delta);");
|
||||
@ -5039,7 +5166,7 @@ void CompilerMSL::emit_custom_functions()
|
||||
statement("template<>");
|
||||
statement("inline bool spvSubgroupShuffleDown(bool value, ushort delta)");
|
||||
begin_scope();
|
||||
if (msl_options.is_ios())
|
||||
if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
|
||||
statement("return !!quad_shuffle_down((ushort)value, delta);");
|
||||
else
|
||||
statement("return !!simd_shuffle_down((ushort)value, delta);");
|
||||
@ -5048,7 +5175,7 @@ void CompilerMSL::emit_custom_functions()
|
||||
statement("template<uint N>");
|
||||
statement("inline vec<bool, N> spvSubgroupShuffleDown(vec<bool, N> value, ushort delta)");
|
||||
begin_scope();
|
||||
if (msl_options.is_ios())
|
||||
if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
|
||||
statement("return (vec<bool, N>)quad_shuffle_down((vec<ushort, N>)value, delta);");
|
||||
else
|
||||
statement("return (vec<bool, N>)simd_shuffle_down((vec<ushort, N>)value, delta);");
|
||||
@ -7677,6 +7804,10 @@ void CompilerMSL::emit_barrier(uint32_t id_exe_scope, uint32_t id_mem_scope, uin
|
||||
// Use the wider of the two scopes (smaller value)
|
||||
exe_scope = min(exe_scope, mem_scope);
|
||||
|
||||
if (msl_options.emulate_subgroups && exe_scope >= ScopeSubgroup && !id_mem_sem)
|
||||
// In this case, we assume a "subgroup" size of 1. The barrier, then, is a noop.
|
||||
return;
|
||||
|
||||
string bar_stmt;
|
||||
if ((msl_options.is_ios() && msl_options.supports_msl_version(1, 2)) || msl_options.supports_msl_version(2))
|
||||
bar_stmt = exe_scope < ScopeSubgroup ? "threadgroup_barrier" : "simdgroup_barrier";
|
||||
@ -9852,9 +9983,11 @@ string CompilerMSL::member_attribute_qualifier(const SPIRType &type, uint32_t in
|
||||
case BuiltInPrimitiveId:
|
||||
if (msl_options.multi_patch_workgroup)
|
||||
return "";
|
||||
/* fallthrough */
|
||||
return string(" [[") + builtin_qualifier(builtin) + "]]" + (mbr_type.array.empty() ? "" : " ");
|
||||
case BuiltInSubgroupLocalInvocationId: // FIXME: Should work in any stage
|
||||
case BuiltInSubgroupSize: // FIXME: Should work in any stage
|
||||
if (msl_options.emulate_subgroups)
|
||||
return "";
|
||||
return string(" [[") + builtin_qualifier(builtin) + "]]" + (mbr_type.array.empty() ? "" : " ");
|
||||
case BuiltInPatchVertices:
|
||||
return "";
|
||||
@ -10062,15 +10195,18 @@ string CompilerMSL::member_attribute_qualifier(const SPIRType &type, uint32_t in
|
||||
{
|
||||
switch (builtin)
|
||||
{
|
||||
case BuiltInNumSubgroups:
|
||||
case BuiltInSubgroupId:
|
||||
case BuiltInSubgroupLocalInvocationId: // FIXME: Should work in any stage
|
||||
case BuiltInSubgroupSize: // FIXME: Should work in any stage
|
||||
if (msl_options.emulate_subgroups)
|
||||
break;
|
||||
/* fallthrough */
|
||||
case BuiltInGlobalInvocationId:
|
||||
case BuiltInWorkgroupId:
|
||||
case BuiltInNumWorkgroups:
|
||||
case BuiltInLocalInvocationId:
|
||||
case BuiltInLocalInvocationIndex:
|
||||
case BuiltInNumSubgroups:
|
||||
case BuiltInSubgroupId:
|
||||
case BuiltInSubgroupLocalInvocationId: // FIXME: Should work in any stage
|
||||
case BuiltInSubgroupSize: // FIXME: Should work in any stage
|
||||
return string(" [[") + builtin_qualifier(builtin) + "]]";
|
||||
|
||||
default:
|
||||
@ -10336,6 +10472,10 @@ bool CompilerMSL::is_direct_input_builtin(BuiltIn bi_type)
|
||||
case BuiltInViewIndex:
|
||||
return get_execution_model() == ExecutionModelFragment && msl_options.multiview &&
|
||||
msl_options.multiview_layered_rendering;
|
||||
// Compute function in
|
||||
case BuiltInSubgroupId:
|
||||
case BuiltInNumSubgroups:
|
||||
return !msl_options.emulate_subgroups;
|
||||
// Any stage function in
|
||||
case BuiltInDeviceIndex:
|
||||
case BuiltInSubgroupEqMask:
|
||||
@ -10344,10 +10484,12 @@ bool CompilerMSL::is_direct_input_builtin(BuiltIn bi_type)
|
||||
case BuiltInSubgroupLeMask:
|
||||
case BuiltInSubgroupLtMask:
|
||||
return false;
|
||||
case BuiltInSubgroupLocalInvocationId:
|
||||
case BuiltInSubgroupSize:
|
||||
return get_execution_model() == ExecutionModelGLCompute ||
|
||||
(get_execution_model() == ExecutionModelFragment && msl_options.supports_msl_version(2, 2));
|
||||
if (msl_options.fixed_subgroup_size != 0)
|
||||
return false;
|
||||
/* fallthrough */
|
||||
case BuiltInSubgroupLocalInvocationId:
|
||||
return !msl_options.emulate_subgroups;
|
||||
default:
|
||||
return true;
|
||||
}
|
||||
@ -10958,67 +11100,78 @@ void CompilerMSL::fix_up_shader_inputs_outputs()
|
||||
entry_func.fixup_hooks_in.push_back([=]() { statement(tc, ".y = 1.0 - ", tc, ".y;"); });
|
||||
}
|
||||
break;
|
||||
case BuiltInSubgroupLocalInvocationId:
|
||||
// This is natively supported in compute shaders.
|
||||
if (get_execution_model() == ExecutionModelGLCompute)
|
||||
case BuiltInSubgroupId:
|
||||
if (!msl_options.emulate_subgroups)
|
||||
break;
|
||||
|
||||
// This is natively supported in fragment shaders in MSL 2.2.
|
||||
if (get_execution_model() == ExecutionModelFragment && msl_options.supports_msl_version(2, 2))
|
||||
break;
|
||||
|
||||
if (msl_options.is_ios())
|
||||
SPIRV_CROSS_THROW(
|
||||
"SubgroupLocalInvocationId cannot be used outside of compute shaders before MSL 2.2 on iOS.");
|
||||
|
||||
if (!msl_options.supports_msl_version(2, 1))
|
||||
SPIRV_CROSS_THROW(
|
||||
"SubgroupLocalInvocationId cannot be used outside of compute shaders before MSL 2.1.");
|
||||
|
||||
// Shaders other than compute shaders don't support the SIMD-group
|
||||
// builtins directly, but we can emulate them using the SIMD-group
|
||||
// functions. This might break if some of the subgroup terminated
|
||||
// before reaching the entry point.
|
||||
// For subgroup emulation, this is the same as the local invocation index.
|
||||
entry_func.fixup_hooks_in.push_back([=]() {
|
||||
statement(builtin_type_decl(bi_type), " ", to_expression(var_id),
|
||||
" = simd_prefix_exclusive_sum(1);");
|
||||
statement(builtin_type_decl(bi_type), " ", to_expression(var_id), " = ",
|
||||
to_expression(builtin_local_invocation_index_id), ";");
|
||||
});
|
||||
break;
|
||||
case BuiltInSubgroupSize:
|
||||
// This is natively supported in compute shaders.
|
||||
if (get_execution_model() == ExecutionModelGLCompute)
|
||||
case BuiltInNumSubgroups:
|
||||
if (!msl_options.emulate_subgroups)
|
||||
break;
|
||||
|
||||
// This is natively supported in fragment shaders in MSL 2.2.
|
||||
if (get_execution_model() == ExecutionModelFragment && msl_options.supports_msl_version(2, 2))
|
||||
// For subgroup emulation, this is the same as the workgroup size.
|
||||
entry_func.fixup_hooks_in.push_back([=]() {
|
||||
auto &type = expression_type(builtin_workgroup_size_id);
|
||||
string size_expr = to_expression(builtin_workgroup_size_id);
|
||||
if (type.vecsize >= 3)
|
||||
size_expr = join(size_expr, ".x * ", size_expr, ".y * ", size_expr, ".z");
|
||||
else if (type.vecsize == 2)
|
||||
size_expr = join(size_expr, ".x * ", size_expr, ".y");
|
||||
statement(builtin_type_decl(bi_type), " ", to_expression(var_id), " = ", size_expr, ";");
|
||||
});
|
||||
break;
|
||||
case BuiltInSubgroupLocalInvocationId:
|
||||
if (!msl_options.emulate_subgroups)
|
||||
break;
|
||||
|
||||
if (msl_options.is_ios())
|
||||
SPIRV_CROSS_THROW("SubgroupSize cannot be used outside of compute shaders on iOS.");
|
||||
|
||||
if (!msl_options.supports_msl_version(2, 1))
|
||||
SPIRV_CROSS_THROW("SubgroupSize cannot be used outside of compute shaders before Metal 2.1.");
|
||||
|
||||
// For subgroup emulation, assume subgroups of size 1.
|
||||
entry_func.fixup_hooks_in.push_back(
|
||||
[=]() { statement(builtin_type_decl(bi_type), " ", to_expression(var_id), " = simd_sum(1);"); });
|
||||
[=]() { statement(builtin_type_decl(bi_type), " ", to_expression(var_id), " = 0;"); });
|
||||
break;
|
||||
case BuiltInSubgroupSize:
|
||||
if (msl_options.emulate_subgroups)
|
||||
{
|
||||
// For subgroup emulation, assume subgroups of size 1.
|
||||
entry_func.fixup_hooks_in.push_back(
|
||||
[=]() { statement(builtin_type_decl(bi_type), " ", to_expression(var_id), " = 1;"); });
|
||||
}
|
||||
else if (msl_options.fixed_subgroup_size != 0)
|
||||
{
|
||||
entry_func.fixup_hooks_in.push_back([=]() {
|
||||
statement(builtin_type_decl(bi_type), " ", to_expression(var_id), " = ",
|
||||
msl_options.fixed_subgroup_size, ";");
|
||||
});
|
||||
}
|
||||
break;
|
||||
case BuiltInSubgroupEqMask:
|
||||
if (msl_options.is_ios())
|
||||
SPIRV_CROSS_THROW("Subgroup ballot functionality is unavailable on iOS.");
|
||||
if (msl_options.is_ios() && !msl_options.supports_msl_version(2, 2))
|
||||
SPIRV_CROSS_THROW("Subgroup ballot functionality requires Metal 2.2 on iOS.");
|
||||
if (!msl_options.supports_msl_version(2, 1))
|
||||
SPIRV_CROSS_THROW("Subgroup ballot functionality requires Metal 2.1.");
|
||||
entry_func.fixup_hooks_in.push_back([=]() {
|
||||
statement(builtin_type_decl(bi_type), " ", to_expression(var_id), " = ",
|
||||
to_expression(builtin_subgroup_invocation_id_id), " >= 32 ? uint4(0, (1 << (",
|
||||
to_expression(builtin_subgroup_invocation_id_id), " - 32)), uint2(0)) : uint4(1 << ",
|
||||
to_expression(builtin_subgroup_invocation_id_id), ", uint3(0));");
|
||||
if (msl_options.is_ios())
|
||||
{
|
||||
statement(builtin_type_decl(bi_type), " ", to_expression(var_id), " = ", "uint4(1 << ",
|
||||
to_expression(builtin_subgroup_invocation_id_id), ", uint3(0));");
|
||||
}
|
||||
else
|
||||
{
|
||||
statement(builtin_type_decl(bi_type), " ", to_expression(var_id), " = ",
|
||||
to_expression(builtin_subgroup_invocation_id_id), " >= 32 ? uint4(0, (1 << (",
|
||||
to_expression(builtin_subgroup_invocation_id_id), " - 32)), uint2(0)) : uint4(1 << ",
|
||||
to_expression(builtin_subgroup_invocation_id_id), ", uint3(0));");
|
||||
}
|
||||
});
|
||||
break;
|
||||
case BuiltInSubgroupGeMask:
|
||||
if (msl_options.is_ios())
|
||||
SPIRV_CROSS_THROW("Subgroup ballot functionality is unavailable on iOS.");
|
||||
if (msl_options.is_ios() && !msl_options.supports_msl_version(2, 2))
|
||||
SPIRV_CROSS_THROW("Subgroup ballot functionality requires Metal 2.2 on iOS.");
|
||||
if (!msl_options.supports_msl_version(2, 1))
|
||||
SPIRV_CROSS_THROW("Subgroup ballot functionality requires Metal 2.1.");
|
||||
if (msl_options.fixed_subgroup_size != 0)
|
||||
add_spv_func_and_recompile(SPVFuncImplSubgroupBallot);
|
||||
entry_func.fixup_hooks_in.push_back([=]() {
|
||||
// Case where index < 32, size < 32:
|
||||
// mask0 = bfi(0, 0xFFFFFFFF, index, size - index);
|
||||
@ -11034,60 +11187,150 @@ void CompilerMSL::fix_up_shader_inputs_outputs()
|
||||
// This is further complicated by the fact that if you attempt
|
||||
// to bfi/bfe out-of-bounds on Metal, undefined behavior is the
|
||||
// result.
|
||||
statement(builtin_type_decl(bi_type), " ", to_expression(var_id),
|
||||
" = uint4(insert_bits(0u, 0xFFFFFFFF, min(",
|
||||
to_expression(builtin_subgroup_invocation_id_id), ", 32u), (uint)max(min((int)",
|
||||
to_expression(builtin_subgroup_size_id), ", 32) - (int)",
|
||||
to_expression(builtin_subgroup_invocation_id_id),
|
||||
", 0)), insert_bits(0u, 0xFFFFFFFF, (uint)max((int)",
|
||||
to_expression(builtin_subgroup_invocation_id_id), " - 32, 0), (uint)max((int)",
|
||||
to_expression(builtin_subgroup_size_id), " - (int)max(",
|
||||
to_expression(builtin_subgroup_invocation_id_id), ", 32u), 0)), uint2(0));");
|
||||
if (msl_options.fixed_subgroup_size > 32)
|
||||
{
|
||||
// Don't use the subgroup size variable with fixed subgroup sizes,
|
||||
// since the variables could be defined in the wrong order.
|
||||
statement(builtin_type_decl(bi_type), " ", to_expression(var_id),
|
||||
" = uint4(insert_bits(0u, 0xFFFFFFFF, min(",
|
||||
to_expression(builtin_subgroup_invocation_id_id), ", 32u), (uint)max(32 - (int)",
|
||||
to_expression(builtin_subgroup_invocation_id_id),
|
||||
", 0)), insert_bits(0u, 0xFFFFFFFF,"
|
||||
" (uint)max((int)",
|
||||
to_expression(builtin_subgroup_invocation_id_id), " - 32, 0), ",
|
||||
msl_options.fixed_subgroup_size, " - max(",
|
||||
to_expression(builtin_subgroup_invocation_id_id),
|
||||
", 32u)), uint2(0)) & spvSubgroupBallot(true);");
|
||||
}
|
||||
else if (msl_options.fixed_subgroup_size != 0)
|
||||
{
|
||||
statement(builtin_type_decl(bi_type), " ", to_expression(var_id),
|
||||
" = uint4(insert_bits(0u, 0xFFFFFFFF, ",
|
||||
to_expression(builtin_subgroup_invocation_id_id), ", ",
|
||||
msl_options.fixed_subgroup_size, " - ",
|
||||
to_expression(builtin_subgroup_invocation_id_id),
|
||||
"), uint3(0)) & spvSubgroupBallot(true);");
|
||||
}
|
||||
else if (msl_options.is_ios())
|
||||
{
|
||||
// On iOS, the SIMD-group size will currently never exceed 32.
|
||||
statement(builtin_type_decl(bi_type), " ", to_expression(var_id),
|
||||
" = uint4(insert_bits(0u, 0xFFFFFFFF, ",
|
||||
to_expression(builtin_subgroup_invocation_id_id), ", ",
|
||||
to_expression(builtin_subgroup_size_id), " - ",
|
||||
to_expression(builtin_subgroup_invocation_id_id), "), uint3(0));");
|
||||
}
|
||||
else
|
||||
{
|
||||
statement(builtin_type_decl(bi_type), " ", to_expression(var_id),
|
||||
" = uint4(insert_bits(0u, 0xFFFFFFFF, min(",
|
||||
to_expression(builtin_subgroup_invocation_id_id), ", 32u), (uint)max(min((int)",
|
||||
to_expression(builtin_subgroup_size_id), ", 32) - (int)",
|
||||
to_expression(builtin_subgroup_invocation_id_id),
|
||||
", 0)), insert_bits(0u, 0xFFFFFFFF, (uint)max((int)",
|
||||
to_expression(builtin_subgroup_invocation_id_id), " - 32, 0), (uint)max((int)",
|
||||
to_expression(builtin_subgroup_size_id), " - (int)max(",
|
||||
to_expression(builtin_subgroup_invocation_id_id), ", 32u), 0)), uint2(0));");
|
||||
}
|
||||
});
|
||||
break;
|
||||
case BuiltInSubgroupGtMask:
|
||||
if (msl_options.is_ios())
|
||||
SPIRV_CROSS_THROW("Subgroup ballot functionality is unavailable on iOS.");
|
||||
if (msl_options.is_ios() && !msl_options.supports_msl_version(2, 2))
|
||||
SPIRV_CROSS_THROW("Subgroup ballot functionality requires Metal 2.2 on iOS.");
|
||||
if (!msl_options.supports_msl_version(2, 1))
|
||||
SPIRV_CROSS_THROW("Subgroup ballot functionality requires Metal 2.1.");
|
||||
add_spv_func_and_recompile(SPVFuncImplSubgroupBallot);
|
||||
entry_func.fixup_hooks_in.push_back([=]() {
|
||||
// The same logic applies here, except now the index is one
|
||||
// more than the subgroup invocation ID.
|
||||
statement(builtin_type_decl(bi_type), " ", to_expression(var_id),
|
||||
" = uint4(insert_bits(0u, 0xFFFFFFFF, min(",
|
||||
to_expression(builtin_subgroup_invocation_id_id), " + 1, 32u), (uint)max(min((int)",
|
||||
to_expression(builtin_subgroup_size_id), ", 32) - (int)",
|
||||
to_expression(builtin_subgroup_invocation_id_id),
|
||||
" - 1, 0)), insert_bits(0u, 0xFFFFFFFF, (uint)max((int)",
|
||||
to_expression(builtin_subgroup_invocation_id_id), " + 1 - 32, 0), (uint)max((int)",
|
||||
to_expression(builtin_subgroup_size_id), " - (int)max(",
|
||||
to_expression(builtin_subgroup_invocation_id_id), " + 1, 32u), 0)), uint2(0));");
|
||||
if (msl_options.fixed_subgroup_size > 32)
|
||||
{
|
||||
statement(builtin_type_decl(bi_type), " ", to_expression(var_id),
|
||||
" = uint4(insert_bits(0u, 0xFFFFFFFF, min(",
|
||||
to_expression(builtin_subgroup_invocation_id_id), " + 1, 32u), (uint)max(32 - (int)",
|
||||
to_expression(builtin_subgroup_invocation_id_id),
|
||||
" - 1, 0)), insert_bits(0u, 0xFFFFFFFF, (uint)max((int)",
|
||||
to_expression(builtin_subgroup_invocation_id_id), " + 1 - 32, 0), ",
|
||||
msl_options.fixed_subgroup_size, " - max(",
|
||||
to_expression(builtin_subgroup_invocation_id_id),
|
||||
" + 1, 32u)), uint2(0)) & "
|
||||
"spvSubgroupBallot(true);");
|
||||
}
|
||||
else if (msl_options.fixed_subgroup_size != 0)
|
||||
{
|
||||
statement(builtin_type_decl(bi_type), " ", to_expression(var_id),
|
||||
" = uint4(insert_bits(0u, 0xFFFFFFFF, ",
|
||||
to_expression(builtin_subgroup_invocation_id_id), " + 1, ",
|
||||
msl_options.fixed_subgroup_size, " - ",
|
||||
to_expression(builtin_subgroup_invocation_id_id),
|
||||
" - 1), uint3(0)) & spvSubgroupBallot(true);");
|
||||
}
|
||||
else if (msl_options.is_ios())
|
||||
{
|
||||
statement(builtin_type_decl(bi_type), " ", to_expression(var_id),
|
||||
" = uint4(insert_bits(0u, 0xFFFFFFFF, ",
|
||||
to_expression(builtin_subgroup_invocation_id_id), " + 1, ",
|
||||
to_expression(builtin_subgroup_size_id), " - ",
|
||||
to_expression(builtin_subgroup_invocation_id_id), " - 1), uint3(0));");
|
||||
}
|
||||
else
|
||||
{
|
||||
statement(builtin_type_decl(bi_type), " ", to_expression(var_id),
|
||||
" = uint4(insert_bits(0u, 0xFFFFFFFF, min(",
|
||||
to_expression(builtin_subgroup_invocation_id_id), " + 1, 32u), (uint)max(min((int)",
|
||||
to_expression(builtin_subgroup_size_id), ", 32) - (int)",
|
||||
to_expression(builtin_subgroup_invocation_id_id),
|
||||
" - 1, 0)), insert_bits(0u, 0xFFFFFFFF, (uint)max((int)",
|
||||
to_expression(builtin_subgroup_invocation_id_id), " + 1 - 32, 0), (uint)max((int)",
|
||||
to_expression(builtin_subgroup_size_id), " - (int)max(",
|
||||
to_expression(builtin_subgroup_invocation_id_id), " + 1, 32u), 0)), uint2(0));");
|
||||
}
|
||||
});
|
||||
break;
|
||||
case BuiltInSubgroupLeMask:
|
||||
if (msl_options.is_ios())
|
||||
SPIRV_CROSS_THROW("Subgroup ballot functionality is unavailable on iOS.");
|
||||
if (msl_options.is_ios() && !msl_options.supports_msl_version(2, 2))
|
||||
SPIRV_CROSS_THROW("Subgroup ballot functionality requires Metal 2.2 on iOS.");
|
||||
if (!msl_options.supports_msl_version(2, 1))
|
||||
SPIRV_CROSS_THROW("Subgroup ballot functionality requires Metal 2.1.");
|
||||
add_spv_func_and_recompile(SPVFuncImplSubgroupBallot);
|
||||
entry_func.fixup_hooks_in.push_back([=]() {
|
||||
statement(builtin_type_decl(bi_type), " ", to_expression(var_id),
|
||||
" = uint4(extract_bits(0xFFFFFFFF, 0, min(",
|
||||
to_expression(builtin_subgroup_invocation_id_id),
|
||||
" + 1, 32u)), extract_bits(0xFFFFFFFF, 0, (uint)max((int)",
|
||||
to_expression(builtin_subgroup_invocation_id_id), " + 1 - 32, 0)), uint2(0));");
|
||||
if (msl_options.is_ios())
|
||||
{
|
||||
statement(builtin_type_decl(bi_type), " ", to_expression(var_id),
|
||||
" = uint4(extract_bits(0xFFFFFFFF, 0, ",
|
||||
to_expression(builtin_subgroup_invocation_id_id), " + 1), uint3(0));");
|
||||
}
|
||||
else
|
||||
{
|
||||
statement(builtin_type_decl(bi_type), " ", to_expression(var_id),
|
||||
" = uint4(extract_bits(0xFFFFFFFF, 0, min(",
|
||||
to_expression(builtin_subgroup_invocation_id_id),
|
||||
" + 1, 32u)), extract_bits(0xFFFFFFFF, 0, (uint)max((int)",
|
||||
to_expression(builtin_subgroup_invocation_id_id), " + 1 - 32, 0)), uint2(0));");
|
||||
}
|
||||
});
|
||||
break;
|
||||
case BuiltInSubgroupLtMask:
|
||||
if (msl_options.is_ios())
|
||||
SPIRV_CROSS_THROW("Subgroup ballot functionality is unavailable on iOS.");
|
||||
if (msl_options.is_ios() && !msl_options.supports_msl_version(2, 2))
|
||||
SPIRV_CROSS_THROW("Subgroup ballot functionality requires Metal 2.2 on iOS.");
|
||||
if (!msl_options.supports_msl_version(2, 1))
|
||||
SPIRV_CROSS_THROW("Subgroup ballot functionality requires Metal 2.1.");
|
||||
add_spv_func_and_recompile(SPVFuncImplSubgroupBallot);
|
||||
entry_func.fixup_hooks_in.push_back([=]() {
|
||||
statement(builtin_type_decl(bi_type), " ", to_expression(var_id),
|
||||
" = uint4(extract_bits(0xFFFFFFFF, 0, min(",
|
||||
to_expression(builtin_subgroup_invocation_id_id),
|
||||
", 32u)), extract_bits(0xFFFFFFFF, 0, (uint)max((int)",
|
||||
to_expression(builtin_subgroup_invocation_id_id), " - 32, 0)), uint2(0));");
|
||||
if (msl_options.is_ios())
|
||||
{
|
||||
statement(builtin_type_decl(bi_type), " ", to_expression(var_id),
|
||||
" = uint4(extract_bits(0xFFFFFFFF, 0, ",
|
||||
to_expression(builtin_subgroup_invocation_id_id), "), uint3(0));");
|
||||
}
|
||||
else
|
||||
{
|
||||
statement(builtin_type_decl(bi_type), " ", to_expression(var_id),
|
||||
" = uint4(extract_bits(0xFFFFFFFF, 0, min(",
|
||||
to_expression(builtin_subgroup_invocation_id_id),
|
||||
", 32u)), extract_bits(0xFFFFFFFF, 0, (uint)max((int)",
|
||||
to_expression(builtin_subgroup_invocation_id_id), " - 32, 0)), uint2(0));");
|
||||
}
|
||||
});
|
||||
break;
|
||||
case BuiltInViewIndex:
|
||||
@ -12376,10 +12619,23 @@ void CompilerMSL::emit_subgroup_op(const Instruction &i)
|
||||
const uint32_t *ops = stream(i);
|
||||
auto op = static_cast<Op>(i.op);
|
||||
|
||||
// Metal 2.0 is required. iOS only supports quad ops. macOS only supports
|
||||
// broadcast and shuffle on 10.13 (2.0), with full support in 10.14 (2.1).
|
||||
// Note that iOS makes no distinction between a quad-group and a subgroup;
|
||||
// all subgroups are quad-groups there.
|
||||
if (msl_options.emulate_subgroups)
|
||||
{
|
||||
// In this mode, only the GroupNonUniform cap is supported. The only op
|
||||
// we need to handle, then, is OpGroupNonUniformElect.
|
||||
if (op != OpGroupNonUniformElect)
|
||||
SPIRV_CROSS_THROW("Subgroup emulation does not support operations other than Elect.");
|
||||
// In this mode, the subgroup size is assumed to be one, so every invocation
|
||||
// is elected.
|
||||
emit_op(ops[0], ops[1], "true", true);
|
||||
return;
|
||||
}
|
||||
|
||||
// Metal 2.0 is required. iOS only supports quad ops on 11.0 (2.0), with
|
||||
// full support in 13.0 (2.2). macOS only supports broadcast and shuffle on
|
||||
// 10.13 (2.0), with full support in 10.14 (2.1).
|
||||
// Note that Apple GPUs before A13 make no distinction between a quad-group
|
||||
// and a SIMD-group; all SIMD-groups are quad-groups on those.
|
||||
if (!msl_options.supports_msl_version(2))
|
||||
SPIRV_CROSS_THROW("Subgroups are only supported in Metal 2.0 and up.");
|
||||
|
||||
@ -12388,12 +12644,32 @@ void CompilerMSL::emit_subgroup_op(const Instruction &i)
|
||||
auto int_type = to_signed_basetype(integer_width);
|
||||
auto uint_type = to_unsigned_basetype(integer_width);
|
||||
|
||||
if (msl_options.is_ios())
|
||||
if (msl_options.is_ios() && (!msl_options.supports_msl_version(2, 3) || !msl_options.ios_use_simdgroup_functions))
|
||||
{
|
||||
switch (op)
|
||||
{
|
||||
default:
|
||||
SPIRV_CROSS_THROW("iOS only supports quad-group operations.");
|
||||
SPIRV_CROSS_THROW("Subgroup ops beyond broadcast, ballot, and shuffle on iOS require Metal 2.3 and up.");
|
||||
case OpGroupNonUniformBroadcastFirst:
|
||||
if (!msl_options.supports_msl_version(2, 2))
|
||||
SPIRV_CROSS_THROW("BroadcastFirst on iOS requires Metal 2.2 and up.");
|
||||
break;
|
||||
case OpGroupNonUniformElect:
|
||||
if (!msl_options.supports_msl_version(2, 2))
|
||||
SPIRV_CROSS_THROW("Elect on iOS requires Metal 2.2 and up.");
|
||||
break;
|
||||
case OpGroupNonUniformAny:
|
||||
case OpGroupNonUniformAll:
|
||||
case OpGroupNonUniformAllEqual:
|
||||
case OpGroupNonUniformBallot:
|
||||
case OpGroupNonUniformInverseBallot:
|
||||
case OpGroupNonUniformBallotBitExtract:
|
||||
case OpGroupNonUniformBallotFindLSB:
|
||||
case OpGroupNonUniformBallotFindMSB:
|
||||
case OpGroupNonUniformBallotBitCount:
|
||||
if (!msl_options.supports_msl_version(2, 2))
|
||||
SPIRV_CROSS_THROW("Ballot ops on iOS requires Metal 2.2 and up.");
|
||||
break;
|
||||
case OpGroupNonUniformBroadcast:
|
||||
case OpGroupNonUniformShuffle:
|
||||
case OpGroupNonUniformShuffleXor:
|
||||
@ -12430,7 +12706,10 @@ void CompilerMSL::emit_subgroup_op(const Instruction &i)
|
||||
switch (op)
|
||||
{
|
||||
case OpGroupNonUniformElect:
|
||||
emit_op(result_type, id, "simd_is_first()", true);
|
||||
if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
|
||||
emit_op(result_type, id, "quad_is_first()", false);
|
||||
else
|
||||
emit_op(result_type, id, "simd_is_first()", false);
|
||||
break;
|
||||
|
||||
case OpGroupNonUniformBroadcast:
|
||||
@ -12501,11 +12780,17 @@ void CompilerMSL::emit_subgroup_op(const Instruction &i)
|
||||
break;
|
||||
|
||||
case OpGroupNonUniformAll:
|
||||
emit_unary_func_op(result_type, id, ops[3], "simd_all");
|
||||
if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
|
||||
emit_unary_func_op(result_type, id, ops[3], "quad_all");
|
||||
else
|
||||
emit_unary_func_op(result_type, id, ops[3], "simd_all");
|
||||
break;
|
||||
|
||||
case OpGroupNonUniformAny:
|
||||
emit_unary_func_op(result_type, id, ops[3], "simd_any");
|
||||
if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
|
||||
emit_unary_func_op(result_type, id, ops[3], "quad_any");
|
||||
else
|
||||
emit_unary_func_op(result_type, id, ops[3], "simd_any");
|
||||
break;
|
||||
|
||||
case OpGroupNonUniformAllEqual:
|
||||
@ -12969,6 +13254,9 @@ string CompilerMSL::builtin_qualifier(BuiltIn builtin)
|
||||
return "thread_index_in_threadgroup";
|
||||
|
||||
case BuiltInSubgroupSize:
|
||||
if (msl_options.emulate_subgroups || msl_options.fixed_subgroup_size != 0)
|
||||
// Shouldn't be reached.
|
||||
SPIRV_CROSS_THROW("Emitting threads_per_simdgroup attribute with fixed subgroup size??");
|
||||
if (execution.model == ExecutionModelFragment)
|
||||
{
|
||||
if (!msl_options.supports_msl_version(2, 2))
|
||||
@ -12983,16 +13271,25 @@ string CompilerMSL::builtin_qualifier(BuiltIn builtin)
|
||||
}
|
||||
|
||||
case BuiltInNumSubgroups:
|
||||
if (msl_options.emulate_subgroups)
|
||||
// Shouldn't be reached.
|
||||
SPIRV_CROSS_THROW("NumSubgroups is handled specially with emulation.");
|
||||
if (!msl_options.supports_msl_version(2))
|
||||
SPIRV_CROSS_THROW("Subgroup builtins require Metal 2.0.");
|
||||
return msl_options.is_ios() ? "quadgroups_per_threadgroup" : "simdgroups_per_threadgroup";
|
||||
|
||||
case BuiltInSubgroupId:
|
||||
if (msl_options.emulate_subgroups)
|
||||
// Shouldn't be reached.
|
||||
SPIRV_CROSS_THROW("SubgroupId is handled specially with emulation.");
|
||||
if (!msl_options.supports_msl_version(2))
|
||||
SPIRV_CROSS_THROW("Subgroup builtins require Metal 2.0.");
|
||||
return msl_options.is_ios() ? "quadgroup_index_in_threadgroup" : "simdgroup_index_in_threadgroup";
|
||||
|
||||
case BuiltInSubgroupLocalInvocationId:
|
||||
if (msl_options.emulate_subgroups)
|
||||
// Shouldn't be reached.
|
||||
SPIRV_CROSS_THROW("SubgroupLocalInvocationId is handled specially with emulation.");
|
||||
if (execution.model == ExecutionModelFragment)
|
||||
{
|
||||
if (!msl_options.supports_msl_version(2, 2))
|
||||
|
||||
@ -364,6 +364,28 @@ public:
|
||||
// and will be addressed using the current ViewIndex.
|
||||
bool arrayed_subpass_input = false;
|
||||
|
||||
// Whether to use SIMD-group or quadgroup functions to implement group nnon-uniform
|
||||
// operations. Some GPUs on iOS do not support the SIMD-group functions, only the
|
||||
// quadgroup functions.
|
||||
bool ios_use_simdgroup_functions = false;
|
||||
|
||||
// If set, the subgroup size will be assumed to be one, and subgroup-related
|
||||
// builtins and operations will be emitted accordingly. This mode is intended to
|
||||
// be used by MoltenVK on hardware/software configurations which do not provide
|
||||
// sufficient support for subgroups.
|
||||
bool emulate_subgroups = false;
|
||||
|
||||
// If nonzero, a fixed subgroup size to assume. Metal, similarly to VK_EXT_subgroup_size_control,
|
||||
// allows the SIMD-group size (aka thread execution width) to vary depending on
|
||||
// register usage and requirements. In certain circumstances--for example, a pipeline
|
||||
// in MoltenVK without VK_PIPELINE_SHADER_STAGE_CREATE_ALLOW_VARYING_SUBGROUP_SIZE_BIT_EXT--
|
||||
// this is undesirable. This fixes the value of the SubgroupSize builtin, instead of
|
||||
// mapping it to the Metal builtin [[thread_execution_width]]. If the thread
|
||||
// execution width is reduced, the extra invocations will appear to be inactive.
|
||||
// If zero, the SubgroupSize will be allowed to vary, and the builtin will be mapped
|
||||
// to the Metal [[thread_execution_width]] builtin.
|
||||
uint32_t fixed_subgroup_size = 0;
|
||||
|
||||
enum class IndexType
|
||||
{
|
||||
None = 0,
|
||||
@ -853,6 +875,8 @@ protected:
|
||||
uint32_t builtin_subgroup_size_id = 0;
|
||||
uint32_t builtin_dispatch_base_id = 0;
|
||||
uint32_t builtin_stage_input_size_id = 0;
|
||||
uint32_t builtin_local_invocation_index_id = 0;
|
||||
uint32_t builtin_workgroup_size_id = 0;
|
||||
uint32_t swizzle_buffer_id = 0;
|
||||
uint32_t buffer_size_buffer_id = 0;
|
||||
uint32_t view_mask_buffer_id = 0;
|
||||
|
||||
@ -308,6 +308,14 @@ def cross_compile_msl(shader, spirv, opt, iterations, paths):
|
||||
msl_args.append('--msl-arrayed-subpass-input')
|
||||
if '.1d-as-2d.' in shader:
|
||||
msl_args.append('--msl-texture-1d-as-2d')
|
||||
if '.simd.' in shader:
|
||||
msl_args.append('--msl-ios-use-simdgroup-functions')
|
||||
if '.emulate-subgroup.' in shader:
|
||||
msl_args.append('--msl-emulate-subgroups')
|
||||
if '.fixed-subgroup.' in shader:
|
||||
# Arbitrary for testing purposes.
|
||||
msl_args.append('--msl-fixed-subgroup-size')
|
||||
msl_args.append('32')
|
||||
|
||||
subprocess.check_call(msl_args)
|
||||
|
||||
|
||||
Loading…
Reference in New Issue
Block a user