Attachment 360258 Details for Bug 193877 – WIP

[patch] WIP

bug-193877-20190126172936.patch (text/plain), 64.23 KB, created by Myles C. Maxfield on 2019-01-26 17:29:37 PST

(hide)

Description:

Filename:

MIME Type:

Creator: Myles C. Maxfield

Created: 2019-01-26 17:29:37 PST

Size: 64.23 KB

patch

obsolete

>Subversion Revision: 240456
>diff --git a/Source/WebCore/ChangeLog b/Source/WebCore/ChangeLog
>index a112c3fdb6b29278746a261630269c0e11d6170f..1041707b554e6f968769c30762123a405a7dad73 100644
>--- a/Source/WebCore/ChangeLog
>+++ b/Source/WebCore/ChangeLog
>@@ -1,3 +1,66 @@
>+2019-01-26  Myles C. Maxfield  <mmaxfield@apple.com>
>+
>+        [WHLSL] Pack and unpack data at entry points and exit points
>+        https://bugs.webkit.org/show_bug.cgi?id=193877
>+
>+        Reviewed by NOBODY (OOPS!).
>+
>+        No new tests (OOPS!).
>+
>+        * Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp:
>+        (WebCore::WHLSL::Metal::attributeForSemantic):
>+        (WebCore::WHLSL::Metal::EntryPointScaffolding::EntryPointScaffolding):
>+        (WebCore::WHLSL::Metal::EntryPointScaffolding::resourceHelperTypes):
>+        (WebCore::WHLSL::Metal::EntryPointScaffolding::resourceSignature):
>+        (WebCore::WHLSL::Metal::EntryPointScaffolding::builtInsSignature):
>+        (WebCore::WHLSL::Metal::EntryPointScaffolding::mangledInputPath):
>+        (WebCore::WHLSL::Metal::EntryPointScaffolding::mangledOutputPath):
>+        (WebCore::WHLSL::Metal::EntryPointScaffolding::unpackResourcesAndNamedBuiltIns):
>+        (WebCore::WHLSL::Metal::VertexEntryPointScaffolding::VertexEntryPointScaffolding):
>+        (WebCore::WHLSL::Metal::VertexEntryPointScaffolding::helperTypes):
>+        (WebCore::WHLSL::Metal::VertexEntryPointScaffolding::signature):
>+        (WebCore::WHLSL::Metal::VertexEntryPointScaffolding::unpack):
>+        (WebCore::WHLSL::Metal::VertexEntryPointScaffolding::pack):
>+        (WebCore::WHLSL::Metal::FragmentEntryPointScaffolding::FragmentEntryPointScaffolding):
>+        (WebCore::WHLSL::Metal::FragmentEntryPointScaffolding::helperTypes):
>+        (WebCore::WHLSL::Metal::FragmentEntryPointScaffolding::signature):
>+        (WebCore::WHLSL::Metal::FragmentEntryPointScaffolding::unpack):
>+        (WebCore::WHLSL::Metal::FragmentEntryPointScaffolding::pack):
>+        (WebCore::WHLSL::Metal::ComputeEntryPointScaffolding::ComputeEntryPointScaffolding):
>+        (WebCore::WHLSL::Metal::ComputeEntryPointScaffolding::helperTypes):
>+        (WebCore::WHLSL::Metal::ComputeEntryPointScaffolding::signature):
>+        (WebCore::WHLSL::Metal::ComputeEntryPointScaffolding::unpack):
>+        (WebCore::WHLSL::Metal::ComputeEntryPointScaffolding::pack):
>+        (WebCore::WHLSL::Metal::EntryPointScaffolding::helperTypes): Deleted.
>+        (WebCore::WHLSL::Metal::EntryPointScaffolding::signature): Deleted.
>+        (WebCore::WHLSL::Metal::EntryPointScaffolding::unpack): Deleted.
>+        (WebCore::WHLSL::Metal::EntryPointScaffolding::pack): Deleted.
>+        * Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h:
>+        (WebCore::WHLSL::Metal::EntryPointScaffolding::parameterVariables):
>+        * Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp:
>+        (WebCore::WHLSL::Metal::FunctionDefinitionWriter::FunctionDefinitionWriter):
>+        (WebCore::WHLSL::Metal::metalFunctions):
>+        * Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.h:
>+        * Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.cpp:
>+        (WebCore::WHLSL::Metal::generateMetalCode):
>+        * Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.h:
>+        * Modules/webgpu/WHLSL/WHLSLIntrinsics.h:
>+        (WebCore::WHLSL::Intrinsics::float2Type const):
>+        * Modules/webgpu/WHLSL/WHLSLPipelineDescriptor.h: Added.
>+        * Modules/webgpu/WHLSL/WHLSLSemanticMatcher.cpp: Added.
>+        (WebCore::WHLSL::findEntryPoint):
>+        (WebCore::WHLSL::matchMode):
>+        (WebCore::WHLSL::matchResources):
>+        (WebCore::WHLSL::matchInputsOutputs):
>+        (WebCore::WHLSL::isAcceptableFormat):
>+        (WebCore::WHLSL::matchVertexAttributes):
>+        (WebCore::WHLSL::matchColorAttachments):
>+        (WebCore::WHLSL::matchDepthAttachment):
>+        (WebCore::WHLSL::matchSemantics):
>+        * Modules/webgpu/WHLSL/WHLSLSemanticMatcher.h: Copied from Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp.
>+        * Sources.txt:
>+        * WebCore.xcodeproj/project.pbxproj:
>+
> 2019-01-24  Charles Vazac  <cvazac@akamai.com>
> 
>         Implement PerformanceObserver.supportedEntryTypes
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp
>index 48fcf235f251640c5a782210e9d2684f4ee3c07d..3979525a16fc17f1594146bd7893379928eecf56 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp
>@@ -28,7 +28,17 @@
> 
> #if ENABLE(WEBGPU)
> 
>+#include "WHLSLBuiltInSemantic.h"
>+#include "WHLSLFunctionDefinition.h"
> #include "WHLSLGatherEntryPointItems.h"
>+#include "WHLSLPipelineDescriptor.h"
>+#include "WHLSLResourceSemantic.h"
>+#include "WHLSLStageInOutSemantic.h"
>+#include "WHLSLStructureDefinition.h"
>+#include "WHLSLTypeNamer.h"
>+#include <wtf/text/StringBuilder.h>
>+#include <wtf/text/StringConcatenate.h>
>+#include <wtf/text/StringConcatenateNumbers.h>
> 
> namespace WebCore {
> 
>@@ -36,35 +46,459 @@ namespace WHLSL {
> 
> namespace Metal {
> 
>-EntryPointScaffolding::EntryPointScaffolding(AST::FunctionDefinition& functionDefinition, Intrinsics& intrinsics)
>-    : m_functionDefinition(&functionDefinition)
>-    , m_intrinsics(&intrinsics)
>+static String attributeForSemantic(AST::BuiltInSemantic& builtInSemantic)
> {
>-    // FIXME: Implement this.
>-    gatherEntryPointItems(*m_intrinsics, *m_functionDefinition);
>+    switch (builtInSemantic.variable()) {
>+    case AST::BuiltInSemantic::Variable::SVInstanceID:
>+        return "[[instance_id]]"_str;
>+    case AST::BuiltInSemantic::Variable::SVVertexID:
>+        return "[[vertex_id]]"_str;
>+    case AST::BuiltInSemantic::Variable::PSize:
>+        return "[[point_size]]"_str;
>+    case AST::BuiltInSemantic::Variable::SVPosition:
>+        return "[[position]]"_str;
>+    case AST::BuiltInSemantic::Variable::SVIsFrontFace:
>+        return "[[front_facing]]"_str;
>+    case AST::BuiltInSemantic::Variable::SVSampleIndex:
>+        return "[[sample_id]]"_str;
>+    case AST::BuiltInSemantic::Variable::SVInnerCoverage:
>+        return "[[sample_mask]]"_str;
>+    case AST::BuiltInSemantic::Variable::SVTarget:
>+        return makeString("[[color(", *builtInSemantic.targetIndex(), ")]]");
>+    case AST::BuiltInSemantic::Variable::SVDepth:
>+        return "[[depth(any)]]"_str;
>+    case AST::BuiltInSemantic::Variable::SVCoverage:
>+        return "[[sample_mask]]"_str;
>+    case AST::BuiltInSemantic::Variable::SVDispatchThreadID:
>+        return "[[thread_position_in_grid]]"_str;
>+    case AST::BuiltInSemantic::Variable::SVGroupID:
>+        return "[[threadgroup_position_in_grid]]"_str;
>+    case AST::BuiltInSemantic::Variable::SVGroupIndex:
>+        return "[[thread_index_in_threadgroup]]"_str;
>+    default:
>+        ASSERT(builtInSemantic.variable() == AST::BuiltInSemantic::Variable::SVGroupThreadID);
>+        return "[[thread_position_in_threadgroup]]"_str;
>+    }
> }
> 
>-String EntryPointScaffolding::helperTypes()
>+static String attributeForSemantic(AST::Semantic& semantic)
> {
>-    // FIXME: Implement this.
>-    return String();
>+    if (WTF::holds_alternative<AST::BuiltInSemantic>(semantic))
>+        return attributeForSemantic(WTF::get<AST::BuiltInSemantic>(semantic));
>+    auto& stageInOutSemantic = WTF::get<AST::StageInOutSemantic>(semantic);
>+    return makeString("[[user(", stageInOutSemantic.index(), ")]]");
> }
> 
>-String EntryPointScaffolding::signature()
>+EntryPointScaffolding::EntryPointScaffolding(AST::FunctionDefinition& functionDefinition, Intrinsics& intrinsics, TypeNamer& typeNamer, EntryPointItems& entryPointItems, HashMap<Binding*, size_t>& resourceMap, Layout& layout, std::function<String()>&& generateNextVariableName)
>+    : m_functionDefinition(functionDefinition)
>+    , m_intrinsics(intrinsics)
>+    , m_typeNamer(typeNamer)
>+    , m_entryPointItems(entryPointItems)
>+    , m_resourceMap(resourceMap)
>+    , m_layout(layout)
>+    , m_generateNextVariableName(generateNextVariableName)
> {
>-    // FIXME: Implement this.
>-    return String();
>+    unsigned argumentBufferIndex = 0;
>+    m_namedBindGroups.reserveInitialCapacity(m_layout.size());
>+    for (size_t i = 0; i < m_layout.size(); ++i) {
>+        NamedBindGroup namedBindGroup;
>+        namedBindGroup.structName = m_typeNamer.generateNextTypeName();
>+        namedBindGroup.variableName = m_generateNextVariableName();
>+        namedBindGroup.argumentBufferIndex = argumentBufferIndex++;
>+        namedBindGroup.namedBindings.reserveInitialCapacity(m_layout[i].bindings.size());
>+        unsigned index = 0;
>+        for (size_t j = 0; j < m_layout[i].bindings.size(); ++j) {
>+            NamedBinding namedBinding;
>+            namedBinding.elementName = m_typeNamer.generateNextStructureElementName();
>+            namedBinding.index = index++;
>+            namedBindGroup.namedBindings.uncheckedAppend(WTFMove(namedBinding));
>+        }
>+        m_namedBindGroups.uncheckedAppend(WTFMove(namedBindGroup));
>+    }
>+
>+    for (size_t i = 0; i < m_entryPointItems.inputs.size(); ++i) {
>+        if (!WTF::holds_alternative<AST::BuiltInSemantic>(*m_entryPointItems.inputs[i].semantic))
>+            continue;
>+        NamedBuiltIn namedBuiltIn;
>+        namedBuiltIn.indexInEntryPointItems = i;
>+        namedBuiltIn.variableName = m_generateNextVariableName();
>+        m_namedBuiltIns.append(WTFMove(namedBuiltIn));
>+    }
>+
>+    m_parameterVariables.reserveInitialCapacity(m_functionDefinition.parameters().size());
>+    for (size_t i = 0; i < m_functionDefinition.parameters().size(); ++i)
>+        m_parameterVariables.uncheckedAppend(m_generateNextVariableName());
> }
> 
>-String EntryPointScaffolding::unpack()
>+String EntryPointScaffolding::resourceHelperTypes()
> {
>-    // FIXME: Implement this.
>-    return String();
>+    StringBuilder stringBuilder;
>+    for (size_t i = 0; i < m_layout.size(); ++i) {
>+        stringBuilder.append(makeString("struct ", m_namedBindGroups[i].structName, " {\n"));
>+        for (size_t j = 0; j < m_layout[i].bindings.size(); ++j) {
>+            auto iterator = m_resourceMap.find(&m_layout[i].bindings[j]);
>+            if (iterator == m_resourceMap.end())
>+                continue;
>+            auto mangledTypeName = m_typeNamer.mangledNameForType(*m_entryPointItems.inputs[iterator->value].unnamedType);
>+            auto elementName = m_namedBindGroups[i].namedBindings[j].elementName;
>+            auto index = m_namedBindGroups[i].namedBindings[j].index;
>+            stringBuilder.append(makeString("    ", mangledTypeName, ' ', elementName, " [[id(", index, ")]];\n"));
>+        }
>+        stringBuilder.append("}\n\n");
>+    }
>+    return stringBuilder.toString();
>+}
>+
>+Optional<String> EntryPointScaffolding::resourceSignature()
>+{
>+    if (!m_layout.size())
>+        return WTF::nullopt;
>+
>+    StringBuilder stringBuilder;
>+    for (size_t i = 0; i < m_layout.size(); ++i) {
>+        if (i)
>+            stringBuilder.append(", ");
>+        auto& namedBindGroup = m_namedBindGroups[i];
>+        stringBuilder.append(makeString(namedBindGroup.structName, "& ", namedBindGroup.variableName, " [[buffer(", namedBindGroup.argumentBufferIndex, ")]]"));
>+    }
>+    return stringBuilder.toString();
>+}
>+
>+Optional<String> EntryPointScaffolding::builtInsSignature()
>+{
>+    if (!m_namedBuiltIns.size())
>+        return WTF::nullopt;
>+
>+    StringBuilder stringBuilder;
>+    for (size_t i = 0; i < m_namedBuiltIns.size(); ++i) {
>+        if (i)
>+            stringBuilder.append(", ");
>+        auto& namedBuiltIn = m_namedBuiltIns[i];
>+        auto& item = m_entryPointItems.inputs[namedBuiltIn.indexInEntryPointItems];
>+        auto& builtInSemantic = WTF::get<AST::BuiltInSemantic>(*item.semantic);
>+        auto mangledTypeName = m_typeNamer.mangledNameForType(*item.unnamedType);
>+        auto variableName = namedBuiltIn.variableName;
>+        stringBuilder.append(makeString(mangledTypeName, ' ', variableName, ' ', attributeForSemantic(builtInSemantic)));
>+    }
>+    return stringBuilder.toString();
>+}
>+
>+String EntryPointScaffolding::mangledInputPath(Vector<String>& path)
>+{
>+    ASSERT(!path.isEmpty());
>+    StringBuilder stringBuilder;
>+    bool found = false;
>+    AST::StructureDefinition* structureDefinition = nullptr;
>+    for (size_t i = 0; i < m_functionDefinition.parameters().size(); ++i) {
>+        if (m_functionDefinition.parameters()[i].name() == path[0]) {
>+            stringBuilder.append(m_parameterVariables[i]);
>+            auto& unifyNode = m_functionDefinition.parameters()[i].type()->unifyNode();
>+            if (is<AST::NamedType>(unifyNode)) {
>+                auto& namedType = downcast<AST::NamedType>(unifyNode);
>+                if (is<AST::StructureDefinition>(namedType))
>+                    structureDefinition = &downcast<AST::StructureDefinition>(namedType);
>+            }
>+            found = true;
>+            break;
>+        }
>+    }
>+    ASSERT(found);
>+    for (size_t i = 1; i < path.size(); ++i) {
>+        ASSERT(structureDefinition);
>+        auto* next = structureDefinition->find(path[i]);
>+        ASSERT(next);
>+        stringBuilder.append(makeString('.', m_typeNamer.mangledNameForStructureElement(*next)));
>+        structureDefinition = nullptr;
>+        auto& unifyNode = next->type().unifyNode();
>+        if (is<AST::NamedType>(unifyNode)) {
>+            auto& namedType = downcast<AST::NamedType>(unifyNode);
>+            if (is<AST::StructureDefinition>(namedType))
>+                structureDefinition = &downcast<AST::StructureDefinition>(namedType);
>+        }
>+    }
>+
>+    return stringBuilder.toString();
>+}
>+
>+String EntryPointScaffolding::mangledOutputPath(Vector<String>& path)
>+{
>+    StringBuilder stringBuilder;
>+
>+    AST::StructureDefinition* structureDefinition = nullptr;
>+    auto& unifyNode = m_functionDefinition.type().unifyNode();
>+    ASSERT(is<AST::NamedType>(unifyNode));
>+    auto& namedType = downcast<AST::NamedType>(unifyNode);
>+    ASSERT(is<AST::StructureDefinition>(namedType));
>+    structureDefinition = &downcast<AST::StructureDefinition>(namedType);
>+    for (auto& component : path) {
>+        ASSERT(structureDefinition);
>+        auto* next = structureDefinition->find(component);
>+        ASSERT(next);
>+        stringBuilder.append(makeString('.', m_typeNamer.mangledNameForStructureElement(*next)));
>+        structureDefinition = nullptr;
>+        auto& unifyNode = next->type().unifyNode();
>+        if (is<AST::NamedType>(unifyNode)) {
>+            auto& namedType = downcast<AST::NamedType>(unifyNode);
>+            if (is<AST::StructureDefinition>(namedType))
>+                structureDefinition = &downcast<AST::StructureDefinition>(namedType);
>+        }
>+    }
>+
>+    return stringBuilder.toString();
>+}
>+
>+String EntryPointScaffolding::unpackResourcesAndNamedBuiltIns()
>+{
>+    StringBuilder stringBuilder;
>+    for (size_t i = 0; i < m_functionDefinition.parameters().size(); ++i)
>+        stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*m_functionDefinition.parameters()[i].type()), ' ', m_parameterVariables[i], ";\n"));
>+
>+    for (size_t i = 0; i < m_layout.size(); ++i) {
>+        auto variableName = m_namedBindGroups[i].variableName;
>+        for (size_t j = 0; j < m_layout[i].bindings.size(); ++j) {
>+            auto iterator = m_resourceMap.find(&m_layout[i].bindings[j]);
>+            if (iterator == m_resourceMap.end())
>+                continue;
>+            auto& path = m_entryPointItems.inputs[iterator->value].path;
>+            auto elementName = m_namedBindGroups[i].namedBindings[j].elementName;
>+            stringBuilder.append(makeString(mangledInputPath(path), " = ", variableName, '.', elementName, ";\n"));
>+        }
>+    }
>+
>+    for (auto& namedBuiltIn : m_namedBuiltIns) {
>+        auto& path = m_entryPointItems.inputs[namedBuiltIn.indexInEntryPointItems].path;
>+        auto& variableName = namedBuiltIn.variableName;
>+        stringBuilder.append(makeString(mangledInputPath(path), " = ", variableName, ";\n"));
>+    }
>+    return stringBuilder.toString();
>+}
>+
>+VertexEntryPointScaffolding::VertexEntryPointScaffolding(AST::FunctionDefinition& functionDefinition, Intrinsics& intrinsics, TypeNamer& typeNamer, EntryPointItems& entryPointItems, HashMap<Binding*, size_t>& resourceMap, Layout& layout, std::function<String()>&& generateNextVariableName, HashMap<VertexAttribute*, size_t>& matchedVertexAttributes)
>+    : EntryPointScaffolding(functionDefinition, intrinsics, typeNamer, entryPointItems, resourceMap, layout, WTFMove(generateNextVariableName))
>+    , m_matchedVertexAttributes(matchedVertexAttributes)
>+    , m_stageInStructName(typeNamer.generateNextTypeName())
>+    , m_returnStructName(typeNamer.generateNextTypeName())
>+    , m_stageInParameterName(m_generateNextVariableName())
>+{
>+    m_namedStageIns.reserveInitialCapacity(m_matchedVertexAttributes.size());
>+    for (auto& keyValuePair : m_matchedVertexAttributes) {
>+        NamedStageIn namedStageIn;
>+        namedStageIn.indexInEntryPointItems = keyValuePair.value;
>+        namedStageIn.elementName = m_typeNamer.generateNextStructureElementName();
>+        namedStageIn.attributeIndex = keyValuePair.key->name;
>+        m_namedStageIns.uncheckedAppend(WTFMove(namedStageIn));
>+    }
>+
>+    m_namedOutputs.reserveInitialCapacity(m_entryPointItems.outputs.size());
>+    for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
>+        NamedOutput namedOutput;
>+        namedOutput.elementName = m_typeNamer.generateNextStructureElementName();
>+        m_namedOutputs.uncheckedAppend(WTFMove(namedOutput));
>+    }
>+}
>+
>+String VertexEntryPointScaffolding::helperTypes()
>+{
>+    StringBuilder stringBuilder;
>+
>+    stringBuilder.append(makeString("struct ", m_stageInStructName, " {\n"));
>+    for (auto& namedStageIn : m_namedStageIns) {
>+        auto mangledTypeName = m_typeNamer.mangledNameForType(*m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].unnamedType);
>+        auto elementName = namedStageIn.elementName;
>+        auto attributeIndex = namedStageIn.elementName;
>+        stringBuilder.append(makeString("    ", mangledTypeName, ' ', elementName, " [[attribute(", attributeIndex, ")]];\n"));
>+    }
>+    stringBuilder.append("}\n\n");
>+
>+    stringBuilder.append(makeString("struct ", m_returnStructName, " {\n"));
>+    for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
>+        auto& outputItem = m_entryPointItems.outputs[i];
>+        auto mangledTypeName = m_typeNamer.mangledNameForType(*outputItem.unnamedType);
>+        auto elementName = m_namedOutputs[i].elementName;
>+        auto attribute = attributeForSemantic(*outputItem.semantic);
>+        stringBuilder.append(makeString("    ", mangledTypeName, ' ', elementName, ' ', attribute, ";\n"));
>+    }
>+    stringBuilder.append("}\n\n");
>+
>+    stringBuilder.append(resourceHelperTypes());
>+
>+    return stringBuilder.toString();
>+}
>+
>+String VertexEntryPointScaffolding::signature(String& functionName)
>+{
>+    StringBuilder stringBuilder;
>+
>+    stringBuilder.append(makeString("vertex ", m_returnStructName, ' ', functionName, '(', m_stageInStructName, ' ', m_stageInParameterName, " [[stage_in]]"));
>+    if (auto resourceSignature = this->resourceSignature())
>+        stringBuilder.append(makeString(", ", *resourceSignature));
>+    if (auto builtInsSignature = this->builtInsSignature())
>+        stringBuilder.append(makeString(", ", *builtInsSignature));
>+    stringBuilder.append(")");
>+
>+    return stringBuilder.toString();
>+}
>+
>+String VertexEntryPointScaffolding::unpack()
>+{
>+    StringBuilder stringBuilder;
>+
>+    stringBuilder.append(unpackResourcesAndNamedBuiltIns());
>+
>+    for (auto& namedStageIn : m_namedStageIns) {
>+        auto& path = m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].path;
>+        auto& elementName = namedStageIn.elementName;
>+        stringBuilder.append(makeString(mangledInputPath(path), " = ", m_stageInStructName, '.', elementName, ";\n"));
>+    }
>+
>+    return stringBuilder.toString();
>+}
>+
>+String VertexEntryPointScaffolding::pack(const String& inputVariableName, const String& outputVariableName)
>+{
>+    StringBuilder stringBuilder;
>+    stringBuilder.append(makeString(m_returnStructName, ' ', outputVariableName));
>+    for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
>+        auto& elementName = m_namedOutputs[i].elementName;
>+        auto& path = m_entryPointItems.outputs[i].path;
>+        stringBuilder.append(makeString(outputVariableName, '.', elementName, " = ", inputVariableName, mangledOutputPath(path), ";\n"));
>+    }
>+    return stringBuilder.toString();
>+}
>+
>+FragmentEntryPointScaffolding::FragmentEntryPointScaffolding(AST::FunctionDefinition& functionDefinition, Intrinsics& intrinsics, TypeNamer& typeNamer, EntryPointItems& entryPointItems, HashMap<Binding*, size_t>& resourceMap, Layout& layout, std::function<String()>&& generateNextVariableName, HashMap<AttachmentDescriptor*, size_t>&)
>+    : EntryPointScaffolding(functionDefinition, intrinsics, typeNamer, entryPointItems, resourceMap, layout, WTFMove(generateNextVariableName))
>+    , m_stageInStructName(typeNamer.generateNextTypeName())
>+    , m_returnStructName(typeNamer.generateNextTypeName())
>+    , m_stageInParameterName(m_generateNextVariableName())
>+{
>+    for (size_t i = 0; i < m_entryPointItems.inputs.size(); ++i) {
>+        auto& inputItem = m_entryPointItems.inputs[i];
>+        if (!WTF::holds_alternative<AST::StageInOutSemantic>(*inputItem.semantic))
>+            continue;
>+        auto& stageInOutSemantic = WTF::get<AST::StageInOutSemantic>(*inputItem.semantic);
>+        NamedStageIn namedStageIn;
>+        namedStageIn.indexInEntryPointItems = i;
>+        namedStageIn.elementName = m_typeNamer.generateNextStructureElementName();
>+        namedStageIn.attributeIndex = stageInOutSemantic.index();
>+        m_namedStageIns.append(WTFMove(namedStageIn));
>+    }
>+
>+    m_namedOutputs.reserveInitialCapacity(m_entryPointItems.outputs.size());
>+    for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
>+        NamedOutput namedOutput;
>+        namedOutput.elementName = m_typeNamer.generateNextStructureElementName();
>+        m_namedOutputs.uncheckedAppend(WTFMove(namedOutput));
>+    }
>+}
>+
>+String FragmentEntryPointScaffolding::helperTypes()
>+{
>+    StringBuilder stringBuilder;
>+
>+    stringBuilder.append(makeString("struct ", m_stageInStructName, " {\n"));
>+    for (auto& namedStageIn : m_namedStageIns) {
>+        auto mangledTypeName = m_typeNamer.mangledNameForType(*m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].unnamedType);
>+        auto elementName = namedStageIn.elementName;
>+        auto attributeIndex = namedStageIn.elementName;
>+        stringBuilder.append(makeString("    ", mangledTypeName, ' ', elementName, " [[user(", attributeIndex, ")]];\n"));
>+    }
>+    stringBuilder.append("}\n\n");
>+
>+    stringBuilder.append(makeString("struct ", m_returnStructName, " {\n"));
>+    for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
>+        auto& outputItem = m_entryPointItems.outputs[i];
>+        auto mangledTypeName = m_typeNamer.mangledNameForType(*outputItem.unnamedType);
>+        auto elementName = m_namedOutputs[i].elementName;
>+        auto attribute = attributeForSemantic(*outputItem.semantic);
>+        stringBuilder.append(makeString("    ", mangledTypeName, ' ', elementName, ' ', attribute, ";\n"));
>+    }
>+    stringBuilder.append("}\n\n");
>+
>+    stringBuilder.append(resourceHelperTypes());
>+
>+    return stringBuilder.toString();
>+}
>+
>+String FragmentEntryPointScaffolding::signature(String& functionName)
>+{
>+    StringBuilder stringBuilder;
>+
>+    stringBuilder.append(makeString("fragment ", m_returnStructName, ' ', functionName, '(', m_stageInStructName, ' ', m_stageInParameterName, " [[stage_in]]"));
>+    if (auto resourceSignature = this->resourceSignature())
>+        stringBuilder.append(makeString(", ", *resourceSignature));
>+    if (auto builtInsSignature = this->builtInsSignature())
>+        stringBuilder.append(makeString(", ", *builtInsSignature));
>+    stringBuilder.append(")");
>+
>+    return stringBuilder.toString();
>+}
>+
>+String FragmentEntryPointScaffolding::unpack()
>+{
>+    StringBuilder stringBuilder;
>+
>+    stringBuilder.append(unpackResourcesAndNamedBuiltIns());
>+
>+    for (auto& namedStageIn : m_namedStageIns) {
>+        auto& path = m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].path;
>+        auto& elementName = namedStageIn.elementName;
>+        stringBuilder.append(makeString(mangledInputPath(path), " = ", m_stageInStructName, '.', elementName, ";\n"));
>+    }
>+
>+    return stringBuilder.toString();
>+}
>+
>+String FragmentEntryPointScaffolding::pack(const String& inputVariableName, const String& outputVariableName)
>+{
>+    StringBuilder stringBuilder;
>+    stringBuilder.append(makeString(m_returnStructName, ' ', outputVariableName));
>+    for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
>+        auto& elementName = m_namedOutputs[i].elementName;
>+        auto& path = m_entryPointItems.outputs[i].path;
>+        stringBuilder.append(makeString(outputVariableName, '.', elementName, " = ", inputVariableName, mangledOutputPath(path), ";\n"));
>+    }
>+    return stringBuilder.toString();
>+}
>+
>+ComputeEntryPointScaffolding::ComputeEntryPointScaffolding(AST::FunctionDefinition& functionDefinition, Intrinsics& intrinsics, TypeNamer& typeNamer, EntryPointItems& entryPointItems, HashMap<Binding*, size_t>& resourceMap, Layout& layout, std::function<String()>&& generateNextVariableName)
>+    : EntryPointScaffolding(functionDefinition, intrinsics, typeNamer, entryPointItems, resourceMap, layout, WTFMove(generateNextVariableName))
>+{
>+}
>+
>+String ComputeEntryPointScaffolding::helperTypes()
>+{
>+    return resourceHelperTypes();
>+}
>+
>+String ComputeEntryPointScaffolding::signature(String& functionName)
>+{
>+    StringBuilder stringBuilder;
>+
>+    stringBuilder.append(makeString("compute void ", functionName, '('));
>+    bool empty = true;
>+    if (auto resourceSignature = this->resourceSignature()) {
>+        empty = false;
>+        stringBuilder.append(makeString(*resourceSignature));
>+    }
>+    if (auto builtInsSignature = this->builtInsSignature()) {
>+        if (!empty)
>+            stringBuilder.append(", ");
>+        stringBuilder.append(*builtInsSignature);
>+    }
>+    stringBuilder.append(")");
>+
>+    return stringBuilder.toString();
>+}
>+
>+String ComputeEntryPointScaffolding::unpack()
>+{
>+    return unpackResourcesAndNamedBuiltIns();
> }
> 
>-String EntryPointScaffolding::pack(const String&, const String&)
>+String ComputeEntryPointScaffolding::pack(const String&, const String&)
> {
>-    // FIXME: Implement this.
>+    ASSERT_NOT_REACHED();
>     return String();
> }
> 
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h
>index 8d65fad207df539e9ea7404505ea6e87214c94e0..7f358a00afc59d1c214d2d2cc01f6998db831700 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h
>@@ -27,6 +27,11 @@
> 
> #if ENABLE(WEBGPU)
> 
>+#include "WHLSLPipelineDescriptor.h"
>+#include <wtf/HashMap.h>
>+#include <wtf/Optional.h>
>+#include <wtf/text/WTFString.h>
>+
> namespace WebCore {
> 
> namespace WHLSL {
>@@ -37,23 +42,132 @@ class FunctionDefinition;
> 
> }
> 
>+struct EntryPointItems;
> class Intrinsics;
> 
> namespace Metal {
> 
>+class TypeNamer;
>+
> // FIXME: This needs to know about the pipeline state object to emit function prologues and epilogues.
> class EntryPointScaffolding {
> public:
>-    EntryPointScaffolding(AST::FunctionDefinition&, Intrinsics&);
>+    EntryPointScaffolding(AST::FunctionDefinition&, Intrinsics&, TypeNamer&, EntryPointItems&, HashMap<Binding*, size_t>& resourceMap, Layout&, std::function<String()>&& generateNextVariableName);
>+    virtual ~EntryPointScaffolding() = default;
>+
>+    virtual String helperTypes() = 0;
>+    virtual String signature(String& functionName) = 0;
>+    virtual String unpack() = 0;
>+    virtual String pack(const String& existingVariableName, const String& variableName) = 0;
>+
>+    Vector<String>& parameterVariables() { return m_parameterVariables; }
>+
>+protected:
>+    String resourceHelperTypes();
>+    Optional<String> resourceSignature();
>+    Optional<String> builtInsSignature();
>+
>+    String mangledInputPath(Vector<String>& path);
>+    String mangledOutputPath(Vector<String>& path);
>+    String unpackResourcesAndNamedBuiltIns();
>+
>+    AST::FunctionDefinition& m_functionDefinition;
>+    Intrinsics& m_intrinsics;
>+    TypeNamer& m_typeNamer;
>+    EntryPointItems& m_entryPointItems;
>+    HashMap<Binding*, size_t>& m_resourceMap;
>+    Layout& m_layout;
>+    std::function<String()> m_generateNextVariableName;
>+
>+    struct NamedBinding {
>+        String elementName;
>+        unsigned index;
>+    };
>+    struct NamedBindGroup {
>+        String structName;
>+        String variableName;
>+        Vector<NamedBinding> namedBindings;
>+        unsigned argumentBufferIndex;
>+    };
>+    Vector<NamedBindGroup> m_namedBindGroups; // Parallel to m_layout
>+
>+    struct NamedBuiltIn {
>+        size_t indexInEntryPointItems;
>+        String variableName;
>+    };
>+    Vector<NamedBuiltIn> m_namedBuiltIns;
>+
>+    Vector<String> m_parameterVariables;
>+};
>+
>+class VertexEntryPointScaffolding : public EntryPointScaffolding {
>+public:
>+    VertexEntryPointScaffolding(AST::FunctionDefinition&, Intrinsics&, TypeNamer&, EntryPointItems&, HashMap<Binding*, size_t>& resourceMap, Layout&, std::function<String()>&& generateNextVariableName, HashMap<VertexAttribute*, size_t>& matchedVertexAttributes);
>+    virtual ~VertexEntryPointScaffolding() = default;
>+
>+    String helperTypes() override;
>+    String signature(String& functionName) override;
>+    String unpack() override;
>+    String pack(const String& existingVariableName, const String& variableName) override;
>+
>+private:
>+    HashMap<VertexAttribute*, size_t>& m_matchedVertexAttributes;
>+    String m_stageInStructName;
>+    String m_returnStructName;
>+    String m_stageInParameterName;
>+
>+    struct NamedStageIn {
>+        size_t indexInEntryPointItems;
>+        String elementName;
>+        unsigned attributeIndex;
>+    };
>+    Vector<NamedStageIn> m_namedStageIns;
>+
>+    struct NamedOutput {
>+        String elementName;
>+    };
>+    Vector<NamedOutput> m_namedOutputs;
>+};
>+
>+class FragmentEntryPointScaffolding : public EntryPointScaffolding {
>+public:
>+    FragmentEntryPointScaffolding(AST::FunctionDefinition&, Intrinsics&, TypeNamer&, EntryPointItems&, HashMap<Binding*, size_t>& resourceMap, Layout&, std::function<String()>&& generateNextVariableName, HashMap<AttachmentDescriptor*, size_t>& matchedColorAttachments);
>+    virtual ~FragmentEntryPointScaffolding() = default;
>+
>+    String helperTypes() override;
>+    String signature(String& functionName) override;
>+    String unpack() override;
>+    String pack(const String& existingVariableName, const String& variableName) override;
>+
>+private:
>+    String m_stageInStructName;
>+    String m_returnStructName;
>+    String m_stageInParameterName;
>+
>+    struct NamedStageIn {
>+        size_t indexInEntryPointItems;
>+        String elementName;
>+        unsigned attributeIndex;
>+    };
>+    Vector<NamedStageIn> m_namedStageIns;
>+
>+    struct NamedOutput {
>+        String elementName;
>+    };
>+    Vector<NamedOutput> m_namedOutputs;
>+};
>+
>+class ComputeEntryPointScaffolding : public EntryPointScaffolding {
>+public:
>+    ComputeEntryPointScaffolding(AST::FunctionDefinition&, Intrinsics&, TypeNamer&, EntryPointItems&, HashMap<Binding*, size_t>& resourceMap, Layout&, std::function<String()>&& generateNextVariableName);
>+    virtual ~ComputeEntryPointScaffolding() = default;
> 
>-    String helperTypes();
>-    String signature();
>-    String unpack();
>-    String pack(const String& existingVariableName, const String& variableName);
>+    String helperTypes() override;
>+    String signature(String& functionName) override;
>+    String unpack() override;
>+    String pack(const String& existingVariableName, const String& variableName) override;
> 
> private:
>-    AST::FunctionDefinition* m_functionDefinition;
>-    Intrinsics* m_intrinsics;
> };
> 
> }
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp
>index 019a267bbc8be6fd27ea0e108f81b79428d50cf6..84981e393d9acf350a1c41f35fec6c4069a6fc09 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp
>@@ -110,10 +110,12 @@ private:
> 
> class FunctionDefinitionWriter : public Visitor {
> public:
>-    FunctionDefinitionWriter(Intrinsics& intrinsics, TypeNamer& typeNamer, HashMap<AST::FunctionDeclaration*, String>& functionMapping)
>+    FunctionDefinitionWriter(Intrinsics& intrinsics, TypeNamer& typeNamer, HashMap<AST::FunctionDeclaration*, String>& functionMapping, Variant<MatchedRenderSemantics, MatchedComputeSemantics>&& matchedSemantics, Layout& layout)
>         : m_intrinsics(intrinsics)
>         , m_typeNamer(typeNamer)
>         , m_functionMapping(functionMapping)
>+        , m_matchedSemantics(WTFMove(matchedSemantics))
>+        , m_layout(layout)
>     {
>     }
> 
>@@ -133,16 +135,34 @@ public:
>         auto iterator = m_functionMapping.find(&functionDefinition);
>         ASSERT(iterator != m_functionMapping.end());
>         if (functionDefinition.entryPointType()) {
>-            m_entryPointScaffolding = EntryPointScaffolding(functionDefinition, m_intrinsics);
>+            auto generateNextVariableName = [this]() -> String {
>+                return this->generateNextVariableName();
>+            };
>+            WTF::visit(WTF::makeVisitor([&](MatchedRenderSemantics& matchedRenderSemantics) {
>+                if (&functionDefinition == matchedRenderSemantics.vertexShader)
>+                    m_entryPointScaffolding = std::make_unique<VertexEntryPointScaffolding>(functionDefinition, m_intrinsics, m_typeNamer, matchedRenderSemantics.vertexShaderEntryPointItems, matchedRenderSemantics.vertexShaderResourceMap, m_layout, WTFMove(generateNextVariableName), matchedRenderSemantics.matchedVertexAttributes);
>+                else if (&functionDefinition == matchedRenderSemantics.fragmentShader)
>+                    m_entryPointScaffolding = std::make_unique<FragmentEntryPointScaffolding>(functionDefinition, m_intrinsics, m_typeNamer, matchedRenderSemantics.fragmentShaderEntryPointItems, matchedRenderSemantics.fragmentShaderResourceMap, m_layout, WTFMove(generateNextVariableName), matchedRenderSemantics.matchedColorAttachments);
>+            }, [&](MatchedComputeSemantics& matchedComputeSemantics) {
>+                if (&functionDefinition == matchedComputeSemantics.shader)
>+                    m_entryPointScaffolding = std::make_unique<ComputeEntryPointScaffolding>(functionDefinition, m_intrinsics, m_typeNamer, matchedComputeSemantics.entryPointItems, matchedComputeSemantics.resourceMap, m_layout, WTFMove(generateNextVariableName));
>+            }), m_matchedSemantics);
>+            if (!m_entryPointScaffolding)
>+                return;
>             m_stringBuilder.append(m_entryPointScaffolding->helperTypes());
>             m_stringBuilder.append('\n');
>-            m_stringBuilder.append(makeString(m_entryPointScaffolding->signature(), " {"));
>+            m_stringBuilder.append(makeString(m_entryPointScaffolding->signature(iterator->value), " {"));
>             m_stringBuilder.append(m_entryPointScaffolding->unpack());
>+            for (size_t i = 0; i < functionDefinition.parameters().size(); ++i) {
>+                auto addResult = m_variableMapping.add(&functionDefinition.parameters()[i], m_entryPointScaffolding->parameterVariables()[i]);
>+                ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+            }
>             checkErrorAndVisit(functionDefinition.block());
>             ASSERT(m_stack.isEmpty());
>             m_stringBuilder.append("}\n");
>+            m_entryPointScaffolding = nullptr;
>         } else {
>-            m_entryPointScaffolding = WTF::nullopt;
>+            ASSERT(m_entryPointScaffolding == nullptr);
>             m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(functionDefinition.type()), ' ', iterator->value, '('));
>             for (size_t i = 0; i < functionDefinition.parameters().size(); ++i) {
>                 auto& parameter = functionDefinition.parameters()[i];
>@@ -579,11 +599,13 @@ private:
>     HashMap<AST::VariableDeclaration*, String> m_variableMapping;
>     StringBuilder m_stringBuilder;
>     Vector<String> m_stack;
>-    Optional<EntryPointScaffolding> m_entryPointScaffolding;
>+    std::unique_ptr<EntryPointScaffolding> m_entryPointScaffolding;
>+    Variant<MatchedRenderSemantics, MatchedComputeSemantics> m_matchedSemantics;
>+    Layout& m_layout;
>     unsigned m_variableCount { 0 };
> };
> 
>-String metalFunctions(Program& program, TypeNamer& typeNamer)
>+String metalFunctions(Program& program, TypeNamer& typeNamer, Variant<MatchedRenderSemantics, MatchedComputeSemantics>&& matchedSemantics, Layout& layout)
> {
>     StringBuilder stringBuilder;
> 
>@@ -612,7 +634,7 @@ String metalFunctions(Program& program, TypeNamer& typeNamer)
>     stringBuilder.append('\n');
> 
>     {
>-        FunctionDefinitionWriter functionDefinitionWriter(program.intrinsics(), typeNamer, functionMapping);
>+        FunctionDefinitionWriter functionDefinitionWriter(program.intrinsics(), typeNamer, functionMapping, WTFMove(matchedSemantics), layout);
>         for (auto& nativeFunctionDeclaration : program.nativeFunctionDeclarations())
>             functionDefinitionWriter.visit(nativeFunctionDeclaration);
>         for (auto& functionDefinition : program.functionDefinitions())
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.h b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.h
>index 9960e5e48d58c24e35e528382456472b7d09b6d7..d7d2b8dab26b5a54a376b8117ab3fd51bd2a4e30 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.h
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.h
>@@ -27,6 +27,8 @@
> 
> #if ENABLE(WEBGPU)
> 
>+#include "WHLSLSemanticMatcher.h"
>+
> namespace WebCore {
> 
> namespace WHLSL {
>@@ -37,7 +39,7 @@ namespace Metal {
> 
> class TypeNamer;
> 
>-String metalFunctions(Program&, TypeNamer&);
>+String metalFunctions(Program&, TypeNamer&, Variant<MatchedRenderSemantics, MatchedComputeSemantics>&&, Layout&);
> 
> }
> 
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.cpp b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.cpp
>index 10c0530e290d7041767f71823270d3b92597b54b..2ec5756d3d47c5b0637d96dc4726b8779e6367a1 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.cpp
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.cpp
>@@ -38,7 +38,7 @@ namespace WHLSL {
> 
> namespace Metal {
> 
>-String generateMetalCode(Program& program)
>+String generateMetalCode(Program& program, Variant<MatchedRenderSemantics, MatchedComputeSemantics>&& matchedSemantics, Layout& layout)
> {
>     StringBuilder stringBuilder;
>     stringBuilder.append("#include <metal_stdlib>\n");
>@@ -53,7 +53,7 @@ String generateMetalCode(Program& program)
> 
>     TypeNamer typeNamer(program);
>     stringBuilder.append(typeNamer.metalTypes());
>-    stringBuilder.append(metalFunctions(program, typeNamer));
>+    stringBuilder.append(metalFunctions(program, typeNamer, WTFMove(matchedSemantics), layout));
>     return stringBuilder.toString();
> }
> 
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.h b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.h
>index e088e010449948411075674347491f73d36a234f..7ea73b8d3c17142dee2e351259419bd9a7027a37 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.h
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.h
>@@ -27,6 +27,7 @@
> 
> #if ENABLE(WEBGPU)
> 
>+#include <wtf/HashSet.h>
> #include <wtf/text/WTFString.h>
> 
> namespace WebCore {
>@@ -38,7 +39,7 @@ class Program;
> namespace Metal {
> 
> // FIXME: This needs to know about the pipeline state object to emit function prologues and epilogues.
>-String generateMetalCode(Program&);
>+String generateMetalCode(Program&, Variant<MatchedRenderSemantics, MatchedComputeSemantics>&& matchedSemantics, Layout&);
> 
> }
> 
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLIntrinsics.h b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLIntrinsics.h
>index 8410c8b7392c7f00144a901d5ec462d496aba81c..fc47d3293540df6cca536631d4c9fbd2ece8979f 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLIntrinsics.h
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLIntrinsics.h
>@@ -82,6 +82,12 @@ public:
>         return *m_floatType;
>     }
> 
>+    AST::NativeTypeDeclaration& float2Type() const
>+    {
>+        ASSERT(m_vectorFloat[0]);
>+        return *m_vectorFloat[0];
>+    }
>+
>     AST::NativeTypeDeclaration& float3Type() const
>     {
>         ASSERT(m_vectorFloat[1]);
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPipelineDescriptor.h b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPipelineDescriptor.h
>new file mode 100644
>index 0000000000000000000000000000000000000000..0f28692c184bbf6a061fc7942525245092122ab6
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPipelineDescriptor.h
>@@ -0,0 +1,116 @@
>+/*
>+ * Copyright (C) 2019 Apple Inc. All rights reserved.
>+ *
>+ * Redistribution and use in source and binary forms, with or without
>+ * modification, are permitted provided that the following conditions
>+ * are met:
>+ * 1. Redistributions of source code must retain the above copyright
>+ *    notice, this list of conditions and the following disclaimer.
>+ * 2. Redistributions in binary form must reproduce the above copyright
>+ *    notice, this list of conditions and the following disclaimer in the
>+ *    documentation and/or other materials provided with the distribution.
>+ *
>+ * THIS SOFTWARE IS PROVIDED BY APPLE INC. AND ITS CONTRIBUTORS ``AS IS''
>+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO,
>+ * THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
>+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL APPLE INC. OR ITS CONTRIBUTORS
>+ * BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
>+ * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
>+ * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
>+ * INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
>+ * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
>+ * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
>+ * THE POSSIBILITY OF SUCH DAMAGE.
>+ */
>+
>+#pragma once
>+
>+#if ENABLE(WEBGPU)
>+
>+#include <wtf/OptionSet.h>
>+#include <wtf/Variant.h>
>+#include <wtf/Vector.h>
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+enum class VertexFormat : uint8_t {
>+    FloatR32G32B32A32,
>+    FloatR32G32B32,
>+    FloatR32G32,
>+    FloatR32
>+};
>+
>+struct VertexAttribute {
>+    VertexFormat vertexFormat;
>+    unsigned name;
>+};
>+
>+using VertexAttributes = Vector<VertexAttribute>;
>+
>+enum class TextureFormat {
>+    R8G8B8A8Unorm,
>+    R8G8B8A8Uint,
>+    B8G8R8A8Unorm,
>+    D32FloatS8Uint
>+};
>+
>+struct AttachmentDescriptor {
>+    TextureFormat textureFormat;
>+    unsigned name;
>+};
>+
>+struct AttachmentsStateDescriptor {
>+    Vector<AttachmentDescriptor> attachmentDescriptors;
>+    Optional<AttachmentDescriptor> depthStencilAttachmentDescriptor;
>+};
>+
>+enum class ShaderStage : uint8_t {
>+    None = 0,
>+    Vertex = 1 << 0,
>+    Fragment = 1 << 1,
>+    Compute = 1 << 2
>+};
>+
>+enum class BindingType : uint8_t {
>+    UniformBuffer,
>+    Sampler,
>+    Texture,
>+    StorageBuffer,
>+    // FIXME: Add the dynamic types
>+};
>+
>+struct Binding {
>+    OptionSet<ShaderStage> visibility;
>+    BindingType bindingType;
>+    unsigned name;
>+};
>+
>+struct BindGroup {
>+    Vector<Binding> bindings;
>+    unsigned name;
>+};
>+
>+using Layout = Vector<BindGroup>;
>+
>+struct RenderPipelineDescriptor {
>+    VertexAttributes vertexAttributes;
>+    AttachmentsStateDescriptor attachmentsStateDescriptor;
>+    Layout layout;
>+    String vertexEntryPointName;
>+    String fragmentEntryPointName;
>+};
>+
>+struct ComputePipelineDescriptor {
>+    Layout layout;
>+    String entryPointName;
>+};
>+
>+using PipelineDescriptor = Variant<RenderPipelineDescriptor, ComputePipelineDescriptor>;
>+
>+} // namespace WHLSL
>+
>+} // namespace WebCore
>+
>+#endif // ENABLE(WEBGPU)
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLSemanticMatcher.cpp b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLSemanticMatcher.cpp
>new file mode 100644
>index 0000000000000000000000000000000000000000..a96de947e9ab47d7972835039482b1dcc7084fd3
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLSemanticMatcher.cpp
>@@ -0,0 +1,303 @@
>+/*
>+ * Copyright (C) 2019 Apple Inc. All rights reserved.
>+ *
>+ * Redistribution and use in source and binary forms, with or without
>+ * modification, are permitted provided that the following conditions
>+ * are met:
>+ * 1. Redistributions of source code must retain the above copyright
>+ *    notice, this list of conditions and the following disclaimer.
>+ * 2. Redistributions in binary form must reproduce the above copyright
>+ *    notice, this list of conditions and the following disclaimer in the
>+ *    documentation and/or other materials provided with the distribution.
>+ *
>+ * THIS SOFTWARE IS PROVIDED BY APPLE INC. AND ITS CONTRIBUTORS ``AS IS''
>+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO,
>+ * THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
>+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL APPLE INC. OR ITS CONTRIBUTORS
>+ * BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
>+ * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
>+ * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
>+ * INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
>+ * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
>+ * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
>+ * THE POSSIBILITY OF SUCH DAMAGE.
>+ */
>+
>+#include "config.h"
>+#include "WHLSLSemanticMatcher.h"
>+
>+#if ENABLE(WEBGPU)
>+
>+#include "WHLSLBuiltInSemantic.h"
>+#include "WHLSLFunctionDefinition.h"
>+#include "WHLSLGatherEntryPointItems.h"
>+#include "WHLSLInferTypes.h"
>+#include "WHLSLPipelineDescriptor.h"
>+#include "WHLSLProgram.h"
>+#include "WHLSLResourceSemantic.h"
>+#include "WHLSLStageInOutSemantic.h"
>+#include <wtf/HashMap.h>
>+#include <wtf/HashSet.h>
>+#include <wtf/Optional.h>
>+#include <wtf/text/WTFString.h>
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+static AST::FunctionDefinition* findEntryPoint(Vector<UniqueRef<AST::FunctionDefinition>>& functionDefinitions, String& name)
>+{
>+    auto iterator = std::find_if(functionDefinitions.begin(), functionDefinitions.end(), [&](AST::FunctionDefinition& functionDefinition) {
>+        return functionDefinition.entryPointType() && functionDefinition.name() == name;
>+    });
>+    if (iterator == functionDefinitions.end())
>+        return nullptr;
>+    return &*iterator;
>+};
>+
>+static bool matchMode(BindingType bindingType, AST::ResourceSemantic::Mode mode)
>+{
>+    switch (bindingType) {
>+    case BindingType::UniformBuffer:
>+        return mode == AST::ResourceSemantic::Mode::Buffer;
>+    case BindingType::Sampler:
>+        return mode == AST::ResourceSemantic::Mode::Sampler;
>+    case BindingType::Texture:
>+        return mode == AST::ResourceSemantic::Mode::Texture;
>+    default:
>+        ASSERT(bindingType == BindingType::StorageBuffer);
>+        return mode == AST::ResourceSemantic::Mode::UnorderedAccessView;
>+    }
>+}
>+
>+static Optional<HashMap<Binding*, size_t>> matchResources(Vector<EntryPointItem>& entryPointItems, Layout& layout, ShaderStage shaderStage)
>+{
>+    HashMap<Binding*, size_t> result;
>+    HashSet<size_t> itemIndices;
>+    if (entryPointItems.size() == std::numeric_limits<size_t>::max())
>+        return WTF::nullopt; // Work around the fact that HashSet's keys are restricted.
>+    for (auto& bindGroup : layout) {
>+        auto space = bindGroup.name;
>+        for (auto& binding : bindGroup.bindings) {
>+            if (!binding.visibility.contains(shaderStage))
>+                continue;
>+            for (size_t i = 0; i < entryPointItems.size(); ++i) {
>+                auto& item = entryPointItems[i];
>+                auto& semantic = *item.semantic;
>+                if (!WTF::holds_alternative<AST::ResourceSemantic>(semantic))
>+                    continue;
>+                auto& resourceSemantic = WTF::get<AST::ResourceSemantic>(semantic);
>+                if (!matchMode(binding.bindingType, resourceSemantic.mode()))
>+                    continue;
>+                if (binding.name != resourceSemantic.index())
>+                    continue;
>+                if (space != resourceSemantic.space())
>+                    continue;
>+                result.add(&binding, i);
>+                itemIndices.add(i + 1); // Work around the fact that HashSet's keys are restricted.
>+            }
>+        }
>+    }
>+
>+    for (size_t i = 0; i < entryPointItems.size(); ++i) {
>+        auto& item = entryPointItems[i];
>+        auto& semantic = *item.semantic;
>+        if (!WTF::holds_alternative<AST::ResourceSemantic>(semantic))
>+            continue;
>+        if (!itemIndices.contains(i + 1))
>+            return WTF::nullopt;
>+    }
>+
>+    return result;
>+}
>+
>+static bool matchInputsOutputs(Vector<EntryPointItem>& vertexOutputs, Vector<EntryPointItem>& fragmentInputs)
>+{
>+    for (auto& fragmentInput : fragmentInputs) {
>+        if (!WTF::holds_alternative<AST::StageInOutSemantic>(*fragmentInput.semantic))
>+            continue;
>+        auto& fragmentInputStageInOutSemantic = downcast<AST::StageInOutSemantic>(*fragmentInput.semantic);
>+        bool found = false;
>+        for (auto& vertexOutput : vertexOutputs) {
>+            if (!WTF::holds_alternative<AST::StageInOutSemantic>(*vertexOutput.semantic))
>+                continue;
>+            auto& vertexOutputStageInOutSemantic = downcast<AST::StageInOutSemantic>(*vertexOutput.semantic);
>+            if (fragmentInputStageInOutSemantic.index() == vertexOutputStageInOutSemantic.index()) {
>+                if (matches(*fragmentInput.unnamedType, *vertexOutput.unnamedType)) {
>+                    found = true;
>+                    break;
>+                }
>+                return false;
>+            }
>+        }
>+        if (!found)
>+            return false;
>+    }
>+    return true;
>+}
>+
>+static bool isAcceptableFormat(VertexFormat vertexFormat, AST::UnnamedType& unnamedType, Intrinsics& intrinsics)
>+{
>+    switch (vertexFormat) {
>+    case VertexFormat::FloatR32G32B32A32:
>+        return matches(unnamedType, intrinsics.float4Type());
>+    case VertexFormat::FloatR32G32B32:
>+        return matches(unnamedType, intrinsics.float3Type());
>+    case VertexFormat::FloatR32G32:
>+        return matches(unnamedType, intrinsics.float2Type());
>+    default:
>+        ASSERT(vertexFormat == VertexFormat::FloatR32);
>+        return matches(unnamedType, intrinsics.floatType());
>+    }
>+}
>+
>+static Optional<HashMap<VertexAttribute*, size_t>> matchVertexAttributes(Vector<EntryPointItem>& vertexInputs, VertexAttributes& vertexAttributes, Intrinsics& intrinsics)
>+{
>+    HashMap<VertexAttribute*, size_t> result;
>+    HashSet<size_t> itemIndices;
>+    if (vertexInputs.size() == std::numeric_limits<size_t>::max())
>+        return WTF::nullopt; // Work around the fact that HashSet's keys are restricted.
>+    for (auto& vertexAttribute : vertexAttributes) {
>+        for (size_t i = 0; i < vertexInputs.size(); ++i) {
>+            auto& item = vertexInputs[i];
>+            auto& semantic = *item.semantic;
>+            if (!WTF::holds_alternative<AST::StageInOutSemantic>(semantic))
>+                continue;
>+            auto& stageInOutSemantic = WTF::get<AST::StageInOutSemantic>(semantic);
>+            if (stageInOutSemantic.index() != vertexAttribute.name)
>+                continue;
>+            if (!isAcceptableFormat(vertexAttribute.vertexFormat, *item.unnamedType, intrinsics))
>+                return WTF::nullopt;
>+            result.add(&vertexAttribute, i);
>+            itemIndices.add(i + 1); // Work around the fact that HashSet's keys are restricted.
>+        }
>+    }
>+
>+    for (size_t i = 0; i < vertexInputs.size(); ++i) {
>+        auto& item = vertexInputs[i];
>+        auto& semantic = *item.semantic;
>+        if (!WTF::holds_alternative<AST::StageInOutSemantic>(semantic))
>+            continue;
>+        if (!itemIndices.contains(i + 1))
>+            return WTF::nullopt;
>+    }
>+
>+    return result;
>+}
>+
>+static bool isAcceptableFormat(TextureFormat textureFormat, AST::UnnamedType& unnamedType, Intrinsics& intrinsics, bool isColor)
>+{
>+    if (isColor) {
>+        switch (textureFormat) {
>+        case TextureFormat::R8G8B8A8Unorm:
>+        case TextureFormat::R8G8B8A8Uint:
>+        case TextureFormat::B8G8R8A8Unorm:
>+            return matches(unnamedType, intrinsics.float4Type());
>+        default:
>+            ASSERT(textureFormat == TextureFormat::D32FloatS8Uint);
>+            return false;
>+        }
>+    }
>+    return textureFormat == TextureFormat::D32FloatS8Uint && matches(unnamedType, intrinsics.floatType());
>+}
>+
>+static Optional<HashMap<AttachmentDescriptor*, size_t>> matchColorAttachments(Vector<EntryPointItem>& fragmentOutputs, Vector<AttachmentDescriptor>& attachmentDescriptors, Intrinsics& intrinsics)
>+{
>+    HashMap<AttachmentDescriptor*, size_t> result;
>+    HashSet<size_t> itemIndices;
>+    if (attachmentDescriptors.size() == std::numeric_limits<size_t>::max())
>+        return WTF::nullopt; // Work around the fact that HashSet's keys are restricted.
>+    for (auto& attachmentDescriptor : attachmentDescriptors) {
>+        for (size_t i = 0; i < fragmentOutputs.size(); ++i) {
>+            auto& item = fragmentOutputs[i];
>+            auto& semantic = *item.semantic;
>+            if (!WTF::holds_alternative<AST::StageInOutSemantic>(semantic))
>+                continue;
>+            auto& stageInOutSemantic = WTF::get<AST::StageInOutSemantic>(semantic);
>+            if (stageInOutSemantic.index() != attachmentDescriptor.name)
>+                continue;
>+            if (!isAcceptableFormat(attachmentDescriptor.textureFormat, *item.unnamedType, intrinsics, true))
>+                return WTF::nullopt;
>+            result.add(&attachmentDescriptor, i);
>+            itemIndices.add(i + 1); // Work around the fact that HashSet's keys are restricted.
>+        }
>+    }
>+
>+    for (size_t i = 0; i < fragmentOutputs.size(); ++i) {
>+        auto& item = fragmentOutputs[i];
>+        auto& semantic = *item.semantic;
>+        if (!WTF::holds_alternative<AST::StageInOutSemantic>(semantic))
>+            continue;
>+        if (!itemIndices.contains(i + 1))
>+            return WTF::nullopt;
>+    }
>+
>+    return result;
>+}
>+
>+static bool matchDepthAttachment(Vector<EntryPointItem>& fragmentOutputs, Optional<AttachmentDescriptor>& depthStencilAttachmentDescriptor, Intrinsics& intrinsics)
>+{
>+    auto iterator = std::find_if(fragmentOutputs.begin(), fragmentOutputs.end(), [&](EntryPointItem& item) {
>+        auto& semantic = *item.semantic;
>+        if (!WTF::holds_alternative<AST::BuiltInSemantic>(semantic))
>+            return false;
>+        auto& builtInSemantic = WTF::get<AST::BuiltInSemantic>(semantic);
>+        return builtInSemantic.variable() == AST::BuiltInSemantic::Variable::SVDepth;
>+    });
>+    if (iterator == fragmentOutputs.end())
>+        return true;
>+
>+    if (depthStencilAttachmentDescriptor) {
>+        ASSERT(!depthStencilAttachmentDescriptor->name);
>+        return isAcceptableFormat(depthStencilAttachmentDescriptor->textureFormat, *iterator->unnamedType, intrinsics, false);
>+    }
>+    return false;
>+}
>+
>+Optional<MatchedRenderSemantics> matchSemantics(Program& program, RenderPipelineDescriptor& renderPipelineDescriptor)
>+{
>+    auto vertexShaderEntryPoint = findEntryPoint(program.functionDefinitions(), renderPipelineDescriptor.vertexEntryPointName);
>+    auto fragmentShaderEntryPoint = findEntryPoint(program.functionDefinitions(), renderPipelineDescriptor.fragmentEntryPointName);
>+    if (!vertexShaderEntryPoint || !fragmentShaderEntryPoint)
>+        return WTF::nullopt;
>+    auto vertexShaderEntryPointItems = gatherEntryPointItems(program.intrinsics(), *vertexShaderEntryPoint);
>+    auto fragmentShaderEntryPointItems = gatherEntryPointItems(program.intrinsics(), *fragmentShaderEntryPoint);
>+    if (!vertexShaderEntryPointItems || !fragmentShaderEntryPointItems)
>+        return WTF::nullopt;
>+    auto vertexShaderResourceMap = matchResources(vertexShaderEntryPointItems->inputs, renderPipelineDescriptor.layout, ShaderStage::Vertex);
>+    auto fragmentShaderResourceMap = matchResources(fragmentShaderEntryPointItems->inputs, renderPipelineDescriptor.layout, ShaderStage::Fragment);
>+    if (!vertexShaderResourceMap || !fragmentShaderResourceMap)
>+        return WTF::nullopt;
>+    if (!matchInputsOutputs(vertexShaderEntryPointItems->outputs, fragmentShaderEntryPointItems->inputs))
>+        return WTF::nullopt;
>+    auto matchedVertexAttributes = matchVertexAttributes(vertexShaderEntryPointItems->inputs, renderPipelineDescriptor.vertexAttributes, program.intrinsics());
>+    if (!matchedVertexAttributes)
>+        return WTF::nullopt;
>+    auto matchedColorAttachments = matchColorAttachments(fragmentShaderEntryPointItems->outputs, renderPipelineDescriptor.attachmentsStateDescriptor.attachmentDescriptors, program.intrinsics());
>+    if (!matchedColorAttachments)
>+        return WTF::nullopt;
>+    if (!matchDepthAttachment(fragmentShaderEntryPointItems->outputs, renderPipelineDescriptor.attachmentsStateDescriptor.depthStencilAttachmentDescriptor, program.intrinsics()))
>+        return WTF::nullopt;
>+    return {{ vertexShaderEntryPoint, fragmentShaderEntryPoint, *vertexShaderEntryPointItems, *fragmentShaderEntryPointItems, *vertexShaderResourceMap, *fragmentShaderResourceMap, *matchedVertexAttributes, *matchedColorAttachments }};
>+}
>+
>+Optional<MatchedComputeSemantics> matchSemantics(Program& program, ComputePipelineDescriptor& computePipelineDescriptor)
>+{
>+    auto entryPoint = findEntryPoint(program.functionDefinitions(), computePipelineDescriptor.entryPointName);
>+    if (!entryPoint)
>+        return WTF::nullopt;
>+    auto entryPointItems = gatherEntryPointItems(program.intrinsics(), *entryPoint);
>+    if (!entryPointItems)
>+        return WTF::nullopt;
>+    ASSERT(entryPointItems->outputs.isEmpty());
>+    auto resourceMap = matchResources(entryPointItems->inputs, computePipelineDescriptor.layout, ShaderStage::Compute);
>+    if (!resourceMap)
>+        return WTF::nullopt;
>+    return {{ entryPoint, *entryPointItems, *resourceMap }};
>+}
>+
>+} // namespace WHLSL
>+
>+} // namespace WebCore
>+
>+#endif // ENABLE(WEBGPU)
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLSemanticMatcher.h b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLSemanticMatcher.h
>new file mode 100644
>index 0000000000000000000000000000000000000000..4ff8c8d13ff0cd774e2654de1cf6f0cf25838667
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLSemanticMatcher.h
>@@ -0,0 +1,64 @@
>+/*
>+ * Copyright (C) 2019 Apple Inc. All rights reserved.
>+ *
>+ * Redistribution and use in source and binary forms, with or without
>+ * modification, are permitted provided that the following conditions
>+ * are met:
>+ * 1. Redistributions of source code must retain the above copyright
>+ *    notice, this list of conditions and the following disclaimer.
>+ * 2. Redistributions in binary form must reproduce the above copyright
>+ *    notice, this list of conditions and the following disclaimer in the
>+ *    documentation and/or other materials provided with the distribution.
>+ *
>+ * THIS SOFTWARE IS PROVIDED BY APPLE INC. AND ITS CONTRIBUTORS ``AS IS''
>+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO,
>+ * THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
>+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL APPLE INC. OR ITS CONTRIBUTORS
>+ * BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
>+ * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
>+ * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
>+ * INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
>+ * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
>+ * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
>+ * THE POSSIBILITY OF SUCH DAMAGE.
>+ */
>+
>+#pragma once
>+
>+#if ENABLE(WEBGPU)
>+
>+#include "WHLSLGatherEntryPointItems.h"
>+#include "WHLSLPipelineDescriptor.h"
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+class Program;
>+
>+struct MatchedRenderSemantics {
>+    AST::FunctionDefinition* vertexShader;
>+    AST::FunctionDefinition* fragmentShader;
>+    EntryPointItems vertexShaderEntryPointItems;
>+    EntryPointItems fragmentShaderEntryPointItems;
>+    HashMap<Binding*, size_t> vertexShaderResourceMap;
>+    HashMap<Binding*, size_t> fragmentShaderResourceMap;
>+    HashMap<VertexAttribute*, size_t> matchedVertexAttributes;
>+    HashMap<AttachmentDescriptor*, size_t> matchedColorAttachments;
>+};
>+
>+Optional<MatchedRenderSemantics> matchSemantics(Program&, RenderPipelineDescriptor&);
>+
>+struct MatchedComputeSemantics {
>+    AST::FunctionDefinition* shader;
>+    EntryPointItems entryPointItems;
>+    HashMap<Binding*, size_t> resourceMap;
>+};
>+
>+Optional<MatchedComputeSemantics> matchSemantics(Program&, ComputePipelineDescriptor&);
>+
>+} // namespace WHLSL
>+
>+} // namespace WebCore
>+
>+#endif // ENABLE(WEBGPU)
>diff --git a/Source/WebCore/Sources.txt b/Source/WebCore/Sources.txt
>index 40fa474a44eb8d2eab80b0986c5eebbf0d50511b..2749d74a23a229476797c6c79b08b8e6661c4a2d 100644
>--- a/Source/WebCore/Sources.txt
>+++ b/Source/WebCore/Sources.txt
>@@ -324,6 +324,7 @@ Modules/webgpu/WHLSL/WHLSLNameResolver.cpp
> Modules/webgpu/WHLSL/WHLSLResolveOverloadImpl.cpp
> Modules/webgpu/WHLSL/WHLSLRecursionChecker.cpp
> Modules/webgpu/WHLSL/WHLSLVisitor.cpp
>+Modules/webgpu/WHLSL/WHLSLSemanticMatcher.cpp
> Modules/webgpu/WHLSL/WHLSLLiteralTypeChecker.cpp
> Modules/webgpu/WHLSL/WHLSLHighZombieFinder.cpp
> Modules/webgpu/WHLSL/Metal/WHLSLTypeNamer.cpp
>diff --git a/Source/WebCore/WebCore.xcodeproj/project.pbxproj b/Source/WebCore/WebCore.xcodeproj/project.pbxproj
>index c3d8fdb0456340c9ccf04a02f7af143988642c64..dd619737ce95df2a95f8ee43551dc09aa5964fa0 100644
>--- a/Source/WebCore/WebCore.xcodeproj/project.pbxproj
>+++ b/Source/WebCore/WebCore.xcodeproj/project.pbxproj
>@@ -13391,6 +13391,9 @@
> 		C234A9B621E92CC0003C984D /* WHLSLIntrinsics.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLIntrinsics.h; sourceTree = "<group>"; };
> 		C234A9B721E92CC1003C984D /* WHLSLIntrinsics.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLIntrinsics.cpp; sourceTree = "<group>"; };
> 		C2458E611FE8979E00594759 /* FontCacheCoreText.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = FontCacheCoreText.h; sourceTree = "<group>"; };
>+		C24A57AF21FAD53F004C6DD1 /* WHLSLPipelineDescriptor.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLPipelineDescriptor.h; sourceTree = "<group>"; };
>+		C24A57B221FB8DDA004C6DD1 /* WHLSLSemanticMatcher.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLSemanticMatcher.cpp; sourceTree = "<group>"; };
>+		C24A57B321FB8DDA004C6DD1 /* WHLSLSemanticMatcher.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLSemanticMatcher.h; sourceTree = "<group>"; };
> 		C26017A11C72DC9900F74A16 /* CSSFontFaceSet.cpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.cpp; path = CSSFontFaceSet.cpp; sourceTree = "<group>"; };
> 		C26017A21C72DC9900F74A16 /* CSSFontFaceSet.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = CSSFontFaceSet.h; sourceTree = "<group>"; };
> 		C280833C1C6DB194001451B6 /* FontFace.idl */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = text; path = FontFace.idl; sourceTree = "<group>"; };
>@@ -25602,6 +25605,7 @@
> 				C234A98C21E8883E003C984D /* WHLSLNameResolver.h */,
> 				C21BF73721CD8A0200227979 /* WHLSLParser.cpp */,
> 				C21BF73821CD8A0300227979 /* WHLSLParser.h */,
>+				C24A57AF21FAD53F004C6DD1 /* WHLSLPipelineDescriptor.h */,
> 				C21BF73A21CD8D7000227979 /* WHLSLProgram.h */,
> 				1CA0C2E021EEB5F500A11860 /* WHLSLRecursionChecker.cpp */,
> 				1CA0C2DE21EEB5F400A11860 /* WHLSLRecursionChecker.h */,
>@@ -25610,6 +25614,8 @@
> 				C234A99921E90F29003C984D /* WHLSLResolveOverloadImpl.cpp */,
> 				C234A99721E90F28003C984D /* WHLSLResolveOverloadImpl.h */,
> 				C234A99D21E910BD003C984D /* WHLSLResolvingType.h */,
>+				C24A57B221FB8DDA004C6DD1 /* WHLSLSemanticMatcher.cpp */,
>+				C24A57B321FB8DDA004C6DD1 /* WHLSLSemanticMatcher.h */,
> 				C21BF74521CD969800227979 /* WHLSLStandardLibrary.txt */,
> 				1CECB3A921F2B67300F44542 /* WHLSLStatementBehaviorChecker.cpp */,
> 				1CECB3A821F2B67300F44542 /* WHLSLStatementBehaviorChecker.h */,

Actions: View | Formatted Diff | Diff

Attachments on bug 193877: 360258 | 360298 | 360319 | 363689 | 363790