AuroraMiddleware/SPIRV-Cross
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Merge pull request #1499 from cdavis5e/subgroup-fixes

Browse Source 
Subgroup fixes
This commit is contained in:
Hans-Kristian Arntzen 2020-10-26 09:47:39 +01:00 committed by GitHub
commit dd35821f2f
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
8 changed files with 877 additions and 108 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

216
reference/shaders-msl-no-opt/comp/subgroups.nocompat.invalid.vk.msl21.comp
View File

@ -12,6 +12,42 @@ struct SSBO
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);
@ -27,37 +63,47 @@ inline bool spvSubgroupBallotBitExtract(uint4 ballot, uint bit)
return !!extract_bits(ballot[bit / 32], bit % 32, 1);
}
inline uint spvSubgroupBallotFindLSB(uint4 ballot)
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)
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 spvSubgroupBallotBitCount(uint4 ballot)
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 spvSubgroupBallotBitCount(ballot & mask);
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 spvSubgroupBallotBitCount(ballot & mask);
return spvPopCount4(ballot & mask);
}
template<typename T>
inline bool spvSubgroupAllEqual(T value)
{
return simd_all(value == simd_broadcast_first(value));
return simd_all(all(value == simd_broadcast_first(value)));
}
template<>
@ -66,11 +112,125 @@ 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_SubgroupSize [[thread_execution_width]], uint gl_SubgroupInvocationID [[thread_index_in_simdgroup]])
{
uint4 gl_SubgroupEqMask = gl_SubgroupInvocationID > 32 ? uint4(0, (1 << (gl_SubgroupInvocationID - 32)), uint2(0)) : uint4(1 << gl_SubgroupInvocationID, uint3(0));
uint4 gl_SubgroupGeMask = uint4(extract_bits(0xFFFFFFFF, min(gl_SubgroupInvocationID, 32u), (uint)max(min((int)gl_SubgroupSize, 32) - (int)gl_SubgroupInvocationID, 0)), extract_bits(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(extract_bits(0xFFFFFFFF, min(gl_SubgroupInvocationID + 1, 32u), (uint)max(min((int)gl_SubgroupSize, 32) - (int)gl_SubgroupInvocationID - 1, 0)), extract_bits(0xFFFFFFFF, (uint)max((int)gl_SubgroupInvocationID + 1 - 32, 0), (uint)max((int)gl_SubgroupSize - (int)max(gl_SubgroupInvocationID + 1, 32u), 0)), uint2(0));
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));
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);
@ -88,24 +248,32 @@ kernel void main0(device SSBO& _9 [[buffer(0)]], uint gl_NumSubgroups [[simdgrou
_9.FragColor = float4(gl_SubgroupGtMask).x;
_9.FragColor = float4(gl_SubgroupLeMask).x;
_9.FragColor = float4(gl_SubgroupLtMask).x;
float4 broadcasted = simd_broadcast(float4(10.0), 8u);
float3 first = simd_broadcast_first(float3(20.0));
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);
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);
uint msb = spvSubgroupBallotFindMSB(ballot_value);
uint shuffled = simd_shuffle(10u, 8u);
uint shuffled_xor = simd_shuffle_xor(30u, 8u);
uint shuffled_up = simd_shuffle_up(20u, 4u);
uint shuffled_down = simd_shuffle_down(20u, 4u);
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));
@ -140,9 +308,13 @@ kernel void main0(device SSBO& _9 [[buffer(0)]], uint gl_NumSubgroups [[simdgrou
anded = quad_and(anded);
ored = quad_or(ored);
xored = quad_xor(xored);
float4 swap_horiz = quad_shuffle_xor(float4(20.0), 1u);
float4 swap_vertical = quad_shuffle_xor(float4(20.0), 2u);
float4 swap_diagonal = quad_shuffle_xor(float4(20.0), 3u);
float4 quad_broadcast0 = quad_broadcast(float4(20.0), 3u);
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);
}

134
reference/shaders-msl-no-opt/comp/subgroups.nocompat.invalid.vk.msl21.ios.comp
View File

@ -1,3 +1,5 @@
#pragma clang diagnostic ignored "-Wmissing-prototypes"
#include <metal_stdlib>
#include <simd/simd.h>
@ -10,6 +12,114 @@ struct SSBO
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
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]])
{
_9.FragColor = float(gl_NumSubgroups);
@ -21,13 +131,21 @@ kernel void main0(device SSBO& _9 [[buffer(0)]], uint gl_NumSubgroups [[quadgrou
simdgroup_barrier(mem_flags::mem_device);
simdgroup_barrier(mem_flags::mem_threadgroup);
simdgroup_barrier(mem_flags::mem_texture);
uint shuffled = quad_shuffle(10u, 8u);
uint shuffled_xor = quad_shuffle_xor(30u, 8u);
uint shuffled_up = quad_shuffle_up(20u, 4u);
uint shuffled_down = quad_shuffle_down(20u, 4u);
float4 swap_horiz = quad_shuffle_xor(float4(20.0), 1u);
float4 swap_vertical = quad_shuffle_xor(float4(20.0), 2u);
float4 swap_diagonal = quad_shuffle_xor(float4(20.0), 3u);
float4 quad_broadcast0 = quad_broadcast(float4(20.0), 3u);
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);
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);
}

216
reference/shaders-msl-no-opt/frag/subgroups.nocompat.invalid.vk.msl21.frag
View File

@ -10,6 +10,42 @@ struct main0_out
float FragColor [[color(0)]];
};
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);
@ -25,37 +61,47 @@ inline bool spvSubgroupBallotBitExtract(uint4 ballot, uint bit)
return !!extract_bits(ballot[bit / 32], bit % 32, 1);
}
inline uint spvSubgroupBallotFindLSB(uint4 ballot)
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)
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 spvSubgroupBallotBitCount(uint4 ballot)
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 spvSubgroupBallotBitCount(ballot & mask);
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 spvSubgroupBallotBitCount(ballot & mask);
return spvPopCount4(ballot & mask);
}
template<typename T>
inline bool spvSubgroupAllEqual(T value)
{
return simd_all(value == simd_broadcast_first(value));
return simd_all(all(value == simd_broadcast_first(value)));
}
template<>
@ -64,14 +110,128 @@ 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);
}
fragment main0_out main0()
{
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(extract_bits(0xFFFFFFFF, min(gl_SubgroupInvocationID, 32u), (uint)max(min((int)gl_SubgroupSize, 32) - (int)gl_SubgroupInvocationID, 0)), extract_bits(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(extract_bits(0xFFFFFFFF, min(gl_SubgroupInvocationID + 1, 32u), (uint)max(min((int)gl_SubgroupSize, 32) - (int)gl_SubgroupInvocationID - 1, 0)), extract_bits(0xFFFFFFFF, (uint)max((int)gl_SubgroupInvocationID + 1 - 32, 0), (uint)max((int)gl_SubgroupSize - (int)max(gl_SubgroupInvocationID + 1, 32u), 0)), uint2(0));
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));
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));
out.FragColor = float(gl_SubgroupSize);
@ -82,24 +242,32 @@ fragment main0_out main0()
out.FragColor = float4(gl_SubgroupGtMask).x;
out.FragColor = float4(gl_SubgroupLeMask).x;
out.FragColor = float4(gl_SubgroupLtMask).x;
float4 broadcasted = simd_broadcast(float4(10.0), 8u);
float3 first = simd_broadcast_first(float3(20.0));
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);
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);
uint msb = spvSubgroupBallotFindMSB(ballot_value);
uint shuffled = simd_shuffle(10u, 8u);
uint shuffled_xor = simd_shuffle_xor(30u, 8u);
uint shuffled_up = simd_shuffle_up(20u, 4u);
uint shuffled_down = simd_shuffle_down(20u, 4u);
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));
@ -134,10 +302,14 @@ fragment main0_out main0()
anded = quad_and(anded);
ored = quad_or(ored);
xored = quad_xor(xored);
float4 swap_horiz = quad_shuffle_xor(float4(20.0), 1u);
float4 swap_vertical = quad_shuffle_xor(float4(20.0), 2u);
float4 swap_diagonal = quad_shuffle_xor(float4(20.0), 3u);
float4 quad_broadcast0 = quad_broadcast(float4(20.0), 3u);
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);
return out;
}

12
shaders-msl-no-opt/comp/subgroups.nocompat.invalid.vk.msl21.comp
View File

@ -35,7 +35,9 @@ void main()
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);
@ -47,17 +49,23 @@ void main()
// 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));
@ -120,7 +128,11 @@ void main()
// 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);
}

8
shaders-msl-no-opt/comp/subgroups.nocompat.invalid.vk.msl21.ios.comp
View File

@ -27,15 +27,23 @@ void main()
// 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);
// 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);
}

12
shaders-msl-no-opt/frag/subgroups.nocompat.invalid.vk.msl21.frag
View File

@ -28,7 +28,9 @@ void main()
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);
@ -40,17 +42,23 @@ void main()
// 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));
@ -113,7 +121,11 @@ void main()
// 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);
}

377
spirv_msl.cpp
View File

@ -160,7 +160,7 @@ void CompilerMSL::build_implicit_builtins()
bool need_sample_mask = msl_options.additional_fixed_sample_mask != 0xffffffff;
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_subgroup_invocation_id || need_sample_mask)
needs_subgroup_invocation_id || needs_subgroup_size || need_sample_mask)
{
bool has_frag_coord = false;
bool has_sample_id = false;
@ -287,7 +287,7 @@ void CompilerMSL::build_implicit_builtins()
has_subgroup_invocation_id = true;
}
if (need_subgroup_ge_mask && builtin == BuiltInSubgroupSize)
if ((need_subgroup_ge_mask || needs_subgroup_size) && builtin == BuiltInSubgroupSize)
{
builtin_subgroup_size_id = var.self;
mark_implicit_builtin(StorageClassInput, BuiltInSubgroupSize, var.self);
@ -593,7 +593,7 @@ void CompilerMSL::build_implicit_builtins()
mark_implicit_builtin(StorageClassInput, BuiltInSubgroupLocalInvocationId, var_id);
}
if (!has_subgroup_size && need_subgroup_ge_mask)
if (!has_subgroup_size && (need_subgroup_ge_mask || needs_subgroup_size))
{
uint32_t offset = ir.increase_bound_by(2);
uint32_t type_ptr_id = offset;
@ -1280,6 +1280,8 @@ void CompilerMSL::preprocess_op_codes()
if (preproc.needs_subgroup_invocation_id)
needs_subgroup_invocation_id = true;
if (preproc.needs_subgroup_size)
needs_subgroup_size = true;
}
// Move the Private and Workgroup global variables to the entry function.
@ -4608,6 +4610,59 @@ void CompilerMSL::emit_custom_functions()
statement("");
break;
case SPVFuncImplSubgroupBroadcast:
// Metal doesn't allow broadcasting boolean values directly, but we can work around that by broadcasting
// them as integers.
statement("template<typename T>");
statement("inline T spvSubgroupBroadcast(T value, ushort lane)");
begin_scope();
if (msl_options.is_ios())
statement("return quad_broadcast(value, lane);");
else
statement("return simd_broadcast(value, lane);");
end_scope();
statement("");
statement("template<>");
statement("inline bool spvSubgroupBroadcast(bool value, ushort lane)");
begin_scope();
if (msl_options.is_ios())
statement("return !!quad_broadcast((ushort)value, lane);");
else
statement("return !!simd_broadcast((ushort)value, lane);");
end_scope();
statement("");
statement("template<uint N>");
statement("inline vec<bool, N> spvSubgroupBroadcast(vec<bool, N> value, ushort lane)");
begin_scope();
if (msl_options.is_ios())
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);");
end_scope();
statement("");
break;
case SPVFuncImplSubgroupBroadcastFirst:
statement("template<typename T>");
statement("inline T spvSubgroupBroadcastFirst(T value)");
begin_scope();
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);");
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);");
end_scope();
statement("");
break;
case SPVFuncImplSubgroupBallot:
statement("inline uint4 spvSubgroupBallot(bool value)");
begin_scope();
@ -4631,8 +4686,11 @@ void CompilerMSL::emit_custom_functions()
break;
case SPVFuncImplSubgroupBallotFindLSB:
statement("inline uint spvSubgroupBallotFindLSB(uint4 ballot)");
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));");
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);");
end_scope();
@ -4640,8 +4698,11 @@ void CompilerMSL::emit_custom_functions()
break;
case SPVFuncImplSubgroupBallotFindMSB:
statement("inline uint spvSubgroupBallotFindMSB(uint4 ballot)");
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));");
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), "
"ballot.z == 0), ballot.w == 0);");
@ -4650,24 +4711,31 @@ void CompilerMSL::emit_custom_functions()
break;
case SPVFuncImplSubgroupBallotBitCount:
statement("inline uint spvSubgroupBallotBitCount(uint4 ballot)");
statement("inline uint spvPopCount4(uint4 ballot)");
begin_scope();
statement("return popcount(ballot.x) + popcount(ballot.y) + popcount(ballot.z) + popcount(ballot.w);");
end_scope();
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));");
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));");
statement("return spvSubgroupBallotBitCount(ballot & mask);");
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));");
statement("return spvSubgroupBallotBitCount(ballot & mask);");
statement("return spvPopCount4(ballot & mask);");
end_scope();
statement("");
break;
@ -4680,7 +4748,7 @@ void CompilerMSL::emit_custom_functions()
statement("template<typename T>");
statement("inline bool spvSubgroupAllEqual(T value)");
begin_scope();
statement("return simd_all(value == simd_broadcast_first(value));");
statement("return simd_all(all(value == simd_broadcast_first(value)));");
end_scope();
statement("");
statement("template<>");
@ -4689,6 +4757,184 @@ void CompilerMSL::emit_custom_functions()
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)));");
end_scope();
statement("");
break;
case SPVFuncImplSubgroupShuffle:
statement("template<typename T>");
statement("inline T spvSubgroupShuffle(T value, ushort lane)");
begin_scope();
if (msl_options.is_ios())
statement("return quad_shuffle(value, lane);");
else
statement("return simd_shuffle(value, lane);");
end_scope();
statement("");
statement("template<>");
statement("inline bool spvSubgroupShuffle(bool value, ushort lane)");
begin_scope();
if (msl_options.is_ios())
statement("return !!quad_shuffle((ushort)value, lane);");
else
statement("return !!simd_shuffle((ushort)value, lane);");
end_scope();
statement("");
statement("template<uint N>");
statement("inline vec<bool, N> spvSubgroupShuffle(vec<bool, N> value, ushort lane)");
begin_scope();
if (msl_options.is_ios())
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);");
end_scope();
statement("");
break;
case SPVFuncImplSubgroupShuffleXor:
statement("template<typename T>");
statement("inline T spvSubgroupShuffleXor(T value, ushort mask)");
begin_scope();
if (msl_options.is_ios())
statement("return quad_shuffle_xor(value, mask);");
else
statement("return simd_shuffle_xor(value, mask);");
end_scope();
statement("");
statement("template<>");
statement("inline bool spvSubgroupShuffleXor(bool value, ushort mask)");
begin_scope();
if (msl_options.is_ios())
statement("return !!quad_shuffle_xor((ushort)value, mask);");
else
statement("return !!simd_shuffle_xor((ushort)value, mask);");
end_scope();
statement("");
statement("template<uint N>");
statement("inline vec<bool, N> spvSubgroupShuffleXor(vec<bool, N> value, ushort mask)");
begin_scope();
if (msl_options.is_ios())
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);");
end_scope();
statement("");
break;
case SPVFuncImplSubgroupShuffleUp:
statement("template<typename T>");
statement("inline T spvSubgroupShuffleUp(T value, ushort delta)");
begin_scope();
if (msl_options.is_ios())
statement("return quad_shuffle_up(value, delta);");
else
statement("return simd_shuffle_up(value, delta);");
end_scope();
statement("");
statement("template<>");
statement("inline bool spvSubgroupShuffleUp(bool value, ushort delta)");
begin_scope();
if (msl_options.is_ios())
statement("return !!quad_shuffle_up((ushort)value, delta);");
else
statement("return !!simd_shuffle_up((ushort)value, delta);");
end_scope();
statement("");
statement("template<uint N>");
statement("inline vec<bool, N> spvSubgroupShuffleUp(vec<bool, N> value, ushort delta)");
begin_scope();
if (msl_options.is_ios())
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);");
end_scope();
statement("");
break;
case SPVFuncImplSubgroupShuffleDown:
statement("template<typename T>");
statement("inline T spvSubgroupShuffleDown(T value, ushort delta)");
begin_scope();
if (msl_options.is_ios())
statement("return quad_shuffle_down(value, delta);");
else
statement("return simd_shuffle_down(value, delta);");
end_scope();
statement("");
statement("template<>");
statement("inline bool spvSubgroupShuffleDown(bool value, ushort delta)");
begin_scope();
if (msl_options.is_ios())
statement("return !!quad_shuffle_down((ushort)value, delta);");
else
statement("return !!simd_shuffle_down((ushort)value, delta);");
end_scope();
statement("");
statement("template<uint N>");
statement("inline vec<bool, N> spvSubgroupShuffleDown(vec<bool, N> value, ushort delta)");
begin_scope();
if (msl_options.is_ios())
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);");
end_scope();
statement("");
break;
case SPVFuncImplQuadBroadcast:
statement("template<typename T>");
statement("inline T spvQuadBroadcast(T value, uint lane)");
begin_scope();
statement("return quad_broadcast(value, lane);");
end_scope();
statement("");
statement("template<>");
statement("inline bool spvQuadBroadcast(bool value, uint lane)");
begin_scope();
statement("return !!quad_broadcast((ushort)value, lane);");
end_scope();
statement("");
statement("template<uint N>");
statement("inline vec<bool, N> spvQuadBroadcast(vec<bool, N> value, uint lane)");
begin_scope();
statement("return (vec<bool, N>)quad_broadcast((vec<ushort, N>)value, lane);");
end_scope();
statement("");
break;
case SPVFuncImplQuadSwap:
// We can implement this easily based on the following table giving
// the target lane ID from the direction and current lane ID:
// Direction
// | 0 | 1 | 2 |
// ---+---+---+---+
// L 0 | 1 2 3
// a 1 | 0 3 2
// n 2 | 3 0 1
// e 3 | 2 1 0
// Notice that target = source ^ (direction + 1).
statement("template<typename T>");
statement("inline T spvQuadSwap(T value, uint dir)");
begin_scope();
statement("return quad_shuffle_xor(value, dir + 1);");
end_scope();
statement("");
statement("template<>");
statement("inline bool spvQuadSwap(bool value, uint dir)");
begin_scope();
statement("return !!quad_shuffle_xor((ushort)value, dir + 1);");
end_scope();
statement("");
statement("template<uint N>");
statement("inline vec<bool, N> spvQuadSwap(vec<bool, N> value, uint dir)");
begin_scope();
statement("return (vec<bool, N>)quad_shuffle_xor((vec<ushort, N>)value, dir + 1);");
end_scope();
statement("");
break;
case SPVFuncImplReflectScalar:
@ -10463,7 +10709,7 @@ void CompilerMSL::fix_up_shader_inputs_outputs()
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 ? uint4(0, (1 << (",
to_expression(builtin_subgroup_invocation_id_id), " - 32)), uint2(0)) : uint4(1 << ",
to_expression(builtin_subgroup_invocation_id_id), ", uint3(0));");
});
@ -10475,25 +10721,25 @@ void CompilerMSL::fix_up_shader_inputs_outputs()
SPIRV_CROSS_THROW("Subgroup ballot functionality requires Metal 2.1.");
entry_func.fixup_hooks_in.push_back([=]() {
// Case where index < 32, size < 32:
// mask0 = bfe(0xFFFFFFFF, index, size - index);
// mask1 = bfe(0xFFFFFFFF, 0, 0); // Gives 0
// mask0 = bfi(0, 0xFFFFFFFF, index, size - index);
// mask1 = bfi(0, 0xFFFFFFFF, 0, 0); // Gives 0
// Case where index < 32 but size >= 32:
// mask0 = bfe(0xFFFFFFFF, index, 32 - index);
// mask1 = bfe(0xFFFFFFFF, 0, size - 32);
// mask0 = bfi(0, 0xFFFFFFFF, index, 32 - index);
// mask1 = bfi(0, 0xFFFFFFFF, 0, size - 32);
// Case where index >= 32:
// mask0 = bfe(0xFFFFFFFF, 32, 0); // Gives 0
// mask1 = bfe(0xFFFFFFFF, index - 32, size - index);
// mask0 = bfi(0, 0xFFFFFFFF, 32, 0); // Gives 0
// mask1 = bfi(0, 0xFFFFFFFF, index - 32, size - index);
// This is expressed without branches to avoid divergent
// control flow--hence the complicated min/max expressions.
// This is further complicated by the fact that if you attempt
// to bfe out-of-bounds on Metal, undefined behavior is the
// to bfi/bfe out-of-bounds on Metal, undefined behavior is the
// result.
statement(builtin_type_decl(bi_type), " ", to_expression(var_id),
" = uint4(extract_bits(0xFFFFFFFF, min(",
" = 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)), extract_bits(0xFFFFFFFF, (uint)max((int)",
", 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));");
@ -10508,11 +10754,11 @@ void CompilerMSL::fix_up_shader_inputs_outputs()
// 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(extract_bits(0xFFFFFFFF, min(",
" = 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)), extract_bits(0xFFFFFFFF, (uint)max((int)",
" - 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));");
@ -11864,12 +12110,11 @@ void CompilerMSL::emit_subgroup_op(const Instruction &i)
break;
case OpGroupNonUniformBroadcast:
emit_binary_func_op(result_type, id, ops[3], ops[4],
msl_options.is_ios() ? "quad_broadcast" : "simd_broadcast");
emit_binary_func_op(result_type, id, ops[3], ops[4], "spvSubgroupBroadcast");
break;
case OpGroupNonUniformBroadcastFirst:
emit_unary_func_op(result_type, id, ops[3], "simd_broadcast_first");
emit_unary_func_op(result_type, id, ops[3], "spvSubgroupBroadcastFirst");
break;
case OpGroupNonUniformBallot:
@ -11885,46 +12130,50 @@ void CompilerMSL::emit_subgroup_op(const Instruction &i)
break;
case OpGroupNonUniformBallotFindLSB:
emit_unary_func_op(result_type, id, ops[3], "spvSubgroupBallotFindLSB");
emit_binary_func_op(result_type, id, ops[3], builtin_subgroup_size_id, "spvSubgroupBallotFindLSB");
break;
case OpGroupNonUniformBallotFindMSB:
emit_unary_func_op(result_type, id, ops[3], "spvSubgroupBallotFindMSB");
emit_binary_func_op(result_type, id, ops[3], builtin_subgroup_size_id, "spvSubgroupBallotFindMSB");
break;
case OpGroupNonUniformBallotBitCount:
{
auto operation = static_cast<GroupOperation>(ops[3]);
if (operation == GroupOperationReduce)
emit_unary_func_op(result_type, id, ops[4], "spvSubgroupBallotBitCount");
else if (operation == GroupOperationInclusiveScan)
switch (operation)
{
case GroupOperationReduce:
emit_binary_func_op(result_type, id, ops[4], builtin_subgroup_size_id, "spvSubgroupBallotBitCount");
break;
case GroupOperationInclusiveScan:
emit_binary_func_op(result_type, id, ops[4], builtin_subgroup_invocation_id_id,
"spvSubgroupBallotInclusiveBitCount");
else if (operation == GroupOperationExclusiveScan)
break;
case GroupOperationExclusiveScan:
emit_binary_func_op(result_type, id, ops[4], builtin_subgroup_invocation_id_id,
"spvSubgroupBallotExclusiveBitCount");
else
break;
default:
SPIRV_CROSS_THROW("Invalid BitCount operation.");
break;
}
break;
}
case OpGroupNonUniformShuffle:
emit_binary_func_op(result_type, id, ops[3], ops[4], msl_options.is_ios() ? "quad_shuffle" : "simd_shuffle");
emit_binary_func_op(result_type, id, ops[3], ops[4], "spvSubgroupShuffle");
break;
case OpGroupNonUniformShuffleXor:
emit_binary_func_op(result_type, id, ops[3], ops[4],
msl_options.is_ios() ? "quad_shuffle_xor" : "simd_shuffle_xor");
emit_binary_func_op(result_type, id, ops[3], ops[4], "spvSubgroupShuffleXor");
break;
case OpGroupNonUniformShuffleUp:
emit_binary_func_op(result_type, id, ops[3], ops[4],
msl_options.is_ios() ? "quad_shuffle_up" : "simd_shuffle_up");
emit_binary_func_op(result_type, id, ops[3], ops[4], "spvSubgroupShuffleUp");
break;
case OpGroupNonUniformShuffleDown:
emit_binary_func_op(result_type, id, ops[3], ops[4],
msl_options.is_ios() ? "quad_shuffle_down" : "simd_shuffle_down");
emit_binary_func_op(result_type, id, ops[3], ops[4], "spvSubgroupShuffleDown");
break;
case OpGroupNonUniformAll:
@ -12032,26 +12281,11 @@ case OpGroupNonUniform##op: \
#undef MSL_GROUP_OP_CAST
case OpGroupNonUniformQuadSwap:
{
// We can implement this easily based on the following table giving
// the target lane ID from the direction and current lane ID:
// Direction
// | 0 | 1 | 2 |
// ---+---+---+---+
// L 0 | 1 2 3
// a 1 | 0 3 2
// n 2 | 3 0 1
// e 3 | 2 1 0
// Notice that target = source ^ (direction + 1).
uint32_t mask = evaluate_constant_u32(ops[4]) + 1;
uint32_t mask_id = ir.increase_bound_by(1);
set<SPIRConstant>(mask_id, expression_type_id(ops[4]), mask, false);
emit_binary_func_op(result_type, id, ops[3], mask_id, "quad_shuffle_xor");
emit_binary_func_op(result_type, id, ops[3], ops[4], "spvQuadSwap");
break;
}
case OpGroupNonUniformQuadBroadcast:
emit_binary_func_op(result_type, id, ops[3], ops[4], "quad_broadcast");
emit_binary_func_op(result_type, id, ops[3], ops[4], "spvQuadBroadcast");
break;
default:
@ -13004,8 +13238,15 @@ bool CompilerMSL::OpCodePreprocessor::handle(Op opcode, const uint32_t *args, ui
needs_subgroup_invocation_id = true;
break;
case OpGroupNonUniformBallotFindLSB:
case OpGroupNonUniformBallotFindMSB:
needs_subgroup_size = true;
break;
case OpGroupNonUniformBallotBitCount:
if (args[3] != GroupOperationReduce)
if (args[3] == GroupOperationReduce)
needs_subgroup_size = true;
else
needs_subgroup_invocation_id = true;
break;
@ -13188,6 +13429,12 @@ CompilerMSL::SPVFuncImpl CompilerMSL::OpCodePreprocessor::get_spv_func_impl(Op o
break;
}
case OpGroupNonUniformBroadcast:
return SPVFuncImplSubgroupBroadcast;
case OpGroupNonUniformBroadcastFirst:
return SPVFuncImplSubgroupBroadcastFirst;
case OpGroupNonUniformBallot:
return SPVFuncImplSubgroupBallot;
@ -13207,6 +13454,24 @@ CompilerMSL::SPVFuncImpl CompilerMSL::OpCodePreprocessor::get_spv_func_impl(Op o
case OpGroupNonUniformAllEqual:
return SPVFuncImplSubgroupAllEqual;
case OpGroupNonUniformShuffle:
return SPVFuncImplSubgroupShuffle;
case OpGroupNonUniformShuffleXor:
return SPVFuncImplSubgroupShuffleXor;
case OpGroupNonUniformShuffleUp:
return SPVFuncImplSubgroupShuffleUp;
case OpGroupNonUniformShuffleDown:
return SPVFuncImplSubgroupShuffleDown;
case OpGroupNonUniformQuadBroadcast:
return SPVFuncImplQuadBroadcast;
case OpGroupNonUniformQuadSwap:
return SPVFuncImplQuadSwap;
default:
break;
}

10
spirv_msl.hpp
View File

@ -600,12 +600,20 @@ protected:
SPVFuncImplTextureSwizzle,
SPVFuncImplGatherSwizzle,
SPVFuncImplGatherCompareSwizzle,
SPVFuncImplSubgroupBroadcast,
SPVFuncImplSubgroupBroadcastFirst,
SPVFuncImplSubgroupBallot,
SPVFuncImplSubgroupBallotBitExtract,
SPVFuncImplSubgroupBallotFindLSB,
SPVFuncImplSubgroupBallotFindMSB,
SPVFuncImplSubgroupBallotBitCount,
SPVFuncImplSubgroupAllEqual,
SPVFuncImplSubgroupShuffle,
SPVFuncImplSubgroupShuffleXor,
SPVFuncImplSubgroupShuffleUp,
SPVFuncImplSubgroupShuffleDown,
SPVFuncImplQuadBroadcast,
SPVFuncImplQuadSwap,
SPVFuncImplReflectScalar,
SPVFuncImplRefractScalar,
SPVFuncImplFaceForwardScalar,
@ -913,6 +921,7 @@ protected:
bool used_swizzle_buffer = false;
bool added_builtin_tess_level = false;
bool needs_subgroup_invocation_id = false;
bool needs_subgroup_size = false;
std::string qual_pos_var_name;
std::string stage_in_var_name = "in";
std::string stage_out_var_name = "out";
@ -984,6 +993,7 @@ protected:
bool uses_atomics = false;
bool uses_resource_write = false;
bool needs_subgroup_invocation_id = false;
bool needs_subgroup_size = false;
};
// OpcodeHandler that scans for uses of sampled images