MSL: Drop stores to nonexistent tess levels.

In SPIR-V, there are always two inner levels and four outer levels, even
if the input patch isn't a quad patch. But in MSL, due to requirements
imposed by Metal, only one inner level and three outer levels exist when
the input patch is a triangle patch. We must explicitly ignore any write
to the nonexistent second inner and fourth outer levels in this case.

reference/opt/shaders-msl/asm/tesc/tess-level-overrun.asm.tesc

@@ -0,0 +1,23 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct TessLevels
+{
+    float inner0;
+    float inner1;
+    float outer0;
+    float outer1;
+    float outer2;
+    float outer3;
+};
+
+kernel void main0(const device TessLevels& sb_levels [[buffer(0)]], uint gl_InvocationID [[thread_index_in_threadgroup]], uint gl_PrimitiveID [[threadgroup_position_in_grid]], device uint* spvIndirectParams [[buffer(29)]], device MTLTriangleTessellationFactorsHalf* spvTessLevel [[buffer(26)]])
+{
+    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);
+    spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[2] = half(sb_levels.outer2);
+}
+

reference/shaders-msl/asm/tesc/tess-level-overrun.asm.tesc

@@ -0,0 +1,23 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct TessLevels
+{
+    float inner0;
+    float inner1;
+    float outer0;
+    float outer1;
+    float outer2;
+    float outer3;
+};
+
+kernel void main0(const device TessLevels& sb_levels [[buffer(0)]], uint gl_InvocationID [[thread_index_in_threadgroup]], uint gl_PrimitiveID [[threadgroup_position_in_grid]], device uint* spvIndirectParams [[buffer(29)]], device MTLTriangleTessellationFactorsHalf* spvTessLevel [[buffer(26)]])
+{
+    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);
+    spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[2] = half(sb_levels.outer2);
+}
+

shaders-msl/asm/tesc/tess-level-overrun.asm.tesc

@@ -0,0 +1,102 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Khronos Glslang Reference Front End; 7
+; Bound: 46
+; Schema: 0
+               OpCapability Tessellation
+          %1 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint TessellationControl %main "main" %gl_TessLevelInner %gl_TessLevelOuter
+               OpExecutionMode %main OutputVertices 1
+               OpExecutionMode %main Triangles
+               OpSource ESSL 310
+               OpSourceExtension "GL_EXT_shader_io_blocks"
+               OpSourceExtension "GL_EXT_tessellation_shader"
+               OpName %main "main"
+               OpName %gl_TessLevelInner "gl_TessLevelInner"
+               OpName %TessLevels "TessLevels"
+               OpMemberName %TessLevels 0 "inner0"
+               OpMemberName %TessLevels 1 "inner1"
+               OpMemberName %TessLevels 2 "outer0"
+               OpMemberName %TessLevels 3 "outer1"
+               OpMemberName %TessLevels 4 "outer2"
+               OpMemberName %TessLevels 5 "outer3"
+               OpName %sb_levels "sb_levels"
+               OpName %gl_TessLevelOuter "gl_TessLevelOuter"
+               OpDecorate %gl_TessLevelInner Patch
+               OpDecorate %gl_TessLevelInner BuiltIn TessLevelInner
+               OpMemberDecorate %TessLevels 0 Restrict
+               OpMemberDecorate %TessLevels 0 NonWritable
+               OpMemberDecorate %TessLevels 0 Offset 0
+               OpMemberDecorate %TessLevels 1 Restrict
+               OpMemberDecorate %TessLevels 1 NonWritable
+               OpMemberDecorate %TessLevels 1 Offset 4
+               OpMemberDecorate %TessLevels 2 Restrict
+               OpMemberDecorate %TessLevels 2 NonWritable
+               OpMemberDecorate %TessLevels 2 Offset 8
+               OpMemberDecorate %TessLevels 3 Restrict
+               OpMemberDecorate %TessLevels 3 NonWritable
+               OpMemberDecorate %TessLevels 3 Offset 12
+               OpMemberDecorate %TessLevels 4 Restrict
+               OpMemberDecorate %TessLevels 4 NonWritable
+               OpMemberDecorate %TessLevels 4 Offset 16
+               OpMemberDecorate %TessLevels 5 Restrict
+               OpMemberDecorate %TessLevels 5 NonWritable
+               OpMemberDecorate %TessLevels 5 Offset 20
+               OpDecorate %TessLevels Block
+               OpDecorate %sb_levels DescriptorSet 0
+               OpDecorate %sb_levels Binding 0
+               OpDecorate %gl_TessLevelOuter Patch
+               OpDecorate %gl_TessLevelOuter BuiltIn TessLevelOuter
+       %void = OpTypeVoid
+          %3 = OpTypeFunction %void
+      %float = OpTypeFloat 32
+       %uint = OpTypeInt 32 0
+     %uint_2 = OpConstant %uint 2
+%_arr_float_uint_2 = OpTypeArray %float %uint_2
+%_ptr_Output__arr_float_uint_2 = OpTypePointer Output %_arr_float_uint_2
+%gl_TessLevelInner = OpVariable %_ptr_Output__arr_float_uint_2 Output
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+ %TessLevels = OpTypeStruct %float %float %float %float %float %float
+%_ptr_StorageBuffer_TessLevels = OpTypePointer StorageBuffer %TessLevels
+  %sb_levels = OpVariable %_ptr_StorageBuffer_TessLevels StorageBuffer
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+%_ptr_Output_float = OpTypePointer Output %float
+      %int_1 = OpConstant %int 1
+     %uint_4 = OpConstant %uint 4
+%_arr_float_uint_4 = OpTypeArray %float %uint_4
+%_ptr_Output__arr_float_uint_4 = OpTypePointer Output %_arr_float_uint_4
+%gl_TessLevelOuter = OpVariable %_ptr_Output__arr_float_uint_4 Output
+      %int_2 = OpConstant %int 2
+      %int_3 = OpConstant %int 3
+      %int_4 = OpConstant %int 4
+      %int_5 = OpConstant %int 5
+       %main = OpFunction %void None %3
+          %5 = OpLabel
+         %18 = OpAccessChain %_ptr_StorageBuffer_float %sb_levels %int_0
+         %19 = OpLoad %float %18
+         %21 = OpAccessChain %_ptr_Output_float %gl_TessLevelInner %int_0
+               OpStore %21 %19
+         %23 = OpAccessChain %_ptr_StorageBuffer_float %sb_levels %int_1
+         %24 = OpLoad %float %23
+         %25 = OpAccessChain %_ptr_Output_float %gl_TessLevelInner %int_1
+               OpStore %25 %24
+         %31 = OpAccessChain %_ptr_StorageBuffer_float %sb_levels %int_2
+         %32 = OpLoad %float %31
+         %33 = OpAccessChain %_ptr_Output_float %gl_TessLevelOuter %int_0
+               OpStore %33 %32
+         %35 = OpAccessChain %_ptr_StorageBuffer_float %sb_levels %int_3
+         %36 = OpLoad %float %35
+         %37 = OpAccessChain %_ptr_Output_float %gl_TessLevelOuter %int_1
+               OpStore %37 %36
+         %39 = OpAccessChain %_ptr_StorageBuffer_float %sb_levels %int_4
+         %40 = OpLoad %float %39
+         %41 = OpAccessChain %_ptr_Output_float %gl_TessLevelOuter %int_2
+               OpStore %41 %40
+         %43 = OpAccessChain %_ptr_StorageBuffer_float %sb_levels %int_5
+         %44 = OpLoad %float %43
+         %45 = OpAccessChain %_ptr_Output_float %gl_TessLevelOuter %int_3
+               OpStore %45 %44
+               OpReturn
+               OpFunctionEnd

spirv_msl.cpp

@@ -3099,10 +3099,14 @@ bool CompilerMSL::emit_tessellation_access_chain(const uint32_t *ops, uint32_t l
 	// drop the last index. It isn't an array in this case, so we can't have an
 	// array reference here. We need to make this ID a variable instead of an
 	// expression so we don't try to dereference it as a variable pointer.
+	// Don't do this if the index is a constant 1, though. We need to drop stores
+	// to that one.
 	auto *m = ir.find_meta(var ? var->self : 0);
 	if (get_execution_model() == ExecutionModelTessellationControl && var && m &&
 	    m->decoration.builtin_type == BuiltInTessLevelInner && get_entry_point().flags.get(ExecutionModeTriangles))
 	{
+		if (auto *c = maybe_get<SPIRConstant>(ops[3]))
+			if (c->scalar() == 1) return false;
 		auto &dest_var = set<SPIRVariable>(ops[1], *var);
 		dest_var.basetype = ops[0];
 		ir.meta[ops[1]] = ir.meta[ops[2]];
@@ -3113,6 +3117,28 @@ bool CompilerMSL::emit_tessellation_access_chain(const uint32_t *ops, uint32_t l
 	return false;
 }
 
+bool CompilerMSL::is_out_of_bounds_tessellation_level(uint32_t id_lhs) {
+	if (!get_entry_point().flags.get(ExecutionModeTriangles))
+		return false;
+
+	// In SPIR-V, TessLevelInner always has two elements and TessLevelOuter always has
+	// four. This is true even if we are tessellating triangles. This allows clients
+	// to use a single tessellation control shader with multiple tessellation evaluation
+	// shaders.
+	// In Metal, however, only the first element of TessLevelInner and the first three
+	// of TessLevelOuter are accessible. This stems from how in Metal, the tessellation
+	// levels must be stored to a dedicated buffer in a particular format that depends
+	// on the patch type. Therefore, in Triangles mode, any access to the second
+	// inner level or the fourth outer level must be dropped.
+	const auto *e = maybe_get<SPIRExpression>(id_lhs);
+	if (!e || !e->access_chain) return false;
+	BuiltIn builtin = BuiltIn(get_decoration(e->loaded_from, DecorationBuiltIn));
+	if (builtin != BuiltInTessLevelInner && builtin != BuiltInTessLevelOuter) return false;
+	auto *c = maybe_get<SPIRConstant>(e->implied_read_expressions[1]);
+	if (!c) return false;
+	return (builtin == BuiltInTessLevelInner && c->scalar() == 1) || (builtin == BuiltInTessLevelOuter && c->scalar() == 3);
+}
+
 // Override for MSL-specific syntax instructions
 void CompilerMSL::emit_instruction(const Instruction &instruction)
 {
@@ -3594,6 +3620,9 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
 		break;
 
 	case OpStore:
+		if (is_out_of_bounds_tessellation_level(ops[0]))
+			break;
+
 		if (maybe_emit_array_assignment(ops[0], ops[1]))
 			break;
 

spirv_msl.hpp

@@ -489,6 +489,7 @@ protected:
 	void analyze_sampled_image_usage();
 
 	bool emit_tessellation_access_chain(const uint32_t *ops, uint32_t length);
+	bool is_out_of_bounds_tessellation_level(uint32_t id_lhs);
 
 	Options msl_options;
 	std::set<SPVFuncImpl> spv_function_implementations;