Attachment 360319 Details for Bug 193877 – Patch

[patch] Patch

bug-193877-20190127224533.patch (text/plain), 125.19 KB, created by Myles C. Maxfield on 2019-01-27 22:45:34 PST

(hide)

Description:

Filename:

MIME Type:

Creator: Myles C. Maxfield

Created: 2019-01-27 22:45:34 PST

Size: 125.19 KB

patch

obsolete

>Subversion Revision: 240456
>diff --git a/Source/WebCore/ChangeLog b/Source/WebCore/ChangeLog
>index a112c3fdb6b29278746a261630269c0e11d6170f..e9e5dd6d78ce4214acb1a40aae0f496fa78d84a8 100644
>--- a/Source/WebCore/ChangeLog
>+++ b/Source/WebCore/ChangeLog
>@@ -1,3 +1,86 @@
>+2019-01-27  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::mappedBindGroups const):
>+        (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::FunctionDeclarationWriter::toString):
>+        (WebCore::WHLSL::Metal::FunctionDeclarationWriter::visit):
>+        (WebCore::WHLSL::Metal::FunctionDefinitionWriter::FunctionDefinitionWriter):
>+        (WebCore::WHLSL::Metal::FunctionDefinitionWriter::generateNextVariableName):
>+        (WebCore::WHLSL::Metal::FunctionDefinitionWriter::visit):
>+        (WebCore::WHLSL::Metal::FunctionDefinitionWriter::constantExpressionString):
>+        (WebCore::WHLSL::Metal::RenderFunctionDefinitionWriter::RenderFunctionDefinitionWriter):
>+        (WebCore::WHLSL::Metal::RenderFunctionDefinitionWriter::takeVertexMappedBindGroups):
>+        (WebCore::WHLSL::Metal::RenderFunctionDefinitionWriter::takeFragmentMappedBindGroups):
>+        (WebCore::WHLSL::Metal::RenderFunctionDefinitionWriter::createEntryPointScaffolding):
>+        (WebCore::WHLSL::Metal::ComputeFunctionDefinitionWriter::ComputeFunctionDefinitionWriter):
>+        (WebCore::WHLSL::Metal::ComputeFunctionDefinitionWriter::takeMappedBindGroups):
>+        (WebCore::WHLSL::Metal::ComputeFunctionDefinitionWriter::createEntryPointScaffolding):
>+        (WebCore::WHLSL::Metal::sharedMetalFunctions):
>+        (WebCore::WHLSL::Metal::metalFunctions):
>+        * Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.h:
>+        * Modules/webgpu/WHLSL/Metal/WHLSLMappedBindings.h: Copied from Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.h.
>+        * Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.cpp:
>+        (WebCore::WHLSL::Metal::generateMetalCodeShared):
>+        (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/WHLSLPrepare.cpp: Added.
>+        (WebCore::WHLSL::prepareShared):
>+        (WebCore::WHLSL::prepare):
>+        * Modules/webgpu/WHLSL/WHLSLPrepare.h: Copied from Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.h.
>+        * 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..035284ca4fcf32f876810bb7372fb8152a3e7a12 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,474 @@ 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()
>+MappedBindGroups EntryPointScaffolding::mappedBindGroups() const
> {
>-    // FIXME: Implement this.
>-    return String();
>+    MappedBindGroups result;
>+    result.reserveInitialCapacity(m_layout.size());
>+    for (auto& namedBindGroup : m_namedBindGroups) {
>+        MappedBindGroup mappedBindGroup;
>+        mappedBindGroup.argumentBufferIndex = namedBindGroup.argumentBufferIndex;
>+        mappedBindGroup.bindingIndices.reserveInitialCapacity(namedBindGroup.namedBindings.size());
>+        for (auto& namedBinding : namedBindGroup.namedBindings)
>+            mappedBindGroup.bindingIndices.uncheckedAppend(namedBinding.index);
>+        result.uncheckedAppend(WTFMove(mappedBindGroup));
>+    }
>+    return result;
>+}
>+
>+String EntryPointScaffolding::resourceHelperTypes()
>+{
>+    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..d87f6e9f54e77164fcb41fe6a65e0c5c50589a81 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h
>@@ -27,6 +27,12 @@
> 
> #if ENABLE(WEBGPU)
> 
>+#include "WHLSLMappedBindings.h"
>+#include "WHLSLPipelineDescriptor.h"
>+#include <wtf/HashMap.h>
>+#include <wtf/Optional.h>
>+#include <wtf/text/WTFString.h>
>+
> namespace WebCore {
> 
> namespace WHLSL {
>@@ -37,23 +43,132 @@ class FunctionDefinition;
> 
> }
> 
>+struct EntryPointItems;
> class Intrinsics;
> 
> namespace Metal {
> 
>-// FIXME: This needs to know about the pipeline state object to emit function prologues and epilogues.
>+class TypeNamer;
>+
> 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;
>+
>+    MappedBindGroups mappedBindGroups() const;
>+    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..1b7a42abcde62c029539dc4af1b0f3bd6aaa465c 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp
>@@ -50,6 +50,7 @@
> #include "WHLSLLogicalNotExpression.h"
> #include "WHLSLMakeArrayReferenceExpression.h"
> #include "WHLSLMakePointerExpression.h"
>+#include "WHLSLMappedBindings.h"
> #include "WHLSLNativeFunctionDeclaration.h"
> #include "WHLSLNativeFunctionWriter.h"
> #include "WHLSLNativeTypeDeclaration.h"
>@@ -86,21 +87,8 @@ public:
>     virtual ~FunctionDeclarationWriter() = default;
> 
>     String toString() { return m_stringBuilder.toString(); }
>-    void visit(AST::FunctionDeclaration& functionDeclaration) override
>-    {
>-        if (functionDeclaration.entryPointType())
>-            return;
> 
>-        auto iterator = m_functionMapping.find(&functionDeclaration);
>-        ASSERT(iterator != m_functionMapping.end());
>-        m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(functionDeclaration.type()), ' ', iterator->value, '('));
>-        for (size_t i = 0; i < functionDeclaration.parameters().size(); ++i) {
>-            if (i)
>-                m_stringBuilder.append(", ");
>-            m_stringBuilder.append(m_typeNamer.mangledNameForType(*functionDeclaration.parameters()[i].type()));
>-        }
>-        m_stringBuilder.append(");\n");
>-    }
>+    void visit(AST::FunctionDeclaration&) override;
> 
> private:
>     TypeNamer& m_typeNamer;
>@@ -108,12 +96,29 @@ private:
>     StringBuilder m_stringBuilder;
> };
> 
>+void FunctionDeclarationWriter::visit(AST::FunctionDeclaration& functionDeclaration)
>+{
>+    if (functionDeclaration.entryPointType())
>+        return;
>+
>+    auto iterator = m_functionMapping.find(&functionDeclaration);
>+    ASSERT(iterator != m_functionMapping.end());
>+    m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(functionDeclaration.type()), ' ', iterator->value, '('));
>+    for (size_t i = 0; i < functionDeclaration.parameters().size(); ++i) {
>+        if (i)
>+            m_stringBuilder.append(", ");
>+        m_stringBuilder.append(m_typeNamer.mangledNameForType(*functionDeclaration.parameters()[i].type()));
>+    }
>+    m_stringBuilder.append(");\n");
>+}
>+
> 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, Layout& layout)
>         : m_intrinsics(intrinsics)
>         , m_typeNamer(typeNamer)
>         , m_functionMapping(functionMapping)
>+        , m_layout(layout)
>     {
>     }
> 
>@@ -121,469 +126,610 @@ public:
> 
>     String toString() { return m_stringBuilder.toString(); }
> 
>-    void visit(AST::NativeFunctionDeclaration& nativeFunctionDeclaration) override
>-    {
>-        auto iterator = m_functionMapping.find(&nativeFunctionDeclaration);
>-        ASSERT(iterator != m_functionMapping.end());
>-        m_stringBuilder.append(writeNativeFunction(nativeFunctionDeclaration, iterator->value, m_typeNamer));
>-    }
>+    void visit(AST::NativeFunctionDeclaration&) override;
>+    void visit(AST::FunctionDefinition&) override;
>+
>+protected:
>+    virtual std::unique_ptr<EntryPointScaffolding> createEntryPointScaffolding(AST::FunctionDefinition&) = 0;
>+
>+    void visit(AST::FunctionDeclaration&) override;
>+    void visit(AST::Statement&) override;
>+    void visit(AST::Block&) override;
>+    void visit(AST::Break&) override;
>+    void visit(AST::Continue&) override;
>+    void visit(AST::DoWhileLoop&) override;
>+    void visit(AST::EffectfulExpressionStatement&) override;
>+    void visit(AST::Fallthrough&) override;
>+    void visit(AST::ForLoop&) override;
>+    void visit(AST::IfStatement&) override;
>+    void visit(AST::Return&) override;
>+    void visit(AST::SwitchStatement&) override;
>+    void visit(AST::SwitchCase&) override;
>+    void visit(AST::Trap&) override;
>+    void visit(AST::VariableDeclarationsStatement&) override;
>+    void visit(AST::WhileLoop&) override;
>+    void visit(AST::IntegerLiteral&) override;
>+    void visit(AST::UnsignedIntegerLiteral&) override;
>+    void visit(AST::FloatLiteral&) override;
>+    void visit(AST::NullLiteral&) override;
>+    void visit(AST::BooleanLiteral&) override;
>+    void visit(AST::EnumerationMemberLiteral&) override;
>+    void visit(AST::Expression&) override;
>+    void visit(AST::DotExpression&) override;
>+    void visit(AST::IndexExpression&) override;
>+    void visit(AST::PropertyAccessExpression&) override;
>+    void visit(AST::VariableDeclaration&) override;
>+    void visit(AST::AssignmentExpression&) override;
>+    void visit(AST::CallExpression&) override;
>+    void visit(AST::CommaExpression&) override;
>+    void visit(AST::DereferenceExpression&) override;
>+    void visit(AST::LogicalExpression&) override;
>+    void visit(AST::LogicalNotExpression&) override;
>+    void visit(AST::MakeArrayReferenceExpression&) override;
>+    void visit(AST::MakePointerExpression&) override;
>+    void visit(AST::ReadModifyWriteExpression&) override;
>+    void visit(AST::TernaryExpression&) override;
>+    void visit(AST::VariableReference&) override;
>+
>+    String constantExpressionString(AST::ConstantExpression&);
> 
>-    void visit(AST::FunctionDefinition& functionDefinition) override
>+    String generateNextVariableName()
>     {
>-        auto iterator = m_functionMapping.find(&functionDefinition);
>-        ASSERT(iterator != m_functionMapping.end());
>-        if (functionDefinition.entryPointType()) {
>-            m_entryPointScaffolding = EntryPointScaffolding(functionDefinition, m_intrinsics);
>-            m_stringBuilder.append(m_entryPointScaffolding->helperTypes());
>-            m_stringBuilder.append('\n');
>-            m_stringBuilder.append(makeString(m_entryPointScaffolding->signature(), " {"));
>-            m_stringBuilder.append(m_entryPointScaffolding->unpack());
>-            checkErrorAndVisit(functionDefinition.block());
>-            ASSERT(m_stack.isEmpty());
>-            m_stringBuilder.append("}\n");
>-        } else {
>-            m_entryPointScaffolding = WTF::nullopt;
>-            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];
>-                if (i)
>-                    m_stringBuilder.append(", ");
>-                auto parameterName = generateNextVariableName();
>-                auto addResult = m_variableMapping.add(&parameter, parameterName);
>-                ASSERT_UNUSED(addResult, addResult.isNewEntry);
>-                m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*parameter.type()), ' ', parameterName));
>-            }
>-            m_stringBuilder.append(") {\n");
>-            checkErrorAndVisit(functionDefinition.block());
>-            ASSERT(m_stack.isEmpty());
>-            m_stringBuilder.append("}\n");
>-        }
>+        return makeString("variable", m_variableCount++);
>     }
> 
>-private:
>-    void visit(AST::FunctionDeclaration&) override
>-    {
>-        ASSERT_NOT_REACHED();
>-    }
>+    Intrinsics& m_intrinsics;
>+    TypeNamer& m_typeNamer;
>+    HashMap<AST::FunctionDeclaration*, String>& m_functionMapping;
>+    HashMap<AST::VariableDeclaration*, String> m_variableMapping;
>+    StringBuilder m_stringBuilder;
>+    Vector<String> m_stack;
>+    std::unique_ptr<EntryPointScaffolding> m_entryPointScaffolding;
>+    Layout& m_layout;
>+    unsigned m_variableCount { 0 };
>+};
> 
>-    void visit(AST::Statement& statement) override
>-    {
>-        Visitor::visit(statement);
>-    }
>+void FunctionDefinitionWriter::visit(AST::NativeFunctionDeclaration& nativeFunctionDeclaration)
>+{
>+    auto iterator = m_functionMapping.find(&nativeFunctionDeclaration);
>+    ASSERT(iterator != m_functionMapping.end());
>+    m_stringBuilder.append(writeNativeFunction(nativeFunctionDeclaration, iterator->value, m_typeNamer));
>+}
> 
>-    void visit(AST::Block& block) override
>-    {
>-        m_stringBuilder.append("{\n");
>-        for (auto& statement : block.statements())
>-            checkErrorAndVisit(statement);
>+void FunctionDefinitionWriter::visit(AST::FunctionDefinition& functionDefinition)
>+{
>+    auto iterator = m_functionMapping.find(&functionDefinition);
>+    ASSERT(iterator != m_functionMapping.end());
>+    if (functionDefinition.entryPointType()) {
>+        auto entryPointScaffolding = createEntryPointScaffolding(functionDefinition);
>+        if (!entryPointScaffolding)
>+            return;
>+        m_entryPointScaffolding = WTFMove(entryPointScaffolding);
>+        m_stringBuilder.append(m_entryPointScaffolding->helperTypes());
>+        m_stringBuilder.append('\n');
>+        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 {
>+        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];
>+            if (i)
>+                m_stringBuilder.append(", ");
>+            auto parameterName = generateNextVariableName();
>+            auto addResult = m_variableMapping.add(&parameter, parameterName);
>+            ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+            m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*parameter.type()), ' ', parameterName));
>+        }
>+        m_stringBuilder.append(") {\n");
>+        checkErrorAndVisit(functionDefinition.block());
>+        ASSERT(m_stack.isEmpty());
>         m_stringBuilder.append("}\n");
>     }
>+}
> 
>-    void visit(AST::Break&) override
>-    {
>-        m_stringBuilder.append("break;\n");
>-    }
>+void FunctionDefinitionWriter::visit(AST::FunctionDeclaration&)
>+{
>+    ASSERT_NOT_REACHED();
>+}
> 
>-    void visit(AST::Continue&) override
>-    {
>-        // FIXME: Figure out which loop we're in, and run the increment code
>-        CRASH();
>-    }
>+void FunctionDefinitionWriter::visit(AST::Statement& statement)
>+{
>+    Visitor::visit(statement);
>+}
> 
>-    void visit(AST::DoWhileLoop& doWhileLoop) override
>-    {
>-        m_stringBuilder.append("do {\n");
>-        checkErrorAndVisit(doWhileLoop.body());
>-        checkErrorAndVisit(doWhileLoop.conditional());
>-        m_stringBuilder.append(makeString("if (!", m_stack.takeLast(), ") break;\n"));
>-        m_stringBuilder.append(makeString("} while(true);\n"));
>-    }
>+void FunctionDefinitionWriter::visit(AST::Block& block)
>+{
>+    m_stringBuilder.append("{\n");
>+    for (auto& statement : block.statements())
>+        checkErrorAndVisit(statement);
>+    m_stringBuilder.append("}\n");
>+}
> 
>-    void visit(AST::EffectfulExpressionStatement& effectfulExpressionStatement) override
>-    {
>-        checkErrorAndVisit(effectfulExpressionStatement.effectfulExpression());
>-        m_stack.takeLast(); // The statement is already effectful, so we don't need to do anything with the result.
>-    }
>+void FunctionDefinitionWriter::visit(AST::Break&)
>+{
>+    m_stringBuilder.append("break;\n");
>+}
> 
>-    void visit(AST::Fallthrough&) override
>-    {
>-        m_stringBuilder.append("[[clang::fallthrough]];\n"); // FIXME: Make sure this is okay. Alternatively, we could do nothing and just return here instead.
>-    }
>+void FunctionDefinitionWriter::visit(AST::Continue&)
>+{
>+    // FIXME: Figure out which loop we're in, and run the increment code
>+    CRASH();
>+}
> 
>-    void visit(AST::ForLoop& forLoop) override
>-    {
>-        WTF::visit(WTF::makeVisitor([&](AST::VariableDeclarationsStatement& variableDeclarationsStatement) {
>-            checkErrorAndVisit(variableDeclarationsStatement);
>-        }, [&](UniqueRef<AST::Expression>& expression) {
>-            checkErrorAndVisit(expression);
>-            m_stack.takeLast(); // We don't need to do anything with the result.
>-        }), forLoop.initialization());
>-
>-        m_stringBuilder.append("for ( ; ; ) {\n");
>-        if (forLoop.condition()) {
>-            checkErrorAndVisit(*forLoop.condition());
>-            m_stringBuilder.append(makeString("if (!", m_stack.takeLast(), ") break;\n"));
>-        }
>-        checkErrorAndVisit(forLoop.body());
>-        if (forLoop.increment()) {
>-            checkErrorAndVisit(*forLoop.increment());
>-            m_stack.takeLast();
>-        }
>-        m_stringBuilder.append("}\n");
>+void FunctionDefinitionWriter::visit(AST::DoWhileLoop& doWhileLoop)
>+{
>+    m_stringBuilder.append("do {\n");
>+    checkErrorAndVisit(doWhileLoop.body());
>+    checkErrorAndVisit(doWhileLoop.conditional());
>+    m_stringBuilder.append(makeString("if (!", m_stack.takeLast(), ") break;\n"));
>+    m_stringBuilder.append(makeString("} while(true);\n"));
>+}
>+
>+void FunctionDefinitionWriter::visit(AST::EffectfulExpressionStatement& effectfulExpressionStatement)
>+{
>+    checkErrorAndVisit(effectfulExpressionStatement.effectfulExpression());
>+    m_stack.takeLast(); // The statement is already effectful, so we don't need to do anything with the result.
>+}
>+
>+void FunctionDefinitionWriter::visit(AST::Fallthrough&)
>+{
>+    m_stringBuilder.append("[[clang::fallthrough]];\n"); // FIXME: Make sure this is okay. Alternatively, we could do nothing and just return here instead.
>+}
>+
>+void FunctionDefinitionWriter::visit(AST::ForLoop& forLoop)
>+{
>+    WTF::visit(WTF::makeVisitor([&](AST::VariableDeclarationsStatement& variableDeclarationsStatement) {
>+        checkErrorAndVisit(variableDeclarationsStatement);
>+    }, [&](UniqueRef<AST::Expression>& expression) {
>+        checkErrorAndVisit(expression);
>+        m_stack.takeLast(); // We don't need to do anything with the result.
>+    }), forLoop.initialization());
>+
>+    m_stringBuilder.append("for ( ; ; ) {\n");
>+    if (forLoop.condition()) {
>+        checkErrorAndVisit(*forLoop.condition());
>+        m_stringBuilder.append(makeString("if (!", m_stack.takeLast(), ") break;\n"));
>     }
>+    checkErrorAndVisit(forLoop.body());
>+    if (forLoop.increment()) {
>+        checkErrorAndVisit(*forLoop.increment());
>+        m_stack.takeLast();
>+    }
>+    m_stringBuilder.append("}\n");
>+}
> 
>-    void visit(AST::IfStatement& ifStatement) override
>-    {
>-        checkErrorAndVisit(ifStatement.conditional());
>-        m_stringBuilder.append(makeString("if (", m_stack.takeLast(), ") {\n"));
>-        checkErrorAndVisit(ifStatement.body());
>-        if (ifStatement.elseBody()) {
>-            m_stringBuilder.append("} else {\n");
>-            checkErrorAndVisit(*ifStatement.elseBody());
>-        }
>-        m_stringBuilder.append("}\n");
>+void FunctionDefinitionWriter::visit(AST::IfStatement& ifStatement)
>+{
>+    checkErrorAndVisit(ifStatement.conditional());
>+    m_stringBuilder.append(makeString("if (", m_stack.takeLast(), ") {\n"));
>+    checkErrorAndVisit(ifStatement.body());
>+    if (ifStatement.elseBody()) {
>+        m_stringBuilder.append("} else {\n");
>+        checkErrorAndVisit(*ifStatement.elseBody());
>     }
>+    m_stringBuilder.append("}\n");
>+}
> 
>-    void visit(AST::Return& returnStatement) override
>-    {
>-        if (returnStatement.value()) {
>-            checkErrorAndVisit(*returnStatement.value());
>-            if (m_entryPointScaffolding) {
>-                auto variableName = generateNextVariableName();
>-                m_stringBuilder.append(m_entryPointScaffolding->pack(m_stack.takeLast(), variableName));
>-                m_stringBuilder.append(makeString("return ", variableName, ";\n"));
>-            } else
>-                m_stringBuilder.append(makeString("return ", m_stack.takeLast(), ";\n"));
>+void FunctionDefinitionWriter::visit(AST::Return& returnStatement)
>+{
>+    if (returnStatement.value()) {
>+        checkErrorAndVisit(*returnStatement.value());
>+        if (m_entryPointScaffolding) {
>+            auto variableName = generateNextVariableName();
>+            m_stringBuilder.append(m_entryPointScaffolding->pack(m_stack.takeLast(), variableName));
>+            m_stringBuilder.append(makeString("return ", variableName, ";\n"));
>         } else
>-            m_stringBuilder.append("return;\n");
>-    }
>+            m_stringBuilder.append(makeString("return ", m_stack.takeLast(), ";\n"));
>+    } else
>+        m_stringBuilder.append("return;\n");
>+}
> 
>-    void visit(AST::SwitchStatement& switchStatement) override
>-    {
>-        checkErrorAndVisit(switchStatement.value());
>+void FunctionDefinitionWriter::visit(AST::SwitchStatement& switchStatement)
>+{
>+    checkErrorAndVisit(switchStatement.value());
> 
>-        m_stringBuilder.append(makeString("switch (", m_stack.takeLast(), ") {"));
>-        for (auto& switchCase : switchStatement.switchCases())
>-            checkErrorAndVisit(switchCase);
>-        m_stringBuilder.append("}\n");
>-    }
>+    m_stringBuilder.append(makeString("switch (", m_stack.takeLast(), ") {"));
>+    for (auto& switchCase : switchStatement.switchCases())
>+        checkErrorAndVisit(switchCase);
>+    m_stringBuilder.append("}\n");
>+}
> 
>-    void visit(AST::SwitchCase& switchCase) override
>-    {
>-        if (switchCase.value())
>-            m_stringBuilder.append(makeString("case ", constantExpressionString(*switchCase.value()), ":\n"));
>-        else
>-            m_stringBuilder.append("default:\n");
>-        checkErrorAndVisit(switchCase.block());
>-        // FIXME: Figure out whether we need to break or fallthrough.
>-        CRASH();
>-    }
>+void FunctionDefinitionWriter::visit(AST::SwitchCase& switchCase)
>+{
>+    if (switchCase.value())
>+        m_stringBuilder.append(makeString("case ", constantExpressionString(*switchCase.value()), ":\n"));
>+    else
>+        m_stringBuilder.append("default:\n");
>+    checkErrorAndVisit(switchCase.block());
>+    // FIXME: Figure out whether we need to break or fallthrough.
>+    CRASH();
>+}
> 
>-    void visit(AST::Trap&) override
>-    {
>-        // FIXME: Implement this
>-        CRASH();
>-    }
>+void FunctionDefinitionWriter::visit(AST::Trap&)
>+{
>+    // FIXME: Implement this
>+    CRASH();
>+}
> 
>-    void visit(AST::VariableDeclarationsStatement& variableDeclarationsStatement) override
>-    {
>-        Visitor::visit(variableDeclarationsStatement);
>-    }
>+void FunctionDefinitionWriter::visit(AST::VariableDeclarationsStatement& variableDeclarationsStatement)
>+{
>+    Visitor::visit(variableDeclarationsStatement);
>+}
> 
>-    void visit(AST::WhileLoop& whileLoop) override
>-    {
>-        m_stringBuilder.append(makeString("while (true) {\n"));
>-        checkErrorAndVisit(whileLoop.conditional());
>-        m_stringBuilder.append(makeString("if (!", m_stack.takeLast(), ") break;\n"));
>-        checkErrorAndVisit(whileLoop.body());
>-        m_stringBuilder.append("}\n");
>-    }
>+void FunctionDefinitionWriter::visit(AST::WhileLoop& whileLoop)
>+{
>+    m_stringBuilder.append(makeString("while (true) {\n"));
>+    checkErrorAndVisit(whileLoop.conditional());
>+    m_stringBuilder.append(makeString("if (!", m_stack.takeLast(), ") break;\n"));
>+    checkErrorAndVisit(whileLoop.body());
>+    m_stringBuilder.append("}\n");
>+}
> 
>-    void visit(AST::IntegerLiteral& integerLiteral) override
>-    {
>-        ASSERT(integerLiteral.resolvedType());
>-        auto variableName = generateNextVariableName();
>-        auto mangledTypeName = m_typeNamer.mangledNameForType(*integerLiteral.resolvedType());
>-        m_stringBuilder.append(makeString(mangledTypeName, ' ', variableName, " = static_cast<", mangledTypeName, ">(", integerLiteral.value(), ");\n"));
>-        m_stack.append(variableName);
>-    }
>+void FunctionDefinitionWriter::visit(AST::IntegerLiteral& integerLiteral)
>+{
>+    ASSERT(integerLiteral.resolvedType());
>+    auto variableName = generateNextVariableName();
>+    auto mangledTypeName = m_typeNamer.mangledNameForType(*integerLiteral.resolvedType());
>+    m_stringBuilder.append(makeString(mangledTypeName, ' ', variableName, " = static_cast<", mangledTypeName, ">(", integerLiteral.value(), ");\n"));
>+    m_stack.append(variableName);
>+}
> 
>-    void visit(AST::UnsignedIntegerLiteral& unsignedIntegerLiteral) override
>-    {
>-        ASSERT(unsignedIntegerLiteral.resolvedType());
>-        auto variableName = generateNextVariableName();
>-        auto mangledTypeName = m_typeNamer.mangledNameForType(*unsignedIntegerLiteral.resolvedType());
>-        m_stringBuilder.append(makeString(mangledTypeName, ' ', variableName, " = static_cast<", mangledTypeName, ">(", unsignedIntegerLiteral.value(), ");\n"));
>-        m_stack.append(variableName);
>-    }
>+void FunctionDefinitionWriter::visit(AST::UnsignedIntegerLiteral& unsignedIntegerLiteral)
>+{
>+    ASSERT(unsignedIntegerLiteral.resolvedType());
>+    auto variableName = generateNextVariableName();
>+    auto mangledTypeName = m_typeNamer.mangledNameForType(*unsignedIntegerLiteral.resolvedType());
>+    m_stringBuilder.append(makeString(mangledTypeName, ' ', variableName, " = static_cast<", mangledTypeName, ">(", unsignedIntegerLiteral.value(), ");\n"));
>+    m_stack.append(variableName);
>+}
> 
>-    void visit(AST::FloatLiteral& floatLiteral) override
>-    {
>-        ASSERT(floatLiteral.resolvedType());
>-        auto variableName = generateNextVariableName();
>-        auto mangledTypeName = m_typeNamer.mangledNameForType(*floatLiteral.resolvedType());
>-        m_stringBuilder.append(makeString(mangledTypeName, ' ', variableName, " = static_cast<", mangledTypeName, ">(", floatLiteral.value(), ");\n"));
>-        m_stack.append(variableName);
>-    }
>+void FunctionDefinitionWriter::visit(AST::FloatLiteral& floatLiteral)
>+{
>+    ASSERT(floatLiteral.resolvedType());
>+    auto variableName = generateNextVariableName();
>+    auto mangledTypeName = m_typeNamer.mangledNameForType(*floatLiteral.resolvedType());
>+    m_stringBuilder.append(makeString(mangledTypeName, ' ', variableName, " = static_cast<", mangledTypeName, ">(", floatLiteral.value(), ");\n"));
>+    m_stack.append(variableName);
>+}
> 
>-    void visit(AST::NullLiteral& nullLiteral) override
>-    {
>-        ASSERT(nullLiteral.resolvedType());
>-        auto& unifyNode = nullLiteral.resolvedType()->unifyNode();
>-        ASSERT(is<AST::UnnamedType>(unifyNode));
>-        auto& unnamedType = downcast<AST::UnnamedType>(unifyNode);
>-        bool isArrayReferenceType = is<AST::ArrayReferenceType>(unnamedType);
>-
>-        auto variableName = generateNextVariableName();
>-        m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*nullLiteral.resolvedType()), ' ', variableName, " = "));
>-        if (isArrayReferenceType)
>-            m_stringBuilder.append("{ nullptr, 0 }");
>-        else
>-            m_stringBuilder.append("nullptr");
>-        m_stringBuilder.append(";\n");
>-        m_stack.append(variableName);
>-    }
>+void FunctionDefinitionWriter::visit(AST::NullLiteral& nullLiteral)
>+{
>+    ASSERT(nullLiteral.resolvedType());
>+    auto& unifyNode = nullLiteral.resolvedType()->unifyNode();
>+    ASSERT(is<AST::UnnamedType>(unifyNode));
>+    auto& unnamedType = downcast<AST::UnnamedType>(unifyNode);
>+    bool isArrayReferenceType = is<AST::ArrayReferenceType>(unnamedType);
>+
>+    auto variableName = generateNextVariableName();
>+    m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*nullLiteral.resolvedType()), ' ', variableName, " = "));
>+    if (isArrayReferenceType)
>+        m_stringBuilder.append("{ nullptr, 0 }");
>+    else
>+        m_stringBuilder.append("nullptr");
>+    m_stringBuilder.append(";\n");
>+    m_stack.append(variableName);
>+}
> 
>-    void visit(AST::BooleanLiteral& booleanLiteral) override
>-    {
>-        ASSERT(booleanLiteral.resolvedType());
>-        auto variableName = generateNextVariableName();
>-        auto mangledTypeName = m_typeNamer.mangledNameForType(*booleanLiteral.resolvedType());
>-        m_stringBuilder.append(makeString(mangledTypeName, ' ', variableName, " = static_cast<", mangledTypeName, ">(", booleanLiteral.value() ? "true" : "false", ");\n"));
>-        m_stack.append(variableName);
>-    }
>+void FunctionDefinitionWriter::visit(AST::BooleanLiteral& booleanLiteral)
>+{
>+    ASSERT(booleanLiteral.resolvedType());
>+    auto variableName = generateNextVariableName();
>+    auto mangledTypeName = m_typeNamer.mangledNameForType(*booleanLiteral.resolvedType());
>+    m_stringBuilder.append(makeString(mangledTypeName, ' ', variableName, " = static_cast<", mangledTypeName, ">(", booleanLiteral.value() ? "true" : "false", ");\n"));
>+    m_stack.append(variableName);
>+}
> 
>-    void visit(AST::EnumerationMemberLiteral& enumerationMemberLiteral) override
>-    {
>-        ASSERT(enumerationMemberLiteral.resolvedType());
>-        ASSERT(enumerationMemberLiteral.enumerationDefinition());
>-        ASSERT(enumerationMemberLiteral.enumerationDefinition());
>-        auto variableName = generateNextVariableName();
>-        auto mangledTypeName = m_typeNamer.mangledNameForType(*enumerationMemberLiteral.resolvedType());
>-        m_stringBuilder.append(makeString(mangledTypeName, ' ', variableName, " = ", mangledTypeName, '.', m_typeNamer.mangledNameForEnumerationMember(*enumerationMemberLiteral.enumerationMember()), ";\n"));
>-        m_stack.append(variableName);
>-    }
>+void FunctionDefinitionWriter::visit(AST::EnumerationMemberLiteral& enumerationMemberLiteral)
>+{
>+    ASSERT(enumerationMemberLiteral.resolvedType());
>+    ASSERT(enumerationMemberLiteral.enumerationDefinition());
>+    ASSERT(enumerationMemberLiteral.enumerationDefinition());
>+    auto variableName = generateNextVariableName();
>+    auto mangledTypeName = m_typeNamer.mangledNameForType(*enumerationMemberLiteral.resolvedType());
>+    m_stringBuilder.append(makeString(mangledTypeName, ' ', variableName, " = ", mangledTypeName, '.', m_typeNamer.mangledNameForEnumerationMember(*enumerationMemberLiteral.enumerationMember()), ";\n"));
>+    m_stack.append(variableName);
>+}
> 
>-    void visit(AST::Expression& expression) override
>-    {
>-        Visitor::visit(expression);
>-    }
>+void FunctionDefinitionWriter::visit(AST::Expression& expression)
>+{
>+    Visitor::visit(expression);
>+}
> 
>-    void visit(AST::DotExpression&) override
>-    {
>-        // This should be lowered already.
>-        ASSERT_NOT_REACHED();
>-    }
>+void FunctionDefinitionWriter::visit(AST::DotExpression&)
>+{
>+    // This should be lowered already.
>+    ASSERT_NOT_REACHED();
>+}
> 
>-    void visit(AST::IndexExpression&) override
>-    {
>-        // This should be lowered already.
>-        ASSERT_NOT_REACHED();
>-    }
>+void FunctionDefinitionWriter::visit(AST::IndexExpression&)
>+{
>+    // This should be lowered already.
>+    ASSERT_NOT_REACHED();
>+}
> 
>-    void visit(AST::PropertyAccessExpression&) override
>-    {
>-        ASSERT_NOT_REACHED();
>-    }
>+void FunctionDefinitionWriter::visit(AST::PropertyAccessExpression&)
>+{
>+    ASSERT_NOT_REACHED();
>+}
> 
>-    void visit(AST::VariableDeclaration& variableDeclaration) override
>-    {
>-        ASSERT(variableDeclaration.type());
>-        if (variableDeclaration.initializer())
>-            checkErrorAndVisit(*variableDeclaration.initializer());
>-        else {
>-            // FIXME: Zero-fill the variable.
>-            CRASH();
>-        }
>-        // FIXME: Implement qualifiers.
>-        auto variableName = generateNextVariableName();
>-        auto addResult = m_variableMapping.add(&variableDeclaration, variableName);
>-        ASSERT_UNUSED(addResult, addResult.isNewEntry);
>-        m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*variableDeclaration.type()), ' ', variableName, " = ", m_stack.takeLast(), ";\n"));
>-        m_stack.append(variableName);
>+void FunctionDefinitionWriter::visit(AST::VariableDeclaration& variableDeclaration)
>+{
>+    ASSERT(variableDeclaration.type());
>+    if (variableDeclaration.initializer())
>+        checkErrorAndVisit(*variableDeclaration.initializer());
>+    else {
>+        // FIXME: Zero-fill the variable.
>+        CRASH();
>     }
>+    // FIXME: Implement qualifiers.
>+    auto variableName = generateNextVariableName();
>+    auto addResult = m_variableMapping.add(&variableDeclaration, variableName);
>+    ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+    m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*variableDeclaration.type()), ' ', variableName, " = ", m_stack.takeLast(), ";\n"));
>+    m_stack.append(variableName);
>+}
> 
>-    void visit(AST::AssignmentExpression& assignmentExpression) override
>-    {
>-        checkErrorAndVisit(assignmentExpression.left());
>-        auto leftName = m_stack.takeLast();
>-        checkErrorAndVisit(assignmentExpression.right());
>-        auto rightName = m_stack.takeLast();
>-        m_stringBuilder.append(makeString(leftName, " = ", rightName, ";\n"));
>-    }
>+void FunctionDefinitionWriter::visit(AST::AssignmentExpression& assignmentExpression)
>+{
>+    checkErrorAndVisit(assignmentExpression.left());
>+    auto leftName = m_stack.takeLast();
>+    checkErrorAndVisit(assignmentExpression.right());
>+    auto rightName = m_stack.takeLast();
>+    m_stringBuilder.append(makeString(leftName, " = ", rightName, ";\n"));
>+}
> 
>-    void visit(AST::CallExpression& callExpression) override
>-    {
>-        Vector<String> argumentNames;
>-        for (auto& argument : callExpression.arguments()) {
>-            checkErrorAndVisit(argument);
>-            argumentNames.append(m_stack.takeLast());
>-        }
>-        ASSERT(callExpression.resolvedType());
>-        ASSERT(callExpression.function());
>-        auto iterator = m_functionMapping.find(callExpression.function());
>-        ASSERT(iterator != m_functionMapping.end());
>-        auto variableName = generateNextVariableName();
>-        m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*callExpression.resolvedType()), ' ', variableName, " = ", iterator->value, '('));
>-        for (size_t i = 0; i < argumentNames.size(); ++i) {
>-            if (i)
>-                m_stringBuilder.append(", ");
>-            m_stringBuilder.append(argumentNames[i]);
>-        }
>-        m_stringBuilder.append(");\n");
>-        m_stack.append(variableName);
>-    }
>+void FunctionDefinitionWriter::visit(AST::CallExpression& callExpression)
>+{
>+    Vector<String> argumentNames;
>+    for (auto& argument : callExpression.arguments()) {
>+        checkErrorAndVisit(argument);
>+        argumentNames.append(m_stack.takeLast());
>+    }
>+    ASSERT(callExpression.resolvedType());
>+    ASSERT(callExpression.function());
>+    auto iterator = m_functionMapping.find(callExpression.function());
>+    ASSERT(iterator != m_functionMapping.end());
>+    auto variableName = generateNextVariableName();
>+    m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*callExpression.resolvedType()), ' ', variableName, " = ", iterator->value, '('));
>+    for (size_t i = 0; i < argumentNames.size(); ++i) {
>+        if (i)
>+            m_stringBuilder.append(", ");
>+        m_stringBuilder.append(argumentNames[i]);
>+    }
>+    m_stringBuilder.append(");\n");
>+    m_stack.append(variableName);
>+}
> 
>-    void visit(AST::CommaExpression& commaExpression) override
>-    {
>-        String result;
>-        for (auto& expression : commaExpression.list()) {
>-            checkErrorAndVisit(expression);
>-            result = m_stack.takeLast();
>-        }
>-        m_stack.append(result);
>+void FunctionDefinitionWriter::visit(AST::CommaExpression& commaExpression)
>+{
>+    String result;
>+    for (auto& expression : commaExpression.list()) {
>+        checkErrorAndVisit(expression);
>+        result = m_stack.takeLast();
>     }
>+    m_stack.append(result);
>+}
> 
>-    void visit(AST::DereferenceExpression& dereferenceExpression) override
>-    {
>-        checkErrorAndVisit(dereferenceExpression.pointer());
>-        auto right = m_stack.takeLast();
>-        ASSERT(dereferenceExpression.resolvedType());
>-        auto variableName = generateNextVariableName();
>-        m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*dereferenceExpression.resolvedType()), ' ', variableName, " = *", right, ";\n"));
>-        m_stack.append(variableName);
>-    }
>+void FunctionDefinitionWriter::visit(AST::DereferenceExpression& dereferenceExpression)
>+{
>+    checkErrorAndVisit(dereferenceExpression.pointer());
>+    auto right = m_stack.takeLast();
>+    ASSERT(dereferenceExpression.resolvedType());
>+    auto variableName = generateNextVariableName();
>+    m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*dereferenceExpression.resolvedType()), ' ', variableName, " = *", right, ";\n"));
>+    m_stack.append(variableName);
>+}
> 
>-    void visit(AST::LogicalExpression& logicalExpression) override
>-    {
>-        checkErrorAndVisit(logicalExpression.left());
>-        auto left = m_stack.takeLast();
>-        checkErrorAndVisit(logicalExpression.right());
>-        auto right = m_stack.takeLast();
>-        ASSERT(logicalExpression.resolvedType());
>-        auto variableName = generateNextVariableName();
>-        m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*logicalExpression.resolvedType()), ' ', variableName, " = ", left));
>-        switch (logicalExpression.type()) {
>-        case AST::LogicalExpression::Type::And:
>-            m_stringBuilder.append(" && ");
>-            break;
>-        default:
>-            ASSERT(logicalExpression.type() == AST::LogicalExpression::Type::Or);
>-            m_stringBuilder.append(" || ");
>-            break;
>-        }
>-        m_stringBuilder.append(makeString(right, ";\n"));
>-        m_stack.append(variableName);
>-    }
>+void FunctionDefinitionWriter::visit(AST::LogicalExpression& logicalExpression)
>+{
>+    checkErrorAndVisit(logicalExpression.left());
>+    auto left = m_stack.takeLast();
>+    checkErrorAndVisit(logicalExpression.right());
>+    auto right = m_stack.takeLast();
>+    ASSERT(logicalExpression.resolvedType());
>+    auto variableName = generateNextVariableName();
>+    m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*logicalExpression.resolvedType()), ' ', variableName, " = ", left));
>+    switch (logicalExpression.type()) {
>+    case AST::LogicalExpression::Type::And:
>+        m_stringBuilder.append(" && ");
>+        break;
>+    default:
>+        ASSERT(logicalExpression.type() == AST::LogicalExpression::Type::Or);
>+        m_stringBuilder.append(" || ");
>+        break;
>+    }
>+    m_stringBuilder.append(makeString(right, ";\n"));
>+    m_stack.append(variableName);
>+}
> 
>-    void visit(AST::LogicalNotExpression& logicalNotExpression) override
>-    {
>-        checkErrorAndVisit(logicalNotExpression.operand());
>-        auto operand = m_stack.takeLast();
>-        ASSERT(logicalNotExpression.resolvedType());
>-        auto variableName = generateNextVariableName();
>-        m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*logicalNotExpression.resolvedType()), ' ', variableName, " = !", operand, ";\n"));
>-        m_stack.append(variableName);
>-    }
>+void FunctionDefinitionWriter::visit(AST::LogicalNotExpression& logicalNotExpression)
>+{
>+    checkErrorAndVisit(logicalNotExpression.operand());
>+    auto operand = m_stack.takeLast();
>+    ASSERT(logicalNotExpression.resolvedType());
>+    auto variableName = generateNextVariableName();
>+    m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*logicalNotExpression.resolvedType()), ' ', variableName, " = !", operand, ";\n"));
>+    m_stack.append(variableName);
>+}
> 
>-    void visit(AST::MakeArrayReferenceExpression& makeArrayReferenceExpression) override
>+void FunctionDefinitionWriter::visit(AST::MakeArrayReferenceExpression& makeArrayReferenceExpression)
>+{
>+    checkErrorAndVisit(makeArrayReferenceExpression.lValue());
>+    auto lValue = m_stack.takeLast();
>+    ASSERT(makeArrayReferenceExpression.resolvedType());
>+    auto variableName = generateNextVariableName();
>+    auto mangledTypeName = m_typeNamer.mangledNameForType(*makeArrayReferenceExpression.resolvedType());
>+    if (is<AST::PointerType>(*makeArrayReferenceExpression.resolvedType()))
>+        m_stringBuilder.append(makeString(mangledTypeName, ' ', variableName, " = { ", lValue, ", 1 };\n"));
>+    else if (is<AST::ArrayType>(*makeArrayReferenceExpression.resolvedType())) {
>+        auto& arrayType = downcast<AST::ArrayType>(*makeArrayReferenceExpression.resolvedType());
>+        m_stringBuilder.append(makeString(mangledTypeName, ' ', variableName, " = { &(", lValue, "[0]), ", arrayType.numElements(), " };\n"));
>+    } else
>+        m_stringBuilder.append(makeString(mangledTypeName, ' ', variableName, " = { &", lValue, ", 1 };\n"));
>+    m_stack.append(variableName);
>+}
>+
>+void FunctionDefinitionWriter::visit(AST::MakePointerExpression& makePointerExpression)
>+{
>+    checkErrorAndVisit(makePointerExpression.lValue());
>+    auto lValue = m_stack.takeLast();
>+    ASSERT(makePointerExpression.resolvedType());
>+    auto variableName = generateNextVariableName();
>+    m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*makePointerExpression.resolvedType()), ' ', variableName, " = &", lValue, ";\n"));
>+    m_stack.append(variableName);
>+}
>+
>+void FunctionDefinitionWriter::visit(AST::ReadModifyWriteExpression&)
>+{
>+    // This should be lowered already.
>+    ASSERT_NOT_REACHED();
>+}
>+
>+void FunctionDefinitionWriter::visit(AST::TernaryExpression& ternaryExpression)
>+{
>+    checkErrorAndVisit(ternaryExpression.predicate());
>+    auto check = m_stack.takeLast();
>+
>+    ASSERT(ternaryExpression.resolvedType());
>+    auto variableName = generateNextVariableName();
>+    m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*ternaryExpression.resolvedType()), ' ', variableName, ";\n"));
>+
>+    m_stringBuilder.append(makeString("if (", check, ") {\n"));
>+    checkErrorAndVisit(ternaryExpression.bodyExpression());
>+    m_stringBuilder.append(makeString(variableName, " = ", m_stack.takeLast(), ";\n"));
>+    m_stringBuilder.append("} else {\n");
>+    checkErrorAndVisit(ternaryExpression.elseExpression());
>+    m_stringBuilder.append(makeString(variableName, " = ", m_stack.takeLast(), ";\n"));
>+    m_stringBuilder.append("}\n");
>+    m_stack.append(variableName);
>+}
>+
>+void FunctionDefinitionWriter::visit(AST::VariableReference& variableReference)
>+{
>+    ASSERT(variableReference.variable());
>+    auto iterator = m_variableMapping.find(variableReference.variable());
>+    ASSERT(iterator != m_variableMapping.end());
>+    m_stack.append(iterator->value);
>+}
>+
>+String FunctionDefinitionWriter::constantExpressionString(AST::ConstantExpression& constantExpression)
>+{
>+    String result;
>+    constantExpression.visit(WTF::makeVisitor([&](AST::IntegerLiteral& integerLiteral) {
>+        result = makeString("", integerLiteral.value());
>+    }, [&](AST::UnsignedIntegerLiteral& unsignedIntegerLiteral) {
>+        result = makeString("", unsignedIntegerLiteral.value());
>+    }, [&](AST::FloatLiteral& floatLiteral) {
>+        result = makeString("", floatLiteral.value());
>+    }, [&](AST::NullLiteral&) {
>+        result = "nullptr"_str;
>+    }, [&](AST::BooleanLiteral& booleanLiteral) {
>+        result = booleanLiteral.value() ? "true"_str : "false"_str;
>+    }, [&](AST::EnumerationMemberLiteral& enumerationMemberLiteral) {
>+        ASSERT(enumerationMemberLiteral.enumerationDefinition());
>+        ASSERT(enumerationMemberLiteral.enumerationDefinition());
>+        result = makeString(m_typeNamer.mangledNameForType(*enumerationMemberLiteral.enumerationDefinition()), '.', m_typeNamer.mangledNameForEnumerationMember(*enumerationMemberLiteral.enumerationMember()));
>+    }));
>+    return result;
>+}
>+
>+class RenderFunctionDefinitionWriter : public FunctionDefinitionWriter {
>+public:
>+    RenderFunctionDefinitionWriter(Intrinsics& intrinsics, TypeNamer& typeNamer, HashMap<AST::FunctionDeclaration*, String>& functionMapping, MatchedRenderSemantics&& matchedSemantics, Layout& layout)
>+        : FunctionDefinitionWriter(intrinsics, typeNamer, functionMapping, layout)
>+        , m_matchedSemantics(WTFMove(matchedSemantics))
>     {
>-        checkErrorAndVisit(makeArrayReferenceExpression.lValue());
>-        auto lValue = m_stack.takeLast();
>-        ASSERT(makeArrayReferenceExpression.resolvedType());
>-        auto variableName = generateNextVariableName();
>-        auto mangledTypeName = m_typeNamer.mangledNameForType(*makeArrayReferenceExpression.resolvedType());
>-        if (is<AST::PointerType>(*makeArrayReferenceExpression.resolvedType()))
>-            m_stringBuilder.append(makeString(mangledTypeName, ' ', variableName, " = { ", lValue, ", 1 };\n"));
>-        else if (is<AST::ArrayType>(*makeArrayReferenceExpression.resolvedType())) {
>-            auto& arrayType = downcast<AST::ArrayType>(*makeArrayReferenceExpression.resolvedType());
>-            m_stringBuilder.append(makeString(mangledTypeName, ' ', variableName, " = { &(", lValue, "[0]), ", arrayType.numElements(), " };\n"));
>-        } else
>-            m_stringBuilder.append(makeString(mangledTypeName, ' ', variableName, " = { &", lValue, ", 1 };\n"));
>-        m_stack.append(variableName);
>     }
> 
>-    void visit(AST::MakePointerExpression& makePointerExpression) override
>+    MappedBindGroups&& takeVertexMappedBindGroups()
>     {
>-        checkErrorAndVisit(makePointerExpression.lValue());
>-        auto lValue = m_stack.takeLast();
>-        ASSERT(makePointerExpression.resolvedType());
>-        auto variableName = generateNextVariableName();
>-        m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*makePointerExpression.resolvedType()), ' ', variableName, " = &", lValue, ";\n"));
>-        m_stack.append(variableName);
>+        ASSERT(m_vertexMappedBindGroups);
>+        return WTFMove(*m_vertexMappedBindGroups);
>     }
> 
>-    void visit(AST::ReadModifyWriteExpression&) override
>+    MappedBindGroups&& takeFragmentMappedBindGroups()
>     {
>-        // This should be lowered already.
>-        ASSERT_NOT_REACHED();
>+        ASSERT(m_fragmentMappedBindGroups);
>+        return WTFMove(*m_fragmentMappedBindGroups);
>     }
> 
>-    void visit(AST::TernaryExpression& ternaryExpression) override
>-    {
>-        checkErrorAndVisit(ternaryExpression.predicate());
>-        auto check = m_stack.takeLast();
>+private:
>+    std::unique_ptr<EntryPointScaffolding> createEntryPointScaffolding(AST::FunctionDefinition&) override;
> 
>-        ASSERT(ternaryExpression.resolvedType());
>-        auto variableName = generateNextVariableName();
>-        m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*ternaryExpression.resolvedType()), ' ', variableName, ";\n"));
>+    MatchedRenderSemantics m_matchedSemantics;
>+    Optional<MappedBindGroups> m_vertexMappedBindGroups;
>+    Optional<MappedBindGroups> m_fragmentMappedBindGroups;
>+};
> 
>-        m_stringBuilder.append(makeString("if (", check, ") {\n"));
>-        checkErrorAndVisit(ternaryExpression.bodyExpression());
>-        m_stringBuilder.append(makeString(variableName, " = ", m_stack.takeLast(), ";\n"));
>-        m_stringBuilder.append("} else {\n");
>-        checkErrorAndVisit(ternaryExpression.elseExpression());
>-        m_stringBuilder.append(makeString(variableName, " = ", m_stack.takeLast(), ";\n"));
>-        m_stringBuilder.append("}\n");
>-        m_stack.append(variableName);
>+std::unique_ptr<EntryPointScaffolding> RenderFunctionDefinitionWriter::createEntryPointScaffolding(AST::FunctionDefinition& functionDefinition)
>+{
>+    auto generateNextVariableName = [this]() -> String {
>+        return this->generateNextVariableName();
>+    };
>+    if (&functionDefinition == m_matchedSemantics.vertexShader) {
>+        auto result = std::make_unique<VertexEntryPointScaffolding>(functionDefinition, m_intrinsics, m_typeNamer, m_matchedSemantics.vertexShaderEntryPointItems, m_matchedSemantics.vertexShaderResourceMap, m_layout, WTFMove(generateNextVariableName), m_matchedSemantics.matchedVertexAttributes);
>+        ASSERT(!m_vertexMappedBindGroups);
>+        m_vertexMappedBindGroups = result->mappedBindGroups();
>+        return result;
>     }
>-
>-    void visit(AST::VariableReference& variableReference) override
>-    {
>-        ASSERT(variableReference.variable());
>-        auto iterator = m_variableMapping.find(variableReference.variable());
>-        ASSERT(iterator != m_variableMapping.end());
>-        m_stack.append(iterator->value);
>+    if (&functionDefinition == m_matchedSemantics.fragmentShader) {
>+        auto result = std::make_unique<FragmentEntryPointScaffolding>(functionDefinition, m_intrinsics, m_typeNamer, m_matchedSemantics.fragmentShaderEntryPointItems, m_matchedSemantics.fragmentShaderResourceMap, m_layout, WTFMove(generateNextVariableName), m_matchedSemantics.matchedColorAttachments);
>+        ASSERT(!m_fragmentMappedBindGroups);
>+        m_fragmentMappedBindGroups = result->mappedBindGroups();
>+        return result;
>     }
>+    return nullptr;
>+}
> 
>-    String constantExpressionString(AST::ConstantExpression& constantExpression)
>+class ComputeFunctionDefinitionWriter : public FunctionDefinitionWriter {
>+public:
>+    ComputeFunctionDefinitionWriter(Intrinsics& intrinsics, TypeNamer& typeNamer, HashMap<AST::FunctionDeclaration*, String>& functionMapping, MatchedComputeSemantics&& matchedSemantics, Layout& layout)
>+        : FunctionDefinitionWriter(intrinsics, typeNamer, functionMapping, layout)
>+        , m_matchedSemantics(WTFMove(matchedSemantics))
>     {
>-        String result;
>-        constantExpression.visit(WTF::makeVisitor([&](AST::IntegerLiteral& integerLiteral) {
>-            result = makeString("", integerLiteral.value());
>-        }, [&](AST::UnsignedIntegerLiteral& unsignedIntegerLiteral) {
>-            result = makeString("", unsignedIntegerLiteral.value());
>-        }, [&](AST::FloatLiteral& floatLiteral) {
>-            result = makeString("", floatLiteral.value());
>-        }, [&](AST::NullLiteral&) {
>-            result = "nullptr"_str;
>-        }, [&](AST::BooleanLiteral& booleanLiteral) {
>-            result = booleanLiteral.value() ? "true"_str : "false"_str;
>-        }, [&](AST::EnumerationMemberLiteral& enumerationMemberLiteral) {
>-            ASSERT(enumerationMemberLiteral.enumerationDefinition());
>-            ASSERT(enumerationMemberLiteral.enumerationDefinition());
>-            result = makeString(m_typeNamer.mangledNameForType(*enumerationMemberLiteral.enumerationDefinition()), '.', m_typeNamer.mangledNameForEnumerationMember(*enumerationMemberLiteral.enumerationMember()));
>-        }));
>-        return result;
>     }
> 
>-    String generateNextVariableName()
>+    MappedBindGroups&& takeMappedBindGroups()
>     {
>-        return makeString("variable", m_variableCount++);
>+        ASSERT(m_mappedBindGroups);
>+        return WTFMove(*m_mappedBindGroups);
>     }
> 
> private:
>-    Intrinsics& m_intrinsics;
>-    TypeNamer& m_typeNamer;
>-    HashMap<AST::FunctionDeclaration*, String>& m_functionMapping;
>-    HashMap<AST::VariableDeclaration*, String> m_variableMapping;
>-    StringBuilder m_stringBuilder;
>-    Vector<String> m_stack;
>-    Optional<EntryPointScaffolding> m_entryPointScaffolding;
>-    unsigned m_variableCount { 0 };
>+    std::unique_ptr<EntryPointScaffolding> createEntryPointScaffolding(AST::FunctionDefinition&) override;
>+
>+    MatchedComputeSemantics m_matchedSemantics;
>+    Optional<MappedBindGroups> m_mappedBindGroups;
> };
> 
>-String metalFunctions(Program& program, TypeNamer& typeNamer)
>+std::unique_ptr<EntryPointScaffolding> ComputeFunctionDefinitionWriter::createEntryPointScaffolding(AST::FunctionDefinition& functionDefinition)
>+{
>+    auto generateNextVariableName = [this]() -> String {
>+        return this->generateNextVariableName();
>+    };
>+    if (&functionDefinition == m_matchedSemantics.shader) {
>+        auto result = std::make_unique<ComputeEntryPointScaffolding>(functionDefinition, m_intrinsics, m_typeNamer, m_matchedSemantics.entryPointItems, m_matchedSemantics.resourceMap, m_layout, WTFMove(generateNextVariableName));
>+        ASSERT(!m_mappedBindGroups);
>+        m_mappedBindGroups = result->mappedBindGroups();
>+        return result;
>+    }
>+    return nullptr;
>+}
>+
>+struct SharedMetalFunctionsResult {
>+    HashMap<AST::FunctionDeclaration*, String> functionMapping;
>+    String metalFunctions;
>+};
>+static SharedMetalFunctionsResult sharedMetalFunctions(Program& program, TypeNamer& typeNamer)
> {
>     StringBuilder stringBuilder;
> 
>@@ -610,17 +756,48 @@ String metalFunctions(Program& program, TypeNamer& typeNamer)
>     }
> 
>     stringBuilder.append('\n');
>+    return { WTFMove(functionMapping), stringBuilder.toString() };
>+}
> 
>-    {
>-        FunctionDefinitionWriter functionDefinitionWriter(program.intrinsics(), typeNamer, functionMapping);
>-        for (auto& nativeFunctionDeclaration : program.nativeFunctionDeclarations())
>-            functionDefinitionWriter.visit(nativeFunctionDeclaration);
>-        for (auto& functionDefinition : program.functionDefinitions())
>-            functionDefinitionWriter.visit(functionDefinition);
>-        stringBuilder.append(functionDefinitionWriter.toString());
>-    }
>+RenderMetalFunctions metalFunctions(Program& program, TypeNamer& typeNamer, MatchedRenderSemantics&& matchedSemantics, Layout& layout)
>+{
>+    auto sharedMetalFunctions = Metal::sharedMetalFunctions(program, typeNamer);
>+
>+    StringBuilder stringBuilder;
>+    stringBuilder.append(sharedMetalFunctions.metalFunctions);
>+
>+    RenderFunctionDefinitionWriter functionDefinitionWriter(program.intrinsics(), typeNamer, sharedMetalFunctions.functionMapping, WTFMove(matchedSemantics), layout);
>+    for (auto& nativeFunctionDeclaration : program.nativeFunctionDeclarations())
>+        functionDefinitionWriter.visit(nativeFunctionDeclaration);
>+    for (auto& functionDefinition : program.functionDefinitions())
>+        functionDefinitionWriter.visit(functionDefinition);
>+    stringBuilder.append(functionDefinitionWriter.toString());
>+
>+    RenderMetalFunctions result;
>+    result.metalSource = stringBuilder.toString();
>+    result.vertexMappedBindGroups = functionDefinitionWriter.takeVertexMappedBindGroups();
>+    result.fragmentMappedBindGroups = functionDefinitionWriter.takeFragmentMappedBindGroups();
>+    return result;
>+}
> 
>-    return stringBuilder.toString();
>+ComputeMetalFunctions metalFunctions(Program& program, TypeNamer& typeNamer, MatchedComputeSemantics&& matchedSemantics, Layout& layout)
>+{
>+    auto sharedMetalFunctions = Metal::sharedMetalFunctions(program, typeNamer);
>+
>+    StringBuilder stringBuilder;
>+    stringBuilder.append(sharedMetalFunctions.metalFunctions);
>+
>+    ComputeFunctionDefinitionWriter functionDefinitionWriter(program.intrinsics(), typeNamer, sharedMetalFunctions.functionMapping, WTFMove(matchedSemantics), layout);
>+    for (auto& nativeFunctionDeclaration : program.nativeFunctionDeclarations())
>+        functionDefinitionWriter.visit(nativeFunctionDeclaration);
>+    for (auto& functionDefinition : program.functionDefinitions())
>+        functionDefinitionWriter.visit(functionDefinition);
>+    stringBuilder.append(functionDefinitionWriter.toString());
>+
>+    ComputeMetalFunctions result;
>+    result.metalSource = stringBuilder.toString();
>+    result.mappedBindGroups = functionDefinitionWriter.takeMappedBindGroups();
>+    return result;
> }
> 
> } // namespace Metal
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.h b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.h
>index 9960e5e48d58c24e35e528382456472b7d09b6d7..062821fab89ec5f07bd7c7f48df2a20ac547acce 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.h
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.h
>@@ -27,6 +27,9 @@
> 
> #if ENABLE(WEBGPU)
> 
>+#include "WHLSLMappedBindings.h"
>+#include "WHLSLSemanticMatcher.h"
>+
> namespace WebCore {
> 
> namespace WHLSL {
>@@ -37,7 +40,18 @@ namespace Metal {
> 
> class TypeNamer;
> 
>-String metalFunctions(Program&, TypeNamer&);
>+struct RenderMetalFunctions {
>+    String metalSource;
>+    MappedBindGroups vertexMappedBindGroups;
>+    MappedBindGroups fragmentMappedBindGroups;
>+};
>+RenderMetalFunctions metalFunctions(Program&, TypeNamer&, MatchedRenderSemantics&&, Layout&);
>+
>+struct ComputeMetalFunctions {
>+    String metalSource;
>+    MappedBindGroups mappedBindGroups;
>+};
>+ComputeMetalFunctions metalFunctions(Program&, TypeNamer&, MatchedComputeSemantics&&, Layout&);
> 
> }
> 
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMappedBindings.h b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMappedBindings.h
>new file mode 100644
>index 0000000000000000000000000000000000000000..db3c52bf496398f3d040d15a8bd8b7e844e1d136
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMappedBindings.h
>@@ -0,0 +1,51 @@
>+/*
>+ * 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/Vector.h>
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+namespace Metal {
>+
>+struct MappedBindGroup {
>+    unsigned argumentBufferIndex;
>+    Vector<unsigned> bindingIndices;
>+};
>+
>+using MappedBindGroups = Vector<MappedBindGroup>; // Parallel to the input resource Layout.
>+
>+} // namespace Metal
>+
>+} // namespace WHLSL
>+
>+} // namespace WebCore
>+
>+#endif
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.cpp b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.cpp
>index 10c0530e290d7041767f71823270d3b92597b54b..88cfe73017d602f3411eeeadf57ba5e044f3cd71 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)
>+static String generateMetalCodeShared(String&& metalTypes, String&& metalFunctions)
> {
>     StringBuilder stringBuilder;
>     stringBuilder.append("#include <metal_stdlib>\n");
>@@ -51,12 +51,29 @@ String generateMetalCode(Program& program)
>     stringBuilder.append("using namespace metal;\n");
>     stringBuilder.append("\n");
> 
>-    TypeNamer typeNamer(program);
>-    stringBuilder.append(typeNamer.metalTypes());
>-    stringBuilder.append(metalFunctions(program, typeNamer));
>+    stringBuilder.append(WTFMove(metalTypes));
>+    stringBuilder.append(WTFMove(metalFunctions));
>     return stringBuilder.toString();
> }
> 
>+RenderMetalCode generateMetalCode(Program& program, MatchedRenderSemantics&& matchedSemantics, Layout& layout)
>+{
>+    TypeNamer typeNamer(program);
>+    auto metalTypes = typeNamer.metalTypes();
>+    auto metalFunctions = Metal::metalFunctions(program, typeNamer, WTFMove(matchedSemantics), layout);
>+    auto metalCode = generateMetalCodeShared(WTFMove(metalTypes), WTFMove(metalFunctions.metalSource));
>+    return { WTFMove(metalCode), WTFMove(metalFunctions.vertexMappedBindGroups), WTFMove(metalFunctions.fragmentMappedBindGroups) };
>+}
>+
>+ComputeMetalCode generateMetalCode(Program& program, MatchedComputeSemantics&& matchedSemantics, Layout& layout)
>+{
>+    TypeNamer typeNamer(program);
>+    auto metalTypes = typeNamer.metalTypes();
>+    auto metalFunctions = Metal::metalFunctions(program, typeNamer, WTFMove(matchedSemantics), layout);
>+    auto metalCode = generateMetalCodeShared(WTFMove(metalTypes), WTFMove(metalFunctions.metalSource));
>+    return { WTFMove(metalCode), WTFMove(metalFunctions.mappedBindGroups) };
>+}
>+
> }
> 
> }
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.h b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.h
>index e088e010449948411075674347491f73d36a234f..a69463df4af3512da06beaae152d48d694447813 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.h
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.h
>@@ -27,6 +27,10 @@
> 
> #if ENABLE(WEBGPU)
> 
>+#include "WHLSLMappedBindings.h"
>+#include "WHLSLPipelineDescriptor.h"
>+#include "WHLSLSemanticMatcher.h"
>+#include <wtf/Variant.h>
> #include <wtf/text/WTFString.h>
> 
> namespace WebCore {
>@@ -37,8 +41,18 @@ class Program;
> 
> namespace Metal {
> 
>-// FIXME: This needs to know about the pipeline state object to emit function prologues and epilogues.
>-String generateMetalCode(Program&);
>+struct RenderMetalCode {
>+    String metalSource;
>+    MappedBindGroups vertexMappedBindGroups;
>+    MappedBindGroups fragmentMappedBindGroups;
>+};
>+RenderMetalCode generateMetalCode(Program&, MatchedRenderSemantics&& matchedSemantics, Layout&);
>+
>+struct ComputeMetalCode {
>+    String metalSource;
>+    MappedBindGroups bindGroups;
>+};
>+ComputeMetalCode generateMetalCode(Program&, 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/WHLSLPrepare.cpp b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.cpp
>new file mode 100644
>index 0000000000000000000000000000000000000000..4ec2bd2815e71e8285d45c8dc6877611aabc1bb0
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.cpp
>@@ -0,0 +1,130 @@
>+/*
>+ * 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 "WHLSLPrepare.h"
>+
>+#if ENABLE(WEBGPU)
>+
>+#include "WHLSLCheckDuplicateFunctions.h"
>+#include "WHLSLChecker.h"
>+#include "WHLSLFunctionStageChecker.h"
>+#include "WHLSLHighZombieFinder.h"
>+#include "WHLSLLiteralTypeChecker.h"
>+#include "WHLSLMetalCodeGenerator.h"
>+#include "WHLSLNameResolver.h"
>+#include "WHLSLParser.h"
>+#include "WHLSLProgram.h"
>+#include "WHLSLRecursionChecker.h"
>+#include "WHLSLRecursiveTypeChecker.h"
>+#include "WHLSLSemanticMatcher.h"
>+#include "WHLSLStandardLibrary.h"
>+#include "WHLSLStatementBehaviorChecker.h"
>+#include "WHLSLSynthesizeArrayOperatorLength.h"
>+#include "WHLSLSynthesizeConstructors.h"
>+#include "WHLSLSynthesizeEnumerationFunctions.h"
>+#include "WHLSLSynthesizeStructureAccessors.h"
>+#include <wtf/Variant.h>
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+static Optional<Program> prepareShared(String& whlslSource)
>+{
>+    Program program;
>+    Parser parser;
>+    auto standardLibrary = String::fromUTF8(WHLSLStandardLibrary, sizeof(WHLSLStandardLibrary));
>+    auto failure = static_cast<bool>(parser.parse(program, standardLibrary, Parser::Mode::StandardLibrary));
>+    ASSERT_UNUSED(failure, !failure);
>+    if (parser.parse(program, whlslSource, Parser::Mode::User))
>+        return WTF::nullopt;
>+    NameResolver nameResolver(program.nameContext());
>+    if (!resolveNamesInTypes(program, nameResolver))
>+        return WTF::nullopt;
>+    if (!checkRecursiveTypes(program))
>+        return WTF::nullopt;
>+    synthesizeStructureAccessors(program);
>+    synthesizeEnumerationFunctions(program);
>+    synthesizeArrayOperatorLength(program);
>+    synthesizeConstructors(program);
>+    resolveNamesInFunctions(program, nameResolver);
>+    if (!checkDuplicateFunctions(program))
>+        return WTF::nullopt;
>+
>+    if (!check(program))
>+        return WTF::nullopt;
>+    checkLiteralTypes(program);
>+    // resolveProperties(program);
>+    findHighZombies(program);
>+    if (!checkStatementBehavior(program))
>+        return WTF::nullopt;
>+    if (!checkRecursion(program))
>+        return WTF::nullopt;
>+    if (!checkFunctionStages(program))
>+        return WTF::nullopt;
>+    return program;
>+}
>+
>+Optional<RenderPrepareResult> prepare(String& whlslSource, RenderPipelineDescriptor& renderPipelineDescriptor)
>+{
>+    auto program = prepareShared(whlslSource);
>+    if (!program)
>+        return WTF::nullopt;
>+    auto matchedSemantics = matchSemantics(*program, renderPipelineDescriptor);
>+    if (!matchedSemantics)
>+        return WTF::nullopt;
>+
>+    auto generatedCode = Metal::generateMetalCode(*program, WTFMove(*matchedSemantics), renderPipelineDescriptor.layout);
>+
>+    RenderPrepareResult result;
>+    result.metalSource = WTFMove(generatedCode.metalSource);
>+    result.vertexMappedBindGroups = WTFMove(generatedCode.vertexMappedBindGroups);
>+    result.fragmentMappedBindGroups = WTFMove(generatedCode.fragmentMappedBindGroups);
>+    return result;
>+}
>+
>+Optional<ComputePrepareResult> prepare(String& whlslSource, ComputePipelineDescriptor& computePipelineDescriptor)
>+{
>+    auto program = prepareShared(whlslSource);
>+    if (!program)
>+        return WTF::nullopt;
>+    auto matchedSemantics = matchSemantics(*program, computePipelineDescriptor);
>+    if (!matchedSemantics)
>+        return WTF::nullopt;
>+
>+    auto generatedCode = Metal::generateMetalCode(*program, WTFMove(*matchedSemantics), computePipelineDescriptor.layout);
>+
>+    ComputePrepareResult result;
>+    result.metalSource = WTFMove(generatedCode.metalSource);
>+    result.mappedBindGroups = WTFMove(generatedCode.bindGroups);
>+    return result;
>+}
>+
>+} // namespace WHLSL
>+
>+} // namespace WebCore
>+
>+#endif // ENABLE(WEBGPU)
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.h b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.h
>new file mode 100644
>index 0000000000000000000000000000000000000000..0f2c8466bf565b7b1e50c4aaecfdb353f2dc8cf0
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.h
>@@ -0,0 +1,57 @@
>+/*
>+ * 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 "WHLSLMappedBindings.h"
>+#include "WHLSLPipelineDescriptor.h"
>+#include <wtf/Optional.h>
>+#include <wtf/text/WTFString.h>
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+// FIXME: Generate descriptive error messages and return them here.
>+struct RenderPrepareResult {
>+    String metalSource;
>+    Metal::MappedBindGroups vertexMappedBindGroups;
>+    Metal::MappedBindGroups fragmentMappedBindGroups;
>+};
>+Optional<RenderPrepareResult> prepare(String& whlslSource, RenderPipelineDescriptor&);
>+
>+struct ComputePrepareResult {
>+    String metalSource;
>+    Metal::MappedBindGroups mappedBindGroups;
>+};
>+Optional<ComputePrepareResult> prepare(String& whlslSource, 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..0bdcc0ce8f1da576c5bf734b936e6c6a9793298a
>--- /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 = WTF::get<AST::StageInOutSemantic>(*fragmentInput.semantic);
>+        bool found = false;
>+        for (auto& vertexOutput : vertexOutputs) {
>+            if (!WTF::holds_alternative<AST::StageInOutSemantic>(*vertexOutput.semantic))
>+                continue;
>+            auto& vertexOutputStageInOutSemantic = WTF::get<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..d1991ee4f2613632b02c4ea7d76e74e2ebb94d76 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
>@@ -345,6 +346,7 @@ Modules/webgpu/WHLSL/AST/WHLSLUnsignedIntegerLiteralType.cpp
> Modules/webgpu/WHLSL/AST/WHLSLTypeReference.cpp
> Modules/webgpu/WHLSL/AST/WHLSLIntegerLiteral.cpp
> Modules/webgpu/WHLSL/AST/WHLSLUnsignedIntegerLiteral.cpp
>+Modules/webgpu/WHLSL/WHLSLPrepare.cpp
> Modules/webgpu/WebGPU.cpp
> Modules/webgpu/WebGPUBindGroup.cpp
> Modules/webgpu/WebGPUAdapter.cpp
>diff --git a/Source/WebCore/WebCore.xcodeproj/project.pbxproj b/Source/WebCore/WebCore.xcodeproj/project.pbxproj
>index c3d8fdb0456340c9ccf04a02f7af143988642c64..f68144302b5b39e7a3463129f0dd8b239e2fde93 100644
>--- a/Source/WebCore/WebCore.xcodeproj/project.pbxproj
>+++ b/Source/WebCore/WebCore.xcodeproj/project.pbxproj
>@@ -13391,6 +13391,12 @@
> 		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>"; };
>+		C24A57BA21FEAFEA004C6DD1 /* WHLSLPrepare.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLPrepare.cpp; sourceTree = "<group>"; };
>+		C24A57BB21FEAFEA004C6DD1 /* WHLSLPrepare.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLPrepare.h; sourceTree = "<group>"; };
>+		C24A57BE21FEC65C004C6DD1 /* WHLSLMappedBindings.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLMappedBindings.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>"; };
>@@ -17167,6 +17173,7 @@
> 				1CECB3BB21F511AA00F44542 /* WHLSLEntryPointScaffolding.h */,
> 				1CECB3AF21F2B98400F44542 /* WHLSLFunctionWriter.cpp */,
> 				1CECB3B221F2B98600F44542 /* WHLSLFunctionWriter.h */,
>+				C24A57BE21FEC65C004C6DD1 /* WHLSLMappedBindings.h */,
> 				1CECB3B521F50AC700F44542 /* WHLSLMetalCodeGenerator.cpp */,
> 				1CECB3B621F50AC700F44542 /* WHLSLMetalCodeGenerator.h */,
> 				1CECB3B821F50D1000F44542 /* WHLSLNativeFunctionWriter.cpp */,
>@@ -25602,6 +25609,9 @@
> 				C234A98C21E8883E003C984D /* WHLSLNameResolver.h */,
> 				C21BF73721CD8A0200227979 /* WHLSLParser.cpp */,
> 				C21BF73821CD8A0300227979 /* WHLSLParser.h */,
>+				C24A57AF21FAD53F004C6DD1 /* WHLSLPipelineDescriptor.h */,
>+				C24A57BA21FEAFEA004C6DD1 /* WHLSLPrepare.cpp */,
>+				C24A57BB21FEAFEA004C6DD1 /* WHLSLPrepare.h */,
> 				C21BF73A21CD8D7000227979 /* WHLSLProgram.h */,
> 				1CA0C2E021EEB5F500A11860 /* WHLSLRecursionChecker.cpp */,
> 				1CA0C2DE21EEB5F400A11860 /* WHLSLRecursionChecker.h */,
>@@ -25610,6 +25620,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