diff --git a/reference/shaders-msl-no-opt/frag/volatile-helper-invocation.msl23.spv16.frag b/reference/shaders-msl-no-opt/frag/volatile-helper-invocation.msl23.spv16.frag
index f42aeb876..80399adb9 100644
--- a/reference/shaders-msl-no-opt/frag/volatile-helper-invocation.msl23.spv16.frag
+++ b/reference/shaders-msl-no-opt/frag/volatile-helper-invocation.msl23.spv16.frag
@@ -1,3 +1,5 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
 #include <metal_stdlib>
 #include <simd/simd.h>
 
@@ -8,18 +10,24 @@ struct main0_out
     float FragColor [[color(0)]];
 };
 
+static inline __attribute__((always_inline))
+void func(thread float& FragColor, volatile thread bool& gl_HelperInvocation)
+{
+    bool _14 = gl_HelperInvocation;
+    float _17 = float(_14);
+    FragColor = _17;
+    gl_HelperInvocation = true, discard_fragment();
+    bool _18 = gl_HelperInvocation;
+    float _19 = float(_18);
+    FragColor = _19;
+}
+
 fragment main0_out main0()
 {
     main0_out out = {};
     bool gl_HelperInvocation = {};
     gl_HelperInvocation = simd_is_helper_thread();
-    bool _12 = gl_HelperInvocation;
-    float _15 = float(_12);
-    out.FragColor = _15;
-    gl_HelperInvocation = true, discard_fragment();
-    bool _16 = gl_HelperInvocation;
-    float _17 = float(_16);
-    out.FragColor = _17;
+    func(out.FragColor, gl_HelperInvocation);
     return out;
 }
 
diff --git a/shaders-msl-no-opt/frag/volatile-helper-invocation.msl23.spv16.frag b/shaders-msl-no-opt/frag/volatile-helper-invocation.msl23.spv16.frag
index 9a8d9d20b..44cd4f117 100644
--- a/shaders-msl-no-opt/frag/volatile-helper-invocation.msl23.spv16.frag
+++ b/shaders-msl-no-opt/frag/volatile-helper-invocation.msl23.spv16.frag
@@ -3,9 +3,14 @@
 
 layout(location = 0) out float FragColor;
 
-void main()
+void func()
 {
 	FragColor = float(gl_HelperInvocation);
 	demote;
 	FragColor = float(gl_HelperInvocation);
 }
+
+void main()
+{
+	func();
+}
diff --git a/spirv_msl.cpp b/spirv_msl.cpp
index 383ce688e..3dd941587 100644
--- a/spirv_msl.cpp
+++ b/spirv_msl.cpp
@@ -271,13 +271,14 @@ void CompilerMSL::build_implicit_builtins()
 	     active_input_builtins.get(BuiltInInstanceIndex) || active_input_builtins.get(BuiltInBaseInstance));
 	bool need_local_invocation_index = msl_options.emulate_subgroups && active_input_builtins.get(BuiltInSubgroupId);
 	bool need_workgroup_size = msl_options.emulate_subgroups && active_input_builtins.get(BuiltInNumSubgroups);
+	bool needs_helper_invocation_ = needs_helper_invocation || active_input_builtins.get(BuiltInHelperInvocation);
 	bool force_frag_depth_passthrough =
 	    get_execution_model() == ExecutionModelFragment && !uses_explicit_early_fragment_test() && need_subpass_input &&
 	    msl_options.enable_frag_depth_builtin && msl_options.input_attachment_is_ds_attachment;
 
 	if (need_subpass_input || need_sample_pos || need_subgroup_mask || need_vertex_params || need_tesc_params ||
 	    need_tese_params || need_multiview || need_dispatch_base || need_vertex_base_params || need_grid_params ||
-	    needs_sample_id || needs_subgroup_invocation_id || needs_subgroup_size || needs_helper_invocation ||
+	    needs_sample_id || needs_subgroup_invocation_id || needs_subgroup_size || needs_helper_invocation_ ||
 		has_additional_fixed_sample_mask() || need_local_invocation_index || need_workgroup_size || force_frag_depth_passthrough)
 	{
 		bool has_frag_coord = false;
@@ -457,7 +458,7 @@ void CompilerMSL::build_implicit_builtins()
 				}
 			}
 
-			if (needs_helper_invocation && builtin == BuiltInHelperInvocation)
+			if (needs_helper_invocation_ && builtin == BuiltInHelperInvocation)
 			{
 				builtin_helper_invocation_id = var.self;
 				mark_implicit_builtin(StorageClassInput, BuiltInHelperInvocation, var.self);
@@ -840,7 +841,7 @@ void CompilerMSL::build_implicit_builtins()
 			mark_implicit_builtin(StorageClassOutput, BuiltInSampleMask, var_id);
 		}
 
-		if (!has_helper_invocation && needs_helper_invocation)
+		if (!has_helper_invocation && needs_helper_invocation_)
 		{
 			uint32_t offset = ir.increase_bound_by(3);
 			uint32_t type_id = offset;