mirror of
https://github.com/KhronosGroup/SPIRV-Cross.git
synced 2024-11-09 22:00:05 +00:00
Move all .invalid shaders into no-opt folders.
This commit is contained in:
Dan Sinclair
parent
e5af41255c
commit
d409210ee5
@ -1,32 +0,0 @@
|
|||||||
#ifndef SPIRV_CROSS_CONSTANT_ID_0
|
|
||||||
#define SPIRV_CROSS_CONSTANT_ID_0 1u
|
|
||||||
#endif
|
|
||||||
static const uint _3 = SPIRV_CROSS_CONSTANT_ID_0;
|
|
||||||
#ifndef SPIRV_CROSS_CONSTANT_ID_2
|
|
||||||
#define SPIRV_CROSS_CONSTANT_ID_2 3u
|
|
||||||
#endif
|
|
||||||
static const uint _4 = SPIRV_CROSS_CONSTANT_ID_2;
|
|
||||||
static const uint3 gl_WorkGroupSize = uint3(_3, 2u, _4);
|
|
||||||
|
|
||||||
RWByteAddressBuffer _8 : register(u0);
|
|
||||||
RWByteAddressBuffer _9 : register(u1);
|
|
||||||
|
|
||||||
static uint3 gl_WorkGroupID;
|
|
||||||
struct SPIRV_Cross_Input
|
|
||||||
{
|
|
||||||
uint3 gl_WorkGroupID : SV_GroupID;
|
|
||||||
};
|
|
||||||
|
|
||||||
static uint3 _22 = gl_WorkGroupSize;
|
|
||||||
|
|
||||||
void comp_main()
|
|
||||||
{
|
|
||||||
_8.Store(gl_WorkGroupID.x * 4 + 0, asuint(asfloat(_9.Load(gl_WorkGroupID.x * 4 + 0)) + asfloat(_8.Load(gl_WorkGroupID.x * 4 + 0))));
|
|
||||||
}
|
|
||||||
|
|
||||||
[numthreads(SPIRV_CROSS_CONSTANT_ID_0, 2, SPIRV_CROSS_CONSTANT_ID_2)]
|
|
||||||
void main(SPIRV_Cross_Input stage_input)
|
|
||||||
{
|
|
||||||
gl_WorkGroupID = stage_input.gl_WorkGroupID;
|
|
||||||
comp_main();
|
|
||||||
}
|
|
||||||
@ -1,95 +0,0 @@
|
|||||||
static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u);
|
|
||||||
|
|
||||||
RWByteAddressBuffer _9 : register(u0, space0);
|
|
||||||
|
|
||||||
static uint4 gl_SubgroupEqMask;
|
|
||||||
static uint4 gl_SubgroupGeMask;
|
|
||||||
static uint4 gl_SubgroupGtMask;
|
|
||||||
static uint4 gl_SubgroupLeMask;
|
|
||||||
static uint4 gl_SubgroupLtMask;
|
|
||||||
void comp_main()
|
|
||||||
{
|
|
||||||
_9.Store(0, asuint(float(WaveGetLaneCount())));
|
|
||||||
_9.Store(0, asuint(float(WaveGetLaneIndex())));
|
|
||||||
bool elected = WaveIsFirstLane();
|
|
||||||
_9.Store(0, asuint(float4(gl_SubgroupEqMask).x));
|
|
||||||
_9.Store(0, asuint(float4(gl_SubgroupGeMask).x));
|
|
||||||
_9.Store(0, asuint(float4(gl_SubgroupGtMask).x));
|
|
||||||
_9.Store(0, asuint(float4(gl_SubgroupLeMask).x));
|
|
||||||
_9.Store(0, asuint(float4(gl_SubgroupLtMask).x));
|
|
||||||
float4 broadcasted = WaveReadLaneAt(10.0f.xxxx, 8u);
|
|
||||||
float3 first = WaveReadLaneFirst(20.0f.xxx);
|
|
||||||
uint4 ballot_value = WaveActiveBallot(true);
|
|
||||||
uint bit_count = countbits(ballot_value.x) + countbits(ballot_value.y) + countbits(ballot_value.z) + countbits(ballot_value.w);
|
|
||||||
bool has_all = WaveActiveAllTrue(true);
|
|
||||||
bool has_any = WaveActiveAnyTrue(true);
|
|
||||||
bool has_equal = WaveActiveAllEqualBool(true);
|
|
||||||
float4 added = WaveActiveSum(20.0f.xxxx);
|
|
||||||
int4 iadded = WaveActiveSum(int4(20, 20, 20, 20));
|
|
||||||
float4 multiplied = WaveActiveProduct(20.0f.xxxx);
|
|
||||||
int4 imultiplied = WaveActiveProduct(int4(20, 20, 20, 20));
|
|
||||||
float4 lo = WaveActiveMin(20.0f.xxxx);
|
|
||||||
float4 hi = WaveActiveMax(20.0f.xxxx);
|
|
||||||
int4 slo = WaveActiveMin(int4(20, 20, 20, 20));
|
|
||||||
int4 shi = WaveActiveMax(int4(20, 20, 20, 20));
|
|
||||||
uint4 ulo = WaveActiveMin(uint4(20u, 20u, 20u, 20u));
|
|
||||||
uint4 uhi = WaveActiveMax(uint4(20u, 20u, 20u, 20u));
|
|
||||||
uint4 anded = WaveActiveBitAnd(ballot_value);
|
|
||||||
uint4 ored = WaveActiveBitOr(ballot_value);
|
|
||||||
uint4 xored = WaveActiveBitXor(ballot_value);
|
|
||||||
added = WavePrefixSum(added) + added;
|
|
||||||
iadded = WavePrefixSum(iadded) + iadded;
|
|
||||||
multiplied = WavePrefixProduct(multiplied) * multiplied;
|
|
||||||
imultiplied = WavePrefixProduct(imultiplied) * imultiplied;
|
|
||||||
added = WavePrefixSum(multiplied);
|
|
||||||
multiplied = WavePrefixProduct(multiplied);
|
|
||||||
iadded = WavePrefixSum(imultiplied);
|
|
||||||
imultiplied = WavePrefixProduct(imultiplied);
|
|
||||||
float4 swap_horiz = QuadReadAcrossX(20.0f.xxxx);
|
|
||||||
float4 swap_vertical = QuadReadAcrossY(20.0f.xxxx);
|
|
||||||
float4 swap_diagonal = QuadReadAcrossDiagonal(20.0f.xxxx);
|
|
||||||
float4 quad_broadcast = QuadReadLaneAt(20.0f.xxxx, 3u);
|
|
||||||
}
|
|
||||||
|
|
||||||
[numthreads(1, 1, 1)]
|
|
||||||
void main()
|
|
||||||
{
|
|
||||||
gl_SubgroupEqMask = 1u << (WaveGetLaneIndex() - uint4(0, 32, 64, 96));
|
|
||||||
if (WaveGetLaneIndex() >= 32) gl_SubgroupEqMask.x = 0;
|
|
||||||
if (WaveGetLaneIndex() >= 64 || WaveGetLaneIndex() < 32) gl_SubgroupEqMask.y = 0;
|
|
||||||
if (WaveGetLaneIndex() >= 96 || WaveGetLaneIndex() < 64) gl_SubgroupEqMask.z = 0;
|
|
||||||
if (WaveGetLaneIndex() < 96) gl_SubgroupEqMask.w = 0;
|
|
||||||
gl_SubgroupGeMask = ~((1u << (WaveGetLaneIndex() - uint4(0, 32, 64, 96))) - 1u);
|
|
||||||
if (WaveGetLaneIndex() >= 32) gl_SubgroupGeMask.x = 0u;
|
|
||||||
if (WaveGetLaneIndex() >= 64) gl_SubgroupGeMask.y = 0u;
|
|
||||||
if (WaveGetLaneIndex() >= 96) gl_SubgroupGeMask.z = 0u;
|
|
||||||
if (WaveGetLaneIndex() < 32) gl_SubgroupGeMask.y = ~0u;
|
|
||||||
if (WaveGetLaneIndex() < 64) gl_SubgroupGeMask.z = ~0u;
|
|
||||||
if (WaveGetLaneIndex() < 96) gl_SubgroupGeMask.w = ~0u;
|
|
||||||
uint gt_lane_index = WaveGetLaneIndex() + 1;
|
|
||||||
gl_SubgroupGtMask = ~((1u << (gt_lane_index - uint4(0, 32, 64, 96))) - 1u);
|
|
||||||
if (gt_lane_index >= 32) gl_SubgroupGtMask.x = 0u;
|
|
||||||
if (gt_lane_index >= 64) gl_SubgroupGtMask.y = 0u;
|
|
||||||
if (gt_lane_index >= 96) gl_SubgroupGtMask.z = 0u;
|
|
||||||
if (gt_lane_index >= 128) gl_SubgroupGtMask.w = 0u;
|
|
||||||
if (gt_lane_index < 32) gl_SubgroupGtMask.y = ~0u;
|
|
||||||
if (gt_lane_index < 64) gl_SubgroupGtMask.z = ~0u;
|
|
||||||
if (gt_lane_index < 96) gl_SubgroupGtMask.w = ~0u;
|
|
||||||
uint le_lane_index = WaveGetLaneIndex() + 1;
|
|
||||||
gl_SubgroupLeMask = (1u << (le_lane_index - uint4(0, 32, 64, 96))) - 1u;
|
|
||||||
if (le_lane_index >= 32) gl_SubgroupLeMask.x = ~0u;
|
|
||||||
if (le_lane_index >= 64) gl_SubgroupLeMask.y = ~0u;
|
|
||||||
if (le_lane_index >= 96) gl_SubgroupLeMask.z = ~0u;
|
|
||||||
if (le_lane_index >= 128) gl_SubgroupLeMask.w = ~0u;
|
|
||||||
if (le_lane_index < 32) gl_SubgroupLeMask.y = 0u;
|
|
||||||
if (le_lane_index < 64) gl_SubgroupLeMask.z = 0u;
|
|
||||||
if (le_lane_index < 96) gl_SubgroupLeMask.w = 0u;
|
|
||||||
gl_SubgroupLtMask = (1u << (WaveGetLaneIndex() - uint4(0, 32, 64, 96))) - 1u;
|
|
||||||
if (WaveGetLaneIndex() >= 32) gl_SubgroupLtMask.x = ~0u;
|
|
||||||
if (WaveGetLaneIndex() >= 64) gl_SubgroupLtMask.y = ~0u;
|
|
||||||
if (WaveGetLaneIndex() >= 96) gl_SubgroupLtMask.z = ~0u;
|
|
||||||
if (WaveGetLaneIndex() < 32) gl_SubgroupLtMask.y = 0u;
|
|
||||||
if (WaveGetLaneIndex() < 64) gl_SubgroupLtMask.z = 0u;
|
|
||||||
if (WaveGetLaneIndex() < 96) gl_SubgroupLtMask.w = 0u;
|
|
||||||
comp_main();
|
|
||||||
}
|
|
||||||
@ -1,44 +0,0 @@
|
|||||||
struct CBO_1
|
|
||||||
{
|
|
||||||
float4 a;
|
|
||||||
float4 b;
|
|
||||||
float4 c;
|
|
||||||
float4 d;
|
|
||||||
};
|
|
||||||
|
|
||||||
ConstantBuffer<CBO_1> cbo[2][4] : register(b4, space0);
|
|
||||||
cbuffer PushMe
|
|
||||||
{
|
|
||||||
float4 push_a : packoffset(c0);
|
|
||||||
float4 push_b : packoffset(c1);
|
|
||||||
float4 push_c : packoffset(c2);
|
|
||||||
float4 push_d : packoffset(c3);
|
|
||||||
};
|
|
||||||
|
|
||||||
|
|
||||||
static float4 FragColor;
|
|
||||||
|
|
||||||
struct SPIRV_Cross_Output
|
|
||||||
{
|
|
||||||
float4 FragColor : SV_Target0;
|
|
||||||
};
|
|
||||||
|
|
||||||
void frag_main()
|
|
||||||
{
|
|
||||||
FragColor = cbo[1][2].a;
|
|
||||||
FragColor += cbo[1][2].b;
|
|
||||||
FragColor += cbo[1][2].c;
|
|
||||||
FragColor += cbo[1][2].d;
|
|
||||||
FragColor += push_a;
|
|
||||||
FragColor += push_b;
|
|
||||||
FragColor += push_c;
|
|
||||||
FragColor += push_d;
|
|
||||||
}
|
|
||||||
|
|
||||||
SPIRV_Cross_Output main()
|
|
||||||
{
|
|
||||||
frag_main();
|
|
||||||
SPIRV_Cross_Output stage_output;
|
|
||||||
stage_output.FragColor = FragColor;
|
|
||||||
return stage_output;
|
|
||||||
}
|
|
||||||
@ -1,179 +0,0 @@
|
|||||||
static min16float4 v4;
|
|
||||||
static min16float3 v3;
|
|
||||||
static min16float v1;
|
|
||||||
static min16float2 v2;
|
|
||||||
static float o1;
|
|
||||||
static float2 o2;
|
|
||||||
static float3 o3;
|
|
||||||
static float4 o4;
|
|
||||||
|
|
||||||
struct SPIRV_Cross_Input
|
|
||||||
{
|
|
||||||
min16float v1 : TEXCOORD0;
|
|
||||||
min16float2 v2 : TEXCOORD1;
|
|
||||||
min16float3 v3 : TEXCOORD2;
|
|
||||||
min16float4 v4 : TEXCOORD3;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct SPIRV_Cross_Output
|
|
||||||
{
|
|
||||||
float o1 : SV_Target0;
|
|
||||||
float2 o2 : SV_Target1;
|
|
||||||
float3 o3 : SV_Target2;
|
|
||||||
float4 o4 : SV_Target3;
|
|
||||||
};
|
|
||||||
|
|
||||||
float mod(float x, float y)
|
|
||||||
{
|
|
||||||
return x - y * floor(x / y);
|
|
||||||
}
|
|
||||||
|
|
||||||
float2 mod(float2 x, float2 y)
|
|
||||||
{
|
|
||||||
return x - y * floor(x / y);
|
|
||||||
}
|
|
||||||
|
|
||||||
float3 mod(float3 x, float3 y)
|
|
||||||
{
|
|
||||||
return x - y * floor(x / y);
|
|
||||||
}
|
|
||||||
|
|
||||||
float4 mod(float4 x, float4 y)
|
|
||||||
{
|
|
||||||
return x - y * floor(x / y);
|
|
||||||
}
|
|
||||||
|
|
||||||
uint SPIRV_Cross_packFloat2x16(min16float2 value)
|
|
||||||
{
|
|
||||||
uint2 Packed = f32tof16(value);
|
|
||||||
return Packed.x | (Packed.y << 16);
|
|
||||||
}
|
|
||||||
|
|
||||||
min16float2 SPIRV_Cross_unpackFloat2x16(uint value)
|
|
||||||
{
|
|
||||||
return min16float2(f16tof32(uint2(value & 0xffff, value >> 16)));
|
|
||||||
}
|
|
||||||
|
|
||||||
void test_constants()
|
|
||||||
{
|
|
||||||
min16float a = min16float(1.0);
|
|
||||||
min16float b = min16float(1.5);
|
|
||||||
min16float c = min16float(-1.5);
|
|
||||||
min16float d = min16float(0.0 / 0.0);
|
|
||||||
min16float e = min16float(1.0 / 0.0);
|
|
||||||
min16float f = min16float(-1.0 / 0.0);
|
|
||||||
min16float g = min16float(1014.0);
|
|
||||||
min16float h = min16float(9.5367431640625e-07);
|
|
||||||
}
|
|
||||||
|
|
||||||
min16float test_result()
|
|
||||||
{
|
|
||||||
return min16float(1.0);
|
|
||||||
}
|
|
||||||
|
|
||||||
void test_conversions()
|
|
||||||
{
|
|
||||||
min16float one = test_result();
|
|
||||||
int a = int(one);
|
|
||||||
uint b = uint(one);
|
|
||||||
bool c = one != min16float(0.0);
|
|
||||||
float d = float(one);
|
|
||||||
double e = double(one);
|
|
||||||
min16float a2 = min16float(a);
|
|
||||||
min16float b2 = min16float(b);
|
|
||||||
min16float c2 = min16float(c);
|
|
||||||
min16float d2 = min16float(d);
|
|
||||||
min16float e2 = min16float(e);
|
|
||||||
}
|
|
||||||
|
|
||||||
void test_builtins()
|
|
||||||
{
|
|
||||||
min16float4 res = radians(v4);
|
|
||||||
res = degrees(v4);
|
|
||||||
res = sin(v4);
|
|
||||||
res = cos(v4);
|
|
||||||
res = tan(v4);
|
|
||||||
res = asin(v4);
|
|
||||||
res = atan2(v4, v3.xyzz);
|
|
||||||
res = atan(v4);
|
|
||||||
res = sinh(v4);
|
|
||||||
res = cosh(v4);
|
|
||||||
res = tanh(v4);
|
|
||||||
res = pow(v4, v4);
|
|
||||||
res = exp(v4);
|
|
||||||
res = log(v4);
|
|
||||||
res = exp2(v4);
|
|
||||||
res = log2(v4);
|
|
||||||
res = sqrt(v4);
|
|
||||||
res = rsqrt(v4);
|
|
||||||
res = abs(v4);
|
|
||||||
res = sign(v4);
|
|
||||||
res = floor(v4);
|
|
||||||
res = trunc(v4);
|
|
||||||
res = round(v4);
|
|
||||||
res = ceil(v4);
|
|
||||||
res = frac(v4);
|
|
||||||
res = mod(v4, v4);
|
|
||||||
min16float4 tmp;
|
|
||||||
min16float4 _144 = modf(v4, tmp);
|
|
||||||
res = _144;
|
|
||||||
res = min(v4, v4);
|
|
||||||
res = max(v4, v4);
|
|
||||||
res = clamp(v4, v4, v4);
|
|
||||||
res = lerp(v4, v4, v4);
|
|
||||||
bool4 _164 = bool4(v4.x < v4.x, v4.y < v4.y, v4.z < v4.z, v4.w < v4.w);
|
|
||||||
res = min16float4(_164.x ? v4.x : v4.x, _164.y ? v4.y : v4.y, _164.z ? v4.z : v4.z, _164.w ? v4.w : v4.w);
|
|
||||||
res = step(v4, v4);
|
|
||||||
res = smoothstep(v4, v4, v4);
|
|
||||||
bool4 btmp = isnan(v4);
|
|
||||||
btmp = isinf(v4);
|
|
||||||
res = mad(v4, v4, v4);
|
|
||||||
uint pack0 = SPIRV_Cross_packFloat2x16(v4.xy);
|
|
||||||
uint pack1 = SPIRV_Cross_packFloat2x16(v4.zw);
|
|
||||||
res = min16float4(SPIRV_Cross_unpackFloat2x16(pack0), SPIRV_Cross_unpackFloat2x16(pack1));
|
|
||||||
min16float t0 = length(v4);
|
|
||||||
t0 = distance(v4, v4);
|
|
||||||
t0 = dot(v4, v4);
|
|
||||||
min16float3 res3 = cross(v3, v3);
|
|
||||||
res = normalize(v4);
|
|
||||||
res = faceforward(v4, v4, v4);
|
|
||||||
res = reflect(v4, v4);
|
|
||||||
res = refract(v4, v4, v1);
|
|
||||||
btmp = bool4(v4.x < v4.x, v4.y < v4.y, v4.z < v4.z, v4.w < v4.w);
|
|
||||||
btmp = bool4(v4.x <= v4.x, v4.y <= v4.y, v4.z <= v4.z, v4.w <= v4.w);
|
|
||||||
btmp = bool4(v4.x > v4.x, v4.y > v4.y, v4.z > v4.z, v4.w > v4.w);
|
|
||||||
btmp = bool4(v4.x >= v4.x, v4.y >= v4.y, v4.z >= v4.z, v4.w >= v4.w);
|
|
||||||
btmp = bool4(v4.x == v4.x, v4.y == v4.y, v4.z == v4.z, v4.w == v4.w);
|
|
||||||
btmp = bool4(v4.x != v4.x, v4.y != v4.y, v4.z != v4.z, v4.w != v4.w);
|
|
||||||
res = ddx(v4);
|
|
||||||
res = ddy(v4);
|
|
||||||
res = ddx_fine(v4);
|
|
||||||
res = ddy_fine(v4);
|
|
||||||
res = ddx_coarse(v4);
|
|
||||||
res = ddy_coarse(v4);
|
|
||||||
res = fwidth(v4);
|
|
||||||
res = fwidth(v4);
|
|
||||||
res = fwidth(v4);
|
|
||||||
}
|
|
||||||
|
|
||||||
void frag_main()
|
|
||||||
{
|
|
||||||
test_constants();
|
|
||||||
test_conversions();
|
|
||||||
test_builtins();
|
|
||||||
}
|
|
||||||
|
|
||||||
SPIRV_Cross_Output main(SPIRV_Cross_Input stage_input)
|
|
||||||
{
|
|
||||||
v4 = stage_input.v4;
|
|
||||||
v3 = stage_input.v3;
|
|
||||||
v1 = stage_input.v1;
|
|
||||||
v2 = stage_input.v2;
|
|
||||||
frag_main();
|
|
||||||
SPIRV_Cross_Output stage_output;
|
|
||||||
stage_output.o1 = o1;
|
|
||||||
stage_output.o2 = o2;
|
|
||||||
stage_output.o3 = o3;
|
|
||||||
stage_output.o4 = o4;
|
|
||||||
return stage_output;
|
|
||||||
}
|
|
||||||
@ -1,22 +0,0 @@
|
|||||||
#include <metal_stdlib>
|
|
||||||
#include <simd/simd.h>
|
|
||||||
|
|
||||||
using namespace metal;
|
|
||||||
|
|
||||||
struct _6
|
|
||||||
{
|
|
||||||
float _m0[1];
|
|
||||||
};
|
|
||||||
|
|
||||||
constant uint _3_tmp [[function_constant(0)]];
|
|
||||||
constant uint _3 = is_function_constant_defined(_3_tmp) ? _3_tmp : 1u;
|
|
||||||
constant uint _4_tmp [[function_constant(2)]];
|
|
||||||
constant uint _4 = is_function_constant_defined(_4_tmp) ? _4_tmp : 3u;
|
|
||||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(_3, 2u, _4);
|
|
||||||
|
|
||||||
kernel void main0(device _6& _8 [[buffer(0)]], device _6& _9 [[buffer(1)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
|
|
||||||
{
|
|
||||||
_8._m0[gl_WorkGroupID.x] = _9._m0[gl_WorkGroupID.x] + _8._m0[gl_WorkGroupID.x];
|
|
||||||
uint3 _23 = gl_WorkGroupSize;
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,69 +0,0 @@
|
|||||||
#include <metal_stdlib>
|
|
||||||
#include <simd/simd.h>
|
|
||||||
|
|
||||||
using namespace metal;
|
|
||||||
|
|
||||||
struct main0_in
|
|
||||||
{
|
|
||||||
float v1 [[user(locn0)]];
|
|
||||||
float2 v2 [[user(locn1)]];
|
|
||||||
float3 v3 [[user(locn2)]];
|
|
||||||
float4 v4 [[user(locn3)]];
|
|
||||||
half h1 [[user(locn4)]];
|
|
||||||
half2 h2 [[user(locn5)]];
|
|
||||||
half3 h3 [[user(locn6)]];
|
|
||||||
half4 h4 [[user(locn7)]];
|
|
||||||
};
|
|
||||||
|
|
||||||
fragment void main0(main0_in in [[stage_in]])
|
|
||||||
{
|
|
||||||
float res = fast::min(in.v1, in.v1);
|
|
||||||
res = fast::max(in.v1, in.v1);
|
|
||||||
res = fast::clamp(in.v1, in.v1, in.v1);
|
|
||||||
res = precise::min(in.v1, in.v1);
|
|
||||||
res = precise::max(in.v1, in.v1);
|
|
||||||
res = precise::clamp(in.v1, in.v1, in.v1);
|
|
||||||
float2 res2 = fast::min(in.v2, in.v2);
|
|
||||||
res2 = fast::max(in.v2, in.v2);
|
|
||||||
res2 = fast::clamp(in.v2, in.v2, in.v2);
|
|
||||||
res2 = precise::min(in.v2, in.v2);
|
|
||||||
res2 = precise::max(in.v2, in.v2);
|
|
||||||
res2 = precise::clamp(in.v2, in.v2, in.v2);
|
|
||||||
float3 res3 = fast::min(in.v3, in.v3);
|
|
||||||
res3 = fast::max(in.v3, in.v3);
|
|
||||||
res3 = fast::clamp(in.v3, in.v3, in.v3);
|
|
||||||
res3 = precise::min(in.v3, in.v3);
|
|
||||||
res3 = precise::max(in.v3, in.v3);
|
|
||||||
res3 = precise::clamp(in.v3, in.v3, in.v3);
|
|
||||||
float4 res4 = fast::min(in.v4, in.v4);
|
|
||||||
res4 = fast::max(in.v4, in.v4);
|
|
||||||
res4 = fast::clamp(in.v4, in.v4, in.v4);
|
|
||||||
res4 = precise::min(in.v4, in.v4);
|
|
||||||
res4 = precise::max(in.v4, in.v4);
|
|
||||||
res4 = precise::clamp(in.v4, in.v4, in.v4);
|
|
||||||
half hres = min(in.h1, in.h1);
|
|
||||||
hres = max(in.h1, in.h1);
|
|
||||||
hres = clamp(in.h1, in.h1, in.h1);
|
|
||||||
hres = min(in.h1, in.h1);
|
|
||||||
hres = max(in.h1, in.h1);
|
|
||||||
hres = clamp(in.h1, in.h1, in.h1);
|
|
||||||
half2 hres2 = min(in.h2, in.h2);
|
|
||||||
hres2 = max(in.h2, in.h2);
|
|
||||||
hres2 = clamp(in.h2, in.h2, in.h2);
|
|
||||||
hres2 = min(in.h2, in.h2);
|
|
||||||
hres2 = max(in.h2, in.h2);
|
|
||||||
hres2 = clamp(in.h2, in.h2, in.h2);
|
|
||||||
half3 hres3 = min(in.h3, in.h3);
|
|
||||||
hres3 = max(in.h3, in.h3);
|
|
||||||
hres3 = clamp(in.h3, in.h3, in.h3);
|
|
||||||
hres3 = min(in.h3, in.h3);
|
|
||||||
hres3 = max(in.h3, in.h3);
|
|
||||||
hres3 = clamp(in.h3, in.h3, in.h3);
|
|
||||||
half4 hres4 = min(in.h4, in.h4);
|
|
||||||
hres4 = max(in.h4, in.h4);
|
|
||||||
hres4 = clamp(in.h4, in.h4, in.h4);
|
|
||||||
hres4 = min(in.h4, in.h4);
|
|
||||||
hres4 = max(in.h4, in.h4);
|
|
||||||
hres4 = clamp(in.h4, in.h4, in.h4);
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,141 +0,0 @@
|
|||||||
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
|
||||||
#pragma clang diagnostic ignored "-Wmissing-braces"
|
|
||||||
|
|
||||||
#include <metal_stdlib>
|
|
||||||
#include <simd/simd.h>
|
|
||||||
|
|
||||||
using namespace metal;
|
|
||||||
|
|
||||||
template<typename T, size_t Num>
|
|
||||||
struct spvUnsafeArray
|
|
||||||
{
|
|
||||||
T elements[Num ? Num : 1];
|
|
||||||
|
|
||||||
thread T& operator [] (size_t pos) thread
|
|
||||||
{
|
|
||||||
return elements[pos];
|
|
||||||
}
|
|
||||||
constexpr const thread T& operator [] (size_t pos) const thread
|
|
||||||
{
|
|
||||||
return elements[pos];
|
|
||||||
}
|
|
||||||
|
|
||||||
device T& operator [] (size_t pos) device
|
|
||||||
{
|
|
||||||
return elements[pos];
|
|
||||||
}
|
|
||||||
constexpr const device T& operator [] (size_t pos) const device
|
|
||||||
{
|
|
||||||
return elements[pos];
|
|
||||||
}
|
|
||||||
|
|
||||||
constexpr const constant T& operator [] (size_t pos) const constant
|
|
||||||
{
|
|
||||||
return elements[pos];
|
|
||||||
}
|
|
||||||
|
|
||||||
threadgroup T& operator [] (size_t pos) threadgroup
|
|
||||||
{
|
|
||||||
return elements[pos];
|
|
||||||
}
|
|
||||||
constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
|
|
||||||
{
|
|
||||||
return elements[pos];
|
|
||||||
}
|
|
||||||
};
|
|
||||||
|
|
||||||
struct VertexOutput
|
|
||||||
{
|
|
||||||
float4 pos;
|
|
||||||
float2 uv;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct HSOut
|
|
||||||
{
|
|
||||||
float4 pos;
|
|
||||||
float2 uv;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct HSConstantOut
|
|
||||||
{
|
|
||||||
spvUnsafeArray<float, 3> EdgeTess;
|
|
||||||
float InsideTess;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct VertexOutput_1
|
|
||||||
{
|
|
||||||
float2 uv;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct HSOut_1
|
|
||||||
{
|
|
||||||
float2 uv;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct main0_out
|
|
||||||
{
|
|
||||||
HSOut_1 _entryPointOutput;
|
|
||||||
float4 gl_Position;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct main0_in
|
|
||||||
{
|
|
||||||
float2 VertexOutput_uv [[attribute(0)]];
|
|
||||||
float4 gl_Position [[attribute(1)]];
|
|
||||||
};
|
|
||||||
|
|
||||||
static inline __attribute__((always_inline))
|
|
||||||
HSOut _hs_main(thread const spvUnsafeArray<VertexOutput, 3> (&p), thread const uint& i)
|
|
||||||
{
|
|
||||||
HSOut _output;
|
|
||||||
_output.pos = p[i].pos;
|
|
||||||
_output.uv = p[i].uv;
|
|
||||||
return _output;
|
|
||||||
}
|
|
||||||
|
|
||||||
static inline __attribute__((always_inline))
|
|
||||||
HSConstantOut PatchHS(thread const spvUnsafeArray<VertexOutput, 3> (&_patch))
|
|
||||||
{
|
|
||||||
HSConstantOut _output;
|
|
||||||
_output.EdgeTess[0] = (float2(1.0) + _patch[0].uv).x;
|
|
||||||
_output.EdgeTess[1] = (float2(1.0) + _patch[0].uv).x;
|
|
||||||
_output.EdgeTess[2] = (float2(1.0) + _patch[0].uv).x;
|
|
||||||
_output.InsideTess = (float2(1.0) + _patch[0].uv).x;
|
|
||||||
return _output;
|
|
||||||
}
|
|
||||||
|
|
||||||
kernel void main0(main0_in in [[stage_in]], uint gl_InvocationID [[thread_index_in_threadgroup]], uint gl_PrimitiveID [[threadgroup_position_in_grid]], device main0_out* spvOut [[buffer(28)]], constant uint* spvIndirectParams [[buffer(29)]], device MTLTriangleTessellationFactorsHalf* spvTessLevel [[buffer(26)]], threadgroup main0_in* gl_in [[threadgroup(0)]])
|
|
||||||
{
|
|
||||||
device main0_out* gl_out = &spvOut[gl_PrimitiveID * 3];
|
|
||||||
if (gl_InvocationID < spvIndirectParams[0])
|
|
||||||
gl_in[gl_InvocationID] = in;
|
|
||||||
threadgroup_barrier(mem_flags::mem_threadgroup);
|
|
||||||
if (gl_InvocationID >= 3)
|
|
||||||
return;
|
|
||||||
spvUnsafeArray<VertexOutput, 3> p;
|
|
||||||
p[0].pos = gl_in[0].gl_Position;
|
|
||||||
p[0].uv = gl_in[0].VertexOutput_uv;
|
|
||||||
p[1].pos = gl_in[1].gl_Position;
|
|
||||||
p[1].uv = gl_in[1].VertexOutput_uv;
|
|
||||||
p[2].pos = gl_in[2].gl_Position;
|
|
||||||
p[2].uv = gl_in[2].VertexOutput_uv;
|
|
||||||
uint i = gl_InvocationID;
|
|
||||||
spvUnsafeArray<VertexOutput, 3> param;
|
|
||||||
param = p;
|
|
||||||
uint param_1 = i;
|
|
||||||
HSOut flattenTemp = _hs_main(param, param_1);
|
|
||||||
gl_out[gl_InvocationID].gl_Position = flattenTemp.pos;
|
|
||||||
gl_out[gl_InvocationID]._entryPointOutput.uv = flattenTemp.uv;
|
|
||||||
threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup);
|
|
||||||
if (int(gl_InvocationID) == 0)
|
|
||||||
{
|
|
||||||
spvUnsafeArray<VertexOutput, 3> param_2;
|
|
||||||
param_2 = p;
|
|
||||||
HSConstantOut _patchConstantResult = PatchHS(param_2);
|
|
||||||
spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0] = half(_patchConstantResult.EdgeTess[0]);
|
|
||||||
spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[1] = half(_patchConstantResult.EdgeTess[1]);
|
|
||||||
spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[2] = half(_patchConstantResult.EdgeTess[2]);
|
|
||||||
spvTessLevel[gl_PrimitiveID].insideTessellationFactor = half(_patchConstantResult.InsideTess);
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,90 +0,0 @@
|
|||||||
#include <metal_stdlib>
|
|
||||||
#include <simd/simd.h>
|
|
||||||
|
|
||||||
using namespace metal;
|
|
||||||
|
|
||||||
struct Baz
|
|
||||||
{
|
|
||||||
int e;
|
|
||||||
int f;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct Foo
|
|
||||||
{
|
|
||||||
int a;
|
|
||||||
int b;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct Bar
|
|
||||||
{
|
|
||||||
int c;
|
|
||||||
int d;
|
|
||||||
};
|
|
||||||
|
|
||||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(3u, 3u, 2u);
|
|
||||||
|
|
||||||
struct spvDescriptorSetBuffer0
|
|
||||||
{
|
|
||||||
constant Foo* m_34 [[id(0)]];
|
|
||||||
constant Bar* m_40 [[id(1)]];
|
|
||||||
};
|
|
||||||
|
|
||||||
struct spvDescriptorSetBuffer1
|
|
||||||
{
|
|
||||||
device Baz* baz [[id(0)]][3][3][2];
|
|
||||||
};
|
|
||||||
|
|
||||||
kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], constant spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], constant uint* spvDynamicOffsets [[buffer(23)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
|
||||||
{
|
|
||||||
constant auto& _34 = *(constant Foo* )((constant char* )spvDescriptorSet0.m_34 + spvDynamicOffsets[0]);
|
|
||||||
device Baz* baz[3][3][2] =
|
|
||||||
{
|
|
||||||
{
|
|
||||||
{
|
|
||||||
(device Baz* )((device char* )spvDescriptorSet1.baz[0][0][0] + spvDynamicOffsets[1]),
|
|
||||||
(device Baz* )((device char* )spvDescriptorSet1.baz[0][0][1] + spvDynamicOffsets[2]),
|
|
||||||
},
|
|
||||||
{
|
|
||||||
(device Baz* )((device char* )spvDescriptorSet1.baz[0][1][0] + spvDynamicOffsets[3]),
|
|
||||||
(device Baz* )((device char* )spvDescriptorSet1.baz[0][1][1] + spvDynamicOffsets[4]),
|
|
||||||
},
|
|
||||||
{
|
|
||||||
(device Baz* )((device char* )spvDescriptorSet1.baz[0][2][0] + spvDynamicOffsets[5]),
|
|
||||||
(device Baz* )((device char* )spvDescriptorSet1.baz[0][2][1] + spvDynamicOffsets[6]),
|
|
||||||
},
|
|
||||||
},
|
|
||||||
{
|
|
||||||
{
|
|
||||||
(device Baz* )((device char* )spvDescriptorSet1.baz[1][0][0] + spvDynamicOffsets[7]),
|
|
||||||
(device Baz* )((device char* )spvDescriptorSet1.baz[1][0][1] + spvDynamicOffsets[8]),
|
|
||||||
},
|
|
||||||
{
|
|
||||||
(device Baz* )((device char* )spvDescriptorSet1.baz[1][1][0] + spvDynamicOffsets[9]),
|
|
||||||
(device Baz* )((device char* )spvDescriptorSet1.baz[1][1][1] + spvDynamicOffsets[10]),
|
|
||||||
},
|
|
||||||
{
|
|
||||||
(device Baz* )((device char* )spvDescriptorSet1.baz[1][2][0] + spvDynamicOffsets[11]),
|
|
||||||
(device Baz* )((device char* )spvDescriptorSet1.baz[1][2][1] + spvDynamicOffsets[12]),
|
|
||||||
},
|
|
||||||
},
|
|
||||||
{
|
|
||||||
{
|
|
||||||
(device Baz* )((device char* )spvDescriptorSet1.baz[2][0][0] + spvDynamicOffsets[13]),
|
|
||||||
(device Baz* )((device char* )spvDescriptorSet1.baz[2][0][1] + spvDynamicOffsets[14]),
|
|
||||||
},
|
|
||||||
{
|
|
||||||
(device Baz* )((device char* )spvDescriptorSet1.baz[2][1][0] + spvDynamicOffsets[15]),
|
|
||||||
(device Baz* )((device char* )spvDescriptorSet1.baz[2][1][1] + spvDynamicOffsets[16]),
|
|
||||||
},
|
|
||||||
{
|
|
||||||
(device Baz* )((device char* )spvDescriptorSet1.baz[2][2][0] + spvDynamicOffsets[17]),
|
|
||||||
(device Baz* )((device char* )spvDescriptorSet1.baz[2][2][1] + spvDynamicOffsets[18]),
|
|
||||||
},
|
|
||||||
},
|
|
||||||
};
|
|
||||||
|
|
||||||
uint3 coords = gl_GlobalInvocationID;
|
|
||||||
baz[coords.x][coords.y][coords.z]->e = _34.a + (*spvDescriptorSet0.m_40).c;
|
|
||||||
baz[coords.x][coords.y][coords.z]->f = _34.b * (*spvDescriptorSet0.m_40).d;
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,26 +0,0 @@
|
|||||||
#include <metal_stdlib>
|
|
||||||
#include <simd/simd.h>
|
|
||||||
|
|
||||||
using namespace metal;
|
|
||||||
|
|
||||||
struct SSBO0
|
|
||||||
{
|
|
||||||
short4 inputs[1];
|
|
||||||
};
|
|
||||||
|
|
||||||
struct SSBO1
|
|
||||||
{
|
|
||||||
int4 outputs[1];
|
|
||||||
};
|
|
||||||
|
|
||||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
|
||||||
|
|
||||||
kernel void main0(device SSBO0& _25 [[buffer(0)]], device SSBO1& _39 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
|
||||||
{
|
|
||||||
uint ident = gl_GlobalInvocationID.x;
|
|
||||||
half2 a = as_type<half2>(_25.inputs[ident].xy);
|
|
||||||
_39.outputs[ident].x = int(as_type<uint>(a + half2(half(1.0))));
|
|
||||||
_39.outputs[ident].y = as_type<int>(_25.inputs[ident].zw);
|
|
||||||
_39.outputs[ident].z = int(as_type<uint>(ushort2(_25.inputs[ident].xy)));
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,31 +0,0 @@
|
|||||||
#include <metal_stdlib>
|
|
||||||
#include <simd/simd.h>
|
|
||||||
|
|
||||||
using namespace metal;
|
|
||||||
|
|
||||||
struct SSBO1
|
|
||||||
{
|
|
||||||
short4 outputs[1];
|
|
||||||
};
|
|
||||||
|
|
||||||
struct SSBO0
|
|
||||||
{
|
|
||||||
int4 inputs[1];
|
|
||||||
};
|
|
||||||
|
|
||||||
struct UBO
|
|
||||||
{
|
|
||||||
half4 const0;
|
|
||||||
};
|
|
||||||
|
|
||||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
|
||||||
|
|
||||||
kernel void main0(device SSBO1& _21 [[buffer(0)]], device SSBO0& _29 [[buffer(1)]], constant UBO& _40 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
|
||||||
{
|
|
||||||
uint ident = gl_GlobalInvocationID.x;
|
|
||||||
short2 _47 = as_type<short2>(_29.inputs[ident].x) + as_type<short2>(_40.const0.xy);
|
|
||||||
_21.outputs[ident] = short4(_47.x, _47.y, _21.outputs[ident].z, _21.outputs[ident].w);
|
|
||||||
short2 _66 = short2(as_type<ushort2>(uint(_29.inputs[ident].y)) - as_type<ushort2>(_40.const0.zw));
|
|
||||||
_21.outputs[ident] = short4(_21.outputs[ident].x, _21.outputs[ident].y, _66.x, _66.y);
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,106 +0,0 @@
|
|||||||
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
|
||||||
#pragma clang diagnostic ignored "-Wmissing-braces"
|
|
||||||
|
|
||||||
#include <metal_stdlib>
|
|
||||||
#include <simd/simd.h>
|
|
||||||
|
|
||||||
using namespace metal;
|
|
||||||
|
|
||||||
template<typename T, size_t Num>
|
|
||||||
struct spvUnsafeArray
|
|
||||||
{
|
|
||||||
T elements[Num ? Num : 1];
|
|
||||||
|
|
||||||
thread T& operator [] (size_t pos) thread
|
|
||||||
{
|
|
||||||
return elements[pos];
|
|
||||||
}
|
|
||||||
constexpr const thread T& operator [] (size_t pos) const thread
|
|
||||||
{
|
|
||||||
return elements[pos];
|
|
||||||
}
|
|
||||||
|
|
||||||
device T& operator [] (size_t pos) device
|
|
||||||
{
|
|
||||||
return elements[pos];
|
|
||||||
}
|
|
||||||
constexpr const device T& operator [] (size_t pos) const device
|
|
||||||
{
|
|
||||||
return elements[pos];
|
|
||||||
}
|
|
||||||
|
|
||||||
constexpr const constant T& operator [] (size_t pos) const constant
|
|
||||||
{
|
|
||||||
return elements[pos];
|
|
||||||
}
|
|
||||||
|
|
||||||
threadgroup T& operator [] (size_t pos) threadgroup
|
|
||||||
{
|
|
||||||
return elements[pos];
|
|
||||||
}
|
|
||||||
constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
|
|
||||||
{
|
|
||||||
return elements[pos];
|
|
||||||
}
|
|
||||||
};
|
|
||||||
|
|
||||||
struct M0
|
|
||||||
{
|
|
||||||
long v;
|
|
||||||
spvUnsafeArray<long2, 2> b;
|
|
||||||
ulong c;
|
|
||||||
spvUnsafeArray<ulong, 5> d;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct SSBO0_Type
|
|
||||||
{
|
|
||||||
long4 a;
|
|
||||||
M0 m0;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct SSBO1_Type
|
|
||||||
{
|
|
||||||
ulong4 b;
|
|
||||||
M0 m0;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct SSBO2_Type
|
|
||||||
{
|
|
||||||
spvUnsafeArray<long, 4> a;
|
|
||||||
spvUnsafeArray<long2, 4> b;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct SSBO3_Type
|
|
||||||
{
|
|
||||||
spvUnsafeArray<long, 4> a;
|
|
||||||
spvUnsafeArray<long2, 4> b;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct SSBO
|
|
||||||
{
|
|
||||||
int s32;
|
|
||||||
uint u32;
|
|
||||||
};
|
|
||||||
|
|
||||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
|
||||||
|
|
||||||
kernel void main0(device SSBO& _96 [[buffer(0)]])
|
|
||||||
{
|
|
||||||
SSBO0_Type ssbo_0;
|
|
||||||
ssbo_0.a += long4(10l, 20l, 30l, 40l);
|
|
||||||
SSBO1_Type ssbo_1;
|
|
||||||
ssbo_1.b += ulong4(999999999999999999ul, 8888888888888888ul, 77777777777777777ul, 6666666666666666ul);
|
|
||||||
ssbo_0.a += long4(20l);
|
|
||||||
ssbo_0.a = abs(ssbo_0.a + long4(ssbo_1.b));
|
|
||||||
ssbo_0.a += long4(1l);
|
|
||||||
ssbo_1.b += ulong4(long4(1l));
|
|
||||||
ssbo_0.a -= long4(1l);
|
|
||||||
ssbo_1.b -= ulong4(long4(1l));
|
|
||||||
SSBO2_Type ssbo_2;
|
|
||||||
ssbo_2.a[0] += 1l;
|
|
||||||
SSBO3_Type ssbo_3;
|
|
||||||
ssbo_3.a[0] += 2l;
|
|
||||||
_96.s32 = int(uint(((ulong(ssbo_0.a.x) + ssbo_1.b.y) + ulong(ssbo_2.a[1])) + ulong(ssbo_3.a[2])));
|
|
||||||
_96.u32 = uint(((ulong(ssbo_0.a.y) + ssbo_1.b.z) + ulong(ssbo_2.a[0])) + ulong(ssbo_3.a[1]));
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,21 +0,0 @@
|
|||||||
#include <metal_stdlib>
|
|
||||||
#include <simd/simd.h>
|
|
||||||
|
|
||||||
using namespace metal;
|
|
||||||
|
|
||||||
struct main0_out
|
|
||||||
{
|
|
||||||
half foo [[color(0)]];
|
|
||||||
short bar [[color(1)]];
|
|
||||||
ushort baz [[color(2)]];
|
|
||||||
};
|
|
||||||
|
|
||||||
fragment main0_out main0()
|
|
||||||
{
|
|
||||||
main0_out out = {};
|
|
||||||
out.foo = half(1.0);
|
|
||||||
out.bar = short(2);
|
|
||||||
out.baz = ushort(3);
|
|
||||||
return out;
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,185 +0,0 @@
|
|||||||
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
|
||||||
|
|
||||||
#include <metal_stdlib>
|
|
||||||
#include <simd/simd.h>
|
|
||||||
|
|
||||||
using namespace metal;
|
|
||||||
|
|
||||||
struct ResType
|
|
||||||
{
|
|
||||||
half4 _m0;
|
|
||||||
int4 _m1;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct main0_in
|
|
||||||
{
|
|
||||||
half v1 [[user(locn0)]];
|
|
||||||
half2 v2 [[user(locn1)]];
|
|
||||||
half3 v3 [[user(locn2)]];
|
|
||||||
half4 v4 [[user(locn3)]];
|
|
||||||
};
|
|
||||||
|
|
||||||
// Implementation of the GLSL mod() function, which is slightly different than Metal fmod()
|
|
||||||
template<typename Tx, typename Ty>
|
|
||||||
inline Tx mod(Tx x, Ty y)
|
|
||||||
{
|
|
||||||
return x - y * floor(x / y);
|
|
||||||
}
|
|
||||||
|
|
||||||
// Implementation of the GLSL radians() function
|
|
||||||
template<typename T>
|
|
||||||
inline T radians(T d)
|
|
||||||
{
|
|
||||||
return d * T(0.01745329251);
|
|
||||||
}
|
|
||||||
|
|
||||||
// Implementation of the GLSL degrees() function
|
|
||||||
template<typename T>
|
|
||||||
inline T degrees(T r)
|
|
||||||
{
|
|
||||||
return r * T(57.2957795131);
|
|
||||||
}
|
|
||||||
|
|
||||||
static inline __attribute__((always_inline))
|
|
||||||
half2x2 test_mat2(thread const half2& a, thread const half2& b, thread const half2& c, thread const half2& d)
|
|
||||||
{
|
|
||||||
return half2x2(half2(a), half2(b)) * half2x2(half2(c), half2(d));
|
|
||||||
}
|
|
||||||
|
|
||||||
static inline __attribute__((always_inline))
|
|
||||||
half3x3 test_mat3(thread const half3& a, thread const half3& b, thread const half3& c, thread const half3& d, thread const half3& e, thread const half3& f)
|
|
||||||
{
|
|
||||||
return half3x3(half3(a), half3(b), half3(c)) * half3x3(half3(d), half3(e), half3(f));
|
|
||||||
}
|
|
||||||
|
|
||||||
static inline __attribute__((always_inline))
|
|
||||||
void test_constants()
|
|
||||||
{
|
|
||||||
half a = half(1.0);
|
|
||||||
half b = half(1.5);
|
|
||||||
half c = half(-1.5);
|
|
||||||
half d = half(0.0 / 0.0);
|
|
||||||
half e = half(1.0 / 0.0);
|
|
||||||
half f = half(-1.0 / 0.0);
|
|
||||||
half g = half(1014.0);
|
|
||||||
half h = half(9.5367431640625e-07);
|
|
||||||
}
|
|
||||||
|
|
||||||
static inline __attribute__((always_inline))
|
|
||||||
half test_result()
|
|
||||||
{
|
|
||||||
return half(1.0);
|
|
||||||
}
|
|
||||||
|
|
||||||
static inline __attribute__((always_inline))
|
|
||||||
void test_conversions()
|
|
||||||
{
|
|
||||||
half one = test_result();
|
|
||||||
int a = int(one);
|
|
||||||
uint b = uint(one);
|
|
||||||
bool c = one != half(0.0);
|
|
||||||
float d = float(one);
|
|
||||||
half a2 = half(a);
|
|
||||||
half b2 = half(b);
|
|
||||||
half c2 = half(c);
|
|
||||||
half d2 = half(d);
|
|
||||||
}
|
|
||||||
|
|
||||||
static inline __attribute__((always_inline))
|
|
||||||
void test_builtins(thread half4& v4, thread half3& v3, thread half& v1)
|
|
||||||
{
|
|
||||||
half4 res = radians(v4);
|
|
||||||
res = degrees(v4);
|
|
||||||
res = sin(v4);
|
|
||||||
res = cos(v4);
|
|
||||||
res = tan(v4);
|
|
||||||
res = asin(v4);
|
|
||||||
res = atan2(v4, v3.xyzz);
|
|
||||||
res = atan(v4);
|
|
||||||
res = sinh(v4);
|
|
||||||
res = cosh(v4);
|
|
||||||
res = tanh(v4);
|
|
||||||
res = asinh(v4);
|
|
||||||
res = acosh(v4);
|
|
||||||
res = atanh(v4);
|
|
||||||
res = pow(v4, v4);
|
|
||||||
res = exp(v4);
|
|
||||||
res = log(v4);
|
|
||||||
res = exp2(v4);
|
|
||||||
res = log2(v4);
|
|
||||||
res = sqrt(v4);
|
|
||||||
res = rsqrt(v4);
|
|
||||||
res = abs(v4);
|
|
||||||
res = sign(v4);
|
|
||||||
res = floor(v4);
|
|
||||||
res = trunc(v4);
|
|
||||||
res = round(v4);
|
|
||||||
res = rint(v4);
|
|
||||||
res = ceil(v4);
|
|
||||||
res = fract(v4);
|
|
||||||
res = mod(v4, v4);
|
|
||||||
half4 tmp;
|
|
||||||
half4 _223 = modf(v4, tmp);
|
|
||||||
res = _223;
|
|
||||||
res = min(v4, v4);
|
|
||||||
res = max(v4, v4);
|
|
||||||
res = clamp(v4, v4, v4);
|
|
||||||
res = mix(v4, v4, v4);
|
|
||||||
res = select(v4, v4, v4 < v4);
|
|
||||||
res = step(v4, v4);
|
|
||||||
res = smoothstep(v4, v4, v4);
|
|
||||||
bool4 btmp = isnan(v4);
|
|
||||||
btmp = isinf(v4);
|
|
||||||
res = fma(v4, v4, v4);
|
|
||||||
ResType _267;
|
|
||||||
_267._m0 = frexp(v4, _267._m1);
|
|
||||||
int4 itmp = _267._m1;
|
|
||||||
res = _267._m0;
|
|
||||||
res = ldexp(res, itmp);
|
|
||||||
uint pack0 = as_type<uint>(v4.xy);
|
|
||||||
uint pack1 = as_type<uint>(v4.zw);
|
|
||||||
res = half4(as_type<half2>(pack0), as_type<half2>(pack1));
|
|
||||||
half t0 = length(v4);
|
|
||||||
t0 = distance(v4, v4);
|
|
||||||
t0 = dot(v4, v4);
|
|
||||||
half3 res3 = cross(v3, v3);
|
|
||||||
res = normalize(v4);
|
|
||||||
res = faceforward(v4, v4, v4);
|
|
||||||
res = reflect(v4, v4);
|
|
||||||
res = refract(v4, v4, v1);
|
|
||||||
btmp = v4 < v4;
|
|
||||||
btmp = v4 <= v4;
|
|
||||||
btmp = v4 > v4;
|
|
||||||
btmp = v4 >= v4;
|
|
||||||
btmp = v4 == v4;
|
|
||||||
btmp = v4 != v4;
|
|
||||||
res = dfdx(v4);
|
|
||||||
res = dfdy(v4);
|
|
||||||
res = dfdx(v4);
|
|
||||||
res = dfdy(v4);
|
|
||||||
res = dfdx(v4);
|
|
||||||
res = dfdy(v4);
|
|
||||||
res = fwidth(v4);
|
|
||||||
res = fwidth(v4);
|
|
||||||
res = fwidth(v4);
|
|
||||||
}
|
|
||||||
|
|
||||||
fragment void main0(main0_in in [[stage_in]])
|
|
||||||
{
|
|
||||||
half2 param = in.v2;
|
|
||||||
half2 param_1 = in.v2;
|
|
||||||
half2 param_2 = in.v3.xy;
|
|
||||||
half2 param_3 = in.v3.xy;
|
|
||||||
half2x2 m0 = test_mat2(param, param_1, param_2, param_3);
|
|
||||||
half3 param_4 = in.v3;
|
|
||||||
half3 param_5 = in.v3;
|
|
||||||
half3 param_6 = in.v3;
|
|
||||||
half3 param_7 = in.v4.xyz;
|
|
||||||
half3 param_8 = in.v4.xyz;
|
|
||||||
half3 param_9 = in.v4.yzw;
|
|
||||||
half3x3 m1 = test_mat3(param_4, param_5, param_6, param_7, param_8, param_9);
|
|
||||||
test_constants();
|
|
||||||
test_conversions();
|
|
||||||
test_builtins(in.v4, in.v3, in.v1);
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,57 +0,0 @@
|
|||||||
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
|
||||||
|
|
||||||
#include <metal_stdlib>
|
|
||||||
#include <simd/simd.h>
|
|
||||||
|
|
||||||
using namespace metal;
|
|
||||||
|
|
||||||
struct main0_out
|
|
||||||
{
|
|
||||||
float FragColor [[color(0)]];
|
|
||||||
};
|
|
||||||
|
|
||||||
struct main0_in
|
|
||||||
{
|
|
||||||
float3 vUV [[user(locn0)]];
|
|
||||||
};
|
|
||||||
|
|
||||||
static inline __attribute__((always_inline))
|
|
||||||
float Samp(thread const float3& uv, thread depth2d<float> uTex, thread sampler uSamp)
|
|
||||||
{
|
|
||||||
return uTex.sample_compare(uSamp, uv.xy, uv.z);
|
|
||||||
}
|
|
||||||
|
|
||||||
static inline __attribute__((always_inline))
|
|
||||||
float Samp2(thread const float3& uv, thread depth2d<float> uSampler, thread const sampler uSamplerSmplr, thread float3& vUV)
|
|
||||||
{
|
|
||||||
return uSampler.sample_compare(uSamplerSmplr, vUV.xy, vUV.z);
|
|
||||||
}
|
|
||||||
|
|
||||||
static inline __attribute__((always_inline))
|
|
||||||
float Samp3(thread const depth2d<float> uT, thread const sampler uS, thread const float3& uv, thread float3& vUV)
|
|
||||||
{
|
|
||||||
return uT.sample_compare(uS, vUV.xy, vUV.z);
|
|
||||||
}
|
|
||||||
|
|
||||||
static inline __attribute__((always_inline))
|
|
||||||
float Samp4(thread const depth2d<float> uS, thread const sampler uSSmplr, thread const float3& uv, thread float3& vUV)
|
|
||||||
{
|
|
||||||
return uS.sample_compare(uSSmplr, vUV.xy, vUV.z);
|
|
||||||
}
|
|
||||||
|
|
||||||
fragment main0_out main0(main0_in in [[stage_in]], depth2d<float> uTex [[texture(0)]], depth2d<float> uSampler [[texture(1)]], sampler uSamp [[sampler(0)]], sampler uSamplerSmplr [[sampler(1)]])
|
|
||||||
{
|
|
||||||
main0_out out = {};
|
|
||||||
out.FragColor = uSampler.sample_compare(uSamplerSmplr, in.vUV.xy, in.vUV.z);
|
|
||||||
out.FragColor += uTex.sample_compare(uSamp, in.vUV.xy, in.vUV.z);
|
|
||||||
float3 param = in.vUV;
|
|
||||||
out.FragColor += Samp(param, uTex, uSamp);
|
|
||||||
float3 param_1 = in.vUV;
|
|
||||||
out.FragColor += Samp2(param_1, uSampler, uSamplerSmplr, in.vUV);
|
|
||||||
float3 param_2 = in.vUV;
|
|
||||||
out.FragColor += Samp3(uTex, uSamp, param_2, in.vUV);
|
|
||||||
float3 param_3 = in.vUV;
|
|
||||||
out.FragColor += Samp4(uSampler, uSamplerSmplr, param_3, in.vUV);
|
|
||||||
return out;
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,24 +0,0 @@
|
|||||||
#include <metal_stdlib>
|
|
||||||
#include <simd/simd.h>
|
|
||||||
|
|
||||||
using namespace metal;
|
|
||||||
|
|
||||||
struct main0_out
|
|
||||||
{
|
|
||||||
float4 gl_Position [[position]];
|
|
||||||
uint gl_Layer [[render_target_array_index]];
|
|
||||||
};
|
|
||||||
|
|
||||||
struct main0_in
|
|
||||||
{
|
|
||||||
float4 coord [[attribute(0)]];
|
|
||||||
};
|
|
||||||
|
|
||||||
vertex main0_out main0(main0_in in [[stage_in]])
|
|
||||||
{
|
|
||||||
main0_out out = {};
|
|
||||||
out.gl_Position = in.coord;
|
|
||||||
out.gl_Layer = uint(int(in.coord.z));
|
|
||||||
return out;
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,24 +0,0 @@
|
|||||||
#include <metal_stdlib>
|
|
||||||
#include <simd/simd.h>
|
|
||||||
|
|
||||||
using namespace metal;
|
|
||||||
|
|
||||||
struct main0_out
|
|
||||||
{
|
|
||||||
float4 gl_Position [[position]];
|
|
||||||
uint gl_ViewportIndex [[viewport_array_index]];
|
|
||||||
};
|
|
||||||
|
|
||||||
struct main0_in
|
|
||||||
{
|
|
||||||
float4 coord [[attribute(0)]];
|
|
||||||
};
|
|
||||||
|
|
||||||
vertex main0_out main0(main0_in in [[stage_in]])
|
|
||||||
{
|
|
||||||
main0_out out = {};
|
|
||||||
out.gl_Position = in.coord;
|
|
||||||
out.gl_ViewportIndex = uint(int(in.coord.z));
|
|
||||||
return out;
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,150 +0,0 @@
|
|||||||
#include <metal_stdlib>
|
|
||||||
#include <simd/simd.h>
|
|
||||||
|
|
||||||
using namespace metal;
|
|
||||||
|
|
||||||
typedef packed_float2 packed_float2x2[2];
|
|
||||||
typedef packed_float3 packed_float2x3[2];
|
|
||||||
typedef packed_float3 packed_rm_float3x2[2];
|
|
||||||
|
|
||||||
struct S0
|
|
||||||
{
|
|
||||||
packed_float2 a[1];
|
|
||||||
float b;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct S1
|
|
||||||
{
|
|
||||||
packed_float3 a;
|
|
||||||
float b;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct S2
|
|
||||||
{
|
|
||||||
packed_float3 a[1];
|
|
||||||
float b;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct S3
|
|
||||||
{
|
|
||||||
packed_float2 a;
|
|
||||||
float b;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct Content
|
|
||||||
{
|
|
||||||
S0 m0s[1];
|
|
||||||
S1 m1s[1];
|
|
||||||
S2 m2s[1];
|
|
||||||
S0 m0;
|
|
||||||
S1 m1;
|
|
||||||
S2 m2;
|
|
||||||
S3 m3;
|
|
||||||
float m4;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct SSBO1
|
|
||||||
{
|
|
||||||
Content content;
|
|
||||||
Content content1[2];
|
|
||||||
Content content2;
|
|
||||||
float2x2 m0;
|
|
||||||
float2x2 m1;
|
|
||||||
packed_float2x3 m2[4];
|
|
||||||
float3x2 m3;
|
|
||||||
float2x2 m4;
|
|
||||||
float2x2 m5[9];
|
|
||||||
float3x2 m6[4][2];
|
|
||||||
packed_rm_float3x2 m7;
|
|
||||||
float array[1];
|
|
||||||
};
|
|
||||||
|
|
||||||
struct S0_1
|
|
||||||
{
|
|
||||||
float4 a[1];
|
|
||||||
float b;
|
|
||||||
char _m0_final_padding[12];
|
|
||||||
};
|
|
||||||
|
|
||||||
struct S1_1
|
|
||||||
{
|
|
||||||
packed_float3 a;
|
|
||||||
float b;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct S2_1
|
|
||||||
{
|
|
||||||
float3 a[1];
|
|
||||||
float b;
|
|
||||||
char _m0_final_padding[12];
|
|
||||||
};
|
|
||||||
|
|
||||||
struct S3_1
|
|
||||||
{
|
|
||||||
float2 a;
|
|
||||||
float b;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct Content_1
|
|
||||||
{
|
|
||||||
S0_1 m0s[1];
|
|
||||||
S1_1 m1s[1];
|
|
||||||
S2_1 m2s[1];
|
|
||||||
S0_1 m0;
|
|
||||||
S1_1 m1;
|
|
||||||
S2_1 m2;
|
|
||||||
S3_1 m3;
|
|
||||||
float m4;
|
|
||||||
char _m0_final_padding[12];
|
|
||||||
};
|
|
||||||
|
|
||||||
struct SSBO0
|
|
||||||
{
|
|
||||||
Content_1 content;
|
|
||||||
Content_1 content1[2];
|
|
||||||
Content_1 content2;
|
|
||||||
float2x4 m0;
|
|
||||||
float2x4 m1;
|
|
||||||
float2x3 m2[4];
|
|
||||||
float3x4 m3;
|
|
||||||
float2x4 m4;
|
|
||||||
float2x4 m5[9];
|
|
||||||
float3x4 m6[4][2];
|
|
||||||
float2x3 m7;
|
|
||||||
float4 array[1];
|
|
||||||
};
|
|
||||||
|
|
||||||
struct SSBO2
|
|
||||||
{
|
|
||||||
float m0;
|
|
||||||
packed_float2x2 m1;
|
|
||||||
packed_rm_float3x2 m2;
|
|
||||||
};
|
|
||||||
|
|
||||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
|
||||||
|
|
||||||
kernel void main0(device SSBO1& ssbo_scalar [[buffer(0)]], device SSBO0& ssbo_140 [[buffer(1)]], device SSBO2& ssbo_scalar2 [[buffer(2)]])
|
|
||||||
{
|
|
||||||
ssbo_scalar.content.m0s[0].a[0] = ssbo_140.content.m0s[0].a[0].xy;
|
|
||||||
ssbo_scalar.content.m0s[0].b = ssbo_140.content.m0s[0].b;
|
|
||||||
ssbo_scalar.content.m1s[0].a = float3(ssbo_140.content.m1s[0].a);
|
|
||||||
ssbo_scalar.content.m1s[0].b = ssbo_140.content.m1s[0].b;
|
|
||||||
ssbo_scalar.content.m2s[0].a[0] = ssbo_140.content.m2s[0].a[0];
|
|
||||||
ssbo_scalar.content.m2s[0].b = ssbo_140.content.m2s[0].b;
|
|
||||||
ssbo_scalar.content.m0.a[0] = ssbo_140.content.m0.a[0].xy;
|
|
||||||
ssbo_scalar.content.m0.b = ssbo_140.content.m0.b;
|
|
||||||
ssbo_scalar.content.m1.a = float3(ssbo_140.content.m1.a);
|
|
||||||
ssbo_scalar.content.m1.b = ssbo_140.content.m1.b;
|
|
||||||
ssbo_scalar.content.m2.a[0] = ssbo_140.content.m2.a[0];
|
|
||||||
ssbo_scalar.content.m2.b = ssbo_140.content.m2.b;
|
|
||||||
ssbo_scalar.content.m3.a = ssbo_140.content.m3.a;
|
|
||||||
ssbo_scalar.content.m3.b = ssbo_140.content.m3.b;
|
|
||||||
ssbo_scalar.content.m4 = ssbo_140.content.m4;
|
|
||||||
ssbo_scalar.content.m1.a = float2x3(float3(ssbo_scalar.m2[1][0]), float3(ssbo_scalar.m2[1][1])) * float2(ssbo_scalar.content.m0.a[0]);
|
|
||||||
ssbo_scalar.m0 = float2x2(float2(ssbo_scalar2.m1[0]), float2(ssbo_scalar2.m1[1]));
|
|
||||||
ssbo_scalar2.m1[0] = float2(ssbo_scalar.m4[0][0], ssbo_scalar.m4[1][0]);
|
|
||||||
ssbo_scalar2.m1[1] = float2(ssbo_scalar.m4[0][1], ssbo_scalar.m4[1][1]);
|
|
||||||
ssbo_scalar2.m2[0] = float3(ssbo_scalar.m3[0][0], ssbo_scalar.m3[1][0], ssbo_scalar.m3[2][0]);
|
|
||||||
ssbo_scalar2.m2[1] = float3(ssbo_scalar.m3[0][1], ssbo_scalar.m3[1][1], ssbo_scalar.m3[2][1]);
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,148 +0,0 @@
|
|||||||
#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);
|
|
||||||
|
|
||||||
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)
|
|
||||||
{
|
|
||||||
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)
|
|
||||||
{
|
|
||||||
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)
|
|
||||||
{
|
|
||||||
return popcount(ballot.x) + popcount(ballot.y) + popcount(ballot.z) + popcount(ballot.w);
|
|
||||||
}
|
|
||||||
|
|
||||||
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);
|
|
||||||
}
|
|
||||||
|
|
||||||
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);
|
|
||||||
}
|
|
||||||
|
|
||||||
template<typename T>
|
|
||||||
inline bool spvSubgroupAllEqual(T value)
|
|
||||||
{
|
|
||||||
return simd_all(value == simd_broadcast_first(value));
|
|
||||||
}
|
|
||||||
|
|
||||||
template<>
|
|
||||||
inline bool spvSubgroupAllEqual(bool value)
|
|
||||||
{
|
|
||||||
return simd_all(value) || !simd_any(value);
|
|
||||||
}
|
|
||||||
|
|
||||||
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_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 elected = simd_is_first();
|
|
||||||
_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 = simd_broadcast(float4(10.0), 8u);
|
|
||||||
float3 first = simd_broadcast_first(float3(20.0));
|
|
||||||
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 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);
|
|
||||||
bool has_all = simd_all(true);
|
|
||||||
bool has_any = simd_any(true);
|
|
||||||
bool has_equal = spvSubgroupAllEqual(0);
|
|
||||||
has_equal = spvSubgroupAllEqual(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 = 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);
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,33 +0,0 @@
|
|||||||
#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_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);
|
|
||||||
_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);
|
|
||||||
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);
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,11 +0,0 @@
|
|||||||
#include <metal_stdlib>
|
|
||||||
#include <simd/simd.h>
|
|
||||||
|
|
||||||
using namespace metal;
|
|
||||||
|
|
||||||
fragment void main0()
|
|
||||||
{
|
|
||||||
bool _9 = simd_is_helper_thread();
|
|
||||||
bool helper = _9;
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,36 +0,0 @@
|
|||||||
#include <metal_stdlib>
|
|
||||||
#include <simd/simd.h>
|
|
||||||
|
|
||||||
using namespace metal;
|
|
||||||
|
|
||||||
struct UBO
|
|
||||||
{
|
|
||||||
float a[1];
|
|
||||||
float2 b[2];
|
|
||||||
};
|
|
||||||
|
|
||||||
struct UBOEnhancedLayout
|
|
||||||
{
|
|
||||||
float c[1];
|
|
||||||
float2 d[2];
|
|
||||||
char _m2_pad[9976];
|
|
||||||
float e;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct main0_out
|
|
||||||
{
|
|
||||||
float FragColor [[color(0)]];
|
|
||||||
};
|
|
||||||
|
|
||||||
struct main0_in
|
|
||||||
{
|
|
||||||
int vIndex [[user(locn0)]];
|
|
||||||
};
|
|
||||||
|
|
||||||
fragment main0_out main0(main0_in in [[stage_in]], constant UBO& _17 [[buffer(0)]], constant UBOEnhancedLayout& _30 [[buffer(1)]])
|
|
||||||
{
|
|
||||||
main0_out out = {};
|
|
||||||
out.FragColor = (_17.a[in.vIndex] + _30.c[in.vIndex]) + _30.e;
|
|
||||||
return out;
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,143 +0,0 @@
|
|||||||
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
|
||||||
|
|
||||||
#include <metal_stdlib>
|
|
||||||
#include <simd/simd.h>
|
|
||||||
|
|
||||||
using namespace metal;
|
|
||||||
|
|
||||||
struct main0_out
|
|
||||||
{
|
|
||||||
float FragColor [[color(0)]];
|
|
||||||
};
|
|
||||||
|
|
||||||
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)
|
|
||||||
{
|
|
||||||
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)
|
|
||||||
{
|
|
||||||
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)
|
|
||||||
{
|
|
||||||
return popcount(ballot.x) + popcount(ballot.y) + popcount(ballot.z) + popcount(ballot.w);
|
|
||||||
}
|
|
||||||
|
|
||||||
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);
|
|
||||||
}
|
|
||||||
|
|
||||||
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);
|
|
||||||
}
|
|
||||||
|
|
||||||
template<typename T>
|
|
||||||
inline bool spvSubgroupAllEqual(T value)
|
|
||||||
{
|
|
||||||
return simd_all(value == simd_broadcast_first(value));
|
|
||||||
}
|
|
||||||
|
|
||||||
template<>
|
|
||||||
inline bool spvSubgroupAllEqual(bool value)
|
|
||||||
{
|
|
||||||
return simd_all(value) || !simd_any(value);
|
|
||||||
}
|
|
||||||
|
|
||||||
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_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);
|
|
||||||
out.FragColor = float(gl_SubgroupInvocationID);
|
|
||||||
bool elected = simd_is_first();
|
|
||||||
out.FragColor = float4(gl_SubgroupEqMask).x;
|
|
||||||
out.FragColor = float4(gl_SubgroupGeMask).x;
|
|
||||||
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));
|
|
||||||
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 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);
|
|
||||||
bool has_all = simd_all(true);
|
|
||||||
bool has_any = simd_any(true);
|
|
||||||
bool has_equal = spvSubgroupAllEqual(0);
|
|
||||||
has_equal = spvSubgroupAllEqual(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 = 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);
|
|
||||||
return out;
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,11 +0,0 @@
|
|||||||
#version 450
|
|
||||||
#extension GL_AMD_shader_fragment_mask : require
|
|
||||||
|
|
||||||
layout(input_attachment_index = 0, set = 0, binding = 0) uniform subpassInputMS t;
|
|
||||||
|
|
||||||
void main()
|
|
||||||
{
|
|
||||||
vec4 test2 = fragmentFetchAMD(t, 4u);
|
|
||||||
uint testi2 = fragmentMaskFetchAMD(t);
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,15 +0,0 @@
|
|||||||
#version 450
|
|
||||||
#extension GL_AMD_shader_fragment_mask : require
|
|
||||||
#extension GL_AMD_shader_explicit_vertex_parameter : require
|
|
||||||
|
|
||||||
layout(binding = 0) uniform sampler2DMS texture1;
|
|
||||||
|
|
||||||
layout(location = 0) __explicitInterpAMD in vec4 vary;
|
|
||||||
|
|
||||||
void main()
|
|
||||||
{
|
|
||||||
uint testi1 = fragmentMaskFetchAMD(texture1, ivec2(0));
|
|
||||||
vec4 test1 = fragmentFetchAMD(texture1, ivec2(1), 2u);
|
|
||||||
vec4 pos = interpolateAtVertexAMD(vary, 0u);
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,11 +0,0 @@
|
|||||||
#version 450
|
|
||||||
#extension GL_AMD_shader_ballot : require
|
|
||||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
|
||||||
|
|
||||||
void main()
|
|
||||||
{
|
|
||||||
float addInvocations = addInvocationsNonUniformAMD(0.0);
|
|
||||||
int minInvocations = minInvocationsNonUniformAMD(1);
|
|
||||||
uint maxInvocations = uint(maxInvocationsNonUniformAMD(4));
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,28 +0,0 @@
|
|||||||
#version 450
|
|
||||||
|
|
||||||
#ifndef SPIRV_CROSS_CONSTANT_ID_0
|
|
||||||
#define SPIRV_CROSS_CONSTANT_ID_0 1u
|
|
||||||
#endif
|
|
||||||
#ifndef SPIRV_CROSS_CONSTANT_ID_2
|
|
||||||
#define SPIRV_CROSS_CONSTANT_ID_2 3u
|
|
||||||
#endif
|
|
||||||
|
|
||||||
layout(local_size_x = SPIRV_CROSS_CONSTANT_ID_0, local_size_y = 2, local_size_z = SPIRV_CROSS_CONSTANT_ID_2) in;
|
|
||||||
|
|
||||||
layout(binding = 0, std430) buffer _6_8
|
|
||||||
{
|
|
||||||
float _m0[];
|
|
||||||
} _8;
|
|
||||||
|
|
||||||
layout(binding = 1, std430) buffer _6_9
|
|
||||||
{
|
|
||||||
float _m0[];
|
|
||||||
} _9;
|
|
||||||
|
|
||||||
uvec3 _22 = gl_WorkGroupSize;
|
|
||||||
|
|
||||||
void main()
|
|
||||||
{
|
|
||||||
_8._m0[gl_WorkGroupID.x] = _9._m0[gl_WorkGroupID.x] + _8._m0[gl_WorkGroupID.x];
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,18 +0,0 @@
|
|||||||
#version 430
|
|
||||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
|
||||||
|
|
||||||
layout(binding = 0, std430) buffer _8_9
|
|
||||||
{
|
|
||||||
float _m0[];
|
|
||||||
} _9;
|
|
||||||
|
|
||||||
layout(binding = 1, std430) buffer _8_10
|
|
||||||
{
|
|
||||||
float _m0[];
|
|
||||||
} _10;
|
|
||||||
|
|
||||||
void main()
|
|
||||||
{
|
|
||||||
_10._m0[gl_GlobalInvocationID.x] = -_9._m0[gl_GlobalInvocationID.x];
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,41 +0,0 @@
|
|||||||
#version 450
|
|
||||||
layout(triangles) in;
|
|
||||||
layout(max_vertices = 3, triangle_strip) out;
|
|
||||||
|
|
||||||
struct VertexOutput
|
|
||||||
{
|
|
||||||
vec4 pos;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct GeometryOutput
|
|
||||||
{
|
|
||||||
vec4 pos;
|
|
||||||
uint layer;
|
|
||||||
};
|
|
||||||
|
|
||||||
void _main(VertexOutput _input[3], GeometryOutput stream)
|
|
||||||
{
|
|
||||||
GeometryOutput _output;
|
|
||||||
_output.layer = 1u;
|
|
||||||
for (int v = 0; v < 3; v++)
|
|
||||||
{
|
|
||||||
_output.pos = _input[v].pos;
|
|
||||||
gl_Position = _output.pos;
|
|
||||||
gl_Layer = int(_output.layer);
|
|
||||||
EmitVertex();
|
|
||||||
}
|
|
||||||
EndPrimitive();
|
|
||||||
}
|
|
||||||
|
|
||||||
void main()
|
|
||||||
{
|
|
||||||
VertexOutput _input[3];
|
|
||||||
_input[0].pos = gl_in[0].gl_Position;
|
|
||||||
_input[1].pos = gl_in[1].gl_Position;
|
|
||||||
_input[2].pos = gl_in[2].gl_Position;
|
|
||||||
VertexOutput param[3] = _input;
|
|
||||||
GeometryOutput param_1;
|
|
||||||
_main(param, param_1);
|
|
||||||
GeometryOutput stream = param_1;
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,79 +0,0 @@
|
|||||||
#version 450
|
|
||||||
layout(vertices = 3) out;
|
|
||||||
|
|
||||||
struct VertexOutput
|
|
||||||
{
|
|
||||||
vec4 pos;
|
|
||||||
vec2 uv;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct HSOut
|
|
||||||
{
|
|
||||||
vec4 pos;
|
|
||||||
vec2 uv;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct HSConstantOut
|
|
||||||
{
|
|
||||||
float EdgeTess[3];
|
|
||||||
float InsideTess;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct VertexOutput_1
|
|
||||||
{
|
|
||||||
vec2 uv;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct HSOut_1
|
|
||||||
{
|
|
||||||
vec2 uv;
|
|
||||||
};
|
|
||||||
|
|
||||||
layout(location = 0) in VertexOutput_1 p[];
|
|
||||||
layout(location = 0) out HSOut_1 _entryPointOutput[3];
|
|
||||||
|
|
||||||
HSOut _hs_main(VertexOutput p_1[3], uint i)
|
|
||||||
{
|
|
||||||
HSOut _output;
|
|
||||||
_output.pos = p_1[i].pos;
|
|
||||||
_output.uv = p_1[i].uv;
|
|
||||||
return _output;
|
|
||||||
}
|
|
||||||
|
|
||||||
HSConstantOut PatchHS(VertexOutput _patch[3])
|
|
||||||
{
|
|
||||||
HSConstantOut _output;
|
|
||||||
_output.EdgeTess[0] = (vec2(1.0) + _patch[0].uv).x;
|
|
||||||
_output.EdgeTess[1] = (vec2(1.0) + _patch[0].uv).x;
|
|
||||||
_output.EdgeTess[2] = (vec2(1.0) + _patch[0].uv).x;
|
|
||||||
_output.InsideTess = (vec2(1.0) + _patch[0].uv).x;
|
|
||||||
return _output;
|
|
||||||
}
|
|
||||||
|
|
||||||
void main()
|
|
||||||
{
|
|
||||||
VertexOutput p_1[3];
|
|
||||||
p_1[0].pos = gl_in[0].gl_Position;
|
|
||||||
p_1[0].uv = p[0].uv;
|
|
||||||
p_1[1].pos = gl_in[1].gl_Position;
|
|
||||||
p_1[1].uv = p[1].uv;
|
|
||||||
p_1[2].pos = gl_in[2].gl_Position;
|
|
||||||
p_1[2].uv = p[2].uv;
|
|
||||||
uint i = gl_InvocationID;
|
|
||||||
VertexOutput param[3] = p_1;
|
|
||||||
uint param_1 = i;
|
|
||||||
HSOut flattenTemp = _hs_main(param, param_1);
|
|
||||||
gl_out[gl_InvocationID].gl_Position = flattenTemp.pos;
|
|
||||||
_entryPointOutput[gl_InvocationID].uv = flattenTemp.uv;
|
|
||||||
barrier();
|
|
||||||
if (int(gl_InvocationID) == 0)
|
|
||||||
{
|
|
||||||
VertexOutput param_2[3] = p_1;
|
|
||||||
HSConstantOut _patchConstantResult = PatchHS(param_2);
|
|
||||||
gl_TessLevelOuter[0] = _patchConstantResult.EdgeTess[0];
|
|
||||||
gl_TessLevelOuter[1] = _patchConstantResult.EdgeTess[1];
|
|
||||||
gl_TessLevelOuter[2] = _patchConstantResult.EdgeTess[2];
|
|
||||||
gl_TessLevelInner[0] = _patchConstantResult.InsideTess;
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,34 +0,0 @@
|
|||||||
#version 450
|
|
||||||
#if defined(GL_AMD_gpu_shader_half_float)
|
|
||||||
#extension GL_AMD_gpu_shader_half_float : require
|
|
||||||
#elif defined(GL_NV_gpu_shader5)
|
|
||||||
#extension GL_NV_gpu_shader5 : require
|
|
||||||
#else
|
|
||||||
#error No extension available for FP16.
|
|
||||||
#endif
|
|
||||||
#if defined(GL_AMD_gpu_shader_int16)
|
|
||||||
#extension GL_AMD_gpu_shader_int16 : require
|
|
||||||
#else
|
|
||||||
#error No extension available for Int16.
|
|
||||||
#endif
|
|
||||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
|
||||||
|
|
||||||
layout(binding = 0, std430) buffer SSBO0
|
|
||||||
{
|
|
||||||
i16vec4 inputs[];
|
|
||||||
} _25;
|
|
||||||
|
|
||||||
layout(binding = 1, std430) buffer SSBO1
|
|
||||||
{
|
|
||||||
ivec4 outputs[];
|
|
||||||
} _39;
|
|
||||||
|
|
||||||
void main()
|
|
||||||
{
|
|
||||||
uint ident = gl_GlobalInvocationID.x;
|
|
||||||
f16vec2 a = int16BitsToFloat16(_25.inputs[ident].xy);
|
|
||||||
_39.outputs[ident].x = int(packFloat2x16(a + f16vec2(float16_t(1.0))));
|
|
||||||
_39.outputs[ident].y = packInt2x16(_25.inputs[ident].zw);
|
|
||||||
_39.outputs[ident].z = int(packUint2x16(u16vec2(_25.inputs[ident].xy)));
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,39 +0,0 @@
|
|||||||
#version 450
|
|
||||||
#if defined(GL_AMD_gpu_shader_int16)
|
|
||||||
#extension GL_AMD_gpu_shader_int16 : require
|
|
||||||
#else
|
|
||||||
#error No extension available for Int16.
|
|
||||||
#endif
|
|
||||||
#if defined(GL_AMD_gpu_shader_half_float)
|
|
||||||
#extension GL_AMD_gpu_shader_half_float : require
|
|
||||||
#elif defined(GL_NV_gpu_shader5)
|
|
||||||
#extension GL_NV_gpu_shader5 : require
|
|
||||||
#else
|
|
||||||
#error No extension available for FP16.
|
|
||||||
#endif
|
|
||||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
|
||||||
|
|
||||||
layout(binding = 1, std430) buffer SSBO1
|
|
||||||
{
|
|
||||||
i16vec4 outputs[];
|
|
||||||
} _21;
|
|
||||||
|
|
||||||
layout(binding = 0, std430) buffer SSBO0
|
|
||||||
{
|
|
||||||
ivec4 inputs[];
|
|
||||||
} _29;
|
|
||||||
|
|
||||||
layout(binding = 2, std140) uniform UBO
|
|
||||||
{
|
|
||||||
f16vec4 const0;
|
|
||||||
} _40;
|
|
||||||
|
|
||||||
void main()
|
|
||||||
{
|
|
||||||
uint ident = gl_GlobalInvocationID.x;
|
|
||||||
i16vec2 _47 = unpackInt2x16(_29.inputs[ident].x) + float16BitsToInt16(_40.const0.xy);
|
|
||||||
_21.outputs[ident] = i16vec4(_47.x, _47.y, _21.outputs[ident].z, _21.outputs[ident].w);
|
|
||||||
i16vec2 _66 = i16vec2(unpackUint2x16(uint(_29.inputs[ident].y)) - float16BitsToUint16(_40.const0.zw));
|
|
||||||
_21.outputs[ident] = i16vec4(_21.outputs[ident].x, _21.outputs[ident].y, _66.x, _66.y);
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,65 +0,0 @@
|
|||||||
#version 310 es
|
|
||||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
|
||||||
|
|
||||||
struct Foo
|
|
||||||
{
|
|
||||||
vec4 a;
|
|
||||||
vec4 b;
|
|
||||||
vec4 c;
|
|
||||||
vec4 d;
|
|
||||||
};
|
|
||||||
|
|
||||||
layout(binding = 1, std430) readonly buffer SSBO2
|
|
||||||
{
|
|
||||||
vec4 data[];
|
|
||||||
} indata;
|
|
||||||
|
|
||||||
layout(binding = 0, std430) writeonly buffer SSBO
|
|
||||||
{
|
|
||||||
vec4 data[];
|
|
||||||
} outdata;
|
|
||||||
|
|
||||||
layout(binding = 2, std430) readonly buffer SSBO3
|
|
||||||
{
|
|
||||||
Foo foos[];
|
|
||||||
} foobar;
|
|
||||||
|
|
||||||
void baz(inout Foo foo)
|
|
||||||
{
|
|
||||||
uint ident = gl_GlobalInvocationID.x;
|
|
||||||
foo.a = indata.data[(4u * ident) + 0u];
|
|
||||||
foo.b = indata.data[(4u * ident) + 1u];
|
|
||||||
foo.c = indata.data[(4u * ident) + 2u];
|
|
||||||
foo.d = indata.data[(4u * ident) + 3u];
|
|
||||||
}
|
|
||||||
|
|
||||||
void meow(inout Foo foo)
|
|
||||||
{
|
|
||||||
foo.a += vec4(10.0);
|
|
||||||
foo.b += vec4(20.0);
|
|
||||||
foo.c += vec4(30.0);
|
|
||||||
foo.d += vec4(40.0);
|
|
||||||
}
|
|
||||||
|
|
||||||
vec4 bar(Foo foo)
|
|
||||||
{
|
|
||||||
return ((foo.a + foo.b) + foo.c) + foo.d;
|
|
||||||
}
|
|
||||||
|
|
||||||
void main()
|
|
||||||
{
|
|
||||||
Foo param;
|
|
||||||
baz(param);
|
|
||||||
Foo foo = param;
|
|
||||||
Foo param_1 = foo;
|
|
||||||
meow(param_1);
|
|
||||||
foo = param_1;
|
|
||||||
Foo param_2 = foo;
|
|
||||||
Foo param_3;
|
|
||||||
param_3.a = foobar.foos[gl_GlobalInvocationID.x].a;
|
|
||||||
param_3.b = foobar.foos[gl_GlobalInvocationID.x].b;
|
|
||||||
param_3.c = foobar.foos[gl_GlobalInvocationID.x].c;
|
|
||||||
param_3.d = foobar.foos[gl_GlobalInvocationID.x].d;
|
|
||||||
outdata.data[gl_GlobalInvocationID.x] = bar(param_2) + bar(param_3);
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,159 +0,0 @@
|
|||||||
#version 450
|
|
||||||
#if defined(GL_AMD_gpu_shader_half_float)
|
|
||||||
#extension GL_AMD_gpu_shader_half_float : require
|
|
||||||
#elif defined(GL_NV_gpu_shader5)
|
|
||||||
#extension GL_NV_gpu_shader5 : require
|
|
||||||
#else
|
|
||||||
#error No extension available for FP16.
|
|
||||||
#endif
|
|
||||||
|
|
||||||
struct ResType
|
|
||||||
{
|
|
||||||
f16vec4 _m0;
|
|
||||||
ivec4 _m1;
|
|
||||||
};
|
|
||||||
|
|
||||||
layout(location = 3) in f16vec4 v4;
|
|
||||||
layout(location = 2) in f16vec3 v3;
|
|
||||||
layout(location = 0) in float16_t v1;
|
|
||||||
layout(location = 1) in f16vec2 v2;
|
|
||||||
|
|
||||||
f16mat2 test_mat2(f16vec2 a, f16vec2 b, f16vec2 c, f16vec2 d)
|
|
||||||
{
|
|
||||||
return f16mat2(f16vec2(a), f16vec2(b)) * f16mat2(f16vec2(c), f16vec2(d));
|
|
||||||
}
|
|
||||||
|
|
||||||
f16mat3 test_mat3(f16vec3 a, f16vec3 b, f16vec3 c, f16vec3 d, f16vec3 e, f16vec3 f)
|
|
||||||
{
|
|
||||||
return f16mat3(f16vec3(a), f16vec3(b), f16vec3(c)) * f16mat3(f16vec3(d), f16vec3(e), f16vec3(f));
|
|
||||||
}
|
|
||||||
|
|
||||||
void test_constants()
|
|
||||||
{
|
|
||||||
float16_t a = float16_t(1.0);
|
|
||||||
float16_t b = float16_t(1.5);
|
|
||||||
float16_t c = float16_t(-1.5);
|
|
||||||
float16_t d = float16_t(0.0 / 0.0);
|
|
||||||
float16_t e = float16_t(1.0 / 0.0);
|
|
||||||
float16_t f = float16_t(-1.0 / 0.0);
|
|
||||||
float16_t g = float16_t(1014.0);
|
|
||||||
float16_t h = float16_t(9.5367431640625e-07);
|
|
||||||
}
|
|
||||||
|
|
||||||
float16_t test_result()
|
|
||||||
{
|
|
||||||
return float16_t(1.0);
|
|
||||||
}
|
|
||||||
|
|
||||||
void test_conversions()
|
|
||||||
{
|
|
||||||
float16_t one = test_result();
|
|
||||||
int a = int(one);
|
|
||||||
uint b = uint(one);
|
|
||||||
bool c = one != float16_t(0.0);
|
|
||||||
float d = float(one);
|
|
||||||
double e = double(one);
|
|
||||||
float16_t a2 = float16_t(a);
|
|
||||||
float16_t b2 = float16_t(b);
|
|
||||||
float16_t c2 = float16_t(c);
|
|
||||||
float16_t d2 = float16_t(d);
|
|
||||||
float16_t e2 = float16_t(e);
|
|
||||||
}
|
|
||||||
|
|
||||||
void test_builtins()
|
|
||||||
{
|
|
||||||
f16vec4 res = radians(v4);
|
|
||||||
res = degrees(v4);
|
|
||||||
res = sin(v4);
|
|
||||||
res = cos(v4);
|
|
||||||
res = tan(v4);
|
|
||||||
res = asin(v4);
|
|
||||||
res = atan(v4, v3.xyzz);
|
|
||||||
res = atan(v4);
|
|
||||||
res = sinh(v4);
|
|
||||||
res = cosh(v4);
|
|
||||||
res = tanh(v4);
|
|
||||||
res = asinh(v4);
|
|
||||||
res = acosh(v4);
|
|
||||||
res = atanh(v4);
|
|
||||||
res = pow(v4, v4);
|
|
||||||
res = exp(v4);
|
|
||||||
res = log(v4);
|
|
||||||
res = exp2(v4);
|
|
||||||
res = log2(v4);
|
|
||||||
res = sqrt(v4);
|
|
||||||
res = inversesqrt(v4);
|
|
||||||
res = abs(v4);
|
|
||||||
res = sign(v4);
|
|
||||||
res = floor(v4);
|
|
||||||
res = trunc(v4);
|
|
||||||
res = round(v4);
|
|
||||||
res = roundEven(v4);
|
|
||||||
res = ceil(v4);
|
|
||||||
res = fract(v4);
|
|
||||||
res = mod(v4, v4);
|
|
||||||
f16vec4 tmp;
|
|
||||||
f16vec4 _231 = modf(v4, tmp);
|
|
||||||
res = _231;
|
|
||||||
res = min(v4, v4);
|
|
||||||
res = max(v4, v4);
|
|
||||||
res = clamp(v4, v4, v4);
|
|
||||||
res = mix(v4, v4, v4);
|
|
||||||
res = mix(v4, v4, lessThan(v4, v4));
|
|
||||||
res = step(v4, v4);
|
|
||||||
res = smoothstep(v4, v4, v4);
|
|
||||||
bvec4 btmp = isnan(v4);
|
|
||||||
btmp = isinf(v4);
|
|
||||||
res = fma(v4, v4, v4);
|
|
||||||
ResType _275;
|
|
||||||
_275._m0 = frexp(v4, _275._m1);
|
|
||||||
ivec4 itmp = _275._m1;
|
|
||||||
res = _275._m0;
|
|
||||||
res = ldexp(res, itmp);
|
|
||||||
uint pack0 = packFloat2x16(v4.xy);
|
|
||||||
uint pack1 = packFloat2x16(v4.zw);
|
|
||||||
res = f16vec4(unpackFloat2x16(pack0), unpackFloat2x16(pack1));
|
|
||||||
float16_t t0 = length(v4);
|
|
||||||
t0 = distance(v4, v4);
|
|
||||||
t0 = dot(v4, v4);
|
|
||||||
f16vec3 res3 = cross(v3, v3);
|
|
||||||
res = normalize(v4);
|
|
||||||
res = faceforward(v4, v4, v4);
|
|
||||||
res = reflect(v4, v4);
|
|
||||||
res = refract(v4, v4, v1);
|
|
||||||
btmp = lessThan(v4, v4);
|
|
||||||
btmp = lessThanEqual(v4, v4);
|
|
||||||
btmp = greaterThan(v4, v4);
|
|
||||||
btmp = greaterThanEqual(v4, v4);
|
|
||||||
btmp = equal(v4, v4);
|
|
||||||
btmp = notEqual(v4, v4);
|
|
||||||
res = dFdx(v4);
|
|
||||||
res = dFdy(v4);
|
|
||||||
res = dFdxFine(v4);
|
|
||||||
res = dFdyFine(v4);
|
|
||||||
res = dFdxCoarse(v4);
|
|
||||||
res = dFdyCoarse(v4);
|
|
||||||
res = fwidth(v4);
|
|
||||||
res = fwidthFine(v4);
|
|
||||||
res = fwidthCoarse(v4);
|
|
||||||
}
|
|
||||||
|
|
||||||
void main()
|
|
||||||
{
|
|
||||||
f16vec2 param = v2;
|
|
||||||
f16vec2 param_1 = v2;
|
|
||||||
f16vec2 param_2 = v3.xy;
|
|
||||||
f16vec2 param_3 = v3.xy;
|
|
||||||
f16mat2 m0 = test_mat2(param, param_1, param_2, param_3);
|
|
||||||
f16vec3 param_4 = v3;
|
|
||||||
f16vec3 param_5 = v3;
|
|
||||||
f16vec3 param_6 = v3;
|
|
||||||
f16vec3 param_7 = v4.xyz;
|
|
||||||
f16vec3 param_8 = v4.xyz;
|
|
||||||
f16vec3 param_9 = v4.yzw;
|
|
||||||
f16mat3 m1 = test_mat3(param_4, param_5, param_6, param_7, param_8, param_9);
|
|
||||||
test_constants();
|
|
||||||
test_conversions();
|
|
||||||
test_builtins();
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,24 +0,0 @@
|
|||||||
#version 450
|
|
||||||
|
|
||||||
layout(binding = 0) uniform sampler2D uTextures[2 * 3 * 1];
|
|
||||||
|
|
||||||
layout(location = 1) in vec2 vUV;
|
|
||||||
layout(location = 0) out vec4 FragColor;
|
|
||||||
layout(location = 0) flat in int vIndex;
|
|
||||||
|
|
||||||
void main()
|
|
||||||
{
|
|
||||||
vec4 values3[2 * 3 * 1];
|
|
||||||
for (int z = 0; z < 2; z++)
|
|
||||||
{
|
|
||||||
for (int y = 0; y < 3; y++)
|
|
||||||
{
|
|
||||||
for (int x = 0; x < 1; x++)
|
|
||||||
{
|
|
||||||
values3[z * 3 * 1 + y * 1 + x] = texture(uTextures[z * 3 * 1 + y * 1 + x], vUV);
|
|
||||||
}
|
|
||||||
}
|
|
||||||
}
|
|
||||||
FragColor = (values3[1 * 3 * 1 + 2 * 1 + 0] + values3[0 * 3 * 1 + 2 * 1 + 0]) + values3[(vIndex + 1) * 3 * 1 + 2 * 1 + vIndex];
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,25 +0,0 @@
|
|||||||
#version 450
|
|
||||||
#if defined(GL_AMD_gpu_shader_half_float)
|
|
||||||
#extension GL_AMD_gpu_shader_half_float : require
|
|
||||||
#elif defined(GL_NV_gpu_shader5)
|
|
||||||
#extension GL_NV_gpu_shader5 : require
|
|
||||||
#else
|
|
||||||
#error No extension available for FP16.
|
|
||||||
#endif
|
|
||||||
#if defined(GL_AMD_gpu_shader_int16)
|
|
||||||
#extension GL_AMD_gpu_shader_int16 : require
|
|
||||||
#else
|
|
||||||
#error No extension available for Int16.
|
|
||||||
#endif
|
|
||||||
|
|
||||||
layout(location = 0) out float16_t foo;
|
|
||||||
layout(location = 1) out int16_t bar;
|
|
||||||
layout(location = 2) out uint16_t baz;
|
|
||||||
|
|
||||||
void main()
|
|
||||||
{
|
|
||||||
foo = float16_t(1.0);
|
|
||||||
bar = 2s;
|
|
||||||
baz = 3us;
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,147 +0,0 @@
|
|||||||
#version 310 es
|
|
||||||
#extension GL_EXT_scalar_block_layout : require
|
|
||||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
|
||||||
|
|
||||||
struct S0
|
|
||||||
{
|
|
||||||
vec2 a[1];
|
|
||||||
float b;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct S1
|
|
||||||
{
|
|
||||||
vec3 a;
|
|
||||||
float b;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct S2
|
|
||||||
{
|
|
||||||
vec3 a[1];
|
|
||||||
float b;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct S3
|
|
||||||
{
|
|
||||||
vec2 a;
|
|
||||||
float b;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct S4
|
|
||||||
{
|
|
||||||
vec2 c;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct Content
|
|
||||||
{
|
|
||||||
S0 m0s[1];
|
|
||||||
S1 m1s[1];
|
|
||||||
S2 m2s[1];
|
|
||||||
S0 m0;
|
|
||||||
S1 m1;
|
|
||||||
S2 m2;
|
|
||||||
S3 m3;
|
|
||||||
float m4;
|
|
||||||
S4 m3s[8];
|
|
||||||
};
|
|
||||||
|
|
||||||
struct S0_1
|
|
||||||
{
|
|
||||||
vec2 a[1];
|
|
||||||
float b;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct S1_1
|
|
||||||
{
|
|
||||||
vec3 a;
|
|
||||||
float b;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct S2_1
|
|
||||||
{
|
|
||||||
vec3 a[1];
|
|
||||||
float b;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct S3_1
|
|
||||||
{
|
|
||||||
vec2 a;
|
|
||||||
float b;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct S4_1
|
|
||||||
{
|
|
||||||
vec2 c;
|
|
||||||
};
|
|
||||||
|
|
||||||
struct Content_1
|
|
||||||
{
|
|
||||||
S0_1 m0s[1];
|
|
||||||
S1_1 m1s[1];
|
|
||||||
S2_1 m2s[1];
|
|
||||||
S0_1 m0;
|
|
||||||
S1_1 m1;
|
|
||||||
S2_1 m2;
|
|
||||||
S3_1 m3;
|
|
||||||
float m4;
|
|
||||||
S4_1 m3s[8];
|
|
||||||
};
|
|
||||||
|
|
||||||
layout(set = 0, binding = 1, scalar) restrict buffer SSBO1
|
|
||||||
{
|
|
||||||
Content content;
|
|
||||||
Content content1[2];
|
|
||||||
Content content2;
|
|
||||||
mat2 m0;
|
|
||||||
mat2 m1;
|
|
||||||
mat2x3 m2[4];
|
|
||||||
mat3x2 m3;
|
|
||||||
layout(row_major) mat2 m4;
|
|
||||||
layout(row_major) mat2 m5[9];
|
|
||||||
layout(row_major) mat2x3 m6[4][2];
|
|
||||||
layout(row_major) mat3x2 m7;
|
|
||||||
float array[];
|
|
||||||
} ssbo_430;
|
|
||||||
|
|
||||||
layout(set = 0, binding = 0, std140) restrict buffer SSBO0
|
|
||||||
{
|
|
||||||
Content_1 content;
|
|
||||||
Content_1 content1[2];
|
|
||||||
Content_1 content2;
|
|
||||||
mat2 m0;
|
|
||||||
mat2 m1;
|
|
||||||
mat2x3 m2[4];
|
|
||||||
mat3x2 m3;
|
|
||||||
layout(row_major) mat2 m4;
|
|
||||||
layout(row_major) mat2 m5[9];
|
|
||||||
layout(row_major) mat2x3 m6[4][2];
|
|
||||||
layout(row_major) mat3x2 m7;
|
|
||||||
float array[];
|
|
||||||
} ssbo_140;
|
|
||||||
|
|
||||||
void main()
|
|
||||||
{
|
|
||||||
ssbo_430.content.m0s[0].a[0] = ssbo_140.content.m0s[0].a[0];
|
|
||||||
ssbo_430.content.m0s[0].b = ssbo_140.content.m0s[0].b;
|
|
||||||
ssbo_430.content.m1s[0].a = ssbo_140.content.m1s[0].a;
|
|
||||||
ssbo_430.content.m1s[0].b = ssbo_140.content.m1s[0].b;
|
|
||||||
ssbo_430.content.m2s[0].a[0] = ssbo_140.content.m2s[0].a[0];
|
|
||||||
ssbo_430.content.m2s[0].b = ssbo_140.content.m2s[0].b;
|
|
||||||
ssbo_430.content.m0.a[0] = ssbo_140.content.m0.a[0];
|
|
||||||
ssbo_430.content.m0.b = ssbo_140.content.m0.b;
|
|
||||||
ssbo_430.content.m1.a = ssbo_140.content.m1.a;
|
|
||||||
ssbo_430.content.m1.b = ssbo_140.content.m1.b;
|
|
||||||
ssbo_430.content.m2.a[0] = ssbo_140.content.m2.a[0];
|
|
||||||
ssbo_430.content.m2.b = ssbo_140.content.m2.b;
|
|
||||||
ssbo_430.content.m3.a = ssbo_140.content.m3.a;
|
|
||||||
ssbo_430.content.m3.b = ssbo_140.content.m3.b;
|
|
||||||
ssbo_430.content.m4 = ssbo_140.content.m4;
|
|
||||||
ssbo_430.content.m3s[0].c = ssbo_140.content.m3s[0].c;
|
|
||||||
ssbo_430.content.m3s[1].c = ssbo_140.content.m3s[1].c;
|
|
||||||
ssbo_430.content.m3s[2].c = ssbo_140.content.m3s[2].c;
|
|
||||||
ssbo_430.content.m3s[3].c = ssbo_140.content.m3s[3].c;
|
|
||||||
ssbo_430.content.m3s[4].c = ssbo_140.content.m3s[4].c;
|
|
||||||
ssbo_430.content.m3s[5].c = ssbo_140.content.m3s[5].c;
|
|
||||||
ssbo_430.content.m3s[6].c = ssbo_140.content.m3s[6].c;
|
|
||||||
ssbo_430.content.m3s[7].c = ssbo_140.content.m3s[7].c;
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,110 +0,0 @@
|
|||||||
#version 450
|
|
||||||
#extension GL_KHR_shader_subgroup_basic : require
|
|
||||||
#extension GL_KHR_shader_subgroup_ballot : require
|
|
||||||
#extension GL_KHR_shader_subgroup_shuffle : require
|
|
||||||
#extension GL_KHR_shader_subgroup_shuffle_relative : require
|
|
||||||
#extension GL_KHR_shader_subgroup_vote : 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, local_size_y = 1, local_size_z = 1) in;
|
|
||||||
|
|
||||||
layout(set = 0, binding = 0, std430) buffer SSBO
|
|
||||||
{
|
|
||||||
float FragColor;
|
|
||||||
} _9;
|
|
||||||
|
|
||||||
void main()
|
|
||||||
{
|
|
||||||
_9.FragColor = float(gl_NumSubgroups);
|
|
||||||
_9.FragColor = float(gl_SubgroupID);
|
|
||||||
_9.FragColor = float(gl_SubgroupSize);
|
|
||||||
_9.FragColor = float(gl_SubgroupInvocationID);
|
|
||||||
subgroupMemoryBarrier();
|
|
||||||
subgroupBarrier();
|
|
||||||
subgroupMemoryBarrier();
|
|
||||||
subgroupMemoryBarrierBuffer();
|
|
||||||
subgroupMemoryBarrierShared();
|
|
||||||
subgroupMemoryBarrierImage();
|
|
||||||
bool elected = subgroupElect();
|
|
||||||
_9.FragColor = vec4(gl_SubgroupEqMask).x;
|
|
||||||
_9.FragColor = vec4(gl_SubgroupGeMask).x;
|
|
||||||
_9.FragColor = vec4(gl_SubgroupGtMask).x;
|
|
||||||
_9.FragColor = vec4(gl_SubgroupLeMask).x;
|
|
||||||
_9.FragColor = vec4(gl_SubgroupLtMask).x;
|
|
||||||
vec4 broadcasted = subgroupBroadcast(vec4(10.0), 8u);
|
|
||||||
vec3 first = subgroupBroadcastFirst(vec3(20.0));
|
|
||||||
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);
|
|
||||||
uint shuffled = subgroupShuffle(10u, 8u);
|
|
||||||
uint shuffled_xor = subgroupShuffleXor(30u, 8u);
|
|
||||||
uint shuffled_up = subgroupShuffleUp(20u, 4u);
|
|
||||||
uint shuffled_down = subgroupShuffleDown(20u, 4u);
|
|
||||||
bool has_all = subgroupAll(true);
|
|
||||||
bool has_any = subgroupAny(true);
|
|
||||||
bool has_equal = subgroupAllEqual(true);
|
|
||||||
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(20u));
|
|
||||||
uvec4 uhi = subgroupMax(uvec4(20u));
|
|
||||||
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);
|
|
||||||
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);
|
|
||||||
hi = subgroupExclusiveMax(hi);
|
|
||||||
ulo = subgroupExclusiveMin(ulo);
|
|
||||||
uhi = subgroupExclusiveMax(uhi);
|
|
||||||
slo = subgroupExclusiveMin(slo);
|
|
||||||
shi = subgroupExclusiveMax(shi);
|
|
||||||
anded = subgroupExclusiveAnd(anded);
|
|
||||||
ored = subgroupExclusiveOr(ored);
|
|
||||||
xored = subgroupExclusiveXor(ored);
|
|
||||||
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);
|
|
||||||
vec4 swap_horiz = subgroupQuadSwapHorizontal(vec4(20.0));
|
|
||||||
vec4 swap_vertical = subgroupQuadSwapVertical(vec4(20.0));
|
|
||||||
vec4 swap_diagonal = subgroupQuadSwapDiagonal(vec4(20.0));
|
|
||||||
vec4 quad_broadcast = subgroupQuadBroadcast(vec4(20.0), 3u);
|
|
||||||
}
|
|
||||||
|
|
||||||
@ -1,24 +0,0 @@
|
|||||||
#version 450
|
|
||||||
#extension GL_EXT_scalar_block_layout : require
|
|
||||||
|
|
||||||
layout(set = 0, binding = 0, std430) uniform UBO
|
|
||||||
{
|
|
||||||
float a[1024];
|
|
||||||
vec3 b[2];
|
|
||||||
} _17;
|
|
||||||
|
|
||||||
layout(set = 0, binding = 1, std430) uniform UBOEnhancedLayout
|
|
||||||
{
|
|
||||||
layout(offset = 0) float c[1024];
|
|
||||||
layout(offset = 4096) vec3 d[2];
|
|
||||||
layout(offset = 10000) float e;
|
|
||||||
} _30;
|
|
||||||
|
|
||||||
layout(location = 0) out float FragColor;
|
|
||||||
layout(location = 0) flat in int vIndex;
|
|
||||||
|
|
||||||
void main()
|
|
||||||
{
|
|
||||||
FragColor = (_17.a[vIndex] + _30.c[vIndex]) + _30.e;
|
|
||||||
}
|
|
||||||
|
|
||||||
Some files were not shown because too many files have changed in this diff Show More
Loading…
Reference in New Issue
Block a user