Metal system value overhaul. (shader-slang#4308)

aroidzap · Jun 8, 2024 · 65928af · 65928af
1 parent e39ceab
commit 65928af
Show file tree

Hide file tree

Showing 9 changed files with 473 additions and 89 deletions.
diff --git a/source/slang/slang-diagnostic-defs.h b/source/slang/slang-diagnostic-defs.h
@@ -828,8 +828,11 @@ DIAGNOSTIC(55102, Error, invalidTorchKernelParamType, "'$0' is not a valid param
 
 DIAGNOSTIC(55200, Error, unsupportedBuiltinType, "'$0' is not a supported builtin type for the target.")
 DIAGNOSTIC(55201, Error, unsupportedRecursion, "recursion detected in call to '$0', but the current code generation target does not allow recursion.")
+DIAGNOSTIC(55202, Error, systemValueAttributeNotSupported, "system value semantic '$0' is not supported for the current target.")
+DIAGNOSTIC(55203, Error, systemValueTypeIncompatible, "system value semantic '$0' should have type '$1' or convertible to type '$1'.")
 DIAGNOSTIC(56001, Error, unableToAutoMapCUDATypeToHostType, "Could not automatically map '$0' to a host type. Automatic binding generation failed for '$1'")
 
+
 DIAGNOSTIC(57001, Warning, spirvOptFailed, "spirv-opt failed. $0")
 
 // GLSL Compatibility

diff --git a/source/slang/slang-emit-metal.cpp b/source/slang/slang-emit-metal.cpp
@@ -170,8 +170,11 @@ void MetalSourceEmitter::emitFuncParamLayoutImpl(IRInst* param)
             break;
         }
     }
-    if (auto sysSemanticAttr = layout->findSystemValueSemanticAttr())
-       _emitSystemSemantic(sysSemanticAttr->getName(), sysSemanticAttr->getIndex());
+    if (!maybeEmitSystemSemantic(param))
+    {
+        if (auto sysSemanticAttr = layout->findSystemValueSemanticAttr())
+            _emitUserSemantic(sysSemanticAttr->getName(), sysSemanticAttr->getIndex());
+    }
 }
 
 void MetalSourceEmitter::emitParameterGroupImpl(IRGlobalParam* varDecl, IRUniformParameterGroupType* type)
@@ -741,84 +744,24 @@ void MetalSourceEmitter::_emitType(IRType* type, DeclaratorInfo* declarator)
     }
 }
 
-void MetalSourceEmitter::_emitSystemSemantic(UnownedStringSlice semanticName, IRIntegerValue semanticIndex)
+bool MetalSourceEmitter::maybeEmitSystemSemantic(IRInst* inst)
 {
-    if (semanticName.caseInsensitiveEquals(toSlice("SV_POSITION")))
-    {
-        m_writer->emit(" [[position]]");
-    }
-    else if (semanticName.caseInsensitiveEquals(toSlice("SV_VERTEXID")))
-    {
-        m_writer->emit(" [[vertex_id]]");
-    }
-    else if (semanticName.caseInsensitiveEquals(toSlice("SV_INSTANCEID")))
-    {
-        m_writer->emit(" [[instance_id]]");
-    }
-    else if (semanticName.caseInsensitiveEquals(toSlice("SV_Target")))
-    {
-        m_writer->emit(" [[color(");
-        m_writer->emit(semanticIndex);
-        m_writer->emit(")]]");
-    }
-    else if (semanticName.caseInsensitiveEquals(toSlice("SV_PRIMITIVEID")))
+    if (auto sysSemanticDecor = inst->findDecoration<IRTargetSystemValueDecoration>())
     {
-        m_writer->emit(" [[primitive_id]]");
-    }
-    else if (semanticName.caseInsensitiveEquals(toSlice("SV_GROUPID")))
-    {
-        // TODO: not supported by metal.
-        // We need to implement the transformation logic in slang-ir-metal-legalize.cpp
-        // to convert SV_GroupID to something like METAL_threadgroup_position_in_grid.
-    }
-    else if (semanticName.caseInsensitiveEquals(toSlice("SV_GROUPINDEX")))
-    {
-        // TODO: not supported by metal.
-    }
-    else if (semanticName.caseInsensitiveEquals(toSlice("SV_DISPATCHTHREADID")))
-    {
-        m_writer->emit(" [[thread_position_in_grid]]");
-    }
-    else if (semanticName.caseInsensitiveEquals(toSlice("SV_GROUPTHREADID")))
-    {
-        m_writer->emit(" [[thread_position_in_threadgroup]]");
-    }
-    else if (semanticName.caseInsensitiveEquals(toSlice("SV_CLIPDISTANCE")))
-    {
-        m_writer->emit(" [[clip_distance]]");
-    }
-    else if (semanticName.caseInsensitiveEquals(toSlice("SV_RENDERTARGETARRAYINDEX")))
-    {
-        m_writer->emit(" [[render_target_array_index]]");
-    }
-    else if (semanticName.caseInsensitiveEquals(toSlice("SV_VIEWPORTARRAYINDEX")))
-    {
-        m_writer->emit(" [[viewport_array_index]]");
-    }
-    else if (semanticName.caseInsensitiveEquals(toSlice("SV_Depth")))
-    {
-        m_writer->emit(" [[depth(any)]]");
-    }
-    else if (semanticName.caseInsensitiveEquals(toSlice("SV_DepthGreaterEqual")))
-    {
-        m_writer->emit(" [[depth(greater)]]");
-    }
-    else if (semanticName.caseInsensitiveEquals(toSlice("SV_DepthLessEqual")))
-    {
-        m_writer->emit(" [[depth(less)]]");
-    }
-    else if (semanticName.caseInsensitiveEquals(toSlice("SV_Coverage")))
-    {
-        m_writer->emit(" [[sample_mask]]");
-    }
-    else if (semanticName.caseInsensitiveEquals(toSlice("SV_StencilRef")))
-    {
-        m_writer->emit(" [[stencil]]");
+        m_writer->emit(" [[");
+        m_writer->emit(sysSemanticDecor->getSemantic());
+        m_writer->emit("]]");
+        return true;
     }
-    else
+    return false;
+}
+
+void MetalSourceEmitter::_emitUserSemantic(UnownedStringSlice semanticName, IRIntegerValue semanticIndex)
+{
+    if (!semanticName.startsWithCaseInsensitive(toSlice("SV_")))
     {
         m_writer->emit(" [[user(");
-        m_writer->emit(semanticName);
+        m_writer->emit(String(semanticName).toUpper());
         if (semanticIndex != 0)
         {
             m_writer->emit("_");
@@ -834,6 +777,10 @@ void MetalSourceEmitter::emitSemanticsImpl(IRInst* inst, bool allowOffsets)
     if (inst->getOp() == kIROp_StructKey)
     {
         // Only emit [[attribute(n)]] on struct keys.
+
+        if (maybeEmitSystemSemantic(inst))
+            return;
+
         bool hasSemanticFromLayout = false;
         if (auto varLayout = findVarLayout(inst))
         {
@@ -851,7 +798,7 @@ void MetalSourceEmitter::emitSemanticsImpl(IRInst* inst, bool allowOffsets)
                 else if (auto semanticAttr = as<IRSemanticAttr>(attr))
                 {
                     auto semanticName = String(semanticAttr->getName()).toUpper();
-                    _emitSystemSemantic(semanticAttr->getName(), semanticAttr->getIndex());
+                    _emitUserSemantic(semanticAttr->getName(), semanticAttr->getIndex());
                     hasSemanticFromLayout = true;
                 }
             }
@@ -861,7 +808,7 @@ void MetalSourceEmitter::emitSemanticsImpl(IRInst* inst, bool allowOffsets)
         {
             if (auto semanticDecor = inst->findDecoration<IRSemanticDecoration>())
             {
-                _emitSystemSemantic(semanticDecor->getSemanticName(), semanticDecor->getSemanticIndex());
+                _emitUserSemantic(semanticDecor->getSemanticName(), semanticDecor->getSemanticIndex());
             }
         }
     }

diff --git a/source/slang/slang-emit-metal.h b/source/slang/slang-emit-metal.h
@@ -76,7 +76,8 @@ class MetalSourceEmitter : public CLikeSourceEmitter
     void _emitHLSLDecorationSingleInt(const char* name, IRFunc* entryPoint, IRIntLit* val);
 
     void _emitStageAccessSemantic(IRStageAccessDecoration* decoration, const char* name);
-    void _emitSystemSemantic(UnownedStringSlice semanticName, IRIntegerValue semanticIndex);
+    void _emitUserSemantic(UnownedStringSlice semanticName, IRIntegerValue semanticIndex);
+    bool maybeEmitSystemSemantic(IRInst* inst);
 };
 
 }

diff --git a/source/slang/slang-ir-inst-defs.h b/source/slang/slang-ir-inst-defs.h
@@ -715,6 +715,8 @@ INST_RANGE(BindingQuery, GetRegisterIndex, GetRegisterSpace)
     INST_RANGE(TargetSpecificDecoration, TargetDecoration, RequirePreludeDecoration)
     INST(GLSLOuterArrayDecoration,          glslOuterArray,         1, 0)
 
+    INST(TargetSystemValueDecoration,       TargetSystemValue,      2, 0)
+
     INST(InterpolationModeDecoration,       interpolationMode,      1, 0)
     INST(NameHintDecoration,                nameHint,               1, 0)
 

diff --git a/source/slang/slang-ir-insts.h b/source/slang/slang-ir-insts.h
@@ -118,6 +118,19 @@ struct IRTargetDecoration : IRTargetSpecificDefinitionDecoration
     IR_LEAF_ISA(TargetDecoration)
 };
 
+struct IRTargetSystemValueDecoration : IRDecoration
+{
+    enum { kOp = kIROp_TargetSystemValueDecoration };
+    IR_LEAF_ISA(TargetSystemValueDecoration)
+
+    IRStringLit* getSemanticOperand() { return cast<IRStringLit>(getOperand(0)); }
+
+    UnownedStringSlice getSemantic()
+    {
+        return getSemanticOperand()->getStringSlice();
+    }
+};
+
 struct IRTargetIntrinsicDecoration : IRTargetSpecificDefinitionDecoration
 {
     enum { kOp = kIROp_TargetIntrinsicDecoration };
@@ -4351,6 +4364,12 @@ struct IRBuilder
 
     void addHighLevelDeclDecoration(IRInst* value, Decl* decl);
 
+    IRDecoration* addTargetSystemValueDecoration(IRInst* value, UnownedStringSlice sysValName, UInt index = 0)
+    {
+        IRInst* operands[] = { getStringValue(sysValName), getIntValue(getIntType(), index)};
+        return addDecoration(value, kIROp_TargetSystemValueDecoration, operands, SLANG_COUNT_OF(operands));
+    }
+
 //    void addLayoutDecoration(IRInst* value, Layout* layout);
     void addLayoutDecoration(IRInst* value, IRLayout* layout);