MSL: SPIR-V assembly fixes

reference/opt/shaders-hlsl/comp/rayquery.nofxc.fxconly.comp

@@ -1,11 +1,11 @@
+static float3x4 _362;
+static float4x3 _364;
+
 RWByteAddressBuffer _17 : register(u0);
 uniform RaytracingAccelerationStructure rtas : register(t1);
 
 static RayQuery<RAY_FLAG_NONE> rayQuery;
 
-static float3x4 _362;
-static float4x3 _364;
-
 void comp_main()
 {
     RayDesc _1ident = {0.0f.xxx, 0.0f, float3(1.0f, 0.0f, 0.0f), 9999.0f};

reference/opt/shaders-hlsl/frag/bvec-operations.frag

@@ -1,3 +1,5 @@
+static bool _47;
+
 static float2 value;
 static float4 FragColor;
 
@@ -11,8 +13,6 @@ struct SPIRV_Cross_Output
     float4 FragColor : SV_Target0;
 };
 
-static bool _47;
-
 void frag_main()
 {
     bool2 _25 = bool2(value.x == 0.0f, _47);

reference/opt/shaders-msl/asm/comp/copy-object-ssbo-to-ssbo.asm.comp

@@ -0,0 +1,23 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct _19
+{
+};
+struct _5
+{
+    int _m0;
+    _19 _m1;
+    char _m2_pad[4];
+    _19 _m2;
+    char _m3_pad[4];
+    int _m3;
+};
+
+kernel void main0(device _5& _3 [[buffer(0)]], device _5& _4 [[buffer(1)]])
+{
+    _4 = _3;
+}
+

reference/opt/shaders-msl/asm/comp/copy-object-ubo-to-ssbo.asm.comp

@@ -0,0 +1,24 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct _19
+{
+};
+struct _5
+{
+    int _m0;
+    char _m1_pad[12];
+    _19 _m1;
+    char _m2_pad[16];
+    _19 _m2;
+    char _m3_pad[16];
+    int _m3;
+};
+
+kernel void main0(constant _5& _3 [[buffer(0)]], device _5& _4 [[buffer(1)]])
+{
+    _4 = _3;
+}
+

reference/opt/shaders-msl/asm/comp/op-spec-constant-op-vector-related.asm.comp

@@ -54,6 +54,7 @@ struct _7
     int _m0[1];
 };
 
+constant int3 _32 = {};
 constant int _3_tmp [[function_constant(0)]];
 constant int _3 = is_function_constant_defined(_3_tmp) ? _3_tmp : 0;
 constant int _4_tmp [[function_constant(1)]];
@@ -102,8 +103,6 @@ constant int _74 = (_73 * _72);
 constant spvUnsafeArray<int, 3> _33 = spvUnsafeArray<int, 3>({ 0, 0, 0 });
 constant spvUnsafeArray<spvUnsafeArray<int, 3>, 3> _34 = spvUnsafeArray<spvUnsafeArray<int, 3>, 3>({ spvUnsafeArray<int, 3>({ 0, 0, 0 }), spvUnsafeArray<int, 3>({ 0, 0, 0 }), spvUnsafeArray<int, 3>({ 0, 0, 0 }) });
 
-constant int3 _32 = {};
-
 kernel void main0(device _7& _8 [[buffer(0)]], device _7& _9 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
     _9._m0[gl_GlobalInvocationID.x] = _8._m0[gl_GlobalInvocationID.x] + ((((1 - _59) * _60) * (_61 - 1)) * _74);

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

@@ -0,0 +1,25 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct _4
+{
+    uint _m0[1];
+};
+
+struct _20
+{
+    uint _m0;
+    uint _m1;
+};
+
+kernel void main0(device _4& _5 [[buffer(0)]], device _4& _6 [[buffer(1)]], device _4& _7 [[buffer(2)]], device _4& _8 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+{
+    _20 _28;
+    _28._m0 = uint(int(_5._m0[gl_GlobalInvocationID.x]) * int(_6._m0[gl_GlobalInvocationID.x]));
+    _28._m1 = uint(mulhi(int(_5._m0[gl_GlobalInvocationID.x]), int(_6._m0[gl_GlobalInvocationID.x])));
+    _7._m0[gl_GlobalInvocationID.x] = _28._m0;
+    _8._m0[gl_GlobalInvocationID.x] = _28._m1;
+}
+

reference/opt/shaders-msl/asm/comp/undefined-constant-composite.asm.comp

@@ -0,0 +1,28 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct _20
+{
+    int _m0;
+    int _m1;
+};
+
+struct _5
+{
+    int _m0[10];
+};
+
+struct _7
+{
+    int _m0[10];
+};
+
+constant int _28 = {};
+
+kernel void main0(device _5& _6 [[buffer(0)]], device _7& _8 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+{
+    _6._m0[gl_GlobalInvocationID.x] = _8._m0[gl_GlobalInvocationID.x] + (_20{ _28, 200 })._m1;
+}
+

reference/opt/shaders-msl/asm/comp/undefined-spec-constant-composite.asm.comp

@@ -0,0 +1,31 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct _21
+{
+    int _m0;
+    int _m1;
+};
+
+struct _5
+{
+    int _m0[10];
+};
+
+struct _7
+{
+    int _m0[10];
+};
+
+constant int _29 = {};
+constant int _9_tmp [[function_constant(0)]];
+constant int _9 = is_function_constant_defined(_9_tmp) ? _9_tmp : 0;
+constant _21 _30 = _21{ _9, _29 };
+
+kernel void main0(device _5& _6 [[buffer(0)]], device _7& _8 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+{
+    _6._m0[gl_GlobalInvocationID.x] = (_8._m0[gl_GlobalInvocationID.x] + _30._m0) + (_21{ _29, 200 })._m1;
+}
+

reference/opt/shaders-ue4/asm/frag/texture-atomics.asm.argument.msl2.frag

@@ -16,6 +16,8 @@ struct type_Globals
     uint2 ShadowTileListGroupSize;
 };
 
+constant float3 _70 = {};
+
 struct spvDescriptorSetBuffer0
 {
     const device type_StructuredBuffer_v4float* CulledObjectBoxBounds [[id(0)]];
@@ -24,8 +26,6 @@ struct spvDescriptorSetBuffer0
     device atomic_uint* RWShadowTileNumCulledObjects_atomic [[id(3)]];
 };
 
-constant float3 _70 = {};
-
 struct main0_out
 {
     float4 out_var_SV_Target0 [[color(0)]];

reference/opt/shaders/asm/comp/undefined-constant-composite.asm.comp

@@ -0,0 +1,26 @@
+#version 450
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+struct _20
+{
+    int _m0;
+    int _m1;
+};
+
+int _28;
+
+layout(binding = 1, std430) buffer _5_6
+{
+    int _m0[10];
+} _6;
+
+layout(binding = 0, std430) buffer _7_8
+{
+    int _m0[10];
+} _8;
+
+void main()
+{
+    _6._m0[gl_GlobalInvocationID.x] = _8._m0[gl_GlobalInvocationID.x] + _20(_28, 200)._m1;
+}
+

reference/opt/shaders/asm/frag/default-member-names.asm.frag

@@ -1,9 +1,9 @@
 #version 450
 
-layout(location = 0) out vec4 _3;
-
 float _49;
 
+layout(location = 0) out vec4 _3;
+
 void main()
 {
     _3 = vec4(_49);

reference/opt/shaders/asm/frag/loop-body-dominator-continue-access.asm.frag

@@ -12,6 +12,8 @@
 #define SPIRV_CROSS_LOOP
 #endif
 
+int _231;
+
 layout(binding = 0, std140) uniform Foo
 {
     layout(row_major) mat4 lightVP[64];
@@ -22,8 +24,6 @@ layout(binding = 0, std140) uniform Foo
 layout(location = 0) in vec3 fragWorld;
 layout(location = 0) out int _entryPointOutput;
 
-int _231;
-
 mat4 spvWorkaroundRowMajor(mat4 wrap) { return wrap; }
 
 void main()

reference/opt/shaders/asm/frag/struct-composite-extract-swizzle.asm.frag

@@ -8,12 +8,12 @@ struct Foo
     float var2;
 };
 
+Foo _22;
+
 layout(binding = 0) uniform mediump sampler2D uSampler;
 
 layout(location = 0) out vec4 FragColor;
 
-Foo _22;
-
 void main()
 {
     FragColor = texture(uSampler, vec2(_22.var1, _22.var2));

reference/opt/shaders/comp/cfg.comp

@@ -1,13 +1,13 @@
 #version 310 es
 layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 
+float _188;
+
 layout(binding = 0, std430) buffer SSBO
 {
     float data;
 } _11;
 
-float _188;
-
 void main()
 {
     if (_11.data != 0.0)

reference/opt/shaders/frag/loop-dominator-and-switch-default.frag

@@ -2,10 +2,10 @@
 precision mediump float;
 precision highp int;
 
-layout(location = 0) out vec4 fragColor;
-
 vec4 _80;
 
+layout(location = 0) out vec4 fragColor;
+
 void main()
 {
     mediump int _18 = int(_80.x);

reference/shaders-hlsl-no-opt/asm/comp/constant-composite-undef.asm.comp

@@ -1,7 +1,7 @@
-RWByteAddressBuffer block : register(u0);
-
 static float _15;
 
+RWByteAddressBuffer block : register(u0);
+
 void comp_main()
 {
     block.Store4(0, asuint(float4(0.100000001490116119384765625f, 0.20000000298023223876953125f, 0.300000011920928955078125f, 0.0f)));

reference/shaders-hlsl-no-opt/asm/frag/composite-insert-inheritance.asm.frag

@@ -1,3 +1,5 @@
+static float4 _32;
+
 static const float4 _34[2] = { 0.0f.xxxx, 0.0f.xxxx };
 
 static float4 vInput;
@@ -13,8 +15,6 @@ struct SPIRV_Cross_Output
     float4 FragColor : SV_Target0;
 };
 
-static float4 _32;
-
 void frag_main()
 {
     float4 _37 = vInput;

reference/shaders-hlsl-no-opt/asm/frag/phi.zero-initialize.asm.frag

@@ -1,8 +1,14 @@
+static int uninit_int = 0;
+static int4 uninit_vector = int4(0, 0, 0, 0);
+static float4x4 uninit_matrix = float4x4(0.0f.xxxx, 0.0f.xxxx, 0.0f.xxxx, 0.0f.xxxx);
+
 struct Foo
 {
     int a;
 };
 
+static Foo uninit_foo = { 0 };
+
 static float4 vColor;
 static float4 FragColor;
 
@@ -16,11 +22,6 @@ struct SPIRV_Cross_Output
     float4 FragColor : SV_Target0;
 };
 
-static int uninit_int = 0;
-static int4 uninit_vector = int4(0, 0, 0, 0);
-static float4x4 uninit_matrix = float4x4(0.0f.xxxx, 0.0f.xxxx, 0.0f.xxxx, 0.0f.xxxx);
-static Foo uninit_foo = { 0 };
-
 void frag_main()
 {
     int _39 = 0;

reference/shaders-hlsl-no-opt/asm/frag/vector-shuffle-undef-index.asm.frag

@@ -1,3 +1,5 @@
+static float4 undef;
+
 static float4 FragColor;
 static float4 vFloat;
 
@@ -11,8 +13,6 @@ struct SPIRV_Cross_Output
     float4 FragColor : SV_Target0;
 };
 
-static float4 undef;
-
 void frag_main()
 {
     FragColor = float4(undef.x, vFloat.y, 0.0f, vFloat.w) + float4(vFloat.z, vFloat.y, 0.0f, vFloat.w);

reference/shaders-hlsl/asm/frag/unreachable.asm.frag

@@ -1,3 +1,5 @@
+static float4 _21;
+
 static int counter;
 static float4 FragColor;
 
@@ -11,8 +13,6 @@ struct SPIRV_Cross_Output
     float4 FragColor : SV_Target0;
 };
 
-static float4 _21;
-
 void frag_main()
 {
     float4 _24;

reference/shaders-msl-no-opt/asm/frag/composite-insert-inheritance.asm.frag

@@ -44,10 +44,10 @@ struct spvUnsafeArray
     }
 };
 
-constant spvUnsafeArray<float4, 2> _34 = spvUnsafeArray<float4, 2>({ float4(0.0), float4(0.0) });
-
 constant float4 _32 = {};
 
+constant spvUnsafeArray<float4, 2> _34 = spvUnsafeArray<float4, 2>({ float4(0.0), float4(0.0) });
+
 struct main0_out
 {
     float4 FragColor [[color(0)]];