Updated spirv-cross.

3rdparty/spirv-cross/main.cpp

@@ -516,8 +516,10 @@ struct CLIArguments
 	bool msl_texture_buffer_native = false;
 	bool msl_multiview = false;
 	bool msl_view_index_from_device_index = false;
+	bool msl_dispatch_base = false;
 	bool glsl_emit_push_constant_as_ubo = false;
 	bool glsl_emit_ubo_as_plain_uniforms = false;
+	bool vulkan_glsl_disable_ext_samplerless_texture_functions = false;
 	bool emit_line_directives = false;
 	SmallVector<uint32_t> msl_discrete_descriptor_sets;
 	SmallVector<PLSArg> pls_in;
@@ -584,6 +586,7 @@ static void print_help()
 	                "\t[--cpp-interface-name <name>]\n"
 	                "\t[--glsl-emit-push-constant-as-ubo]\n"
 	                "\t[--glsl-emit-ubo-as-plain-uniforms]\n"
+	                "\t[--vulkan-glsl-disable-ext-samplerless-texture-functions]\n"
 	                "\t[--msl]\n"
 	                "\t[--msl-version <MMmmpp>]\n"
 	                "\t[--msl-capture-output]\n"
@@ -596,6 +599,7 @@ static void print_help()
 	                "\t[--msl-discrete-descriptor-set <index>]\n"
 	                "\t[--msl-multiview]\n"
 	                "\t[--msl-view-index-from-device-index]\n"
+	                "\t[--msl-dispatch-base]\n"
 	                "\t[--hlsl]\n"
 	                "\t[--reflect]\n"
 	                "\t[--shader-model]\n"
@@ -756,6 +760,7 @@ static string compile_iteration(const CLIArguments &args, std::vector<uint32_t>
 		msl_opts.texture_buffer_native = args.msl_texture_buffer_native;
 		msl_opts.multiview = args.msl_multiview;
 		msl_opts.view_index_from_device_index = args.msl_view_index_from_device_index;
+		msl_opts.dispatch_base = args.msl_dispatch_base;
 		msl_comp->set_msl_options(msl_opts);
 		for (auto &v : args.msl_discrete_descriptor_sets)
 			msl_comp->add_discrete_descriptor_set(v);
@@ -765,7 +770,7 @@ static string compile_iteration(const CLIArguments &args, std::vector<uint32_t>
 	else
 	{
 		combined_image_samplers = !args.vulkan_semantics;
-		if (!args.vulkan_semantics)
+		if (!args.vulkan_semantics || args.vulkan_glsl_disable_ext_samplerless_texture_functions)
 			build_dummy_sampler = true;
 		compiler.reset(new CompilerGLSL(move(spirv_parser.get_parsed_ir())));
 	}
@@ -1058,6 +1063,8 @@ static int main_inner(int argc, char *argv[])
 	cbs.add("--metal", [&args](CLIParser &) { args.msl = true; }); // Legacy compatibility
 	cbs.add("--glsl-emit-push-constant-as-ubo", [&args](CLIParser &) { args.glsl_emit_push_constant_as_ubo = true; });
 	cbs.add("--glsl-emit-ubo-as-plain-uniforms", [&args](CLIParser &) { args.glsl_emit_ubo_as_plain_uniforms = true; });
+	cbs.add("--vulkan-glsl-disable-ext-samplerless-texture-functions",
+	        [&args](CLIParser &) { args.vulkan_glsl_disable_ext_samplerless_texture_functions = true; });
 	cbs.add("--msl", [&args](CLIParser &) { args.msl = true; });
 	cbs.add("--hlsl", [&args](CLIParser &) { args.hlsl = true; });
 	cbs.add("--hlsl-enable-compat", [&args](CLIParser &) { args.hlsl_compat = true; });
@@ -1078,6 +1085,7 @@ static int main_inner(int argc, char *argv[])
 	cbs.add("--msl-multiview", [&args](CLIParser &) { args.msl_multiview = true; });
 	cbs.add("--msl-view-index-from-device-index",
 	        [&args](CLIParser &) { args.msl_view_index_from_device_index = true; });
+	cbs.add("--msl-dispatch-base", [&args](CLIParser &) { args.msl_dispatch_base = 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();

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

@@ -14,14 +14,12 @@ struct SPIRV_Cross_Output
 #line 8 "test.frag"
 void frag_main()
 {
-    float _80;
 #line 8 "test.frag"
     FragColor = 1.0f;
 #line 9 "test.frag"
     FragColor = 2.0f;
 #line 10 "test.frag"
-    _80 = vColor;
-    if (_80 < 0.0f)
+    if (vColor < 0.0f)
     {
 #line 12 "test.frag"
         FragColor = 3.0f;
@@ -31,16 +29,16 @@ void frag_main()
 #line 16 "test.frag"
         FragColor = 4.0f;
     }
-    for (int _126 = 0; float(_126) < (40.0f + _80); )
+    for (int _126 = 0; float(_126) < (40.0f + vColor); )
     {
 #line 21 "test.frag"
         FragColor += 0.20000000298023223876953125f;
 #line 22 "test.frag"
         FragColor += 0.300000011920928955078125f;
-        _126 += (int(_80) + 5);
+        _126 += (int(vColor) + 5);
         continue;
     }
-    switch (int(_80))
+    switch (int(vColor))
     {
         case 0:
         {
@@ -66,7 +64,7 @@ void frag_main()
     }
     for (;;)
     {
-        FragColor += (10.0f + _80);
+        FragColor += (10.0f + vColor);
 #line 43 "test.frag"
         if (FragColor < 100.0f)
         {

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

@@ -17,14 +17,12 @@ struct main0_in
 fragment main0_out main0(main0_in in [[stage_in]])
 {
     main0_out out = {};
-    float _80;
 #line 8 "test.frag"
     out.FragColor = 1.0;
 #line 9 "test.frag"
     out.FragColor = 2.0;
 #line 10 "test.frag"
-    _80 = in.vColor;
-    if (_80 < 0.0)
+    if (in.vColor < 0.0)
     {
 #line 12 "test.frag"
         out.FragColor = 3.0;
@@ -34,16 +32,16 @@ fragment main0_out main0(main0_in in [[stage_in]])
 #line 16 "test.frag"
         out.FragColor = 4.0;
     }
-    for (int _126 = 0; float(_126) < (40.0 + _80); )
+    for (int _126 = 0; float(_126) < (40.0 + in.vColor); )
     {
 #line 21 "test.frag"
         out.FragColor += 0.20000000298023223876953125;
 #line 22 "test.frag"
         out.FragColor += 0.300000011920928955078125;
-        _126 += (int(_80) + 5);
+        _126 += (int(in.vColor) + 5);
         continue;
     }
-    switch (int(_80))
+    switch (int(in.vColor))
     {
         case 0:
         {
@@ -69,7 +67,7 @@ fragment main0_out main0(main0_in in [[stage_in]])
     }
     for (;;)
     {
-        out.FragColor += (10.0 + _80);
+        out.FragColor += (10.0 + in.vColor);
 #line 43 "test.frag"
         if (out.FragColor < 100.0)
         {

3rdparty/spirv-cross/reference/opt/shaders-msl/comp/basic.dispatchbase.comp

@@ -0,0 +1,38 @@
+#pragma clang diagnostic ignored "-Wunused-variable"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+#include <metal_atomic>
+
+using namespace metal;
+
+struct SSBO
+{
+    float4 in_data[1];
+};
+
+struct SSBO2
+{
+    float4 out_data[1];
+};
+
+struct SSBO3
+{
+    uint counter;
+};
+
+constant uint _59_tmp [[function_constant(10)]];
+constant uint _59 = is_function_constant_defined(_59_tmp) ? _59_tmp : 1u;
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(_59, 1u, 1u);
+
+kernel void main0(const device SSBO& _27 [[buffer(0)]], device SSBO2& _49 [[buffer(1)]], device SSBO3& _52 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 spvDispatchBase [[grid_origin]])
+{
+    gl_GlobalInvocationID += spvDispatchBase * gl_WorkGroupSize;
+    float4 _33 = _27.in_data[gl_GlobalInvocationID.x];
+    if (dot(_33, float4(1.0, 5.0, 6.0, 2.0)) > 8.19999980926513671875)
+    {
+        uint _56 = atomic_fetch_add_explicit((device atomic_uint*)&_52.counter, 1u, memory_order_relaxed);
+        _49.out_data[_56] = _33;
+    }
+}
+

3rdparty/spirv-cross/reference/opt/shaders-msl/comp/basic.dispatchbase.msl11.comp

@@ -0,0 +1,34 @@
+#pragma clang diagnostic ignored "-Wunused-variable"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+#include <metal_atomic>
+
+using namespace metal;
+
+struct SSBO
+{
+    float4 in_data[1];
+};
+
+struct SSBO2
+{
+    float4 out_data[1];
+};
+
+struct SSBO3
+{
+    uint counter;
+};
+
+kernel void main0(constant uint3& spvDispatchBase [[buffer(29)]], const device SSBO& _27 [[buffer(0)]], device SSBO2& _49 [[buffer(1)]], device SSBO3& _52 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+{
+    gl_GlobalInvocationID += spvDispatchBase * uint3(1, 1, 1);
+    float4 _33 = _27.in_data[gl_GlobalInvocationID.x];
+    if (dot(_33, float4(1.0, 5.0, 6.0, 2.0)) > 8.19999980926513671875)
+    {
+        uint _56 = atomic_fetch_add_explicit((device atomic_uint*)&_52.counter, 1u, memory_order_relaxed);
+        _49.out_data[_56] = _33;
+    }
+}
+

3rdparty/spirv-cross/reference/opt/shaders/asm/comp/phi-temporary-copy-loop-variable.asm.comp

@@ -14,10 +14,6 @@ void main()
         {
             break;
         }
-        else
-        {
-            continue;
-        }
         continue;
     }
     imageStore(outImageTexture, ivec2(gl_GlobalInvocationID.xy), vec4(float(_30 - 1), float(_30), 1.0, 1.0));

3rdparty/spirv-cross/reference/opt/shaders/asm/frag/image-fetch-no-sampler.no-samplerless.asm.vk.frag

@@ -0,0 +1,13 @@
+#version 450
+
+uniform sampler2D SPIRV_Cross_CombinedSampledImageSPIRV_Cross_DummySampler;
+uniform sampler2D SPIRV_Cross_CombinedSampledImageSampler;
+
+layout(location = 0) out vec4 _entryPointOutput;
+
+void main()
+{
+    ivec2 _152 = ivec3(int(gl_FragCoord.x * 1280.0), int(gl_FragCoord.y * 720.0), 0).xy;
+    _entryPointOutput = ((texelFetch(SPIRV_Cross_CombinedSampledImageSPIRV_Cross_DummySampler, _152, 0) + texelFetch(SPIRV_Cross_CombinedSampledImageSPIRV_Cross_DummySampler, _152, 0)) + texture(SPIRV_Cross_CombinedSampledImageSampler, gl_FragCoord.xy)) + texture(SPIRV_Cross_CombinedSampledImageSampler, gl_FragCoord.xy);
+}
+

3rdparty/spirv-cross/reference/opt/shaders/asm/frag/image-fetch-no-sampler.no-samplerless.asm.vk.frag.vk

@@ -0,0 +1,14 @@
+#version 450
+
+layout(set = 0, binding = 0) uniform sampler Sampler;
+layout(set = 0, binding = 0) uniform texture2D SampledImage;
+layout(set = 0, binding = 0) uniform sampler SPIRV_Cross_DummySampler;
+
+layout(location = 0) out vec4 _entryPointOutput;
+
+void main()
+{
+    ivec2 _152 = ivec3(int(gl_FragCoord.x * 1280.0), int(gl_FragCoord.y * 720.0), 0).xy;
+    _entryPointOutput = ((texelFetch(sampler2D(SampledImage, SPIRV_Cross_DummySampler), _152, 0) + texelFetch(sampler2D(SampledImage, SPIRV_Cross_DummySampler), _152, 0)) + texture(sampler2D(SampledImage, Sampler), gl_FragCoord.xy)) + texture(sampler2D(SampledImage, Sampler), gl_FragCoord.xy);
+}
+

3rdparty/spirv-cross/reference/opt/shaders/asm/frag/image-query-no-sampler.no-samplerless.vk.asm.frag

@@ -0,0 +1,6 @@
+#version 450
+
+void main()
+{
+}
+

3rdparty/spirv-cross/reference/opt/shaders/asm/frag/image-query-no-sampler.no-samplerless.vk.asm.frag.vk

@@ -0,0 +1,6 @@
+#version 450
+
+void main()
+{
+}
+

3rdparty/spirv-cross/reference/opt/shaders/asm/frag/line-directive.line.asm.frag

@@ -7,14 +7,12 @@ layout(location = 0) in float vColor;
 #line 8 "test.frag"
 void main()
 {
-    float _80;
 #line 8 "test.frag"
     FragColor = 1.0;
 #line 9 "test.frag"
     FragColor = 2.0;
 #line 10 "test.frag"
-    _80 = vColor;
-    if (_80 < 0.0)
+    if (vColor < 0.0)
     {
 #line 12 "test.frag"
         FragColor = 3.0;
@@ -24,16 +22,16 @@ void main()
 #line 16 "test.frag"
         FragColor = 4.0;
     }
-    for (int _126 = 0; float(_126) < (40.0 + _80); )
+    for (int _126 = 0; float(_126) < (40.0 + vColor); )
     {
 #line 21 "test.frag"
         FragColor += 0.20000000298023223876953125;
 #line 22 "test.frag"
         FragColor += 0.300000011920928955078125;
-        _126 += (int(_80) + 5);
+        _126 += (int(vColor) + 5);
         continue;
     }
-    switch (int(_80))
+    switch (int(vColor))
     {
         case 0:
         {
@@ -59,7 +57,7 @@ void main()
     }
     for (;;)
     {
-        FragColor += (10.0 + _80);
+        FragColor += (10.0 + vColor);
 #line 43 "test.frag"
         if (FragColor < 100.0)
         {

3rdparty/spirv-cross/reference/opt/shaders/asm/frag/loop-body-dominator-continue-access.asm.frag

@@ -46,14 +46,7 @@ void main()
                     _231 = true;
                     break;
                 }
-                else
-                {
-                    uint _204 = _227 + uint(1);
-                    _227 = _204;
-                    continue;
-                }
-                uint _204 = _227 + uint(1);
-                _227 = _204;
+                _227++;
                 continue;
             }
             else

3rdparty/spirv-cross/reference/shaders-msl-no-opt/asm/frag/inliner-dominator-inside-loop.asm.frag

@@ -137,7 +137,7 @@ fragment main0_out main0(main0_in in [[stage_in]], constant CB0& _19 [[buffer(0)
     float2 _166 = in.IN_Uv_EdgeDistance1.xy * 1.0;
     bool _173;
     float4 _193;
-    do
+    for (;;)
     {
         _173 = 0.0 == 0.0;
         if (_173)
@@ -153,9 +153,9 @@ fragment main0_out main0(main0_in in [[stage_in]], constant CB0& _19 [[buffer(0)
         }
         _193 = _192;
         break;
-    } while (false);
+    }
     float4 _220;
-    do
+    for (;;)
     {
         if (_173)
         {
@@ -170,7 +170,7 @@ fragment main0_out main0(main0_in in [[stage_in]], constant CB0& _19 [[buffer(0)
         }
         _220 = _219;
         break;
-    } while (false);
+    }
     float2 _223 = float2(1.0);
     float2 _224 = (_220.wy * 2.0) - _223;
     float3 _232 = float3(_224, sqrt(fast::clamp(1.0 + dot(-_224, _224), 0.0, 1.0)));
@@ -181,7 +181,7 @@ fragment main0_out main0(main0_in in [[stage_in]], constant CB0& _19 [[buffer(0)
     float3 _256 = float3(_255.x, _255.y, _253.z);
     float3 _271 = ((in.IN_Color.xyz * (_193 * 1.0).xyz) * (1.0 + (_256.x * 0.300000011920928955078125))) * (StudsMapTexture.sample(StudsMapSampler, _156.UvStuds).x * 2.0);
     float4 _298;
-    do
+    for (;;)
     {
         if (0.75 == 0.0)
         {
@@ -196,7 +196,7 @@ fragment main0_out main0(main0_in in [[stage_in]], constant CB0& _19 [[buffer(0)
         }
         _298 = _297;
         break;
-    } while (false);
+    }
     float2 _303 = mix(float2(0.800000011920928955078125, 120.0), (_298.xy * float2(2.0, 256.0)) + float2(0.0, 0.00999999977648258209228515625), float2(_165));
     Surface _304 = _125;
     _304.albedo = _271;

3rdparty/spirv-cross/reference/shaders-msl-no-opt/asm/frag/modf-frexp-scalar-access-chain-output.asm.frag

@@ -0,0 +1,17 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+fragment void main0()
+{
+    float3 col;
+    int2 _18;
+    float _23;
+    float _21 = modf(0.1500000059604644775390625, _23);
+    col.x = _23;
+    int _24;
+    float _22 = frexp(0.1500000059604644775390625, _24);
+    _18.y = _24;
+}
+

3rdparty/spirv-cross/reference/shaders-msl/comp/basic.dispatchbase.comp

@@ -0,0 +1,41 @@
+#pragma clang diagnostic ignored "-Wunused-variable"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+#include <metal_atomic>
+
+using namespace metal;
+
+struct SSBO
+{
+    float4 in_data[1];
+};
+
+struct SSBO2
+{
+    float4 out_data[1];
+};
+
+struct SSBO3
+{
+    uint counter;
+};
+
+constant uint _59_tmp [[function_constant(10)]];
+constant uint _59 = is_function_constant_defined(_59_tmp) ? _59_tmp : 1u;
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(_59, 1u, 1u);
+
+kernel void main0(const device SSBO& _27 [[buffer(0)]], device SSBO2& _49 [[buffer(1)]], device SSBO3& _52 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 spvDispatchBase [[grid_origin]])
+{
+    gl_GlobalInvocationID += spvDispatchBase * gl_WorkGroupSize;
+    gl_WorkGroupID += spvDispatchBase;
+    uint ident = gl_GlobalInvocationID.x;
+    uint workgroup = gl_WorkGroupID.x;
+    float4 idata = _27.in_data[ident];
+    if (dot(idata, float4(1.0, 5.0, 6.0, 2.0)) > 8.19999980926513671875)
+    {
+        uint _56 = atomic_fetch_add_explicit((device atomic_uint*)&_52.counter, 1u, memory_order_relaxed);
+        _49.out_data[_56] = idata;
+    }
+}
+

3rdparty/spirv-cross/reference/shaders-msl/comp/basic.dispatchbase.msl11.comp

@@ -0,0 +1,37 @@
+#pragma clang diagnostic ignored "-Wunused-variable"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+#include <metal_atomic>
+
+using namespace metal;
+
+struct SSBO
+{
+    float4 in_data[1];
+};
+
+struct SSBO2
+{
+    float4 out_data[1];
+};
+
+struct SSBO3
+{
+    uint counter;
+};
+
+kernel void main0(constant uint3& spvDispatchBase [[buffer(29)]], const device SSBO& _27 [[buffer(0)]], device SSBO2& _49 [[buffer(1)]], device SSBO3& _52 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
+{
+    gl_GlobalInvocationID += spvDispatchBase * uint3(1, 1, 1);
+    gl_WorkGroupID += spvDispatchBase;
+    uint ident = gl_GlobalInvocationID.x;
+    uint workgroup = gl_WorkGroupID.x;
+    float4 idata = _27.in_data[ident];
+    if (dot(idata, float4(1.0, 5.0, 6.0, 2.0)) > 8.19999980926513671875)
+    {
+        uint _56 = atomic_fetch_add_explicit((device atomic_uint*)&_52.counter, 1u, memory_order_relaxed);
+        _49.out_data[_56] = idata;
+    }
+}
+

3rdparty/spirv-cross/reference/shaders-no-opt/asm/frag/do-while-continue-phi.asm.frag

@@ -0,0 +1,37 @@
+#version 310 es
+precision mediump float;
+precision highp int;
+
+layout(location = 0) out highp vec4 _GLF_color;
+
+void main()
+{
+    for (;;)
+    {
+        bool _32;
+        for (;;)
+        {
+            if (gl_FragCoord.x != gl_FragCoord.x)
+            {
+                _32 = true;
+                break;
+            }
+            if (false)
+            {
+                continue;
+            }
+            else
+            {
+                _32 = false;
+                break;
+            }
+        }
+        if (_32)
+        {
+            break;
+        }
+        _GLF_color = vec4(1.0, 0.0, 0.0, 1.0);
+        break;
+    }
+}
+

3rdparty/spirv-cross/reference/shaders-no-opt/asm/frag/inliner-dominator-inside-loop.asm.frag

@@ -130,7 +130,7 @@ void main()
     vec2 _166 = IN_Uv_EdgeDistance1.xy * 1.0;
     bool _173;
     vec4 _193;
-    do
+    for (;;)
     {
         _173 = 0.0 == 0.0;
         if (_173)
@@ -146,9 +146,9 @@ void main()
         }
         _193 = _192;
         break;
-    } while (false);
+    }
     vec4 _220;
-    do
+    for (;;)
     {
         if (_173)
         {
@@ -163,7 +163,7 @@ void main()
         }
         _220 = _219;
         break;
-    } while (false);
+    }
     vec2 _223 = vec2(1.0);
     vec2 _224 = (_220.wy * 2.0) - _223;
     vec3 _232 = vec3(_224, sqrt(clamp(1.0 + dot(-_224, _224), 0.0, 1.0)));
@@ -174,7 +174,7 @@ void main()
     vec3 _256 = vec3(_255.x, _255.y, _253.z);
     vec3 _271 = ((IN_Color.xyz * (_193 * 1.0).xyz) * (1.0 + (_256.x * 0.300000011920928955078125))) * (texture(SPIRV_Cross_CombinedStudsMapTextureStudsMapSampler, _156.UvStuds).x * 2.0);
     vec4 _298;
-    do
+    for (;;)
     {
         if (0.75 == 0.0)
         {
@@ -189,7 +189,7 @@ void main()
         }
         _298 = _297;
         break;
-    } while (false);
+    }
     vec2 _303 = mix(vec2(0.800000011920928955078125, 120.0), (_298.xy * vec2(2.0, 256.0)) + vec2(0.0, 0.00999999977648258209228515625), vec2(_165));
     Surface _304 = _125;
     _304.albedo = _271;

3rdparty/spirv-cross/reference/shaders/asm/comp/phi-temporary-copy-loop-variable.asm.comp

@@ -14,10 +14,6 @@ void main()
         {
             break;
         }
-        else
-        {
-            continue;
-        }
         continue;
     }
     imageStore(outImageTexture, ivec2(gl_GlobalInvocationID.xy), vec4(float(_30 - 1), float(_30), 1.0, 1.0));

3rdparty/spirv-cross/reference/shaders/asm/frag/image-fetch-no-sampler.no-samplerless.asm.vk.frag

@@ -0,0 +1,38 @@
+#version 450
+
+uniform sampler2D SPIRV_Cross_CombinedparamSPIRV_Cross_DummySampler;
+uniform sampler2D SPIRV_Cross_CombinedSampledImageSPIRV_Cross_DummySampler;
+uniform sampler2D SPIRV_Cross_CombinedparamSampler;
+uniform sampler2D SPIRV_Cross_CombinedSampledImageSampler;
+
+layout(location = 0) out vec4 _entryPointOutput;
+
+vec4 sample_fetch(ivec3 UV, sampler2D SPIRV_Cross_CombinedtexSPIRV_Cross_DummySampler)
+{
+    return texelFetch(SPIRV_Cross_CombinedtexSPIRV_Cross_DummySampler, UV.xy, UV.z);
+}
+
+vec4 sample_sampler(vec2 UV, sampler2D SPIRV_Cross_CombinedtexSampler)
+{
+    return texture(SPIRV_Cross_CombinedtexSampler, UV);
+}
+
+vec4 _main(vec4 xIn)
+{
+    ivec3 coord = ivec3(int(xIn.x * 1280.0), int(xIn.y * 720.0), 0);
+    ivec3 param = coord;
+    vec4 value = sample_fetch(param, SPIRV_Cross_CombinedparamSPIRV_Cross_DummySampler);
+    value += texelFetch(SPIRV_Cross_CombinedSampledImageSPIRV_Cross_DummySampler, coord.xy, coord.z);
+    vec2 param_1 = xIn.xy;
+    value += sample_sampler(param_1, SPIRV_Cross_CombinedparamSampler);
+    value += texture(SPIRV_Cross_CombinedSampledImageSampler, xIn.xy);
+    return value;
+}
+
+void main()
+{
+    vec4 xIn = gl_FragCoord;
+    vec4 param = xIn;
+    _entryPointOutput = _main(param);
+}
+

3rdparty/spirv-cross/reference/shaders/asm/frag/image-fetch-no-sampler.no-samplerless.asm.vk.frag.vk

@@ -0,0 +1,37 @@
+#version 450
+
+layout(set = 0, binding = 0) uniform sampler Sampler;
+layout(set = 0, binding = 0) uniform texture2D SampledImage;
+layout(set = 0, binding = 0) uniform sampler SPIRV_Cross_DummySampler;
+
+layout(location = 0) out vec4 _entryPointOutput;
+
+vec4 sample_fetch(texture2D tex, ivec3 UV)
+{
+    return texelFetch(sampler2D(tex, SPIRV_Cross_DummySampler), UV.xy, UV.z);
+}
+
+vec4 sample_sampler(texture2D tex, vec2 UV)
+{
+    return texture(sampler2D(tex, Sampler), UV);
+}
+
+vec4 _main(vec4 xIn)
+{
+    ivec3 coord = ivec3(int(xIn.x * 1280.0), int(xIn.y * 720.0), 0);
+    ivec3 param = coord;
+    vec4 value = sample_fetch(SampledImage, param);
+    value += texelFetch(sampler2D(SampledImage, SPIRV_Cross_DummySampler), coord.xy, coord.z);
+    vec2 param_1 = xIn.xy;
+    value += sample_sampler(SampledImage, param_1);
+    value += texture(sampler2D(SampledImage, Sampler), xIn.xy);
+    return value;
+}
+
+void main()
+{
+    vec4 xIn = gl_FragCoord;
+    vec4 param = xIn;
+    _entryPointOutput = _main(param);
+}
+

3rdparty/spirv-cross/reference/shaders/asm/frag/image-query-no-sampler.no-samplerless.vk.asm.frag

@@ -0,0 +1,13 @@
+#version 450
+
+uniform sampler2D SPIRV_Cross_CombineduSampler2DSPIRV_Cross_DummySampler;
+uniform sampler2DMS SPIRV_Cross_CombineduSampler2DMSSPIRV_Cross_DummySampler;
+
+void main()
+{
+    ivec2 b = textureSize(SPIRV_Cross_CombineduSampler2DSPIRV_Cross_DummySampler, 0);
+    ivec2 c = textureSize(SPIRV_Cross_CombineduSampler2DMSSPIRV_Cross_DummySampler);
+    int l1 = textureQueryLevels(SPIRV_Cross_CombineduSampler2DSPIRV_Cross_DummySampler);
+    int s0 = textureSamples(SPIRV_Cross_CombineduSampler2DMSSPIRV_Cross_DummySampler);
+}
+

3rdparty/spirv-cross/reference/shaders/asm/frag/image-query-no-sampler.no-samplerless.vk.asm.frag.vk

@@ -0,0 +1,14 @@
+#version 450
+
+layout(set = 0, binding = 0) uniform texture2D uSampler2D;
+layout(set = 0, binding = 0) uniform texture2DMS uSampler2DMS;
+layout(set = 0, binding = 0) uniform sampler SPIRV_Cross_DummySampler;
+
+void main()
+{
+    ivec2 b = textureSize(sampler2D(uSampler2D, SPIRV_Cross_DummySampler), 0);
+    ivec2 c = textureSize(sampler2DMS(uSampler2DMS, SPIRV_Cross_DummySampler));
+    int l1 = textureQueryLevels(sampler2D(uSampler2D, SPIRV_Cross_DummySampler));
+    int s0 = textureSamples(sampler2DMS(uSampler2DMS, SPIRV_Cross_DummySampler));
+}
+

3rdparty/spirv-cross/shaders-msl-no-opt/asm/frag/modf-frexp-scalar-access-chain-output.asm.frag

@@ -0,0 +1,36 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos Glslang Reference Front End; 7
+; Bound: 17
+; Schema: 0
+               OpCapability Shader
+          %1 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %main "main"
+               OpExecutionMode %main OriginUpperLeft
+               OpSource ESSL 310
+               OpName %main "main"
+               OpName %col "col"
+       %void = OpTypeVoid
+          %3 = OpTypeFunction %void
+      %float = OpTypeFloat 32
+%_ptr_Function_float = OpTypePointer Function %float
+%float_0_150000006 = OpConstant %float 0.150000006
+    %v3float = OpTypeVector %float 3
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+       %int = OpTypeInt 32 1
+     %int_0 = OpConstant %int 0
+     %int_1 = OpConstant %int 1
+	 %v2int = OpTypeVector %int 2
+%_ptr_Function_v2int = OpTypePointer Function %v2int
+%_ptr_Function_int = OpTypePointer Function %int
+       %main = OpFunction %void None %3
+          %5 = OpLabel
+        %col = OpVariable %_ptr_Function_v3float Function
+        %icol = OpVariable %_ptr_Function_v2int Function
+         %ptr_x = OpAccessChain %_ptr_Function_float %col %int_0
+         %ptr_y = OpAccessChain %_ptr_Function_int %icol %int_1
+         %16 = OpExtInst %float %1 Modf %float_0_150000006 %ptr_x
+         %17 = OpExtInst %float %1 Frexp %float_0_150000006 %ptr_y
+               OpReturn
+               OpFunctionEnd

3rdparty/spirv-cross/shaders-msl/comp/basic.dispatchbase.comp

@@ -0,0 +1,29 @@
+#version 310 es
+layout(local_size_x_id = 10) in;
+
+layout(std430, binding = 0) readonly buffer SSBO
+{
+    vec4 in_data[];
+};
+
+layout(std430, binding = 1) writeonly buffer SSBO2
+{
+    vec4 out_data[];
+};
+
+layout(std430, binding = 2) buffer SSBO3
+{
+    uint counter;
+};
+
+void main()
+{
+    uint ident = gl_GlobalInvocationID.x;
+    uint workgroup = gl_WorkGroupID.x;
+    vec4 idata = in_data[ident];
+    if (dot(idata, vec4(1.0, 5.0, 6.0, 2.0)) > 8.2)
+    {
+        out_data[atomicAdd(counter, 1u)] = idata;
+    }
+}
+

3rdparty/spirv-cross/shaders-msl/comp/basic.dispatchbase.msl11.comp

@@ -0,0 +1,29 @@
+#version 310 es
+layout(local_size_x = 1) in;
+
+layout(std430, binding = 0) readonly buffer SSBO
+{
+    vec4 in_data[];
+};
+
+layout(std430, binding = 1) writeonly buffer SSBO2
+{
+    vec4 out_data[];
+};
+
+layout(std430, binding = 2) buffer SSBO3
+{
+    uint counter;
+};
+
+void main()
+{
+    uint ident = gl_GlobalInvocationID.x;
+    uint workgroup = gl_WorkGroupID.x;
+    vec4 idata = in_data[ident];
+    if (dot(idata, vec4(1.0, 5.0, 6.0, 2.0)) > 8.2)
+    {
+        out_data[atomicAdd(counter, 1u)] = idata;
+    }
+}
+

3rdparty/spirv-cross/shaders-no-opt/asm/frag/do-while-continue-phi.asm.frag

@@ -0,0 +1,64 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos Glslang Reference Front End; 7
+; Bound: 42
+; Schema: 0
+               OpCapability Shader
+          %1 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %main "main" %gl_FragCoord %_GLF_color
+               OpExecutionMode %main OriginUpperLeft
+               OpSource ESSL 310
+               OpName %main "main"
+               OpName %gl_FragCoord "gl_FragCoord"
+               OpName %_GLF_color "_GLF_color"
+               OpDecorate %gl_FragCoord BuiltIn FragCoord
+               OpDecorate %_GLF_color Location 0
+       %void = OpTypeVoid
+          %3 = OpTypeFunction %void
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Input_v4float = OpTypePointer Input %v4float
+%gl_FragCoord = OpVariable %_ptr_Input_v4float Input
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_Input_float = OpTypePointer Input %float
+       %bool = OpTypeBool
+      %false = OpConstantFalse %bool
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+ %_GLF_color = OpVariable %_ptr_Output_v4float Output
+    %float_1 = OpConstant %float 1
+    %float_0 = OpConstant %float 0
+         %31 = OpConstantComposite %v4float %float_1 %float_0 %float_0 %float_1
+       %true = OpConstantTrue %bool
+       %main = OpFunction %void None %3
+          %5 = OpLabel
+               OpBranch %33
+         %33 = OpLabel
+               OpLoopMerge %32 %35 None
+               OpBranch %6
+          %6 = OpLabel
+               OpLoopMerge %8 %24 None
+               OpBranch %7
+          %7 = OpLabel
+         %17 = OpAccessChain %_ptr_Input_float %gl_FragCoord %uint_0
+         %18 = OpLoad %float %17
+         %22 = OpFOrdNotEqual %bool %18 %18
+               OpSelectionMerge %24 None
+               OpBranchConditional %22 %23 %24
+         %23 = OpLabel
+               OpBranch %8
+         %24 = OpLabel
+               OpBranchConditional %false %6 %8
+          %8 = OpLabel
+         %41 = OpPhi %bool %true %23 %false %24
+               OpSelectionMerge %39 None
+               OpBranchConditional %41 %32 %39
+         %39 = OpLabel
+               OpStore %_GLF_color %31
+               OpBranch %32
+         %35 = OpLabel
+               OpBranch %33
+         %32 = OpLabel
+               OpReturn
+               OpFunctionEnd

3rdparty/spirv-cross/shaders/asm/frag/image-fetch-no-sampler.no-samplerless.asm.vk.frag

@@ -0,0 +1,163 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos Glslang Reference Front End; 2
+; Bound: 113
+; Schema: 0
+               OpCapability Shader
+          %1 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %main "main" %xIn_1 %_entryPointOutput
+               OpExecutionMode %main OriginUpperLeft
+               OpSource HLSL 500
+               OpName %main "main"
+               OpName %sample_fetch_t21_vi3_ "sample_fetch(t21;vi3;"
+               OpName %tex "tex"
+               OpName %UV "UV"
+               OpName %sample_sampler_t21_vf2_ "sample_sampler(t21;vf2;"
+               OpName %tex_0 "tex"
+               OpName %UV_0 "UV"
+               OpName %_main_vf4_ "@main(vf4;"
+               OpName %xIn "xIn"
+               OpName %Sampler "Sampler"
+               OpName %coord "coord"
+               OpName %value "value"
+               OpName %SampledImage "SampledImage"
+               OpName %param "param"
+               OpName %param_0 "param"
+               OpName %param_1 "param"
+               OpName %param_2 "param"
+               OpName %xIn_0 "xIn"
+               OpName %xIn_1 "xIn"
+               OpName %_entryPointOutput "@entryPointOutput"
+               OpName %param_3 "param"
+               OpDecorate %Sampler DescriptorSet 0
+               OpDecorate %Sampler Binding 0
+               OpDecorate %SampledImage DescriptorSet 0
+               OpDecorate %SampledImage Binding 0
+               OpDecorate %xIn_1 BuiltIn FragCoord
+               OpDecorate %_entryPointOutput Location 0
+       %void = OpTypeVoid
+          %3 = OpTypeFunction %void
+      %float = OpTypeFloat 32
+          %7 = OpTypeImage %float 2D 0 0 0 1 Unknown
+%_ptr_Function_7 = OpTypePointer Function %7
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+    %v4float = OpTypeVector %float 4
+         %13 = OpTypeFunction %v4float %_ptr_Function_7 %_ptr_Function_v3int
+    %v2float = OpTypeVector %float 2
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+         %20 = OpTypeFunction %v4float %_ptr_Function_7 %_ptr_Function_v2float
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+         %26 = OpTypeFunction %v4float %_ptr_Function_v4float
+      %v2int = OpTypeVector %int 2
+       %uint = OpTypeInt 32 0
+     %uint_2 = OpConstant %uint 2
+%_ptr_Function_int = OpTypePointer Function %int
+         %43 = OpTypeSampler
+%_ptr_UniformConstant_43 = OpTypePointer UniformConstant %43
+    %Sampler = OpVariable %_ptr_UniformConstant_43 UniformConstant
+         %47 = OpTypeSampledImage %7
+     %uint_0 = OpConstant %uint 0
+%_ptr_Function_float = OpTypePointer Function %float
+ %float_1280 = OpConstant %float 1280
+     %uint_1 = OpConstant %uint 1
+  %float_720 = OpConstant %float 720
+      %int_0 = OpConstant %int 0
+%_ptr_UniformConstant_7 = OpTypePointer UniformConstant %7
+%SampledImage = OpVariable %_ptr_UniformConstant_7 UniformConstant
+%_ptr_Input_v4float = OpTypePointer Input %v4float
+      %xIn_1 = OpVariable %_ptr_Input_v4float Input
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+%_entryPointOutput = OpVariable %_ptr_Output_v4float Output
+       %main = OpFunction %void None %3
+          %5 = OpLabel
+      %xIn_0 = OpVariable %_ptr_Function_v4float Function
+    %param_3 = OpVariable %_ptr_Function_v4float Function
+        %107 = OpLoad %v4float %xIn_1
+               OpStore %xIn_0 %107
+        %111 = OpLoad %v4float %xIn_0
+               OpStore %param_3 %111
+        %112 = OpFunctionCall %v4float %_main_vf4_ %param_3
+               OpStore %_entryPointOutput %112
+               OpReturn
+               OpFunctionEnd
+%sample_fetch_t21_vi3_ = OpFunction %v4float None %13
+        %tex = OpFunctionParameter %_ptr_Function_7
+         %UV = OpFunctionParameter %_ptr_Function_v3int
+         %17 = OpLabel
+         %30 = OpLoad %7 %tex
+         %32 = OpLoad %v3int %UV
+         %33 = OpVectorShuffle %v2int %32 %32 0 1
+         %37 = OpAccessChain %_ptr_Function_int %UV %uint_2
+         %38 = OpLoad %int %37
+         %39 = OpImageFetch %v4float %30 %33 Lod %38
+               OpReturnValue %39
+               OpFunctionEnd
+%sample_sampler_t21_vf2_ = OpFunction %v4float None %20
+      %tex_0 = OpFunctionParameter %_ptr_Function_7
+       %UV_0 = OpFunctionParameter %_ptr_Function_v2float
+         %24 = OpLabel
+         %42 = OpLoad %7 %tex_0
+         %46 = OpLoad %43 %Sampler
+         %48 = OpSampledImage %47 %42 %46
+         %49 = OpLoad %v2float %UV_0
+         %50 = OpImageSampleImplicitLod %v4float %48 %49
+               OpReturnValue %50
+               OpFunctionEnd
+ %_main_vf4_ = OpFunction %v4float None %26
+        %xIn = OpFunctionParameter %_ptr_Function_v4float
+         %29 = OpLabel
+      %coord = OpVariable %_ptr_Function_v3int Function
+      %value = OpVariable %_ptr_Function_v4float Function
+      %param = OpVariable %_ptr_Function_7 Function
+    %param_0 = OpVariable %_ptr_Function_v3int Function
+    %param_1 = OpVariable %_ptr_Function_7 Function
+    %param_2 = OpVariable %_ptr_Function_v2float Function
+         %56 = OpAccessChain %_ptr_Function_float %xIn %uint_0
+         %57 = OpLoad %float %56
+         %59 = OpFMul %float %57 %float_1280
+         %60 = OpConvertFToS %int %59
+         %62 = OpAccessChain %_ptr_Function_float %xIn %uint_1
+         %63 = OpLoad %float %62
+         %65 = OpFMul %float %63 %float_720
+         %66 = OpConvertFToS %int %65
+         %68 = OpCompositeConstruct %v3int %60 %66 %int_0
+               OpStore %coord %68
+         %73 = OpLoad %7 %SampledImage
+               OpStore %param %73
+         %75 = OpLoad %v3int %coord
+               OpStore %param_0 %75
+         %76 = OpFunctionCall %v4float %sample_fetch_t21_vi3_ %param %param_0
+               OpStore %value %76
+         %77 = OpLoad %7 %SampledImage
+         %78 = OpLoad %v3int %coord
+         %79 = OpVectorShuffle %v2int %78 %78 0 1
+         %80 = OpAccessChain %_ptr_Function_int %coord %uint_2
+         %81 = OpLoad %int %80
+         %82 = OpImageFetch %v4float %77 %79 Lod %81
+         %83 = OpLoad %v4float %value
+         %84 = OpFAdd %v4float %83 %82
+               OpStore %value %84
+         %86 = OpLoad %7 %SampledImage
+               OpStore %param_1 %86
+         %88 = OpLoad %v4float %xIn
+         %89 = OpVectorShuffle %v2float %88 %88 0 1
+               OpStore %param_2 %89
+         %90 = OpFunctionCall %v4float %sample_sampler_t21_vf2_ %param_1 %param_2
+         %91 = OpLoad %v4float %value
+         %92 = OpFAdd %v4float %91 %90
+               OpStore %value %92
+         %93 = OpLoad %7 %SampledImage
+         %94 = OpLoad %43 %Sampler
+         %95 = OpSampledImage %47 %93 %94
+         %96 = OpLoad %v4float %xIn
+         %97 = OpVectorShuffle %v2float %96 %96 0 1
+         %98 = OpImageSampleImplicitLod %v4float %95 %97
+         %99 = OpLoad %v4float %value
+        %100 = OpFAdd %v4float %99 %98
+               OpStore %value %100
+        %101 = OpLoad %v4float %value
+               OpReturnValue %101
+               OpFunctionEnd

3rdparty/spirv-cross/shaders/asm/frag/image-query-no-sampler.no-samplerless.vk.asm.frag

@@ -0,0 +1,57 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos Glslang Reference Front End; 6
+; Bound: 36
+; Schema: 0
+               OpCapability Shader
+               OpCapability ImageQuery
+          %1 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %main "main"
+               OpExecutionMode %main OriginUpperLeft
+               OpSource GLSL 450
+               OpName %main "main"
+               OpName %b "b"
+               OpName %uSampler2D "uSampler2D"
+               OpName %c "c"
+               OpName %uSampler2DMS "uSampler2DMS"
+               OpName %l1 "l1"
+               OpName %s0 "s0"
+               OpDecorate %uSampler2D DescriptorSet 0
+               OpDecorate %uSampler2D Binding 0
+               OpDecorate %uSampler2DMS DescriptorSet 0
+               OpDecorate %uSampler2DMS Binding 0
+       %void = OpTypeVoid
+          %3 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+%_ptr_Function_v2int = OpTypePointer Function %v2int
+      %float = OpTypeFloat 32
+         %11 = OpTypeImage %float 2D 0 0 0 1 Unknown
+%_ptr_UniformConstant_12 = OpTypePointer UniformConstant %11
+ %uSampler2D = OpVariable %_ptr_UniformConstant_12 UniformConstant
+      %int_0 = OpConstant %int 0
+         %20 = OpTypeImage %float 2D 0 0 1 1 Unknown
+%_ptr_UniformConstant_21 = OpTypePointer UniformConstant %20
+%uSampler2DMS = OpVariable %_ptr_UniformConstant_21 UniformConstant
+%_ptr_Function_int = OpTypePointer Function %int
+       %main = OpFunction %void None %3
+          %5 = OpLabel
+          %b = OpVariable %_ptr_Function_v2int Function
+          %c = OpVariable %_ptr_Function_v2int Function
+         %l1 = OpVariable %_ptr_Function_int Function
+         %s0 = OpVariable %_ptr_Function_int Function
+         %15 = OpLoad %11 %uSampler2D
+         %18 = OpImageQuerySizeLod %v2int %15 %int_0
+               OpStore %b %18
+         %24 = OpLoad %20 %uSampler2DMS
+         %26 = OpImageQuerySize %v2int %24
+               OpStore %c %26
+         %29 = OpLoad %11 %uSampler2D
+         %31 = OpImageQueryLevels %int %29
+               OpStore %l1 %31
+         %33 = OpLoad %20 %uSampler2DMS
+         %35 = OpImageQuerySamples %int %33
+               OpStore %s0 %35
+               OpReturn
+               OpFunctionEnd

3rdparty/spirv-cross/spirv_common.hpp

@@ -1433,6 +1433,10 @@ enum ExtendedDecorations
 	// Marks a buffer block for using explicit offsets (GLSL/HLSL).
 	SPIRVCrossDecorationExplicitOffset,
 
+	// Apply to a variable in the Input storage class; marks it as holding the base group passed to vkCmdDispatchBase().
+	// In MSL, this is used to adjust the WorkgroupId and GlobalInvocationId variables.
+	SPIRVCrossDecorationBuiltInDispatchBase,
+
 	SPIRVCrossDecorationCount
 };
 

3rdparty/spirv-cross/spirv_cross.cpp

@@ -1501,6 +1501,12 @@ SPIRBlock::ContinueBlockType Compiler::continue_block_type(const SPIRBlock &bloc
 		const auto *true_block = maybe_get<SPIRBlock>(block.true_block);
 		const auto *merge_block = maybe_get<SPIRBlock>(dominator.merge_block);
 
+		// If we need to flush Phi in this block, we cannot have a DoWhile loop.
+		bool flush_phi_to_false = false_block && flush_phi_required(block.self, block.false_block);
+		bool flush_phi_to_true = true_block && flush_phi_required(block.self, block.true_block);
+		if (flush_phi_to_false || flush_phi_to_true)
+			return SPIRBlock::ComplexLoop;
+
 		bool positive_do_while = block.true_block == dominator.self &&
 		                         (block.false_block == dominator.merge_block ||
 		                          (false_block && merge_block && execution_is_noop(*false_block, *merge_block)));
@@ -3288,10 +3294,11 @@ void Compiler::analyze_variable_scope(SPIRFunction &entry, AnalyzeVariableScopeA
 		{
 			builder.add_block(block);
 
-			// If a temporary is used in more than one block, we might have to lift continue block
-			// access up to loop header like we did for variables.
 			if (blocks.size() != 1 && is_continue(block))
 			{
+				// The risk here is that inner loop can dominate the continue block.
+				// Any temporary we access in the continue block must be declared before the loop.
+				// This is moot for complex loops however.
 				auto &loop_header_block = get<SPIRBlock>(ir.continue_block_to_loop_header[block]);
 				assert(loop_header_block.merge == SPIRBlock::MergeLoop);
 
@@ -3299,14 +3306,17 @@ void Compiler::analyze_variable_scope(SPIRFunction &entry, AnalyzeVariableScopeA
 				if (!loop_header_block.complex_continue)
 					builder.add_block(loop_header_block.self);
 			}
-			else if (blocks.size() != 1 && is_single_block_loop(block))
-			{
-				// Awkward case, because the loop header is also the continue block.
-				force_temporary = true;
-			}
 		}
 
 		uint32_t dominating_block = builder.get_dominator();
+
+		if (blocks.size() != 1 && is_single_block_loop(dominating_block))
+		{
+			// Awkward case, because the loop header is also the continue block,
+			// so hoisting to loop header does not help.
+			force_temporary = true;
+		}
+
 		if (dominating_block)
 		{
 			// If we touch a variable in the dominating block, this is the expected setup.
@@ -4246,3 +4256,12 @@ bool Compiler::type_is_array_of_pointers(const SPIRType &type) const
 	// If parent type has same pointer depth, we must have an array of pointers.
 	return type.pointer_depth == get<SPIRType>(type.parent_type).pointer_depth;
 }
+
+bool Compiler::flush_phi_required(uint32_t from, uint32_t to) const
+{
+	auto &child = get<SPIRBlock>(to);
+	for (auto &phi : child.phi_variables)
+		if (phi.parent == from)
+			return true;
+	return false;
+}

3rdparty/spirv-cross/spirv_cross.hpp

@@ -816,6 +816,7 @@ protected:
 	std::unordered_set<uint32_t> forwarded_temporaries;
 	std::unordered_set<uint32_t> suppressed_usage_tracking;
 	std::unordered_set<uint32_t> hoisted_temporaries;
+	std::unordered_set<uint32_t> forced_invariant_temporaries;
 
 	Bitset active_input_builtins;
 	Bitset active_output_builtins;
@@ -972,6 +973,8 @@ protected:
 	bool reflection_ssbo_instance_name_is_significant() const;
 	std::string get_remapped_declared_block_name(uint32_t id, bool fallback_prefer_instance_name) const;
 
+	bool flush_phi_required(uint32_t from, uint32_t to) const;
+
 private:
 	// Used only to implement the old deprecated get_entry_point() interface.
 	const SPIREntryPoint &get_first_entry_point(const std::string &name) const;

3rdparty/spirv-cross/spirv_glsl.cpp

@@ -4654,16 +4654,16 @@ void CompilerGLSL::emit_sampled_image_op(uint32_t result_type, uint32_t result_i
 	{
 		emit_binary_func_op(result_type, result_id, image_id, samp_id,
 		                    type_to_glsl(get<SPIRType>(result_type), result_id).c_str());
-
-		// Make sure to suppress usage tracking and any expression invalidation.
-		// It is illegal to create temporaries of opaque types.
-		forwarded_temporaries.erase(result_id);
 	}
 	else
 	{
 		// Make sure to suppress usage tracking. It is illegal to create temporaries of opaque types.
 		emit_op(result_type, result_id, to_combined_image_sampler(image_id, samp_id), true, true);
 	}
+
+	// Make sure to suppress usage tracking and any expression invalidation.
+	// It is illegal to create temporaries of opaque types.
+	forwarded_temporaries.erase(result_id);
 }
 
 static inline bool image_opcode_is_sample_no_dref(Op op)
@@ -4976,10 +4976,19 @@ std::string CompilerGLSL::convert_separate_image_to_expression(uint32_t id)
 		{
 			if (options.vulkan_semantics)
 			{
-				// Newer glslang supports this extension to deal with texture2D as argument to texture functions.
 				if (dummy_sampler_id)
-					SPIRV_CROSS_THROW("Vulkan GLSL should not have a dummy sampler for combining.");
-				require_extension_internal("GL_EXT_samplerless_texture_functions");
+				{
+					// Don't need to consider Shadow state since the dummy sampler is always non-shadow.
+					auto sampled_type = type;
+					sampled_type.basetype = SPIRType::SampledImage;
+					return join(type_to_glsl(sampled_type), "(", to_expression(id), ", ",
+					            to_expression(dummy_sampler_id), ")");
+				}
+				else
+				{
+					// Newer glslang supports this extension to deal with texture2D as argument to texture functions.
+					require_extension_internal("GL_EXT_samplerless_texture_functions");
+				}
 			}
 			else
 			{
@@ -5277,7 +5286,6 @@ void CompilerGLSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop,
 
 	case GLSLstd450ModfStruct:
 	{
-		forced_temporaries.insert(id);
 		auto &type = get<SPIRType>(result_type);
 		emit_uninitialized_temporary_expression(result_type, id);
 		statement(to_expression(id), ".", to_member_name(type, 0), " = ", "modf(", to_expression(args[0]), ", ",
@@ -5417,7 +5425,6 @@ void CompilerGLSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop,
 
 	case GLSLstd450FrexpStruct:
 	{
-		forced_temporaries.insert(id);
 		auto &type = get<SPIRType>(result_type);
 		emit_uninitialized_temporary_expression(result_type, id);
 		statement(to_expression(id), ".", to_member_name(type, 0), " = ", "frexp(", to_expression(args[0]), ", ",
@@ -7544,14 +7551,16 @@ void CompilerGLSL::disallow_forwarding_in_expression_chain(const SPIRExpression
 	// Allow trivially forwarded expressions like OpLoad or trivial shuffles,
 	// these will be marked as having suppressed usage tracking.
 	// Our only concern is to make sure arithmetic operations are done in similar ways.
-	if (expression_is_forwarded(expr.self) && !expression_suppresses_usage_tracking(expr.self))
+	if (expression_is_forwarded(expr.self) && !expression_suppresses_usage_tracking(expr.self) &&
+	    forced_invariant_temporaries.count(expr.self) == 0)
 	{
 		forced_temporaries.insert(expr.self);
+		forced_invariant_temporaries.insert(expr.self);
 		force_recompile();
-	}
 
-	for (auto &dependent : expr.expression_dependencies)
-		disallow_forwarding_in_expression_chain(get<SPIRExpression>(dependent));
+		for (auto &dependent : expr.expression_dependencies)
+			disallow_forwarding_in_expression_chain(get<SPIRExpression>(dependent));
+	}
 }
 
 void CompilerGLSL::handle_store_to_invariant_variable(uint32_t store_id, uint32_t value_id)
@@ -8525,7 +8534,6 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
 		uint32_t result_id = ops[1];
 		uint32_t op0 = ops[2];
 		uint32_t op1 = ops[3];
-		forced_temporaries.insert(result_id);
 		auto &type = get<SPIRType>(result_type);
 		emit_uninitialized_temporary_expression(result_type, result_id);
 		const char *op = opcode == OpUMulExtended ? "umulExtended" : "imulExtended";
@@ -11060,15 +11068,6 @@ void CompilerGLSL::emit_fixup()
 	}
 }
 
-bool CompilerGLSL::flush_phi_required(uint32_t from, uint32_t to)
-{
-	auto &child = get<SPIRBlock>(to);
-	for (auto &phi : child.phi_variables)
-		if (phi.parent == from)
-			return true;
-	return false;
-}
-
 void CompilerGLSL::flush_phi(uint32_t from, uint32_t to)
 {
 	auto &child = get<SPIRBlock>(to);
@@ -11238,10 +11237,17 @@ void CompilerGLSL::branch(uint32_t from, uint32_t to)
 
 void CompilerGLSL::branch(uint32_t from, uint32_t cond, uint32_t true_block, uint32_t false_block)
 {
-	// If we branch directly to a selection merge target, we don't really need a code path.
+	auto &from_block = get<SPIRBlock>(from);
+	uint32_t merge_block = from_block.merge == SPIRBlock::MergeSelection ? from_block.next_block : 0;
+
+	// If we branch directly to a selection merge target, we don't need a code path.
+	// This covers both merge out of if () / else () as well as a break for switch blocks.
 	bool true_sub = !is_conditional(true_block);
 	bool false_sub = !is_conditional(false_block);
 
+	bool true_block_is_selection_merge = true_block == merge_block;
+	bool false_block_is_selection_merge = false_block == merge_block;
+
 	if (true_sub)
 	{
 		emit_block_hints(get<SPIRBlock>(from));
@@ -11250,7 +11256,11 @@ void CompilerGLSL::branch(uint32_t from, uint32_t cond, uint32_t true_block, uin
 		branch(from, true_block);
 		end_scope();
 
-		if (false_sub || is_continue(false_block) || is_break(false_block))
+		// If we merge to continue, we handle that explicitly in emit_block_chain(),
+		// so there is no need to branch to it directly here.
+		// break; is required to handle ladder fallthrough cases, so keep that in for now, even
+		// if we could potentially handle it in emit_block_chain().
+		if (false_sub || (!false_block_is_selection_merge && is_continue(false_block)) || is_break(false_block))
 		{
 			statement("else");
 			begin_scope();
@@ -11265,7 +11275,7 @@ void CompilerGLSL::branch(uint32_t from, uint32_t cond, uint32_t true_block, uin
 			end_scope();
 		}
 	}
-	else if (false_sub && !true_sub)
+	else if (false_sub)
 	{
 		// Only need false path, use negative conditional.
 		emit_block_hints(get<SPIRBlock>(from));
@@ -11274,7 +11284,7 @@ void CompilerGLSL::branch(uint32_t from, uint32_t cond, uint32_t true_block, uin
 		branch(from, false_block);
 		end_scope();
 
-		if (is_continue(true_block) || is_break(true_block))
+		if ((!true_block_is_selection_merge && is_continue(true_block)) || is_break(true_block))
 		{
 			statement("else");
 			begin_scope();

3rdparty/spirv-cross/spirv_glsl.hpp

@@ -436,7 +436,6 @@ protected:
 	void branch_to_continue(uint32_t from, uint32_t to);
 	void branch(uint32_t from, uint32_t cond, uint32_t true_block, uint32_t false_block);
 	void flush_phi(uint32_t from, uint32_t to);
-	bool flush_phi_required(uint32_t from, uint32_t to);
 	void flush_variable_declaration(uint32_t id);
 	void flush_undeclared_variables(SPIRBlock &block);
 	void emit_variable_temporary_copies(const SPIRVariable &var);

3rdparty/spirv-cross/spirv_msl.cpp

@@ -107,8 +107,11 @@ void CompilerMSL::build_implicit_builtins()
 	                                                       active_input_builtins.get(BuiltInSubgroupGtMask));
 	bool need_multiview = get_execution_model() == ExecutionModelVertex && !msl_options.view_index_from_device_index &&
 	                      (msl_options.multiview || active_input_builtins.get(BuiltInViewIndex));
+	bool need_dispatch_base =
+	    msl_options.dispatch_base && get_execution_model() == ExecutionModelGLCompute &&
+	    (active_input_builtins.get(BuiltInWorkgroupId) || active_input_builtins.get(BuiltInGlobalInvocationId));
 	if (need_subpass_input || need_sample_pos || need_subgroup_mask || need_vertex_params || need_tesc_params ||
-	    need_multiview || needs_subgroup_invocation_id)
+	    need_multiview || need_dispatch_base || needs_subgroup_invocation_id)
 	{
 		bool has_frag_coord = false;
 		bool has_sample_id = false;
@@ -121,6 +124,7 @@ void CompilerMSL::build_implicit_builtins()
 		bool has_subgroup_invocation_id = false;
 		bool has_subgroup_size = false;
 		bool has_view_idx = false;
+		uint32_t workgroup_id_type = 0;
 
 		ir.for_each_typed_id<SPIRVariable>([&](uint32_t, SPIRVariable &var) {
 			if (var.storage != StorageClassInput || !ir.meta[var.self].decoration.builtin)
@@ -208,6 +212,13 @@ void CompilerMSL::build_implicit_builtins()
 					has_view_idx = true;
 				}
 			}
+
+			// The base workgroup needs to have the same type and vector size
+			// as the workgroup or invocation ID, so keep track of the type that
+			// was used.
+			if (need_dispatch_base && workgroup_id_type == 0 &&
+			    (builtin == BuiltInWorkgroupId || builtin == BuiltInGlobalInvocationId))
+				workgroup_id_type = var.basetype;
 		});
 
 		if (!has_frag_coord && need_subpass_input)
@@ -457,6 +468,42 @@ void CompilerMSL::build_implicit_builtins()
 			builtin_subgroup_size_id = var_id;
 			mark_implicit_builtin(StorageClassInput, BuiltInSubgroupSize, var_id);
 		}
+
+		if (need_dispatch_base)
+		{
+			uint32_t var_id;
+			if (msl_options.supports_msl_version(1, 2))
+			{
+				// If we have MSL 1.2, we can (ab)use the [[grid_origin]] builtin
+				// to convey this information and save a buffer slot.
+				uint32_t offset = ir.increase_bound_by(1);
+				var_id = offset;
+
+				set<SPIRVariable>(var_id, workgroup_id_type, StorageClassInput);
+				set_extended_decoration(var_id, SPIRVCrossDecorationBuiltInDispatchBase);
+				get_entry_point().interface_variables.push_back(var_id);
+			}
+			else
+			{
+				// Otherwise, we need to fall back to a good ol' fashioned buffer.
+				uint32_t offset = ir.increase_bound_by(2);
+				var_id = offset;
+				uint32_t type_id = offset + 1;
+
+				SPIRType var_type = get<SPIRType>(workgroup_id_type);
+				var_type.storage = StorageClassUniform;
+				set<SPIRType>(type_id, var_type);
+
+				set<SPIRVariable>(var_id, type_id, StorageClassUniform);
+				// This should never match anything.
+				set_decoration(var_id, DecorationDescriptorSet, ~(5u));
+				set_decoration(var_id, DecorationBinding, msl_options.indirect_params_buffer_index);
+				set_extended_decoration(var_id, SPIRVCrossDecorationResourceIndexPrimary,
+				                        msl_options.indirect_params_buffer_index);
+			}
+			set_name(var_id, "spvDispatchBase");
+			builtin_dispatch_base_id = var_id;
+		}
 	}
 
 	if (needs_swizzle_buffer_def)
@@ -802,6 +849,8 @@ string CompilerMSL::compile()
 		active_interface_variables.insert(view_mask_buffer_id);
 	if (builtin_layer_id)
 		active_interface_variables.insert(builtin_layer_id);
+	if (builtin_dispatch_base_id && !msl_options.supports_msl_version(1, 2))
+		active_interface_variables.insert(builtin_dispatch_base_id);
 
 	// Create structs to hold input, output and uniform variables.
 	// Do output first to ensure out. is declared at top of entry function.
@@ -4600,10 +4649,8 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
 		uint32_t result_id = ops[1];
 		uint32_t op0 = ops[2];
 		uint32_t op1 = ops[3];
-		forced_temporaries.insert(result_id);
 		auto &type = get<SPIRType>(result_type);
-		statement(variable_decl(type, to_name(result_id)), ";");
-		set<SPIRExpression>(result_id, to_name(result_id), result_type, true);
+		emit_uninitialized_temporary_expression(result_type, result_id);
 
 		auto &res_type = get<SPIRType>(type.member_types[1]);
 		if (opcode == OpIAddCarry)
@@ -4632,10 +4679,8 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
 		uint32_t result_id = ops[1];
 		uint32_t op0 = ops[2];
 		uint32_t op1 = ops[3];
-		forced_temporaries.insert(result_id);
 		auto &type = get<SPIRType>(result_type);
-		statement(variable_decl(type, to_name(result_id)), ";");
-		set<SPIRExpression>(result_id, to_name(result_id), result_type, true);
+		emit_uninitialized_temporary_expression(result_type, result_id);
 
 		statement(to_expression(result_id), ".", to_member_name(type, 0), " = ", to_enclosed_expression(op0), " * ",
 		          to_enclosed_expression(op1), ";");
@@ -4917,8 +4962,6 @@ void CompilerMSL::emit_atomic_func_op(uint32_t result_type, uint32_t result_id,
                                       uint32_t mem_order_2, bool has_mem_order_2, uint32_t obj, uint32_t op1,
                                       bool op1_is_pointer, bool op1_is_literal, uint32_t op2)
 {
-	forced_temporaries.insert(result_id);
-
 	string exp = string(op) + "(";
 
 	auto &type = get_pointee_type(expression_type(obj));
@@ -4957,12 +5000,11 @@ void CompilerMSL::emit_atomic_func_op(uint32_t result_type, uint32_t result_id,
 		// the CAS loop, otherwise it will loop infinitely, with the comparison test always failing.
 		// The function updates the comparitor value from the memory value, so the additional
 		// comparison test evaluates the memory value against the expected value.
-		statement(variable_decl(type, to_name(result_id)), ";");
+		emit_uninitialized_temporary_expression(result_type, result_id);
 		statement("do");
 		begin_scope();
 		statement(to_name(result_id), " = ", to_expression(op1), ";");
 		end_scope_decl(join("while (!", exp, " && ", to_name(result_id), " == ", to_enclosed_expression(op1), ")"));
-		set<SPIRExpression>(result_id, to_name(result_id), result_type, true);
 	}
 	else
 	{
@@ -5211,6 +5253,32 @@ void CompilerMSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop,
 			CompilerGLSL::emit_glsl_op(result_type, id, eop, args, count);
 		break;
 
+	case GLSLstd450Modf:
+	case GLSLstd450Frexp:
+	{
+		// Special case. If the variable is a scalar access chain, we cannot use it directly. We have to emit a temporary.
+		auto *ptr = maybe_get<SPIRExpression>(args[1]);
+		if (ptr && ptr->access_chain && is_scalar(expression_type(args[1])))
+		{
+			register_call_out_argument(args[1]);
+			forced_temporaries.insert(id);
+
+			// Need to create temporaries and copy over to access chain after.
+			// We cannot directly take the reference of a vector swizzle in MSL, even if it's scalar ...
+			uint32_t &tmp_id = extra_sub_expressions[id];
+			if (!tmp_id)
+				tmp_id = ir.increase_bound_by(1);
+
+			uint32_t tmp_type_id = get_pointee_type_id(ptr->expression_type);
+			emit_uninitialized_temporary_expression(tmp_type_id, tmp_id);
+			emit_binary_func_op(result_type, id, args[0], tmp_id, eop == GLSLstd450Modf ? "modf" : "frexp");
+			statement(to_expression(args[1]), " = ", to_expression(tmp_id), ";");
+		}
+		else
+			CompilerGLSL::emit_glsl_op(result_type, id, eop, args, count);
+		break;
+	}
+
 	default:
 		CompilerGLSL::emit_glsl_op(result_type, id, eop, args, count);
 		break;
@@ -6527,14 +6595,21 @@ string CompilerMSL::func_type_decl(SPIRType &type)
 string CompilerMSL::get_argument_address_space(const SPIRVariable &argument)
 {
 	const auto &type = get<SPIRType>(argument.basetype);
+	return get_type_address_space(type, argument.self, true);
+}
+
+string CompilerMSL::get_type_address_space(const SPIRType &type, uint32_t id, bool argument)
+{
+	// This can be called for variable pointer contexts as well, so be very careful about which method we choose.
 	Bitset flags;
-	if (type.basetype == SPIRType::Struct &&
+	auto *var = maybe_get<SPIRVariable>(id);
+	if (var && type.basetype == SPIRType::Struct &&
 	    (has_decoration(type.self, DecorationBlock) || has_decoration(type.self, DecorationBufferBlock)))
-		flags = ir.get_buffer_block_flags(argument);
+		flags = get_buffer_block_flags(id);
 	else
-		flags = get_decoration_bitset(argument.self);
-	const char *addr_space = nullptr;
+		flags = get_decoration_bitset(id);
 
+	const char *addr_space = nullptr;
 	switch (type.storage)
 	{
 	case StorageClassWorkgroup:
@@ -6546,7 +6621,7 @@ string CompilerMSL::get_argument_address_space(const SPIRVariable &argument)
 		// For arguments from variable pointers, we use the write count deduction, so
 		// we should not assume any constness here. Only for global SSBOs.
 		bool readonly = false;
-		if (has_decoration(type.self, DecorationBlock))
+		if (!var || has_decoration(type.self, DecorationBlock))
 			readonly = flags.get(DecorationNonWritable);
 
 		addr_space = readonly ? "const device" : "device";
@@ -6560,24 +6635,21 @@ string CompilerMSL::get_argument_address_space(const SPIRVariable &argument)
 		{
 			bool ssbo = has_decoration(type.self, DecorationBufferBlock);
 			if (ssbo)
-			{
-				bool readonly = flags.get(DecorationNonWritable);
-				addr_space = readonly ? "const device" : "device";
-			}
+				addr_space = flags.get(DecorationNonWritable) ? "const device" : "device";
 			else
 				addr_space = "constant";
-			break;
 		}
+		else if (!argument)
+			addr_space = "constant";
 		break;
 
 	case StorageClassFunction:
 	case StorageClassGeneric:
-		// No address space for plain values.
-		addr_space = type.pointer ? "thread" : "";
 		break;
 
 	case StorageClassInput:
-		if (get_execution_model() == ExecutionModelTessellationControl && argument.basevariable == stage_in_ptr_var_id)
+		if (get_execution_model() == ExecutionModelTessellationControl && var &&
+		    var->basevariable == stage_in_ptr_var_id)
 			addr_space = "threadgroup";
 		break;
 
@@ -6591,64 +6663,8 @@ string CompilerMSL::get_argument_address_space(const SPIRVariable &argument)
 	}
 
 	if (!addr_space)
-		addr_space = "thread";
-
-	return join(flags.get(DecorationVolatile) || flags.get(DecorationCoherent) ? "volatile " : "", addr_space);
-}
-
-string CompilerMSL::get_type_address_space(const SPIRType &type, uint32_t id)
-{
-	// This can be called for variable pointer contexts as well, so be very careful about which method we choose.
-	Bitset flags;
-	if (ir.ids[id].get_type() == TypeVariable && type.basetype == SPIRType::Struct &&
-	    (has_decoration(type.self, DecorationBlock) || has_decoration(type.self, DecorationBufferBlock)))
-		flags = get_buffer_block_flags(id);
-	else
-		flags = get_decoration_bitset(id);
-
-	const char *addr_space = nullptr;
-	switch (type.storage)
-	{
-	case StorageClassWorkgroup:
-		addr_space = "threadgroup";
-		break;
-
-	case StorageClassStorageBuffer:
-		addr_space = flags.get(DecorationNonWritable) ? "const device" : "device";
-		break;
-
-	case StorageClassUniform:
-	case StorageClassUniformConstant:
-	case StorageClassPushConstant:
-		if (type.basetype == SPIRType::Struct)
-		{
-			bool ssbo = has_decoration(type.self, DecorationBufferBlock);
-			if (ssbo)
-				addr_space = flags.get(DecorationNonWritable) ? "const device" : "device";
-			else
-				addr_space = "constant";
-		}
-		else
-			addr_space = "constant";
-		break;
-
-	case StorageClassFunction:
-	case StorageClassGeneric:
 		// No address space for plain values.
-		addr_space = type.pointer ? "thread" : "";
-		break;
-
-	case StorageClassOutput:
-		if (capture_output_to_buffer)
-			addr_space = "device";
-		break;
-
-	default:
-		break;
-	}
-
-	if (!addr_space)
-		addr_space = "thread";
+		addr_space = type.pointer || (argument && type.basetype == SPIRType::ControlPointArray) ? "thread" : "";
 
 	return join(flags.get(DecorationVolatile) || flags.get(DecorationCoherent) ? "volatile " : "", addr_space);
 }
@@ -6748,6 +6764,19 @@ void CompilerMSL::entry_point_args_builtin(string &ep_args)
 				ep_args += "]]";
 			}
 		}
+
+		if (var.storage == StorageClassInput &&
+		    has_extended_decoration(var_id, SPIRVCrossDecorationBuiltInDispatchBase))
+		{
+			// This is a special implicit builtin, not corresponding to any SPIR-V builtin,
+			// which holds the base that was passed to vkCmdDispatchBase(). If it's present,
+			// assume we emitted it for a good reason.
+			assert(msl_options.supports_msl_version(1, 2));
+			if (!ep_args.empty())
+				ep_args += ", ";
+
+			ep_args += type_to_glsl(get_variable_data_type(var)) + " " + to_expression(var_id) + " [[grid_origin]]";
+		}
 	});
 
 	// Correct the types of all encountered active builtins. We couldn't do this before
@@ -7023,7 +7052,11 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args)
 		default:
 			if (!ep_args.empty())
 				ep_args += ", ";
-			ep_args += type_to_glsl(type, var_id) + " " + r.name;
+			if (!type.pointer)
+				ep_args += get_type_address_space(get<SPIRType>(var.basetype), var_id) + " " +
+				           type_to_glsl(type, var_id) + "& " + r.name;
+			else
+				ep_args += type_to_glsl(type, var_id) + " " + r.name;
 			ep_args += " [[buffer(" + convert_to_string(r.index) + ")]]";
 			break;
 		}
@@ -7343,6 +7376,35 @@ void CompilerMSL::fix_up_shader_inputs_outputs()
 					          msl_options.device_index, ";");
 				});
 				break;
+			case BuiltInWorkgroupId:
+				if (!msl_options.dispatch_base || !active_input_builtins.get(BuiltInWorkgroupId))
+					break;
+
+				// The vkCmdDispatchBase() command lets the client set the base value
+				// of WorkgroupId. Metal has no direct equivalent; we must make this
+				// adjustment ourselves.
+				entry_func.fixup_hooks_in.push_back([=]() {
+					statement(to_expression(var_id), " += ", to_dereferenced_expression(builtin_dispatch_base_id), ";");
+				});
+				break;
+			case BuiltInGlobalInvocationId:
+				if (!msl_options.dispatch_base || !active_input_builtins.get(BuiltInGlobalInvocationId))
+					break;
+
+				// GlobalInvocationId is defined as LocalInvocationId + WorkgroupId * WorkgroupSize.
+				// This needs to be adjusted too.
+				entry_func.fixup_hooks_in.push_back([=]() {
+					auto &execution = this->get_entry_point();
+					uint32_t workgroup_size_id = execution.workgroup_size.constant;
+					if (workgroup_size_id)
+						statement(to_expression(var_id), " += ", to_dereferenced_expression(builtin_dispatch_base_id),
+						          " * ", to_expression(workgroup_size_id), ";");
+					else
+						statement(to_expression(var_id), " += ", to_dereferenced_expression(builtin_dispatch_base_id),
+						          " * uint3(", execution.workgroup_size.x, ", ", execution.workgroup_size.y, ", ",
+						          execution.workgroup_size.z, ");");
+				});
+				break;
 			default:
 				break;
 			}

3rdparty/spirv-cross/spirv_msl.hpp

@@ -198,6 +198,7 @@ public:
 		bool tess_domain_origin_lower_left = false;
 		bool multiview = false;
 		bool view_index_from_device_index = false;
+		bool dispatch_base = false;
 
 		// Enable use of MSL 2.0 indirect argument buffers.
 		// MSL 2.0 must also be enabled.
@@ -225,7 +226,7 @@ public:
 			msl_version = make_msl_version(major, minor, patch);
 		}
 
-		bool supports_msl_version(uint32_t major, uint32_t minor = 0, uint32_t patch = 0)
+		bool supports_msl_version(uint32_t major, uint32_t minor = 0, uint32_t patch = 0) const
 		{
 			return msl_version >= make_msl_version(major, minor, patch);
 		}
@@ -276,6 +277,13 @@ public:
 		return msl_options.multiview && !msl_options.view_index_from_device_index;
 	}
 
+	// Provide feedback to calling API to allow it to pass a buffer
+	// containing the dispatch base workgroup ID.
+	bool needs_dispatch_base_buffer() const
+	{
+		return msl_options.dispatch_base && !msl_options.supports_msl_version(1, 2);
+	}
+
 	// Provide feedback to calling API to allow it to pass an output
 	// buffer if the shader needs it.
 	bool needs_output_buffer() const
@@ -533,7 +541,7 @@ protected:
 	void ensure_member_packing_rules_msl(SPIRType &ib_type, uint32_t index);
 	bool validate_member_packing_rules_msl(const SPIRType &type, uint32_t index) const;
 	std::string get_argument_address_space(const SPIRVariable &argument);
-	std::string get_type_address_space(const SPIRType &type, uint32_t id);
+	std::string get_type_address_space(const SPIRType &type, uint32_t id, bool argument = false);
 	const char *to_restrict(uint32_t id, bool space = true);
 	SPIRType &get_stage_in_struct_type();
 	SPIRType &get_stage_out_struct_type();
@@ -563,6 +571,7 @@ protected:
 	uint32_t builtin_primitive_id_id = 0;
 	uint32_t builtin_subgroup_invocation_id_id = 0;
 	uint32_t builtin_subgroup_size_id = 0;
+	uint32_t builtin_dispatch_base_id = 0;
 	uint32_t swizzle_buffer_id = 0;
 	uint32_t buffer_size_buffer_id = 0;
 	uint32_t view_mask_buffer_id = 0;

3rdparty/spirv-cross/test_shaders.py

@@ -207,6 +207,8 @@ def cross_compile_msl(shader, spirv, opt, iterations, paths):
         msl_args.append('--msl-multiview')
     if '.viewfromdev.' in shader:
         msl_args.append('--msl-view-index-from-device-index')
+    if '.dispatchbase.' in shader:
+        msl_args.append('--msl-dispatch-base')
 
     subprocess.check_call(msl_args)
 
@@ -383,6 +385,8 @@ def cross_compile(shader, vulkan, spirv, invalid_spirv, eliminate, is_legacy, fl
         extra_args += ['--glsl-emit-push-constant-as-ubo']
     if '.line.' in shader:
         extra_args += ['--emit-line-directives']
+    if '.no-samplerless.' in shader:
+        extra_args += ['--vulkan-glsl-disable-ext-samplerless-texture-functions']
 
     spirv_cross_path = paths.spirv_cross