diff --git a/reference/opt/shaders-msl/asm/tesc/tess-fixed-input-array-builtin-array.invalid.asm.tesc b/reference/opt/shaders-msl/asm/tesc/tess-fixed-input-array-builtin-array.invalid.asm.tesc
index e256409d..0a2bb6bf 100644
--- a/reference/opt/shaders-msl/asm/tesc/tess-fixed-input-array-builtin-array.invalid.asm.tesc
+++ b/reference/opt/shaders-msl/asm/tesc/tess-fixed-input-array-builtin-array.invalid.asm.tesc
@@ -52,6 +52,8 @@ kernel void main0(main0_in in [[stage_in]], uint gl_InvocationID [[thread_index_
     if (gl_InvocationID < spvIndirectParams[0])
         gl_in[gl_InvocationID] = in;
     threadgroup_barrier(mem_flags::mem_threadgroup);
+    if (gl_InvocationID >= 3)
+        return;
     VertexOutput _223[3] = { VertexOutput{ gl_in[0].gl_Position, gl_in[0].VertexOutput_uv }, VertexOutput{ gl_in[1].gl_Position, gl_in[1].VertexOutput_uv }, VertexOutput{ gl_in[2].gl_Position, gl_in[2].VertexOutput_uv } };
     VertexOutput param[3];
     spvArrayCopyFromStack1(param, _223);
diff --git a/reference/opt/shaders-msl/desktop-only/tesc/arrayed-output.desktop.sso.tesc b/reference/opt/shaders-msl/desktop-only/tesc/arrayed-output.desktop.sso.tesc
index b9949290..69949361 100644
--- a/reference/opt/shaders-msl/desktop-only/tesc/arrayed-output.desktop.sso.tesc
+++ b/reference/opt/shaders-msl/desktop-only/tesc/arrayed-output.desktop.sso.tesc
@@ -25,6 +25,8 @@ kernel void main0(main0_in in [[stage_in]], uint gl_InvocationID [[thread_index_
     if (gl_InvocationID < spvIndirectParams[0])
         gl_in[gl_InvocationID] = in;
     threadgroup_barrier(mem_flags::mem_threadgroup);
+    if (gl_InvocationID >= 4)
+        return;
     gl_out[gl_InvocationID].vVertex = gl_in[gl_InvocationID].vInput + gl_in[gl_InvocationID ^ 1].vInput;
     threadgroup_barrier(mem_flags::mem_device);
     if (gl_InvocationID == 0)
diff --git a/reference/opt/shaders-msl/desktop-only/tesc/basic.desktop.sso.tesc b/reference/opt/shaders-msl/desktop-only/tesc/basic.desktop.sso.tesc
index 4d932d58..53093848 100644
--- a/reference/opt/shaders-msl/desktop-only/tesc/basic.desktop.sso.tesc
+++ b/reference/opt/shaders-msl/desktop-only/tesc/basic.desktop.sso.tesc
@@ -25,6 +25,8 @@ kernel void main0(main0_in in [[stage_in]], uint gl_InvocationID [[thread_index_
     if (gl_InvocationID < spvIndirectParams[0])
         gl_in[gl_InvocationID] = in;
     threadgroup_barrier(mem_flags::mem_threadgroup);
+    if (gl_InvocationID >= 1)
+        return;
     spvTessLevel[gl_PrimitiveID].insideTessellationFactor[0] = half(8.8999996185302734375);
     spvTessLevel[gl_PrimitiveID].insideTessellationFactor[1] = half(6.900000095367431640625);
     spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0] = half(8.8999996185302734375);
diff --git a/reference/opt/shaders-msl/desktop-only/tesc/struct-copy.desktop.sso.tesc b/reference/opt/shaders-msl/desktop-only/tesc/struct-copy.desktop.sso.tesc
index 3496d04c..2e63aee1 100644
--- a/reference/opt/shaders-msl/desktop-only/tesc/struct-copy.desktop.sso.tesc
+++ b/reference/opt/shaders-msl/desktop-only/tesc/struct-copy.desktop.sso.tesc
@@ -26,6 +26,8 @@ kernel void main0(main0_in in [[stage_in]], uint gl_InvocationID [[thread_index_
     if (gl_InvocationID < spvIndirectParams[0])
         gl_in[gl_InvocationID] = in;
     threadgroup_barrier(mem_flags::mem_threadgroup);
+    if (gl_InvocationID >= 4)
+        return;
     Boo vInput_24;
     vInput_24.a = gl_in[gl_InvocationID].Boo_a;
     vInput_24.b = gl_in[gl_InvocationID].Boo_b;
diff --git a/reference/opt/shaders-msl/tesc/water_tess.tesc b/reference/opt/shaders-msl/tesc/water_tess.tesc
index e89e6570..b67ae276 100644
--- a/reference/opt/shaders-msl/tesc/water_tess.tesc
+++ b/reference/opt/shaders-msl/tesc/water_tess.tesc
@@ -30,6 +30,8 @@ kernel void main0(main0_in in [[stage_in]], constant UBO& _41 [[buffer(0)]], uin
     if (gl_InvocationID < spvIndirectParams[0])
         gl_in[gl_InvocationID] = in;
     threadgroup_barrier(mem_flags::mem_threadgroup);
+    if (gl_InvocationID >= 1)
+        return;
     float2 _430 = (gl_in[0].vPatchPosBase - float2(10.0)) * _41.uScale.xy;
     float2 _440 = ((gl_in[0].vPatchPosBase + _41.uPatchSize) + float2(10.0)) * _41.uScale.xy;
     float3 _445 = float3(_430.x, -10.0, _430.y);
diff --git a/reference/shaders-msl/asm/tesc/tess-fixed-input-array-builtin-array.invalid.asm.tesc b/reference/shaders-msl/asm/tesc/tess-fixed-input-array-builtin-array.invalid.asm.tesc
index ef9ef7cc..a8e3af26 100644
--- a/reference/shaders-msl/asm/tesc/tess-fixed-input-array-builtin-array.invalid.asm.tesc
+++ b/reference/shaders-msl/asm/tesc/tess-fixed-input-array-builtin-array.invalid.asm.tesc
@@ -82,6 +82,8 @@ kernel void main0(main0_in in [[stage_in]], uint gl_InvocationID [[thread_index_
     if (gl_InvocationID < spvIndirectParams[0])
         gl_in[gl_InvocationID] = in;
     threadgroup_barrier(mem_flags::mem_threadgroup);
+    if (gl_InvocationID >= 3)
+        return;
     VertexOutput p[3];
     p[0].pos = gl_in[0].gl_Position;
     p[0].uv = gl_in[0].VertexOutput_uv;
diff --git a/reference/shaders-msl/desktop-only/tesc/arrayed-output.desktop.sso.tesc b/reference/shaders-msl/desktop-only/tesc/arrayed-output.desktop.sso.tesc
index b9949290..69949361 100644
--- a/reference/shaders-msl/desktop-only/tesc/arrayed-output.desktop.sso.tesc
+++ b/reference/shaders-msl/desktop-only/tesc/arrayed-output.desktop.sso.tesc
@@ -25,6 +25,8 @@ kernel void main0(main0_in in [[stage_in]], uint gl_InvocationID [[thread_index_
     if (gl_InvocationID < spvIndirectParams[0])
         gl_in[gl_InvocationID] = in;
     threadgroup_barrier(mem_flags::mem_threadgroup);
+    if (gl_InvocationID >= 4)
+        return;
     gl_out[gl_InvocationID].vVertex = gl_in[gl_InvocationID].vInput + gl_in[gl_InvocationID ^ 1].vInput;
     threadgroup_barrier(mem_flags::mem_device);
     if (gl_InvocationID == 0)
diff --git a/reference/shaders-msl/desktop-only/tesc/basic.desktop.sso.tesc b/reference/shaders-msl/desktop-only/tesc/basic.desktop.sso.tesc
index 9c4e3b3a..eacb3428 100644
--- a/reference/shaders-msl/desktop-only/tesc/basic.desktop.sso.tesc
+++ b/reference/shaders-msl/desktop-only/tesc/basic.desktop.sso.tesc
@@ -32,6 +32,8 @@ kernel void main0(main0_in in [[stage_in]], uint gl_InvocationID [[thread_index_
     if (gl_InvocationID < spvIndirectParams[0])
         gl_in[gl_InvocationID] = in;
     threadgroup_barrier(mem_flags::mem_threadgroup);
+    if (gl_InvocationID >= 1)
+        return;
     spvTessLevel[gl_PrimitiveID].insideTessellationFactor[0] = half(8.8999996185302734375);
     spvTessLevel[gl_PrimitiveID].insideTessellationFactor[1] = half(6.900000095367431640625);
     spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0] = half(8.8999996185302734375);
diff --git a/reference/shaders-msl/desktop-only/tesc/struct-copy.desktop.sso.tesc b/reference/shaders-msl/desktop-only/tesc/struct-copy.desktop.sso.tesc
index 3496d04c..2e63aee1 100644
--- a/reference/shaders-msl/desktop-only/tesc/struct-copy.desktop.sso.tesc
+++ b/reference/shaders-msl/desktop-only/tesc/struct-copy.desktop.sso.tesc
@@ -26,6 +26,8 @@ kernel void main0(main0_in in [[stage_in]], uint gl_InvocationID [[thread_index_
     if (gl_InvocationID < spvIndirectParams[0])
         gl_in[gl_InvocationID] = in;
     threadgroup_barrier(mem_flags::mem_threadgroup);
+    if (gl_InvocationID >= 4)
+        return;
     Boo vInput_24;
     vInput_24.a = gl_in[gl_InvocationID].Boo_a;
     vInput_24.b = gl_in[gl_InvocationID].Boo_b;
diff --git a/reference/shaders-msl/tesc/water_tess.tesc b/reference/shaders-msl/tesc/water_tess.tesc
index 48b3566f..53eab0a1 100644
--- a/reference/shaders-msl/tesc/water_tess.tesc
+++ b/reference/shaders-msl/tesc/water_tess.tesc
@@ -112,6 +112,8 @@ kernel void main0(main0_in in [[stage_in]], constant UBO& v_41 [[buffer(0)]], ui
     if (gl_InvocationID < spvIndirectParams[0])
         gl_in[gl_InvocationID] = in;
     threadgroup_barrier(mem_flags::mem_threadgroup);
+    if (gl_InvocationID >= 1)
+        return;
     float2 p0 = gl_in[0].vPatchPosBase;
     float2 param = p0;
     if (!frustum_cull(param, v_41))
diff --git a/spirv_msl.cpp b/spirv_msl.cpp
index df30267f..1c720c44 100644
--- a/spirv_msl.cpp
+++ b/spirv_msl.cpp
@@ -1916,6 +1916,8 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage, bool patch)
 				statement("    ", input_wg_var_name, "[", to_expression(builtin_invocation_id_id), "] = ", ib_var_ref,
 				          ";");
 				statement("threadgroup_barrier(mem_flags::mem_threadgroup);");
+				statement("if (", to_expression(builtin_invocation_id_id), " >= ", get_entry_point().output_vertices, ")");
+				statement("    return;");
 			});
 		}
 		break;