MSL: Fix '--msl-multi-patch-workgroup' out of bounds reads when dispatching more threads than control points (#1662)

* Fix '--msl-multi-patch-workgroup' cases where thread count exceeds data bounds

*Fix gl_PrimitiveID off by one error when computing last valid index
*Point gl_out to the last patch's data when threads exceed input data bounds
*Point patchOut to the last patch's data when threads exceed input data bounds

* Update MSL test expectations.

* Undo change to MSL multi-patch hull output bound checks

* Update MSL multi-patch test expectations.

20 files changed, 20 insertions, 20 deletions

diff --git a/reference/opt/shaders-msl/asm/tesc/tess-level-overrun.multi-patch.asm.tesc b/reference/opt/shaders-msl/asm/tesc/tess-level-overrun.multi-patch.asm.tesc
index a5f316dd..79395a4b 100644
--- a/reference/opt/shaders-msl/asm/tesc/tess-level-overrun.multi-patch.asm.tesc
+++ b/reference/opt/shaders-msl/asm/tesc/tess-level-overrun.multi-patch.asm.tesc
@@ -15,7 +15,7 @@ struct TessLevels
 
 kernel void main0(const device TessLevels& sb_levels [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], constant uint* spvIndirectParams [[buffer(29)]], device MTLTriangleTessellationFactorsHalf* spvTessLevel [[buffer(26)]])
 {
-    uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 1, spvIndirectParams[1]);
+    uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 1, spvIndirectParams[1] - 1);
     spvTessLevel[gl_PrimitiveID].insideTessellationFactor = half(sb_levels.inner0);
     spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0] = half(sb_levels.outer0);
     spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[1] = half(sb_levels.outer1);
diff --git a/reference/opt/shaders-msl/desktop-only/tesc/basic.desktop.sso.multi-patch.tesc b/reference/opt/shaders-msl/desktop-only/tesc/basic.desktop.sso.multi-patch.tesc
index 863a32a8..8ebde9d9 100644
--- a/reference/opt/shaders-msl/desktop-only/tesc/basic.desktop.sso.multi-patch.tesc
+++ b/reference/opt/shaders-msl/desktop-only/tesc/basic.desktop.sso.multi-patch.tesc
@@ -26,7 +26,7 @@ kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], devic
     device main0_patchOut& patchOut = spvPatchOut[gl_GlobalInvocationID.x / 1];
     device main0_in* gl_in = &spvIn[min(gl_GlobalInvocationID.x / 1, spvIndirectParams[1] - 1) * spvIndirectParams[0]];
     uint gl_InvocationID = gl_GlobalInvocationID.x % 1;
-    uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 1, spvIndirectParams[1]);
+    uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 1, spvIndirectParams[1] - 1);
     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.multi-patch.tesc b/reference/opt/shaders-msl/desktop-only/tesc/struct-copy.desktop.sso.multi-patch.tesc
index c5e309ad..184a4a6f 100644
--- a/reference/opt/shaders-msl/desktop-only/tesc/struct-copy.desktop.sso.multi-patch.tesc
+++ b/reference/opt/shaders-msl/desktop-only/tesc/struct-copy.desktop.sso.multi-patch.tesc
@@ -24,7 +24,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_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]);
+    uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 4, spvIndirectParams[1] - 1);
     gl_out[gl_InvocationID].vVertex = gl_in[gl_InvocationID].vInput;
     spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0] = half(1.0);
     spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[1] = half(2.0);
diff --git a/reference/opt/shaders-msl/masking/copy-arrays.mask-location-0.msl2.multi-patch.tesc b/reference/opt/shaders-msl/masking/copy-arrays.mask-location-0.msl2.multi-patch.tesc
index 4f9134ee..406feed9 100644
--- a/reference/opt/shaders-msl/masking/copy-arrays.mask-location-0.msl2.multi-patch.tesc
+++ b/reference/opt/shaders-msl/masking/copy-arrays.mask-location-0.msl2.multi-patch.tesc
@@ -176,7 +176,7 @@ kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], devic
     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]);
+    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);
     if (gl_InvocationID == 0)
diff --git a/reference/opt/shaders-msl/masking/copy-arrays.mask-location-1.msl2.multi-patch.tesc b/reference/opt/shaders-msl/masking/copy-arrays.mask-location-1.msl2.multi-patch.tesc
index a2ad0102..abc95ca8 100644
--- a/reference/opt/shaders-msl/masking/copy-arrays.mask-location-1.msl2.multi-patch.tesc
+++ b/reference/opt/shaders-msl/masking/copy-arrays.mask-location-1.msl2.multi-patch.tesc
@@ -67,7 +67,7 @@ kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], devic
     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]);
+    uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 4, spvIndirectParams[1] - 1);
     gl_out[gl_InvocationID].gl_Position = float4(1.0);
     gl_out[gl_InvocationID].Foo = gl_in[gl_InvocationID].iFoo;
     if (gl_InvocationID == 0)
diff --git a/reference/opt/shaders-msl/masking/write-outputs-block.mask-location-0.multi-patch.msl2.tesc b/reference/opt/shaders-msl/masking/write-outputs-block.mask-location-0.multi-patch.msl2.tesc
index ca025cdb..ef9877b1 100644
--- a/reference/opt/shaders-msl/masking/write-outputs-block.mask-location-0.multi-patch.msl2.tesc
+++ b/reference/opt/shaders-msl/masking/write-outputs-block.mask-location-0.multi-patch.msl2.tesc
@@ -34,7 +34,7 @@ kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], devic
     threadgroup P spvStorage_11[8];
     threadgroup P (&_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]);
+    uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 4, spvIndirectParams[1] - 1);
     _11.a = 1.0;
     patchOut.P_b = 2.0;
     gl_out[gl_InvocationID].C_a = 3.0;
diff --git a/reference/opt/shaders-msl/masking/write-outputs-block.mask-location-1.multi-patch.msl2.tesc b/reference/opt/shaders-msl/masking/write-outputs-block.mask-location-1.multi-patch.msl2.tesc
index 700e3fc5..7b72ded6 100644
--- a/reference/opt/shaders-msl/masking/write-outputs-block.mask-location-1.multi-patch.msl2.tesc
+++ b/reference/opt/shaders-msl/masking/write-outputs-block.mask-location-1.multi-patch.msl2.tesc
@@ -34,7 +34,7 @@ kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], devic
     threadgroup C (&c)[4] = 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]);
+    uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 4, spvIndirectParams[1] - 1);
     patchOut.P_a = 1.0;
     patchOut.P_b = 2.0;
     c[gl_InvocationID].a = 3.0;
diff --git a/reference/opt/shaders-msl/masking/write-outputs.mask-location-0.multi-patch.tesc b/reference/opt/shaders-msl/masking/write-outputs.mask-location-0.multi-patch.tesc
index d20b7d78..7a5e183a 100644
--- a/reference/opt/shaders-msl/masking/write-outputs.mask-location-0.multi-patch.tesc
+++ b/reference/opt/shaders-msl/masking/write-outputs.mask-location-0.multi-patch.tesc
@@ -63,7 +63,7 @@ kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], devic
     threadgroup float4 (&v0)[4] = 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]);
+    uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 4, spvIndirectParams[1] - 1);
     v0[gl_InvocationID] = float4(1.0);
     v0[gl_InvocationID].z = 3.0;
     if (gl_InvocationID == 0)
diff --git a/reference/opt/shaders-msl/masking/write-outputs.mask-location-1.multi-patch.tesc b/reference/opt/shaders-msl/masking/write-outputs.mask-location-1.multi-patch.tesc
index 2831008f..28ec0be0 100644
--- a/reference/opt/shaders-msl/masking/write-outputs.mask-location-1.multi-patch.tesc
+++ b/reference/opt/shaders-msl/masking/write-outputs.mask-location-1.multi-patch.tesc
@@ -22,7 +22,7 @@ kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], devic
     threadgroup float4 spvStoragev1[8][2];
     threadgroup float4 (&v1)[2] = spvStoragev1[(gl_GlobalInvocationID.x / 4) % 8];
     uint gl_InvocationID = gl_GlobalInvocationID.x % 4;
-    uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 4, spvIndirectParams[1]);
+    uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 4, spvIndirectParams[1] - 1);
     gl_out[gl_InvocationID].v0 = float4(1.0);
     gl_out[gl_InvocationID].v0.z = 3.0;
     if (gl_InvocationID == 0)
diff --git a/reference/opt/shaders-msl/masking/write-outputs.mask-point-size.multi-patch.tesc b/reference/opt/shaders-msl/masking/write-outputs.mask-point-size.multi-patch.tesc
index 3b336780..694cdbb7 100644
--- a/reference/opt/shaders-msl/masking/write-outputs.mask-point-size.multi-patch.tesc
+++ b/reference/opt/shaders-msl/masking/write-outputs.mask-point-size.multi-patch.tesc
@@ -71,7 +71,7 @@ kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], devic
     threadgroup gl_PerVertex (&gl_out_masked)[4] = 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]);
+    uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 4, spvIndirectParams[1] - 1);
     gl_out[gl_InvocationID].v0 = float4(1.0);
     gl_out[gl_InvocationID].v0.z = 3.0;
     if (gl_InvocationID == 0)
diff --git a/reference/opt/shaders-msl/masking/write-outputs.mask-position.multi-patch.tesc b/reference/opt/shaders-msl/masking/write-outputs.mask-position.multi-patch.tesc
index 7ef25b9f..c55e3376 100644
--- a/reference/opt/shaders-msl/masking/write-outputs.mask-position.multi-patch.tesc
+++ b/reference/opt/shaders-msl/masking/write-outputs.mask-position.multi-patch.tesc
@@ -71,7 +71,7 @@ kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], devic
     threadgroup gl_PerVertex (&gl_out_masked)[4] = 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]);
+    uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 4, spvIndirectParams[1] - 1);
     gl_out[gl_InvocationID].v0 = float4(1.0);
     gl_out[gl_InvocationID].v0.z = 3.0;
     if (gl_InvocationID == 0)
diff --git a/reference/opt/shaders-msl/tesc/basic.multi-patch.tesc b/reference/opt/shaders-msl/tesc/basic.multi-patch.tesc
index b1403a3f..fe268316 100644
--- a/reference/opt/shaders-msl/tesc/basic.multi-patch.tesc
+++ b/reference/opt/shaders-msl/tesc/basic.multi-patch.tesc
@@ -11,7 +11,7 @@ struct main0_patchOut
 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_patchOut& patchOut = spvPatchOut[gl_GlobalInvocationID.x / 1];
-    uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 1, spvIndirectParams[1]);
+    uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 1, spvIndirectParams[1] - 1);
     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/tesc/complex-control-point-inout-types.multi-patch.tesc b/reference/opt/shaders-msl/tesc/complex-control-point-inout-types.multi-patch.tesc
index b8f4235c..d3b4738d 100644
--- a/reference/opt/shaders-msl/tesc/complex-control-point-inout-types.multi-patch.tesc
+++ b/reference/opt/shaders-msl/tesc/complex-control-point-inout-types.multi-patch.tesc
@@ -102,7 +102,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_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]);
+    uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 4, spvIndirectParams[1] - 1);
     gl_out[gl_InvocationID].gl_Position = float4(1.0);
     gl_out[gl_InvocationID].a[0] = gl_in[gl_InvocationID].in_a[0];
     gl_out[gl_InvocationID].a[1] = gl_in[gl_InvocationID].in_a[1];
diff --git a/reference/opt/shaders-msl/tesc/load-control-point-array-of-matrix.multi-patch.tesc b/reference/opt/shaders-msl/tesc/load-control-point-array-of-matrix.multi-patch.tesc
index 4f4cf0b0..28effad2 100644
--- a/reference/opt/shaders-msl/tesc/load-control-point-array-of-matrix.multi-patch.tesc
+++ b/reference/opt/shaders-msl/tesc/load-control-point-array-of-matrix.multi-patch.tesc
@@ -59,7 +59,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_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]);
+    uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 4, spvIndirectParams[1] - 1);
     spvUnsafeArray<float4x4, 32> _16 = spvUnsafeArray<float4x4, 32>({ gl_in[0].vInputs, gl_in[1].vInputs, gl_in[2].vInputs, gl_in[3].vInputs, gl_in[4].vInputs, gl_in[5].vInputs, gl_in[6].vInputs, gl_in[7].vInputs, gl_in[8].vInputs, gl_in[9].vInputs, gl_in[10].vInputs, gl_in[11].vInputs, gl_in[12].vInputs, gl_in[13].vInputs, gl_in[14].vInputs, gl_in[15].vInputs, gl_in[16].vInputs, gl_in[17].vInputs, gl_in[18].vInputs, gl_in[19].vInputs, gl_in[20].vInputs, gl_in[21].vInputs, gl_in[22].vInputs, gl_in[23].vInputs, gl_in[24].vInputs, gl_in[25].vInputs, gl_in[26].vInputs, gl_in[27].vInputs, gl_in[28].vInputs, gl_in[29].vInputs, gl_in[30].vInputs, gl_in[31].vInputs });
     spvUnsafeArray<float4x4, 32> tmp;
     tmp = _16;
diff --git a/reference/opt/shaders-msl/tesc/load-control-point-array-of-struct.multi-patch.tesc b/reference/opt/shaders-msl/tesc/load-control-point-array-of-struct.multi-patch.tesc
index add59f69..e6bd83e8 100644
--- a/reference/opt/shaders-msl/tesc/load-control-point-array-of-struct.multi-patch.tesc
+++ b/reference/opt/shaders-msl/tesc/load-control-point-array-of-struct.multi-patch.tesc
@@ -66,7 +66,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_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]);
+    uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 4, spvIndirectParams[1] - 1);
     spvUnsafeArray<VertexData, 32> _19 = spvUnsafeArray<VertexData, 32>({ gl_in[0].vInputs, gl_in[1].vInputs, gl_in[2].vInputs, gl_in[3].vInputs, gl_in[4].vInputs, gl_in[5].vInputs, gl_in[6].vInputs, gl_in[7].vInputs, gl_in[8].vInputs, gl_in[9].vInputs, gl_in[10].vInputs, gl_in[11].vInputs, gl_in[12].vInputs, gl_in[13].vInputs, gl_in[14].vInputs, gl_in[15].vInputs, gl_in[16].vInputs, gl_in[17].vInputs, gl_in[18].vInputs, gl_in[19].vInputs, gl_in[20].vInputs, gl_in[21].vInputs, gl_in[22].vInputs, gl_in[23].vInputs, gl_in[24].vInputs, gl_in[25].vInputs, gl_in[26].vInputs, gl_in[27].vInputs, gl_in[28].vInputs, gl_in[29].vInputs, gl_in[30].vInputs, gl_in[31].vInputs });
     spvUnsafeArray<VertexData, 32> tmp;
     tmp = _19;
diff --git a/reference/opt/shaders-msl/tesc/load-control-point-array.multi-patch.tesc b/reference/opt/shaders-msl/tesc/load-control-point-array.multi-patch.tesc
index 416bc681..45baadb6 100644
--- a/reference/opt/shaders-msl/tesc/load-control-point-array.multi-patch.tesc
+++ b/reference/opt/shaders-msl/tesc/load-control-point-array.multi-patch.tesc
@@ -60,7 +60,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_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]);
+    uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 4, spvIndirectParams[1] - 1);
     spvUnsafeArray<float4, 32> _15 = spvUnsafeArray<float4, 32>({ gl_in[0].vInputs, gl_in[1].vInputs, gl_in[2].vInputs, gl_in[3].vInputs, gl_in[4].vInputs, gl_in[5].vInputs, gl_in[6].vInputs, gl_in[7].vInputs, gl_in[8].vInputs, gl_in[9].vInputs, gl_in[10].vInputs, gl_in[11].vInputs, gl_in[12].vInputs, gl_in[13].vInputs, gl_in[14].vInputs, gl_in[15].vInputs, gl_in[16].vInputs, gl_in[17].vInputs, gl_in[18].vInputs, gl_in[19].vInputs, gl_in[20].vInputs, gl_in[21].vInputs, gl_in[22].vInputs, gl_in[23].vInputs, gl_in[24].vInputs, gl_in[25].vInputs, gl_in[26].vInputs, gl_in[27].vInputs, gl_in[28].vInputs, gl_in[29].vInputs, gl_in[30].vInputs, gl_in[31].vInputs });
     spvUnsafeArray<float4, 32> tmp;
     tmp = _15;
diff --git a/reference/opt/shaders-msl/tesc/matrix-output.multi-patch.tesc b/reference/opt/shaders-msl/tesc/matrix-output.multi-patch.tesc
index b3b06d18..28fff015 100644
--- a/reference/opt/shaders-msl/tesc/matrix-output.multi-patch.tesc
+++ b/reference/opt/shaders-msl/tesc/matrix-output.multi-patch.tesc
@@ -21,7 +21,7 @@ kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], devic
     device main0_out* gl_out = &spvOut[gl_GlobalInvocationID.x - gl_GlobalInvocationID.x % 3];
     device main0_in* gl_in = &spvIn[min(gl_GlobalInvocationID.x / 3, spvIndirectParams[1] - 1) * spvIndirectParams[0]];
     uint gl_InvocationID = gl_GlobalInvocationID.x % 3;
-    uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 3, spvIndirectParams[1]);
+    uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 3, spvIndirectParams[1] - 1);
     float _15 = float(gl_InvocationID);
     float3 _18 = float3(_15, 0.0, 0.0);
     float3 _19 = float3(0.0, _15, 0.0);
diff --git a/reference/opt/shaders-msl/tesc/reload-tess-level.multi-patch.tesc b/reference/opt/shaders-msl/tesc/reload-tess-level.multi-patch.tesc
index a55755e0..ae33de51 100644
--- a/reference/opt/shaders-msl/tesc/reload-tess-level.multi-patch.tesc
+++ b/reference/opt/shaders-msl/tesc/reload-tess-level.multi-patch.tesc
@@ -20,7 +20,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_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]);
+    uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 4, spvIndirectParams[1] - 1);
     if (gl_InvocationID == 0)
     {
         spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0] = half(2.0);
diff --git a/reference/opt/shaders-msl/tesc/struct-output.multi-patch.tesc b/reference/opt/shaders-msl/tesc/struct-output.multi-patch.tesc
index eb59973c..6c526546 100644
--- a/reference/opt/shaders-msl/tesc/struct-output.multi-patch.tesc
+++ b/reference/opt/shaders-msl/tesc/struct-output.multi-patch.tesc
@@ -28,7 +28,7 @@ kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], devic
     device main0_out* gl_out = &spvOut[gl_GlobalInvocationID.x - gl_GlobalInvocationID.x % 3];
     device main0_in* gl_in = &spvIn[min(gl_GlobalInvocationID.x / 3, spvIndirectParams[1] - 1) * spvIndirectParams[0]];
     uint gl_InvocationID = gl_GlobalInvocationID.x % 3;
-    uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 3, spvIndirectParams[1]);
+    uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 3, spvIndirectParams[1] - 1);
     float _15 = float(gl_InvocationID);
     int _18 = gl_InvocationID + 1;
     float _19 = float(_18);
diff --git a/reference/opt/shaders-msl/tesc/water_tess.multi-patch.tesc b/reference/opt/shaders-msl/tesc/water_tess.multi-patch.tesc
index 6264a8ef..99e094c9 100644
--- a/reference/opt/shaders-msl/tesc/water_tess.multi-patch.tesc
+++ b/reference/opt/shaders-msl/tesc/water_tess.multi-patch.tesc
@@ -29,7 +29,7 @@ kernel void main0(constant UBO& _41 [[buffer(0)]], uint3 gl_GlobalInvocationID [
 {
     device main0_patchOut& patchOut = spvPatchOut[gl_GlobalInvocationID.x / 1];
     device main0_in* gl_in = &spvIn[min(gl_GlobalInvocationID.x / 1, spvIndirectParams[1] - 1) * spvIndirectParams[0]];
-    uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 1, spvIndirectParams[1]);
+    uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 1, spvIndirectParams[1] - 1);
     float2 _431 = (gl_in[0].vPatchPosBase.xy - float2(10.0)) * _41.uScale.xy;
     float2 _441 = ((gl_in[0].vPatchPosBase.xy + _41.uPatchSize) + float2(10.0)) * _41.uScale.xy;
     float3 _446 = float3(_431.x, -10.0, _431.y);