From 8095434dc45f36f5790e425024bfbe15bf5eb166 Mon Sep 17 00:00:00 2001 From: Chip Davis Date: Wed, 20 Feb 2019 00:33:46 -0600 Subject: [PATCH] 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. --- .../asm/tesc/tess-level-overrun.asm.tesc | 23 ++++ .../asm/tesc/tess-level-overrun.asm.tesc | 23 ++++ .../asm/tesc/tess-level-overrun.asm.tesc | 102 ++++++++++++++++++ spirv_msl.cpp | 29 +++++ spirv_msl.hpp | 1 + 5 files changed, 178 insertions(+) create mode 100644 reference/opt/shaders-msl/asm/tesc/tess-level-overrun.asm.tesc create mode 100644 reference/shaders-msl/asm/tesc/tess-level-overrun.asm.tesc create mode 100644 shaders-msl/asm/tesc/tess-level-overrun.asm.tesc diff --git a/reference/opt/shaders-msl/asm/tesc/tess-level-overrun.asm.tesc b/reference/opt/shaders-msl/asm/tesc/tess-level-overrun.asm.tesc new file mode 100644 index 00000000..ac15a1fa --- /dev/null +++ b/reference/opt/shaders-msl/asm/tesc/tess-level-overrun.asm.tesc @@ -0,0 +1,23 @@ +#include +#include + +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); +} + diff --git a/reference/shaders-msl/asm/tesc/tess-level-overrun.asm.tesc b/reference/shaders-msl/asm/tesc/tess-level-overrun.asm.tesc new file mode 100644 index 00000000..ac15a1fa --- /dev/null +++ b/reference/shaders-msl/asm/tesc/tess-level-overrun.asm.tesc @@ -0,0 +1,23 @@ +#include +#include + +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); +} + diff --git a/shaders-msl/asm/tesc/tess-level-overrun.asm.tesc b/shaders-msl/asm/tesc/tess-level-overrun.asm.tesc new file mode 100644 index 00000000..b21a2d3d --- /dev/null +++ b/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 diff --git a/spirv_msl.cpp b/spirv_msl.cpp index a191bd05..72b66510 100644 --- a/spirv_msl.cpp +++ b/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(ops[3])) + if (c->scalar() == 1) return false; auto &dest_var = set(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(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(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; diff --git a/spirv_msl.hpp b/spirv_msl.hpp index a98bcec3..017cddf5 100644 --- a/spirv_msl.hpp +++ b/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 spv_function_implementations;