Test glsl.std450 more exhaustively.
Make sure to test everything with scalar as well to catch any weird edge cases. Not all opcodes are covered here, just the arithmetic ones. FP64 packing is also ignored.
This commit is contained in:
Hans-Kristian Arntzen
parent
3319a3bfd1
commit
c7eda1bce9
@ -0,0 +1,31 @@
|
||||
struct _8
|
||||
{
|
||||
float _m0;
|
||||
float _m1;
|
||||
};
|
||||
|
||||
struct _15
|
||||
{
|
||||
float _m0;
|
||||
int _m1;
|
||||
};
|
||||
|
||||
RWByteAddressBuffer _4 : register(u0);
|
||||
|
||||
void comp_main()
|
||||
{
|
||||
_8 _23;
|
||||
_23._m0 = modf(20.0f, _23._m1);
|
||||
_15 _24;
|
||||
_24._m0 = frexp(40.0f, _24._m1);
|
||||
_4.Store(0, asuint(_23._m0));
|
||||
_4.Store(0, asuint(_23._m1));
|
||||
_4.Store(0, asuint(_24._m0));
|
||||
_4.Store(4, uint(_24._m1));
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void main()
|
||||
{
|
||||
comp_main();
|
||||
}
|
||||
295
reference/shaders-hlsl-no-opt/comp/glsl.std450.fxconly.comp
Normal file
295
reference/shaders-hlsl-no-opt/comp/glsl.std450.fxconly.comp
Normal file
@ -0,0 +1,295 @@
|
||||
struct ResType
|
||||
{
|
||||
float _m0;
|
||||
int _m1;
|
||||
};
|
||||
|
||||
RWByteAddressBuffer _19 : register(u0);
|
||||
|
||||
uint SPIRV_Cross_packHalf2x16(float2 value)
|
||||
{
|
||||
uint2 Packed = f32tof16(value);
|
||||
return Packed.x | (Packed.y << 16);
|
||||
}
|
||||
|
||||
float2 SPIRV_Cross_unpackHalf2x16(uint value)
|
||||
{
|
||||
return f16tof32(uint2(value & 0xffff, value >> 16));
|
||||
}
|
||||
|
||||
uint SPIRV_Cross_packUnorm4x8(float4 value)
|
||||
{
|
||||
uint4 Packed = uint4(round(saturate(value) * 255.0));
|
||||
return Packed.x | (Packed.y << 8) | (Packed.z << 16) | (Packed.w << 24);
|
||||
}
|
||||
|
||||
float4 SPIRV_Cross_unpackUnorm4x8(uint value)
|
||||
{
|
||||
uint4 Packed = uint4(value & 0xff, (value >> 8) & 0xff, (value >> 16) & 0xff, value >> 24);
|
||||
return float4(Packed) / 255.0;
|
||||
}
|
||||
|
||||
uint SPIRV_Cross_packSnorm4x8(float4 value)
|
||||
{
|
||||
int4 Packed = int4(round(clamp(value, -1.0, 1.0) * 127.0)) & 0xff;
|
||||
return uint(Packed.x | (Packed.y << 8) | (Packed.z << 16) | (Packed.w << 24));
|
||||
}
|
||||
|
||||
float4 SPIRV_Cross_unpackSnorm4x8(uint value)
|
||||
{
|
||||
int SignedValue = int(value);
|
||||
int4 Packed = int4(SignedValue << 24, SignedValue << 16, SignedValue << 8, SignedValue) >> 24;
|
||||
return clamp(float4(Packed) / 127.0, -1.0, 1.0);
|
||||
}
|
||||
|
||||
uint SPIRV_Cross_packUnorm2x16(float2 value)
|
||||
{
|
||||
uint2 Packed = uint2(round(saturate(value) * 65535.0));
|
||||
return Packed.x | (Packed.y << 16);
|
||||
}
|
||||
|
||||
float2 SPIRV_Cross_unpackUnorm2x16(uint value)
|
||||
{
|
||||
uint2 Packed = uint2(value & 0xffff, value >> 16);
|
||||
return float2(Packed) / 65535.0;
|
||||
}
|
||||
|
||||
uint SPIRV_Cross_packSnorm2x16(float2 value)
|
||||
{
|
||||
int2 Packed = int2(round(clamp(value, -1.0, 1.0) * 32767.0)) & 0xffff;
|
||||
return uint(Packed.x | (Packed.y << 16));
|
||||
}
|
||||
|
||||
float2 SPIRV_Cross_unpackSnorm2x16(uint value)
|
||||
{
|
||||
int SignedValue = int(value);
|
||||
int2 Packed = int2(SignedValue << 16, SignedValue) >> 16;
|
||||
return clamp(float2(Packed) / 32767.0, -1.0, 1.0);
|
||||
}
|
||||
|
||||
// Returns the inverse of a matrix, by using the algorithm of calculating the classical
|
||||
// adjoint and dividing by the determinant. The contents of the matrix are changed.
|
||||
float2x2 SPIRV_Cross_Inverse(float2x2 m)
|
||||
{
|
||||
float2x2 adj; // The adjoint matrix (inverse after dividing by determinant)
|
||||
|
||||
// Create the transpose of the cofactors, as the classical adjoint of the matrix.
|
||||
adj[0][0] = m[1][1];
|
||||
adj[0][1] = -m[0][1];
|
||||
|
||||
adj[1][0] = -m[1][0];
|
||||
adj[1][1] = m[0][0];
|
||||
|
||||
// Calculate the determinant as a combination of the cofactors of the first row.
|
||||
float det = (adj[0][0] * m[0][0]) + (adj[0][1] * m[1][0]);
|
||||
|
||||
// Divide the classical adjoint matrix by the determinant.
|
||||
// If determinant is zero, matrix is not invertable, so leave it unchanged.
|
||||
return (det != 0.0f) ? (adj * (1.0f / det)) : m;
|
||||
}
|
||||
|
||||
// Returns the determinant of a 2x2 matrix.
|
||||
float SPIRV_Cross_Det2x2(float a1, float a2, float b1, float b2)
|
||||
{
|
||||
return a1 * b2 - b1 * a2;
|
||||
}
|
||||
|
||||
// Returns the inverse of a matrix, by using the algorithm of calculating the classical
|
||||
// adjoint and dividing by the determinant. The contents of the matrix are changed.
|
||||
float3x3 SPIRV_Cross_Inverse(float3x3 m)
|
||||
{
|
||||
float3x3 adj; // The adjoint matrix (inverse after dividing by determinant)
|
||||
|
||||
// Create the transpose of the cofactors, as the classical adjoint of the matrix.
|
||||
adj[0][0] = SPIRV_Cross_Det2x2(m[1][1], m[1][2], m[2][1], m[2][2]);
|
||||
adj[0][1] = -SPIRV_Cross_Det2x2(m[0][1], m[0][2], m[2][1], m[2][2]);
|
||||
adj[0][2] = SPIRV_Cross_Det2x2(m[0][1], m[0][2], m[1][1], m[1][2]);
|
||||
|
||||
adj[1][0] = -SPIRV_Cross_Det2x2(m[1][0], m[1][2], m[2][0], m[2][2]);
|
||||
adj[1][1] = SPIRV_Cross_Det2x2(m[0][0], m[0][2], m[2][0], m[2][2]);
|
||||
adj[1][2] = -SPIRV_Cross_Det2x2(m[0][0], m[0][2], m[1][0], m[1][2]);
|
||||
|
||||
adj[2][0] = SPIRV_Cross_Det2x2(m[1][0], m[1][1], m[2][0], m[2][1]);
|
||||
adj[2][1] = -SPIRV_Cross_Det2x2(m[0][0], m[0][1], m[2][0], m[2][1]);
|
||||
adj[2][2] = SPIRV_Cross_Det2x2(m[0][0], m[0][1], m[1][0], m[1][1]);
|
||||
|
||||
// Calculate the determinant as a combination of the cofactors of the first row.
|
||||
float det = (adj[0][0] * m[0][0]) + (adj[0][1] * m[1][0]) + (adj[0][2] * m[2][0]);
|
||||
|
||||
// Divide the classical adjoint matrix by the determinant.
|
||||
// If determinant is zero, matrix is not invertable, so leave it unchanged.
|
||||
return (det != 0.0f) ? (adj * (1.0f / det)) : m;
|
||||
}
|
||||
|
||||
// Returns the determinant of a 3x3 matrix.
|
||||
float SPIRV_Cross_Det3x3(float a1, float a2, float a3, float b1, float b2, float b3, float c1, float c2, float c3)
|
||||
{
|
||||
return a1 * SPIRV_Cross_Det2x2(b2, b3, c2, c3) - b1 * SPIRV_Cross_Det2x2(a2, a3, c2, c3) + c1 * SPIRV_Cross_Det2x2(a2, a3, b2, b3);
|
||||
}
|
||||
|
||||
// Returns the inverse of a matrix, by using the algorithm of calculating the classical
|
||||
// adjoint and dividing by the determinant. The contents of the matrix are changed.
|
||||
float4x4 SPIRV_Cross_Inverse(float4x4 m)
|
||||
{
|
||||
float4x4 adj; // The adjoint matrix (inverse after dividing by determinant)
|
||||
|
||||
// Create the transpose of the cofactors, as the classical adjoint of the matrix.
|
||||
adj[0][0] = SPIRV_Cross_Det3x3(m[1][1], m[1][2], m[1][3], m[2][1], m[2][2], m[2][3], m[3][1], m[3][2], m[3][3]);
|
||||
adj[0][1] = -SPIRV_Cross_Det3x3(m[0][1], m[0][2], m[0][3], m[2][1], m[2][2], m[2][3], m[3][1], m[3][2], m[3][3]);
|
||||
adj[0][2] = SPIRV_Cross_Det3x3(m[0][1], m[0][2], m[0][3], m[1][1], m[1][2], m[1][3], m[3][1], m[3][2], m[3][3]);
|
||||
adj[0][3] = -SPIRV_Cross_Det3x3(m[0][1], m[0][2], m[0][3], m[1][1], m[1][2], m[1][3], m[2][1], m[2][2], m[2][3]);
|
||||
|
||||
adj[1][0] = -SPIRV_Cross_Det3x3(m[1][0], m[1][2], m[1][3], m[2][0], m[2][2], m[2][3], m[3][0], m[3][2], m[3][3]);
|
||||
adj[1][1] = SPIRV_Cross_Det3x3(m[0][0], m[0][2], m[0][3], m[2][0], m[2][2], m[2][3], m[3][0], m[3][2], m[3][3]);
|
||||
adj[1][2] = -SPIRV_Cross_Det3x3(m[0][0], m[0][2], m[0][3], m[1][0], m[1][2], m[1][3], m[3][0], m[3][2], m[3][3]);
|
||||
adj[1][3] = SPIRV_Cross_Det3x3(m[0][0], m[0][2], m[0][3], m[1][0], m[1][2], m[1][3], m[2][0], m[2][2], m[2][3]);
|
||||
|
||||
adj[2][0] = SPIRV_Cross_Det3x3(m[1][0], m[1][1], m[1][3], m[2][0], m[2][1], m[2][3], m[3][0], m[3][1], m[3][3]);
|
||||
adj[2][1] = -SPIRV_Cross_Det3x3(m[0][0], m[0][1], m[0][3], m[2][0], m[2][1], m[2][3], m[3][0], m[3][1], m[3][3]);
|
||||
adj[2][2] = SPIRV_Cross_Det3x3(m[0][0], m[0][1], m[0][3], m[1][0], m[1][1], m[1][3], m[3][0], m[3][1], m[3][3]);
|
||||
adj[2][3] = -SPIRV_Cross_Det3x3(m[0][0], m[0][1], m[0][3], m[1][0], m[1][1], m[1][3], m[2][0], m[2][1], m[2][3]);
|
||||
|
||||
adj[3][0] = -SPIRV_Cross_Det3x3(m[1][0], m[1][1], m[1][2], m[2][0], m[2][1], m[2][2], m[3][0], m[3][1], m[3][2]);
|
||||
adj[3][1] = SPIRV_Cross_Det3x3(m[0][0], m[0][1], m[0][2], m[2][0], m[2][1], m[2][2], m[3][0], m[3][1], m[3][2]);
|
||||
adj[3][2] = -SPIRV_Cross_Det3x3(m[0][0], m[0][1], m[0][2], m[1][0], m[1][1], m[1][2], m[3][0], m[3][1], m[3][2]);
|
||||
adj[3][3] = SPIRV_Cross_Det3x3(m[0][0], m[0][1], m[0][2], m[1][0], m[1][1], m[1][2], m[2][0], m[2][1], m[2][2]);
|
||||
|
||||
// Calculate the determinant as a combination of the cofactors of the first row.
|
||||
float det = (adj[0][0] * m[0][0]) + (adj[0][1] * m[1][0]) + (adj[0][2] * m[2][0]) + (adj[0][3] * m[3][0]);
|
||||
|
||||
// Divide the classical adjoint matrix by the determinant.
|
||||
// If determinant is zero, matrix is not invertable, so leave it unchanged.
|
||||
return (det != 0.0f) ? (adj * (1.0f / det)) : m;
|
||||
}
|
||||
|
||||
float SPIRV_Cross_Reflect(float i, float n)
|
||||
{
|
||||
return i - 2.0 * dot(n, i) * n;
|
||||
}
|
||||
|
||||
float SPIRV_Cross_Refract(float i, float n, float eta)
|
||||
{
|
||||
float NoI = n * i;
|
||||
float NoI2 = NoI * NoI;
|
||||
float k = 1.0 - eta * eta * (1.0 - NoI2);
|
||||
if (k < 0.0)
|
||||
{
|
||||
return 0.0;
|
||||
}
|
||||
else
|
||||
{
|
||||
return eta * i - (eta * NoI + sqrt(k)) * n;
|
||||
}
|
||||
}
|
||||
|
||||
float SPIRV_Cross_FaceForward(float n, float i, float nref)
|
||||
{
|
||||
return i * nref < 0.0 ? n : -n;
|
||||
}
|
||||
|
||||
void comp_main()
|
||||
{
|
||||
_19.Store(0, asuint(round(asfloat(_19.Load(16)))));
|
||||
_19.Store(0, asuint(trunc(asfloat(_19.Load(16)))));
|
||||
_19.Store(0, asuint(abs(asfloat(_19.Load(16)))));
|
||||
_19.Store(4, uint(abs(int(_19.Load(32)))));
|
||||
_19.Store(0, asuint(sign(asfloat(_19.Load(16)))));
|
||||
_19.Store(4, uint(sign(int(_19.Load(32)))));
|
||||
_19.Store(0, asuint(floor(asfloat(_19.Load(16)))));
|
||||
_19.Store(0, asuint(ceil(asfloat(_19.Load(16)))));
|
||||
_19.Store(0, asuint(frac(asfloat(_19.Load(16)))));
|
||||
_19.Store(0, asuint(radians(asfloat(_19.Load(16)))));
|
||||
_19.Store(0, asuint(degrees(asfloat(_19.Load(16)))));
|
||||
_19.Store(0, asuint(sin(asfloat(_19.Load(16)))));
|
||||
_19.Store(0, asuint(cos(asfloat(_19.Load(16)))));
|
||||
_19.Store(0, asuint(tan(asfloat(_19.Load(16)))));
|
||||
_19.Store(0, asuint(asin(asfloat(_19.Load(16)))));
|
||||
_19.Store(0, asuint(acos(asfloat(_19.Load(16)))));
|
||||
_19.Store(0, asuint(atan(asfloat(_19.Load(16)))));
|
||||
_19.Store(0, asuint(sinh(asfloat(_19.Load(16)))));
|
||||
_19.Store(0, asuint(cosh(asfloat(_19.Load(16)))));
|
||||
_19.Store(0, asuint(tanh(asfloat(_19.Load(16)))));
|
||||
_19.Store(0, asuint(atan2(asfloat(_19.Load(16)), asfloat(_19.Load(20)))));
|
||||
_19.Store(0, asuint(pow(asfloat(_19.Load(16)), asfloat(_19.Load(20)))));
|
||||
_19.Store(0, asuint(exp(asfloat(_19.Load(16)))));
|
||||
_19.Store(0, asuint(log(asfloat(_19.Load(16)))));
|
||||
_19.Store(0, asuint(exp2(asfloat(_19.Load(16)))));
|
||||
_19.Store(0, asuint(log2(asfloat(_19.Load(16)))));
|
||||
_19.Store(0, asuint(sqrt(asfloat(_19.Load(16)))));
|
||||
_19.Store(0, asuint(rsqrt(asfloat(_19.Load(16)))));
|
||||
_19.Store(0, asuint(length(asfloat(_19.Load(16)))));
|
||||
_19.Store(0, asuint(distance(asfloat(_19.Load(16)), asfloat(_19.Load(20)))));
|
||||
_19.Store(0, asuint(sign(asfloat(_19.Load(16)))));
|
||||
_19.Store(0, asuint(SPIRV_Cross_FaceForward(asfloat(_19.Load(16)), asfloat(_19.Load(20)), asfloat(_19.Load(24)))));
|
||||
_19.Store(0, asuint(SPIRV_Cross_Reflect(asfloat(_19.Load(16)), asfloat(_19.Load(20)))));
|
||||
_19.Store(0, asuint(SPIRV_Cross_Refract(asfloat(_19.Load(16)), asfloat(_19.Load(20)), asfloat(_19.Load(24)))));
|
||||
_19.Store(0, asuint(length(asfloat(_19.Load4(16)).xy)));
|
||||
_19.Store(0, asuint(distance(asfloat(_19.Load4(16)).xy, asfloat(_19.Load4(16)).zw)));
|
||||
float2 v2 = normalize(asfloat(_19.Load4(16)).xy);
|
||||
v2 = faceforward(asfloat(_19.Load4(16)).xy, asfloat(_19.Load4(16)).yz, asfloat(_19.Load4(16)).zw);
|
||||
v2 = reflect(asfloat(_19.Load4(16)).xy, asfloat(_19.Load4(16)).zw);
|
||||
v2 = refract(asfloat(_19.Load4(16)).xy, asfloat(_19.Load4(16)).yz, asfloat(_19.Load(28)));
|
||||
float3 v3 = cross(asfloat(_19.Load4(16)).xyz, asfloat(_19.Load4(16)).yzw);
|
||||
float2x2 _240 = asfloat(uint2x2(_19.Load2(64), _19.Load2(72)));
|
||||
_19.Store(0, asuint(determinant(_240)));
|
||||
float3x3 _246 = asfloat(uint3x3(_19.Load3(80), _19.Load3(96), _19.Load3(112)));
|
||||
_19.Store(0, asuint(determinant(_246)));
|
||||
float4x4 _252 = asfloat(uint4x4(_19.Load4(128), _19.Load4(144), _19.Load4(160), _19.Load4(176)));
|
||||
_19.Store(0, asuint(determinant(_252)));
|
||||
float2x2 _256 = asfloat(uint2x2(_19.Load2(64), _19.Load2(72)));
|
||||
float2x2 _257 = SPIRV_Cross_Inverse(_256);
|
||||
_19.Store2(64, asuint(_257[0]));
|
||||
_19.Store2(72, asuint(_257[1]));
|
||||
float3x3 _260 = asfloat(uint3x3(_19.Load3(80), _19.Load3(96), _19.Load3(112)));
|
||||
float3x3 _261 = SPIRV_Cross_Inverse(_260);
|
||||
_19.Store3(80, asuint(_261[0]));
|
||||
_19.Store3(96, asuint(_261[1]));
|
||||
_19.Store3(112, asuint(_261[2]));
|
||||
float4x4 _264 = asfloat(uint4x4(_19.Load4(128), _19.Load4(144), _19.Load4(160), _19.Load4(176)));
|
||||
float4x4 _265 = SPIRV_Cross_Inverse(_264);
|
||||
_19.Store4(128, asuint(_265[0]));
|
||||
_19.Store4(144, asuint(_265[1]));
|
||||
_19.Store4(160, asuint(_265[2]));
|
||||
_19.Store4(176, asuint(_265[3]));
|
||||
float tmp;
|
||||
float _271 = modf(asfloat(_19.Load(16)), tmp);
|
||||
_19.Store(0, asuint(_271));
|
||||
_19.Store(0, asuint(min(asfloat(_19.Load(16)), asfloat(_19.Load(20)))));
|
||||
_19.Store(8, min(_19.Load(48), _19.Load(52)));
|
||||
_19.Store(4, uint(min(int(_19.Load(32)), int(_19.Load(36)))));
|
||||
_19.Store(0, asuint(max(asfloat(_19.Load(16)), asfloat(_19.Load(20)))));
|
||||
_19.Store(8, max(_19.Load(48), _19.Load(52)));
|
||||
_19.Store(4, uint(max(int(_19.Load(32)), int(_19.Load(36)))));
|
||||
_19.Store(0, asuint(clamp(asfloat(_19.Load(16)), asfloat(_19.Load(20)), asfloat(_19.Load(24)))));
|
||||
_19.Store(8, clamp(_19.Load(48), _19.Load(52), _19.Load(56)));
|
||||
_19.Store(4, uint(clamp(int(_19.Load(32)), int(_19.Load(36)), int(_19.Load(40)))));
|
||||
_19.Store(0, asuint(lerp(asfloat(_19.Load(16)), asfloat(_19.Load(20)), asfloat(_19.Load(24)))));
|
||||
_19.Store(0, asuint(step(asfloat(_19.Load(16)), asfloat(_19.Load(20)))));
|
||||
_19.Store(0, asuint(smoothstep(asfloat(_19.Load(16)), asfloat(_19.Load(20)), asfloat(_19.Load(24)))));
|
||||
_19.Store(0, asuint(mad(asfloat(_19.Load(16)), asfloat(_19.Load(20)), asfloat(_19.Load(24)))));
|
||||
ResType _371;
|
||||
_371._m0 = frexp(asfloat(_19.Load(16)), _371._m1);
|
||||
int itmp = _371._m1;
|
||||
_19.Store(0, asuint(_371._m0));
|
||||
_19.Store(0, asuint(ldexp(asfloat(_19.Load(16)), itmp)));
|
||||
_19.Store(8, SPIRV_Cross_packSnorm4x8(asfloat(_19.Load4(16))));
|
||||
_19.Store(8, SPIRV_Cross_packUnorm4x8(asfloat(_19.Load4(16))));
|
||||
_19.Store(8, SPIRV_Cross_packSnorm2x16(asfloat(_19.Load4(16)).xy));
|
||||
_19.Store(8, SPIRV_Cross_packUnorm2x16(asfloat(_19.Load4(16)).xy));
|
||||
_19.Store(8, SPIRV_Cross_packHalf2x16(asfloat(_19.Load4(16)).xy));
|
||||
v2 = SPIRV_Cross_unpackSnorm2x16(_19.Load(48));
|
||||
v2 = SPIRV_Cross_unpackUnorm2x16(_19.Load(48));
|
||||
v2 = SPIRV_Cross_unpackHalf2x16(_19.Load(48));
|
||||
float4 v4 = SPIRV_Cross_unpackSnorm4x8(_19.Load(48));
|
||||
v4 = SPIRV_Cross_unpackUnorm4x8(_19.Load(48));
|
||||
_19.Store4(32, uint4(firstbitlow(int4(_19.Load4(32)))));
|
||||
_19.Store4(32, uint4(int4(firstbitlow(_19.Load4(48)))));
|
||||
_19.Store4(32, uint4(firstbithigh(int4(_19.Load4(32)))));
|
||||
_19.Store4(32, uint4(int4(firstbithigh(_19.Load4(48)))));
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void main()
|
||||
{
|
||||
comp_main();
|
||||
}
|
||||
@ -0,0 +1,35 @@
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct _8
|
||||
{
|
||||
float _m0;
|
||||
float _m1;
|
||||
};
|
||||
|
||||
struct _15
|
||||
{
|
||||
float _m0;
|
||||
int _m1;
|
||||
};
|
||||
|
||||
struct _3
|
||||
{
|
||||
float _m0;
|
||||
int _m1;
|
||||
};
|
||||
|
||||
kernel void main0(device _3& _4 [[buffer(0)]])
|
||||
{
|
||||
_8 _23;
|
||||
_23._m0 = modf(20.0, _23._m1);
|
||||
_15 _24;
|
||||
_24._m0 = frexp(40.0, _24._m1);
|
||||
_4._m0 = _23._m0;
|
||||
_4._m0 = _23._m1;
|
||||
_4._m0 = _24._m0;
|
||||
_4._m1 = _24._m1;
|
||||
}
|
||||
|
||||
282
reference/shaders-msl-no-opt/comp/glsl.std450.comp
Normal file
282
reference/shaders-msl-no-opt/comp/glsl.std450.comp
Normal file
@ -0,0 +1,282 @@
|
||||
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
||||
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
float res;
|
||||
int ires;
|
||||
uint ures;
|
||||
float4 f32;
|
||||
int4 s32;
|
||||
uint4 u32;
|
||||
float2x2 m2;
|
||||
float3x3 m3;
|
||||
float4x4 m4;
|
||||
};
|
||||
|
||||
struct ResType
|
||||
{
|
||||
float _m0;
|
||||
int _m1;
|
||||
};
|
||||
|
||||
// Implementation of the GLSL radians() function
|
||||
template<typename T>
|
||||
T radians(T d)
|
||||
{
|
||||
return d * T(0.01745329251);
|
||||
}
|
||||
|
||||
// Implementation of the GLSL degrees() function
|
||||
template<typename T>
|
||||
T degrees(T r)
|
||||
{
|
||||
return r * T(57.2957795131);
|
||||
}
|
||||
|
||||
// Implementation of the GLSL findLSB() function
|
||||
template<typename T>
|
||||
T spvFindLSB(T x)
|
||||
{
|
||||
return select(ctz(x), T(-1), x == T(0));
|
||||
}
|
||||
|
||||
// Implementation of the signed GLSL findMSB() function
|
||||
template<typename T>
|
||||
T spvFindSMSB(T x)
|
||||
{
|
||||
T v = select(x, T(-1) - x, x < T(0));
|
||||
return select(clz(T(0)) - (clz(v) + T(1)), T(-1), v == T(0));
|
||||
}
|
||||
|
||||
// Implementation of the unsigned GLSL findMSB() function
|
||||
template<typename T>
|
||||
T spvFindUMSB(T x)
|
||||
{
|
||||
return select(clz(T(0)) - (clz(x) + T(1)), T(-1), x == T(0));
|
||||
}
|
||||
|
||||
// Implementation of the GLSL sign() function for integer types
|
||||
template<typename T, typename E = typename enable_if<is_integral<T>::value>::type>
|
||||
T sign(T x)
|
||||
{
|
||||
return select(select(select(x, T(0), x == T(0)), T(1), x > T(0)), T(-1), x < T(0));
|
||||
}
|
||||
|
||||
// Returns the determinant of a 2x2 matrix.
|
||||
inline float spvDet2x2(float a1, float a2, float b1, float b2)
|
||||
{
|
||||
return a1 * b2 - b1 * a2;
|
||||
}
|
||||
|
||||
// Returns the determinant of a 3x3 matrix.
|
||||
inline float spvDet3x3(float a1, float a2, float a3, float b1, float b2, float b3, float c1, float c2, float c3)
|
||||
{
|
||||
return a1 * spvDet2x2(b2, b3, c2, c3) - b1 * spvDet2x2(a2, a3, c2, c3) + c1 * spvDet2x2(a2, a3, b2, b3);
|
||||
}
|
||||
|
||||
// Returns the inverse of a matrix, by using the algorithm of calculating the classical
|
||||
// adjoint and dividing by the determinant. The contents of the matrix are changed.
|
||||
float4x4 spvInverse4x4(float4x4 m)
|
||||
{
|
||||
float4x4 adj; // The adjoint matrix (inverse after dividing by determinant)
|
||||
|
||||
// Create the transpose of the cofactors, as the classical adjoint of the matrix.
|
||||
adj[0][0] = spvDet3x3(m[1][1], m[1][2], m[1][3], m[2][1], m[2][2], m[2][3], m[3][1], m[3][2], m[3][3]);
|
||||
adj[0][1] = -spvDet3x3(m[0][1], m[0][2], m[0][3], m[2][1], m[2][2], m[2][3], m[3][1], m[3][2], m[3][3]);
|
||||
adj[0][2] = spvDet3x3(m[0][1], m[0][2], m[0][3], m[1][1], m[1][2], m[1][3], m[3][1], m[3][2], m[3][3]);
|
||||
adj[0][3] = -spvDet3x3(m[0][1], m[0][2], m[0][3], m[1][1], m[1][2], m[1][3], m[2][1], m[2][2], m[2][3]);
|
||||
|
||||
adj[1][0] = -spvDet3x3(m[1][0], m[1][2], m[1][3], m[2][0], m[2][2], m[2][3], m[3][0], m[3][2], m[3][3]);
|
||||
adj[1][1] = spvDet3x3(m[0][0], m[0][2], m[0][3], m[2][0], m[2][2], m[2][3], m[3][0], m[3][2], m[3][3]);
|
||||
adj[1][2] = -spvDet3x3(m[0][0], m[0][2], m[0][3], m[1][0], m[1][2], m[1][3], m[3][0], m[3][2], m[3][3]);
|
||||
adj[1][3] = spvDet3x3(m[0][0], m[0][2], m[0][3], m[1][0], m[1][2], m[1][3], m[2][0], m[2][2], m[2][3]);
|
||||
|
||||
adj[2][0] = spvDet3x3(m[1][0], m[1][1], m[1][3], m[2][0], m[2][1], m[2][3], m[3][0], m[3][1], m[3][3]);
|
||||
adj[2][1] = -spvDet3x3(m[0][0], m[0][1], m[0][3], m[2][0], m[2][1], m[2][3], m[3][0], m[3][1], m[3][3]);
|
||||
adj[2][2] = spvDet3x3(m[0][0], m[0][1], m[0][3], m[1][0], m[1][1], m[1][3], m[3][0], m[3][1], m[3][3]);
|
||||
adj[2][3] = -spvDet3x3(m[0][0], m[0][1], m[0][3], m[1][0], m[1][1], m[1][3], m[2][0], m[2][1], m[2][3]);
|
||||
|
||||
adj[3][0] = -spvDet3x3(m[1][0], m[1][1], m[1][2], m[2][0], m[2][1], m[2][2], m[3][0], m[3][1], m[3][2]);
|
||||
adj[3][1] = spvDet3x3(m[0][0], m[0][1], m[0][2], m[2][0], m[2][1], m[2][2], m[3][0], m[3][1], m[3][2]);
|
||||
adj[3][2] = -spvDet3x3(m[0][0], m[0][1], m[0][2], m[1][0], m[1][1], m[1][2], m[3][0], m[3][1], m[3][2]);
|
||||
adj[3][3] = spvDet3x3(m[0][0], m[0][1], m[0][2], m[1][0], m[1][1], m[1][2], m[2][0], m[2][1], m[2][2]);
|
||||
|
||||
// Calculate the determinant as a combination of the cofactors of the first row.
|
||||
float det = (adj[0][0] * m[0][0]) + (adj[0][1] * m[1][0]) + (adj[0][2] * m[2][0]) + (adj[0][3] * m[3][0]);
|
||||
|
||||
// Divide the classical adjoint matrix by the determinant.
|
||||
// If determinant is zero, matrix is not invertable, so leave it unchanged.
|
||||
return (det != 0.0f) ? (adj * (1.0f / det)) : m;
|
||||
}
|
||||
|
||||
// Returns the inverse of a matrix, by using the algorithm of calculating the classical
|
||||
// adjoint and dividing by the determinant. The contents of the matrix are changed.
|
||||
float3x3 spvInverse3x3(float3x3 m)
|
||||
{
|
||||
float3x3 adj; // The adjoint matrix (inverse after dividing by determinant)
|
||||
|
||||
// Create the transpose of the cofactors, as the classical adjoint of the matrix.
|
||||
adj[0][0] = spvDet2x2(m[1][1], m[1][2], m[2][1], m[2][2]);
|
||||
adj[0][1] = -spvDet2x2(m[0][1], m[0][2], m[2][1], m[2][2]);
|
||||
adj[0][2] = spvDet2x2(m[0][1], m[0][2], m[1][1], m[1][2]);
|
||||
|
||||
adj[1][0] = -spvDet2x2(m[1][0], m[1][2], m[2][0], m[2][2]);
|
||||
adj[1][1] = spvDet2x2(m[0][0], m[0][2], m[2][0], m[2][2]);
|
||||
adj[1][2] = -spvDet2x2(m[0][0], m[0][2], m[1][0], m[1][2]);
|
||||
|
||||
adj[2][0] = spvDet2x2(m[1][0], m[1][1], m[2][0], m[2][1]);
|
||||
adj[2][1] = -spvDet2x2(m[0][0], m[0][1], m[2][0], m[2][1]);
|
||||
adj[2][2] = spvDet2x2(m[0][0], m[0][1], m[1][0], m[1][1]);
|
||||
|
||||
// Calculate the determinant as a combination of the cofactors of the first row.
|
||||
float det = (adj[0][0] * m[0][0]) + (adj[0][1] * m[1][0]) + (adj[0][2] * m[2][0]);
|
||||
|
||||
// Divide the classical adjoint matrix by the determinant.
|
||||
// If determinant is zero, matrix is not invertable, so leave it unchanged.
|
||||
return (det != 0.0f) ? (adj * (1.0f / det)) : m;
|
||||
}
|
||||
|
||||
// Returns the inverse of a matrix, by using the algorithm of calculating the classical
|
||||
// adjoint and dividing by the determinant. The contents of the matrix are changed.
|
||||
float2x2 spvInverse2x2(float2x2 m)
|
||||
{
|
||||
float2x2 adj; // The adjoint matrix (inverse after dividing by determinant)
|
||||
|
||||
// Create the transpose of the cofactors, as the classical adjoint of the matrix.
|
||||
adj[0][0] = m[1][1];
|
||||
adj[0][1] = -m[0][1];
|
||||
|
||||
adj[1][0] = -m[1][0];
|
||||
adj[1][1] = m[0][0];
|
||||
|
||||
// Calculate the determinant as a combination of the cofactors of the first row.
|
||||
float det = (adj[0][0] * m[0][0]) + (adj[0][1] * m[1][0]);
|
||||
|
||||
// Divide the classical adjoint matrix by the determinant.
|
||||
// If determinant is zero, matrix is not invertable, so leave it unchanged.
|
||||
return (det != 0.0f) ? (adj * (1.0f / det)) : m;
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline T spvReflect(T i, T n)
|
||||
{
|
||||
return i - T(2) * i * n * n;
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline T spvRefract(T i, T n, T eta)
|
||||
{
|
||||
T NoI = n * i;
|
||||
T NoI2 = NoI * NoI;
|
||||
T k = T(1) - eta * eta * (T(1) - NoI2);
|
||||
if (k < T(0))
|
||||
{
|
||||
return T(0);
|
||||
}
|
||||
else
|
||||
{
|
||||
return eta * i - (eta * NoI + sqrt(k)) * n;
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline T spvFaceForward(T n, T i, T nref)
|
||||
{
|
||||
return i * nref < T(0) ? n : -n;
|
||||
}
|
||||
|
||||
kernel void main0(device SSBO& _19 [[buffer(0)]])
|
||||
{
|
||||
_19.res = round(_19.f32.x);
|
||||
_19.res = rint(_19.f32.x);
|
||||
_19.res = trunc(_19.f32.x);
|
||||
_19.res = abs(_19.f32.x);
|
||||
_19.ires = abs(_19.s32.x);
|
||||
_19.res = sign(_19.f32.x);
|
||||
_19.ires = sign(_19.s32.x);
|
||||
_19.res = floor(_19.f32.x);
|
||||
_19.res = ceil(_19.f32.x);
|
||||
_19.res = fract(_19.f32.x);
|
||||
_19.res = radians(_19.f32.x);
|
||||
_19.res = degrees(_19.f32.x);
|
||||
_19.res = sin(_19.f32.x);
|
||||
_19.res = cos(_19.f32.x);
|
||||
_19.res = tan(_19.f32.x);
|
||||
_19.res = asin(_19.f32.x);
|
||||
_19.res = acos(_19.f32.x);
|
||||
_19.res = atan(_19.f32.x);
|
||||
_19.res = sinh(_19.f32.x);
|
||||
_19.res = cosh(_19.f32.x);
|
||||
_19.res = tanh(_19.f32.x);
|
||||
_19.res = asinh(_19.f32.x);
|
||||
_19.res = acosh(_19.f32.x);
|
||||
_19.res = atanh(_19.f32.x);
|
||||
_19.res = atan2(_19.f32.x, _19.f32.y);
|
||||
_19.res = pow(_19.f32.x, _19.f32.y);
|
||||
_19.res = exp(_19.f32.x);
|
||||
_19.res = log(_19.f32.x);
|
||||
_19.res = exp2(_19.f32.x);
|
||||
_19.res = log2(_19.f32.x);
|
||||
_19.res = sqrt(_19.f32.x);
|
||||
_19.res = rsqrt(_19.f32.x);
|
||||
_19.res = abs(_19.f32.x);
|
||||
_19.res = abs(_19.f32.x - _19.f32.y);
|
||||
_19.res = sign(_19.f32.x);
|
||||
_19.res = spvFaceForward(_19.f32.x, _19.f32.y, _19.f32.z);
|
||||
_19.res = spvReflect(_19.f32.x, _19.f32.y);
|
||||
_19.res = spvRefract(_19.f32.x, _19.f32.y, _19.f32.z);
|
||||
_19.res = length(_19.f32.xy);
|
||||
_19.res = distance(_19.f32.xy, _19.f32.zw);
|
||||
float2 v2 = normalize(_19.f32.xy);
|
||||
v2 = faceforward(_19.f32.xy, _19.f32.yz, _19.f32.zw);
|
||||
v2 = reflect(_19.f32.xy, _19.f32.zw);
|
||||
v2 = refract(_19.f32.xy, _19.f32.yz, _19.f32.w);
|
||||
float3 v3 = cross(_19.f32.xyz, _19.f32.yzw);
|
||||
_19.res = determinant(_19.m2);
|
||||
_19.res = determinant(_19.m3);
|
||||
_19.res = determinant(_19.m4);
|
||||
_19.m2 = spvInverse2x2(_19.m2);
|
||||
_19.m3 = spvInverse3x3(_19.m3);
|
||||
_19.m4 = spvInverse4x4(_19.m4);
|
||||
float tmp;
|
||||
float _287 = modf(_19.f32.x, tmp);
|
||||
_19.res = _287;
|
||||
_19.res = fast::min(_19.f32.x, _19.f32.y);
|
||||
_19.ures = min(_19.u32.x, _19.u32.y);
|
||||
_19.ires = min(_19.s32.x, _19.s32.y);
|
||||
_19.res = fast::max(_19.f32.x, _19.f32.y);
|
||||
_19.ures = max(_19.u32.x, _19.u32.y);
|
||||
_19.ires = max(_19.s32.x, _19.s32.y);
|
||||
_19.res = fast::clamp(_19.f32.x, _19.f32.y, _19.f32.z);
|
||||
_19.ures = clamp(_19.u32.x, _19.u32.y, _19.u32.z);
|
||||
_19.ires = clamp(_19.s32.x, _19.s32.y, _19.s32.z);
|
||||
_19.res = mix(_19.f32.x, _19.f32.y, _19.f32.z);
|
||||
_19.res = step(_19.f32.x, _19.f32.y);
|
||||
_19.res = smoothstep(_19.f32.x, _19.f32.y, _19.f32.z);
|
||||
_19.res = fma(_19.f32.x, _19.f32.y, _19.f32.z);
|
||||
ResType _387;
|
||||
_387._m0 = frexp(_19.f32.x, _387._m1);
|
||||
int itmp = _387._m1;
|
||||
_19.res = _387._m0;
|
||||
_19.res = ldexp(_19.f32.x, itmp);
|
||||
_19.ures = pack_float_to_snorm4x8(_19.f32);
|
||||
_19.ures = pack_float_to_unorm4x8(_19.f32);
|
||||
_19.ures = pack_float_to_snorm2x16(_19.f32.xy);
|
||||
_19.ures = pack_float_to_unorm2x16(_19.f32.xy);
|
||||
_19.ures = as_type<uint>(half2(_19.f32.xy));
|
||||
v2 = unpack_snorm2x16_to_float(_19.u32.x);
|
||||
v2 = unpack_unorm2x16_to_float(_19.u32.x);
|
||||
v2 = float2(as_type<half2>(_19.u32.x));
|
||||
float4 v4 = unpack_snorm4x8_to_float(_19.u32.x);
|
||||
v4 = unpack_unorm4x8_to_float(_19.u32.x);
|
||||
_19.s32 = spvFindLSB(_19.s32);
|
||||
_19.s32 = int4(spvFindLSB(_19.u32));
|
||||
_19.s32 = spvFindSMSB(_19.s32);
|
||||
_19.s32 = int4(spvFindUMSB(_19.u32));
|
||||
}
|
||||
|
||||
@ -0,0 +1,33 @@
|
||||
#version 450
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
||||
struct _8
|
||||
{
|
||||
float _m0;
|
||||
float _m1;
|
||||
};
|
||||
|
||||
struct _15
|
||||
{
|
||||
float _m0;
|
||||
int _m1;
|
||||
};
|
||||
|
||||
layout(binding = 0, std430) buffer _3_4
|
||||
{
|
||||
float _m0;
|
||||
int _m1;
|
||||
} _4;
|
||||
|
||||
void main()
|
||||
{
|
||||
_8 _23;
|
||||
_23._m0 = modf(20.0, _23._m1);
|
||||
_15 _24;
|
||||
_24._m0 = frexp(40.0, _24._m1);
|
||||
_4._m0 = _23._m0;
|
||||
_4._m0 = _23._m1;
|
||||
_4._m0 = _24._m0;
|
||||
_4._m1 = _24._m1;
|
||||
}
|
||||
|
||||
112
reference/shaders-no-opt/comp/glsl.std450.comp
Normal file
112
reference/shaders-no-opt/comp/glsl.std450.comp
Normal file
@ -0,0 +1,112 @@
|
||||
#version 450
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
||||
struct ResType
|
||||
{
|
||||
float _m0;
|
||||
int _m1;
|
||||
};
|
||||
|
||||
layout(binding = 0, std430) buffer SSBO
|
||||
{
|
||||
float res;
|
||||
int ires;
|
||||
uint ures;
|
||||
vec4 f32;
|
||||
ivec4 s32;
|
||||
uvec4 u32;
|
||||
mat2 m2;
|
||||
mat3 m3;
|
||||
mat4 m4;
|
||||
} _19;
|
||||
|
||||
void main()
|
||||
{
|
||||
_19.res = round(_19.f32.x);
|
||||
_19.res = roundEven(_19.f32.x);
|
||||
_19.res = trunc(_19.f32.x);
|
||||
_19.res = abs(_19.f32.x);
|
||||
_19.ires = abs(_19.s32.x);
|
||||
_19.res = sign(_19.f32.x);
|
||||
_19.ires = sign(_19.s32.x);
|
||||
_19.res = floor(_19.f32.x);
|
||||
_19.res = ceil(_19.f32.x);
|
||||
_19.res = fract(_19.f32.x);
|
||||
_19.res = radians(_19.f32.x);
|
||||
_19.res = degrees(_19.f32.x);
|
||||
_19.res = sin(_19.f32.x);
|
||||
_19.res = cos(_19.f32.x);
|
||||
_19.res = tan(_19.f32.x);
|
||||
_19.res = asin(_19.f32.x);
|
||||
_19.res = acos(_19.f32.x);
|
||||
_19.res = atan(_19.f32.x);
|
||||
_19.res = sinh(_19.f32.x);
|
||||
_19.res = cosh(_19.f32.x);
|
||||
_19.res = tanh(_19.f32.x);
|
||||
_19.res = asinh(_19.f32.x);
|
||||
_19.res = acosh(_19.f32.x);
|
||||
_19.res = atanh(_19.f32.x);
|
||||
_19.res = atan(_19.f32.x, _19.f32.y);
|
||||
_19.res = pow(_19.f32.x, _19.f32.y);
|
||||
_19.res = exp(_19.f32.x);
|
||||
_19.res = log(_19.f32.x);
|
||||
_19.res = exp2(_19.f32.x);
|
||||
_19.res = log2(_19.f32.x);
|
||||
_19.res = sqrt(_19.f32.x);
|
||||
_19.res = inversesqrt(_19.f32.x);
|
||||
_19.res = length(_19.f32.x);
|
||||
_19.res = distance(_19.f32.x, _19.f32.y);
|
||||
_19.res = normalize(_19.f32.x);
|
||||
_19.res = faceforward(_19.f32.x, _19.f32.y, _19.f32.z);
|
||||
_19.res = reflect(_19.f32.x, _19.f32.y);
|
||||
_19.res = refract(_19.f32.x, _19.f32.y, _19.f32.z);
|
||||
_19.res = length(_19.f32.xy);
|
||||
_19.res = distance(_19.f32.xy, _19.f32.zw);
|
||||
vec2 v2 = normalize(_19.f32.xy);
|
||||
v2 = faceforward(_19.f32.xy, _19.f32.yz, _19.f32.zw);
|
||||
v2 = reflect(_19.f32.xy, _19.f32.zw);
|
||||
v2 = refract(_19.f32.xy, _19.f32.yz, _19.f32.w);
|
||||
vec3 v3 = cross(_19.f32.xyz, _19.f32.yzw);
|
||||
_19.res = determinant(_19.m2);
|
||||
_19.res = determinant(_19.m3);
|
||||
_19.res = determinant(_19.m4);
|
||||
_19.m2 = inverse(_19.m2);
|
||||
_19.m3 = inverse(_19.m3);
|
||||
_19.m4 = inverse(_19.m4);
|
||||
float tmp;
|
||||
float _287 = modf(_19.f32.x, tmp);
|
||||
_19.res = _287;
|
||||
_19.res = min(_19.f32.x, _19.f32.y);
|
||||
_19.ures = min(_19.u32.x, _19.u32.y);
|
||||
_19.ires = min(_19.s32.x, _19.s32.y);
|
||||
_19.res = max(_19.f32.x, _19.f32.y);
|
||||
_19.ures = max(_19.u32.x, _19.u32.y);
|
||||
_19.ires = max(_19.s32.x, _19.s32.y);
|
||||
_19.res = clamp(_19.f32.x, _19.f32.y, _19.f32.z);
|
||||
_19.ures = clamp(_19.u32.x, _19.u32.y, _19.u32.z);
|
||||
_19.ires = clamp(_19.s32.x, _19.s32.y, _19.s32.z);
|
||||
_19.res = mix(_19.f32.x, _19.f32.y, _19.f32.z);
|
||||
_19.res = step(_19.f32.x, _19.f32.y);
|
||||
_19.res = smoothstep(_19.f32.x, _19.f32.y, _19.f32.z);
|
||||
_19.res = fma(_19.f32.x, _19.f32.y, _19.f32.z);
|
||||
ResType _387;
|
||||
_387._m0 = frexp(_19.f32.x, _387._m1);
|
||||
int itmp = _387._m1;
|
||||
_19.res = _387._m0;
|
||||
_19.res = ldexp(_19.f32.x, itmp);
|
||||
_19.ures = packSnorm4x8(_19.f32);
|
||||
_19.ures = packUnorm4x8(_19.f32);
|
||||
_19.ures = packSnorm2x16(_19.f32.xy);
|
||||
_19.ures = packUnorm2x16(_19.f32.xy);
|
||||
_19.ures = packHalf2x16(_19.f32.xy);
|
||||
v2 = unpackSnorm2x16(_19.u32.x);
|
||||
v2 = unpackUnorm2x16(_19.u32.x);
|
||||
v2 = unpackHalf2x16(_19.u32.x);
|
||||
vec4 v4 = unpackSnorm4x8(_19.u32.x);
|
||||
v4 = unpackUnorm4x8(_19.u32.x);
|
||||
_19.s32 = findLSB(_19.s32);
|
||||
_19.s32 = findLSB(_19.u32);
|
||||
_19.s32 = findMSB(_19.s32);
|
||||
_19.s32 = findMSB(_19.u32);
|
||||
}
|
||||
|
||||
@ -0,0 +1,55 @@
|
||||
; SPIR-V
|
||||
; Version: 1.0
|
||||
; Generator: Khronos Glslang Reference Front End; 7
|
||||
; Bound: 45
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main"
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpSource GLSL 450
|
||||
OpMemberDecorate %SSBO 0 Offset 0
|
||||
OpMemberDecorate %SSBO 1 Offset 4
|
||||
OpDecorate %SSBO BufferBlock
|
||||
OpDecorate %_ DescriptorSet 0
|
||||
OpDecorate %_ Binding 0
|
||||
%void = OpTypeVoid
|
||||
%3 = OpTypeFunction %void
|
||||
%float = OpTypeFloat 32
|
||||
%ResTypeMod = OpTypeStruct %float %float
|
||||
%_ptr_Function_ResTypeMod = OpTypePointer Function %ResTypeMod
|
||||
%int = OpTypeInt 32 1
|
||||
%int_0 = OpConstant %int 0
|
||||
%float_20 = OpConstant %float 20
|
||||
%int_1 = OpConstant %int 1
|
||||
%_ptr_Function_float = OpTypePointer Function %float
|
||||
%ResTypeFrexp = OpTypeStruct %float %int
|
||||
%_ptr_Function_ResTypeFrexp = OpTypePointer Function %ResTypeFrexp
|
||||
%float_40 = OpConstant %float 40
|
||||
%_ptr_Function_int = OpTypePointer Function %int
|
||||
%SSBO = OpTypeStruct %float %int
|
||||
%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO
|
||||
%_ = OpVariable %_ptr_Uniform_SSBO Uniform
|
||||
%_ptr_Uniform_float = OpTypePointer Uniform %float
|
||||
%_ptr_Uniform_int = OpTypePointer Uniform %int
|
||||
%main = OpFunction %void None %3
|
||||
%5 = OpLabel
|
||||
%modres = OpExtInst %ResTypeMod %1 ModfStruct %float_20
|
||||
%frexpres = OpExtInst %ResTypeFrexp %1 FrexpStruct %float_40
|
||||
|
||||
%modres_f = OpCompositeExtract %float %modres 0
|
||||
%modres_i = OpCompositeExtract %float %modres 1
|
||||
%frexpres_f = OpCompositeExtract %float %frexpres 0
|
||||
%frexpres_i = OpCompositeExtract %int %frexpres 1
|
||||
|
||||
%float_ptr = OpAccessChain %_ptr_Uniform_float %_ %int_0
|
||||
%int_ptr = OpAccessChain %_ptr_Uniform_int %_ %int_1
|
||||
|
||||
OpStore %float_ptr %modres_f
|
||||
OpStore %float_ptr %modres_i
|
||||
OpStore %float_ptr %frexpres_f
|
||||
OpStore %int_ptr %frexpres_i
|
||||
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
130
shaders-hlsl-no-opt/comp/glsl.std450.fxconly.comp
Normal file
130
shaders-hlsl-no-opt/comp/glsl.std450.fxconly.comp
Normal file
@ -0,0 +1,130 @@
|
||||
#version 450
|
||||
layout(local_size_x = 1) in;
|
||||
|
||||
layout(binding = 0, std430) buffer SSBO
|
||||
{
|
||||
float res;
|
||||
int ires;
|
||||
uint ures;
|
||||
|
||||
vec4 f32;
|
||||
ivec4 s32;
|
||||
uvec4 u32;
|
||||
|
||||
mat2 m2;
|
||||
mat3 m3;
|
||||
mat4 m4;
|
||||
};
|
||||
|
||||
void main()
|
||||
{
|
||||
float tmp;
|
||||
vec2 v2;
|
||||
vec3 v3;
|
||||
vec4 v4;
|
||||
int itmp;
|
||||
|
||||
res = round(f32.x);
|
||||
//res = roundEven(f32.x);
|
||||
res = trunc(f32.x);
|
||||
res = abs(f32.x);
|
||||
ires = abs(s32.x);
|
||||
res = sign(f32.x);
|
||||
ires = sign(s32.x);
|
||||
res = floor(f32.x);
|
||||
res = ceil(f32.x);
|
||||
res = fract(f32.x);
|
||||
res = radians(f32.x);
|
||||
res = degrees(f32.x);
|
||||
res = sin(f32.x);
|
||||
res = cos(f32.x);
|
||||
res = tan(f32.x);
|
||||
res = asin(f32.x);
|
||||
res = acos(f32.x);
|
||||
res = atan(f32.x);
|
||||
res = sinh(f32.x);
|
||||
res = cosh(f32.x);
|
||||
res = tanh(f32.x);
|
||||
//res = asinh(f32.x);
|
||||
//res = acosh(f32.x);
|
||||
//res = atanh(f32.x);
|
||||
res = atan(f32.x, f32.y);
|
||||
res = pow(f32.x, f32.y);
|
||||
res = exp(f32.x);
|
||||
res = log(f32.x);
|
||||
res = exp2(f32.x);
|
||||
res = log2(f32.x);
|
||||
res = sqrt(f32.x);
|
||||
res = inversesqrt(f32.x);
|
||||
|
||||
res = length(f32.x);
|
||||
res = distance(f32.x, f32.y);
|
||||
res = normalize(f32.x);
|
||||
res = faceforward(f32.x, f32.y, f32.z);
|
||||
res = reflect(f32.x, f32.y);
|
||||
res = refract(f32.x, f32.y, f32.z);
|
||||
|
||||
res = length(f32.xy);
|
||||
res = distance(f32.xy, f32.zw);
|
||||
v2 = normalize(f32.xy);
|
||||
v2 = faceforward(f32.xy, f32.yz, f32.zw);
|
||||
v2 = reflect(f32.xy, f32.zw);
|
||||
v2 = refract(f32.xy, f32.yz, f32.w);
|
||||
|
||||
v3 = cross(f32.xyz, f32.yzw);
|
||||
|
||||
res = determinant(m2);
|
||||
res = determinant(m3);
|
||||
res = determinant(m4);
|
||||
m2 = inverse(m2);
|
||||
m3 = inverse(m3);
|
||||
m4 = inverse(m4);
|
||||
|
||||
res = modf(f32.x, tmp);
|
||||
// ModfStruct
|
||||
|
||||
res = min(f32.x, f32.y);
|
||||
ures = min(u32.x, u32.y);
|
||||
ires = min(s32.x, s32.y);
|
||||
res = max(f32.x, f32.y);
|
||||
ures = max(u32.x, u32.y);
|
||||
ires = max(s32.x, s32.y);
|
||||
|
||||
res = clamp(f32.x, f32.y, f32.z);
|
||||
ures = clamp(u32.x, u32.y, u32.z);
|
||||
ires = clamp(s32.x, s32.y, s32.z);
|
||||
|
||||
res = mix(f32.x, f32.y, f32.z);
|
||||
res = step(f32.x, f32.y);
|
||||
res = smoothstep(f32.x, f32.y, f32.z);
|
||||
res = fma(f32.x, f32.y, f32.z);
|
||||
|
||||
res = frexp(f32.x, itmp);
|
||||
|
||||
// FrexpStruct
|
||||
res = ldexp(f32.x, itmp);
|
||||
|
||||
ures = packSnorm4x8(f32);
|
||||
ures = packUnorm4x8(f32);
|
||||
ures = packSnorm2x16(f32.xy);
|
||||
ures = packUnorm2x16(f32.xy);
|
||||
ures = packHalf2x16(f32.xy);
|
||||
// packDouble2x32
|
||||
|
||||
v2 = unpackSnorm2x16(u32.x);
|
||||
v2 = unpackUnorm2x16(u32.x);
|
||||
v2 = unpackHalf2x16(u32.x);
|
||||
v4 = unpackSnorm4x8(u32.x);
|
||||
v4 = unpackUnorm4x8(u32.x);
|
||||
// unpackDouble2x32
|
||||
|
||||
s32 = findLSB(s32);
|
||||
s32 = findLSB(u32);
|
||||
s32 = findMSB(s32);
|
||||
s32 = findMSB(u32);
|
||||
|
||||
// interpolateAtSample
|
||||
// interpolateAtOffset
|
||||
|
||||
// NMin, NMax, NClamp
|
||||
}
|
||||
@ -0,0 +1,55 @@
|
||||
; SPIR-V
|
||||
; Version: 1.0
|
||||
; Generator: Khronos Glslang Reference Front End; 7
|
||||
; Bound: 45
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main"
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpSource GLSL 450
|
||||
OpMemberDecorate %SSBO 0 Offset 0
|
||||
OpMemberDecorate %SSBO 1 Offset 4
|
||||
OpDecorate %SSBO BufferBlock
|
||||
OpDecorate %_ DescriptorSet 0
|
||||
OpDecorate %_ Binding 0
|
||||
%void = OpTypeVoid
|
||||
%3 = OpTypeFunction %void
|
||||
%float = OpTypeFloat 32
|
||||
%ResTypeMod = OpTypeStruct %float %float
|
||||
%_ptr_Function_ResTypeMod = OpTypePointer Function %ResTypeMod
|
||||
%int = OpTypeInt 32 1
|
||||
%int_0 = OpConstant %int 0
|
||||
%float_20 = OpConstant %float 20
|
||||
%int_1 = OpConstant %int 1
|
||||
%_ptr_Function_float = OpTypePointer Function %float
|
||||
%ResTypeFrexp = OpTypeStruct %float %int
|
||||
%_ptr_Function_ResTypeFrexp = OpTypePointer Function %ResTypeFrexp
|
||||
%float_40 = OpConstant %float 40
|
||||
%_ptr_Function_int = OpTypePointer Function %int
|
||||
%SSBO = OpTypeStruct %float %int
|
||||
%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO
|
||||
%_ = OpVariable %_ptr_Uniform_SSBO Uniform
|
||||
%_ptr_Uniform_float = OpTypePointer Uniform %float
|
||||
%_ptr_Uniform_int = OpTypePointer Uniform %int
|
||||
%main = OpFunction %void None %3
|
||||
%5 = OpLabel
|
||||
%modres = OpExtInst %ResTypeMod %1 ModfStruct %float_20
|
||||
%frexpres = OpExtInst %ResTypeFrexp %1 FrexpStruct %float_40
|
||||
|
||||
%modres_f = OpCompositeExtract %float %modres 0
|
||||
%modres_i = OpCompositeExtract %float %modres 1
|
||||
%frexpres_f = OpCompositeExtract %float %frexpres 0
|
||||
%frexpres_i = OpCompositeExtract %int %frexpres 1
|
||||
|
||||
%float_ptr = OpAccessChain %_ptr_Uniform_float %_ %int_0
|
||||
%int_ptr = OpAccessChain %_ptr_Uniform_int %_ %int_1
|
||||
|
||||
OpStore %float_ptr %modres_f
|
||||
OpStore %float_ptr %modres_i
|
||||
OpStore %float_ptr %frexpres_f
|
||||
OpStore %int_ptr %frexpres_i
|
||||
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
129
shaders-msl-no-opt/comp/glsl.std450.comp
Normal file
129
shaders-msl-no-opt/comp/glsl.std450.comp
Normal file
@ -0,0 +1,129 @@
|
||||
#version 450
|
||||
layout(local_size_x = 1) in;
|
||||
|
||||
layout(binding = 0, std430) buffer SSBO
|
||||
{
|
||||
float res;
|
||||
int ires;
|
||||
uint ures;
|
||||
|
||||
vec4 f32;
|
||||
ivec4 s32;
|
||||
uvec4 u32;
|
||||
|
||||
mat2 m2;
|
||||
mat3 m3;
|
||||
mat4 m4;
|
||||
};
|
||||
|
||||
void main()
|
||||
{
|
||||
float tmp;
|
||||
vec2 v2;
|
||||
vec3 v3;
|
||||
vec4 v4;
|
||||
int itmp;
|
||||
|
||||
res = round(f32.x);
|
||||
res = roundEven(f32.x);
|
||||
res = trunc(f32.x);
|
||||
res = abs(f32.x);
|
||||
ires = abs(s32.x);
|
||||
res = sign(f32.x);
|
||||
ires = sign(s32.x);
|
||||
res = floor(f32.x);
|
||||
res = ceil(f32.x);
|
||||
res = fract(f32.x);
|
||||
res = radians(f32.x);
|
||||
res = degrees(f32.x);
|
||||
res = sin(f32.x);
|
||||
res = cos(f32.x);
|
||||
res = tan(f32.x);
|
||||
res = asin(f32.x);
|
||||
res = acos(f32.x);
|
||||
res = atan(f32.x);
|
||||
res = sinh(f32.x);
|
||||
res = cosh(f32.x);
|
||||
res = tanh(f32.x);
|
||||
res = asinh(f32.x);
|
||||
res = acosh(f32.x);
|
||||
res = atanh(f32.x);
|
||||
res = atan(f32.x, f32.y);
|
||||
res = pow(f32.x, f32.y);
|
||||
res = exp(f32.x);
|
||||
res = log(f32.x);
|
||||
res = exp2(f32.x);
|
||||
res = log2(f32.x);
|
||||
res = sqrt(f32.x);
|
||||
res = inversesqrt(f32.x);
|
||||
|
||||
res = length(f32.x);
|
||||
res = distance(f32.x, f32.y);
|
||||
res = normalize(f32.x);
|
||||
res = faceforward(f32.x, f32.y, f32.z);
|
||||
res = reflect(f32.x, f32.y);
|
||||
res = refract(f32.x, f32.y, f32.z);
|
||||
|
||||
res = length(f32.xy);
|
||||
res = distance(f32.xy, f32.zw);
|
||||
v2 = normalize(f32.xy);
|
||||
v2 = faceforward(f32.xy, f32.yz, f32.zw);
|
||||
v2 = reflect(f32.xy, f32.zw);
|
||||
v2 = refract(f32.xy, f32.yz, f32.w);
|
||||
|
||||
v3 = cross(f32.xyz, f32.yzw);
|
||||
|
||||
res = determinant(m2);
|
||||
res = determinant(m3);
|
||||
res = determinant(m4);
|
||||
m2 = inverse(m2);
|
||||
m3 = inverse(m3);
|
||||
m4 = inverse(m4);
|
||||
|
||||
res = modf(f32.x, tmp);
|
||||
// ModfStruct
|
||||
|
||||
res = min(f32.x, f32.y);
|
||||
ures = min(u32.x, u32.y);
|
||||
ires = min(s32.x, s32.y);
|
||||
res = max(f32.x, f32.y);
|
||||
ures = max(u32.x, u32.y);
|
||||
ires = max(s32.x, s32.y);
|
||||
|
||||
res = clamp(f32.x, f32.y, f32.z);
|
||||
ures = clamp(u32.x, u32.y, u32.z);
|
||||
ires = clamp(s32.x, s32.y, s32.z);
|
||||
|
||||
res = mix(f32.x, f32.y, f32.z);
|
||||
res = step(f32.x, f32.y);
|
||||
res = smoothstep(f32.x, f32.y, f32.z);
|
||||
res = fma(f32.x, f32.y, f32.z);
|
||||
|
||||
res = frexp(f32.x, itmp);
|
||||
// FrexpStruct
|
||||
res = ldexp(f32.x, itmp);
|
||||
|
||||
ures = packSnorm4x8(f32);
|
||||
ures = packUnorm4x8(f32);
|
||||
ures = packSnorm2x16(f32.xy);
|
||||
ures = packUnorm2x16(f32.xy);
|
||||
ures = packHalf2x16(f32.xy);
|
||||
// packDouble2x32
|
||||
|
||||
v2 = unpackSnorm2x16(u32.x);
|
||||
v2 = unpackUnorm2x16(u32.x);
|
||||
v2 = unpackHalf2x16(u32.x);
|
||||
v4 = unpackSnorm4x8(u32.x);
|
||||
v4 = unpackUnorm4x8(u32.x);
|
||||
// unpackDouble2x32
|
||||
|
||||
s32 = findLSB(s32);
|
||||
s32 = findLSB(u32);
|
||||
s32 = findMSB(s32);
|
||||
s32 = findMSB(u32);
|
||||
|
||||
// interpolateAtSample
|
||||
// interpolateAtOffset
|
||||
|
||||
// NMin, NMax, NClamp
|
||||
}
|
||||
@ -0,0 +1,55 @@
|
||||
; SPIR-V
|
||||
; Version: 1.0
|
||||
; Generator: Khronos Glslang Reference Front End; 7
|
||||
; Bound: 45
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main"
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpSource GLSL 450
|
||||
OpMemberDecorate %SSBO 0 Offset 0
|
||||
OpMemberDecorate %SSBO 1 Offset 4
|
||||
OpDecorate %SSBO BufferBlock
|
||||
OpDecorate %_ DescriptorSet 0
|
||||
OpDecorate %_ Binding 0
|
||||
%void = OpTypeVoid
|
||||
%3 = OpTypeFunction %void
|
||||
%float = OpTypeFloat 32
|
||||
%ResTypeMod = OpTypeStruct %float %float
|
||||
%_ptr_Function_ResTypeMod = OpTypePointer Function %ResTypeMod
|
||||
%int = OpTypeInt 32 1
|
||||
%int_0 = OpConstant %int 0
|
||||
%float_20 = OpConstant %float 20
|
||||
%int_1 = OpConstant %int 1
|
||||
%_ptr_Function_float = OpTypePointer Function %float
|
||||
%ResTypeFrexp = OpTypeStruct %float %int
|
||||
%_ptr_Function_ResTypeFrexp = OpTypePointer Function %ResTypeFrexp
|
||||
%float_40 = OpConstant %float 40
|
||||
%_ptr_Function_int = OpTypePointer Function %int
|
||||
%SSBO = OpTypeStruct %float %int
|
||||
%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO
|
||||
%_ = OpVariable %_ptr_Uniform_SSBO Uniform
|
||||
%_ptr_Uniform_float = OpTypePointer Uniform %float
|
||||
%_ptr_Uniform_int = OpTypePointer Uniform %int
|
||||
%main = OpFunction %void None %3
|
||||
%5 = OpLabel
|
||||
%modres = OpExtInst %ResTypeMod %1 ModfStruct %float_20
|
||||
%frexpres = OpExtInst %ResTypeFrexp %1 FrexpStruct %float_40
|
||||
|
||||
%modres_f = OpCompositeExtract %float %modres 0
|
||||
%modres_i = OpCompositeExtract %float %modres 1
|
||||
%frexpres_f = OpCompositeExtract %float %frexpres 0
|
||||
%frexpres_i = OpCompositeExtract %int %frexpres 1
|
||||
|
||||
%float_ptr = OpAccessChain %_ptr_Uniform_float %_ %int_0
|
||||
%int_ptr = OpAccessChain %_ptr_Uniform_int %_ %int_1
|
||||
|
||||
OpStore %float_ptr %modres_f
|
||||
OpStore %float_ptr %modres_i
|
||||
OpStore %float_ptr %frexpres_f
|
||||
OpStore %int_ptr %frexpres_i
|
||||
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
129
shaders-no-opt/comp/glsl.std450.comp
Normal file
129
shaders-no-opt/comp/glsl.std450.comp
Normal file
@ -0,0 +1,129 @@
|
||||
#version 450
|
||||
layout(local_size_x = 1) in;
|
||||
|
||||
layout(binding = 0, std430) buffer SSBO
|
||||
{
|
||||
float res;
|
||||
int ires;
|
||||
uint ures;
|
||||
|
||||
vec4 f32;
|
||||
ivec4 s32;
|
||||
uvec4 u32;
|
||||
|
||||
mat2 m2;
|
||||
mat3 m3;
|
||||
mat4 m4;
|
||||
};
|
||||
|
||||
void main()
|
||||
{
|
||||
float tmp;
|
||||
vec2 v2;
|
||||
vec3 v3;
|
||||
vec4 v4;
|
||||
int itmp;
|
||||
|
||||
res = round(f32.x);
|
||||
res = roundEven(f32.x);
|
||||
res = trunc(f32.x);
|
||||
res = abs(f32.x);
|
||||
ires = abs(s32.x);
|
||||
res = sign(f32.x);
|
||||
ires = sign(s32.x);
|
||||
res = floor(f32.x);
|
||||
res = ceil(f32.x);
|
||||
res = fract(f32.x);
|
||||
res = radians(f32.x);
|
||||
res = degrees(f32.x);
|
||||
res = sin(f32.x);
|
||||
res = cos(f32.x);
|
||||
res = tan(f32.x);
|
||||
res = asin(f32.x);
|
||||
res = acos(f32.x);
|
||||
res = atan(f32.x);
|
||||
res = sinh(f32.x);
|
||||
res = cosh(f32.x);
|
||||
res = tanh(f32.x);
|
||||
res = asinh(f32.x);
|
||||
res = acosh(f32.x);
|
||||
res = atanh(f32.x);
|
||||
res = atan(f32.x, f32.y);
|
||||
res = pow(f32.x, f32.y);
|
||||
res = exp(f32.x);
|
||||
res = log(f32.x);
|
||||
res = exp2(f32.x);
|
||||
res = log2(f32.x);
|
||||
res = sqrt(f32.x);
|
||||
res = inversesqrt(f32.x);
|
||||
|
||||
res = length(f32.x);
|
||||
res = distance(f32.x, f32.y);
|
||||
res = normalize(f32.x);
|
||||
res = faceforward(f32.x, f32.y, f32.z);
|
||||
res = reflect(f32.x, f32.y);
|
||||
res = refract(f32.x, f32.y, f32.z);
|
||||
|
||||
res = length(f32.xy);
|
||||
res = distance(f32.xy, f32.zw);
|
||||
v2 = normalize(f32.xy);
|
||||
v2 = faceforward(f32.xy, f32.yz, f32.zw);
|
||||
v2 = reflect(f32.xy, f32.zw);
|
||||
v2 = refract(f32.xy, f32.yz, f32.w);
|
||||
|
||||
v3 = cross(f32.xyz, f32.yzw);
|
||||
|
||||
res = determinant(m2);
|
||||
res = determinant(m3);
|
||||
res = determinant(m4);
|
||||
m2 = inverse(m2);
|
||||
m3 = inverse(m3);
|
||||
m4 = inverse(m4);
|
||||
|
||||
res = modf(f32.x, tmp);
|
||||
// ModfStruct
|
||||
|
||||
res = min(f32.x, f32.y);
|
||||
ures = min(u32.x, u32.y);
|
||||
ires = min(s32.x, s32.y);
|
||||
res = max(f32.x, f32.y);
|
||||
ures = max(u32.x, u32.y);
|
||||
ires = max(s32.x, s32.y);
|
||||
|
||||
res = clamp(f32.x, f32.y, f32.z);
|
||||
ures = clamp(u32.x, u32.y, u32.z);
|
||||
ires = clamp(s32.x, s32.y, s32.z);
|
||||
|
||||
res = mix(f32.x, f32.y, f32.z);
|
||||
res = step(f32.x, f32.y);
|
||||
res = smoothstep(f32.x, f32.y, f32.z);
|
||||
res = fma(f32.x, f32.y, f32.z);
|
||||
|
||||
res = frexp(f32.x, itmp);
|
||||
// FrexpStruct
|
||||
res = ldexp(f32.x, itmp);
|
||||
|
||||
ures = packSnorm4x8(f32);
|
||||
ures = packUnorm4x8(f32);
|
||||
ures = packSnorm2x16(f32.xy);
|
||||
ures = packUnorm2x16(f32.xy);
|
||||
ures = packHalf2x16(f32.xy);
|
||||
// packDouble2x32
|
||||
|
||||
v2 = unpackSnorm2x16(u32.x);
|
||||
v2 = unpackUnorm2x16(u32.x);
|
||||
v2 = unpackHalf2x16(u32.x);
|
||||
v4 = unpackSnorm4x8(u32.x);
|
||||
v4 = unpackUnorm4x8(u32.x);
|
||||
// unpackDouble2x32
|
||||
|
||||
s32 = findLSB(s32);
|
||||
s32 = findLSB(u32);
|
||||
s32 = findMSB(s32);
|
||||
s32 = findMSB(u32);
|
||||
|
||||
// interpolateAtSample
|
||||
// interpolateAtOffset
|
||||
|
||||
// NMin, NMax, NClamp
|
||||
}
|
||||
@ -1772,6 +1772,16 @@ void CompilerHLSL::emit_resources()
|
||||
end_scope();
|
||||
statement("");
|
||||
}
|
||||
|
||||
if (requires_scalar_faceforward)
|
||||
{
|
||||
// FP16/FP64? No templates in HLSL.
|
||||
statement("float SPIRV_Cross_FaceForward(float n, float i, float nref)");
|
||||
begin_scope();
|
||||
statement("return i * nref < 0.0 ? n : -n;");
|
||||
end_scope();
|
||||
statement("");
|
||||
}
|
||||
}
|
||||
|
||||
string CompilerHLSL::layout_for_member(const SPIRType &type, uint32_t index)
|
||||
@ -3321,6 +3331,20 @@ void CompilerHLSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop,
|
||||
CompilerGLSL::emit_glsl_op(result_type, id, eop, args, count);
|
||||
break;
|
||||
|
||||
case GLSLstd450FaceForward:
|
||||
if (get<SPIRType>(result_type).vecsize == 1)
|
||||
{
|
||||
if (!requires_scalar_faceforward)
|
||||
{
|
||||
requires_scalar_faceforward = true;
|
||||
force_recompile();
|
||||
}
|
||||
emit_trinary_func_op(result_type, id, args[0], args[1], args[2], "SPIRV_Cross_FaceForward");
|
||||
}
|
||||
else
|
||||
CompilerGLSL::emit_glsl_op(result_type, id, eop, args, count);
|
||||
break;
|
||||
|
||||
default:
|
||||
CompilerGLSL::emit_glsl_op(result_type, id, eop, args, count);
|
||||
break;
|
||||
|
||||
@ -183,6 +183,7 @@ private:
|
||||
bool requires_inverse_4x4 = false;
|
||||
bool requires_scalar_reflect = false;
|
||||
bool requires_scalar_refract = false;
|
||||
bool requires_scalar_faceforward = false;
|
||||
uint64_t required_textureSizeVariants = 0;
|
||||
void require_texture_query_variant(const SPIRType &type);
|
||||
|
||||
|
||||
@ -3345,6 +3345,16 @@ void CompilerMSL::emit_custom_functions()
|
||||
statement("");
|
||||
break;
|
||||
|
||||
case SPVFuncImplFaceForwardScalar:
|
||||
// Metal does not support scalar versions of these functions.
|
||||
statement("template<typename T>");
|
||||
statement("inline T spvFaceForward(T n, T i, T nref)");
|
||||
begin_scope();
|
||||
statement("return i * nref < T(0) ? n : -n;");
|
||||
end_scope();
|
||||
statement("");
|
||||
break;
|
||||
|
||||
default:
|
||||
break;
|
||||
}
|
||||
@ -4938,6 +4948,13 @@ void CompilerMSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop,
|
||||
CompilerGLSL::emit_glsl_op(result_type, id, eop, args, count);
|
||||
break;
|
||||
|
||||
case GLSLstd450FaceForward:
|
||||
if (get<SPIRType>(result_type).vecsize == 1)
|
||||
emit_trinary_func_op(result_type, id, args[0], args[1], args[2], "spvFaceForward");
|
||||
else
|
||||
CompilerGLSL::emit_glsl_op(result_type, id, eop, args, count);
|
||||
break;
|
||||
|
||||
default:
|
||||
CompilerGLSL::emit_glsl_op(result_type, id, eop, args, count);
|
||||
break;
|
||||
@ -9061,6 +9078,14 @@ CompilerMSL::SPVFuncImpl CompilerMSL::OpCodePreprocessor::get_spv_func_impl(Op o
|
||||
else
|
||||
return SPVFuncImplNone;
|
||||
}
|
||||
case GLSLstd450FaceForward:
|
||||
{
|
||||
auto &type = compiler.get<SPIRType>(args[0]);
|
||||
if (type.vecsize == 1)
|
||||
return SPVFuncImplFaceForwardScalar;
|
||||
else
|
||||
return SPVFuncImplNone;
|
||||
}
|
||||
case GLSLstd450MatrixInverse:
|
||||
{
|
||||
auto &mat_type = compiler.get<SPIRType>(args[0]);
|
||||
|
||||
@ -405,6 +405,7 @@ protected:
|
||||
SPVFuncImplSubgroupAllEqual,
|
||||
SPVFuncImplReflectScalar,
|
||||
SPVFuncImplRefractScalar,
|
||||
SPVFuncImplFaceForwardScalar,
|
||||
SPVFuncImplArrayCopyMultidimMax = 6
|
||||
};
|
||||
|
||||
|
||||
@ -247,8 +247,13 @@ def shader_to_win_path(shader):
|
||||
|
||||
ignore_fxc = False
|
||||
def validate_shader_hlsl(shader, force_no_external_validation, paths):
|
||||
if not '.nonuniformresource' in shader:
|
||||
# glslang HLSL does not support this, so rely on fxc to test it.
|
||||
test_glslang = True
|
||||
if '.nonuniformresource.' in shader:
|
||||
test_glslang = False
|
||||
if '.fxconly.' in shader:
|
||||
test_glslang = False
|
||||
|
||||
if test_glslang:
|
||||
subprocess.check_call([paths.glslang, '-e', 'main', '-D', '--target-env', 'vulkan1.1', '-V', shader])
|
||||
is_no_fxc = '.nofxc.' in shader
|
||||
global ignore_fxc
|
||||
|
||||
Loading…
Reference in New Issue
Block a user