MSL: Support synthetic functions in function constants.

Emit synthetic functions before function constants.
Support use of spvQuantizeToF16() in function constants for numerical
behavior consistency with the op code.
Ensure subnormal results from OpQuantizeToF16 are flushed to zero per SPIR-V spec.

Adjust SPIRV-Cross unit test reference shaders to accommodate these changes.
Any MSL reference shader that inclues a synthetic function is affected,
since the location it is emitted has changed.

reference/opt/shaders-msl/asm/comp/atomic-decrement.asm.comp

@@ -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);

reference/opt/shaders-msl/asm/comp/atomic-increment.asm.comp

@@ -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);

reference/opt/shaders-msl/asm/comp/buffer-write-relative-addr.asm.comp

@@ -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;

reference/opt/shaders-msl/asm/comp/buffer-write.asm.comp

@@ -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)));

reference/opt/shaders-msl/asm/comp/quantize.asm.comp

@@ -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);

reference/opt/shaders-msl/asm/frag/single-function-private-lut.asm.frag

@@ -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 } });

reference/opt/shaders-msl/comp/composite-array-initialization.force-native-array.comp

@@ -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 } };

reference/opt/shaders-msl/comp/global-invocation-id-writable-ssbo-in-function.comp

@@ -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;

reference/opt/shaders-msl/comp/global-invocation-id.comp

@@ -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;

reference/opt/shaders-msl/comp/image-atomic-automatic-bindings.argument.msl2.comp

@@ -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);

reference/opt/shaders-msl/comp/image-atomic-automatic-bindings.comp

@@ -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);

reference/opt/shaders-msl/comp/inverse.comp

@@ -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);

reference/opt/shaders-msl/comp/local-invocation-id.comp

@@ -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;

reference/opt/shaders-msl/comp/local-invocation-index.comp

@@ -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;

reference/opt/shaders-msl/comp/mod.comp

@@ -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]);

reference/opt/shaders-msl/comp/writable-ssbo.comp

@@ -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;

reference/opt/shaders-msl/frag/array-of-texture-swizzle-nonconstant-uniform.msl2.argument.discrete.swizzle.frag

@@ -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 = {};

reference/opt/shaders-msl/frag/array-of-texture-swizzle-nonconstant-uniform.msl2.swizzle.frag

@@ -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 = {};

reference/opt/shaders-msl/frag/array-of-texture-swizzle.msl2.argument.discrete.swizzle.frag

@@ -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 = {};

reference/opt/shaders-msl/frag/array-of-texture-swizzle.msl2.swizzle.frag

@@ -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 = {};

reference/opt/shaders-msl/frag/buffer-read-write.frag

@@ -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 = {};

reference/opt/shaders-msl/frag/mrt-array.frag

@@ -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 = {};

reference/opt/shaders-msl/frag/pixel-interlock-ordered.msl2.argument.frag

@@ -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;

reference/opt/shaders-msl/frag/pixel-interlock-ordered.msl2.frag

@@ -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;

reference/opt/shaders-msl/frag/scalar-refract-reflect.frag

@@ -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 = {};

reference/opt/shaders-msl/frag/subgroup-globals-extract.msl22.frag

@@ -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 = {};

reference/opt/shaders-msl/frag/texture-cube-array.ios.emulate-cube-array.frag

@@ -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 = {};

reference/opt/shaders-msl/masking/copy-arrays.mask-location-0.msl2.multi-patch.tesc

@@ -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];

reference/opt/shaders-msl/masking/copy-arrays.mask-location-0.msl2.tesc

@@ -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];

reference/opt/shaders-msl/vert/float-math.invariant-float-math.vert

@@ -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 = {};

reference/opt/shaders-msl/vert/functions.vert

@@ -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 = {};

reference/opt/shaders-msl/vert/no-contraction.vert

@@ -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 = {};

reference/opt/shaders-msl/vert/sign-int-types.vert

@@ -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 = {};

reference/opt/shaders-msl/vert/texture_buffer.vert

@@ -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 = {};

reference/opt/shaders-ue4/asm/vert/texture-buffer.asm.vert

@@ -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 = {};

reference/shaders-msl-no-opt/asm/comp/bitscan.asm.comp

@@ -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;

reference/shaders-msl-no-opt/asm/comp/device-array-load-temporary.asm.comp

@@ -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];

reference/shaders-msl-no-opt/asm/comp/device-array-load-temporary.force-native-array.asm.comp

@@ -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];

reference/shaders-msl-no-opt/asm/comp/device-constant-array-load-store.asm.comp

@@ -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];

reference/shaders-msl-no-opt/asm/comp/device-constant-array-load-store.force-native-array.asm.comp

@@ -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];

reference/shaders-msl-no-opt/asm/comp/glsl-signed-operations.asm.comp

@@ -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;

reference/shaders-msl-no-opt/asm/frag/copy-memory-block-like-thread-local.asm.frag

@@ -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;

reference/shaders-msl-no-opt/comp/array-copy-threadgroup-memory.comp

@@ -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];

reference/shaders-msl-no-opt/comp/glsl.std450.comp

@@ -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]);

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

@@ -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));

reference/shaders-msl-no-opt/comp/subgroups.nocompat.invalid.vk.msl21.fixed-subgroup.comp

@@ -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;

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

@@ -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);

reference/shaders-msl-no-opt/comp/subgroups.nocompat.invalid.vk.msl22.ios.comp

@@ -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));

reference/shaders-msl-no-opt/comp/subgroups.nocompat.invalid.vk.msl23.ios.simd.comp

@@ -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));

reference/shaders-msl-no-opt/frag/fp16.desktop.invalid.frag

@@ -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)
 {

reference/shaders-msl-no-opt/frag/subgroups.nocompat.invalid.vk.msl22.frag

@@ -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 = {};

reference/shaders-msl-no-opt/vert/functions_nested.vert

@@ -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)
 {

reference/shaders-msl-no-opt/vert/pass-array-by-value.force-native-array.vert

@@ -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)
 {

reference/shaders-msl-no-opt/vulkan/frag/texture-access-function.swizzle.vk.frag

@@ -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)
 {

reference/shaders-msl/asm/comp/atomic-decrement.asm.comp

@@ -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);

reference/shaders-msl/asm/comp/atomic-increment.asm.comp

@@ -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);

reference/shaders-msl/asm/comp/buffer-write-relative-addr.asm.comp

@@ -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;

reference/shaders-msl/asm/comp/buffer-write.asm.comp

@@ -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)));

reference/shaders-msl/asm/comp/quantize.asm.comp

@@ -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);

reference/shaders-msl/asm/frag/single-function-private-lut.asm.frag

@@ -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 } });

reference/shaders-msl/asm/vert/spec-constant-op-composite.asm.vert

@@ -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
 {

reference/shaders-msl/comp/composite-array-initialization.force-native-array.comp

@@ -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)
 {

reference/shaders-msl/comp/copy-array-of-arrays.force-native-array.comp

@@ -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];

reference/shaders-msl/comp/global-invocation-id-writable-ssbo-in-function.comp

@@ -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)
 {

reference/shaders-msl/comp/global-invocation-id.comp

@@ -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;

reference/shaders-msl/comp/image-atomic-automatic-bindings.argument.msl2.comp

@@ -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);

reference/shaders-msl/comp/image-atomic-automatic-bindings.comp

@@ -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);

reference/shaders-msl/comp/inverse.comp

@@ -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);

reference/shaders-msl/comp/local-invocation-id.comp

@@ -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;

reference/shaders-msl/comp/local-invocation-index.comp

@@ -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;

reference/shaders-msl/comp/mod.comp

@@ -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;

reference/shaders-msl/comp/writable-ssbo.comp

@@ -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;

reference/shaders-msl/frag/array-of-texture-swizzle-nonconstant-uniform.msl2.argument.discrete.swizzle.frag

@@ -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)
 {

reference/shaders-msl/frag/array-of-texture-swizzle-nonconstant-uniform.msl2.swizzle.frag

@@ -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)
 {

reference/shaders-msl/frag/array-of-texture-swizzle.msl2.argument.discrete.swizzle.frag

@@ -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)
 {

reference/shaders-msl/frag/array-of-texture-swizzle.msl2.swizzle.frag

@@ -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)
 {

reference/shaders-msl/frag/buffer-read-write.frag

@@ -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 = {};

reference/shaders-msl/frag/mrt-array.frag

@@ -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)
 {

reference/shaders-msl/frag/pixel-interlock-ordered.msl2.argument.frag

@@ -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;

reference/shaders-msl/frag/pixel-interlock-ordered.msl2.frag

@@ -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;

reference/shaders-msl/frag/scalar-refract-reflect.frag

@@ -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 = {};

reference/shaders-msl/frag/subgroup-globals-extract.msl22.frag

@@ -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)
 {

reference/shaders-msl/frag/texture-cube-array.ios.emulate-cube-array.frag

@@ -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 = {};

reference/shaders-msl/masking/copy-arrays.mask-location-0.msl2.multi-patch.tesc

@@ -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];

reference/shaders-msl/masking/copy-arrays.mask-location-0.msl2.tesc

@@ -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];

reference/shaders-msl/vert/float-math.invariant-float-math.vert

@@ -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 = {};

reference/shaders-msl/vert/functions.vert

@@ -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 = {};

reference/shaders-msl/vert/no-contraction.vert

@@ -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 = {};

reference/shaders-msl/vert/return-array.force-native-array.vert

@@ -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])
 {

reference/shaders-msl/vert/sign-int-types.vert

@@ -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 = {};

reference/shaders-msl/vert/texture_buffer.vert

@@ -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 = {};

reference/shaders-ue4/asm/vert/texture-buffer.asm.vert

@@ -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 = {};

spirv_msl.cpp

@@ -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++;
@@ -5004,13 +5004,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;
@@ -13331,31 +13335,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)
+		add_spv_func_and_recompile(SPVFuncImplQuantizeToF16);
-		{
+		return join("spvQuantizeToF16(", to_expression(cop.arguments[0]), ")");
-		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]), "))");
-
 	default:
 		return CompilerGLSL::constant_op_expression(cop);
 	}