[WHLSL] Pack and unpack data at entry points and exit points
authormmaxfield@apple.com <mmaxfield@apple.com@268f45cc-cd09-0410-ab3c-d52691b4dbfc>
Wed, 6 Mar 2019 23:08:47 +0000 (23:08 +0000)
committermmaxfield@apple.com <mmaxfield@apple.com@268f45cc-cd09-0410-ab3c-d52691b4dbfc>
Wed, 6 Mar 2019 23:08:47 +0000 (23:08 +0000)
https://bugs.webkit.org/show_bug.cgi?id=193877

Reviewed by Dean Jackson.

This implements the bulk of the rest of the Metal code generation piece.
Specifically, this patch adds support for the packing and unpacking steps
at the beginning of the entry point functions and all returns from those
functions. It does this by creating an object, EntryPointScaffolding, that
knows how to emit all the helper types, packing / unpacking code, and the
function signature. This patch also accepts parts of the
PipelineStateDescriptor so we know the data layout we should be packing and
unpacking.

This is the last patch before enough of the compiler is present that we can
start testing it. The next patch will start porting the test suite.

* 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:

git-svn-id: https://svn.webkit.org/repository/webkit/trunk@242571 268f45cc-cd09-0410-ab3c-d52691b4dbfc

18 files changed:
Source/WebCore/ChangeLog
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.h
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMappedBindings.h [new file with mode: 0644]
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.cpp
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.h
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLVertexBufferIndexCalculator.cpp
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLVertexBufferIndexCalculator.h
Source/WebCore/Modules/webgpu/WHLSL/WHLSLIntrinsics.h
Source/WebCore/Modules/webgpu/WHLSL/WHLSLPipelineDescriptor.h [new file with mode: 0644]
Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.cpp [new file with mode: 0644]
Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.h [new file with mode: 0644]
Source/WebCore/Modules/webgpu/WHLSL/WHLSLSemanticMatcher.cpp [new file with mode: 0644]
Source/WebCore/Modules/webgpu/WHLSL/WHLSLSemanticMatcher.h [new file with mode: 0644]
Source/WebCore/Sources.txt
Source/WebCore/WebCore.xcodeproj/project.pbxproj

index 901211c..409284e 100644 (file)
@@ -1,3 +1,96 @@
+2019-03-06  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 Dean Jackson.
+
+        This implements the bulk of the rest of the Metal code generation piece.
+        Specifically, this patch adds support for the packing and unpacking steps
+        at the beginning of the entry point functions and all returns from those
+        functions. It does this by creating an object, EntryPointScaffolding, that
+        knows how to emit all the helper types, packing / unpacking code, and the
+        function signature. This patch also accepts parts of the
+        PipelineStateDescriptor so we know the data layout we should be packing and
+        unpacking.
+
+        This is the last patch before enough of the compiler is present that we can
+        start testing it. The next patch will start porting the test suite.
+
+        * 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-03-06  Wenson Hsieh  <wenson_hsieh@apple.com>
 
         Move RenderObject::isTransparentOrFullyClippedRespectingParentFrames() to RenderLayer
index 48fcf23..035284c 100644 (file)
 
 #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();
 }
 
index 8d65fad..d87f6e9 100644 (file)
 
 #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;
 };
 
 }
index 019a267..1b7a42a 100644 (file)
@@ -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
index 9960e5e..062821f 100644 (file)
@@ -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 (file)
index 0000000..db3c52b
--- /dev/null
@@ -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
index 10c0530..305f338 100644 (file)
@@ -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");
@@ -48,15 +48,32 @@ String generateMetalCode(Program& program)
     stringBuilder.append("#include <metal_compute>\n");
     stringBuilder.append("#include <metal_texture>\n");
     stringBuilder.append("\n");
-    stringBuilder.append("using namespace metal;\n");
+    stringBuilder.append("using namespace metal;\n"); // FIXME: Probably should qualify all calls to built-in functions, instead of using this line.
     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) };
+}
+
 }
 
 }
index e088e01..a6b3d31 100644 (file)
 
 #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,20 @@ 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;
+};
+// Can't fail. Any failure checks need to be done earlier, in the backend-agnostic part of the compiler.
+RenderMetalCode generateMetalCode(Program&, MatchedRenderSemantics&& matchedSemantics, Layout&);
+
+struct ComputeMetalCode {
+    String metalSource;
+    MappedBindGroups bindGroups;
+};
+// Can't fail. Any failure checks need to be done earlier, in the backend-agnostic part of the compiler.
+ComputeMetalCode generateMetalCode(Program&, MatchedComputeSemantics&& matchedSemantics, Layout&);
 
 }
 
index e26831e..a8091e9 100644 (file)
@@ -40,10 +40,10 @@ unsigned long calculateVertexBufferIndex(unsigned long index)
     return index + 8;
 }
 
-} // namespace Metal
+}
 
-} // namespace WHLSL
+}
 
-} // namespace WebCore
+}
 
 #endif // ENABLE(WEBGPU)
index bb55b89..9f8c253 100644 (file)
@@ -35,10 +35,10 @@ namespace Metal {
 
 unsigned long calculateVertexBufferIndex(unsigned long);
 
-} // namespace Metal
+}
 
-} // namespace WHLSL
+}
 
-} // namespace WebCore
+}
 
 #endif // ENABLE(WEBGPU)
index 8410c8b..fc47d32 100644 (file)
@@ -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 (file)
index 0000000..0f28692
--- /dev/null
@@ -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 (file)
index 0000000..4ec2bd2
--- /dev/null
@@ -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 (file)
index 0000000..0f2c846
--- /dev/null
@@ -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 (file)
index 0000000..0bdcc0c
--- /dev/null
@@ -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 (file)
index 0000000..bdff4bb
--- /dev/null
@@ -0,0 +1,65 @@
+/*
+ * 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"
+#include <wtf/HashMap.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)
index b575e68..8415770 100644 (file)
@@ -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
@@ -346,6 +347,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/WebGPUBindGroupDescriptor.cpp
index 8f0909b..1f26c39 100644 (file)
                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>"; };
                                1CECB3BB21F511AA00F44542 /* WHLSLEntryPointScaffolding.h */,
                                1CECB3AF21F2B98400F44542 /* WHLSLFunctionWriter.cpp */,
                                1CECB3B221F2B98600F44542 /* WHLSLFunctionWriter.h */,
+                               C24A57BE21FEC65C004C6DD1 /* WHLSLMappedBindings.h */,
                                1CECB3B521F50AC700F44542 /* WHLSLMetalCodeGenerator.cpp */,
                                1CECB3B621F50AC700F44542 /* WHLSLMetalCodeGenerator.h */,
                                1CECB3B821F50D1000F44542 /* WHLSLNativeFunctionWriter.cpp */,
                                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 */,
                                C234A99921E90F29003C984D /* WHLSLResolveOverloadImpl.cpp */,
                                C234A99721E90F28003C984D /* WHLSLResolveOverloadImpl.h */,
                                C234A99D21E910BD003C984D /* WHLSLResolvingType.h */,
+                               C24A57B221FB8DDA004C6DD1 /* WHLSLSemanticMatcher.cpp */,
+                               C24A57B321FB8DDA004C6DD1 /* WHLSLSemanticMatcher.h */,
                                C21BF74521CD969800227979 /* WHLSLStandardLibrary.txt */,
                                1CECB3A921F2B67300F44542 /* WHLSLStatementBehaviorChecker.cpp */,
                                1CECB3A821F2B67300F44542 /* WHLSLStatementBehaviorChecker.h */,