MSL: Terminate function with return value using return if ending in u…

…nreachable.

reference/shaders-msl/asm/frag/unreachable-return.msl23.spv14.asm.frag

@@ -0,0 +1,26 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct buff_t
+{
+    int m0[1024];
+};
+
+struct main0_out
+{
+    float4 frag_clr [[color(0)]];
+};
+
+fragment main0_out main0(device buff_t& buff [[buffer(0)]], float4 gl_FragCoord [[position]])
+{
+    main0_out out = {};
+    int2 frag_coord = int4(gl_FragCoord).xy;
+    int buff_idx = (frag_coord.y * 32) + frag_coord.x;
+    out.frag_clr = float4(0.0, 0.0, 1.0, 1.0);
+    buff.m0[buff_idx] = 1;
+    discard_fragment();
+    return out;
+}
+

reference/shaders-msl/comp/coherent-block.comp

@@ -10,7 +10,7 @@ struct SSBO
 
 constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
 
-kernel void main0(volatile device SSBO& _10 [[buffer(0)]])
+kernel void main0(volatile device SSBO& __restrict _10 [[buffer(0)]])
 {
     _10.value = float4(20.0);
 }

reference/shaders-msl/comp/coherent-image.comp

@@ -10,7 +10,7 @@ struct SSBO
 
 constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
 
-kernel void main0(volatile device SSBO& _10 [[buffer(0)]], texture2d<int> uImage [[texture(0)]])
+kernel void main0(volatile device SSBO& __restrict _10 [[buffer(0)]], texture2d<int> uImage [[texture(0)]])
 {
     _10.value = uImage.read(uint2(int2(10)));
 }

reference/shaders-msl/comp/read-write-only.comp

@@ -23,7 +23,7 @@ struct SSBO1
 
 constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
 
-kernel void main0(device SSBO2& _10 [[buffer(0)]], const device SSBO0& _15 [[buffer(1)]], device SSBO1& _21 [[buffer(2)]])
+kernel void main0(device SSBO2& __restrict _10 [[buffer(0)]], const device SSBO0& _15 [[buffer(1)]], device SSBO1& __restrict _21 [[buffer(2)]])
 {
     _10.data4 = _15.data0 + _21.data2;
     _10.data5 = _15.data1 + _21.data3;

reference/shaders-msl/frag/in_mat.frag

@@ -29,7 +29,7 @@ fragment main0_out main0(main0_in in [[stage_in]], texturecube<float> samplerCol
     inInvModelView[3] = in.inInvModelView_3;
     float3 cI = fast::normalize(in.inPos);
     float3 cR = reflect(cI, fast::normalize(in.inNormal));
-    cR = float3((inInvModelView * float4(cR, 0.0)).xyz);
+    cR = (inInvModelView * float4(cR, 0.0)).xyz;
     cR.x *= (-1.0);
     out.outFragColor = samplerColor.sample(samplerColorSmplr, cR, bias(in.inLodBias));
     return out;

reference/shaders-msl/frag/return-value-after-discard-terminator.frag

@@ -16,7 +16,7 @@ struct main0_out
 fragment main0_out main0(device buff_t& buff [[buffer(0)]], float4 gl_FragCoord [[position]])
 {
     main0_out out = {};
-    int2 frag_coord = int2(int4(gl_FragCoord).xy);
+    int2 frag_coord = int4(gl_FragCoord).xy;
     int buff_idx = (frag_coord.y * 32) + frag_coord.x;
     out.frag_clr = float4(0.0, 0.0, 1.0, 1.0);
     buff.m0[buff_idx] = 1;

reference/shaders-msl/vert/in_out_array_mat.vert

@@ -105,7 +105,7 @@ vertex main0_out main0(main0_in in [[stage_in]], constant UBO& ubo [[buffer(0)]]
     inViewMat[2] = in.inViewMat_2;
     inViewMat[3] = in.inViewMat_3;
     out.gl_Position = (ubo.projection * ubo.model) * float4(in.inPos, 1.0);
-    out.outPos = float3((ubo.model * float4(in.inPos, 1.0)).xyz);
+    out.outPos = (ubo.model * float4(in.inPos, 1.0)).xyz;
     out.outNormal = float3x3(float3(float3(ubo.model[0].x, ubo.model[0].y, ubo.model[0].z)), float3(float3(ubo.model[1].x, ubo.model[1].y, ubo.model[1].z)), float3(float3(ubo.model[2].x, ubo.model[2].y, ubo.model[2].z))) * in.inNormal;
     out.outLodBias = ubo.lodBias;
     outTransModel = transpose(ubo.model) * inViewMat;

shaders-msl/asm/frag/unreachable-return.msl23.spv14.asm.frag

@@ -0,0 +1,78 @@
+; SPIR-V
+; Version: 1.5
+; Generator: Khronos Glslang Reference Front End; 11
+; Bound: 46
+; Schema: 0
+               OpCapability Shader
+               OpCapability DemoteToHelperInvocation
+               OpExtension "SPV_EXT_demote_to_helper_invocation"
+          %1 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %main "main" %gl_FragCoord %frag_clr %buff
+               OpExecutionMode %main OriginUpperLeft
+               OpSource GLSL 450
+               OpName %main "main"
+               OpName %frag_coord "frag_coord"
+               OpName %gl_FragCoord "gl_FragCoord"
+               OpName %buff_idx "buff_idx"
+               OpName %frag_clr "frag_clr"
+               OpName %buff_t "buff_t"
+               OpMemberName %buff_t 0 "m0"
+               OpName %buff "buff"
+               OpDecorate %gl_FragCoord BuiltIn FragCoord
+               OpDecorate %frag_clr Location 0
+               OpDecorate %_arr_int_uint_1024 ArrayStride 4
+               OpDecorate %buff_t Block
+               OpMemberDecorate %buff_t 0 Offset 0
+               OpDecorate %buff Binding 0
+               OpDecorate %buff DescriptorSet 0
+       %void = OpTypeVoid
+          %3 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+%_ptr_Function_v2int = OpTypePointer Function %v2int
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Input_v4float = OpTypePointer Input %v4float
+%gl_FragCoord = OpVariable %_ptr_Input_v4float Input
+      %v4int = OpTypeVector %int 4
+%_ptr_Function_int = OpTypePointer Function %int
+       %uint = OpTypeInt 32 0
+     %uint_1 = OpConstant %uint 1
+     %int_32 = OpConstant %int 32
+     %uint_0 = OpConstant %uint 0
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+   %frag_clr = OpVariable %_ptr_Output_v4float Output
+    %float_0 = OpConstant %float 0
+    %float_1 = OpConstant %float 1
+         %34 = OpConstantComposite %v4float %float_0 %float_0 %float_1 %float_1
+  %uint_1024 = OpConstant %uint 1024
+%_arr_int_uint_1024 = OpTypeArray %int %uint_1024
+     %buff_t = OpTypeStruct %_arr_int_uint_1024
+%_ptr_StorageBuffer_buff_t = OpTypePointer StorageBuffer %buff_t
+       %buff = OpVariable %_ptr_StorageBuffer_buff_t StorageBuffer
+      %int_0 = OpConstant %int 0
+      %int_1 = OpConstant %int 1
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+       %main = OpFunction %void None %3
+          %5 = OpLabel
+ %frag_coord = OpVariable %_ptr_Function_v2int Function
+   %buff_idx = OpVariable %_ptr_Function_int Function
+         %14 = OpLoad %v4float %gl_FragCoord
+         %16 = OpConvertFToS %v4int %14
+         %17 = OpVectorShuffle %v2int %16 %16 0 1
+               OpStore %frag_coord %17
+         %22 = OpAccessChain %_ptr_Function_int %frag_coord %uint_1
+         %23 = OpLoad %int %22
+         %25 = OpIMul %int %23 %int_32
+         %27 = OpAccessChain %_ptr_Function_int %frag_coord %uint_0
+         %28 = OpLoad %int %27
+         %29 = OpIAdd %int %25 %28
+               OpStore %buff_idx %29
+               OpStore %frag_clr %34
+         %41 = OpLoad %int %buff_idx
+         %44 = OpAccessChain %_ptr_StorageBuffer_int %buff %int_0 %41
+               OpStore %44 %int_1
+               OpDemoteToHelperInvocation
+               OpUnreachable
+               OpFunctionEnd

spirv_glsl.cpp

@@ -17869,6 +17869,14 @@ void CompilerGLSL::emit_block_chain(SPIRBlock &block)
 
 	case SPIRBlock::Unreachable:
 	{
+		// If the entry point ends with unreachable and has a return value, insert a return
+		// statement to avoid potential compiler errors from non-void functions without a return value.
+		if (processing_entry_point && block.return_value)
+		{
+			statement("return ", to_unpacked_expression(block.return_value), ";");
+			break;
+		}
+
 		// Avoid emitting false fallthrough, which can happen for
 		// if (cond) break; else discard; inside a case label.
 		// Discard is not always implementable as a terminator.

spirv_msl.cpp

@@ -4219,8 +4219,9 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage, bool patch)
 		// If the entry point should return the output struct, set the entry function
 		// to return the output interface struct, otherwise to return nothing.
 		// Watch out for the rare case where the terminator of the last entry point block is a
-		// Kill, instead of a Return. Based on SPIR-V's block-domination rules, we assume that
-		// any block that has a Kill will also have a terminating Return, except the last block.
+		// Kill or Unreachable, instead of a Return. Based on SPIR-V's block-domination rules,
+		// we assume that any block that has a Kill will also have a terminating Return, except
+		// the last block.
 		// Indicate the output var requires early initialization.
 		bool ep_should_return_output = !get_is_rasterization_disabled();
 		uint32_t rtn_id = ep_should_return_output ? ib_var_id : 0;
@@ -4230,7 +4231,8 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage, bool patch)
 			for (auto &blk_id : entry_func.blocks)
 			{
 				auto &blk = get<SPIRBlock>(blk_id);
-				if (blk.terminator == SPIRBlock::Return || (blk.terminator == SPIRBlock::Kill && blk_id == entry_func.blocks.back()))
+				auto last_blk_return = blk.terminator == SPIRBlock::Kill || blk.terminator == SPIRBlock::Unreachable;
+				if (blk.terminator == SPIRBlock::Return || (last_blk_return && blk_id == entry_func.blocks.back()))
 					blk.return_value = rtn_id;
 			}
 			vars_needing_early_declaration.push_back(ib_var_id);