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

MSL: Add test for logical subgroup arith ops.

Browse Source
This commit is contained in:
Hans-Kristian Arntzen 2021-03-08 12:57:37 +01:00
parent d6c2c1b39a
commit aea6d29aa8
2 changed files with 13 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -288,6 +288,9 @@ kernel void main0(device SSBO& _9 [[buffer(0)]], uint gl_NumSubgroups [[simdgrou
uint4 anded = simd_and(ballot_value);
uint4 ored = simd_or(ballot_value);
uint4 xored = simd_xor(ballot_value);
bool4 anded_b = simd_and(ballot_value == uint4(42u));
bool4 ored_b = simd_or(ballot_value == uint4(42u));
bool4 xored_b = simd_xor(ballot_value == uint4(42u));
added = simd_prefix_inclusive_sum(added);
iadded = simd_prefix_inclusive_sum(iadded);
multiplied = simd_prefix_inclusive_product(multiplied);
@ -309,6 +312,9 @@ kernel void main0(device SSBO& _9 [[buffer(0)]], uint gl_NumSubgroups [[simdgrou
anded = quad_and(anded);
ored = quad_or(ored);
xored = quad_xor(xored);
anded_b = quad_and(anded == uint4(2u));
ored_b = quad_or(ored == uint4(3u));
xored_b = quad_xor(xored == uint4(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);

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

@ -81,6 +81,9 @@ void main()
uvec4 anded = subgroupAnd(ballot_value);
uvec4 ored = subgroupOr(ballot_value);
uvec4 xored = subgroupXor(ballot_value);
bvec4 anded_b = subgroupAnd(equal(ballot_value, uvec4(42)));
bvec4 ored_b = subgroupOr(equal(ballot_value, uvec4(42)));
bvec4 xored_b = subgroupXor(equal(ballot_value, uvec4(42)));
added = subgroupInclusiveAdd(added);
iadded = subgroupInclusiveAdd(iadded);
@ -126,6 +129,10 @@ void main()
ored = subgroupClusteredOr(ored, 4u);
xored = subgroupClusteredXor(xored, 4u);
anded_b = subgroupClusteredAnd(equal(anded, uvec4(2u)), 4u);
ored_b = subgroupClusteredOr(equal(ored, uvec4(3u)), 4u);
xored_b = subgroupClusteredXor(equal(xored, uvec4(4u)), 4u);
// quad
vec4 swap_horiz = subgroupQuadSwapHorizontal(vec4(20.0));
bvec4 swap_horiz_bool = subgroupQuadSwapHorizontal(bvec4(true));