Update SPIR-V Dependencies to match Dawn requirements

third_party/spirv-cross/.github/workflows/main.yml

@@ -26,15 +26,24 @@ jobs:
         with:
           python-version: '3.x'
 
+      - name: Cache glslang / SPIRV-Tools
+        id: cache-externals
+        uses: actions/cache@v3
+        with:
+          path: external/*-build/output
+          key: ${{ matrix.platform }} externals ${{ hashFiles('checkout_glslang_spirv_tools.sh', 'build_glslang_spirv_tools.sh') }}
+
       - name: Pull glslang / SPIRV-Tools
+        if: steps.cache-externals.outputs.cache-hit != 'true'
         shell: bash
         working-directory: ${{github.workspace}}
         run: ./checkout_glslang_spirv_tools.sh
 
       - name: Build glslang / SPIRV-Tools
+        if: steps.cache-externals.outputs.cache-hit != 'true'
         shell: bash
         working-directory: ${{github.workspace}}
-        run: ./build_glslang_spirv_tools.sh Release
+        run: ./build_glslang_spirv_tools.sh Release 4
 
       - name: Configure SPIRV-Cross
         shell: bash
@@ -47,7 +56,7 @@ jobs:
         shell: bash
         working-directory: ${{github.workspace}}/build
         run: |
-          cmake --build . --config Release
+          cmake --build . --config Release --parallel 4
           cmake --build . --config Release --target install
 
       - uses: actions/upload-artifact@v3

third_party/spirv-cross/.gitignore

@@ -4,6 +4,7 @@
 *.o
 *.d
 *.txt
+/build
 /test
 /spirv-cross
 /obj
@@ -20,5 +21,8 @@
 .vs/
 *.vcxproj.user
 
+# Mac OS X Finder
+.DS_Store
+
 !CMakeLists.txt
 !LICENSES/*.txt

third_party/spirv-cross/CMakeLists.txt

@@ -19,7 +19,7 @@
 #  2. The MIT License, found at <http://opensource.org/licenses/MIT>.
 #
 
-cmake_minimum_required(VERSION 3.0)
+cmake_minimum_required(VERSION 3.6)
 set(CMAKE_CXX_STANDARD 11)
 set(CMAKE_CXX_EXTENSIONS OFF)
 
@@ -77,7 +77,7 @@ set(spirv-cross-build-version "unknown")
 find_package(Git)
 if (GIT_FOUND)
 	execute_process(
-		COMMAND ${GIT_EXECUTABLE} rev-parse --short HEAD
+		COMMAND ${GIT_EXECUTABLE} describe --always --tags --dirty=+
 		WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}
 		OUTPUT_VARIABLE spirv-cross-build-version
 		ERROR_QUIET
@@ -243,7 +243,7 @@ set(spirv-cross-util-sources
 		${CMAKE_CURRENT_SOURCE_DIR}/spirv_cross_util.hpp)
 
 set(spirv-cross-abi-major 0)
-set(spirv-cross-abi-minor 57)
+set(spirv-cross-abi-minor 59)
 set(spirv-cross-abi-patch 0)
 set(SPIRV_CROSS_VERSION ${spirv-cross-abi-major}.${spirv-cross-abi-minor}.${spirv-cross-abi-patch})
 
@@ -465,7 +465,19 @@ if (SPIRV_CROSS_CLI)
 		#  - Update the reference files
 		#  - Get cycle counts from malisc
 		#  - Keep failing outputs
-		find_package(PythonInterp)
+		if (${CMAKE_VERSION} VERSION_GREATER "3.12")
+			find_package(Python3)
+			if (${PYTHON3_FOUND})
+				set(PYTHONINTERP_FOUND ON)
+				set(PYTHON_VERSION_MAJOR 3)
+				set(PYTHON_EXECUTABLE ${Python3_EXECUTABLE})
+			else()
+				set(PYTHONINTERP_FOUND OFF)
+			endif()
+		else()
+			find_package(PythonInterp)
+		endif()
+
 		find_program(spirv-cross-glslang NAMES glslangValidator
 				PATHS ${CMAKE_CURRENT_SOURCE_DIR}/external/glslang-build/output/bin
 				NO_DEFAULT_PATH)

third_party/spirv-cross/build_glslang_spirv_tools.sh

@@ -15,7 +15,7 @@ fi
 echo "Building glslang."
 mkdir -p external/glslang-build
 cd external/glslang-build
-cmake ../glslang -DCMAKE_BUILD_TYPE=$PROFILE -DCMAKE_INSTALL_PREFIX=output
+cmake ../glslang -DCMAKE_BUILD_TYPE=$PROFILE -DCMAKE_INSTALL_PREFIX=output -DENABLE_OPT=OFF
 cmake --build . --config $PROFILE --target install ${NPROC}
 cd ../..
 

third_party/spirv-cross/checkout_glslang_spirv_tools.sh

@@ -2,9 +2,9 @@
 # Copyright 2016-2021 The Khronos Group Inc.
 # SPDX-License-Identifier: Apache-2.0
 
-GLSLANG_REV=06a7078ce74ab5c7801a165b8145859678831fb8
-SPIRV_TOOLS_REV=f62e121b0df5374d1f043d1fbda98467406af0b1
-SPIRV_HEADERS_REV=d13b52222c39a7e9a401b44646f0ca3a640fbd47
+GLSLANG_REV=a7785ea1ff5b10bfc2d8ca77fdad5929562897b7
+SPIRV_TOOLS_REV=afaf8fda2ad0364655909b56c8b634ce89095bb5
+SPIRV_HEADERS_REV=e867c06631767a2d96424cbec530f9ee5e78180f
 PROTOCOL=https
 
 if [ -d external/glslang ]; then

third_party/spirv-cross/main.cpp

@@ -676,6 +676,10 @@ struct CLIArguments
 	bool msl_manual_helper_invocation_updates = true;
 	bool msl_check_discarded_frag_stores = false;
 	bool msl_sample_dref_lod_array_as_grad = false;
+	bool msl_runtime_array_rich_descriptor = false;
+	bool msl_replace_recursive_inputs = false;
+	bool msl_readwrite_texture_fences = true;
+	bool msl_agx_manual_cube_grad_fixup = false;
 	const char *msl_combined_sampler_suffix = nullptr;
 	bool glsl_emit_push_constant_as_ubo = false;
 	bool glsl_emit_ubo_as_plain_uniforms = false;
@@ -864,7 +868,9 @@ static void print_help_msl()
 	                "\t\tRequires MSL 2.0 to be enabled.\n"
 	                "\t[--msl-argument-buffer-tier]:\n\t\tWhen using Metal argument buffers, indicate the Metal argument buffer tier level supported by the Metal platform.\n"
 	                "\t\tUses same values as Metal MTLArgumentBuffersTier enumeration (0 = Tier1, 1 = Tier2).\n"
-	                "\t\tSetting this value also enables msl-argument-buffers.\n"
+	                "\t\tNOTE: Setting this value no longer enables msl-argument-buffers implicitly.\n"
+	                "\t[--msl-runtime-array-rich-descriptor]:\n\t\tWhen declaring a runtime array of SSBOs, declare an array of {ptr, len} pairs to support OpArrayLength.\n"
+	                "\t[--msl-replace-recursive-inputs]:\n\t\tWorks around a Metal 3.1 regression bug, which causes an infinite recursion crash during Metal's analysis of an entry point input structure that itself contains internal recursion.\n"
 	                "\t[--msl-texture-buffer-native]:\n\t\tEnable native support for texel buffers. Otherwise, it is emulated as a normal texture.\n"
 	                "\t[--msl-framebuffer-fetch]:\n\t\tImplement subpass inputs with frame buffer fetch.\n"
 	                "\t\tEmits [[color(N)]] inputs in fragment stage.\n"
@@ -954,6 +960,14 @@ static void print_help_msl()
 	                "\t\tSome Metal devices have a bug where the level() argument to\n"
 	                "\t\tdepth2d_array<T>::sample_compare() in a fragment shader is biased by some\n"
 	                "\t\tunknown amount. This prevents the bias from being added.\n"
+	                "\t[--msl-no-readwrite-texture-fences]:\n\t\tDo not insert fences before each read of a\n"
+	                "\t\tread_write texture. MSL does not guarantee coherence between writes and later reads\n"
+	                "\t\tof read_write textures. If you don't rely on this, you can disable this for a\n"
+	                "\t\tpossible performance improvement.\n"
+	                "\t[--msl-agx-manual-cube-grad-fixup]:\n\t\tManually transform cube texture gradients.\n"
+	                "\t\tAll released Apple Silicon GPUs to date ignore one of the three partial derivatives\n"
+	                "\t\tbased on the selected major axis, and expect the remaining derivatives to be\n"
+	                "\t\tpartially transformed. This fixup gives correct results on Apple Silicon.\n"
 	                "\t[--msl-combined-sampler-suffix <suffix>]:\n\t\tUses a custom suffix for combined samplers.\n");
 	// clang-format on
 }
@@ -1230,6 +1244,10 @@ static string compile_iteration(const CLIArguments &args, std::vector<uint32_t>
 		msl_opts.check_discarded_frag_stores = args.msl_check_discarded_frag_stores;
 		msl_opts.sample_dref_lod_array_as_grad = args.msl_sample_dref_lod_array_as_grad;
 		msl_opts.ios_support_base_vertex_instance = true;
+		msl_opts.runtime_array_rich_descriptor = args.msl_runtime_array_rich_descriptor;
+		msl_opts.replace_recursive_inputs = args.msl_replace_recursive_inputs;
+		msl_opts.readwrite_texture_fences = args.msl_readwrite_texture_fences;
+		msl_opts.agx_manual_cube_grad_fixup = args.msl_agx_manual_cube_grad_fixup;
 		msl_comp->set_msl_options(msl_opts);
 		for (auto &v : args.msl_discrete_descriptor_sets)
 			msl_comp->add_discrete_descriptor_set(v);
@@ -1636,10 +1654,8 @@ static int main_inner(int argc, char *argv[])
 	cbs.add("--msl-pad-fragment-output", [&args](CLIParser &) { args.msl_pad_fragment_output = true; });
 	cbs.add("--msl-domain-lower-left", [&args](CLIParser &) { args.msl_domain_lower_left = true; });
 	cbs.add("--msl-argument-buffers", [&args](CLIParser &) { args.msl_argument_buffers = true; });
-	cbs.add("--msl-argument-buffer-tier", [&args](CLIParser &parser) {
-		args.msl_argument_buffers_tier = parser.next_uint();
-		args.msl_argument_buffers = true;
-	});
+	cbs.add("--msl-argument-buffer-tier",
+	        [&args](CLIParser &parser) { args.msl_argument_buffers_tier = parser.next_uint(); });
 	cbs.add("--msl-discrete-descriptor-set",
 	        [&args](CLIParser &parser) { args.msl_discrete_descriptor_sets.push_back(parser.next_uint()); });
 	cbs.add("--msl-device-argument-buffer",
@@ -1786,9 +1802,15 @@ static int main_inner(int argc, char *argv[])
 	cbs.add("--msl-check-discarded-frag-stores", [&args](CLIParser &) { args.msl_check_discarded_frag_stores = true; });
 	cbs.add("--msl-sample-dref-lod-array-as-grad",
 	        [&args](CLIParser &) { args.msl_sample_dref_lod_array_as_grad = true; });
+	cbs.add("--msl-no-readwrite-texture-fences", [&args](CLIParser &) { args.msl_readwrite_texture_fences = false; });
+	cbs.add("--msl-agx-manual-cube-grad-fixup", [&args](CLIParser &) { args.msl_agx_manual_cube_grad_fixup = true; });
 	cbs.add("--msl-combined-sampler-suffix", [&args](CLIParser &parser) {
 		args.msl_combined_sampler_suffix = parser.next_string();
 	});
+	cbs.add("--msl-runtime-array-rich-descriptor",
+	        [&args](CLIParser &) { args.msl_runtime_array_rich_descriptor = true; });
+	cbs.add("--msl-replace-recursive-inputs",
+	        [&args](CLIParser &) { args.msl_replace_recursive_inputs = true; });
 	cbs.add("--extension", [&args](CLIParser &parser) { args.extensions.push_back(parser.next_string()); });
 	cbs.add("--rename-entry-point", [&args](CLIParser &parser) {
 		auto old_name = parser.next_string();

third_party/spirv-cross/reference/opt/shaders-hlsl/asm/comp/access-chain-invalidate.asm.comp

@@ -1,15 +1,15 @@
-RWByteAddressBuffer _4 : register(u0);
+RWByteAddressBuffer _6 : register(u0);
 
 void comp_main()
 {
-    uint _21 = _4.Load(_4.Load(0) * 4 + 4);
-    for (uint _23 = 0u; _23 < 64u; )
+    uint _21 = _6.Load(_6.Load(0) * 4 + 4);
+    for (uint _40 = 0u; _40 < 64u; )
     {
-        _4.Store(_23 * 4 + 4, 0u);
-        _23++;
+        _6.Store(_40 * 4 + 4, 0u);
+        _40++;
         continue;
     }
-    _4.Store(_4.Load(0) * 4 + 4, _21);
+    _6.Store(_6.Load(0) * 4 + 4, _21);
 }
 
 [numthreads(1, 1, 1)]

third_party/spirv-cross/reference/opt/shaders-hlsl/asm/comp/atomic-decrement.asm.comp

@@ -9,9 +9,9 @@ struct SPIRV_Cross_Input
 
 void comp_main()
 {
-    uint _29;
-    u0_counter.InterlockedAdd(0, -1, _29);
-    u0[asint(asfloat(_29))] = uint(int(gl_GlobalInvocationID.x)).x;
+    uint _24;
+    u0_counter.InterlockedAdd(0, -1, _24);
+    u0[asint(asfloat(_24))] = uint(int(gl_GlobalInvocationID.x)).x;
 }
 
 [numthreads(4, 1, 1)]

third_party/spirv-cross/reference/opt/shaders-hlsl/asm/comp/atomic-increment.asm.comp

@@ -9,9 +9,9 @@ struct SPIRV_Cross_Input
 
 void comp_main()
 {
-    uint _29;
-    u0_counter.InterlockedAdd(0, 1, _29);
-    u0[asint(asfloat(_29))] = uint(int(gl_GlobalInvocationID.x)).x;
+    uint _24;
+    u0_counter.InterlockedAdd(0, 1, _24);
+    u0[asint(asfloat(_24))] = uint(int(gl_GlobalInvocationID.x)).x;
 }
 
 [numthreads(4, 1, 1)]

third_party/spirv-cross/reference/opt/shaders-hlsl/asm/frag/cbuffer-stripped.asm.frag

@@ -1,26 +1,26 @@
-cbuffer _4_5 : register(b0)
+cbuffer _6_7 : register(b0)
 {
-    column_major float2x4 _5_m0 : packoffset(c0);
-    float4 _5_m1 : packoffset(c4);
+    column_major float2x4 _7_m0 : packoffset(c0);
+    float4 _7_m1 : packoffset(c4);
 };
 
 
-static float2 _3;
+static float2 _4;
 
 struct SPIRV_Cross_Output
 {
-    float2 _3 : SV_Target0;
+    float2 _4 : SV_Target0;
 };
 
 void frag_main()
 {
-    _3 = mul(_5_m0, _5_m1);
+    _4 = mul(_7_m0, _7_m1);
 }
 
 SPIRV_Cross_Output main()
 {
     frag_main();
     SPIRV_Cross_Output stage_output;
-    stage_output._3 = _3;
+    stage_output._4 = _4;
     return stage_output;
 }

third_party/spirv-cross/reference/opt/shaders-hlsl/asm/frag/image-extract-reuse.asm.frag

@@ -17,9 +17,9 @@ uint2 spvTextureSize(Texture2D<float4> Tex, uint Level, out uint Param)
 
 void frag_main()
 {
+    uint _18_dummy_parameter;
     uint _19_dummy_parameter;
-    uint _20_dummy_parameter;
-    Size = int2(spvTextureSize(uTexture, uint(0), _19_dummy_parameter)) + int2(spvTextureSize(uTexture, uint(1), _20_dummy_parameter));
+    Size = int2(spvTextureSize(uTexture, uint(0), _18_dummy_parameter)) + int2(spvTextureSize(uTexture, uint(1), _19_dummy_parameter));
 }
 
 SPIRV_Cross_Output main()

third_party/spirv-cross/reference/opt/shaders-hlsl/asm/frag/implicit-read-dep-phi.asm.frag

@@ -17,22 +17,22 @@ struct SPIRV_Cross_Output
 void frag_main()
 {
     float phi;
-    float4 _36;
-    int _51;
-    _51 = 0;
+    float4 _45;
+    int _57;
+    _57 = 0;
     phi = 1.0f;
-    _36 = float4(1.0f, 2.0f, 1.0f, 2.0f);
+    _45 = float4(1.0f, 2.0f, 1.0f, 2.0f);
     for (;;)
     {
-        FragColor = _36;
-        if (_51 < 4)
+        FragColor = _45;
+        if (_57 < 4)
         {
-            if (v0[_51] > 0.0f)
+            if (v0[_57] > 0.0f)
             {
-                float2 _48 = phi.xx;
-                _51++;
+                float2 _43 = phi.xx;
+                _57++;
                 phi += 2.0f;
-                _36 = uImage.SampleLevel(_uImage_sampler, _48, 0.0f);
+                _45 = uImage.SampleLevel(_uImage_sampler, _43, 0.0f);
                 continue;
             }
             else

third_party/spirv-cross/reference/opt/shaders-hlsl/asm/frag/line-directive.line.asm.frag

@@ -30,14 +30,14 @@ void frag_main()
         FragColor = 4.0f;
     }
 #line 19 "test.frag"
-    for (int _127 = 0; float(_127) < (40.0f + vColor); )
+    for (int _131 = 0; float(_131) < (40.0f + vColor); )
     {
 #line 21 "test.frag"
         FragColor += 0.20000000298023223876953125f;
 #line 22 "test.frag"
         FragColor += 0.300000011920928955078125f;
 #line 19 "test.frag"
-        _127 += (int(vColor) + 5);
+        _131 += (int(vColor) + 5);
         continue;
     }
 #line 25 "test.frag"

third_party/spirv-cross/reference/opt/shaders-hlsl/asm/frag/lut-promotion-initializer.asm.frag

@@ -1,6 +1,6 @@
-static const float _46[16] = { 1.0f, 2.0f, 3.0f, 4.0f, 1.0f, 2.0f, 3.0f, 4.0f, 1.0f, 2.0f, 3.0f, 4.0f, 1.0f, 2.0f, 3.0f, 4.0f };
-static const float4 _76[4] = { 0.0f.xxxx, 1.0f.xxxx, 8.0f.xxxx, 5.0f.xxxx };
-static const float4 _90[4] = { 20.0f.xxxx, 30.0f.xxxx, 50.0f.xxxx, 60.0f.xxxx };
+static const float _16[16] = { 1.0f, 2.0f, 3.0f, 4.0f, 1.0f, 2.0f, 3.0f, 4.0f, 1.0f, 2.0f, 3.0f, 4.0f, 1.0f, 2.0f, 3.0f, 4.0f };
+static const float4 _60[4] = { 0.0f.xxxx, 1.0f.xxxx, 8.0f.xxxx, 5.0f.xxxx };
+static const float4 _104[4] = { 20.0f.xxxx, 30.0f.xxxx, 50.0f.xxxx, 60.0f.xxxx };
 
 static float FragColor;
 static int index;
@@ -17,34 +17,34 @@ struct SPIRV_Cross_Output
 
 void frag_main()
 {
-    float4 foobar[4] = _76;
-    float4 baz[4] = _76;
-    FragColor = _46[index];
+    float4 foobar[4] = _60;
+    float4 baz[4] = _60;
+    FragColor = _16[index];
     if (index < 10)
     {
-        FragColor += _46[index ^ 1];
+        FragColor += _16[index ^ 1];
     }
     else
     {
-        FragColor += _46[index & 1];
+        FragColor += _16[index & 1];
     }
-    bool _99 = index > 30;
-    if (_99)
+    bool _63 = index > 30;
+    if (_63)
     {
-        FragColor += _76[index & 3].y;
+        FragColor += _60[index & 3].y;
     }
     else
     {
-        FragColor += _76[index & 1].x;
+        FragColor += _60[index & 1].x;
     }
-    if (_99)
+    if (_63)
     {
         foobar[1].z = 20.0f;
     }
-    int _37 = index & 3;
-    FragColor += foobar[_37].z;
-    baz = _90;
-    FragColor += baz[_37].z;
+    int _91 = index & 3;
+    FragColor += foobar[_91].z;
+    baz = _104;
+    FragColor += baz[_91].z;
 }
 
 SPIRV_Cross_Output main(SPIRV_Cross_Input stage_input)

third_party/spirv-cross/reference/opt/shaders-hlsl/asm/frag/unreachable.asm.frag

@@ -13,21 +13,21 @@ struct SPIRV_Cross_Output
 
 void frag_main()
 {
-    float4 _33;
+    float4 _46;
     for (;;)
     {
         if (counter == 10)
         {
-            _33 = 10.0f.xxxx;
+            _46 = 10.0f.xxxx;
             break;
         }
         else
         {
-            _33 = 30.0f.xxxx;
+            _46 = 30.0f.xxxx;
             break;
         }
     }
-    FragColor = _33;
+    FragColor = _46;
 }
 
 SPIRV_Cross_Output main(SPIRV_Cross_Input stage_input)

third_party/spirv-cross/reference/opt/shaders-hlsl/asm/vert/spec-constant-op-composite.asm.vert

@@ -1,35 +1,35 @@
 #ifndef SPIRV_CROSS_CONSTANT_ID_201
 #define SPIRV_CROSS_CONSTANT_ID_201 -10
 #endif
-static const int _7 = SPIRV_CROSS_CONSTANT_ID_201;
-static const int _20 = (_7 + 2);
+static const int _13 = SPIRV_CROSS_CONSTANT_ID_201;
+static const int _15 = (_13 + 2);
 #ifndef SPIRV_CROSS_CONSTANT_ID_202
 #define SPIRV_CROSS_CONSTANT_ID_202 100u
 #endif
-static const uint _8 = SPIRV_CROSS_CONSTANT_ID_202;
-static const uint _25 = (_8 % 5u);
-static const int4 _30 = int4(20, 30, _20, _20);
-static const int2 _32 = int2(_30.y, _30.x);
-static const int _33 = _30.y;
+static const uint _24 = SPIRV_CROSS_CONSTANT_ID_202;
+static const uint _26 = (_24 % 5u);
+static const int4 _36 = int4(20, 30, _15, _15);
+static const int2 _41 = int2(_36.y, _36.x);
+static const int _60 = _36.y;
 
 static float4 gl_Position;
-static int _4;
+static int _58;
 
 struct SPIRV_Cross_Output
 {
-    nointerpolation int _4 : TEXCOORD0;
+    nointerpolation int _58 : TEXCOORD0;
     float4 gl_Position : SV_Position;
 };
 
 void vert_main()
 {
     float4 _63 = 0.0f.xxxx;
-    _63.y = float(_20);
-    _63.z = float(_25);
-    float4 _52 = _63 + float4(_30);
-    float2 _56 = _52.xy + float2(_32);
-    gl_Position = float4(_56.x, _56.y, _52.z, _52.w);
-    _4 = _33;
+    _63.y = float(_15);
+    _63.z = float(_26);
+    float4 _39 = _63 + float4(_36);
+    float2 _46 = _39.xy + float2(_41);
+    gl_Position = float4(_46.x, _46.y, _39.z, _39.w);
+    _58 = _60;
 }
 
 SPIRV_Cross_Output main()
@@ -37,6 +37,6 @@ SPIRV_Cross_Output main()
     vert_main();
     SPIRV_Cross_Output stage_output;
     stage_output.gl_Position = gl_Position;
-    stage_output._4 = _4;
+    stage_output._58 = _58;
     return stage_output;
 }

third_party/spirv-cross/reference/opt/shaders-hlsl/comp/ssbo-store-array.comp

@@ -0,0 +1,24 @@
+struct Data
+{
+    uint arr[3];
+};
+
+static uint _33;
+
+RWByteAddressBuffer _13 : register(u0);
+
+void comp_main()
+{
+    uint _32[3] = { _33, _33, _33 };
+    [unroll]
+    for (int _0ident = 0; _0ident < 3; _0ident++)
+    {
+        _13.Store(_0ident * 4 + 0, _32[_0ident]);
+    }
+}
+
+[numthreads(1, 1, 1)]
+void main()
+{
+    comp_main();
+}

third_party/spirv-cross/reference/opt/shaders-hlsl/frag/sample-mask-in-and-out.frag

@@ -1,5 +1,5 @@
-static int gl_SampleMaskIn;
-static int gl_SampleMask;
+static int gl_SampleMaskIn[1];
+static int gl_SampleMask[1];
 static float4 FragColor;
 
 struct SPIRV_Cross_Input
@@ -16,15 +16,15 @@ struct SPIRV_Cross_Output
 void frag_main()
 {
     FragColor = 1.0f.xxxx;
-    gl_SampleMask = gl_SampleMaskIn;
+    gl_SampleMask[0] = gl_SampleMaskIn[0];
 }
 
 SPIRV_Cross_Output main(SPIRV_Cross_Input stage_input)
 {
-    gl_SampleMaskIn = stage_input.gl_SampleMaskIn;
+    gl_SampleMaskIn[0] = stage_input.gl_SampleMaskIn;
     frag_main();
     SPIRV_Cross_Output stage_output;
-    stage_output.gl_SampleMask = gl_SampleMask;
+    stage_output.gl_SampleMask = gl_SampleMask[0];
     stage_output.FragColor = FragColor;
     return stage_output;
 }

third_party/spirv-cross/reference/opt/shaders-hlsl/frag/sample-mask-in.frag

@@ -1,5 +1,5 @@
 static int gl_SampleID;
-static int gl_SampleMaskIn;
+static int gl_SampleMaskIn[1];
 static float4 FragColor;
 
 struct SPIRV_Cross_Input
@@ -15,7 +15,7 @@ struct SPIRV_Cross_Output
 
 void frag_main()
 {
-    if ((gl_SampleMaskIn & (1 << gl_SampleID)) != 0)
+    if ((gl_SampleMaskIn[0] & (1 << gl_SampleID)) != 0)
     {
         FragColor = 1.0f.xxxx;
     }
@@ -24,7 +24,7 @@ void frag_main()
 SPIRV_Cross_Output main(SPIRV_Cross_Input stage_input)
 {
     gl_SampleID = stage_input.gl_SampleID;
-    gl_SampleMaskIn = stage_input.gl_SampleMaskIn;
+    gl_SampleMaskIn[0] = stage_input.gl_SampleMaskIn;
     frag_main();
     SPIRV_Cross_Output stage_output;
     stage_output.FragColor = FragColor;

third_party/spirv-cross/reference/opt/shaders-hlsl/frag/sample-mask-out.frag

@@ -1,4 +1,4 @@
-static int gl_SampleMask;
+static int gl_SampleMask[1];
 static float4 FragColor;
 
 struct SPIRV_Cross_Output
@@ -10,14 +10,14 @@ struct SPIRV_Cross_Output
 void frag_main()
 {
     FragColor = 1.0f.xxxx;
-    gl_SampleMask = 0;
+    gl_SampleMask[0] = 0;
 }
 
 SPIRV_Cross_Output main()
 {
     frag_main();
     SPIRV_Cross_Output stage_output;
-    stage_output.gl_SampleMask = gl_SampleMask;
+    stage_output.gl_SampleMask = gl_SampleMask[0];
     stage_output.FragColor = FragColor;
     return stage_output;
 }

third_party/spirv-cross/reference/opt/shaders-msl/asm/comp/atomic-decrement.asm.comp

@@ -21,7 +21,7 @@ struct u0_counters
 
 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);
-    u0.write(uint4(uint(int(gl_GlobalInvocationID.x))), spvTexelBufferCoord(as_type<int>(as_type<float>(_29))));
+    uint _24 = atomic_fetch_sub_explicit((device atomic_uint*)&u0_counter.c, 1, memory_order_relaxed);
+    u0.write(uint4(uint(int(gl_GlobalInvocationID.x))), spvTexelBufferCoord(as_type<int>(as_type<float>(_24))));
 }
 

third_party/spirv-cross/reference/opt/shaders-msl/asm/comp/atomic-increment.asm.comp

@@ -21,7 +21,7 @@ struct u0_counters
 
 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);
-    u0.write(uint4(uint(int(gl_GlobalInvocationID.x))), spvTexelBufferCoord(as_type<int>(as_type<float>(_29))));
+    uint _24 = atomic_fetch_add_explicit((device atomic_uint*)&u0_counter.c, 1, memory_order_relaxed);
+    u0.write(uint4(uint(int(gl_GlobalInvocationID.x))), spvTexelBufferCoord(as_type<int>(as_type<float>(_24))));
 }
 

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

@@ -19,11 +19,11 @@ struct cb5_struct
 
 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;
-    uint4 _51 = as_type<uint4>(cb0_5._m0[uint(as_type<int>(as_type<float>(int(gl_LocalInvocationID.x)))) + 1u]);
-    u0.write(_51.xxxx, spvTexelBufferCoord(_44));
-    u0.write(_51.yyyy, spvTexelBufferCoord((_44 + 1u)));
-    u0.write(_51.zzzz, spvTexelBufferCoord((_44 + 2u)));
-    u0.write(_51.wwww, spvTexelBufferCoord((_44 + 3u)));
+    uint _41 = as_type<uint>(as_type<float>(int(gl_LocalInvocationID.x) << 4)) >> 2u;
+    uint4 _50 = as_type<uint4>(cb0_5._m0[uint(as_type<int>(as_type<float>(int(gl_LocalInvocationID.x)))) + 1u]);
+    u0.write(_50.xxxx, spvTexelBufferCoord(_41));
+    u0.write(_50.yyyy, spvTexelBufferCoord((_41 + 1u)));
+    u0.write(_50.zzzz, spvTexelBufferCoord((_41 + 2u)));
+    u0.write(_50.wwww, spvTexelBufferCoord((_41 + 3u)));
 }
 

third_party/spirv-cross/reference/opt/shaders-msl/asm/comp/buffer-write.asm.comp

@@ -17,8 +17,8 @@ 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]])
+kernel void main0(constant cb& _8 [[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)));
+    _buffer.write(float4(_8.value), spvTexelBufferCoord(((32u * gl_WorkGroupID.x) + gl_LocalInvocationIndex)));
 }
 

third_party/spirv-cross/reference/opt/shaders-msl/asm/comp/image-load-store-short-vector.invalid.asm.comp

@@ -8,6 +8,7 @@ using namespace metal;
 static inline __attribute__((always_inline))
 void _main(thread const uint3& id, texture2d<float, access::read_write> TargetTexture)
 {
+    TargetTexture.fence();
     float2 loaded = TargetTexture.read(uint2(id.xy)).xy;
     float2 storeTemp = loaded + float2(1.0);
     TargetTexture.write(storeTemp.xyyy, uint2((id.xy + uint2(1u))));

third_party/spirv-cross/reference/opt/shaders-msl/asm/comp/quantize.asm.comp

@@ -24,11 +24,11 @@ struct SSBO0
     float4 vec4_val;
 };
 
-kernel void main0(device SSBO0& _4 [[buffer(0)]])
+kernel void main0(device SSBO0& _12 [[buffer(0)]])
 {
-    _4.scalar = spvQuantizeToF16(_4.scalar);
-    _4.vec2_val = spvQuantizeToF16(_4.vec2_val);
-    _4.vec3_val = spvQuantizeToF16(_4.vec3_val);
-    _4.vec4_val = spvQuantizeToF16(_4.vec4_val);
+    _12.scalar = spvQuantizeToF16(_12.scalar);
+    _12.vec2_val = spvQuantizeToF16(_12.vec2_val);
+    _12.vec3_val = spvQuantizeToF16(_12.vec3_val);
+    _12.vec4_val = spvQuantizeToF16(_12.vec4_val);
 }
 

third_party/spirv-cross/reference/opt/shaders-msl/asm/comp/relaxed-block-layout.asm.comp

@@ -12,11 +12,11 @@ struct foo
     packed_half2 wibble;
 };
 
-kernel void main0(device foo& _8 [[buffer(0)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_NumWorkGroups [[threadgroups_per_grid]])
+kernel void main0(device foo& _10 [[buffer(0)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_NumWorkGroups [[threadgroups_per_grid]])
 {
-    _8.bar = gl_LocalInvocationID.x;
-    _8.baz = float3(gl_GlobalInvocationID);
-    _8.blah = uchar4(uint4(uint4(uchar4(_8.blah)).xyz + gl_WorkGroupID, 0u));
-    _8.wibble = half2(float2(half2(_8.wibble)) * float2(gl_NumWorkGroups.xy));
+    _10.bar = gl_LocalInvocationID.x;
+    _10.baz = float3(gl_GlobalInvocationID);
+    _10.blah = uchar4(uint4(uint4(uchar4(_10.blah)).xyz + gl_WorkGroupID, 0u));
+    _10.wibble = half2(float2(half2(_10.wibble)) * float2(gl_NumWorkGroups.xy));
 }
 

third_party/spirv-cross/reference/opt/shaders-msl/asm/comp/specialization-constant-workgroup.asm.comp

@@ -8,14 +8,14 @@ struct SSBO
     float a;
 };
 
-constant uint _5_tmp [[function_constant(10)]];
-constant uint _5 = is_function_constant_defined(_5_tmp) ? _5_tmp : 9u;
-constant uint _6_tmp [[function_constant(12)]];
-constant uint _6 = is_function_constant_defined(_6_tmp) ? _6_tmp : 4u;
-constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(_5, 20u, _6);
+constant uint _19_tmp [[function_constant(10)]];
+constant uint _19 = is_function_constant_defined(_19_tmp) ? _19_tmp : 9u;
+constant uint _21_tmp [[function_constant(12)]];
+constant uint _21 = is_function_constant_defined(_21_tmp) ? _21_tmp : 4u;
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(_19, 20u, _21);
 
-kernel void main0(device SSBO& _4 [[buffer(0)]])
+kernel void main0(device SSBO& _6 [[buffer(0)]])
 {
-    _4.a += 1.0;
+    _6.a += 1.0;
 }
 

third_party/spirv-cross/reference/opt/shaders-msl/asm/comp/variable-pointers-2.asm.comp

@@ -20,26 +20,26 @@ kernel void main0(device foo& buf [[buffer(0)]], constant bar& cb [[buffer(1)]],
     bool _71 = cb.d != 0;
     device foo* _72 = _71 ? &buf : nullptr;
     device foo* _67 = _72;
-    device foo* _45 = _72;
+    device foo* _65 = _72;
     thread uint3* _79 = _71 ? &gl_GlobalInvocationID : &gl_LocalInvocationID;
     thread uint3* _74 = _79;
+    device int* _49;
     device int* _52;
-    device int* _55;
-    _52 = &_72->a[0u];
-    _55 = &buf.a[0u];
-    int _57;
-    int _58;
+    _49 = &_72->a[0u];
+    _52 = &buf.a[0u];
+    int _54;
+    int _55;
     for (;;)
     {
-        _57 = *_52;
-        _58 = *_55;
-        if (_57 != _58)
+        _54 = *_49;
+        _55 = *_52;
+        if (_54 != _55)
         {
-            int _66 = (_57 + _58) + int((*_79).x);
-            *_52 = _66;
-            *_55 = _66;
+            int _63 = (_54 + _55) + int((*_79).x);
+            *_49 = _63;
+            *_52 = _63;
+            _49 = &_49[1u];
             _52 = &_52[1u];
-            _55 = &_55[1u];
             continue;
         }
         else

third_party/spirv-cross/reference/opt/shaders-msl/asm/comp/vector-builtin-type-cast-func.asm.comp

@@ -12,14 +12,14 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(16u, 16u, 1u);
 
 kernel void main0(constant cb1_struct& cb0_1 [[buffer(0)]], texture2d<float, access::write> u0 [[texture(0)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
 {
-    int2 _46 = int2(u0.get_width(), u0.get_height()) >> int2(uint2(4u));
+    int2 _37 = int2(u0.get_width(), u0.get_height()) >> int2(uint2(4u));
     int _98;
     _98 = 0;
-    for (; _98 < _46.y; _98++)
+    for (; _98 < _37.y; _98++)
     {
-        for (int _99 = 0; _99 < _46.x; )
+        for (int _99 = 0; _99 < _37.x; )
         {
-            u0.write(cb0_1._RESERVED_IDENTIFIER_FIXUP_m0[0].xxxx, uint2(((_46 * int3(gl_LocalInvocationID).xy) + int2(_98, _99))));
+            u0.write(cb0_1._RESERVED_IDENTIFIER_FIXUP_m0[0].xxxx, uint2(((_37 * int3(gl_LocalInvocationID).xy) + int2(_98, _99))));
             _99++;
             continue;
         }

third_party/spirv-cross/reference/opt/shaders-msl/asm/comp/vector-builtin-type-cast.asm.comp

@@ -12,14 +12,14 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(16u, 16u, 1u);
 
 kernel void main0(constant cb1_struct& cb0_1 [[buffer(0)]], texture2d<float, access::write> u0 [[texture(0)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
 {
-    int2 _40 = int2(u0.get_width(), u0.get_height()) >> int2(uint2(4u));
+    int2 _20 = int2(u0.get_width(), u0.get_height()) >> int2(uint2(4u));
     int _80;
     _80 = 0;
-    for (; _80 < _40.y; _80++)
+    for (; _80 < _20.y; _80++)
     {
-        for (int _81 = 0; _81 < _40.x; )
+        for (int _81 = 0; _81 < _20.x; )
         {
-            u0.write(cb0_1._RESERVED_IDENTIFIER_FIXUP_m0[0].xxxx, uint2(((_40 * int3(gl_LocalInvocationID).xy) + int2(_80, _81))));
+            u0.write(cb0_1._RESERVED_IDENTIFIER_FIXUP_m0[0].xxxx, uint2(((_20 * int3(gl_LocalInvocationID).xy) + int2(_80, _81))));
             _81++;
             continue;
         }

third_party/spirv-cross/reference/opt/shaders-msl/asm/frag/depth-array-texture-lod.lod-as-grad.1d-as-2d.agx-cube-grad.msl23.asm.frag

@@ -0,0 +1,45 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+static inline gradientcube spvGradientCube(float3 P, float3 dPdx, float3 dPdy)
+{
+    // Major axis selection
+    float3 absP = abs(P);
+    bool xMajor = absP.x >= max(absP.y, absP.z);
+    bool yMajor = absP.y >= absP.z;
+    float3 Q = xMajor ? P.yzx : (yMajor ? P.xzy : P);
+    float3 dQdx = xMajor ? dPdx.yzx : (yMajor ? dPdx.xzy : dPdx);
+    float3 dQdy = xMajor ? dPdy.yzx : (yMajor ? dPdy.xzy : dPdy);
+
+    // Skip a couple of operations compared to usual projection
+    float4 d = float4(dQdx.xy, dQdy.xy) - (Q.xy / Q.z).xyxy * float4(dQdx.zz, dQdy.zz);
+
+    // Final swizzle to put the intermediate values into non-ignored components
+    // X major: X and Z
+    // Y major: X and Y
+    // Z major: Y and Z
+    return gradientcube(xMajor ? d.xxy : d.xyx, xMajor ? d.zzw : d.zwz);
+}
+
+struct main0_out
+{
+    float4 o_color [[color(0)]];
+};
+
+struct main0_in
+{
+    float4 v_texCoord [[user(locn0)]];
+    float2 v_drefLodBias [[user(locn1)]];
+};
+
+fragment main0_out main0(main0_in in [[stage_in]], depthcube_array<float> u_sampler [[texture(0)]], sampler u_samplerSmplr [[sampler(0)]])
+{
+    main0_out out = {};
+    out.o_color = float4(u_sampler.sample_compare(u_samplerSmplr, in.v_texCoord.xyz, uint(rint(in.v_texCoord.w)), in.v_drefLodBias.x, spvGradientCube(in.v_texCoord.xyz, exp2(in.v_drefLodBias.y - 0.5) / float3(u_sampler.get_width()), exp2(in.v_drefLodBias.y - 0.5) / float3(u_sampler.get_width()))), 0.0, 0.0, 1.0);
+    return out;
+}
+

third_party/spirv-cross/reference/opt/shaders-msl/asm/frag/descriptor-array-unnamed.asm.frag

@@ -3,46 +3,46 @@
 
 using namespace metal;
 
-struct _4
+struct _6
 {
     float4 _m0;
 };
 
-struct _6
+struct _8
 {
     int _m0;
 };
 
-struct _7
+struct _9
 {
     float4 _m0;
 };
 
 struct main0_out
 {
-    float4 m_3 [[color(0)]];
+    float4 m_4 [[color(0)]];
 };
 
-fragment main0_out main0(const device _4* _5_0 [[buffer(0)]], const device _4* _5_1 [[buffer(1)]], const device _4* _5_2 [[buffer(2)]], const device _4* _5_3 [[buffer(3)]], constant _6& _20 [[buffer(4)]], constant _7* _8_0 [[buffer(5)]], constant _7* _8_1 [[buffer(6)]], constant _7* _8_2 [[buffer(7)]], constant _7* _8_3 [[buffer(8)]])
+fragment main0_out main0(const device _6* _7_0 [[buffer(0)]], const device _6* _7_1 [[buffer(1)]], const device _6* _7_2 [[buffer(2)]], const device _6* _7_3 [[buffer(3)]], constant _8& _21 [[buffer(4)]], constant _9* _10_0 [[buffer(5)]], constant _9* _10_1 [[buffer(6)]], constant _9* _10_2 [[buffer(7)]], constant _9* _10_3 [[buffer(8)]])
 {
-    const device _4* _5[] =
+    const device _6* _7[] =
     {
-        _5_0,
-        _5_1,
-        _5_2,
-        _5_3,
+        _7_0,
+        _7_1,
+        _7_2,
+        _7_3,
     };
 
-    constant _7* _8[] =
+    constant _9* _10[] =
     {
-        _8_0,
-        _8_1,
-        _8_2,
-        _8_3,
+        _10_0,
+        _10_1,
+        _10_2,
+        _10_3,
     };
 
     main0_out out = {};
-    out.m_3 = fma(_8[_20._m0]->_m0, float4(0.20000000298023223876953125), _5[_20._m0]->_m0);
+    out.m_4 = fma(_10[_21._m0]->_m0, float4(0.20000000298023223876953125), _7[_21._m0]->_m0);
     return out;
 }
 

third_party/spirv-cross/reference/opt/shaders-msl/asm/frag/extract-packed-from-composite.asm.frag

@@ -20,11 +20,11 @@ struct main0_out
     float4 _entryPointOutput [[color(0)]];
 };
 
-fragment main0_out main0(constant buf& _11 [[buffer(0)]], float4 gl_FragCoord [[position]])
+fragment main0_out main0(constant buf& _15 [[buffer(0)]], float4 gl_FragCoord [[position]])
 {
     main0_out out = {};
-    int _68 = int(gl_FragCoord.x) % 16;
-    out._entryPointOutput = float4(dot(float3(_11.results[_68].a), _11.bar.xyz), _11.results[_68].b, 0.0, 0.0);
+    int _70 = int(gl_FragCoord.x) % 16;
+    out._entryPointOutput = float4(dot(float3(_15.results[_70].a), _15.bar.xyz), _15.results[_70].b, 0.0, 0.0);
     return out;
 }
 

third_party/spirv-cross/reference/opt/shaders-msl/asm/frag/implicit-read-dep-phi.asm.frag

@@ -17,22 +17,22 @@ fragment main0_out main0(main0_in in [[stage_in]], texture2d<float> uImage [[tex
 {
     main0_out out = {};
     float phi;
-    float4 _36;
-    int _51;
-    _51 = 0;
+    float4 _45;
+    int _57;
+    _57 = 0;
     phi = 1.0;
-    _36 = float4(1.0, 2.0, 1.0, 2.0);
+    _45 = float4(1.0, 2.0, 1.0, 2.0);
     for (;;)
     {
-        out.FragColor = _36;
-        if (_51 < 4)
+        out.FragColor = _45;
+        if (_57 < 4)
         {
-            if (in.v0[_51] > 0.0)
+            if (in.v0[_57] > 0.0)
             {
-                float2 _48 = float2(phi);
-                _51++;
+                float2 _43 = float2(phi);
+                _57++;
                 phi += 2.0;
-                _36 = uImage.sample(uImageSmplr, _48, level(0.0));
+                _45 = uImage.sample(uImageSmplr, _43, level(0.0));
                 continue;
             }
             else

third_party/spirv-cross/reference/opt/shaders-msl/asm/frag/line-directive.line.asm.frag

@@ -33,14 +33,14 @@ fragment main0_out main0(main0_in in [[stage_in]])
         out.FragColor = 4.0;
     }
 #line 19 "test.frag"
-    for (int _127 = 0; float(_127) < (40.0 + in.vColor); )
+    for (int _131 = 0; float(_131) < (40.0 + in.vColor); )
     {
 #line 21 "test.frag"
         out.FragColor += 0.20000000298023223876953125;
 #line 22 "test.frag"
         out.FragColor += 0.300000011920928955078125;
 #line 19 "test.frag"
-        _127 += (int(in.vColor) + 5);
+        _131 += (int(in.vColor) + 5);
         continue;
     }
 #line 25 "test.frag"

third_party/spirv-cross/reference/opt/shaders-msl/asm/frag/locations-components.asm.frag

@@ -10,23 +10,23 @@ struct main0_out
 
 struct main0_in
 {
-    float2 m_2 [[user(locn1)]];
-    float m_3 [[user(locn1_2)]];
-    float m_4 [[user(locn2), flat]];
-    uint m_5 [[user(locn2_1)]];
-    uint m_6 [[user(locn2_2)]];
+    float2 m_8 [[user(locn1)]];
+    float m_16 [[user(locn1_2)]];
+    float m_22 [[user(locn2), flat]];
+    uint m_28 [[user(locn2_1)]];
+    uint m_33 [[user(locn2_2)]];
 };
 
 fragment main0_out main0(main0_in in [[stage_in]])
 {
     main0_out out = {};
     float4 v1;
-    v1 = float4(in.m_2.x, in.m_2.y, v1.z, v1.w);
-    v1.z = in.m_3;
+    v1 = float4(in.m_8.x, in.m_8.y, v1.z, v1.w);
+    v1.z = in.m_16;
     float4 v2;
-    v2.x = in.m_4;
-    v2.y = as_type<float>(in.m_5);
-    v2.z = as_type<float>(in.m_6);
+    v2.x = in.m_22;
+    v2.y = as_type<float>(in.m_28);
+    v2.z = as_type<float>(in.m_33);
     out.o0.y = float(as_type<uint>(as_type<float>(as_type<int>(v2.y) + as_type<int>(v2.z))));
     out.o0.x = v1.y + v2.x;
     out.o0 = float4(out.o0.x, out.o0.y, v1.z, v1.x);

third_party/spirv-cross/reference/opt/shaders-msl/asm/frag/lut-promotion-initializer.asm.frag

@@ -44,9 +44,9 @@ struct spvUnsafeArray
     }
 };
 
-constant spvUnsafeArray<float, 16> _46 = spvUnsafeArray<float, 16>({ 1.0, 2.0, 3.0, 4.0, 1.0, 2.0, 3.0, 4.0, 1.0, 2.0, 3.0, 4.0, 1.0, 2.0, 3.0, 4.0 });
-constant spvUnsafeArray<float4, 4> _76 = spvUnsafeArray<float4, 4>({ float4(0.0), float4(1.0), float4(8.0), float4(5.0) });
-constant spvUnsafeArray<float4, 4> _90 = spvUnsafeArray<float4, 4>({ float4(20.0), float4(30.0), float4(50.0), float4(60.0) });
+constant spvUnsafeArray<float, 16> _16 = spvUnsafeArray<float, 16>({ 1.0, 2.0, 3.0, 4.0, 1.0, 2.0, 3.0, 4.0, 1.0, 2.0, 3.0, 4.0, 1.0, 2.0, 3.0, 4.0 });
+constant spvUnsafeArray<float4, 4> _60 = spvUnsafeArray<float4, 4>({ float4(0.0), float4(1.0), float4(8.0), float4(5.0) });
+constant spvUnsafeArray<float4, 4> _104 = spvUnsafeArray<float4, 4>({ float4(20.0), float4(30.0), float4(50.0), float4(60.0) });
 
 struct main0_out
 {
@@ -63,32 +63,32 @@ fragment main0_out main0(main0_in in [[stage_in]])
     spvUnsafeArray<float4, 4> foobar = spvUnsafeArray<float4, 4>({ float4(0.0), float4(1.0), float4(8.0), float4(5.0) });
     spvUnsafeArray<float4, 4> baz = spvUnsafeArray<float4, 4>({ float4(0.0), float4(1.0), float4(8.0), float4(5.0) });
     main0_out out = {};
-    out.FragColor = _46[in.index];
+    out.FragColor = _16[in.index];
     if (in.index < 10)
     {
-        out.FragColor += _46[in.index ^ 1];
+        out.FragColor += _16[in.index ^ 1];
     }
     else
     {
-        out.FragColor += _46[in.index & 1];
+        out.FragColor += _16[in.index & 1];
     }
-    bool _99 = in.index > 30;
-    if (_99)
+    bool _63 = in.index > 30;
+    if (_63)
     {
-        out.FragColor += _76[in.index & 3].y;
+        out.FragColor += _60[in.index & 3].y;
     }
     else
     {
-        out.FragColor += _76[in.index & 1].x;
+        out.FragColor += _60[in.index & 1].x;
     }
-    if (_99)
+    if (_63)
     {
         foobar[1].z = 20.0;
     }
-    int _37 = in.index & 3;
-    out.FragColor += foobar[_37].z;
-    baz = _90;
-    out.FragColor += baz[_37].z;
+    int _91 = in.index & 3;
+    out.FragColor += foobar[_91].z;
+    baz = _104;
+    out.FragColor += baz[_91].z;
     return out;
 }
 

third_party/spirv-cross/reference/opt/shaders-msl/asm/frag/pull-model-interpolation.asm.msl23.frag

@@ -154,24 +154,24 @@ fragment main0_out main0(main0_in in [[stage_in]], uint gl_SampleID [[sample_id]
     out.FragColor.x += in.s_w_1.interpolate_at_centroid();
     out.FragColor.x += in.s_w_0.interpolate_at_sample(2);
     out.FragColor.x += in.s_w_1.interpolate_at_offset(float2(-0.100000001490116119384765625, 0.100000001490116119384765625) + 0.4375);
-    float2 _328 = out.FragColor.xy + in.baz.interpolate_at_sample(gl_SampleID);
-    out.FragColor = float4(_328.x, _328.y, out.FragColor.z, out.FragColor.w);
+    float2 _329 = out.FragColor.xy + in.baz.interpolate_at_sample(gl_SampleID);
+    out.FragColor = float4(_329.x, _329.y, out.FragColor.z, out.FragColor.w);
     out.FragColor.x += in.baz.interpolate_at_centroid().x;
     out.FragColor.y += in.baz.interpolate_at_sample(3).y;
     out.FragColor.z += in.baz.interpolate_at_offset(float2(-0.100000001490116119384765625, 0.100000001490116119384765625) + 0.4375).y;
-    float2 _353 = out.FragColor.xy + in.a_1.interpolate_at_centroid();
-    out.FragColor = float4(_353.x, _353.y, out.FragColor.z, out.FragColor.w);
-    float2 _360 = out.FragColor.xy + in.a_0.interpolate_at_sample(2);
-    out.FragColor = float4(_360.x, _360.y, out.FragColor.z, out.FragColor.w);
-    float2 _367 = out.FragColor.xy + in.a_1.interpolate_at_offset(float2(-0.100000001490116119384765625, 0.100000001490116119384765625) + 0.4375);
-    out.FragColor = float4(_367.x, _367.y, out.FragColor.z, out.FragColor.w);
+    float2 _354 = out.FragColor.xy + in.a_1.interpolate_at_centroid();
+    out.FragColor = float4(_354.x, _354.y, out.FragColor.z, out.FragColor.w);
+    float2 _361 = out.FragColor.xy + in.a_0.interpolate_at_sample(2);
+    out.FragColor = float4(_361.x, _361.y, out.FragColor.z, out.FragColor.w);
+    float2 _368 = out.FragColor.xy + in.a_1.interpolate_at_offset(float2(-0.100000001490116119384765625, 0.100000001490116119384765625) + 0.4375);
+    out.FragColor = float4(_368.x, _368.y, out.FragColor.z, out.FragColor.w);
     out.FragColor += s.z;
-    float2 _379 = out.FragColor.xy + in.s_z.interpolate_at_centroid().yy;
-    out.FragColor = float4(_379.x, _379.y, out.FragColor.z, out.FragColor.w);
-    float2 _387 = out.FragColor.yz + in.s_z.interpolate_at_sample(3).xy;
-    out.FragColor = float4(out.FragColor.x, _387.x, _387.y, out.FragColor.w);
-    float2 _395 = out.FragColor.zw + in.s_z.interpolate_at_offset(float2(-0.100000001490116119384765625, 0.100000001490116119384765625) + 0.4375).wx;
-    out.FragColor = float4(out.FragColor.x, out.FragColor.y, _395.x, _395.y);
+    float2 _380 = out.FragColor.xy + in.s_z.interpolate_at_centroid().yy;
+    out.FragColor = float4(_380.x, _380.y, out.FragColor.z, out.FragColor.w);
+    float2 _388 = out.FragColor.yz + in.s_z.interpolate_at_sample(3).xy;
+    out.FragColor = float4(out.FragColor.x, _388.x, _388.y, out.FragColor.w);
+    float2 _396 = out.FragColor.zw + in.s_z.interpolate_at_offset(float2(-0.100000001490116119384765625, 0.100000001490116119384765625) + 0.4375).wx;
+    out.FragColor = float4(out.FragColor.x, out.FragColor.y, _396.x, _396.y);
     out.FragColor += s.u[0];
     out.FragColor += in.s_u_1.interpolate_at_centroid();
     out.FragColor += in.s_u_0.interpolate_at_sample(2);

third_party/spirv-cross/reference/opt/shaders-msl/asm/frag/unreachable.asm.frag

@@ -16,21 +16,21 @@ struct main0_in
 fragment main0_out main0(main0_in in [[stage_in]])
 {
     main0_out out = {};
-    float4 _33;
+    float4 _46;
     for (;;)
     {
         if (in.counter == 10)
         {
-            _33 = float4(10.0);
+            _46 = float4(10.0);
             break;
         }
         else
         {
-            _33 = float4(30.0);
+            _46 = float4(30.0);
             break;
         }
     }
-    out.FragColor = _33;
+    out.FragColor = _46;
     return out;
 }
 

third_party/spirv-cross/reference/opt/shaders-msl/asm/frag/vector-shuffle-oom.asm.frag

@@ -3,19 +3,19 @@
 
 using namespace metal;
 
-struct _28
+struct _15
 {
     float4 _m0;
 };
 
-struct _6
+struct _3
 {
     float4 _m0;
     float _m1;
     float4 _m2;
 };
 
-struct _10
+struct _4
 {
     float3 _m0;
     packed_float3 _m1;
@@ -44,7 +44,7 @@ struct _10
     float4 _m24;
 };
 
-struct _18
+struct _7
 {
     float4x4 _m0;
     float4x4 _m1;
@@ -90,188 +90,188 @@ struct _18
 
 struct main0_out
 {
-    float4 m_5 [[color(0)]];
+    float4 m_4317 [[color(0)]];
 };
 
-fragment main0_out main0(constant _6& _7 [[buffer(0)]], constant _10& _11 [[buffer(1)]], constant _18& _19 [[buffer(2)]], texture2d<float> _8 [[texture(0)]], texture2d<float> _12 [[texture(1)]], texture2d<float> _14 [[texture(2)]], sampler _9 [[sampler(0)]], sampler _13 [[sampler(1)]], sampler _15 [[sampler(2)]], float4 gl_FragCoord [[position]])
+fragment main0_out main0(constant _3& _22044 [[buffer(0)]], constant _4& _12348 [[buffer(1)]], constant _7& _15259 [[buffer(2)]], texture2d<float> _5785 [[texture(0)]], texture2d<float> _3312 [[texture(1)]], texture2d<float> _4862 [[texture(2)]], sampler _5688 [[sampler(0)]], sampler _4646 [[sampler(1)]], sampler _3594 [[sampler(2)]], float4 gl_FragCoord [[position]])
 {
     main0_out out = {};
-    float4 _88 = _7._m2 * _7._m0.xyxy;
-    float2 _95 = _88.xy;
-    float2 _96 = _88.zw;
-    float2 _97 = fast::clamp(fma(gl_FragCoord.xy, _19._m23.xy, float2(0.0, -2.0) * _7._m0.xy), _95, _96);
-    float3 _109 = float3(_11._m5) * fast::clamp(_8.sample(_9, _97, level(0.0)).w * _7._m1, 0.0, 1.0);
-    float4 _113 = _12.sample(_13, _97, level(0.0));
-    float _114 = _113.y;
-    float3 _129;
-    if (_114 > 0.0)
+    float4 _17581 = _22044._m2 * _22044._m0.xyxy;
+    float2 _7011 = _17581.xy;
+    float2 _21058 = _17581.zw;
+    float2 _13149 = fast::clamp(fma(gl_FragCoord.xy, _15259._m23.xy, float2(0.0, -2.0) * _22044._m0.xy), _7011, _21058);
+    float3 _12103 = float3(_12348._m5) * fast::clamp(_5785.sample(_5688, _13149, level(0.0)).w * _22044._m1, 0.0, 1.0);
+    float4 _17670 = _3312.sample(_4646, _13149, level(0.0));
+    float _16938 = _17670.y;
+    float3 _7719;
+    if (_16938 > 0.0)
     {
-        _129 = _109 + (_14.sample(_15, _97, level(0.0)).xyz * fast::clamp(_114 * _113.z, 0.0, 1.0));
+        _7719 = _12103 + (_4862.sample(_3594, _13149, level(0.0)).xyz * fast::clamp(_16938 * _17670.z, 0.0, 1.0));
     }
     else
     {
-        _129 = _109;
+        _7719 = _12103;
     }
-    float2 _144 = fast::clamp(fma(gl_FragCoord.xy, _19._m23.xy, float2(-1.0) * _7._m0.xy), _95, _96);
-    float3 _156 = float3(_11._m5) * fast::clamp(_8.sample(_9, _144, level(0.0)).w * _7._m1, 0.0, 1.0);
-    float4 _160 = _12.sample(_13, _144, level(0.0));
-    float _161 = _160.y;
-    float3 _176;
-    if (_161 > 0.0)
+    float2 _13150 = fast::clamp(fma(gl_FragCoord.xy, _15259._m23.xy, float2(-1.0) * _22044._m0.xy), _7011, _21058);
+    float3 _12104 = float3(_12348._m5) * fast::clamp(_5785.sample(_5688, _13150, level(0.0)).w * _22044._m1, 0.0, 1.0);
+    float4 _17671 = _3312.sample(_4646, _13150, level(0.0));
+    float _16939 = _17671.y;
+    float3 _7720;
+    if (_16939 > 0.0)
     {
-        _176 = _156 + (_14.sample(_15, _144, level(0.0)).xyz * fast::clamp(_161 * _160.z, 0.0, 1.0));
+        _7720 = _12104 + (_4862.sample(_3594, _13150, level(0.0)).xyz * fast::clamp(_16939 * _17671.z, 0.0, 1.0));
     }
     else
     {
-        _176 = _156;
+        _7720 = _12104;
     }
-    float2 _191 = fast::clamp(fma(gl_FragCoord.xy, _19._m23.xy, float2(0.0, -1.0) * _7._m0.xy), _95, _96);
-    float3 _203 = float3(_11._m5) * fast::clamp(_8.sample(_9, _191, level(0.0)).w * _7._m1, 0.0, 1.0);
-    float4 _207 = _12.sample(_13, _191, level(0.0));
-    float _208 = _207.y;
-    float3 _223;
-    if (_208 > 0.0)
+    float2 _13151 = fast::clamp(fma(gl_FragCoord.xy, _15259._m23.xy, float2(0.0, -1.0) * _22044._m0.xy), _7011, _21058);
+    float3 _12105 = float3(_12348._m5) * fast::clamp(_5785.sample(_5688, _13151, level(0.0)).w * _22044._m1, 0.0, 1.0);
+    float4 _17672 = _3312.sample(_4646, _13151, level(0.0));
+    float _16940 = _17672.y;
+    float3 _7721;
+    if (_16940 > 0.0)
     {
-        _223 = _203 + (_14.sample(_15, _191, level(0.0)).xyz * fast::clamp(_208 * _207.z, 0.0, 1.0));
+        _7721 = _12105 + (_4862.sample(_3594, _13151, level(0.0)).xyz * fast::clamp(_16940 * _17672.z, 0.0, 1.0));
     }
     else
     {
-        _223 = _203;
+        _7721 = _12105;
     }
-    float2 _238 = fast::clamp(fma(gl_FragCoord.xy, _19._m23.xy, float2(1.0, -1.0) * _7._m0.xy), _95, _96);
-    float3 _250 = float3(_11._m5) * fast::clamp(_8.sample(_9, _238, level(0.0)).w * _7._m1, 0.0, 1.0);
-    float4 _254 = _12.sample(_13, _238, level(0.0));
-    float _255 = _254.y;
-    float3 _270;
-    if (_255 > 0.0)
+    float2 _13152 = fast::clamp(fma(gl_FragCoord.xy, _15259._m23.xy, float2(1.0, -1.0) * _22044._m0.xy), _7011, _21058);
+    float3 _12106 = float3(_12348._m5) * fast::clamp(_5785.sample(_5688, _13152, level(0.0)).w * _22044._m1, 0.0, 1.0);
+    float4 _17673 = _3312.sample(_4646, _13152, level(0.0));
+    float _16941 = _17673.y;
+    float3 _7722;
+    if (_16941 > 0.0)
     {
-        _270 = _250 + (_14.sample(_15, _238, level(0.0)).xyz * fast::clamp(_255 * _254.z, 0.0, 1.0));
+        _7722 = _12106 + (_4862.sample(_3594, _13152, level(0.0)).xyz * fast::clamp(_16941 * _17673.z, 0.0, 1.0));
     }
     else
     {
-        _270 = _250;
+        _7722 = _12106;
     }
-    float2 _285 = fast::clamp(fma(gl_FragCoord.xy, _19._m23.xy, float2(-2.0, 0.0) * _7._m0.xy), _95, _96);
-    float3 _297 = float3(_11._m5) * fast::clamp(_8.sample(_9, _285, level(0.0)).w * _7._m1, 0.0, 1.0);
-    float4 _301 = _12.sample(_13, _285, level(0.0));
-    float _302 = _301.y;
-    float3 _317;
-    if (_302 > 0.0)
+    float2 _13153 = fast::clamp(fma(gl_FragCoord.xy, _15259._m23.xy, float2(-2.0, 0.0) * _22044._m0.xy), _7011, _21058);
+    float3 _12107 = float3(_12348._m5) * fast::clamp(_5785.sample(_5688, _13153, level(0.0)).w * _22044._m1, 0.0, 1.0);
+    float4 _17674 = _3312.sample(_4646, _13153, level(0.0));
+    float _16942 = _17674.y;
+    float3 _7723;
+    if (_16942 > 0.0)
     {
-        _317 = _297 + (_14.sample(_15, _285, level(0.0)).xyz * fast::clamp(_302 * _301.z, 0.0, 1.0));
+        _7723 = _12107 + (_4862.sample(_3594, _13153, level(0.0)).xyz * fast::clamp(_16942 * _17674.z, 0.0, 1.0));
     }
     else
     {
-        _317 = _297;
+        _7723 = _12107;
     }
-    float2 _332 = fast::clamp(fma(gl_FragCoord.xy, _19._m23.xy, float2(-1.0, 0.0) * _7._m0.xy), _95, _96);
-    float3 _344 = float3(_11._m5) * fast::clamp(_8.sample(_9, _332, level(0.0)).w * _7._m1, 0.0, 1.0);
-    float4 _348 = _12.sample(_13, _332, level(0.0));
-    float _349 = _348.y;
-    float3 _364;
-    if (_349 > 0.0)
+    float2 _13154 = fast::clamp(fma(gl_FragCoord.xy, _15259._m23.xy, float2(-1.0, 0.0) * _22044._m0.xy), _7011, _21058);
+    float3 _12108 = float3(_12348._m5) * fast::clamp(_5785.sample(_5688, _13154, level(0.0)).w * _22044._m1, 0.0, 1.0);
+    float4 _17675 = _3312.sample(_4646, _13154, level(0.0));
+    float _16943 = _17675.y;
+    float3 _7724;
+    if (_16943 > 0.0)
     {
-        _364 = _344 + (_14.sample(_15, _332, level(0.0)).xyz * fast::clamp(_349 * _348.z, 0.0, 1.0));
+        _7724 = _12108 + (_4862.sample(_3594, _13154, level(0.0)).xyz * fast::clamp(_16943 * _17675.z, 0.0, 1.0));
     }
     else
     {
-        _364 = _344;
+        _7724 = _12108;
     }
-    float2 _379 = fast::clamp(gl_FragCoord.xy * _19._m23.xy, _95, _96);
-    float3 _391 = float3(_11._m5) * fast::clamp(_8.sample(_9, _379, level(0.0)).w * _7._m1, 0.0, 1.0);
-    float4 _395 = _12.sample(_13, _379, level(0.0));
-    float _396 = _395.y;
-    float3 _411;
-    if (_396 > 0.0)
+    float2 _13155 = fast::clamp(gl_FragCoord.xy * _15259._m23.xy, _7011, _21058);
+    float3 _12109 = float3(_12348._m5) * fast::clamp(_5785.sample(_5688, _13155, level(0.0)).w * _22044._m1, 0.0, 1.0);
+    float4 _17676 = _3312.sample(_4646, _13155, level(0.0));
+    float _16944 = _17676.y;
+    float3 _7725;
+    if (_16944 > 0.0)
     {
-        _411 = _391 + (_14.sample(_15, _379, level(0.0)).xyz * fast::clamp(_396 * _395.z, 0.0, 1.0));
+        _7725 = _12109 + (_4862.sample(_3594, _13155, level(0.0)).xyz * fast::clamp(_16944 * _17676.z, 0.0, 1.0));
     }
     else
     {
-        _411 = _391;
+        _7725 = _12109;
     }
-    float2 _426 = fast::clamp(fma(gl_FragCoord.xy, _19._m23.xy, float2(1.0, 0.0) * _7._m0.xy), _95, _96);
-    float3 _438 = float3(_11._m5) * fast::clamp(_8.sample(_9, _426, level(0.0)).w * _7._m1, 0.0, 1.0);
-    float4 _442 = _12.sample(_13, _426, level(0.0));
-    float _443 = _442.y;
-    float3 _458;
-    if (_443 > 0.0)
+    float2 _13156 = fast::clamp(fma(gl_FragCoord.xy, _15259._m23.xy, float2(1.0, 0.0) * _22044._m0.xy), _7011, _21058);
+    float3 _12110 = float3(_12348._m5) * fast::clamp(_5785.sample(_5688, _13156, level(0.0)).w * _22044._m1, 0.0, 1.0);
+    float4 _17677 = _3312.sample(_4646, _13156, level(0.0));
+    float _16945 = _17677.y;
+    float3 _7726;
+    if (_16945 > 0.0)
     {
-        _458 = _438 + (_14.sample(_15, _426, level(0.0)).xyz * fast::clamp(_443 * _442.z, 0.0, 1.0));
+        _7726 = _12110 + (_4862.sample(_3594, _13156, level(0.0)).xyz * fast::clamp(_16945 * _17677.z, 0.0, 1.0));
     }
     else
     {
-        _458 = _438;
+        _7726 = _12110;
     }
-    float2 _473 = fast::clamp(fma(gl_FragCoord.xy, _19._m23.xy, float2(2.0, 0.0) * _7._m0.xy), _95, _96);
-    float3 _485 = float3(_11._m5) * fast::clamp(_8.sample(_9, _473, level(0.0)).w * _7._m1, 0.0, 1.0);
-    float4 _489 = _12.sample(_13, _473, level(0.0));
-    float _490 = _489.y;
-    float3 _505;
-    if (_490 > 0.0)
+    float2 _13157 = fast::clamp(fma(gl_FragCoord.xy, _15259._m23.xy, float2(2.0, 0.0) * _22044._m0.xy), _7011, _21058);
+    float3 _12111 = float3(_12348._m5) * fast::clamp(_5785.sample(_5688, _13157, level(0.0)).w * _22044._m1, 0.0, 1.0);
+    float4 _17678 = _3312.sample(_4646, _13157, level(0.0));
+    float _16946 = _17678.y;
+    float3 _7727;
+    if (_16946 > 0.0)
     {
-        _505 = _485 + (_14.sample(_15, _473, level(0.0)).xyz * fast::clamp(_490 * _489.z, 0.0, 1.0));
+        _7727 = _12111 + (_4862.sample(_3594, _13157, level(0.0)).xyz * fast::clamp(_16946 * _17678.z, 0.0, 1.0));
     }
     else
     {
-        _505 = _485;
+        _7727 = _12111;
     }
-    float2 _520 = fast::clamp(fma(gl_FragCoord.xy, _19._m23.xy, float2(-1.0, 1.0) * _7._m0.xy), _95, _96);
-    float3 _532 = float3(_11._m5) * fast::clamp(_8.sample(_9, _520, level(0.0)).w * _7._m1, 0.0, 1.0);
-    float4 _536 = _12.sample(_13, _520, level(0.0));
-    float _537 = _536.y;
-    float3 _552;
-    if (_537 > 0.0)
+    float2 _13158 = fast::clamp(fma(gl_FragCoord.xy, _15259._m23.xy, float2(-1.0, 1.0) * _22044._m0.xy), _7011, _21058);
+    float3 _12112 = float3(_12348._m5) * fast::clamp(_5785.sample(_5688, _13158, level(0.0)).w * _22044._m1, 0.0, 1.0);
+    float4 _17679 = _3312.sample(_4646, _13158, level(0.0));
+    float _16947 = _17679.y;
+    float3 _7728;
+    if (_16947 > 0.0)
     {
-        _552 = _532 + (_14.sample(_15, _520, level(0.0)).xyz * fast::clamp(_537 * _536.z, 0.0, 1.0));
+        _7728 = _12112 + (_4862.sample(_3594, _13158, level(0.0)).xyz * fast::clamp(_16947 * _17679.z, 0.0, 1.0));
     }
     else
     {
-        _552 = _532;
+        _7728 = _12112;
     }
-    float2 _567 = fast::clamp(fma(gl_FragCoord.xy, _19._m23.xy, float2(0.0, 1.0) * _7._m0.xy), _95, _96);
-    float3 _579 = float3(_11._m5) * fast::clamp(_8.sample(_9, _567, level(0.0)).w * _7._m1, 0.0, 1.0);
-    float4 _583 = _12.sample(_13, _567, level(0.0));
-    float _584 = _583.y;
-    float3 _599;
-    if (_584 > 0.0)
+    float2 _13159 = fast::clamp(fma(gl_FragCoord.xy, _15259._m23.xy, float2(0.0, 1.0) * _22044._m0.xy), _7011, _21058);
+    float3 _12113 = float3(_12348._m5) * fast::clamp(_5785.sample(_5688, _13159, level(0.0)).w * _22044._m1, 0.0, 1.0);
+    float4 _17680 = _3312.sample(_4646, _13159, level(0.0));
+    float _16948 = _17680.y;
+    float3 _7729;
+    if (_16948 > 0.0)
     {
-        _599 = _579 + (_14.sample(_15, _567, level(0.0)).xyz * fast::clamp(_584 * _583.z, 0.0, 1.0));
+        _7729 = _12113 + (_4862.sample(_3594, _13159, level(0.0)).xyz * fast::clamp(_16948 * _17680.z, 0.0, 1.0));
     }
     else
     {
-        _599 = _579;
+        _7729 = _12113;
     }
-    float2 _614 = fast::clamp(fma(gl_FragCoord.xy, _19._m23.xy, _7._m0.xy), _95, _96);
-    float3 _626 = float3(_11._m5) * fast::clamp(_8.sample(_9, _614, level(0.0)).w * _7._m1, 0.0, 1.0);
-    float4 _630 = _12.sample(_13, _614, level(0.0));
-    float _631 = _630.y;
-    float3 _646;
-    if (_631 > 0.0)
+    float2 _13160 = fast::clamp(fma(gl_FragCoord.xy, _15259._m23.xy, _22044._m0.xy), _7011, _21058);
+    float3 _12114 = float3(_12348._m5) * fast::clamp(_5785.sample(_5688, _13160, level(0.0)).w * _22044._m1, 0.0, 1.0);
+    float4 _17681 = _3312.sample(_4646, _13160, level(0.0));
+    float _16949 = _17681.y;
+    float3 _7730;
+    if (_16949 > 0.0)
     {
-        _646 = _626 + (_14.sample(_15, _614, level(0.0)).xyz * fast::clamp(_631 * _630.z, 0.0, 1.0));
+        _7730 = _12114 + (_4862.sample(_3594, _13160, level(0.0)).xyz * fast::clamp(_16949 * _17681.z, 0.0, 1.0));
     }
     else
     {
-        _646 = _626;
+        _7730 = _12114;
     }
-    float2 _661 = fast::clamp(fma(gl_FragCoord.xy, _19._m23.xy, float2(0.0, 2.0) * _7._m0.xy), _95, _96);
-    float3 _673 = float3(_11._m5) * fast::clamp(_8.sample(_9, _661, level(0.0)).w * _7._m1, 0.0, 1.0);
-    float4 _677 = _12.sample(_13, _661, level(0.0));
-    float _678 = _677.y;
-    float3 _693;
-    if (_678 > 0.0)
+    float2 _13161 = fast::clamp(fma(gl_FragCoord.xy, _15259._m23.xy, float2(0.0, 2.0) * _22044._m0.xy), _7011, _21058);
+    float3 _12115 = float3(_12348._m5) * fast::clamp(_5785.sample(_5688, _13161, level(0.0)).w * _22044._m1, 0.0, 1.0);
+    float4 _17682 = _3312.sample(_4646, _13161, level(0.0));
+    float _16950 = _17682.y;
+    float3 _7731;
+    if (_16950 > 0.0)
     {
-        _693 = _673 + (_14.sample(_15, _661, level(0.0)).xyz * fast::clamp(_678 * _677.z, 0.0, 1.0));
+        _7731 = _12115 + (_4862.sample(_3594, _13161, level(0.0)).xyz * fast::clamp(_16950 * _17682.z, 0.0, 1.0));
     }
     else
     {
-        _693 = _673;
+        _7731 = _12115;
     }
-    float3 _702 = (((((((((((((_129 * 0.5).xyz + (_176 * 0.5)).xyz + (_223 * 0.75)).xyz + (_270 * 0.5)).xyz + (_317 * 0.5)).xyz + (_364 * 0.75)).xyz + (_411 * 1.0)).xyz + (_458 * 0.75)).xyz + (_505 * 0.5)).xyz + (_552 * 0.5)).xyz + (_599 * 0.75)).xyz + (_646 * 0.5)).xyz + (_693 * 0.5)).xyz * float3(0.125);
-    _28 _750 = _28{ float4(_702.x, _702.y, _702.z, float4(0.0).w) };
-    _750._m0.w = 1.0;
-    out.m_5 = _750._m0;
+    float3 _13750 = (((((((((((((_7719 * 0.5).xyz + (_7720 * 0.5)).xyz + (_7721 * 0.75)).xyz + (_7722 * 0.5)).xyz + (_7723 * 0.5)).xyz + (_7724 * 0.75)).xyz + (_7725 * 1.0)).xyz + (_7726 * 0.75)).xyz + (_7727 * 0.5)).xyz + (_7728 * 0.5)).xyz + (_7729 * 0.75)).xyz + (_7730 * 0.5)).xyz + (_7731 * 0.5)).xyz * float3(0.125);
+    _15 _25050 = _15{ float4(_13750.x, _13750.y, _13750.z, float4(0.0).w) };
+    _25050._m0.w = 1.0;
+    out.m_4317 = _25050._m0;
     return out;
 }
 

third_party/spirv-cross/reference/opt/shaders-msl/asm/vert/packed-bool-to-uint.asm.vert

@@ -25,11 +25,11 @@ struct main0_in
     float4 a_position [[attribute(0)]];
 };
 
-vertex main0_out main0(main0_in in [[stage_in]], constant defaultUniformsVS& _9 [[buffer(0)]], uint gl_VertexIndex [[vertex_id]])
+vertex main0_out main0(main0_in in [[stage_in]], constant defaultUniformsVS& _11 [[buffer(0)]], uint gl_VertexIndex [[vertex_id]])
 {
     main0_out out = {};
-    out.gl_Position = _9.umatrix * float4(_9.uquad[int(gl_VertexIndex)].x, _9.uquad[int(gl_VertexIndex)].y, in.a_position.z, in.a_position.w);
-    if (_9.flags.flags[0] != 0u)
+    out.gl_Position = _11.umatrix * float4(_11.uquad[int(gl_VertexIndex)].x, _11.uquad[int(gl_VertexIndex)].y, in.a_position.z, in.a_position.w);
+    if (_11.flags.flags[0] != 0u)
     {
         out.gl_Position.z = 0.0;
     }

third_party/spirv-cross/reference/opt/shaders-msl/asm/vert/packed-bool2-to-packed_uint2.asm.vert

@@ -25,11 +25,11 @@ struct main0_in
     float4 a_position [[attribute(0)]];
 };
 
-vertex main0_out main0(main0_in in [[stage_in]], constant defaultUniformsVS& _9 [[buffer(0)]], uint gl_VertexIndex [[vertex_id]])
+vertex main0_out main0(main0_in in [[stage_in]], constant defaultUniformsVS& _11 [[buffer(0)]], uint gl_VertexIndex [[vertex_id]])
 {
     main0_out out = {};
-    out.gl_Position = _9.umatrix * float4(_9.uquad[int(gl_VertexIndex)].x, _9.uquad[int(gl_VertexIndex)].y, in.a_position.z, in.a_position.w);
-    if (_9.flags.flags[0].x != 0u)
+    out.gl_Position = _11.umatrix * float4(_11.uquad[int(gl_VertexIndex)].x, _11.uquad[int(gl_VertexIndex)].y, in.a_position.z, in.a_position.w);
+    if (_11.flags.flags[0].x != 0u)
     {
         out.gl_Position.z = 0.0;
     }

third_party/spirv-cross/reference/opt/shaders-msl/asm/vert/spec-constant-op-composite.asm.vert

@@ -3,20 +3,20 @@
 
 using namespace metal;
 
-constant int _7_tmp [[function_constant(201)]];
-constant int _7 = is_function_constant_defined(_7_tmp) ? _7_tmp : -10;
-constant int _20 = (_7 + 2);
-constant uint _8_tmp [[function_constant(202)]];
-constant uint _8 = is_function_constant_defined(_8_tmp) ? _8_tmp : 100u;
-constant uint _25 = (_8 % 5u);
-constant int _30 = _7 - (-3) * (_7 / (-3));
-constant int4 _32 = int4(20, 30, _20, _30);
-constant int2 _34 = int2(_32.y, _32.x);
-constant int _35 = _32.y;
+constant int _13_tmp [[function_constant(201)]];
+constant int _13 = is_function_constant_defined(_13_tmp) ? _13_tmp : -10;
+constant int _15 = (_13 + 2);
+constant uint _24_tmp [[function_constant(202)]];
+constant uint _24 = is_function_constant_defined(_24_tmp) ? _24_tmp : 100u;
+constant uint _26 = (_24 % 5u);
+constant int _61 = _13 - (-3) * (_13 / (-3));
+constant int4 _36 = int4(20, 30, _15, _61);
+constant int2 _41 = int2(_36.y, _36.x);
+constant int _62 = _36.y;
 
 struct main0_out
 {
-    int m_4 [[user(locn0)]];
+    int m_58 [[user(locn0)]];
     float4 gl_Position [[position]];
 };
 
@@ -24,12 +24,12 @@ vertex main0_out main0()
 {
     main0_out out = {};
     float4 _66 = float4(0.0);
-    _66.y = float(_20);
-    _66.z = float(_25);
-    float4 _55 = _66 + float4(_32);
-    float2 _59 = _55.xy + float2(_34);
-    out.gl_Position = float4(_59.x, _59.y, _55.z, _55.w);
-    out.m_4 = _35;
+    _66.y = float(_15);
+    _66.z = float(_26);
+    float4 _39 = _66 + float4(_36);
+    float2 _46 = _39.xy + float2(_41);
+    out.gl_Position = float4(_46.x, _46.y, _39.z, _39.w);
+    out.m_58 = _62;
     return out;
 }
 

third_party/spirv-cross/reference/opt/shaders-msl/comp/argument-buffers-runtime-array-buffer.argument.device-argument-buffer.argument-tier-1.msl2.comp

@@ -0,0 +1,77 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+template<typename T>
+struct spvDescriptor
+{
+    T value;
+};
+
+template<typename T>
+struct spvDescriptorArray
+{
+    spvDescriptorArray(const device spvDescriptor<T>* ptr) : ptr(ptr)
+    {
+    }
+    const device T& operator [] (size_t i) const
+    {
+        return ptr[i].value;
+    }
+    const device spvDescriptor<T>* ptr;
+};
+
+struct SSBO
+{
+    float4 v[1];
+};
+
+struct UBO
+{
+    float4 v[1024];
+};
+
+struct UBOs
+{
+    float4 v;
+};
+
+struct SSBOIn
+{
+    float4 v[1024];
+};
+
+struct SSBOIns
+{
+    float4 v;
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u);
+
+struct spvDescriptorSetBuffer0
+{
+    device SSBO* o [[id(0)]];
+    constant UBO* v [[id(1)]];
+    spvDescriptor<constant UBOs *> vs [[id(2)]][1] /* unsized array hack */;
+};
+
+struct spvDescriptorSetBuffer1
+{
+    device SSBOIn* w [[id(0)]];
+    spvDescriptor<device SSBOIns *> ws [[id(1)]][1] /* unsized array hack */;
+};
+
+kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], const device spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
+{
+    spvDescriptorArray<constant UBOs*> vs {spvDescriptorSet0.vs};
+    spvDescriptorArray<device SSBOIns*> ws {spvDescriptorSet1.ws};
+
+    (*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = (*spvDescriptorSet0.v).v[gl_WorkGroupID.x];
+    (*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = vs[gl_WorkGroupID.x]->v;
+    (*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = (*spvDescriptorSet1.w).v[gl_WorkGroupID.x];
+    (*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = ws[gl_WorkGroupID.x]->v;
+}
+

third_party/spirv-cross/reference/opt/shaders-msl/comp/argument-buffers-runtime-array-buffer.rich-descriptor.argument.device-argument-buffer.argument-tier-1.msl2.comp

@@ -0,0 +1,109 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+template<typename T>
+struct spvDescriptor
+{
+    T value;
+};
+
+template<typename T>
+struct spvBufferDescriptor
+{
+    T value;
+    int length;
+    const device T& operator -> () const device
+    {
+        return value;
+    }
+    const device T& operator * () const device
+    {
+        return value;
+    }
+};
+
+template<typename T>
+struct spvDescriptorArray
+{
+    spvDescriptorArray(const device spvDescriptor<T>* ptr) : ptr(ptr)
+    {
+    }
+    const device T& operator [] (size_t i) const
+    {
+        return ptr[i].value;
+    }
+    const device spvDescriptor<T>* ptr;
+};
+
+template<typename T>
+struct spvDescriptorArray<device T*>
+{
+    spvDescriptorArray(const device spvBufferDescriptor<device T*>* ptr) : ptr(ptr)
+    {
+    }
+    const device T* operator [] (size_t i) const
+    {
+        return ptr[i].value;
+    }
+    const int length(int i) const
+    {
+        return ptr[i].length;
+    }
+    const device spvBufferDescriptor<device T*>* ptr;
+};
+
+struct SSBO
+{
+    float4 v[1];
+};
+
+struct UBO
+{
+    float4 v[1024];
+};
+
+struct UBOs
+{
+    float4 v;
+};
+
+struct SSBOIn
+{
+    float4 v[1024];
+};
+
+struct SSBOIns
+{
+    float4 v;
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u);
+
+struct spvDescriptorSetBuffer0
+{
+    device SSBO* o [[id(0)]];
+    constant UBO* v [[id(1)]];
+    spvDescriptor<constant UBOs *> vs [[id(2)]][1] /* unsized array hack */;
+};
+
+struct spvDescriptorSetBuffer1
+{
+    device SSBOIn* w [[id(0)]];
+    spvBufferDescriptor<device SSBOIns *> ws [[id(1)]][1] /* unsized array hack */;
+};
+
+kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], const device spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
+{
+    spvDescriptorArray<constant UBOs*> vs {spvDescriptorSet0.vs};
+    spvDescriptorArray<device SSBOIns*> ws {spvDescriptorSet1.ws};
+
+    (*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = (*spvDescriptorSet0.v).v[gl_WorkGroupID.x];
+    (*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = vs[gl_WorkGroupID.x]->v;
+    (*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = (*spvDescriptorSet1.w).v[gl_WorkGroupID.x];
+    (*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = ws[gl_WorkGroupID.x]->v;
+}
+

third_party/spirv-cross/reference/opt/shaders-msl/comp/argument-buffers-runtime-array.argument.device-argument-buffer.argument-tier-1.msl2.comp

@@ -0,0 +1,55 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+template<typename T>
+struct spvDescriptor
+{
+    T value;
+};
+
+template<typename T>
+struct spvDescriptorArray
+{
+    spvDescriptorArray(const device spvDescriptor<T>* ptr) : ptr(ptr)
+    {
+    }
+    const device T& operator [] (size_t i) const
+    {
+        return ptr[i].value;
+    }
+    const device spvDescriptor<T>* ptr;
+};
+
+struct SSBO
+{
+    float4 v[1];
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u);
+
+struct spvDescriptorSetBuffer0
+{
+    device SSBO* m_13 [[id(0)]];
+    texture2d<float> T [[id(1)]];
+    spvDescriptor<texture2d<float>> Ts [[id(2)]][1] /* unsized array hack */;
+};
+
+struct spvDescriptorSetBuffer1
+{
+    sampler S [[id(0)]];
+    spvDescriptor<sampler> Ss [[id(1)]][1] /* unsized array hack */;
+};
+
+kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], const device spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
+{
+    spvDescriptorArray<texture2d<float>> Ts {spvDescriptorSet0.Ts};
+    spvDescriptorArray<sampler> Ss {spvDescriptorSet1.Ss};
+
+    (*spvDescriptorSet0.m_13).v[gl_GlobalInvocationID.x] = spvDescriptorSet0.T.sample(spvDescriptorSet1.S, float2(0.5), level(0.0));
+    (*spvDescriptorSet0.m_13).v[gl_GlobalInvocationID.x] = Ts[gl_WorkGroupID.x].sample(Ss[gl_WorkGroupID.x], float2(0.5), level(0.0));
+}
+

third_party/spirv-cross/reference/opt/shaders-msl/comp/atomic-image.comp

@@ -0,0 +1,75 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+#pragma clang diagnostic ignored "-Wunused-variable"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+#include <metal_atomic>
+
+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
+{
+    uint u32;
+    int i32;
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
+
+kernel void main0(device SSBO& ssbo [[buffer(2)]], texture2d<uint> uImage [[texture(0)]], device atomic_uint* uImage_atomic [[buffer(0)]], texture2d<int, access::write> iImage [[texture(1)]], device atomic_int* iImage_atomic [[buffer(1)]])
+{
+    uint _19 = atomic_fetch_add_explicit((device atomic_uint*)&uImage_atomic[spvImage2DAtomicCoord(int2(1, 5), uImage)], 1u, memory_order_relaxed);
+    uint _27 = atomic_fetch_add_explicit((device atomic_uint*)&uImage_atomic[spvImage2DAtomicCoord(int2(1, 5), uImage)], 1u, memory_order_relaxed);
+    iImage.write(int4(int(_27)), uint2(int2(1, 6)));
+    uint _32 = atomic_fetch_or_explicit((device atomic_uint*)&uImage_atomic[spvImage2DAtomicCoord(int2(1, 5), uImage)], 1u, memory_order_relaxed);
+    uint _34 = atomic_fetch_xor_explicit((device atomic_uint*)&uImage_atomic[spvImage2DAtomicCoord(int2(1, 5), uImage)], 1u, memory_order_relaxed);
+    uint _36 = atomic_fetch_and_explicit((device atomic_uint*)&uImage_atomic[spvImage2DAtomicCoord(int2(1, 5), uImage)], 1u, memory_order_relaxed);
+    uint _38 = atomic_fetch_min_explicit((device atomic_uint*)&uImage_atomic[spvImage2DAtomicCoord(int2(1, 5), uImage)], 1u, memory_order_relaxed);
+    uint _40 = atomic_fetch_max_explicit((device atomic_uint*)&uImage_atomic[spvImage2DAtomicCoord(int2(1, 5), uImage)], 1u, memory_order_relaxed);
+    uint _44;
+    do
+    {
+        _44 = 10u;
+    } while (!atomic_compare_exchange_weak_explicit((device atomic_uint*)&uImage_atomic[spvImage2DAtomicCoord(int2(1, 5), uImage)], &_44, 2u, memory_order_relaxed, memory_order_relaxed) && _44 == 10u);
+    int _47 = atomic_fetch_add_explicit((device atomic_int*)&iImage_atomic[spvImage2DAtomicCoord(int2(1, 6), iImage)], 1, memory_order_relaxed);
+    int _49 = atomic_fetch_or_explicit((device atomic_int*)&iImage_atomic[spvImage2DAtomicCoord(int2(1, 6), iImage)], 1, memory_order_relaxed);
+    int _51 = atomic_fetch_xor_explicit((device atomic_int*)&iImage_atomic[spvImage2DAtomicCoord(int2(1, 6), iImage)], 1, memory_order_relaxed);
+    int _53 = atomic_fetch_and_explicit((device atomic_int*)&iImage_atomic[spvImage2DAtomicCoord(int2(1, 6), iImage)], 1, memory_order_relaxed);
+    int _55 = atomic_fetch_min_explicit((device atomic_int*)&iImage_atomic[spvImage2DAtomicCoord(int2(1, 6), iImage)], 1, memory_order_relaxed);
+    int _57 = atomic_fetch_max_explicit((device atomic_int*)&iImage_atomic[spvImage2DAtomicCoord(int2(1, 6), iImage)], 1, memory_order_relaxed);
+    int _61;
+    do
+    {
+        _61 = 10;
+    } while (!atomic_compare_exchange_weak_explicit((device atomic_int*)&iImage_atomic[spvImage2DAtomicCoord(int2(1, 5), iImage)], &_61, 2, memory_order_relaxed, memory_order_relaxed) && _61 == 10);
+    uint _68 = atomic_fetch_add_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
+    uint _70 = atomic_fetch_or_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
+    uint _72 = atomic_fetch_xor_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
+    uint _74 = atomic_fetch_and_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
+    uint _76 = atomic_fetch_min_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
+    uint _78 = atomic_fetch_max_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
+    uint _80 = atomic_exchange_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
+    uint _82;
+    do
+    {
+        _82 = 10u;
+    } while (!atomic_compare_exchange_weak_explicit((device atomic_uint*)&ssbo.u32, &_82, 2u, memory_order_relaxed, memory_order_relaxed) && _82 == 10u);
+    int _85 = atomic_fetch_add_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
+    int _87 = atomic_fetch_or_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
+    int _89 = atomic_fetch_xor_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
+    int _91 = atomic_fetch_and_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
+    int _93 = atomic_fetch_min_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
+    int _95 = atomic_fetch_max_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
+    int _97 = atomic_exchange_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
+    int _99;
+    do
+    {
+        _99 = 10;
+    } while (!atomic_compare_exchange_weak_explicit((device atomic_int*)&ssbo.i32, &_99, 2, memory_order_relaxed, memory_order_relaxed) && _99 == 10);
+}
+

third_party/spirv-cross/reference/opt/shaders-msl/comp/atomic-image.msl31.comp

@@ -0,0 +1,72 @@
+#pragma clang diagnostic ignored "-Wunused-variable"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+#include <metal_atomic>
+
+using namespace metal;
+
+struct SSBO
+{
+    uint u32;
+    int i32;
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
+
+kernel void main0(device SSBO& ssbo [[buffer(0)]], texture2d<uint, access::read_write> uImage [[texture(0)]], texture2d<int, access::read_write> iImage [[texture(1)]])
+{
+    uint _19 = uImage.atomic_fetch_add(uint2(int2(1, 5)), 1u).x;
+    uint _27 = uImage.atomic_fetch_add(uint2(int2(1, 5)), 1u).x;
+    iImage.write(int4(int(_27)), uint2(int2(1, 6)));
+    uint _32 = uImage.atomic_fetch_or(uint2(int2(1, 5)), 1u).x;
+    uint _34 = uImage.atomic_fetch_xor(uint2(int2(1, 5)), 1u).x;
+    uint _36 = uImage.atomic_fetch_and(uint2(int2(1, 5)), 1u).x;
+    uint _38 = uImage.atomic_fetch_min(uint2(int2(1, 5)), 1u).x;
+    uint _40 = uImage.atomic_fetch_max(uint2(int2(1, 5)), 1u).x;
+    uint _44;
+    uint4 _102;
+    do
+    {
+        _102.x = 10u;
+    } while (!uImage.atomic_compare_exchange_weak(uint2(int2(1, 5)), &_102, 2u) && _102.x == 10u);
+    _44 = _102.x;
+    int _47 = iImage.atomic_fetch_add(uint2(int2(1, 6)), 1).x;
+    int _49 = iImage.atomic_fetch_or(uint2(int2(1, 6)), 1).x;
+    int _51 = iImage.atomic_fetch_xor(uint2(int2(1, 6)), 1).x;
+    int _53 = iImage.atomic_fetch_and(uint2(int2(1, 6)), 1).x;
+    int _55 = iImage.atomic_fetch_min(uint2(int2(1, 6)), 1).x;
+    int _57 = iImage.atomic_fetch_max(uint2(int2(1, 6)), 1).x;
+    int _61;
+    int4 _104;
+    do
+    {
+        _104.x = 10;
+    } while (!iImage.atomic_compare_exchange_weak(uint2(int2(1, 5)), &_104, 2) && _104.x == 10);
+    _61 = _104.x;
+    uint _68 = atomic_fetch_add_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
+    uint _70 = atomic_fetch_or_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
+    uint _72 = atomic_fetch_xor_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
+    uint _74 = atomic_fetch_and_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
+    uint _76 = atomic_fetch_min_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
+    uint _78 = atomic_fetch_max_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
+    uint _80 = atomic_exchange_explicit((device atomic_uint*)&ssbo.u32, 1u, memory_order_relaxed);
+    uint _82;
+    do
+    {
+        _82 = 10u;
+    } while (!atomic_compare_exchange_weak_explicit((device atomic_uint*)&ssbo.u32, &_82, 2u, memory_order_relaxed, memory_order_relaxed) && _82 == 10u);
+    int _85 = atomic_fetch_add_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
+    int _87 = atomic_fetch_or_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
+    int _89 = atomic_fetch_xor_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
+    int _91 = atomic_fetch_and_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
+    int _93 = atomic_fetch_min_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
+    int _95 = atomic_fetch_max_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
+    int _97 = atomic_exchange_explicit((device atomic_int*)&ssbo.i32, 1, memory_order_relaxed);
+    int _99;
+    do
+    {
+        _99 = 10;
+    } while (!atomic_compare_exchange_weak_explicit((device atomic_int*)&ssbo.i32, &_99, 2, memory_order_relaxed, memory_order_relaxed) && _99 == 10);
+}
+

third_party/spirv-cross/reference/opt/shaders-msl/comp/buffer_device_address-packed-vec-and-cast-to-and-from-uvec2.msl23.comp

@@ -19,10 +19,7 @@ struct SSBO
 kernel void main0(constant UBO& _10 [[buffer(0)]])
 {
     (reinterpret_cast<device SSBO*>(as_type<ulong>(_10.b)))->a1 = float3(1.0, 2.0, 3.0);
-    uint2 _35 = as_type<uint2>(reinterpret_cast<ulong>(reinterpret_cast<device SSBO*>(as_type<ulong>(_10.b + uint2(32u)))));
-    uint2 v2 = _35;
-    device SSBO* _39 = reinterpret_cast<device SSBO*>(as_type<ulong>(_35));
-    float3 v3 = float3(_39->a1);
+    device SSBO* _39 = reinterpret_cast<device SSBO*>(as_type<ulong>(as_type<uint2>(reinterpret_cast<ulong>(reinterpret_cast<device SSBO*>(as_type<ulong>(_10.b + uint2(32u)))))));
     _39->a1 = float3(_39->a1) + float3(1.0);
 }
 

third_party/spirv-cross/reference/opt/shaders-msl/comp/buffer_device_address-recursive-struct-pointers.msl23.comp

@@ -32,65 +32,8 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
 
 kernel void main0(constant t24& u24 [[buffer(0)]], constant t35& u35 [[buffer(1)]], texture2d<uint, access::write> v295 [[texture(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
-    int v8 = 0;
-    int _30 = 0 | (u24.m0[0].x - 0);
-    v8 = _30;
-    int _44 = _30 | (u24.m0[u35.m0[1]].x - 1);
-    v8 = _44;
-    int _50 = _44 | (u24.m1 - 2);
-    v8 = _50;
-    int _60 = _50 | int(u24.m4[0u][0] - 3.0);
-    v8 = _60;
-    int _68 = _60 | int(u24.m4[1u][0] - 5.0);
-    v8 = _68;
-    int _75 = _68 | int(u24.m4[0u][1] - 4.0);
-    v8 = _75;
-    int _82 = _75 | int(u24.m4[1u][1] - 6.0);
-    v8 = _82;
-    int _92 = _82 | (((device t21*)u24.m2[0].x)->m0[0].x - 3);
-    v8 = _92;
-    int _101 = _92 | (((device t21*)u24.m2[0].x)->m0[u35.m0[1]].x - 4);
-    v8 = _101;
-    int _109 = _101 | (((device t21*)u24.m2[0].x)->m1 - 5);
-    v8 = _109;
-    int _118 = _109 | int(((device t21*)u24.m2[0].x)->m4[0u][0] - 6.0);
-    v8 = _118;
-    int _127 = _118 | int(((device t21*)u24.m2[0].x)->m4[1u][0] - 8.0);
-    v8 = _127;
-    int _136 = _127 | int(((device t21*)u24.m2[0].x)->m4[0u][1] - 7.0);
-    v8 = _136;
-    int _145 = _136 | int(((device t21*)u24.m2[0].x)->m4[1u][1] - 9.0);
-    v8 = _145;
-    int _155 = _145 | (((device t21*)u24.m2[u35.m0[1]].x)->m0[0].x - 6);
-    v8 = _155;
-    int _167 = _155 | (((device t21*)u24.m2[u35.m0[1]].x)->m0[u35.m0[1]].x - 7);
-    v8 = _167;
-    int _177 = _167 | (((device t21*)u24.m2[u35.m0[1]].x)->m1 - 8);
-    v8 = _177;
-    int _187 = _177 | int(((device t21*)u24.m2[u35.m0[1]].x)->m4[0u][0] - 9.0);
-    v8 = _187;
-    int _198 = _187 | int(((device t21*)u24.m2[u35.m0[1]].x)->m4[1u][0] - 11.0);
-    v8 = _198;
-    int _209 = _198 | int(((device t21*)u24.m2[u35.m0[1]].x)->m4[0u][1] - 10.0);
-    v8 = _209;
-    int _220 = _209 | int(((device t21*)u24.m2[u35.m0[1]].x)->m4[1u][1] - 12.0);
-    v8 = _220;
-    int _228 = _220 | (u24.m3->m0[0].x - 9);
-    v8 = _228;
-    int _238 = _228 | (u24.m3->m0[u35.m0[1]].x - 10);
-    v8 = _238;
-    int _246 = _238 | (u24.m3->m1 - 11);
-    v8 = _246;
-    int _254 = _246 | int(u24.m3->m4[0u][0] - 12.0);
-    v8 = _254;
-    int _263 = _254 | int(u24.m3->m4[1u][0] - 14.0);
-    v8 = _263;
-    int _272 = _263 | int(u24.m3->m4[0u][1] - 13.0);
-    v8 = _272;
-    int _281 = _272 | int(u24.m3->m4[1u][1] - 15.0);
-    v8 = _281;
-    uint4 _292 = select(uint4(1u, 0u, 0u, 1u), uint4(0u), bool4(_281 != 0));
-    uint4 v284 = _292;
-    v295.write(_292, uint2(int2(gl_GlobalInvocationID.xy)));
+    int _155 = ((((((((((((((0 | (u24.m0[0].x - 0)) | (u24.m0[u35.m0[1]].x - 1)) | (u24.m1 - 2)) | int(u24.m4[0u][0] - 3.0)) | int(u24.m4[1u][0] - 5.0)) | int(u24.m4[0u][1] - 4.0)) | int(u24.m4[1u][1] - 6.0)) | (((device t21*)u24.m2[0].x)->m0[0].x - 3)) | (((device t21*)u24.m2[0].x)->m0[u35.m0[1]].x - 4)) | (((device t21*)u24.m2[0].x)->m1 - 5)) | int(((device t21*)u24.m2[0].x)->m4[0u][0] - 6.0)) | int(((device t21*)u24.m2[0].x)->m4[1u][0] - 8.0)) | int(((device t21*)u24.m2[0].x)->m4[0u][1] - 7.0)) | int(((device t21*)u24.m2[0].x)->m4[1u][1] - 9.0)) | (((device t21*)u24.m2[u35.m0[1]].x)->m0[0].x - 6);
+    bool _289 = (((((((((((((_155 | (((device t21*)u24.m2[u35.m0[1]].x)->m0[u35.m0[1]].x - 7)) | (((device t21*)u24.m2[u35.m0[1]].x)->m1 - 8)) | int(((device t21*)u24.m2[u35.m0[1]].x)->m4[0u][0] - 9.0)) | int(((device t21*)u24.m2[u35.m0[1]].x)->m4[1u][0] - 11.0)) | int(((device t21*)u24.m2[u35.m0[1]].x)->m4[0u][1] - 10.0)) | int(((device t21*)u24.m2[u35.m0[1]].x)->m4[1u][1] - 12.0)) | (u24.m3->m0[0].x - 9)) | (u24.m3->m0[u35.m0[1]].x - 10)) | (u24.m3->m1 - 11)) | int(u24.m3->m4[0u][0] - 12.0)) | int(u24.m3->m4[1u][0] - 14.0)) | int(u24.m3->m4[0u][1] - 13.0)) | int(u24.m3->m4[1u][1] - 15.0)) != 0;
+    v295.write(select(uint4(1u, 0u, 0u, 1u), uint4(0u), bool4(_289)), uint2(int2(gl_GlobalInvocationID.xy)));
 }
 

third_party/spirv-cross/reference/opt/shaders-msl/comp/buffer_device_address.msl2.comp

@@ -27,30 +27,15 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(8u, 8u, 1u);
 kernel void main0(constant Registers& registers [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_NumWorkGroups [[threadgroups_per_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
 {
     uint2 local_offset = gl_GlobalInvocationID.xy;
-    uint _19 = local_offset.y;
-    uint _29 = local_offset.x;
-    uint _30 = ((_19 * 8u) * gl_NumWorkGroups.x) + _29;
-    uint local_index = _30;
-    uint slice = gl_WorkGroupID.z;
-    device Position* __restrict positions = registers.references->buffers[gl_WorkGroupID.z];
-    float _66 = float(gl_WorkGroupID.z);
-    float _70 = fract(fma(_66, 0.100000001490116119384765625, registers.fract_time));
+    float _70 = fract(fma(float(gl_WorkGroupID.z), 0.100000001490116119384765625, registers.fract_time));
     float _71 = 6.283125400543212890625 * _70;
-    float offset = _71;
-    float2 pos = float2(local_offset);
-    float _83 = sin(fma(2.2000000476837158203125, pos.x, _71));
-    pos.x = fma(0.20000000298023223876953125, _83, pos.x);
-    float _97 = sin(fma(2.25, pos.y, _70 * 12.56625080108642578125));
-    pos.y = fma(0.20000000298023223876953125, _97, pos.y);
-    float _111 = cos(fma(1.7999999523162841796875, pos.y, _70 * 18.849376678466796875));
-    pos.x = fma(0.20000000298023223876953125, _111, pos.x);
-    float _125 = cos(fma(2.849999904632568359375, pos.x, _70 * 25.1325016021728515625));
-    pos.y = fma(0.20000000298023223876953125, _125, pos.y);
-    float _133 = sin(_71);
-    pos.x = fma(0.5, _133, pos.x);
-    float _142 = sin(fma(6.283125400543212890625, _70, 0.300000011920928955078125));
-    pos.y = fma(0.5, _142, pos.y);
-    float2 _155 = float2(gl_NumWorkGroups.xy);
-    registers.references->buffers[gl_WorkGroupID.z]->positions[_30] = (pos / fma(float2(8.0), _155, float2(-1.0))) - float2(0.5);
+    float2 pos = float2(gl_GlobalInvocationID.xy);
+    pos.x = fma(0.20000000298023223876953125, sin(fma(2.2000000476837158203125, pos.x, _71)), pos.x);
+    pos.y = fma(0.20000000298023223876953125, sin(fma(2.25, pos.y, _70 * 12.56625080108642578125)), pos.y);
+    pos.x = fma(0.20000000298023223876953125, cos(fma(1.7999999523162841796875, pos.y, _70 * 18.849376678466796875)), pos.x);
+    pos.y = fma(0.20000000298023223876953125, cos(fma(2.849999904632568359375, pos.x, _70 * 25.1325016021728515625)), pos.y);
+    pos.x = fma(0.5, sin(_71), pos.x);
+    pos.y = fma(0.5, sin(fma(6.283125400543212890625, _70, 0.300000011920928955078125)), pos.y);
+    registers.references->buffers[gl_WorkGroupID.z]->positions[((local_offset.y * 8u) * gl_NumWorkGroups.x) + local_offset.x] = (pos / fma(float2(8.0), float2(gl_NumWorkGroups.xy), float2(-1.0))) - float2(0.5);
 }
 

third_party/spirv-cross/reference/opt/shaders-msl/comp/coherent-image-atomic.comp

@@ -0,0 +1,22 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+#pragma clang diagnostic ignored "-Wunused-variable"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+#include <metal_atomic>
+
+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)
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(256u, 1u, 1u);
+
+kernel void main0(texture2d<uint> rw_spd_global_atomic [[texture(0)]], volatile device atomic_uint* rw_spd_global_atomic_atomic [[buffer(0)]])
+{
+    uint _43 = atomic_fetch_add_explicit((volatile device atomic_uint*)&rw_spd_global_atomic_atomic[spvImage2DAtomicCoord(int2(0), rw_spd_global_atomic)], 1u, memory_order_relaxed);
+}
+

third_party/spirv-cross/reference/opt/shaders-msl/comp/coherent-image-atomic.msl2.argument.comp

@@ -0,0 +1,28 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+#pragma clang diagnostic ignored "-Wunused-variable"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+#include <metal_atomic>
+
+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)
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(256u, 1u, 1u);
+
+struct spvDescriptorSetBuffer1
+{
+    texture2d<uint> rw_spd_global_atomic [[id(0)]];
+    volatile device atomic_uint* rw_spd_global_atomic_atomic [[id(1)]];
+};
+
+kernel void main0(constant spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]])
+{
+    uint _43 = atomic_fetch_add_explicit((volatile device atomic_uint*)&spvDescriptorSet1.rw_spd_global_atomic_atomic[spvImage2DAtomicCoord(int2(0), spvDescriptorSet1.rw_spd_global_atomic)], 1u, memory_order_relaxed);
+}
+

third_party/spirv-cross/reference/opt/shaders-msl/comp/coherent-image-atomic.msl2.comp

@@ -0,0 +1,22 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+#pragma clang diagnostic ignored "-Wunused-variable"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+#include <metal_atomic>
+
+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)
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(256u, 1u, 1u);
+
+kernel void main0(texture2d<uint> rw_spd_global_atomic [[texture(0)]], volatile device atomic_uint* rw_spd_global_atomic_atomic [[buffer(0)]])
+{
+    uint _43 = atomic_fetch_add_explicit((volatile device atomic_uint*)&rw_spd_global_atomic_atomic[spvImage2DAtomicCoord(int2(0), rw_spd_global_atomic)], 1u, memory_order_relaxed);
+}
+

third_party/spirv-cross/reference/opt/shaders-msl/comp/coherent-image-atomic.msl31.argument.comp

@@ -0,0 +1,20 @@
+#pragma clang diagnostic ignored "-Wunused-variable"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+#include <metal_atomic>
+
+using namespace metal;
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(256u, 1u, 1u);
+
+struct spvDescriptorSetBuffer1
+{
+    texture2d<uint, access::read_write> rw_spd_global_atomic [[id(0)]];
+};
+
+kernel void main0(constant spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]])
+{
+    uint _43 = spvDescriptorSet1.rw_spd_global_atomic.atomic_fetch_add(uint2(int2(0)), 1u).x;
+}
+

third_party/spirv-cross/reference/opt/shaders-msl/comp/coherent-image-atomic.msl31.comp

@@ -0,0 +1,15 @@
+#pragma clang diagnostic ignored "-Wunused-variable"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+#include <metal_atomic>
+
+using namespace metal;
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(256u, 1u, 1u);
+
+kernel void main0(texture2d<uint, access::read_write> rw_spd_global_atomic [[texture(0)]])
+{
+    uint _43 = rw_spd_global_atomic.atomic_fetch_add(uint2(int2(0)), 1u).x;
+}
+

third_party/spirv-cross/reference/opt/shaders-msl/comp/complex-type-alias.comp

@@ -1,8 +1,49 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+#pragma clang diagnostic ignored "-Wmissing-braces"
+
 #include <metal_stdlib>
 #include <simd/simd.h>
 
 using namespace metal;
 
+template<typename T, size_t Num>
+struct spvUnsafeArray
+{
+    T elements[Num ? Num : 1];
+    
+    thread T& operator [] (size_t pos) thread
+    {
+        return elements[pos];
+    }
+    constexpr const thread T& operator [] (size_t pos) const thread
+    {
+        return elements[pos];
+    }
+    
+    device T& operator [] (size_t pos) device
+    {
+        return elements[pos];
+    }
+    constexpr const device T& operator [] (size_t pos) const device
+    {
+        return elements[pos];
+    }
+    
+    constexpr const constant T& operator [] (size_t pos) const constant
+    {
+        return elements[pos];
+    }
+    
+    threadgroup T& operator [] (size_t pos) threadgroup
+    {
+        return elements[pos];
+    }
+    constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
+    {
+        return elements[pos];
+    }
+};
+
 struct Foo0
 {
     float a;
@@ -44,7 +85,7 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(8u, 8u, 1u);
 
 kernel void main0(device SSBO& _53 [[buffer(0)]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
 {
-    threadgroup Foo2 coeffs[64];
+    threadgroup spvUnsafeArray<Foo2, 64> coeffs;
     coeffs[gl_LocalInvocationIndex] = Foo2{ Foo1{ Foo0{ 0.0 } }, 0.0 };
     threadgroup_barrier(mem_flags::mem_threadgroup);
     if (gl_LocalInvocationIndex == 0u)

third_party/spirv-cross/reference/opt/shaders-msl/comp/metal3_1_regression_patch.replace-recursive-inputs.msl3.comp

@@ -0,0 +1,27 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct recurs;
+
+struct recurs
+{
+    int m1;
+    device recurs* m2;
+};
+
+struct recurs_1
+{
+    int m1;
+    device recurs_1* m2;
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
+
+kernel void main0(device void* nums_vp [[buffer(0)]], texture2d<uint, access::write> tex [[texture(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+{
+    device auto& nums = *(device recurs*)nums_vp;
+    tex.write(uint4(uint((nums.m1 + nums.m2->m1) + nums.m2->m2->m1), 0u, 0u, 1u), uint2(int2(gl_GlobalInvocationID.xy)));
+}
+

third_party/spirv-cross/reference/opt/shaders-msl/comp/ray-query.spv14.vk.ios.msl24..invalid.comp

@@ -1,3 +1,5 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
 #include <metal_stdlib>
 #include <simd/simd.h>
 #if __METAL_VERSION__ >= 230
@@ -7,6 +9,30 @@ using namespace metal::raytracing;
 
 using namespace metal;
 
+intersection_params spvMakeIntersectionParams(uint flags)
+{
+    intersection_params ip;
+    if ((flags & 1) != 0)
+        ip.force_opacity(forced_opacity::opaque);
+    if ((flags & 2) != 0)
+        ip.force_opacity(forced_opacity::non_opaque);
+    if ((flags & 4) != 0)
+        ip.accept_any_intersection(true);
+    if ((flags & 16) != 0)
+        ip.set_triangle_cull_mode(triangle_cull_mode::back);
+    if ((flags & 32) != 0)
+        ip.set_triangle_cull_mode(triangle_cull_mode::front);
+    if ((flags & 64) != 0)
+        ip.set_opacity_cull_mode(opacity_cull_mode::opaque);
+    if ((flags & 128) != 0)
+        ip.set_opacity_cull_mode(opacity_cull_mode::non_opaque);
+    if ((flags & 256) != 0)
+        ip.set_geometry_cull_mode(geometry_cull_mode::triangle);
+    if ((flags & 512) != 0)
+        ip.set_geometry_cull_mode(geometry_cull_mode::bounding_box);
+    return ip;
+}
+
 struct Params
 {
     uint ray_flags;
@@ -22,9 +48,9 @@ struct Params
 kernel void main0(constant Params& _18 [[buffer(1)]], raytracing::acceleration_structure<raytracing::instancing> AS0 [[buffer(0)]], raytracing::acceleration_structure<raytracing::instancing> AS1 [[buffer(2)]])
 {
     raytracing::intersection_query<raytracing::instancing, raytracing::triangle_data> q;
-    q.reset(ray(_18.origin, _18.dir, _18.tmin, _18.tmax), AS0, intersection_params());
+    q.reset(ray(_18.origin, _18.dir, _18.tmin, _18.tmax), AS0, _18.cull_mask, spvMakeIntersectionParams(_18.ray_flags));
     raytracing::intersection_query<raytracing::instancing, raytracing::triangle_data> q2[2];
-    q2[1].reset(ray(_18.origin, _18.dir, _18.tmin, _18.tmax), AS1, intersection_params());
+    q2[1].reset(ray(_18.origin, _18.dir, _18.tmin, _18.tmax), AS1, _18.cull_mask, spvMakeIntersectionParams(_18.ray_flags));
     bool _63 = q.next();
     bool res = _63;
     q2[0].abort();

third_party/spirv-cross/reference/opt/shaders-msl/comp/shader_ballot.msl22.comp

@@ -0,0 +1,77 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+template<typename T>
+inline T spvSubgroupBroadcast(T value, ushort lane)
+{
+    return simd_broadcast(value, lane);
+}
+
+template<>
+inline bool spvSubgroupBroadcast(bool value, ushort lane)
+{
+    return !!simd_broadcast((ushort)value, lane);
+}
+
+template<uint N>
+inline vec<bool, N> spvSubgroupBroadcast(vec<bool, N> value, ushort lane)
+{
+    return (vec<bool, N>)simd_broadcast((vec<ushort, N>)value, lane);
+}
+
+template<typename T>
+inline T spvSubgroupBroadcastFirst(T value)
+{
+    return simd_broadcast_first(value);
+}
+
+template<>
+inline bool spvSubgroupBroadcastFirst(bool value)
+{
+    return !!simd_broadcast_first((ushort)value);
+}
+
+template<uint N>
+inline vec<bool, N> spvSubgroupBroadcastFirst(vec<bool, N> value)
+{
+    return (vec<bool, N>)simd_broadcast_first((vec<ushort, N>)value);
+}
+
+inline uint4 spvSubgroupBallot(bool value)
+{
+    simd_vote vote = simd_ballot(value);
+    // simd_ballot() returns a 64-bit integer-like object, but
+    // SPIR-V callers expect a uint4. We must convert.
+    // FIXME: This won't include higher bits if Apple ever supports
+    // 128 lanes in an SIMD-group.
+    return uint4(as_type<uint2>((simd_vote::vote_t)vote), 0, 0);
+}
+
+struct inputData
+{
+    float inputDataArray[1];
+};
+
+struct outputData
+{
+    float outputDataArray[1];
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u);
+
+kernel void main0(device inputData& _12 [[buffer(0)]], device outputData& _87 [[buffer(1)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint gl_SubgroupInvocationID [[thread_index_in_simdgroup]])
+{
+    uint4 gl_SubgroupLtMask = uint4(extract_bits(0xFFFFFFFF, 0, min(gl_SubgroupInvocationID, 32u)), extract_bits(0xFFFFFFFF, 0, (uint)max((int)gl_SubgroupInvocationID - 32, 0)), uint2(0));
+    bool _31 = _12.inputDataArray[gl_LocalInvocationID.x] > 0.0;
+    uint4 _52 = spvSubgroupBallot(_31);
+    uint4 _66 = uint4(int4(popcount(uint4(as_type<uint2>(as_type<ulong>(uint2(gl_SubgroupLtMask.xy))), 0u, 0u) & uint4(as_type<uint2>(as_type<ulong>(uint2(_52.xy))), 0u, 0u))));
+    if (_31)
+    {
+        _87.outputDataArray[_66.x + _66.y] = _12.inputDataArray[gl_LocalInvocationID.x];
+    }
+}
+

third_party/spirv-cross/reference/opt/shaders-msl/comp/shader_group_vote.msl21.comp

@@ -0,0 +1,37 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+template<typename T>
+inline bool spvSubgroupAllEqual(T value)
+{
+    return simd_all(all(value == simd_broadcast_first(value)));
+}
+
+template<>
+inline bool spvSubgroupAllEqual(bool value)
+{
+    return simd_all(value) || !simd_any(value);
+}
+
+template<uint N>
+inline bool spvSubgroupAllEqual(vec<bool, N> value)
+{
+    return simd_all(all(value == (vec<bool, N>)simd_broadcast_first((vec<ushort, N>)value)));
+}
+
+struct inputData
+{
+    float inputDataArray[1];
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u);
+
+kernel void main0(device inputData& _12 [[buffer(0)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
+{
+    bool _31 = _12.inputDataArray[gl_LocalInvocationID.x] > 0.0;
+}
+

third_party/spirv-cross/reference/opt/shaders-msl/comp/shared-array-of-arrays.comp

@@ -1,8 +1,49 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+#pragma clang diagnostic ignored "-Wmissing-braces"
+
 #include <metal_stdlib>
 #include <simd/simd.h>
 
 using namespace metal;
 
+template<typename T, size_t Num>
+struct spvUnsafeArray
+{
+    T elements[Num ? Num : 1];
+    
+    thread T& operator [] (size_t pos) thread
+    {
+        return elements[pos];
+    }
+    constexpr const thread T& operator [] (size_t pos) const thread
+    {
+        return elements[pos];
+    }
+    
+    device T& operator [] (size_t pos) device
+    {
+        return elements[pos];
+    }
+    constexpr const device T& operator [] (size_t pos) const device
+    {
+        return elements[pos];
+    }
+    
+    constexpr const constant T& operator [] (size_t pos) const constant
+    {
+        return elements[pos];
+    }
+    
+    threadgroup T& operator [] (size_t pos) threadgroup
+    {
+        return elements[pos];
+    }
+    constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
+    {
+        return elements[pos];
+    }
+};
+
 struct SSBO
 {
     float out_data[1];
@@ -12,7 +53,7 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(4u, 4u, 1u);
 
 kernel void main0(device SSBO& _67 [[buffer(0)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
-    threadgroup float foo[4][4];
+    threadgroup spvUnsafeArray<spvUnsafeArray<float, 4>, 4> foo;
     foo[gl_LocalInvocationID.x][gl_LocalInvocationID.y] = float(gl_LocalInvocationIndex);
     threadgroup_barrier(mem_flags::mem_threadgroup);
     _67.out_data[gl_GlobalInvocationID.x] = ((foo[gl_LocalInvocationID.x][0] + foo[gl_LocalInvocationID.x][1]) + foo[gl_LocalInvocationID.x][2]) + foo[gl_LocalInvocationID.x][3];

third_party/spirv-cross/reference/opt/shaders-msl/comp/shared.comp

@@ -1,8 +1,49 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+#pragma clang diagnostic ignored "-Wmissing-braces"
+
 #include <metal_stdlib>
 #include <simd/simd.h>
 
 using namespace metal;
 
+template<typename T, size_t Num>
+struct spvUnsafeArray
+{
+    T elements[Num ? Num : 1];
+    
+    thread T& operator [] (size_t pos) thread
+    {
+        return elements[pos];
+    }
+    constexpr const thread T& operator [] (size_t pos) const thread
+    {
+        return elements[pos];
+    }
+    
+    device T& operator [] (size_t pos) device
+    {
+        return elements[pos];
+    }
+    constexpr const device T& operator [] (size_t pos) const device
+    {
+        return elements[pos];
+    }
+    
+    constexpr const constant T& operator [] (size_t pos) const constant
+    {
+        return elements[pos];
+    }
+    
+    threadgroup T& operator [] (size_t pos) threadgroup
+    {
+        return elements[pos];
+    }
+    constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
+    {
+        return elements[pos];
+    }
+};
+
 struct SSBO
 {
     float in_data[1];
@@ -17,7 +58,7 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(4u, 1u, 1u);
 
 kernel void main0(const device SSBO& _22 [[buffer(0)]], device SSBO2& _44 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
 {
-    threadgroup float sShared[4];
+    threadgroup spvUnsafeArray<float, 4> sShared;
     sShared[gl_LocalInvocationIndex] = _22.in_data[gl_GlobalInvocationID.x];
     threadgroup_barrier(mem_flags::mem_threadgroup);
     _44.out_data[gl_GlobalInvocationID.x] = sShared[3u - gl_LocalInvocationIndex];

third_party/spirv-cross/reference/opt/shaders-msl/comp/threadgroup-boolean-workaround.comp

@@ -1,8 +1,49 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+#pragma clang diagnostic ignored "-Wmissing-braces"
+
 #include <metal_stdlib>
 #include <simd/simd.h>
 
 using namespace metal;
 
+template<typename T, size_t Num>
+struct spvUnsafeArray
+{
+    T elements[Num ? Num : 1];
+    
+    thread T& operator [] (size_t pos) thread
+    {
+        return elements[pos];
+    }
+    constexpr const thread T& operator [] (size_t pos) const thread
+    {
+        return elements[pos];
+    }
+    
+    device T& operator [] (size_t pos) device
+    {
+        return elements[pos];
+    }
+    constexpr const device T& operator [] (size_t pos) const device
+    {
+        return elements[pos];
+    }
+    
+    constexpr const constant T& operator [] (size_t pos) const constant
+    {
+        return elements[pos];
+    }
+    
+    threadgroup T& operator [] (size_t pos) threadgroup
+    {
+        return elements[pos];
+    }
+    constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
+    {
+        return elements[pos];
+    }
+};
+
 struct SSBO
 {
     float4 values[1];
@@ -12,7 +53,7 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(4u, 1u, 1u);
 
 kernel void main0(device SSBO& _23 [[buffer(0)]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
-    threadgroup short4 foo[4];
+    threadgroup spvUnsafeArray<short4, 4> foo;
     foo[gl_LocalInvocationIndex] = short4(_23.values[gl_GlobalInvocationID.x] != float4(10.0));
     threadgroup_barrier(mem_flags::mem_threadgroup);
     _23.values[gl_GlobalInvocationID.x] = select(float4(40.0), float4(30.0), bool4(foo[gl_LocalInvocationIndex ^ 3u]));

third_party/spirv-cross/reference/opt/shaders-msl/desktop-only/frag/image-ms.desktop.frag

@@ -5,6 +5,7 @@ using namespace metal;
 
 fragment void main0(texture2d_ms<float> uImageMS [[texture(0)]], texture2d_array<float, access::read_write> uImageArray [[texture(1)]], texture2d<float, access::write> uImage [[texture(2)]])
 {
+    uImageArray.fence();
     uImage.write(uImageMS.read(uint2(int2(1, 2)), 2), uint2(int2(2, 3)));
     uImageArray.write(uImageArray.read(uint2(int3(1, 2, 4).xy), uint(int3(1, 2, 4).z)), uint2(int3(2, 3, 7).xy), uint(int3(2, 3, 7).z));
 }

third_party/spirv-cross/reference/opt/shaders-msl/desktop-only/tesc/basic.desktop.sso.multi-patch.tesc

@@ -15,8 +15,8 @@ struct main0_patchOut
 
 struct main0_in
 {
-    uint3 m_86;
-    ushort2 m_90;
+    uint3 m_87;
+    ushort2 m_92;
     float4 gl_Position;
 };
 

third_party/spirv-cross/reference/opt/shaders-msl/frag/frag-demote-checks.discard-checks.msl23.frag

@@ -33,19 +33,19 @@ fragment main0_out main0(device foo_t& foo [[buffer(0)]], texture2d<uint, access
     {
         foo.x = 1.0;
     }
-    uint _91 = (!gl_HelperInvocation ? atomic_exchange_explicit((device atomic_uint*)&foo.y, 0u, memory_order_relaxed) : atomic_load_explicit((device atomic_uint*)&foo.y, memory_order_relaxed));
+    uint _91 = (!gl_HelperInvocation ? atomic_exchange_explicit((device atomic_uint*)&foo.y, 0u, memory_order_relaxed) : uint{});
     if (int(gl_FragCoord.x) == 3)
     {
         gl_HelperInvocation = true, discard_fragment();
     }
     int2 _101 = int2(gl_FragCoord.xy);
     (gl_HelperInvocation ? ((void)0) : bar.write(uint4(1u), uint2(_101)));
-    uint _103 = (!gl_HelperInvocation ? atomic_fetch_add_explicit((device atomic_uint*)&foo.y, 42u, memory_order_relaxed) : atomic_load_explicit((device atomic_uint*)&foo.y, memory_order_relaxed));
-    uint _108 = (!gl_HelperInvocation ? atomic_fetch_or_explicit((device atomic_uint*)&bar_atomic[spvImage2DAtomicCoord(_101, bar)], 62u, memory_order_relaxed) : atomic_load_explicit((device atomic_uint*)&bar_atomic[spvImage2DAtomicCoord(_101, bar)], memory_order_relaxed));
-    uint _110 = (!gl_HelperInvocation ? atomic_fetch_and_explicit((device atomic_uint*)&foo.y, 65535u, memory_order_relaxed) : atomic_load_explicit((device atomic_uint*)&foo.y, memory_order_relaxed));
-    uint _112 = (!gl_HelperInvocation ? atomic_fetch_xor_explicit((device atomic_uint*)&foo.y, 4294967040u, memory_order_relaxed) : atomic_load_explicit((device atomic_uint*)&foo.y, memory_order_relaxed));
-    uint _114 = (!gl_HelperInvocation ? atomic_fetch_min_explicit((device atomic_uint*)&foo.y, 1u, memory_order_relaxed) : atomic_load_explicit((device atomic_uint*)&foo.y, memory_order_relaxed));
-    uint _119 = (!gl_HelperInvocation ? atomic_fetch_max_explicit((device atomic_uint*)&bar_atomic[spvImage2DAtomicCoord(_101, bar)], 100u, memory_order_relaxed) : atomic_load_explicit((device atomic_uint*)&bar_atomic[spvImage2DAtomicCoord(_101, bar)], memory_order_relaxed));
+    uint _103 = (!gl_HelperInvocation ? atomic_fetch_add_explicit((device atomic_uint*)&foo.y, 42u, memory_order_relaxed) : uint{});
+    uint _108 = (!gl_HelperInvocation ? atomic_fetch_or_explicit((device atomic_uint*)&bar_atomic[spvImage2DAtomicCoord(_101, bar)], 62u, memory_order_relaxed) : uint{});
+    uint _110 = (!gl_HelperInvocation ? atomic_fetch_and_explicit((device atomic_uint*)&foo.y, 65535u, memory_order_relaxed) : uint{});
+    uint _112 = (!gl_HelperInvocation ? atomic_fetch_xor_explicit((device atomic_uint*)&foo.y, 4294967040u, memory_order_relaxed) : uint{});
+    uint _114 = (!gl_HelperInvocation ? atomic_fetch_min_explicit((device atomic_uint*)&foo.y, 1u, memory_order_relaxed) : uint{});
+    uint _119 = (!gl_HelperInvocation ? atomic_fetch_max_explicit((device atomic_uint*)&bar_atomic[spvImage2DAtomicCoord(_101, bar)], 100u, memory_order_relaxed) : uint{});
     uint _124;
     if (!gl_HelperInvocation)
     {
@@ -56,7 +56,7 @@ fragment main0_out main0(device foo_t& foo [[buffer(0)]], texture2d<uint, access
     }
     else
     {
-        _124 = atomic_load_explicit((device atomic_uint*)&bar_atomic[spvImage2DAtomicCoord(_101, bar)], memory_order_relaxed);
+        _124 = {};
     }
     bool _125 = gl_HelperInvocation;
     out.fragColor = float4(1.0, float(_125), 0.0, 1.0);

third_party/spirv-cross/reference/opt/shaders-msl/frag/frag-demote-checks.discard-checks.msl31.frag

@@ -0,0 +1,60 @@
+#pragma clang diagnostic ignored "-Wunused-variable"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+#include <metal_atomic>
+
+using namespace metal;
+
+struct foo_t
+{
+    float x;
+    uint y;
+};
+
+struct main0_out
+{
+    float4 fragColor [[color(0)]];
+};
+
+fragment main0_out main0(device foo_t& foo [[buffer(0)]], texture2d<uint, access::read_write> bar [[texture(0)]], float4 gl_FragCoord [[position]])
+{
+    main0_out out = {};
+    bool gl_HelperInvocation = {};
+    gl_HelperInvocation = simd_is_helper_thread();
+    if (!gl_HelperInvocation)
+    {
+        foo.x = 1.0;
+    }
+    uint _91 = (!gl_HelperInvocation ? atomic_exchange_explicit((device atomic_uint*)&foo.y, 0u, memory_order_relaxed) : uint{});
+    if (int(gl_FragCoord.x) == 3)
+    {
+        gl_HelperInvocation = true, discard_fragment();
+    }
+    int2 _101 = int2(gl_FragCoord.xy);
+    (gl_HelperInvocation ? ((void)0) : bar.write(uint4(1u), uint2(_101)));
+    uint _103 = (!gl_HelperInvocation ? atomic_fetch_add_explicit((device atomic_uint*)&foo.y, 42u, memory_order_relaxed) : uint{});
+    uint _108 = (!gl_HelperInvocation ? bar.atomic_fetch_or(uint2(_101), 62u).x : uint{});
+    uint _110 = (!gl_HelperInvocation ? atomic_fetch_and_explicit((device atomic_uint*)&foo.y, 65535u, memory_order_relaxed) : uint{});
+    uint _112 = (!gl_HelperInvocation ? atomic_fetch_xor_explicit((device atomic_uint*)&foo.y, 4294967040u, memory_order_relaxed) : uint{});
+    uint _114 = (!gl_HelperInvocation ? atomic_fetch_min_explicit((device atomic_uint*)&foo.y, 1u, memory_order_relaxed) : uint{});
+    uint _119 = (!gl_HelperInvocation ? bar.atomic_fetch_max(uint2(_101), 100u).x : uint{});
+    uint _124;
+    uint4 _135;
+    if (!gl_HelperInvocation)
+    {
+        do
+        {
+            _135.x = 100u;
+        } while (!bar.atomic_compare_exchange_weak(uint2(_101), &_135, 42u) && _135.x == 100u);
+        _124 = _135.x;
+    }
+    else
+    {
+        _124 = {};
+    }
+    bool _125 = gl_HelperInvocation;
+    out.fragColor = float4(1.0, float(_125), 0.0, 1.0);
+    return out;
+}
+

third_party/spirv-cross/reference/opt/shaders-msl/frag/frag-discard-checks-continue-block.discard-checks.msl31.frag

@@ -0,0 +1,40 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct foo
+{
+    int x;
+};
+
+struct main0_out
+{
+    float4 fragColor [[color(0)]];
+};
+
+fragment main0_out main0(device foo& _24 [[buffer(0)]], float4 gl_FragCoord [[position]])
+{
+    main0_out out = {};
+    bool gl_HelperInvocation = {};
+    gl_HelperInvocation = simd_is_helper_thread();
+    if (gl_FragCoord.y == 7.0)
+    {
+        gl_HelperInvocation = true, discard_fragment();
+    }
+    if (!gl_HelperInvocation)
+    {
+        _24.x = 0;
+    }
+    for (; float(_24.x) < gl_FragCoord.x; )
+    {
+        if (!gl_HelperInvocation)
+        {
+            _24.x++;
+        }
+        continue;
+    }
+    out.fragColor = float4(float(_24.x), 0.0, 0.0, 1.0);
+    return out;
+}
+

third_party/spirv-cross/reference/opt/shaders-msl/frag/frag-discard-checks.discard-checks.msl23.frag

@@ -33,19 +33,19 @@ fragment main0_out main0(device foo_t& foo [[buffer(0)]], texture2d<uint, access
     {
         foo.x = 1.0;
     }
-    uint _90 = (!gl_HelperInvocation ? atomic_exchange_explicit((device atomic_uint*)&foo.y, 0u, memory_order_relaxed) : atomic_load_explicit((device atomic_uint*)&foo.y, memory_order_relaxed));
+    uint _90 = (!gl_HelperInvocation ? atomic_exchange_explicit((device atomic_uint*)&foo.y, 0u, memory_order_relaxed) : uint{});
     if (int(gl_FragCoord.x) == 3)
     {
         gl_HelperInvocation = true, discard_fragment();
     }
     int2 _100 = int2(gl_FragCoord.xy);
     (gl_HelperInvocation ? ((void)0) : bar.write(uint4(1u), uint2(_100)));
-    uint _102 = (!gl_HelperInvocation ? atomic_fetch_add_explicit((device atomic_uint*)&foo.y, 42u, memory_order_relaxed) : atomic_load_explicit((device atomic_uint*)&foo.y, memory_order_relaxed));
-    uint _107 = (!gl_HelperInvocation ? atomic_fetch_or_explicit((device atomic_uint*)&bar_atomic[spvImage2DAtomicCoord(_100, bar)], 62u, memory_order_relaxed) : atomic_load_explicit((device atomic_uint*)&bar_atomic[spvImage2DAtomicCoord(_100, bar)], memory_order_relaxed));
-    uint _109 = (!gl_HelperInvocation ? atomic_fetch_and_explicit((device atomic_uint*)&foo.y, 65535u, memory_order_relaxed) : atomic_load_explicit((device atomic_uint*)&foo.y, memory_order_relaxed));
-    uint _111 = (!gl_HelperInvocation ? atomic_fetch_xor_explicit((device atomic_uint*)&foo.y, 4294967040u, memory_order_relaxed) : atomic_load_explicit((device atomic_uint*)&foo.y, memory_order_relaxed));
-    uint _113 = (!gl_HelperInvocation ? atomic_fetch_min_explicit((device atomic_uint*)&foo.y, 1u, memory_order_relaxed) : atomic_load_explicit((device atomic_uint*)&foo.y, memory_order_relaxed));
-    uint _118 = (!gl_HelperInvocation ? atomic_fetch_max_explicit((device atomic_uint*)&bar_atomic[spvImage2DAtomicCoord(_100, bar)], 100u, memory_order_relaxed) : atomic_load_explicit((device atomic_uint*)&bar_atomic[spvImage2DAtomicCoord(_100, bar)], memory_order_relaxed));
+    uint _102 = (!gl_HelperInvocation ? atomic_fetch_add_explicit((device atomic_uint*)&foo.y, 42u, memory_order_relaxed) : uint{});
+    uint _107 = (!gl_HelperInvocation ? atomic_fetch_or_explicit((device atomic_uint*)&bar_atomic[spvImage2DAtomicCoord(_100, bar)], 62u, memory_order_relaxed) : uint{});
+    uint _109 = (!gl_HelperInvocation ? atomic_fetch_and_explicit((device atomic_uint*)&foo.y, 65535u, memory_order_relaxed) : uint{});
+    uint _111 = (!gl_HelperInvocation ? atomic_fetch_xor_explicit((device atomic_uint*)&foo.y, 4294967040u, memory_order_relaxed) : uint{});
+    uint _113 = (!gl_HelperInvocation ? atomic_fetch_min_explicit((device atomic_uint*)&foo.y, 1u, memory_order_relaxed) : uint{});
+    uint _118 = (!gl_HelperInvocation ? atomic_fetch_max_explicit((device atomic_uint*)&bar_atomic[spvImage2DAtomicCoord(_100, bar)], 100u, memory_order_relaxed) : uint{});
     uint _123;
     if (!gl_HelperInvocation)
     {
@@ -56,7 +56,7 @@ fragment main0_out main0(device foo_t& foo [[buffer(0)]], texture2d<uint, access
     }
     else
     {
-        _123 = atomic_load_explicit((device atomic_uint*)&bar_atomic[spvImage2DAtomicCoord(_100, bar)], memory_order_relaxed);
+        _123 = {};
     }
     out.fragColor = float4(1.0, 0.0, 0.0, 1.0);
     return out;

third_party/spirv-cross/reference/opt/shaders-msl/frag/frag-discard-checks.discard-checks.msl31.frag

@@ -0,0 +1,59 @@
+#pragma clang diagnostic ignored "-Wunused-variable"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+#include <metal_atomic>
+
+using namespace metal;
+
+struct foo_t
+{
+    float x;
+    uint y;
+};
+
+struct main0_out
+{
+    float4 fragColor [[color(0)]];
+};
+
+fragment main0_out main0(device foo_t& foo [[buffer(0)]], texture2d<uint, access::read_write> bar [[texture(0)]], float4 gl_FragCoord [[position]])
+{
+    main0_out out = {};
+    bool gl_HelperInvocation = {};
+    gl_HelperInvocation = simd_is_helper_thread();
+    if (!gl_HelperInvocation)
+    {
+        foo.x = 1.0;
+    }
+    uint _90 = (!gl_HelperInvocation ? atomic_exchange_explicit((device atomic_uint*)&foo.y, 0u, memory_order_relaxed) : uint{});
+    if (int(gl_FragCoord.x) == 3)
+    {
+        gl_HelperInvocation = true, discard_fragment();
+    }
+    int2 _100 = int2(gl_FragCoord.xy);
+    (gl_HelperInvocation ? ((void)0) : bar.write(uint4(1u), uint2(_100)));
+    uint _102 = (!gl_HelperInvocation ? atomic_fetch_add_explicit((device atomic_uint*)&foo.y, 42u, memory_order_relaxed) : uint{});
+    uint _107 = (!gl_HelperInvocation ? bar.atomic_fetch_or(uint2(_100), 62u).x : uint{});
+    uint _109 = (!gl_HelperInvocation ? atomic_fetch_and_explicit((device atomic_uint*)&foo.y, 65535u, memory_order_relaxed) : uint{});
+    uint _111 = (!gl_HelperInvocation ? atomic_fetch_xor_explicit((device atomic_uint*)&foo.y, 4294967040u, memory_order_relaxed) : uint{});
+    uint _113 = (!gl_HelperInvocation ? atomic_fetch_min_explicit((device atomic_uint*)&foo.y, 1u, memory_order_relaxed) : uint{});
+    uint _118 = (!gl_HelperInvocation ? bar.atomic_fetch_max(uint2(_100), 100u).x : uint{});
+    uint _123;
+    uint4 _131;
+    if (!gl_HelperInvocation)
+    {
+        do
+        {
+            _131.x = 100u;
+        } while (!bar.atomic_compare_exchange_weak(uint2(_100), &_131, 42u) && _131.x == 100u);
+        _123 = _131.x;
+    }
+    else
+    {
+        _123 = {};
+    }
+    out.fragColor = float4(1.0, 0.0, 0.0, 1.0);
+    return out;
+}
+

third_party/spirv-cross/reference/opt/shaders-msl/frag/pixel-interlock-ordered.msl31.argument.frag

@@ -0,0 +1,45 @@
+#pragma clang diagnostic ignored "-Wunused-variable"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+#include <metal_atomic>
+
+using namespace metal;
+
+struct Buffer3
+{
+    int baz;
+};
+
+struct Buffer
+{
+    int foo;
+    uint bar;
+};
+
+struct Buffer2
+{
+    uint quux;
+};
+
+struct spvDescriptorSetBuffer0
+{
+    device Buffer3* m_9 [[id(0)]];
+    texture2d<float, access::write> img4 [[id(1)]];
+    texture2d<float, access::write> img [[id(2), raster_order_group(0)]];
+    texture2d<float> img3 [[id(3), raster_order_group(0)]];
+    texture2d<uint, access::read_write> img2 [[id(4), raster_order_group(0)]];
+    volatile device Buffer* m_42 [[id(5), raster_order_group(0)]];
+    device Buffer2* m_52 [[id(6), raster_order_group(0)]];
+};
+
+fragment void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]])
+{
+    (*spvDescriptorSet0.m_9).baz = 0;
+    spvDescriptorSet0.img4.write(float4(1.0, 0.0, 0.0, 1.0), uint2(int2(1)));
+    spvDescriptorSet0.img.write(spvDescriptorSet0.img3.read(uint2(int2(0))), uint2(int2(0)));
+    uint _39 = spvDescriptorSet0.img2.atomic_fetch_add(uint2(int2(0)), 1u).x;
+    (*spvDescriptorSet0.m_42).foo += 42;
+    uint _55 = atomic_fetch_and_explicit((volatile device atomic_uint*)&(*spvDescriptorSet0.m_42).bar, (*spvDescriptorSet0.m_52).quux, memory_order_relaxed);
+}
+

third_party/spirv-cross/reference/opt/shaders-msl/frag/pixel-interlock-ordered.msl31.frag

@@ -0,0 +1,34 @@
+#pragma clang diagnostic ignored "-Wunused-variable"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+#include <metal_atomic>
+
+using namespace metal;
+
+struct Buffer3
+{
+    int baz;
+};
+
+struct Buffer
+{
+    int foo;
+    uint bar;
+};
+
+struct Buffer2
+{
+    uint quux;
+};
+
+fragment void main0(device Buffer3& _9 [[buffer(0)]], volatile device Buffer& _42 [[buffer(1), raster_order_group(0)]], device Buffer2& _52 [[buffer(2), 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, access::read_write> img2 [[texture(3), raster_order_group(0)]])
+{
+    _9.baz = 0;
+    img4.write(float4(1.0, 0.0, 0.0, 1.0), uint2(int2(1)));
+    img.write(img3.read(uint2(int2(0))), uint2(int2(0)));
+    uint _39 = img2.atomic_fetch_add(uint2(int2(0)), 1u).x;
+    _42.foo += 42;
+    uint _55 = atomic_fetch_and_explicit((volatile device atomic_uint*)&_42.bar, _52.quux, memory_order_relaxed);
+}
+

third_party/spirv-cross/reference/opt/shaders-msl/frag/post-depth-coverage.ios.msl2.frag

@@ -11,7 +11,7 @@ struct main0_out
 [[ early_fragment_tests ]] fragment main0_out main0(uint gl_SampleMaskIn [[sample_mask, post_depth_coverage]])
 {
     main0_out out = {};
-    out.FragColor = float4(float(gl_SampleMaskIn));
+    out.FragColor = float4(float(int(gl_SampleMaskIn)));
     return out;
 }
 

third_party/spirv-cross/reference/opt/shaders-msl/frag/post-depth-coverage.msl23.frag

@@ -11,7 +11,7 @@ struct main0_out
 [[ early_fragment_tests ]] fragment main0_out main0(uint gl_SampleMaskIn [[sample_mask, post_depth_coverage]])
 {
     main0_out out = {};
-    out.FragColor = float4(float(gl_SampleMaskIn));
+    out.FragColor = float4(float(int(gl_SampleMaskIn)));
     return out;
 }
 

third_party/spirv-cross/reference/opt/shaders-msl/frag/ray-query-object-in-function.spv14.vk.msl24.frag

@@ -1,3 +1,5 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
 #include <metal_stdlib>
 #include <simd/simd.h>
 #if __METAL_VERSION__ >= 230
@@ -7,6 +9,30 @@ using namespace metal::raytracing;
 
 using namespace metal;
 
+intersection_params spvMakeIntersectionParams(uint flags)
+{
+    intersection_params ip;
+    if ((flags & 1) != 0)
+        ip.force_opacity(forced_opacity::opaque);
+    if ((flags & 2) != 0)
+        ip.force_opacity(forced_opacity::non_opaque);
+    if ((flags & 4) != 0)
+        ip.accept_any_intersection(true);
+    if ((flags & 16) != 0)
+        ip.set_triangle_cull_mode(triangle_cull_mode::back);
+    if ((flags & 32) != 0)
+        ip.set_triangle_cull_mode(triangle_cull_mode::front);
+    if ((flags & 64) != 0)
+        ip.set_opacity_cull_mode(opacity_cull_mode::opaque);
+    if ((flags & 128) != 0)
+        ip.set_opacity_cull_mode(opacity_cull_mode::non_opaque);
+    if ((flags & 256) != 0)
+        ip.set_geometry_cull_mode(geometry_cull_mode::triangle);
+    if ((flags & 512) != 0)
+        ip.set_geometry_cull_mode(geometry_cull_mode::bounding_box);
+    return ip;
+}
+
 struct main0_out
 {
     float4 outColor [[color(0)]];
@@ -21,7 +47,7 @@ fragment main0_out main0(main0_in in [[stage_in]], raytracing::acceleration_stru
 {
     main0_out out = {};
     raytracing::intersection_query<raytracing::instancing, raytracing::triangle_data> rayQuery;
-    rayQuery.reset(ray(float3((in.inPos.xy * 4.0) - float2(2.0), 1.0), float3(0.0, 0.0, -1.0), 0.001000000047497451305389404296875, 2.0), topLevelAS, intersection_params());
+    rayQuery.reset(ray(float3((in.inPos.xy * 4.0) - float2(2.0), 1.0), float3(0.0, 0.0, -1.0), 0.001000000047497451305389404296875, 2.0), topLevelAS, 255u, spvMakeIntersectionParams(4u));
     for (;;)
     {
         bool _88 = rayQuery.next();

third_party/spirv-cross/reference/opt/shaders-msl/frag/runtime_array_as_argument_buffer.msl3.argument-tier-1.rich-descriptor.frag

@@ -0,0 +1,159 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+#if __METAL_VERSION__ >= 230
+#include <metal_raytracing>
+using namespace metal::raytracing;
+#endif
+
+using namespace metal;
+
+intersection_params spvMakeIntersectionParams(uint flags)
+{
+    intersection_params ip;
+    if ((flags & 1) != 0)
+        ip.force_opacity(forced_opacity::opaque);
+    if ((flags & 2) != 0)
+        ip.force_opacity(forced_opacity::non_opaque);
+    if ((flags & 4) != 0)
+        ip.accept_any_intersection(true);
+    if ((flags & 16) != 0)
+        ip.set_triangle_cull_mode(triangle_cull_mode::back);
+    if ((flags & 32) != 0)
+        ip.set_triangle_cull_mode(triangle_cull_mode::front);
+    if ((flags & 64) != 0)
+        ip.set_opacity_cull_mode(opacity_cull_mode::opaque);
+    if ((flags & 128) != 0)
+        ip.set_opacity_cull_mode(opacity_cull_mode::non_opaque);
+    if ((flags & 256) != 0)
+        ip.set_geometry_cull_mode(geometry_cull_mode::triangle);
+    if ((flags & 512) != 0)
+        ip.set_geometry_cull_mode(geometry_cull_mode::bounding_box);
+    return ip;
+}
+
+template<typename T>
+struct spvDescriptor
+{
+    T value;
+};
+
+template<typename T>
+struct spvBufferDescriptor
+{
+    T value;
+    int length;
+    const device T& operator -> () const device
+    {
+        return value;
+    }
+    const device T& operator * () const device
+    {
+        return value;
+    }
+};
+
+template<typename T>
+struct spvDescriptorArray
+{
+    spvDescriptorArray(const device spvDescriptor<T>* ptr) : ptr(ptr)
+    {
+    }
+    const device T& operator [] (size_t i) const
+    {
+        return ptr[i].value;
+    }
+    const device spvDescriptor<T>* ptr;
+};
+
+template<typename T>
+struct spvDescriptorArray<device T*>
+{
+    spvDescriptorArray(const device spvBufferDescriptor<device T*>* ptr) : ptr(ptr)
+    {
+    }
+    const device T* operator [] (size_t i) const
+    {
+        return ptr[i].value;
+    }
+    const int length(int i) const
+    {
+        return ptr[i].length;
+    }
+    const device spvBufferDescriptor<device T*>* ptr;
+};
+
+struct Ssbo
+{
+    uint val;
+    uint data[1];
+};
+
+struct Ubo
+{
+    uint val;
+};
+
+struct main0_in
+{
+    uint inputId [[user(locn0)]];
+};
+
+fragment void main0(main0_in in [[stage_in]], const device spvBufferDescriptor<const device Ssbo*>* ssbo_ [[buffer(4)]], const device spvDescriptor<constant Ubo*>* ubo_ [[buffer(5)]], const device spvDescriptor<texture2d<float>>* smp_textures_ [[buffer(0)]], const device spvDescriptor<texture2d<float>>* textures_ [[buffer(2)]], const device spvDescriptor<texture2d<float>>* images_ [[buffer(6)]], const device spvDescriptor<sampler>* smp_texturesSmplr_ [[buffer(1)]], const device spvDescriptor<sampler>* smp_ [[buffer(3)]], const device spvDescriptor<raytracing::acceleration_structure<raytracing::instancing>>* tlas_ [[buffer(7)]])
+{
+    spvDescriptorArray<texture2d<float>> smp_textures {smp_textures_};
+    spvDescriptorArray<sampler> smp_texturesSmplr {smp_texturesSmplr_};
+    spvDescriptorArray<texture2d<float>> textures {textures_};
+    spvDescriptorArray<sampler> smp {smp_};
+    spvDescriptorArray<const device Ssbo*> ssbo {ssbo_};
+    spvDescriptorArray<constant Ubo*> ubo {ubo_};
+    spvDescriptorArray<texture2d<float>> images {images_};
+    spvDescriptorArray<raytracing::acceleration_structure<raytracing::instancing>> tlas {tlas_};
+
+    uint _231 = in.inputId;
+    raytracing::intersection_query<raytracing::instancing, raytracing::triangle_data> rayQuery;
+    raytracing::intersection_query<raytracing::instancing, raytracing::triangle_data> rayQuery_1;
+    if (smp_textures[_231].sample(smp_texturesSmplr[_231], float2(0.0), level(0.0)).w > 0.5)
+    {
+        discard_fragment();
+    }
+    uint _249 = in.inputId + 8u;
+    if (textures[_231].sample(smp[_249], float2(0.0), level(0.0)).w > 0.5)
+    {
+        discard_fragment();
+    }
+    if (ssbo[_231]->val == 2u)
+    {
+        discard_fragment();
+    }
+    if (int((ssbo.length(123) - 4) / 4) == 25)
+    {
+        discard_fragment();
+    }
+    if (ubo[_231]->val == 2u)
+    {
+        discard_fragment();
+    }
+    if (images[_231].read(uint2(int2(0))).w > 0.5)
+    {
+        discard_fragment();
+    }
+    rayQuery.reset(ray(float3(0.0), float3(1.0), 0.00999999977648258209228515625, 1.0), tlas[in.inputId], 255u, spvMakeIntersectionParams(0u));
+    bool _301 = rayQuery.next();
+    if (smp_textures[_231].sample(smp_texturesSmplr[_231], float2(0.0), level(0.0)).w > 0.5)
+    {
+        discard_fragment();
+    }
+    if (textures[_231].sample(smp[_231], float2(0.0), level(0.0)).w > 0.5)
+    {
+        discard_fragment();
+    }
+    if (images[_231].read(uint2(int2(0))).w > 0.5)
+    {
+        discard_fragment();
+    }
+    rayQuery_1.reset(ray(float3(0.0), float3(1.0), 0.00999999977648258209228515625, 1.0), tlas[in.inputId], 255u, spvMakeIntersectionParams(0u));
+    bool _336 = rayQuery_1.next();
+}
+

third_party/spirv-cross/reference/opt/shaders-msl/frag/runtime_array_as_argument_buffer_buf.msl3.argument-tier-1.rich-descriptor.frag

@@ -0,0 +1,68 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+template<typename T>
+struct spvBufferDescriptor
+{
+    T value;
+    int length;
+    const device T& operator -> () const device
+    {
+        return value;
+    }
+    const device T& operator * () const device
+    {
+        return value;
+    }
+};
+
+template<typename T>
+struct spvDescriptorArray;
+
+template<typename T>
+struct spvDescriptorArray<device T*>
+{
+    spvDescriptorArray(const device spvBufferDescriptor<device T*>* ptr) : ptr(ptr)
+    {
+    }
+    const device T* operator [] (size_t i) const
+    {
+        return ptr[i].value;
+    }
+    const int length(int i) const
+    {
+        return ptr[i].length;
+    }
+    const device spvBufferDescriptor<device T*>* ptr;
+};
+
+struct Ssbo
+{
+    uint val;
+    uint data[1];
+};
+
+struct main0_in
+{
+    uint inputId [[user(locn0)]];
+};
+
+fragment void main0(main0_in in [[stage_in]], const device spvBufferDescriptor<const device Ssbo*>* ssbo_ [[buffer(0)]])
+{
+    spvDescriptorArray<const device Ssbo*> ssbo {ssbo_};
+
+    uint _15 = in.inputId;
+    if (ssbo[_15]->val == 2u)
+    {
+        discard_fragment();
+    }
+    if (int((ssbo.length(123) - 4) / 4) == 25)
+    {
+        discard_fragment();
+    }
+}
+

third_party/spirv-cross/reference/opt/shaders-msl/frag/sample-mask-in-and-out.fixed-sample-mask.force-sample.frag

@@ -13,7 +13,7 @@ fragment main0_out main0(uint gl_SampleMaskIn [[sample_mask]], uint gl_SampleID
 {
     main0_out out = {};
     out.FragColor = float4(1.0);
-    out.gl_SampleMask = (gl_SampleMaskIn & 0x22 & (1 << gl_SampleID));
+    out.gl_SampleMask = int((gl_SampleMaskIn & 0x22 & (1 << gl_SampleID)));
     out.gl_SampleMask &= 0x22;
     return out;
 }

third_party/spirv-cross/reference/opt/shaders-msl/frag/sample-mask-in-and-out.fixed-sample-mask.frag

@@ -13,7 +13,7 @@ fragment main0_out main0(uint gl_SampleMaskIn [[sample_mask]])
 {
     main0_out out = {};
     out.FragColor = float4(1.0);
-    out.gl_SampleMask = (gl_SampleMaskIn & 0x22);
+    out.gl_SampleMask = int((gl_SampleMaskIn & 0x22));
     out.gl_SampleMask &= 0x22;
     return out;
 }

third_party/spirv-cross/reference/opt/shaders-msl/frag/sampler-cube-grad.agx-cube-grad.frag

@@ -0,0 +1,44 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+static inline gradientcube spvGradientCube(float3 P, float3 dPdx, float3 dPdy)
+{
+    // Major axis selection
+    float3 absP = abs(P);
+    bool xMajor = absP.x >= max(absP.y, absP.z);
+    bool yMajor = absP.y >= absP.z;
+    float3 Q = xMajor ? P.yzx : (yMajor ? P.xzy : P);
+    float3 dQdx = xMajor ? dPdx.yzx : (yMajor ? dPdx.xzy : dPdx);
+    float3 dQdy = xMajor ? dPdy.yzx : (yMajor ? dPdy.xzy : dPdy);
+
+    // Skip a couple of operations compared to usual projection
+    float4 d = float4(dQdx.xy, dQdy.xy) - (Q.xy / Q.z).xyxy * float4(dQdx.zz, dQdy.zz);
+
+    // Final swizzle to put the intermediate values into non-ignored components
+    // X major: X and Z
+    // Y major: X and Y
+    // Z major: Y and Z
+    return gradientcube(xMajor ? d.xxy : d.xyx, xMajor ? d.zzw : d.zwz);
+}
+
+struct main0_out
+{
+    float4 FragColor [[color(0)]];
+};
+
+struct main0_in
+{
+    float3 vTex [[user(locn0), flat]];
+};
+
+fragment main0_out main0(main0_in in [[stage_in]], texturecube<float> uSampler [[texture(0)]], sampler uSamplerSmplr [[sampler(0)]])
+{
+    main0_out out = {};
+    out.FragColor += uSampler.sample(uSamplerSmplr, in.vTex, spvGradientCube(in.vTex, float3(5.0), float3(8.0)));
+    return out;
+}
+

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

@@ -44,114 +44,6 @@ struct spvUnsafeArray
     }
 };
 
-template<typename T, uint A>
-inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
-{
-    for (uint i = 0; i < A; i++)
-    {
-        dst[i] = src[i];
-    }
-}
-
-template<typename T, uint A>
-inline void spvArrayCopyFromConstantToThreadGroup1(threadgroup T (&dst)[A], constant T (&src)[A])
-{
-    for (uint i = 0; i < A; i++)
-    {
-        dst[i] = src[i];
-    }
-}
-
-template<typename T, uint A>
-inline void spvArrayCopyFromStackToStack1(thread T (&dst)[A], thread const T (&src)[A])
-{
-    for (uint i = 0; i < A; i++)
-    {
-        dst[i] = src[i];
-    }
-}
-
-template<typename T, uint A>
-inline void spvArrayCopyFromStackToThreadGroup1(threadgroup T (&dst)[A], thread const T (&src)[A])
-{
-    for (uint i = 0; i < A; i++)
-    {
-        dst[i] = src[i];
-    }
-}
-
-template<typename T, uint A>
-inline void spvArrayCopyFromThreadGroupToStack1(thread T (&dst)[A], threadgroup const T (&src)[A])
-{
-    for (uint i = 0; i < A; i++)
-    {
-        dst[i] = src[i];
-    }
-}
-
-template<typename T, uint A>
-inline void spvArrayCopyFromThreadGroupToThreadGroup1(threadgroup T (&dst)[A], threadgroup const T (&src)[A])
-{
-    for (uint i = 0; i < A; i++)
-    {
-        dst[i] = src[i];
-    }
-}
-
-template<typename T, uint A>
-inline void spvArrayCopyFromDeviceToDevice1(device T (&dst)[A], device const T (&src)[A])
-{
-    for (uint i = 0; i < A; i++)
-    {
-        dst[i] = src[i];
-    }
-}
-
-template<typename T, uint A>
-inline void spvArrayCopyFromConstantToDevice1(device T (&dst)[A], constant T (&src)[A])
-{
-    for (uint i = 0; i < A; i++)
-    {
-        dst[i] = src[i];
-    }
-}
-
-template<typename T, uint A>
-inline void spvArrayCopyFromStackToDevice1(device T (&dst)[A], thread const T (&src)[A])
-{
-    for (uint i = 0; i < A; i++)
-    {
-        dst[i] = src[i];
-    }
-}
-
-template<typename T, uint A>
-inline void spvArrayCopyFromThreadGroupToDevice1(device T (&dst)[A], threadgroup const T (&src)[A])
-{
-    for (uint i = 0; i < A; i++)
-    {
-        dst[i] = src[i];
-    }
-}
-
-template<typename T, uint A>
-inline void spvArrayCopyFromDeviceToStack1(thread T (&dst)[A], device const T (&src)[A])
-{
-    for (uint i = 0; i < A; i++)
-    {
-        dst[i] = src[i];
-    }
-}
-
-template<typename T, uint A>
-inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device const T (&src)[A])
-{
-    for (uint i = 0; i < A; i++)
-    {
-        dst[i] = src[i];
-    }
-}
-
 struct main0_out
 {
     float4 gl_Position;
@@ -171,14 +63,14 @@ struct main0_in
 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];
-    threadgroup float4 spvStorageFoo[8][4][2];
-    threadgroup float4 (&Foo)[4][2] = spvStorageFoo[(gl_GlobalInvocationID.x / 4) % 8];
+    threadgroup spvUnsafeArray<spvUnsafeArray<float4, 2>, 4> spvStorageFoo[8];
+    threadgroup auto &Foo = spvStorageFoo[(gl_GlobalInvocationID.x / 4) % 8];
     device main0_patchOut& patchOut = spvPatchOut[gl_GlobalInvocationID.x / 4];
     device main0_in* gl_in = &spvIn[min(gl_GlobalInvocationID.x / 4, spvIndirectParams[1] - 1) * spvIndirectParams[0]];
     uint gl_InvocationID = gl_GlobalInvocationID.x % 4;
     uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 4, spvIndirectParams[1] - 1);
     gl_out[gl_InvocationID].gl_Position = float4(1.0);
-    spvArrayCopyFromDeviceToThreadGroup1(Foo[gl_InvocationID], gl_in[gl_InvocationID].iFoo.elements);
+    Foo[gl_InvocationID] = gl_in[gl_InvocationID].iFoo;
     if (gl_InvocationID == 0)
     {
         spvUnsafeArray<float4, 2> _56 = spvUnsafeArray<float4, 2>({ gl_in[0].ipFoo, gl_in[1].ipFoo });

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

@@ -44,114 +44,6 @@ struct spvUnsafeArray
     }
 };
 
-template<typename T, uint A>
-inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
-{
-    for (uint i = 0; i < A; i++)
-    {
-        dst[i] = src[i];
-    }
-}
-
-template<typename T, uint A>
-inline void spvArrayCopyFromConstantToThreadGroup1(threadgroup T (&dst)[A], constant T (&src)[A])
-{
-    for (uint i = 0; i < A; i++)
-    {
-        dst[i] = src[i];
-    }
-}
-
-template<typename T, uint A>
-inline void spvArrayCopyFromStackToStack1(thread T (&dst)[A], thread const T (&src)[A])
-{
-    for (uint i = 0; i < A; i++)
-    {
-        dst[i] = src[i];
-    }
-}
-
-template<typename T, uint A>
-inline void spvArrayCopyFromStackToThreadGroup1(threadgroup T (&dst)[A], thread const T (&src)[A])
-{
-    for (uint i = 0; i < A; i++)
-    {
-        dst[i] = src[i];
-    }
-}
-
-template<typename T, uint A>
-inline void spvArrayCopyFromThreadGroupToStack1(thread T (&dst)[A], threadgroup const T (&src)[A])
-{
-    for (uint i = 0; i < A; i++)
-    {
-        dst[i] = src[i];
-    }
-}
-
-template<typename T, uint A>
-inline void spvArrayCopyFromThreadGroupToThreadGroup1(threadgroup T (&dst)[A], threadgroup const T (&src)[A])
-{
-    for (uint i = 0; i < A; i++)
-    {
-        dst[i] = src[i];
-    }
-}
-
-template<typename T, uint A>
-inline void spvArrayCopyFromDeviceToDevice1(device T (&dst)[A], device const T (&src)[A])
-{
-    for (uint i = 0; i < A; i++)
-    {
-        dst[i] = src[i];
-    }
-}
-
-template<typename T, uint A>
-inline void spvArrayCopyFromConstantToDevice1(device T (&dst)[A], constant T (&src)[A])
-{
-    for (uint i = 0; i < A; i++)
-    {
-        dst[i] = src[i];
-    }
-}
-
-template<typename T, uint A>
-inline void spvArrayCopyFromStackToDevice1(device T (&dst)[A], thread const T (&src)[A])
-{
-    for (uint i = 0; i < A; i++)
-    {
-        dst[i] = src[i];
-    }
-}
-
-template<typename T, uint A>
-inline void spvArrayCopyFromThreadGroupToDevice1(device T (&dst)[A], threadgroup const T (&src)[A])
-{
-    for (uint i = 0; i < A; i++)
-    {
-        dst[i] = src[i];
-    }
-}
-
-template<typename T, uint A>
-inline void spvArrayCopyFromDeviceToStack1(thread T (&dst)[A], device const T (&src)[A])
-{
-    for (uint i = 0; i < A; i++)
-    {
-        dst[i] = src[i];
-    }
-}
-
-template<typename T, uint A>
-inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device const T (&src)[A])
-{
-    for (uint i = 0; i < A; i++)
-    {
-        dst[i] = src[i];
-    }
-}
-
 struct main0_out
 {
     float4 gl_Position;
@@ -171,7 +63,7 @@ struct main0_in
 
 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];
+    threadgroup spvUnsafeArray<spvUnsafeArray<float4, 2>, 4> Foo;
     device main0_out* gl_out = &spvOut[gl_PrimitiveID * 4];
     device main0_patchOut& patchOut = spvPatchOut[gl_PrimitiveID];
     if (gl_InvocationID < spvIndirectParams[0])
@@ -181,7 +73,7 @@ kernel void main0(main0_in in [[stage_in]], uint gl_InvocationID [[thread_index_
         return;
     gl_out[gl_InvocationID].gl_Position = float4(1.0);
     spvUnsafeArray<float4, 2> _38 = spvUnsafeArray<float4, 2>({ gl_in[gl_InvocationID].iFoo_0, gl_in[gl_InvocationID].iFoo_1 });
-    spvArrayCopyFromStackToThreadGroup1(Foo[gl_InvocationID], _38.elements);
+    Foo[gl_InvocationID] = _38;
     if (gl_InvocationID == 0)
     {
         spvUnsafeArray<float4, 2> _56 = spvUnsafeArray<float4, 2>({ gl_in[0].ipFoo, gl_in[1].ipFoo });

third_party/spirv-cross/reference/opt/shaders-msl/masking/write-outputs-block.mask-location-0.multi-patch.msl2.tesc

@@ -32,7 +32,7 @@ kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], devic
     device main0_out* gl_out = &spvOut[gl_GlobalInvocationID.x - gl_GlobalInvocationID.x % 4];
     device main0_patchOut& patchOut = spvPatchOut[gl_GlobalInvocationID.x / 4];
     threadgroup P spvStorage_11[8];
-    threadgroup P (&_11) = spvStorage_11[(gl_GlobalInvocationID.x / 4) % 8];
+    threadgroup auto &_11 = spvStorage_11[(gl_GlobalInvocationID.x / 4) % 8];
     uint gl_InvocationID = gl_GlobalInvocationID.x % 4;
     uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 4, spvIndirectParams[1] - 1);
     _11.a = 1.0;

third_party/spirv-cross/reference/opt/shaders-msl/masking/write-outputs-block.mask-location-1.msl2.tesc

@@ -1,8 +1,49 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+#pragma clang diagnostic ignored "-Wmissing-braces"
+
 #include <metal_stdlib>
 #include <simd/simd.h>
 
 using namespace metal;
 
+template<typename T, size_t Num>
+struct spvUnsafeArray
+{
+    T elements[Num ? Num : 1];
+    
+    thread T& operator [] (size_t pos) thread
+    {
+        return elements[pos];
+    }
+    constexpr const thread T& operator [] (size_t pos) const thread
+    {
+        return elements[pos];
+    }
+    
+    device T& operator [] (size_t pos) device
+    {
+        return elements[pos];
+    }
+    constexpr const device T& operator [] (size_t pos) const device
+    {
+        return elements[pos];
+    }
+    
+    constexpr const constant T& operator [] (size_t pos) const constant
+    {
+        return elements[pos];
+    }
+    
+    threadgroup T& operator [] (size_t pos) threadgroup
+    {
+        return elements[pos];
+    }
+    constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
+    {
+        return elements[pos];
+    }
+};
+
 struct P
 {
     float a;
@@ -29,7 +70,7 @@ struct main0_patchOut
 
 kernel void main0(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 C c[4];
+    threadgroup spvUnsafeArray<C, 4> c;
     device main0_out* gl_out = &spvOut[gl_PrimitiveID * 4];
     device main0_patchOut& patchOut = spvPatchOut[gl_PrimitiveID];
     patchOut.m_11_a = 1.0;

third_party/spirv-cross/reference/opt/shaders-msl/masking/write-outputs-block.mask-location-1.multi-patch.msl2.tesc

@@ -1,8 +1,49 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+#pragma clang diagnostic ignored "-Wmissing-braces"
+
 #include <metal_stdlib>
 #include <simd/simd.h>
 
 using namespace metal;
 
+template<typename T, size_t Num>
+struct spvUnsafeArray
+{
+    T elements[Num ? Num : 1];
+    
+    thread T& operator [] (size_t pos) thread
+    {
+        return elements[pos];
+    }
+    constexpr const thread T& operator [] (size_t pos) const thread
+    {
+        return elements[pos];
+    }
+    
+    device T& operator [] (size_t pos) device
+    {
+        return elements[pos];
+    }
+    constexpr const device T& operator [] (size_t pos) const device
+    {
+        return elements[pos];
+    }
+    
+    constexpr const constant T& operator [] (size_t pos) const constant
+    {
+        return elements[pos];
+    }
+    
+    threadgroup T& operator [] (size_t pos) threadgroup
+    {
+        return elements[pos];
+    }
+    constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
+    {
+        return elements[pos];
+    }
+};
+
 struct P
 {
     float a;
@@ -30,8 +71,8 @@ struct main0_patchOut
 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_out* gl_out = &spvOut[gl_GlobalInvocationID.x - gl_GlobalInvocationID.x % 4];
-    threadgroup C spvStoragec[8][4];
-    threadgroup C (&c)[4] = spvStoragec[(gl_GlobalInvocationID.x / 4) % 8];
+    threadgroup spvUnsafeArray<C, 4> spvStoragec[8];
+    threadgroup auto &c = spvStoragec[(gl_GlobalInvocationID.x / 4) % 8];
     device main0_patchOut& patchOut = spvPatchOut[gl_GlobalInvocationID.x / 4];
     uint gl_InvocationID = gl_GlobalInvocationID.x % 4;
     uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 4, spvIndirectParams[1] - 1);

third_party/spirv-cross/reference/opt/shaders-msl/masking/write-outputs.mask-location-0.msl2.tesc

@@ -1,8 +1,49 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+#pragma clang diagnostic ignored "-Wmissing-braces"
+
 #include <metal_stdlib>
 #include <simd/simd.h>
 
 using namespace metal;
 
+template<typename T, size_t Num>
+struct spvUnsafeArray
+{
+    T elements[Num ? Num : 1];
+    
+    thread T& operator [] (size_t pos) thread
+    {
+        return elements[pos];
+    }
+    constexpr const thread T& operator [] (size_t pos) const thread
+    {
+        return elements[pos];
+    }
+    
+    device T& operator [] (size_t pos) device
+    {
+        return elements[pos];
+    }
+    constexpr const device T& operator [] (size_t pos) const device
+    {
+        return elements[pos];
+    }
+    
+    constexpr const constant T& operator [] (size_t pos) const constant
+    {
+        return elements[pos];
+    }
+    
+    threadgroup T& operator [] (size_t pos) threadgroup
+    {
+        return elements[pos];
+    }
+    constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
+    {
+        return elements[pos];
+    }
+};
+
 struct main0_out
 {
     float4 gl_Position;
@@ -16,7 +57,7 @@ struct main0_patchOut
 
 kernel void main0(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 float4 v0[4];
+    threadgroup spvUnsafeArray<float4, 4> v0;
     device main0_out* gl_out = &spvOut[gl_PrimitiveID * 4];
     device main0_patchOut& patchOut = spvPatchOut[gl_PrimitiveID];
     v0[gl_InvocationID] = float4(1.0);

third_party/spirv-cross/reference/opt/shaders-msl/masking/write-outputs.mask-location-0.multi-patch.tesc

@@ -59,8 +59,8 @@ struct main0_patchOut
 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_out* gl_out = &spvOut[gl_GlobalInvocationID.x - gl_GlobalInvocationID.x % 4];
-    threadgroup float4 spvStoragev0[8][4];
-    threadgroup float4 (&v0)[4] = spvStoragev0[(gl_GlobalInvocationID.x / 4) % 8];
+    threadgroup spvUnsafeArray<float4, 4> spvStoragev0[8];
+    threadgroup auto &v0 = spvStoragev0[(gl_GlobalInvocationID.x / 4) % 8];
     device main0_patchOut& patchOut = spvPatchOut[gl_GlobalInvocationID.x / 4];
     uint gl_InvocationID = gl_GlobalInvocationID.x % 4;
     uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 4, spvIndirectParams[1] - 1);

third_party/spirv-cross/reference/opt/shaders-msl/masking/write-outputs.mask-location-0.tesc

@@ -58,7 +58,7 @@ struct main0_patchOut
 
 kernel void main0(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 float4 v0[4];
+    threadgroup spvUnsafeArray<float4, 4> v0;
     device main0_out* gl_out = &spvOut[gl_PrimitiveID * 4];
     device main0_patchOut& patchOut = spvPatchOut[gl_PrimitiveID];
     v0[gl_InvocationID] = float4(1.0);

third_party/spirv-cross/reference/opt/shaders-msl/masking/write-outputs.mask-location-1.multi-patch.tesc

@@ -1,8 +1,49 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+#pragma clang diagnostic ignored "-Wmissing-braces"
+
 #include <metal_stdlib>
 #include <simd/simd.h>
 
 using namespace metal;
 
+template<typename T, size_t Num>
+struct spvUnsafeArray
+{
+    T elements[Num ? Num : 1];
+    
+    thread T& operator [] (size_t pos) thread
+    {
+        return elements[pos];
+    }
+    constexpr const thread T& operator [] (size_t pos) const thread
+    {
+        return elements[pos];
+    }
+    
+    device T& operator [] (size_t pos) device
+    {
+        return elements[pos];
+    }
+    constexpr const device T& operator [] (size_t pos) const device
+    {
+        return elements[pos];
+    }
+    
+    constexpr const constant T& operator [] (size_t pos) const constant
+    {
+        return elements[pos];
+    }
+    
+    threadgroup T& operator [] (size_t pos) threadgroup
+    {
+        return elements[pos];
+    }
+    constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
+    {
+        return elements[pos];
+    }
+};
+
 struct main0_out
 {
     float4 v0;
@@ -19,8 +60,8 @@ kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], devic
 {
     device main0_out* gl_out = &spvOut[gl_GlobalInvocationID.x - gl_GlobalInvocationID.x % 4];
     device main0_patchOut& patchOut = spvPatchOut[gl_GlobalInvocationID.x / 4];
-    threadgroup float4 spvStoragev1[8][2];
-    threadgroup float4 (&v1)[2] = spvStoragev1[(gl_GlobalInvocationID.x / 4) % 8];
+    threadgroup spvUnsafeArray<float4, 2> spvStoragev1[8];
+    threadgroup auto &v1 = spvStoragev1[(gl_GlobalInvocationID.x / 4) % 8];
     uint gl_InvocationID = gl_GlobalInvocationID.x % 4;
     uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 4, spvIndirectParams[1] - 1);
     gl_out[gl_InvocationID].v0 = float4(1.0);

third_party/spirv-cross/reference/opt/shaders-msl/masking/write-outputs.mask-location-1.tesc

@@ -1,8 +1,49 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+#pragma clang diagnostic ignored "-Wmissing-braces"
+
 #include <metal_stdlib>
 #include <simd/simd.h>
 
 using namespace metal;
 
+template<typename T, size_t Num>
+struct spvUnsafeArray
+{
+    T elements[Num ? Num : 1];
+    
+    thread T& operator [] (size_t pos) thread
+    {
+        return elements[pos];
+    }
+    constexpr const thread T& operator [] (size_t pos) const thread
+    {
+        return elements[pos];
+    }
+    
+    device T& operator [] (size_t pos) device
+    {
+        return elements[pos];
+    }
+    constexpr const device T& operator [] (size_t pos) const device
+    {
+        return elements[pos];
+    }
+    
+    constexpr const constant T& operator [] (size_t pos) const constant
+    {
+        return elements[pos];
+    }
+    
+    threadgroup T& operator [] (size_t pos) threadgroup
+    {
+        return elements[pos];
+    }
+    constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
+    {
+        return elements[pos];
+    }
+};
+
 struct main0_out
 {
     float4 v0;
@@ -17,7 +58,7 @@ struct main0_patchOut
 
 kernel void main0(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 float4 v1[2];
+    threadgroup spvUnsafeArray<float4, 2> v1;
     device main0_out* gl_out = &spvOut[gl_PrimitiveID * 4];
     device main0_patchOut& patchOut = spvPatchOut[gl_PrimitiveID];
     gl_out[gl_InvocationID].v0 = float4(1.0);

third_party/spirv-cross/reference/opt/shaders-msl/masking/write-outputs.mask-point-size.multi-patch.tesc

@@ -67,8 +67,8 @@ struct main0_patchOut
 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_out* gl_out = &spvOut[gl_GlobalInvocationID.x - gl_GlobalInvocationID.x % 4];
-    threadgroup gl_PerVertex spvStoragegl_out_masked[8][4];
-    threadgroup gl_PerVertex (&gl_out_masked)[4] = spvStoragegl_out_masked[(gl_GlobalInvocationID.x / 4) % 8];
+    threadgroup spvUnsafeArray<gl_PerVertex, 4> spvStoragegl_out_masked[8];
+    threadgroup auto &gl_out_masked = spvStoragegl_out_masked[(gl_GlobalInvocationID.x / 4) % 8];
     device main0_patchOut& patchOut = spvPatchOut[gl_GlobalInvocationID.x / 4];
     uint gl_InvocationID = gl_GlobalInvocationID.x % 4;
     uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 4, spvIndirectParams[1] - 1);

third_party/spirv-cross/reference/opt/shaders-msl/masking/write-outputs.mask-point-size.tesc

@@ -66,7 +66,7 @@ struct main0_patchOut
 
 kernel void main0(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 gl_PerVertex gl_out_masked[4];
+    threadgroup spvUnsafeArray<gl_PerVertex, 4> gl_out_masked;
     device main0_out* gl_out = &spvOut[gl_PrimitiveID * 4];
     device main0_patchOut& patchOut = spvPatchOut[gl_PrimitiveID];
     gl_out[gl_InvocationID].v0 = float4(1.0);

third_party/spirv-cross/reference/opt/shaders-msl/masking/write-outputs.mask-position.multi-patch.tesc

@@ -67,8 +67,8 @@ struct main0_patchOut
 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_out* gl_out = &spvOut[gl_GlobalInvocationID.x - gl_GlobalInvocationID.x % 4];
-    threadgroup gl_PerVertex spvStoragegl_out_masked[8][4];
-    threadgroup gl_PerVertex (&gl_out_masked)[4] = spvStoragegl_out_masked[(gl_GlobalInvocationID.x / 4) % 8];
+    threadgroup spvUnsafeArray<gl_PerVertex, 4> spvStoragegl_out_masked[8];
+    threadgroup auto &gl_out_masked = spvStoragegl_out_masked[(gl_GlobalInvocationID.x / 4) % 8];
     device main0_patchOut& patchOut = spvPatchOut[gl_GlobalInvocationID.x / 4];
     uint gl_InvocationID = gl_GlobalInvocationID.x % 4;
     uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 4, spvIndirectParams[1] - 1);

third_party/spirv-cross/reference/opt/shaders-msl/masking/write-outputs.mask-position.tesc

@@ -66,7 +66,7 @@ struct main0_patchOut
 
 kernel void main0(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 gl_PerVertex gl_out_masked[4];
+    threadgroup spvUnsafeArray<gl_PerVertex, 4> gl_out_masked;
     device main0_out* gl_out = &spvOut[gl_PrimitiveID * 4];
     device main0_patchOut& patchOut = spvPatchOut[gl_PrimitiveID];
     gl_out[gl_InvocationID].v0 = float4(1.0);

third_party/spirv-cross/reference/opt/shaders-msl/tesc/arrayed-block-io.multi-patch.tesc

@@ -68,7 +68,7 @@ struct main0_patchOut
 struct main0_in
 {
     float3 in_tc_attr;
-    ushort2 m_196;
+    ushort2 m_197;
 };
 
 kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], constant uint* spvIndirectParams [[buffer(29)]], device main0_patchOut* spvPatchOut [[buffer(27)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], device main0_in* spvIn [[buffer(22)]])

third_party/spirv-cross/reference/opt/shaders-msl/tesc/load-control-point-array.multi-patch.tesc

@@ -11,7 +11,7 @@ struct main0_out
 struct main0_in
 {
     float4 vInputs;
-    ushort2 m_44;
+    ushort2 m_45;
 };
 
 kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], device main0_out* spvOut [[buffer(28)]], constant uint* spvIndirectParams [[buffer(29)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], device main0_in* spvIn [[buffer(22)]])

third_party/spirv-cross/reference/opt/shaders-msl/tesc/matrix-output.multi-patch.tesc

@@ -13,7 +13,7 @@ struct main0_out
 struct main0_in
 {
     float3 in_tc_attr;
-    ushort2 m_104;
+    ushort2 m_105;
 };
 
 kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], device main0_out* spvOut [[buffer(28)]], constant uint* spvIndirectParams [[buffer(29)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], device main0_in* spvIn [[buffer(22)]])

third_party/spirv-cross/reference/opt/shaders-msl/tesc/reload-tess-level.multi-patch.tesc

@@ -10,8 +10,8 @@ struct main0_out
 
 struct main0_in
 {
-    uint3 m_82;
-    ushort2 m_86;
+    uint3 m_83;
+    ushort2 m_88;
     float4 gl_Position;
 };