Merge pull request #1752 from billhollings/function-constants-opquantize
MSL: Support synthetic functions in function constants.
This commit is contained in:
Hans-Kristian Arntzen
GitHub
commit
9462b90067
@ -7,11 +7,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct u0_counters
|
||||
{
|
||||
uint c;
|
||||
};
|
||||
|
||||
// Returns 2D texture coords corresponding to 1D texel buffer coords
|
||||
static inline __attribute__((always_inline))
|
||||
uint2 spvTexelBufferCoord(uint tc)
|
||||
@ -19,6 +14,11 @@ uint2 spvTexelBufferCoord(uint tc)
|
||||
return uint2(tc % 4096, tc / 4096);
|
||||
}
|
||||
|
||||
struct u0_counters
|
||||
{
|
||||
uint c;
|
||||
};
|
||||
|
||||
kernel void main0(device u0_counters& u0_counter [[buffer(0)]], texture2d<uint, access::write> u0 [[texture(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||||
{
|
||||
uint _29 = atomic_fetch_sub_explicit((device atomic_uint*)&u0_counter.c, 1, memory_order_relaxed);
|
||||
|
||||
@ -7,11 +7,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct u0_counters
|
||||
{
|
||||
uint c;
|
||||
};
|
||||
|
||||
// Returns 2D texture coords corresponding to 1D texel buffer coords
|
||||
static inline __attribute__((always_inline))
|
||||
uint2 spvTexelBufferCoord(uint tc)
|
||||
@ -19,6 +14,11 @@ uint2 spvTexelBufferCoord(uint tc)
|
||||
return uint2(tc % 4096, tc / 4096);
|
||||
}
|
||||
|
||||
struct u0_counters
|
||||
{
|
||||
uint c;
|
||||
};
|
||||
|
||||
kernel void main0(device u0_counters& u0_counter [[buffer(0)]], texture2d<uint, access::write> u0 [[texture(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||||
{
|
||||
uint _29 = atomic_fetch_add_explicit((device atomic_uint*)&u0_counter.c, 1, memory_order_relaxed);
|
||||
|
||||
@ -5,11 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct cb5_struct
|
||||
{
|
||||
float4 _m0[5];
|
||||
};
|
||||
|
||||
// Returns 2D texture coords corresponding to 1D texel buffer coords
|
||||
static inline __attribute__((always_inline))
|
||||
uint2 spvTexelBufferCoord(uint tc)
|
||||
@ -17,6 +12,11 @@ uint2 spvTexelBufferCoord(uint tc)
|
||||
return uint2(tc % 4096, tc / 4096);
|
||||
}
|
||||
|
||||
struct cb5_struct
|
||||
{
|
||||
float4 _m0[5];
|
||||
};
|
||||
|
||||
kernel void main0(constant cb5_struct& cb0_5 [[buffer(0)]], texture2d<uint, access::write> u0 [[texture(0)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
|
||||
{
|
||||
uint _44 = as_type<uint>(as_type<float>(int(gl_LocalInvocationID.x) << 4)) >> 2u;
|
||||
|
||||
@ -5,11 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct cb
|
||||
{
|
||||
float value;
|
||||
};
|
||||
|
||||
// Returns 2D texture coords corresponding to 1D texel buffer coords
|
||||
static inline __attribute__((always_inline))
|
||||
uint2 spvTexelBufferCoord(uint tc)
|
||||
@ -17,6 +12,11 @@ uint2 spvTexelBufferCoord(uint tc)
|
||||
return uint2(tc % 4096, tc / 4096);
|
||||
}
|
||||
|
||||
struct cb
|
||||
{
|
||||
float value;
|
||||
};
|
||||
|
||||
kernel void main0(constant cb& _6 [[buffer(0)]], texture2d<float, access::write> _buffer [[texture(0)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
|
||||
{
|
||||
_buffer.write(float4(_6.value), spvTexelBufferCoord(((32u * gl_WorkGroupID.x) + gl_LocalInvocationIndex)));
|
||||
|
||||
@ -5,6 +5,17 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template <typename F> struct SpvHalfTypeSelector;
|
||||
template <> struct SpvHalfTypeSelector<float> { public: using H = half; };
|
||||
template<uint N> struct SpvHalfTypeSelector<vec<float, N>> { using H = vec<half, N>; };
|
||||
template<typename F, typename H = typename SpvHalfTypeSelector<F>::H>
|
||||
[[clang::optnone]] F spvQuantizeToF16(F fval)
|
||||
{
|
||||
H hval = H(fval);
|
||||
hval = select(copysign(H(0), hval), hval, isnormal(hval) || isinf(hval) || isnan(hval));
|
||||
return F(hval);
|
||||
}
|
||||
|
||||
struct SSBO0
|
||||
{
|
||||
float scalar;
|
||||
@ -13,15 +24,6 @@ struct SSBO0
|
||||
float4 vec4_val;
|
||||
};
|
||||
|
||||
template <typename F> struct SpvHalfTypeSelector;
|
||||
template <> struct SpvHalfTypeSelector<float> { public: using H = half; };
|
||||
template<uint N> struct SpvHalfTypeSelector<vec<float, N>> { using H = vec<half, N>; };
|
||||
template<typename F, typename H = typename SpvHalfTypeSelector<F>::H>
|
||||
[[clang::optnone]] F spvQuantizeToF16(F val)
|
||||
{
|
||||
return F(H(val));
|
||||
}
|
||||
|
||||
kernel void main0(device SSBO0& _4 [[buffer(0)]])
|
||||
{
|
||||
_4.scalar = spvQuantizeToF16(_4.scalar);
|
||||
|
||||
@ -44,6 +44,13 @@ struct spvUnsafeArray
|
||||
}
|
||||
};
|
||||
|
||||
// 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);
|
||||
}
|
||||
|
||||
struct myType
|
||||
{
|
||||
float data;
|
||||
@ -54,13 +61,6 @@ struct main0_out
|
||||
float4 o_color [[color(0)]];
|
||||
};
|
||||
|
||||
// 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);
|
||||
}
|
||||
|
||||
fragment main0_out main0(float4 gl_FragCoord [[position]])
|
||||
{
|
||||
spvUnsafeArray<myType, 5> _21 = spvUnsafeArray<myType, 5>({ myType{ 0.0 }, myType{ 1.0 }, myType{ 0.0 }, myType{ 1.0 }, myType{ 0.0 } });
|
||||
|
||||
@ -5,28 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct Data
|
||||
{
|
||||
float a;
|
||||
float b;
|
||||
};
|
||||
|
||||
constant float X_tmp [[function_constant(0)]];
|
||||
constant float X = is_function_constant_defined(X_tmp) ? X_tmp : 4.0;
|
||||
|
||||
struct Data_1
|
||||
{
|
||||
float a;
|
||||
float b;
|
||||
};
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
Data_1 outdata[1];
|
||||
};
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(2u, 1u, 1u);
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
|
||||
{
|
||||
@ -135,6 +113,28 @@ inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device
|
||||
}
|
||||
}
|
||||
|
||||
struct Data
|
||||
{
|
||||
float a;
|
||||
float b;
|
||||
};
|
||||
|
||||
constant float X_tmp [[function_constant(0)]];
|
||||
constant float X = is_function_constant_defined(X_tmp) ? X_tmp : 4.0;
|
||||
|
||||
struct Data_1
|
||||
{
|
||||
float a;
|
||||
float b;
|
||||
};
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
Data_1 outdata[1];
|
||||
};
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(2u, 1u, 1u);
|
||||
|
||||
kernel void main0(device SSBO& _53 [[buffer(0)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
|
||||
{
|
||||
Data _25[2] = { Data{ 1.0, 2.0 }, Data{ 3.0, 4.0 } };
|
||||
|
||||
@ -5,12 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct myBlock
|
||||
{
|
||||
int a;
|
||||
float b[1];
|
||||
};
|
||||
|
||||
// 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)
|
||||
@ -18,6 +12,12 @@ inline Tx mod(Tx x, Ty y)
|
||||
return x - y * floor(x / y);
|
||||
}
|
||||
|
||||
struct myBlock
|
||||
{
|
||||
int a;
|
||||
float b[1];
|
||||
};
|
||||
|
||||
kernel void main0(device myBlock& myStorage [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||||
{
|
||||
myStorage.a = (myStorage.a + 1) % 256;
|
||||
|
||||
@ -5,12 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct myBlock
|
||||
{
|
||||
int a;
|
||||
float b[1];
|
||||
};
|
||||
|
||||
// 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)
|
||||
@ -18,6 +12,12 @@ inline Tx mod(Tx x, Ty y)
|
||||
return x - y * floor(x / y);
|
||||
}
|
||||
|
||||
struct myBlock
|
||||
{
|
||||
int a;
|
||||
float b[1];
|
||||
};
|
||||
|
||||
kernel void main0(device myBlock& myStorage [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||||
{
|
||||
myStorage.a = (myStorage.a + 1) % 256;
|
||||
|
||||
@ -7,6 +7,12 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
// The required alignment of a linear texture of R32Uint format.
|
||||
constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
|
||||
constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
|
||||
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
|
||||
#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() + spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
float4 outdata;
|
||||
@ -23,12 +29,6 @@ struct spvDescriptorSetBuffer0
|
||||
sampler uTextureSmplr [[id(4)]];
|
||||
};
|
||||
|
||||
// The required alignment of a linear texture of R32Uint format.
|
||||
constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
|
||||
constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
|
||||
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
|
||||
#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() + spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
|
||||
|
||||
kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||||
{
|
||||
uint _26 = atomic_fetch_add_explicit((device atomic_uint*)&spvDescriptorSet0.uImage_atomic[spvImage2DAtomicCoord(int2(gl_GlobalInvocationID.xy), spvDescriptorSet0.uImage)], 10u, memory_order_relaxed);
|
||||
|
||||
@ -7,6 +7,12 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
// The required alignment of a linear texture of R32Uint format.
|
||||
constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
|
||||
constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
|
||||
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
|
||||
#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() + spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
float4 outdata;
|
||||
@ -14,12 +20,6 @@ struct SSBO
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
||||
|
||||
// The required alignment of a linear texture of R32Uint format.
|
||||
constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
|
||||
constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
|
||||
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
|
||||
#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() + spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
|
||||
|
||||
kernel void main0(device SSBO& _31 [[buffer(1)]], texture2d<uint> uImage [[texture(0)]], device atomic_uint* uImage_atomic [[buffer(0)]], texture2d<float> uTexture [[texture(1)]], sampler uTextureSmplr [[sampler(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||||
{
|
||||
uint _26 = atomic_fetch_add_explicit((device atomic_uint*)&uImage_atomic[spvImage2DAtomicCoord(int2(gl_GlobalInvocationID.xy), uImage)], 10u, memory_order_relaxed);
|
||||
|
||||
@ -5,22 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct MatrixOut
|
||||
{
|
||||
float2x2 m2out;
|
||||
float3x3 m3out;
|
||||
float4x4 m4out;
|
||||
};
|
||||
|
||||
struct MatrixIn
|
||||
{
|
||||
float2x2 m2in;
|
||||
float3x3 m3in;
|
||||
float4x4 m4in;
|
||||
};
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
||||
|
||||
// Returns the determinant of a 2x2 matrix.
|
||||
static inline __attribute__((always_inline))
|
||||
float spvDet2x2(float a1, float a2, float b1, float b2)
|
||||
@ -121,6 +105,22 @@ float2x2 spvInverse2x2(float2x2 m)
|
||||
return (det != 0.0f) ? (adj * (1.0f / det)) : m;
|
||||
}
|
||||
|
||||
struct MatrixOut
|
||||
{
|
||||
float2x2 m2out;
|
||||
float3x3 m3out;
|
||||
float4x4 m4out;
|
||||
};
|
||||
|
||||
struct MatrixIn
|
||||
{
|
||||
float2x2 m2in;
|
||||
float3x3 m3in;
|
||||
float4x4 m4in;
|
||||
};
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
||||
|
||||
kernel void main0(device MatrixOut& _15 [[buffer(0)]], const device MatrixIn& _20 [[buffer(1)]])
|
||||
{
|
||||
_15.m2out = spvInverse2x2(_20.m2in);
|
||||
|
||||
@ -5,12 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct myBlock
|
||||
{
|
||||
int a;
|
||||
float b[1];
|
||||
};
|
||||
|
||||
// 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)
|
||||
@ -18,6 +12,12 @@ inline Tx mod(Tx x, Ty y)
|
||||
return x - y * floor(x / y);
|
||||
}
|
||||
|
||||
struct myBlock
|
||||
{
|
||||
int a;
|
||||
float b[1];
|
||||
};
|
||||
|
||||
kernel void main0(device myBlock& myStorage [[buffer(0)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
|
||||
{
|
||||
myStorage.a = (myStorage.a + 1) % 256;
|
||||
|
||||
@ -5,12 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct myBlock
|
||||
{
|
||||
int a;
|
||||
float b[1];
|
||||
};
|
||||
|
||||
// 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)
|
||||
@ -18,6 +12,12 @@ inline Tx mod(Tx x, Ty y)
|
||||
return x - y * floor(x / y);
|
||||
}
|
||||
|
||||
struct myBlock
|
||||
{
|
||||
int a;
|
||||
float b[1];
|
||||
};
|
||||
|
||||
kernel void main0(device myBlock& myStorage [[buffer(0)]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
|
||||
{
|
||||
myStorage.a = (myStorage.a + 1) % 256;
|
||||
|
||||
@ -5,6 +5,13 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
// 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);
|
||||
}
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
float4 in_data[1];
|
||||
@ -17,13 +24,6 @@ struct SSBO2
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
||||
|
||||
// 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);
|
||||
}
|
||||
|
||||
kernel void main0(const device SSBO& _23 [[buffer(0)]], device SSBO2& _33 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||||
{
|
||||
_33.out_data[gl_GlobalInvocationID.x] = mod(_23.in_data[gl_GlobalInvocationID.x], _33.out_data[gl_GlobalInvocationID.x]);
|
||||
|
||||
@ -5,12 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct myBlock
|
||||
{
|
||||
int a;
|
||||
float b;
|
||||
};
|
||||
|
||||
// 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)
|
||||
@ -18,6 +12,12 @@ inline Tx mod(Tx x, Ty y)
|
||||
return x - y * floor(x / y);
|
||||
}
|
||||
|
||||
struct myBlock
|
||||
{
|
||||
int a;
|
||||
float b;
|
||||
};
|
||||
|
||||
kernel void main0(device myBlock& myStorage [[buffer(0)]])
|
||||
{
|
||||
myStorage.a = (myStorage.a + 1) % 256;
|
||||
|
||||
@ -5,35 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct UBO
|
||||
{
|
||||
uint index;
|
||||
};
|
||||
|
||||
struct UBO2
|
||||
{
|
||||
uint index2;
|
||||
};
|
||||
|
||||
struct spvDescriptorSetBuffer0
|
||||
{
|
||||
array<texture2d<float>, 4> uSampler [[id(0)]];
|
||||
array<sampler, 4> uSamplerSmplr [[id(4)]];
|
||||
constant UBO* uUBO [[id(8)]];
|
||||
constant UBO2* m_50 [[id(9)]];
|
||||
constant uint* spvSwizzleConstants [[id(10)]];
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float2 vUV [[user(locn0)]];
|
||||
};
|
||||
|
||||
template<typename T> struct spvRemoveReference { typedef T type; };
|
||||
template<typename T> struct spvRemoveReference<thread T&> { typedef T type; };
|
||||
template<typename T> struct spvRemoveReference<thread T&&> { typedef T type; };
|
||||
@ -94,6 +65,35 @@ inline T spvTextureSwizzle(T x, uint s)
|
||||
return spvTextureSwizzle(vec<T, 4>(x, 0, 0, 1), s).x;
|
||||
}
|
||||
|
||||
struct UBO
|
||||
{
|
||||
uint index;
|
||||
};
|
||||
|
||||
struct UBO2
|
||||
{
|
||||
uint index2;
|
||||
};
|
||||
|
||||
struct spvDescriptorSetBuffer0
|
||||
{
|
||||
array<texture2d<float>, 4> uSampler [[id(0)]];
|
||||
array<sampler, 4> uSamplerSmplr [[id(4)]];
|
||||
constant UBO* uUBO [[id(8)]];
|
||||
constant UBO2* m_50 [[id(9)]];
|
||||
constant uint* spvSwizzleConstants [[id(10)]];
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float2 vUV [[user(locn0)]];
|
||||
};
|
||||
|
||||
fragment main0_out main0(main0_in in [[stage_in]], constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], constant uint* spvSwizzleConstants [[buffer(30)]])
|
||||
{
|
||||
main0_out out = {};
|
||||
|
||||
@ -5,26 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct UBO
|
||||
{
|
||||
uint index;
|
||||
};
|
||||
|
||||
struct UBO2
|
||||
{
|
||||
uint index2;
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float2 vUV [[user(locn0)]];
|
||||
};
|
||||
|
||||
template<typename T> struct spvRemoveReference { typedef T type; };
|
||||
template<typename T> struct spvRemoveReference<thread T&> { typedef T type; };
|
||||
template<typename T> struct spvRemoveReference<thread T&&> { typedef T type; };
|
||||
@ -85,6 +65,26 @@ inline T spvTextureSwizzle(T x, uint s)
|
||||
return spvTextureSwizzle(vec<T, 4>(x, 0, 0, 1), s).x;
|
||||
}
|
||||
|
||||
struct UBO
|
||||
{
|
||||
uint index;
|
||||
};
|
||||
|
||||
struct UBO2
|
||||
{
|
||||
uint index2;
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float2 vUV [[user(locn0)]];
|
||||
};
|
||||
|
||||
fragment main0_out main0(main0_in in [[stage_in]], constant uint* spvSwizzleConstants [[buffer(30)]], constant UBO& uUBO [[buffer(0)]], constant UBO2& _50 [[buffer(1)]], array<texture2d<float>, 4> uSampler [[texture(0)]], array<sampler, 4> uSamplerSmplr [[sampler(0)]])
|
||||
{
|
||||
main0_out out = {};
|
||||
|
||||
@ -5,23 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct spvDescriptorSetBuffer0
|
||||
{
|
||||
array<texture2d<float>, 4> uSampler0 [[id(0)]];
|
||||
array<sampler, 4> uSampler0Smplr [[id(4)]];
|
||||
constant uint* spvSwizzleConstants [[id(8)]];
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float2 vUV [[user(locn0)]];
|
||||
};
|
||||
|
||||
template<typename T> struct spvRemoveReference { typedef T type; };
|
||||
template<typename T> struct spvRemoveReference<thread T&> { typedef T type; };
|
||||
template<typename T> struct spvRemoveReference<thread T&&> { typedef T type; };
|
||||
@ -82,6 +65,23 @@ inline T spvTextureSwizzle(T x, uint s)
|
||||
return spvTextureSwizzle(vec<T, 4>(x, 0, 0, 1), s).x;
|
||||
}
|
||||
|
||||
struct spvDescriptorSetBuffer0
|
||||
{
|
||||
array<texture2d<float>, 4> uSampler0 [[id(0)]];
|
||||
array<sampler, 4> uSampler0Smplr [[id(4)]];
|
||||
constant uint* spvSwizzleConstants [[id(8)]];
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float2 vUV [[user(locn0)]];
|
||||
};
|
||||
|
||||
fragment main0_out main0(main0_in in [[stage_in]], constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], constant uint* spvSwizzleConstants [[buffer(30)]], texture2d<float> uSampler1 [[texture(0)]], sampler uSampler1Smplr [[sampler(0)]])
|
||||
{
|
||||
main0_out out = {};
|
||||
|
||||
@ -5,16 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float2 vUV [[user(locn0)]];
|
||||
};
|
||||
|
||||
template<typename T> struct spvRemoveReference { typedef T type; };
|
||||
template<typename T> struct spvRemoveReference<thread T&> { typedef T type; };
|
||||
template<typename T> struct spvRemoveReference<thread T&&> { typedef T type; };
|
||||
@ -75,6 +65,16 @@ inline T spvTextureSwizzle(T x, uint s)
|
||||
return spvTextureSwizzle(vec<T, 4>(x, 0, 0, 1), s).x;
|
||||
}
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float2 vUV [[user(locn0)]];
|
||||
};
|
||||
|
||||
fragment main0_out main0(main0_in in [[stage_in]], constant uint* spvSwizzleConstants [[buffer(30)]], array<texture2d<float>, 4> uSampler [[texture(0)]], array<sampler, 4> uSamplerSmplr [[sampler(0)]])
|
||||
{
|
||||
main0_out out = {};
|
||||
|
||||
@ -5,11 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
// Returns 2D texture coords corresponding to 1D texel buffer coords
|
||||
static inline __attribute__((always_inline))
|
||||
uint2 spvTexelBufferCoord(uint tc)
|
||||
@ -17,6 +12,11 @@ uint2 spvTexelBufferCoord(uint tc)
|
||||
return uint2(tc % 4096, tc / 4096);
|
||||
}
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
fragment main0_out main0(texture2d<float> buf [[texture(0)]], texture2d<float, access::write> bufOut [[texture(1)]], float4 gl_FragCoord [[position]])
|
||||
{
|
||||
main0_out out = {};
|
||||
|
||||
@ -44,6 +44,13 @@ struct spvUnsafeArray
|
||||
}
|
||||
};
|
||||
|
||||
// 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);
|
||||
}
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor_0 [[color(0)]];
|
||||
@ -58,13 +65,6 @@ struct main0_in
|
||||
float4 vB [[user(locn1)]];
|
||||
};
|
||||
|
||||
// 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);
|
||||
}
|
||||
|
||||
fragment main0_out main0(main0_in in [[stage_in]])
|
||||
{
|
||||
main0_out out = {};
|
||||
|
||||
@ -7,6 +7,12 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
// The required alignment of a linear texture of R32Uint format.
|
||||
constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
|
||||
constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
|
||||
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
|
||||
#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() + spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
|
||||
|
||||
struct Buffer3
|
||||
{
|
||||
int baz;
|
||||
@ -35,12 +41,6 @@ struct spvDescriptorSetBuffer0
|
||||
device Buffer2* m_52 [[id(7), raster_order_group(0)]];
|
||||
};
|
||||
|
||||
// The required alignment of a linear texture of R32Uint format.
|
||||
constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
|
||||
constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
|
||||
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
|
||||
#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() + spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
|
||||
|
||||
fragment void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]])
|
||||
{
|
||||
(*spvDescriptorSet0.m_9).baz = 0;
|
||||
|
||||
@ -7,6 +7,12 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
// The required alignment of a linear texture of R32Uint format.
|
||||
constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
|
||||
constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
|
||||
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
|
||||
#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() + spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
|
||||
|
||||
struct Buffer3
|
||||
{
|
||||
int baz;
|
||||
@ -23,12 +29,6 @@ struct Buffer2
|
||||
uint quux;
|
||||
};
|
||||
|
||||
// The required alignment of a linear texture of R32Uint format.
|
||||
constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
|
||||
constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
|
||||
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
|
||||
#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() + spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
|
||||
|
||||
fragment void main0(device Buffer3& _9 [[buffer(0)]], volatile device Buffer& _42 [[buffer(2), raster_order_group(0)]], device Buffer2& _52 [[buffer(3), raster_order_group(0)]], texture2d<float, access::write> img4 [[texture(0)]], texture2d<float, access::write> img [[texture(1), raster_order_group(0)]], texture2d<float> img3 [[texture(2), raster_order_group(0)]], texture2d<uint> img2 [[texture(3), raster_order_group(0)]], device atomic_uint* img2_atomic [[buffer(1), raster_order_group(0)]])
|
||||
{
|
||||
_9.baz = 0;
|
||||
|
||||
@ -5,16 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float3 vRefract [[user(locn0)]];
|
||||
};
|
||||
|
||||
template<typename T>
|
||||
[[clang::optnone]] T spvReflect(T i, T n)
|
||||
{
|
||||
@ -37,6 +27,16 @@ inline T spvRefract(T i, T n, T eta)
|
||||
}
|
||||
}
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float3 vRefract [[user(locn0)]];
|
||||
};
|
||||
|
||||
fragment main0_out main0(main0_in in [[stage_in]])
|
||||
{
|
||||
main0_out out = {};
|
||||
|
||||
@ -5,11 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
uint2 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
inline uint spvSubgroupBallotFindLSB(uint4 ballot, uint gl_SubgroupSize)
|
||||
{
|
||||
uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupSize, 32u)), extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupSize - 32, 0)), uint2(0));
|
||||
@ -47,6 +42,11 @@ inline uint spvSubgroupBallotExclusiveBitCount(uint4 ballot, uint gl_SubgroupInv
|
||||
return spvPopCount4(ballot & mask);
|
||||
}
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
uint2 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
fragment main0_out main0(uint gl_SubgroupInvocationID [[thread_index_in_simdgroup]], uint gl_SubgroupSize [[threads_per_simdgroup]])
|
||||
{
|
||||
main0_out out = {};
|
||||
|
||||
@ -5,16 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float4 vUV [[user(locn0)]];
|
||||
};
|
||||
|
||||
static inline __attribute__((always_inline))
|
||||
float3 spvCubemapTo2DArrayFace(float3 P)
|
||||
{
|
||||
@ -49,6 +39,16 @@ float3 spvCubemapTo2DArrayFace(float3 P)
|
||||
return float3(u, v, CubeFace);
|
||||
}
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float4 vUV [[user(locn0)]];
|
||||
};
|
||||
|
||||
fragment main0_out main0(main0_in in [[stage_in]], texturecube<float> cubeSampler [[texture(0)]], texture2d_array<float> cubeArraySampler [[texture(1)]], texture2d_array<float> texArraySampler [[texture(2)]], sampler cubeSamplerSmplr [[sampler(0)]], sampler cubeArraySamplerSmplr [[sampler(1)]], sampler texArraySamplerSmplr [[sampler(2)]])
|
||||
{
|
||||
main0_out out = {};
|
||||
|
||||
@ -44,22 +44,6 @@ struct spvUnsafeArray
|
||||
}
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 gl_Position;
|
||||
};
|
||||
|
||||
struct main0_patchOut
|
||||
{
|
||||
spvUnsafeArray<float4, 2> pFoo;
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
spvUnsafeArray<float4, 2> iFoo;
|
||||
float4 ipFoo;
|
||||
};
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
|
||||
{
|
||||
@ -168,6 +152,22 @@ inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device
|
||||
}
|
||||
}
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 gl_Position;
|
||||
};
|
||||
|
||||
struct main0_patchOut
|
||||
{
|
||||
spvUnsafeArray<float4, 2> pFoo;
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
spvUnsafeArray<float4, 2> iFoo;
|
||||
float4 ipFoo;
|
||||
};
|
||||
|
||||
kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], device main0_out* spvOut [[buffer(28)]], constant uint* spvIndirectParams [[buffer(29)]], device main0_patchOut* spvPatchOut [[buffer(27)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], device main0_in* spvIn [[buffer(22)]])
|
||||
{
|
||||
device main0_out* gl_out = &spvOut[gl_GlobalInvocationID.x - gl_GlobalInvocationID.x % 4];
|
||||
|
||||
@ -44,23 +44,6 @@ struct spvUnsafeArray
|
||||
}
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 gl_Position;
|
||||
};
|
||||
|
||||
struct main0_patchOut
|
||||
{
|
||||
spvUnsafeArray<float4, 2> pFoo;
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float4 iFoo_0 [[attribute(0)]];
|
||||
float4 iFoo_1 [[attribute(1)]];
|
||||
float4 ipFoo [[attribute(2)]];
|
||||
};
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
|
||||
{
|
||||
@ -169,6 +152,23 @@ inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device
|
||||
}
|
||||
}
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 gl_Position;
|
||||
};
|
||||
|
||||
struct main0_patchOut
|
||||
{
|
||||
spvUnsafeArray<float4, 2> pFoo;
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float4 iFoo_0 [[attribute(0)]];
|
||||
float4 iFoo_1 [[attribute(1)]];
|
||||
float4 ipFoo [[attribute(2)]];
|
||||
};
|
||||
|
||||
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 main0_patchOut* spvPatchOut [[buffer(27)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], threadgroup main0_in* gl_in [[threadgroup(0)]])
|
||||
{
|
||||
threadgroup float4 Foo[4][2];
|
||||
|
||||
@ -44,30 +44,6 @@ struct spvUnsafeArray
|
||||
}
|
||||
};
|
||||
|
||||
struct Matrices
|
||||
{
|
||||
float4x4 vpMatrix;
|
||||
float4x4 wMatrix;
|
||||
float4x3 wMatrix4x3;
|
||||
float3x4 wMatrix3x4;
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float3 OutNormal [[user(locn0)]];
|
||||
float4 OutWorldPos_0 [[user(locn1)]];
|
||||
float4 OutWorldPos_1 [[user(locn2)]];
|
||||
float4 OutWorldPos_2 [[user(locn3)]];
|
||||
float4 OutWorldPos_3 [[user(locn4)]];
|
||||
float4 gl_Position [[position]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float3 InPos [[attribute(0)]];
|
||||
float3 InNormal [[attribute(1)]];
|
||||
};
|
||||
|
||||
template<typename T>
|
||||
[[clang::optnone]] T spvFMul(T l, T r)
|
||||
{
|
||||
@ -117,6 +93,30 @@ template<typename T, int LCols, int LRows, int RCols, int RRows>
|
||||
return res;
|
||||
}
|
||||
|
||||
struct Matrices
|
||||
{
|
||||
float4x4 vpMatrix;
|
||||
float4x4 wMatrix;
|
||||
float4x3 wMatrix4x3;
|
||||
float3x4 wMatrix3x4;
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float3 OutNormal [[user(locn0)]];
|
||||
float4 OutWorldPos_0 [[user(locn1)]];
|
||||
float4 OutWorldPos_1 [[user(locn2)]];
|
||||
float4 OutWorldPos_2 [[user(locn3)]];
|
||||
float4 OutWorldPos_3 [[user(locn4)]];
|
||||
float4 gl_Position [[position]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float3 InPos [[attribute(0)]];
|
||||
float3 InNormal [[attribute(1)]];
|
||||
};
|
||||
|
||||
vertex main0_out main0(main0_in in [[stage_in]], constant Matrices& _22 [[buffer(0)]])
|
||||
{
|
||||
main0_out out = {};
|
||||
|
||||
@ -5,30 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct UBO
|
||||
{
|
||||
float4x4 uMVP;
|
||||
float3 rotDeg;
|
||||
float3 rotRad;
|
||||
int2 bits;
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float3 vNormal [[user(locn0)]];
|
||||
float3 vRotDeg [[user(locn1)]];
|
||||
float3 vRotRad [[user(locn2)]];
|
||||
int2 vLSB [[user(locn3)]];
|
||||
int2 vMSB [[user(locn4)]];
|
||||
float4 gl_Position [[position]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float4 aVertex [[attribute(0)]];
|
||||
float3 aNormal [[attribute(1)]];
|
||||
};
|
||||
|
||||
// Implementation of the GLSL radians() function
|
||||
template<typename T>
|
||||
inline T radians(T d)
|
||||
@ -108,6 +84,30 @@ float4x4 spvInverse4x4(float4x4 m)
|
||||
return (det != 0.0f) ? (adj * (1.0f / det)) : m;
|
||||
}
|
||||
|
||||
struct UBO
|
||||
{
|
||||
float4x4 uMVP;
|
||||
float3 rotDeg;
|
||||
float3 rotRad;
|
||||
int2 bits;
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float3 vNormal [[user(locn0)]];
|
||||
float3 vRotDeg [[user(locn1)]];
|
||||
float3 vRotRad [[user(locn2)]];
|
||||
int2 vLSB [[user(locn3)]];
|
||||
int2 vMSB [[user(locn4)]];
|
||||
float4 gl_Position [[position]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float4 aVertex [[attribute(0)]];
|
||||
float3 aNormal [[attribute(1)]];
|
||||
};
|
||||
|
||||
vertex main0_out main0(main0_in in [[stage_in]], constant UBO& _18 [[buffer(0)]])
|
||||
{
|
||||
main0_out out = {};
|
||||
|
||||
@ -5,18 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 gl_Position [[position]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float4 vA [[attribute(0)]];
|
||||
float4 vB [[attribute(1)]];
|
||||
float4 vC [[attribute(2)]];
|
||||
};
|
||||
|
||||
template<typename T>
|
||||
[[clang::optnone]] T spvFMul(T l, T r)
|
||||
{
|
||||
@ -78,6 +66,18 @@ template<typename T>
|
||||
return fma(T(-1), r, l);
|
||||
}
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 gl_Position [[position]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float4 vA [[attribute(0)]];
|
||||
float4 vB [[attribute(1)]];
|
||||
float4 vC [[attribute(2)]];
|
||||
};
|
||||
|
||||
vertex main0_out main0(main0_in in [[stage_in]])
|
||||
{
|
||||
main0_out out = {};
|
||||
|
||||
@ -5,6 +5,13 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
// Implementation of the GLSL sign() function for integer types
|
||||
template<typename T, typename E = typename enable_if<is_integral<T>::value>::type>
|
||||
inline T sign(T x)
|
||||
{
|
||||
return select(select(select(x, T(0), x == T(0)), T(1), x > T(0)), T(-1), x < T(0));
|
||||
}
|
||||
|
||||
struct UBO
|
||||
{
|
||||
float4x4 uMVP;
|
||||
@ -36,13 +43,6 @@ struct main0_in
|
||||
float4 aVertex [[attribute(0)]];
|
||||
};
|
||||
|
||||
// Implementation of the GLSL sign() function for integer types
|
||||
template<typename T, typename E = typename enable_if<is_integral<T>::value>::type>
|
||||
inline T sign(T x)
|
||||
{
|
||||
return select(select(select(x, T(0), x == T(0)), T(1), x > T(0)), T(-1), x < T(0));
|
||||
}
|
||||
|
||||
vertex main0_out main0(main0_in in [[stage_in]], constant UBO& _21 [[buffer(0)]])
|
||||
{
|
||||
main0_out out = {};
|
||||
|
||||
@ -5,11 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 gl_Position [[position]];
|
||||
};
|
||||
|
||||
// Returns 2D texture coords corresponding to 1D texel buffer coords
|
||||
static inline __attribute__((always_inline))
|
||||
uint2 spvTexelBufferCoord(uint tc)
|
||||
@ -17,6 +12,11 @@ uint2 spvTexelBufferCoord(uint tc)
|
||||
return uint2(tc % 4096, tc / 4096);
|
||||
}
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 gl_Position [[position]];
|
||||
};
|
||||
|
||||
vertex main0_out main0(texture2d<float> uSamp [[texture(0)]], texture2d<float> uSampo [[texture(1)]])
|
||||
{
|
||||
main0_out out = {};
|
||||
|
||||
@ -5,6 +5,13 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
// Returns 2D texture coords corresponding to 1D texel buffer coords
|
||||
static inline __attribute__((always_inline))
|
||||
uint2 spvTexelBufferCoord(uint tc)
|
||||
{
|
||||
return uint2(tc % 4096, tc / 4096);
|
||||
}
|
||||
|
||||
struct type_View
|
||||
{
|
||||
float4x4 View_TranslatedWorldToClip;
|
||||
@ -281,13 +288,6 @@ struct main0_in
|
||||
float2 in_var_ATTRIBUTE0 [[attribute(0)]];
|
||||
};
|
||||
|
||||
// Returns 2D texture coords corresponding to 1D texel buffer coords
|
||||
static inline __attribute__((always_inline))
|
||||
uint2 spvTexelBufferCoord(uint tc)
|
||||
{
|
||||
return uint2(tc % 4096, tc / 4096);
|
||||
}
|
||||
|
||||
vertex main0_out main0(main0_in in [[stage_in]], constant type_View& View [[buffer(0)]], constant type_Primitive& Primitive [[buffer(1)]], constant type_MobileShadowDepthPass& MobileShadowDepthPass [[buffer(2)]], constant type_EmitterDynamicUniforms& EmitterDynamicUniforms [[buffer(3)]], constant type_EmitterUniforms& EmitterUniforms [[buffer(4)]], constant type_Globals& _Globals [[buffer(5)]], texture2d<float> ParticleIndices [[texture(0)]], texture2d<float> PositionTexture [[texture(1)]], texture2d<float> VelocityTexture [[texture(2)]], texture2d<float> AttributesTexture [[texture(3)]], texture2d<float> CurveTexture [[texture(4)]], sampler PositionTextureSampler [[sampler(0)]], sampler VelocityTextureSampler [[sampler(1)]], sampler AttributesTextureSampler [[sampler(2)]], sampler CurveTextureSampler [[sampler(3)]], uint gl_VertexIndex [[vertex_id]], uint gl_InstanceIndex [[instance_id]])
|
||||
{
|
||||
main0_out out = {};
|
||||
|
||||
@ -5,12 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
uint4 u;
|
||||
int4 i;
|
||||
};
|
||||
|
||||
// Implementation of the GLSL findLSB() function
|
||||
template<typename T>
|
||||
inline T spvFindLSB(T x)
|
||||
@ -33,6 +27,12 @@ inline T spvFindUMSB(T x)
|
||||
return select(clz(T(0)) - (clz(x) + T(1)), T(-1), x == T(0));
|
||||
}
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
uint4 u;
|
||||
int4 i;
|
||||
};
|
||||
|
||||
kernel void main0(device SSBO& _4 [[buffer(0)]])
|
||||
{
|
||||
uint4 _19 = _4.u;
|
||||
|
||||
@ -44,17 +44,6 @@ struct spvUnsafeArray
|
||||
}
|
||||
};
|
||||
|
||||
struct Block
|
||||
{
|
||||
uint2 _m0[2];
|
||||
uint2 _m1[2];
|
||||
};
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
Block _m0[3];
|
||||
};
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
|
||||
{
|
||||
@ -163,6 +152,17 @@ inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device
|
||||
}
|
||||
}
|
||||
|
||||
struct Block
|
||||
{
|
||||
uint2 _m0[2];
|
||||
uint2 _m1[2];
|
||||
};
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
Block _m0[3];
|
||||
};
|
||||
|
||||
kernel void main0(device SSBO& ssbo [[buffer(0)]])
|
||||
{
|
||||
threadgroup uint2 _18[2];
|
||||
|
||||
@ -5,17 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct Block
|
||||
{
|
||||
uint2 _m0[2];
|
||||
uint2 _m1[2];
|
||||
};
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
Block _m0[3];
|
||||
};
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
|
||||
{
|
||||
@ -124,6 +113,17 @@ inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device
|
||||
}
|
||||
}
|
||||
|
||||
struct Block
|
||||
{
|
||||
uint2 _m0[2];
|
||||
uint2 _m1[2];
|
||||
};
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
Block _m0[3];
|
||||
};
|
||||
|
||||
kernel void main0(device SSBO& ssbo [[buffer(0)]])
|
||||
{
|
||||
threadgroup uint2 _18[2];
|
||||
|
||||
@ -44,17 +44,6 @@ struct spvUnsafeArray
|
||||
}
|
||||
};
|
||||
|
||||
struct Block
|
||||
{
|
||||
uint2 _m0[2];
|
||||
uint2 _m1[2];
|
||||
};
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
Block _m0[3];
|
||||
};
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
|
||||
{
|
||||
@ -163,6 +152,17 @@ inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device
|
||||
}
|
||||
}
|
||||
|
||||
struct Block
|
||||
{
|
||||
uint2 _m0[2];
|
||||
uint2 _m1[2];
|
||||
};
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
Block _m0[3];
|
||||
};
|
||||
|
||||
kernel void main0(device SSBO& ssbo [[buffer(0)]], constant SSBO& ubo [[buffer(1)]])
|
||||
{
|
||||
threadgroup uint2 _18[2];
|
||||
|
||||
@ -5,17 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct Block
|
||||
{
|
||||
uint2 _m0[2];
|
||||
uint2 _m1[2];
|
||||
};
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
Block _m0[3];
|
||||
};
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
|
||||
{
|
||||
@ -124,6 +113,17 @@ inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device
|
||||
}
|
||||
}
|
||||
|
||||
struct Block
|
||||
{
|
||||
uint2 _m0[2];
|
||||
uint2 _m1[2];
|
||||
};
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
Block _m0[3];
|
||||
};
|
||||
|
||||
kernel void main0(device SSBO& ssbo [[buffer(0)]], constant SSBO& ubo [[buffer(1)]])
|
||||
{
|
||||
threadgroup uint2 _18[2];
|
||||
|
||||
@ -5,12 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
int4 ints;
|
||||
uint4 uints;
|
||||
};
|
||||
|
||||
// Implementation of the signed GLSL findMSB() function
|
||||
template<typename T>
|
||||
inline T spvFindSMSB(T x)
|
||||
@ -33,6 +27,12 @@ inline T sign(T x)
|
||||
return select(select(select(x, T(0), x == T(0)), T(1), x > T(0)), T(-1), x < T(0));
|
||||
}
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
int4 ints;
|
||||
uint4 uints;
|
||||
};
|
||||
|
||||
kernel void main0(device SSBO& _4 [[buffer(0)]])
|
||||
{
|
||||
int4 _19 = _4.ints;
|
||||
|
||||
@ -44,11 +44,6 @@ struct spvUnsafeArray
|
||||
}
|
||||
};
|
||||
|
||||
struct _3
|
||||
{
|
||||
float _m0[4];
|
||||
};
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
|
||||
{
|
||||
@ -157,6 +152,11 @@ inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device
|
||||
}
|
||||
}
|
||||
|
||||
struct _3
|
||||
{
|
||||
float _m0[4];
|
||||
};
|
||||
|
||||
fragment void main0()
|
||||
{
|
||||
spvUnsafeArray<float, 4> _20;
|
||||
|
||||
@ -44,8 +44,6 @@ struct spvUnsafeArray
|
||||
}
|
||||
};
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(8u, 1u, 1u);
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
|
||||
{
|
||||
@ -154,6 +152,8 @@ inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device
|
||||
}
|
||||
}
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(8u, 1u, 1u);
|
||||
|
||||
kernel void main0(uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
|
||||
{
|
||||
threadgroup float shared_group[8][8];
|
||||
|
||||
@ -5,27 +5,6 @@
|
||||
|
||||
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;
|
||||
};
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
||||
|
||||
// Implementation of the GLSL radians() function
|
||||
template<typename T>
|
||||
inline T radians(T d)
|
||||
@ -197,6 +176,27 @@ inline T spvFaceForward(T n, T i, T nref)
|
||||
return i * nref < T(0) ? n : -n;
|
||||
}
|
||||
|
||||
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;
|
||||
};
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
||||
|
||||
kernel void main0(device SSBO& _19 [[buffer(0)]])
|
||||
{
|
||||
_19.res = round(((device float*)&_19.f32)[0u]);
|
||||
|
||||
@ -5,13 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
float FragColor;
|
||||
};
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
||||
|
||||
template<typename T>
|
||||
inline T spvSubgroupBroadcast(T value, ushort lane)
|
||||
{
|
||||
@ -226,6 +219,13 @@ inline vec<bool, N> spvQuadSwap(vec<bool, N> value, uint dir)
|
||||
return (vec<bool, N>)quad_shuffle_xor((vec<ushort, N>)value, dir + 1);
|
||||
}
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
float FragColor;
|
||||
};
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
||||
|
||||
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));
|
||||
|
||||
@ -5,13 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
float FragColor;
|
||||
};
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
||||
|
||||
template<typename T>
|
||||
inline T spvSubgroupBroadcast(T value, ushort lane)
|
||||
{
|
||||
@ -226,6 +219,13 @@ inline vec<bool, N> spvQuadSwap(vec<bool, N> value, uint dir)
|
||||
return (vec<bool, N>)quad_shuffle_xor((vec<ushort, N>)value, dir + 1);
|
||||
}
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
float FragColor;
|
||||
};
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
||||
|
||||
kernel void main0(device SSBO& _9 [[buffer(0)]], uint gl_NumSubgroups [[simdgroups_per_threadgroup]], uint gl_SubgroupID [[simdgroup_index_in_threadgroup]], uint gl_SubgroupInvocationID [[thread_index_in_simdgroup]])
|
||||
{
|
||||
uint gl_SubgroupSize = 32;
|
||||
|
||||
@ -5,13 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
float FragColor;
|
||||
};
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
||||
|
||||
template<typename T>
|
||||
inline T spvSubgroupShuffle(T value, ushort lane)
|
||||
{
|
||||
@ -120,6 +113,13 @@ inline vec<bool, N> spvQuadSwap(vec<bool, N> value, uint dir)
|
||||
return (vec<bool, N>)quad_shuffle_xor((vec<ushort, N>)value, dir + 1);
|
||||
}
|
||||
|
||||
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);
|
||||
|
||||
@ -5,13 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
float FragColor;
|
||||
};
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
||||
|
||||
template<typename T>
|
||||
inline T spvSubgroupBroadcast(T value, ushort lane)
|
||||
{
|
||||
@ -221,6 +214,13 @@ inline vec<bool, N> spvQuadSwap(vec<bool, N> value, uint dir)
|
||||
return (vec<bool, N>)quad_shuffle_xor((vec<ushort, N>)value, dir + 1);
|
||||
}
|
||||
|
||||
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]])
|
||||
{
|
||||
uint4 gl_SubgroupEqMask = uint4(1 << gl_SubgroupInvocationID, uint3(0));
|
||||
|
||||
@ -5,13 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
float FragColor;
|
||||
};
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
||||
|
||||
template<typename T>
|
||||
inline T spvSubgroupBroadcast(T value, ushort lane)
|
||||
{
|
||||
@ -221,6 +214,13 @@ inline vec<bool, N> spvQuadSwap(vec<bool, N> value, uint dir)
|
||||
return (vec<bool, N>)quad_shuffle_xor((vec<ushort, N>)value, dir + 1);
|
||||
}
|
||||
|
||||
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]])
|
||||
{
|
||||
uint4 gl_SubgroupEqMask = uint4(1 << gl_SubgroupInvocationID, uint3(0));
|
||||
|
||||
@ -5,20 +5,6 @@
|
||||
|
||||
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)
|
||||
@ -40,6 +26,20 @@ inline T degrees(T r)
|
||||
return r * T(57.2957795131);
|
||||
}
|
||||
|
||||
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)]];
|
||||
};
|
||||
|
||||
static inline __attribute__((always_inline))
|
||||
half2x2 test_mat2(thread const half2& a, thread const half2& b, thread const half2& c, thread const half2& d)
|
||||
{
|
||||
|
||||
@ -5,11 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
template<typename T>
|
||||
inline T spvSubgroupBroadcast(T value, ushort lane)
|
||||
{
|
||||
@ -224,6 +219,11 @@ inline vec<bool, N> spvQuadSwap(vec<bool, N> value, uint dir)
|
||||
return (vec<bool, N>)quad_shuffle_xor((vec<ushort, N>)value, dir + 1);
|
||||
}
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
fragment main0_out main0(uint gl_SubgroupSize [[threads_per_simdgroup]], uint gl_SubgroupInvocationID [[thread_index_in_simdgroup]])
|
||||
{
|
||||
main0_out out = {};
|
||||
|
||||
@ -5,6 +5,13 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
// Returns 2D texture coords corresponding to 1D texel buffer coords
|
||||
static inline __attribute__((always_inline))
|
||||
uint2 spvTexelBufferCoord(uint tc)
|
||||
{
|
||||
return uint2(tc % 4096, tc / 4096);
|
||||
}
|
||||
|
||||
struct attr_desc
|
||||
{
|
||||
int type;
|
||||
@ -34,13 +41,6 @@ struct main0_out
|
||||
float4 gl_Position [[position]];
|
||||
};
|
||||
|
||||
// Returns 2D texture coords corresponding to 1D texel buffer coords
|
||||
static inline __attribute__((always_inline))
|
||||
uint2 spvTexelBufferCoord(uint tc)
|
||||
{
|
||||
return uint2(tc % 4096, tc / 4096);
|
||||
}
|
||||
|
||||
static inline __attribute__((always_inline))
|
||||
attr_desc fetch_desc(thread const int& location, constant VertexBuffer& v_227)
|
||||
{
|
||||
|
||||
@ -5,19 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
constant float4 _68[4] = { float4(0.0), float4(1.0), float4(2.0), float4(3.0) };
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 gl_Position [[position]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
int Index1 [[attribute(0)]];
|
||||
int Index2 [[attribute(1)]];
|
||||
};
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
|
||||
{
|
||||
@ -126,6 +113,19 @@ inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device
|
||||
}
|
||||
}
|
||||
|
||||
constant float4 _68[4] = { float4(0.0), float4(1.0), float4(2.0), float4(3.0) };
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 gl_Position [[position]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
int Index1 [[attribute(0)]];
|
||||
int Index2 [[attribute(1)]];
|
||||
};
|
||||
|
||||
static inline __attribute__((always_inline))
|
||||
float4 consume_constant_arrays2(thread const float4 (&positions)[4], thread const float4 (&positions2)[4], thread int& Index1, thread int& Index2)
|
||||
{
|
||||
|
||||
@ -5,11 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 fragColor [[color(0)]];
|
||||
};
|
||||
|
||||
// Returns 2D texture coords corresponding to 1D texel buffer coords
|
||||
static inline __attribute__((always_inline))
|
||||
uint2 spvTexelBufferCoord(uint tc)
|
||||
@ -137,6 +132,11 @@ inline vec<T, 4> spvGatherCompareSwizzle(const thread Tex<T>& t, sampler s, uint
|
||||
return t.gather_compare(s, spvForward<Ts>(params)...);
|
||||
}
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 fragColor [[color(0)]];
|
||||
};
|
||||
|
||||
static inline __attribute__((always_inline))
|
||||
float4 do_samples(thread const texture1d<float> t1, thread const sampler t1Smplr, constant uint& t1Swzl, thread const texture2d<float> t2, constant uint& t2Swzl, thread const texture3d<float> t3, thread const sampler t3Smplr, constant uint& t3Swzl, thread const texturecube<float> tc, constant uint& tcSwzl, thread const texture2d_array<float> t2a, thread const sampler t2aSmplr, constant uint& t2aSwzl, thread const texturecube_array<float> tca, thread const sampler tcaSmplr, constant uint& tcaSwzl, thread const texture2d<float> tb, thread const depth2d<float> d2, thread const sampler d2Smplr, constant uint& d2Swzl, thread const depthcube<float> dc, thread const sampler dcSmplr, constant uint& dcSwzl, thread const depth2d_array<float> d2a, constant uint& d2aSwzl, thread const depthcube_array<float> dca, thread const sampler dcaSmplr, constant uint& dcaSwzl, thread sampler defaultSampler, thread sampler shadowSampler)
|
||||
{
|
||||
|
||||
@ -7,11 +7,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct u0_counters
|
||||
{
|
||||
uint c;
|
||||
};
|
||||
|
||||
// Returns 2D texture coords corresponding to 1D texel buffer coords
|
||||
static inline __attribute__((always_inline))
|
||||
uint2 spvTexelBufferCoord(uint tc)
|
||||
@ -19,6 +14,11 @@ uint2 spvTexelBufferCoord(uint tc)
|
||||
return uint2(tc % 4096, tc / 4096);
|
||||
}
|
||||
|
||||
struct u0_counters
|
||||
{
|
||||
uint c;
|
||||
};
|
||||
|
||||
kernel void main0(device u0_counters& u0_counter [[buffer(0)]], texture2d<uint, access::write> u0 [[texture(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||||
{
|
||||
uint _29 = atomic_fetch_sub_explicit((device atomic_uint*)&u0_counter.c, 1, memory_order_relaxed);
|
||||
|
||||
@ -7,11 +7,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct u0_counters
|
||||
{
|
||||
uint c;
|
||||
};
|
||||
|
||||
// Returns 2D texture coords corresponding to 1D texel buffer coords
|
||||
static inline __attribute__((always_inline))
|
||||
uint2 spvTexelBufferCoord(uint tc)
|
||||
@ -19,6 +14,11 @@ uint2 spvTexelBufferCoord(uint tc)
|
||||
return uint2(tc % 4096, tc / 4096);
|
||||
}
|
||||
|
||||
struct u0_counters
|
||||
{
|
||||
uint c;
|
||||
};
|
||||
|
||||
kernel void main0(device u0_counters& u0_counter [[buffer(0)]], texture2d<uint, access::write> u0 [[texture(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||||
{
|
||||
uint _29 = atomic_fetch_add_explicit((device atomic_uint*)&u0_counter.c, 1, memory_order_relaxed);
|
||||
|
||||
@ -5,11 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct cb5_struct
|
||||
{
|
||||
float4 _m0[5];
|
||||
};
|
||||
|
||||
// Returns 2D texture coords corresponding to 1D texel buffer coords
|
||||
static inline __attribute__((always_inline))
|
||||
uint2 spvTexelBufferCoord(uint tc)
|
||||
@ -17,6 +12,11 @@ uint2 spvTexelBufferCoord(uint tc)
|
||||
return uint2(tc % 4096, tc / 4096);
|
||||
}
|
||||
|
||||
struct cb5_struct
|
||||
{
|
||||
float4 _m0[5];
|
||||
};
|
||||
|
||||
kernel void main0(constant cb5_struct& cb0_5 [[buffer(0)]], texture2d<uint, access::write> u0 [[texture(0)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
|
||||
{
|
||||
float4 r0;
|
||||
|
||||
@ -5,11 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct cb
|
||||
{
|
||||
float value;
|
||||
};
|
||||
|
||||
// Returns 2D texture coords corresponding to 1D texel buffer coords
|
||||
static inline __attribute__((always_inline))
|
||||
uint2 spvTexelBufferCoord(uint tc)
|
||||
@ -17,6 +12,11 @@ uint2 spvTexelBufferCoord(uint tc)
|
||||
return uint2(tc % 4096, tc / 4096);
|
||||
}
|
||||
|
||||
struct cb
|
||||
{
|
||||
float value;
|
||||
};
|
||||
|
||||
kernel void main0(constant cb& _6 [[buffer(0)]], texture2d<float, access::write> _buffer [[texture(0)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
|
||||
{
|
||||
_buffer.write(float4(_6.value), spvTexelBufferCoord(((32u * gl_WorkGroupID.x) + gl_LocalInvocationIndex)));
|
||||
|
||||
@ -5,6 +5,17 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template <typename F> struct SpvHalfTypeSelector;
|
||||
template <> struct SpvHalfTypeSelector<float> { public: using H = half; };
|
||||
template<uint N> struct SpvHalfTypeSelector<vec<float, N>> { using H = vec<half, N>; };
|
||||
template<typename F, typename H = typename SpvHalfTypeSelector<F>::H>
|
||||
[[clang::optnone]] F spvQuantizeToF16(F fval)
|
||||
{
|
||||
H hval = H(fval);
|
||||
hval = select(copysign(H(0), hval), hval, isnormal(hval) || isinf(hval) || isnan(hval));
|
||||
return F(hval);
|
||||
}
|
||||
|
||||
struct SSBO0
|
||||
{
|
||||
float scalar;
|
||||
@ -13,15 +24,6 @@ struct SSBO0
|
||||
float4 vec4_val;
|
||||
};
|
||||
|
||||
template <typename F> struct SpvHalfTypeSelector;
|
||||
template <> struct SpvHalfTypeSelector<float> { public: using H = half; };
|
||||
template<uint N> struct SpvHalfTypeSelector<vec<float, N>> { using H = vec<half, N>; };
|
||||
template<typename F, typename H = typename SpvHalfTypeSelector<F>::H>
|
||||
[[clang::optnone]] F spvQuantizeToF16(F val)
|
||||
{
|
||||
return F(H(val));
|
||||
}
|
||||
|
||||
kernel void main0(device SSBO0& _4 [[buffer(0)]])
|
||||
{
|
||||
_4.scalar = spvQuantizeToF16(_4.scalar);
|
||||
|
||||
@ -44,6 +44,13 @@ struct spvUnsafeArray
|
||||
}
|
||||
};
|
||||
|
||||
// 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);
|
||||
}
|
||||
|
||||
struct myType
|
||||
{
|
||||
float data;
|
||||
@ -54,13 +61,6 @@ struct main0_out
|
||||
float4 o_color [[color(0)]];
|
||||
};
|
||||
|
||||
// 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);
|
||||
}
|
||||
|
||||
fragment main0_out main0(float4 gl_FragCoord [[position]])
|
||||
{
|
||||
spvUnsafeArray<myType, 5> _21 = spvUnsafeArray<myType, 5>({ myType{ 0.0 }, myType{ 1.0 }, myType{ 0.0 }, myType{ 1.0 }, myType{ 0.0 } });
|
||||
|
||||
@ -1,8 +1,21 @@
|
||||
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
||||
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template <typename F> struct SpvHalfTypeSelector;
|
||||
template <> struct SpvHalfTypeSelector<float> { public: using H = half; };
|
||||
template<uint N> struct SpvHalfTypeSelector<vec<float, N>> { using H = vec<half, N>; };
|
||||
template<typename F, typename H = typename SpvHalfTypeSelector<F>::H>
|
||||
[[clang::optnone]] F spvQuantizeToF16(F fval)
|
||||
{
|
||||
H hval = H(fval);
|
||||
hval = select(copysign(H(0), hval), hval, isnormal(hval) || isinf(hval) || isnan(hval));
|
||||
return F(hval);
|
||||
}
|
||||
|
||||
constant int _7_tmp [[function_constant(201)]];
|
||||
constant int _7 = is_function_constant_defined(_7_tmp) ? _7_tmp : -10;
|
||||
constant int _20 = (_7 + 2);
|
||||
@ -15,7 +28,7 @@ constant int2 _34 = int2(_32.y, _32.x);
|
||||
constant int _35 = _32.y;
|
||||
constant float _9_tmp [[function_constant(200)]];
|
||||
constant float _9 = is_function_constant_defined(_9_tmp) ? _9_tmp : 3.141590118408203125;
|
||||
constant float _41 = float(half(_9));
|
||||
constant float _41 = spvQuantizeToF16(_9);
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
|
||||
@ -5,30 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct Data
|
||||
{
|
||||
float a;
|
||||
float b;
|
||||
};
|
||||
|
||||
constant float X_tmp [[function_constant(0)]];
|
||||
constant float X = is_function_constant_defined(X_tmp) ? X_tmp : 4.0;
|
||||
|
||||
struct Data_1
|
||||
{
|
||||
float a;
|
||||
float b;
|
||||
};
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
Data_1 outdata[1];
|
||||
};
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(2u, 1u, 1u);
|
||||
|
||||
constant Data _25[2] = { Data{ 1.0, 2.0 }, Data{ 3.0, 4.0 } };
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
|
||||
{
|
||||
@ -137,6 +113,30 @@ inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device
|
||||
}
|
||||
}
|
||||
|
||||
struct Data
|
||||
{
|
||||
float a;
|
||||
float b;
|
||||
};
|
||||
|
||||
constant float X_tmp [[function_constant(0)]];
|
||||
constant float X = is_function_constant_defined(X_tmp) ? X_tmp : 4.0;
|
||||
|
||||
struct Data_1
|
||||
{
|
||||
float a;
|
||||
float b;
|
||||
};
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
Data_1 outdata[1];
|
||||
};
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(2u, 1u, 1u);
|
||||
|
||||
constant Data _25[2] = { Data{ 1.0, 2.0 }, Data{ 3.0, 4.0 } };
|
||||
|
||||
static inline __attribute__((always_inline))
|
||||
Data combine(thread const Data& a, thread const Data& b)
|
||||
{
|
||||
|
||||
@ -5,20 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct BUF
|
||||
{
|
||||
int a;
|
||||
float b;
|
||||
float c;
|
||||
};
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
||||
|
||||
constant float _16[2] = { 1.0, 2.0 };
|
||||
constant float _19[2] = { 3.0, 4.0 };
|
||||
constant float _20[2][2] = { { 1.0, 2.0 }, { 3.0, 4.0 } };
|
||||
constant float _21[2][2][2] = { { { 1.0, 2.0 }, { 3.0, 4.0 } }, { { 1.0, 2.0 }, { 3.0, 4.0 } } };
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
|
||||
{
|
||||
@ -343,6 +329,20 @@ inline void spvArrayCopyFromDeviceToThreadGroup3(threadgroup T (&dst)[A][B][C],
|
||||
}
|
||||
}
|
||||
|
||||
struct BUF
|
||||
{
|
||||
int a;
|
||||
float b;
|
||||
float c;
|
||||
};
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
||||
|
||||
constant float _16[2] = { 1.0, 2.0 };
|
||||
constant float _19[2] = { 3.0, 4.0 };
|
||||
constant float _20[2][2] = { { 1.0, 2.0 }, { 3.0, 4.0 } };
|
||||
constant float _21[2][2][2] = { { { 1.0, 2.0 }, { 3.0, 4.0 } }, { { 1.0, 2.0 }, { 3.0, 4.0 } } };
|
||||
|
||||
kernel void main0(device BUF& o [[buffer(0)]])
|
||||
{
|
||||
float c[2][2][2];
|
||||
|
||||
@ -5,12 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct myBlock
|
||||
{
|
||||
int a;
|
||||
float b[1];
|
||||
};
|
||||
|
||||
// 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)
|
||||
@ -18,6 +12,12 @@ inline Tx mod(Tx x, Ty y)
|
||||
return x - y * floor(x / y);
|
||||
}
|
||||
|
||||
struct myBlock
|
||||
{
|
||||
int a;
|
||||
float b[1];
|
||||
};
|
||||
|
||||
static inline __attribute__((always_inline))
|
||||
float getB(device myBlock& myStorage, thread uint3& gl_GlobalInvocationID)
|
||||
{
|
||||
|
||||
@ -5,12 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct myBlock
|
||||
{
|
||||
int a;
|
||||
float b[1];
|
||||
};
|
||||
|
||||
// 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)
|
||||
@ -18,6 +12,12 @@ inline Tx mod(Tx x, Ty y)
|
||||
return x - y * floor(x / y);
|
||||
}
|
||||
|
||||
struct myBlock
|
||||
{
|
||||
int a;
|
||||
float b[1];
|
||||
};
|
||||
|
||||
kernel void main0(device myBlock& myStorage [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||||
{
|
||||
myStorage.a = (myStorage.a + 1) % 256;
|
||||
|
||||
@ -7,6 +7,12 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
// The required alignment of a linear texture of R32Uint format.
|
||||
constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
|
||||
constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
|
||||
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
|
||||
#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() + spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
float4 outdata;
|
||||
@ -23,12 +29,6 @@ struct spvDescriptorSetBuffer0
|
||||
sampler uTextureSmplr [[id(4)]];
|
||||
};
|
||||
|
||||
// The required alignment of a linear texture of R32Uint format.
|
||||
constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
|
||||
constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
|
||||
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
|
||||
#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() + spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
|
||||
|
||||
kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||||
{
|
||||
uint _26 = atomic_fetch_add_explicit((device atomic_uint*)&spvDescriptorSet0.uImage_atomic[spvImage2DAtomicCoord(int2(gl_GlobalInvocationID.xy), spvDescriptorSet0.uImage)], 10u, memory_order_relaxed);
|
||||
|
||||
@ -7,6 +7,12 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
// The required alignment of a linear texture of R32Uint format.
|
||||
constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
|
||||
constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
|
||||
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
|
||||
#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() + spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
float4 outdata;
|
||||
@ -14,12 +20,6 @@ struct SSBO
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
||||
|
||||
// The required alignment of a linear texture of R32Uint format.
|
||||
constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
|
||||
constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
|
||||
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
|
||||
#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() + spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
|
||||
|
||||
kernel void main0(device SSBO& _31 [[buffer(1)]], texture2d<uint> uImage [[texture(0)]], device atomic_uint* uImage_atomic [[buffer(0)]], texture2d<float> uTexture [[texture(1)]], sampler uTextureSmplr [[sampler(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||||
{
|
||||
uint _26 = atomic_fetch_add_explicit((device atomic_uint*)&uImage_atomic[spvImage2DAtomicCoord(int2(gl_GlobalInvocationID.xy), uImage)], 10u, memory_order_relaxed);
|
||||
|
||||
@ -5,22 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct MatrixOut
|
||||
{
|
||||
float2x2 m2out;
|
||||
float3x3 m3out;
|
||||
float4x4 m4out;
|
||||
};
|
||||
|
||||
struct MatrixIn
|
||||
{
|
||||
float2x2 m2in;
|
||||
float3x3 m3in;
|
||||
float4x4 m4in;
|
||||
};
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
||||
|
||||
// Returns the determinant of a 2x2 matrix.
|
||||
static inline __attribute__((always_inline))
|
||||
float spvDet2x2(float a1, float a2, float b1, float b2)
|
||||
@ -121,6 +105,22 @@ float2x2 spvInverse2x2(float2x2 m)
|
||||
return (det != 0.0f) ? (adj * (1.0f / det)) : m;
|
||||
}
|
||||
|
||||
struct MatrixOut
|
||||
{
|
||||
float2x2 m2out;
|
||||
float3x3 m3out;
|
||||
float4x4 m4out;
|
||||
};
|
||||
|
||||
struct MatrixIn
|
||||
{
|
||||
float2x2 m2in;
|
||||
float3x3 m3in;
|
||||
float4x4 m4in;
|
||||
};
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
||||
|
||||
kernel void main0(device MatrixOut& _15 [[buffer(0)]], const device MatrixIn& _20 [[buffer(1)]])
|
||||
{
|
||||
_15.m2out = spvInverse2x2(_20.m2in);
|
||||
|
||||
@ -5,12 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct myBlock
|
||||
{
|
||||
int a;
|
||||
float b[1];
|
||||
};
|
||||
|
||||
// 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)
|
||||
@ -18,6 +12,12 @@ inline Tx mod(Tx x, Ty y)
|
||||
return x - y * floor(x / y);
|
||||
}
|
||||
|
||||
struct myBlock
|
||||
{
|
||||
int a;
|
||||
float b[1];
|
||||
};
|
||||
|
||||
kernel void main0(device myBlock& myStorage [[buffer(0)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
|
||||
{
|
||||
myStorage.a = (myStorage.a + 1) % 256;
|
||||
|
||||
@ -5,12 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct myBlock
|
||||
{
|
||||
int a;
|
||||
float b[1];
|
||||
};
|
||||
|
||||
// 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)
|
||||
@ -18,6 +12,12 @@ inline Tx mod(Tx x, Ty y)
|
||||
return x - y * floor(x / y);
|
||||
}
|
||||
|
||||
struct myBlock
|
||||
{
|
||||
int a;
|
||||
float b[1];
|
||||
};
|
||||
|
||||
kernel void main0(device myBlock& myStorage [[buffer(0)]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
|
||||
{
|
||||
myStorage.a = (myStorage.a + 1) % 256;
|
||||
|
||||
@ -5,6 +5,13 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
// 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);
|
||||
}
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
float4 in_data[1];
|
||||
@ -17,13 +24,6 @@ struct SSBO2
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
||||
|
||||
// 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);
|
||||
}
|
||||
|
||||
kernel void main0(const device SSBO& _23 [[buffer(0)]], device SSBO2& _33 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||||
{
|
||||
uint ident = gl_GlobalInvocationID.x;
|
||||
|
||||
@ -5,12 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct myBlock
|
||||
{
|
||||
int a;
|
||||
float b;
|
||||
};
|
||||
|
||||
// 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)
|
||||
@ -18,6 +12,12 @@ inline Tx mod(Tx x, Ty y)
|
||||
return x - y * floor(x / y);
|
||||
}
|
||||
|
||||
struct myBlock
|
||||
{
|
||||
int a;
|
||||
float b;
|
||||
};
|
||||
|
||||
kernel void main0(device myBlock& myStorage [[buffer(0)]])
|
||||
{
|
||||
myStorage.a = (myStorage.a + 1) % 256;
|
||||
|
||||
@ -5,35 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct UBO
|
||||
{
|
||||
uint index;
|
||||
};
|
||||
|
||||
struct UBO2
|
||||
{
|
||||
uint index2;
|
||||
};
|
||||
|
||||
struct spvDescriptorSetBuffer0
|
||||
{
|
||||
array<texture2d<float>, 4> uSampler [[id(0)]];
|
||||
array<sampler, 4> uSamplerSmplr [[id(4)]];
|
||||
constant UBO* uUBO [[id(8)]];
|
||||
constant UBO2* m_50 [[id(9)]];
|
||||
constant uint* spvSwizzleConstants [[id(10)]];
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float2 vUV [[user(locn0)]];
|
||||
};
|
||||
|
||||
template<typename T> struct spvRemoveReference { typedef T type; };
|
||||
template<typename T> struct spvRemoveReference<thread T&> { typedef T type; };
|
||||
template<typename T> struct spvRemoveReference<thread T&&> { typedef T type; };
|
||||
@ -94,6 +65,35 @@ inline T spvTextureSwizzle(T x, uint s)
|
||||
return spvTextureSwizzle(vec<T, 4>(x, 0, 0, 1), s).x;
|
||||
}
|
||||
|
||||
struct UBO
|
||||
{
|
||||
uint index;
|
||||
};
|
||||
|
||||
struct UBO2
|
||||
{
|
||||
uint index2;
|
||||
};
|
||||
|
||||
struct spvDescriptorSetBuffer0
|
||||
{
|
||||
array<texture2d<float>, 4> uSampler [[id(0)]];
|
||||
array<sampler, 4> uSamplerSmplr [[id(4)]];
|
||||
constant UBO* uUBO [[id(8)]];
|
||||
constant UBO2* m_50 [[id(9)]];
|
||||
constant uint* spvSwizzleConstants [[id(10)]];
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float2 vUV [[user(locn0)]];
|
||||
};
|
||||
|
||||
static inline __attribute__((always_inline))
|
||||
float4 sample_in_func(thread const array<texture2d<float>, 4> uSampler, thread const array<sampler, 4> uSamplerSmplr, constant uint* uSamplerSwzl, constant UBO& uUBO, thread float2& vUV)
|
||||
{
|
||||
|
||||
@ -5,26 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct UBO
|
||||
{
|
||||
uint index;
|
||||
};
|
||||
|
||||
struct UBO2
|
||||
{
|
||||
uint index2;
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float2 vUV [[user(locn0)]];
|
||||
};
|
||||
|
||||
template<typename T> struct spvRemoveReference { typedef T type; };
|
||||
template<typename T> struct spvRemoveReference<thread T&> { typedef T type; };
|
||||
template<typename T> struct spvRemoveReference<thread T&&> { typedef T type; };
|
||||
@ -85,6 +65,26 @@ inline T spvTextureSwizzle(T x, uint s)
|
||||
return spvTextureSwizzle(vec<T, 4>(x, 0, 0, 1), s).x;
|
||||
}
|
||||
|
||||
struct UBO
|
||||
{
|
||||
uint index;
|
||||
};
|
||||
|
||||
struct UBO2
|
||||
{
|
||||
uint index2;
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float2 vUV [[user(locn0)]];
|
||||
};
|
||||
|
||||
static inline __attribute__((always_inline))
|
||||
float4 sample_in_func(thread const array<texture2d<float>, 4> uSampler, thread const array<sampler, 4> uSamplerSmplr, constant uint* uSamplerSwzl, constant UBO& uUBO, thread float2& vUV)
|
||||
{
|
||||
|
||||
@ -5,23 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct spvDescriptorSetBuffer0
|
||||
{
|
||||
array<texture2d<float>, 4> uSampler0 [[id(0)]];
|
||||
array<sampler, 4> uSampler0Smplr [[id(4)]];
|
||||
constant uint* spvSwizzleConstants [[id(8)]];
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float2 vUV [[user(locn0)]];
|
||||
};
|
||||
|
||||
template<typename T> struct spvRemoveReference { typedef T type; };
|
||||
template<typename T> struct spvRemoveReference<thread T&> { typedef T type; };
|
||||
template<typename T> struct spvRemoveReference<thread T&&> { typedef T type; };
|
||||
@ -82,6 +65,23 @@ inline T spvTextureSwizzle(T x, uint s)
|
||||
return spvTextureSwizzle(vec<T, 4>(x, 0, 0, 1), s).x;
|
||||
}
|
||||
|
||||
struct spvDescriptorSetBuffer0
|
||||
{
|
||||
array<texture2d<float>, 4> uSampler0 [[id(0)]];
|
||||
array<sampler, 4> uSampler0Smplr [[id(4)]];
|
||||
constant uint* spvSwizzleConstants [[id(8)]];
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float2 vUV [[user(locn0)]];
|
||||
};
|
||||
|
||||
static inline __attribute__((always_inline))
|
||||
float4 sample_in_func_1(thread const array<texture2d<float>, 4> uSampler0, thread const array<sampler, 4> uSampler0Smplr, constant uint* uSampler0Swzl, thread float2& vUV)
|
||||
{
|
||||
|
||||
@ -5,16 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float2 vUV [[user(locn0)]];
|
||||
};
|
||||
|
||||
template<typename T> struct spvRemoveReference { typedef T type; };
|
||||
template<typename T> struct spvRemoveReference<thread T&> { typedef T type; };
|
||||
template<typename T> struct spvRemoveReference<thread T&&> { typedef T type; };
|
||||
@ -75,6 +65,16 @@ inline T spvTextureSwizzle(T x, uint s)
|
||||
return spvTextureSwizzle(vec<T, 4>(x, 0, 0, 1), s).x;
|
||||
}
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float2 vUV [[user(locn0)]];
|
||||
};
|
||||
|
||||
static inline __attribute__((always_inline))
|
||||
float4 sample_in_func(thread const array<texture2d<float>, 4> uSampler, thread const array<sampler, 4> uSamplerSmplr, constant uint* uSamplerSwzl, thread float2& vUV)
|
||||
{
|
||||
|
||||
@ -5,11 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
// Returns 2D texture coords corresponding to 1D texel buffer coords
|
||||
static inline __attribute__((always_inline))
|
||||
uint2 spvTexelBufferCoord(uint tc)
|
||||
@ -17,6 +12,11 @@ uint2 spvTexelBufferCoord(uint tc)
|
||||
return uint2(tc % 4096, tc / 4096);
|
||||
}
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
fragment main0_out main0(texture2d<float> buf [[texture(0)]], texture2d<float, access::write> bufOut [[texture(1)]], float4 gl_FragCoord [[position]])
|
||||
{
|
||||
main0_out out = {};
|
||||
|
||||
@ -44,6 +44,13 @@ struct spvUnsafeArray
|
||||
}
|
||||
};
|
||||
|
||||
// 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);
|
||||
}
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor_0 [[color(0)]];
|
||||
@ -58,13 +65,6 @@ struct main0_in
|
||||
float4 vB [[user(locn1)]];
|
||||
};
|
||||
|
||||
// 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);
|
||||
}
|
||||
|
||||
static inline __attribute__((always_inline))
|
||||
void write_deeper_in_function(thread spvUnsafeArray<float4, 4>& FragColor, thread float4& vA, thread float4& vB)
|
||||
{
|
||||
|
||||
@ -7,6 +7,12 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
// The required alignment of a linear texture of R32Uint format.
|
||||
constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
|
||||
constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
|
||||
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
|
||||
#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() + spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
|
||||
|
||||
struct Buffer3
|
||||
{
|
||||
int baz;
|
||||
@ -35,12 +41,6 @@ struct spvDescriptorSetBuffer0
|
||||
device Buffer2* m_52 [[id(7), raster_order_group(0)]];
|
||||
};
|
||||
|
||||
// The required alignment of a linear texture of R32Uint format.
|
||||
constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
|
||||
constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
|
||||
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
|
||||
#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() + spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
|
||||
|
||||
fragment void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]])
|
||||
{
|
||||
(*spvDescriptorSet0.m_9).baz = 0;
|
||||
|
||||
@ -7,6 +7,12 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
// The required alignment of a linear texture of R32Uint format.
|
||||
constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
|
||||
constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
|
||||
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
|
||||
#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() + spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
|
||||
|
||||
struct Buffer3
|
||||
{
|
||||
int baz;
|
||||
@ -23,12 +29,6 @@ struct Buffer2
|
||||
uint quux;
|
||||
};
|
||||
|
||||
// The required alignment of a linear texture of R32Uint format.
|
||||
constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
|
||||
constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
|
||||
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
|
||||
#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() + spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)
|
||||
|
||||
fragment void main0(device Buffer3& _9 [[buffer(0)]], volatile device Buffer& _42 [[buffer(2), raster_order_group(0)]], device Buffer2& _52 [[buffer(3), raster_order_group(0)]], texture2d<float, access::write> img4 [[texture(0)]], texture2d<float, access::write> img [[texture(1), raster_order_group(0)]], texture2d<float> img3 [[texture(2), raster_order_group(0)]], texture2d<uint> img2 [[texture(3), raster_order_group(0)]], device atomic_uint* img2_atomic [[buffer(1), raster_order_group(0)]])
|
||||
{
|
||||
_9.baz = 0;
|
||||
|
||||
@ -5,16 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float3 vRefract [[user(locn0)]];
|
||||
};
|
||||
|
||||
template<typename T>
|
||||
[[clang::optnone]] T spvReflect(T i, T n)
|
||||
{
|
||||
@ -37,6 +27,16 @@ inline T spvRefract(T i, T n, T eta)
|
||||
}
|
||||
}
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float3 vRefract [[user(locn0)]];
|
||||
};
|
||||
|
||||
fragment main0_out main0(main0_in in [[stage_in]])
|
||||
{
|
||||
main0_out out = {};
|
||||
|
||||
@ -5,11 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
uint2 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
inline uint spvSubgroupBallotFindLSB(uint4 ballot, uint gl_SubgroupSize)
|
||||
{
|
||||
uint4 mask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupSize, 32u)), extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupSize - 32, 0)), uint2(0));
|
||||
@ -47,6 +42,11 @@ inline uint spvSubgroupBallotExclusiveBitCount(uint4 ballot, uint gl_SubgroupInv
|
||||
return spvPopCount4(ballot & mask);
|
||||
}
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
uint2 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
static inline __attribute__((always_inline))
|
||||
uint sub1(thread uint& gl_SubgroupSize)
|
||||
{
|
||||
|
||||
@ -5,16 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float4 vUV [[user(locn0)]];
|
||||
};
|
||||
|
||||
static inline __attribute__((always_inline))
|
||||
float3 spvCubemapTo2DArrayFace(float3 P)
|
||||
{
|
||||
@ -49,6 +39,16 @@ float3 spvCubemapTo2DArrayFace(float3 P)
|
||||
return float3(u, v, CubeFace);
|
||||
}
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float4 vUV [[user(locn0)]];
|
||||
};
|
||||
|
||||
fragment main0_out main0(main0_in in [[stage_in]], texturecube<float> cubeSampler [[texture(0)]], texture2d_array<float> cubeArraySampler [[texture(1)]], texture2d_array<float> texArraySampler [[texture(2)]], sampler cubeSamplerSmplr [[sampler(0)]], sampler cubeArraySamplerSmplr [[sampler(1)]], sampler texArraySamplerSmplr [[sampler(2)]])
|
||||
{
|
||||
main0_out out = {};
|
||||
|
||||
@ -44,22 +44,6 @@ struct spvUnsafeArray
|
||||
}
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 gl_Position;
|
||||
};
|
||||
|
||||
struct main0_patchOut
|
||||
{
|
||||
spvUnsafeArray<float4, 2> pFoo;
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
spvUnsafeArray<float4, 2> iFoo;
|
||||
float4 ipFoo;
|
||||
};
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
|
||||
{
|
||||
@ -168,6 +152,22 @@ inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device
|
||||
}
|
||||
}
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 gl_Position;
|
||||
};
|
||||
|
||||
struct main0_patchOut
|
||||
{
|
||||
spvUnsafeArray<float4, 2> pFoo;
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
spvUnsafeArray<float4, 2> iFoo;
|
||||
float4 ipFoo;
|
||||
};
|
||||
|
||||
kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], device main0_out* spvOut [[buffer(28)]], constant uint* spvIndirectParams [[buffer(29)]], device main0_patchOut* spvPatchOut [[buffer(27)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], device main0_in* spvIn [[buffer(22)]])
|
||||
{
|
||||
device main0_out* gl_out = &spvOut[gl_GlobalInvocationID.x - gl_GlobalInvocationID.x % 4];
|
||||
|
||||
@ -44,23 +44,6 @@ struct spvUnsafeArray
|
||||
}
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 gl_Position;
|
||||
};
|
||||
|
||||
struct main0_patchOut
|
||||
{
|
||||
spvUnsafeArray<float4, 2> pFoo;
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float4 iFoo_0 [[attribute(0)]];
|
||||
float4 iFoo_1 [[attribute(1)]];
|
||||
float4 ipFoo [[attribute(2)]];
|
||||
};
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
|
||||
{
|
||||
@ -169,6 +152,23 @@ inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device
|
||||
}
|
||||
}
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 gl_Position;
|
||||
};
|
||||
|
||||
struct main0_patchOut
|
||||
{
|
||||
spvUnsafeArray<float4, 2> pFoo;
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float4 iFoo_0 [[attribute(0)]];
|
||||
float4 iFoo_1 [[attribute(1)]];
|
||||
float4 ipFoo [[attribute(2)]];
|
||||
};
|
||||
|
||||
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 main0_patchOut* spvPatchOut [[buffer(27)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], threadgroup main0_in* gl_in [[threadgroup(0)]])
|
||||
{
|
||||
threadgroup float4 Foo[4][2];
|
||||
|
||||
@ -44,30 +44,6 @@ struct spvUnsafeArray
|
||||
}
|
||||
};
|
||||
|
||||
struct Matrices
|
||||
{
|
||||
float4x4 vpMatrix;
|
||||
float4x4 wMatrix;
|
||||
float4x3 wMatrix4x3;
|
||||
float3x4 wMatrix3x4;
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float3 OutNormal [[user(locn0)]];
|
||||
float4 OutWorldPos_0 [[user(locn1)]];
|
||||
float4 OutWorldPos_1 [[user(locn2)]];
|
||||
float4 OutWorldPos_2 [[user(locn3)]];
|
||||
float4 OutWorldPos_3 [[user(locn4)]];
|
||||
float4 gl_Position [[position]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float3 InPos [[attribute(0)]];
|
||||
float3 InNormal [[attribute(1)]];
|
||||
};
|
||||
|
||||
template<typename T>
|
||||
[[clang::optnone]] T spvFMul(T l, T r)
|
||||
{
|
||||
@ -117,6 +93,30 @@ template<typename T, int LCols, int LRows, int RCols, int RRows>
|
||||
return res;
|
||||
}
|
||||
|
||||
struct Matrices
|
||||
{
|
||||
float4x4 vpMatrix;
|
||||
float4x4 wMatrix;
|
||||
float4x3 wMatrix4x3;
|
||||
float3x4 wMatrix3x4;
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float3 OutNormal [[user(locn0)]];
|
||||
float4 OutWorldPos_0 [[user(locn1)]];
|
||||
float4 OutWorldPos_1 [[user(locn2)]];
|
||||
float4 OutWorldPos_2 [[user(locn3)]];
|
||||
float4 OutWorldPos_3 [[user(locn4)]];
|
||||
float4 gl_Position [[position]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float3 InPos [[attribute(0)]];
|
||||
float3 InNormal [[attribute(1)]];
|
||||
};
|
||||
|
||||
vertex main0_out main0(main0_in in [[stage_in]], constant Matrices& _22 [[buffer(0)]])
|
||||
{
|
||||
main0_out out = {};
|
||||
|
||||
@ -5,30 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct UBO
|
||||
{
|
||||
float4x4 uMVP;
|
||||
float3 rotDeg;
|
||||
float3 rotRad;
|
||||
int2 bits;
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float3 vNormal [[user(locn0)]];
|
||||
float3 vRotDeg [[user(locn1)]];
|
||||
float3 vRotRad [[user(locn2)]];
|
||||
int2 vLSB [[user(locn3)]];
|
||||
int2 vMSB [[user(locn4)]];
|
||||
float4 gl_Position [[position]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float4 aVertex [[attribute(0)]];
|
||||
float3 aNormal [[attribute(1)]];
|
||||
};
|
||||
|
||||
// Implementation of the GLSL radians() function
|
||||
template<typename T>
|
||||
inline T radians(T d)
|
||||
@ -108,6 +84,30 @@ float4x4 spvInverse4x4(float4x4 m)
|
||||
return (det != 0.0f) ? (adj * (1.0f / det)) : m;
|
||||
}
|
||||
|
||||
struct UBO
|
||||
{
|
||||
float4x4 uMVP;
|
||||
float3 rotDeg;
|
||||
float3 rotRad;
|
||||
int2 bits;
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float3 vNormal [[user(locn0)]];
|
||||
float3 vRotDeg [[user(locn1)]];
|
||||
float3 vRotRad [[user(locn2)]];
|
||||
int2 vLSB [[user(locn3)]];
|
||||
int2 vMSB [[user(locn4)]];
|
||||
float4 gl_Position [[position]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float4 aVertex [[attribute(0)]];
|
||||
float3 aNormal [[attribute(1)]];
|
||||
};
|
||||
|
||||
vertex main0_out main0(main0_in in [[stage_in]], constant UBO& _18 [[buffer(0)]])
|
||||
{
|
||||
main0_out out = {};
|
||||
|
||||
@ -5,18 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 gl_Position [[position]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float4 vA [[attribute(0)]];
|
||||
float4 vB [[attribute(1)]];
|
||||
float4 vC [[attribute(2)]];
|
||||
};
|
||||
|
||||
template<typename T>
|
||||
[[clang::optnone]] T spvFMul(T l, T r)
|
||||
{
|
||||
@ -78,6 +66,18 @@ template<typename T>
|
||||
return fma(T(-1), r, l);
|
||||
}
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 gl_Position [[position]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float4 vA [[attribute(0)]];
|
||||
float4 vB [[attribute(1)]];
|
||||
float4 vC [[attribute(2)]];
|
||||
};
|
||||
|
||||
vertex main0_out main0(main0_in in [[stage_in]])
|
||||
{
|
||||
main0_out out = {};
|
||||
|
||||
@ -5,19 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
constant float4 _20[2] = { float4(10.0), float4(20.0) };
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 gl_Position [[position]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float4 vInput0 [[attribute(0)]];
|
||||
float4 vInput1 [[attribute(1)]];
|
||||
};
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
|
||||
{
|
||||
@ -126,6 +113,19 @@ inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device
|
||||
}
|
||||
}
|
||||
|
||||
constant float4 _20[2] = { float4(10.0), float4(20.0) };
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 gl_Position [[position]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
float4 vInput0 [[attribute(0)]];
|
||||
float4 vInput1 [[attribute(1)]];
|
||||
};
|
||||
|
||||
static inline __attribute__((always_inline))
|
||||
void test(thread float4 (&spvReturnValue)[2])
|
||||
{
|
||||
|
||||
@ -5,6 +5,13 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
// Implementation of the GLSL sign() function for integer types
|
||||
template<typename T, typename E = typename enable_if<is_integral<T>::value>::type>
|
||||
inline T sign(T x)
|
||||
{
|
||||
return select(select(select(x, T(0), x == T(0)), T(1), x > T(0)), T(-1), x < T(0));
|
||||
}
|
||||
|
||||
struct UBO
|
||||
{
|
||||
float4x4 uMVP;
|
||||
@ -36,13 +43,6 @@ struct main0_in
|
||||
float4 aVertex [[attribute(0)]];
|
||||
};
|
||||
|
||||
// Implementation of the GLSL sign() function for integer types
|
||||
template<typename T, typename E = typename enable_if<is_integral<T>::value>::type>
|
||||
inline T sign(T x)
|
||||
{
|
||||
return select(select(select(x, T(0), x == T(0)), T(1), x > T(0)), T(-1), x < T(0));
|
||||
}
|
||||
|
||||
vertex main0_out main0(main0_in in [[stage_in]], constant UBO& _21 [[buffer(0)]])
|
||||
{
|
||||
main0_out out = {};
|
||||
|
||||
@ -5,11 +5,6 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 gl_Position [[position]];
|
||||
};
|
||||
|
||||
// Returns 2D texture coords corresponding to 1D texel buffer coords
|
||||
static inline __attribute__((always_inline))
|
||||
uint2 spvTexelBufferCoord(uint tc)
|
||||
@ -17,6 +12,11 @@ uint2 spvTexelBufferCoord(uint tc)
|
||||
return uint2(tc % 4096, tc / 4096);
|
||||
}
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 gl_Position [[position]];
|
||||
};
|
||||
|
||||
vertex main0_out main0(texture2d<float> uSamp [[texture(0)]], texture2d<float> uSampo [[texture(1)]])
|
||||
{
|
||||
main0_out out = {};
|
||||
|
||||
@ -5,6 +5,13 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
// Returns 2D texture coords corresponding to 1D texel buffer coords
|
||||
static inline __attribute__((always_inline))
|
||||
uint2 spvTexelBufferCoord(uint tc)
|
||||
{
|
||||
return uint2(tc % 4096, tc / 4096);
|
||||
}
|
||||
|
||||
struct type_View
|
||||
{
|
||||
float4x4 View_TranslatedWorldToClip;
|
||||
@ -281,13 +288,6 @@ struct main0_in
|
||||
float2 in_var_ATTRIBUTE0 [[attribute(0)]];
|
||||
};
|
||||
|
||||
// Returns 2D texture coords corresponding to 1D texel buffer coords
|
||||
static inline __attribute__((always_inline))
|
||||
uint2 spvTexelBufferCoord(uint tc)
|
||||
{
|
||||
return uint2(tc % 4096, tc / 4096);
|
||||
}
|
||||
|
||||
vertex main0_out main0(main0_in in [[stage_in]], constant type_View& View [[buffer(0)]], constant type_Primitive& Primitive [[buffer(1)]], constant type_MobileShadowDepthPass& MobileShadowDepthPass [[buffer(2)]], constant type_EmitterDynamicUniforms& EmitterDynamicUniforms [[buffer(3)]], constant type_EmitterUniforms& EmitterUniforms [[buffer(4)]], constant type_Globals& _Globals [[buffer(5)]], texture2d<float> ParticleIndices [[texture(0)]], texture2d<float> PositionTexture [[texture(1)]], texture2d<float> VelocityTexture [[texture(2)]], texture2d<float> AttributesTexture [[texture(3)]], texture2d<float> CurveTexture [[texture(4)]], sampler PositionTextureSampler [[sampler(0)]], sampler VelocityTextureSampler [[sampler(1)]], sampler AttributesTextureSampler [[sampler(2)]], sampler CurveTextureSampler [[sampler(3)]], uint gl_VertexIndex [[vertex_id]], uint gl_InstanceIndex [[instance_id]])
|
||||
{
|
||||
main0_out out = {};
|
||||
|
||||
@ -1462,9 +1462,9 @@ string CompilerMSL::compile()
|
||||
|
||||
emit_header();
|
||||
emit_custom_templates();
|
||||
emit_custom_functions();
|
||||
emit_specialization_constants_and_structs();
|
||||
emit_resources();
|
||||
emit_custom_functions();
|
||||
emit_function(get<SPIRFunction>(ir.default_entry_point), Bitset());
|
||||
|
||||
pass_count++;
|
||||
@ -5011,13 +5011,17 @@ void CompilerMSL::emit_custom_functions()
|
||||
case SPVFuncImplQuantizeToF16:
|
||||
// Ensure fast-math is disabled to match Vulkan results.
|
||||
// SpvHalfTypeSelector is used to match the half* template type to the float* template type.
|
||||
// Depending on GPU, MSL does not always flush converted subnormal halfs to zero,
|
||||
// as required by OpQuantizeToF16, so check for subnormals and flush them to zero.
|
||||
statement("template <typename F> struct SpvHalfTypeSelector;");
|
||||
statement("template <> struct SpvHalfTypeSelector<float> { public: using H = half; };");
|
||||
statement("template<uint N> struct SpvHalfTypeSelector<vec<float, N>> { using H = vec<half, N>; };");
|
||||
statement("template<typename F, typename H = typename SpvHalfTypeSelector<F>::H>");
|
||||
statement("[[clang::optnone]] F spvQuantizeToF16(F val)");
|
||||
statement("[[clang::optnone]] F spvQuantizeToF16(F fval)");
|
||||
begin_scope();
|
||||
statement("return F(H(val));");
|
||||
statement("H hval = H(fval);");
|
||||
statement("hval = select(copysign(H(0), hval), hval, isnormal(hval) || isinf(hval) || isnan(hval));");
|
||||
statement("return F(hval);");
|
||||
end_scope();
|
||||
statement("");
|
||||
break;
|
||||
@ -13338,31 +13342,11 @@ string CompilerMSL::type_to_array_glsl(const SPIRType &type)
|
||||
|
||||
string CompilerMSL::constant_op_expression(const SPIRConstantOp &cop)
|
||||
{
|
||||
auto &type = get<SPIRType>(cop.basetype);
|
||||
string op;
|
||||
|
||||
switch (cop.opcode)
|
||||
{
|
||||
case OpQuantizeToF16:
|
||||
switch (type.vecsize)
|
||||
{
|
||||
case 1:
|
||||
op = "float(half(";
|
||||
break;
|
||||
case 2:
|
||||
op = "float2(half2(";
|
||||
break;
|
||||
case 3:
|
||||
op = "float3(half3(";
|
||||
break;
|
||||
case 4:
|
||||
op = "float4(half4(";
|
||||
break;
|
||||
default:
|
||||
SPIRV_CROSS_THROW("Illegal argument to OpSpecConstantOp QuantizeToF16.");
|
||||
}
|
||||
return join(op, to_expression(cop.arguments[0]), "))");
|
||||
|
||||
add_spv_func_and_recompile(SPVFuncImplQuantizeToF16);
|
||||
return join("spvQuantizeToF16(", to_expression(cop.arguments[0]), ")");
|
||||
default:
|
||||
return CompilerGLSL::constant_op_expression(cop);
|
||||
}
|
||||
|
||||
Loading…
Reference in New Issue
Block a user