Metal system value overhaul. (#4308)

6 files changed, 370 insertions, 85 deletions

diff --git a/source/slang/slang-diagnostic-defs.h b/source/slang/slang-diagnostic-defs.h
index 4cf6a899d..b011227fd 100644
--- 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
index 96843e286..366c51840 100644
--- 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
index 55ad3d4cb..8b014d604 100644
--- a/source/slang/slang-emit-metal.h
+++ b/source/slang/slang-emit-metal.h
@@ -76,7 +76,8 @@ protected:
     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
index 4d39eb978..37e3d2064 100644
--- 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
index 5670cad47..301604bf3 100644
--- 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 @@ public:
 
     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);
 
diff --git a/source/slang/slang-ir-metal-legalize.cpp b/source/slang/slang-ir-metal-legalize.cpp
index d4a234515..70f4cbd27 100644
--- a/source/slang/slang-ir-metal-legalize.cpp
+++ b/source/slang/slang-ir-metal-legalize.cpp
@@ -50,6 +50,12 @@ namespace Slang
             auto structType = as<IRStructType>(param->getDataType());
             builder.setInsertBefore(func->getFirstBlock()->getFirstOrdinaryInst());
             auto varLayout = findVarLayout(param);
+
+            // If `param` already has a semantic, we don't want to hoist its fields out.
+            if (varLayout->findSystemValueSemanticAttr() != nullptr ||
+                param->findDecoration<IRSemanticDecoration>())
+                continue;
+
             IRStructTypeLayout* structTypeLayout = nullptr;
             if (varLayout)
                 structTypeLayout = as<IRStructTypeLayout>(varLayout->getTypeLayout());
@@ -214,8 +220,168 @@ namespace Slang
         fixUpFuncType(func);
     }
 
+    struct MetalSystemValueInfo
+    {
+        String metalSystemValueName;
+        IRType* requiredType;
+        IRType* altRequiredType;
+        bool isUnsupported;
+        bool isSpecial;
+    };
 
-    void ensureResultStructHasUserSemantic(IRStructType* structType, IRVarLayout* varLayout)
+    MetalSystemValueInfo getSystemValueInfo(IRBuilder& builder, String semanticName, UInt attrIndex)
+    {
+        SLANG_UNUSED(attrIndex);
+
+        MetalSystemValueInfo result = {};
+
+        semanticName = semanticName.toLower();
+
+        if (semanticName == "sv_position")
+        {
+            result.metalSystemValueName = toSlice("position");
+            result.requiredType = builder.getVectorType(builder.getBasicType(BaseType::Float), builder.getIntValue(builder.getIntType(), 4));
+        }
+        else if (semanticName == "sv_clipdistance")
+        {
+            result.isSpecial = true;
+        }
+        else if (semanticName == "sv_culldistance")
+        {
+            result.isSpecial = true;
+        }
+        else if (semanticName == "sv_coverage")
+        {
+            result.metalSystemValueName = toSlice("sample_mask");
+            result.requiredType = builder.getBasicType(BaseType::UInt);
+        }
+        else if (semanticName == "sv_innercoverage")
+        {
+            result.isSpecial = true;
+
+        }
+        else if (semanticName == "sv_depth")
+        {
+            result.metalSystemValueName = toSlice("depth(any)");
+            result.requiredType = builder.getBasicType(BaseType::Float);
+        }
+        else if (semanticName == "sv_depthgreaterequal")
+        {
+            result.metalSystemValueName = toSlice("depth(greater)");
+            result.requiredType = builder.getBasicType(BaseType::Float);
+        }
+        else if (semanticName == "sv_depthlessequal")
+        {
+            result.metalSystemValueName = toSlice("depth(less)");
+            result.requiredType = builder.getBasicType(BaseType::Float);
+        }
+        else if (semanticName == "sv_dispatchthreadid")
+        {
+            result.metalSystemValueName = toSlice("thread_position_in_grid");
+            result.requiredType = builder.getVectorType(builder.getBasicType(BaseType::UInt), builder.getIntValue(builder.getIntType(), 3));
+        }
+        else if (semanticName == "sv_domainlocation")
+        {
+            result.metalSystemValueName = toSlice("position_in_patch");
+            result.requiredType = builder.getVectorType(builder.getBasicType(BaseType::Float), builder.getIntValue(builder.getIntType(), 3));
+            result.altRequiredType = builder.getVectorType(builder.getBasicType(BaseType::Float), builder.getIntValue(builder.getIntType(), 2));
+        }
+        else if (semanticName == "sv_groupid")
+        {
+            result.isSpecial = true;
+        }
+        else if (semanticName == "sv_groupindex")
+        {
+            result.isSpecial = true;
+        }
+        else if (semanticName == "sv_groupthreadid")
+        {
+            result.metalSystemValueName = toSlice("thread_position_in_threadgroup");
+            result.requiredType = builder.getVectorType(builder.getBasicType(BaseType::UInt), builder.getIntValue(builder.getIntType(), 3));
+        }
+        else if (semanticName == "sv_gsinstanceid")
+        {
+            // Metal does not have geometry shader, so this is invalid.
+            result.isUnsupported = true;
+        }
+        else if (semanticName == "sv_instanceid")
+        {
+            result.metalSystemValueName = toSlice("instance_id");
+            result.requiredType = builder.getBasicType(BaseType::UInt);
+        }
+        else if (semanticName == "sv_isfrontface")
+        {
+            result.metalSystemValueName = toSlice("front_facing");
+            result.requiredType = builder.getBasicType(BaseType::Bool);
+        }
+        else if (semanticName == "sv_outputcontrolpointid")
+        {
+            // In metal, a hull shader is just a compute shader.
+            // This needs to be handled separately, by lowering into an ordinary buffer.
+        }
+        else if (semanticName == "sv_pointsize")
+        {
+            result.metalSystemValueName = toSlice("point_size");
+            result.requiredType = builder.getBasicType(BaseType::Float);
+        }
+        else if (semanticName == "sv_primitiveid")
+        {
+            result.metalSystemValueName = toSlice("patch_id");
+            result.requiredType = builder.getBasicType(BaseType::UInt);
+            result.altRequiredType = builder.getBasicType(BaseType::UInt16);
+        }
+        else if (semanticName == "sv_rendertargetarrayindex")
+        {
+            result.metalSystemValueName = toSlice("render_target_array_index");
+            result.requiredType = builder.getBasicType(BaseType::UInt);
+            result.altRequiredType = builder.getBasicType(BaseType::UInt16);
+        }
+        else if (semanticName == "sv_sampleindex")
+        {
+            result.metalSystemValueName = toSlice("sample_id");
+            result.requiredType = builder.getBasicType(BaseType::UInt);
+        }
+        else if (semanticName == "sv_stencilref")
+        {
+            result.metalSystemValueName = toSlice("stencil");
+            result.requiredType = builder.getBasicType(BaseType::UInt);
+        }
+        else if (semanticName == "sv_tessfactor")
+        {
+            // Tessellation factor outputs should be lowered into a write into a normal buffer.
+        }
+        else if (semanticName == "sv_vertexid")
+        {
+            result.metalSystemValueName = toSlice("vertex_id");
+            result.requiredType = builder.getBasicType(BaseType::UInt);
+        }
+        else if (semanticName == "sv_viewid")
+        {
+            result.isUnsupported = true;
+        }
+        else if (semanticName == "sv_viewportarrayindex")
+        {
+            result.metalSystemValueName = toSlice("viewport_array_index");
+            result.requiredType = builder.getBasicType(BaseType::UInt);
+            result.altRequiredType = builder.getBasicType(BaseType::UInt16);
+        }
+        else if (semanticName == "sv_target")
+        {
+            result.metalSystemValueName = (StringBuilder() << "color(" << String(attrIndex) << ")").produceString();
+        }
+        else
+        {
+            result.isUnsupported = true;
+        }
+        return result;
+    }
+
+    void reportUnsupportedSystemAttribute(DiagnosticSink* sink, IRInst* param, String semanticName)
+    {
+        sink->diagnose(param->sourceLoc, Diagnostics::systemValueAttributeNotSupported, semanticName);
+    }
+
+    void ensureResultStructHasUserSemantic(DiagnosticSink* sink, IRStructType* structType, IRVarLayout* varLayout)
     {
         // Ensure each field in an output struct type has either a system semantic or a user semantic,
         // so that signature matching can happen correctly.
@@ -225,8 +391,21 @@ namespace Slang
         for (auto field : structType->getFields())
         {
             auto key = field->getKey();
-            if (key->findDecoration<IRSemanticDecoration>())
+            if (auto semanticDecor = key->findDecoration<IRSemanticDecoration>())
             {
+                if (semanticDecor->getSemanticName().startsWithCaseInsensitive(toSlice("sv_")))
+                {
+                    auto sysValInfo = getSystemValueInfo(builder, semanticDecor->getSemanticName(), semanticDecor->getSemanticIndex());
+                    if (sysValInfo.isUnsupported || sysValInfo.isSpecial)
+                    {
+                        reportUnsupportedSystemAttribute(sink, field, semanticDecor->getSemanticName());
+                    }
+                    else
+                    {
+                        builder.addTargetSystemValueDecoration(key, sysValInfo.metalSystemValueName.getUnownedSlice());
+                        semanticDecor->removeAndDeallocate();
+                    }
+                }
                 index++;
                 continue;
             }
@@ -245,7 +424,7 @@ namespace Slang
     }
 
 
-    void wrapReturnValueInStruct(EntryPointInfo entryPoint)
+    void wrapReturnValueInStruct(DiagnosticSink* sink, EntryPointInfo entryPoint)
     {
         // Wrap return value into a struct if it is not already a struct.
         // For example, given this entry point:
@@ -275,7 +454,7 @@ namespace Slang
         // If return type is already a struct, just make sure every field has a semantic.
         if (auto returnStructType = as<IRStructType>(returnType))
         {
-            ensureResultStructHasUserSemantic(returnStructType, resultLayout);
+            ensureResultStructHasUserSemantic(sink, returnStructType, resultLayout);
             return;
         }
 
@@ -288,13 +467,14 @@ namespace Slang
         auto key = builder.createStructKey();
         builder.addNameHintDecoration(key, toSlice("output"));
         builder.addLayoutDecoration(key, resultLayout);
+        builder.addTargetSystemValueDecoration(key, toSlice("color(0)"));
         builder.createStructField(structType, key, returnType);
         IRStructTypeLayout::Builder structTypeLayoutBuilder(&builder);
         structTypeLayoutBuilder.addField(key, resultLayout);
         auto typeLayout = structTypeLayoutBuilder.build();
         IRVarLayout::Builder varLayoutBuilder(&builder, typeLayout);
         auto varLayout = varLayoutBuilder.build();
-        ensureResultStructHasUserSemantic(structType, varLayout);
+        ensureResultStructHasUserSemantic(sink, structType, varLayout);
 
         for (auto block : func->getBlocks())
         {
@@ -401,13 +581,146 @@ namespace Slang
             });
     }
 
-    void legalizeEntryPointForMetal(EntryPointInfo entryPoint, DiagnosticSink* sink)
+    IRInst* tryConvertValue(IRBuilder& builder, IRInst* val, IRType* toType)
+    {
+        auto fromType = val->getFullType();
+        if (auto fromVector = as<IRVectorType>(fromType))
+        {
+            if (auto toVector = as<IRVectorType>(toType))
+            {
+                if (fromVector->getElementCount() != toVector->getElementCount())
+                {
+                    fromType = builder.getVectorType(fromVector->getElementType(), toVector->getElementCount());
+                    val = builder.emitVectorReshape(fromType, val);
+                }
+            }
+            else if (as<IRBasicType>(toType))
+            {
+                UInt index = 0;
+                val = builder.emitSwizzle(fromVector->getElementType(), val, 1, &index);
+                if (toType->getOp() == kIROp_VoidType)
+                    return nullptr;
+            }
+        }
+        else if (auto fromBasicType = as<IRBasicType>(fromType))
+        {
+            if (fromBasicType->getOp() == kIROp_VoidType)
+                return nullptr;
+            if (!as<IRBasicType>(toType))
+                return nullptr;
+            if (toType->getOp() == kIROp_VoidType)
+                return nullptr;
+        }
+        else
+        {
+            return nullptr;
+        }
+        return builder.emitCast(toType, val);
+    }
+
+    void legalizeSystemValueParameters(EntryPointInfo entryPoint, DiagnosticSink* sink)
     {
         SLANG_UNUSED(sink);
 
+        struct SystemValLegalizationWorkItem
+        {
+            IRParam* param;
+            String attrName;
+            UInt attrIndex;
+        };
+        List<SystemValLegalizationWorkItem> systemValWorkItems;
+        List<SystemValLegalizationWorkItem> workList;
+
+        IRBuilder builder(entryPoint.entryPointFunc);
+        List<IRParam*> params;
+
+        for (auto param : entryPoint.entryPointFunc->getParams())
+        {
+            if (auto semanticDecoration = param->findDecoration<IRSemanticDecoration>())
+            {
+                if (semanticDecoration->getSemanticName().startsWithCaseInsensitive(toSlice("sv_")))
+                {
+                    systemValWorkItems.add({ param, String(semanticDecoration->getSemanticName()).toLower(), (UInt)semanticDecoration->getSemanticIndex() });
+                    continue;
+                }
+            }
+
+            auto layoutDecor = param->findDecoration<IRLayoutDecoration>();
+            if (!layoutDecor)
+                continue;
+            auto sysValAttr = layoutDecor->findAttr<IRSystemValueSemanticAttr>();
+            if (!sysValAttr)
+                continue;
+            auto semanticName = String(sysValAttr->getName());
+            auto sysAttrIndex = sysValAttr->getIndex();
+            systemValWorkItems.add({ param, semanticName, sysAttrIndex });
+        }
+        for (auto workItem : systemValWorkItems)
+        {
+            auto param = workItem.param;
+            auto semanticName = workItem.attrName;
+            auto sysAttrIndex = workItem.attrIndex;
+
+            auto info = getSystemValueInfo(builder, semanticName, sysAttrIndex);
+            if (info.isSpecial)
+            {
+                if (semanticName == "sv_innercoverage")
+                {
+                    // Metal does not support conservative rasterization, so this is always false.
+                    auto val = builder.getBoolValue(false);
+                    param->replaceUsesWith(val);
+                    param->removeAndDeallocate();
+                }
+                else
+                {
+                    // Process special cases after trivial cases.
+                    workList.add(workItem);
+                }
+            }
+            if (info.isUnsupported)
+            {
+                reportUnsupportedSystemAttribute(sink, param, semanticName);
+                continue;
+            }
+            if (!info.requiredType)
+                continue;
+
+            builder.addTargetSystemValueDecoration(param, info.metalSystemValueName.getUnownedSlice());
+
+            // If the required type is different from the actual type, we need to insert a conversion.
+            if (info.requiredType != param->getFullType() && info.altRequiredType != param->getFullType())
+            {
+                auto targetType = param->getFullType();
+                builder.setInsertBefore(entryPoint.entryPointFunc->getFirstBlock()->getFirstOrdinaryInst());
+                param->setFullType(info.requiredType);
+                List<IRUse*> uses;
+                for (auto use = param->firstUse; use; use = use->nextUse)
+                    uses.add(use);
+                auto convertedValue = tryConvertValue(builder, param, targetType);
+                copyNameHintAndDebugDecorations(convertedValue, param);
+                if (!convertedValue)
+                {
+                    // If we can't convert the value, report an error.
+                    StringBuilder typeNameSB;
+                    getTypeNameHint(typeNameSB, info.requiredType);
+                    sink->diagnose(param->sourceLoc, Diagnostics::systemValueTypeIncompatible, semanticName, typeNameSB.produceString());
+                }
+                else
+                {
+                    for (auto use : uses)
+                        builder.replaceOperand(use, convertedValue);
+                }
+            }
+        }
+        fixUpFuncType(entryPoint.entryPointFunc);
+    }
+
+    void legalizeEntryPointForMetal(EntryPointInfo entryPoint, DiagnosticSink* sink)
+    {
         hoistEntryPointParameterFromStruct(entryPoint);
         packStageInParameters(entryPoint);
-        wrapReturnValueInStruct(entryPoint);
+        legalizeSystemValueParameters(entryPoint, sink);
+        wrapReturnValueInStruct(sink, entryPoint);
         legalizeMeshEntryPoint(entryPoint);
         legalizeDispatchMeshPayloadForMetal(entryPoint);
     }