Attachment 359643 Details for Bug 193531 – WIP

[patch] WIP

bug-193531-20190120130338.patch (text/plain), 148.38 KB, created by Myles C. Maxfield on 2019-01-20 13:03:39 PST

(hide)

Description:

Filename:

MIME Type:

Creator: Myles C. Maxfield

Created: 2019-01-20 13:03:39 PST

Size: 148.38 KB

patch

obsolete

>Subversion Revision: 240191
>diff --git a/Source/WebCore/ChangeLog b/Source/WebCore/ChangeLog
>index c041dc2d1645f44974f0bd49b44573c0b63763d6..842a9d0f39da9fa5c67a1d953d0209985114d77d 100644
>--- a/Source/WebCore/ChangeLog
>+++ b/Source/WebCore/ChangeLog
>@@ -1,3 +1,101 @@
>+2019-01-20  Myles C. Maxfield  <mmaxfield@apple.com>
>+
>+        [WHLSL] Implement Metal code generation
>+        https://bugs.webkit.org/show_bug.cgi?id=193531
>+
>+        Reviewed by NOBODY (OOPS!).
>+
>+        No new tests (OOPS!).
>+
>+        * Modules/webgpu/WHLSL/AST/WHLSLBuiltInSemantic.h:
>+        (WebCore::WHLSL::AST::BuiltInSemantic::targetIndex):
>+        * Modules/webgpu/WHLSL/AST/WHLSLExpression.h:
>+        (WebCore::WHLSL::AST::Expression::resolvedType):
>+        (WebCore::WHLSL::AST::Expression::type): Deleted.
>+        * Modules/webgpu/WHLSL/AST/WHLSLFunctionDeclaration.h:
>+        (WebCore::WHLSL::AST::FunctionDeclaration::name):
>+        * Modules/webgpu/WHLSL/AST/WHLSLStructureDefinition.h:
>+        (WebCore::WHLSL::AST::StructureDefinition::find):
>+        * Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp: Added.
>+        (WebCore::WHLSL::Metal::EntryPointScaffolding::EntryPointScaffolding):
>+        (WebCore::WHLSL::Metal::EntryPointScaffolding::helperTypes):
>+        (WebCore::WHLSL::Metal::EntryPointScaffolding::signature):
>+        (WebCore::WHLSL::Metal::EntryPointScaffolding::unpack):
>+        (WebCore::WHLSL::Metal::EntryPointScaffolding::pack):
>+        * Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h: Copied from Source/WebCore/Modules/webgpu/WHLSL/WHLSLLoopChecker.h.
>+        * Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp: Added.
>+        (WebCore::WHLSL::Metal::FunctionDeclarationWriter::FunctionDeclarationWriter):
>+        (WebCore::WHLSL::Metal::FunctionDeclarationWriter::toString):
>+        (WebCore::WHLSL::Metal::FunctionDefinitionWriter::FunctionDefinitionWriter):
>+        (WebCore::WHLSL::Metal::FunctionDefinitionWriter::toString):
>+        (WebCore::WHLSL::Metal::FunctionDefinitionWriter::constantExpressionString):
>+        (WebCore::WHLSL::Metal::FunctionDefinitionWriter::generateNextVariableName):
>+        (WebCore::WHLSL::Metal::metalFunctions):
>+        * Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.h: Copied from Source/WebCore/Modules/webgpu/WHLSL/WHLSLLoopChecker.h.
>+        * Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.cpp: Copied from Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLStructureDefinition.h.
>+        (WebCore::WHLSL::Metal::generateMetalCode):
>+        * Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.h: Copied from Source/WebCore/Modules/webgpu/WHLSL/WHLSLLoopChecker.h.
>+        * Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.cpp: Added.
>+        (WebCore::WHLSL::Metal::getNativeName):
>+        (WebCore::WHLSL::Metal::mapFunctionName):
>+        (WebCore::WHLSL::Metal::convertAddressSpace):
>+        (WebCore::WHLSL::Metal::writeNativeFunction):
>+        * Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.h: Copied from Source/WebCore/Modules/webgpu/WHLSL/WHLSLLoopChecker.h.
>+        * Modules/webgpu/WHLSL/Metal/WHLSLTypeNamer.cpp: Added.
>+        (WebCore::WHLSL::Metal::BaseTypeNameNode::BaseTypeNameNode):
>+        (WebCore::WHLSL::Metal::BaseTypeNameNode::isArrayTypeNameNode const):
>+        (WebCore::WHLSL::Metal::BaseTypeNameNode::isArrayReferenceTypeNameNode const):
>+        (WebCore::WHLSL::Metal::BaseTypeNameNode::isPointerTypeNameNode const):
>+        (WebCore::WHLSL::Metal::BaseTypeNameNode::isReferenceTypeNameNode const):
>+        (WebCore::WHLSL::Metal::BaseTypeNameNode::children):
>+        (WebCore::WHLSL::Metal::BaseTypeNameNode::append):
>+        (WebCore::WHLSL::Metal::BaseTypeNameNode::parent):
>+        (WebCore::WHLSL::Metal::BaseTypeNameNode::mangledName const):
>+        (WebCore::WHLSL::Metal::ArrayTypeNameNode::ArrayTypeNameNode):
>+        (WebCore::WHLSL::Metal::ArrayTypeNameNode::numElements const):
>+        (WebCore::WHLSL::Metal::ArrayReferenceTypeNameNode::ArrayReferenceTypeNameNode):
>+        (WebCore::WHLSL::Metal::ArrayReferenceTypeNameNode::addressSpace const):
>+        (WebCore::WHLSL::Metal::PointerTypeNameNode::PointerTypeNameNode):
>+        (WebCore::WHLSL::Metal::PointerTypeNameNode::addressSpace const):
>+        (WebCore::WHLSL::Metal::ReferenceTypeNameNode::ReferenceTypeNameNode):
>+        (WebCore::WHLSL::Metal::ReferenceTypeNameNode::namedType):
>+        (WebCore::WHLSL::Metal::TypeNamer::TypeNamer):
>+        (WebCore::WHLSL::Metal::TypeNamer::visit):
>+        (WebCore::WHLSL::Metal::findInVector):
>+        (WebCore::WHLSL::Metal::find):
>+        (WebCore::WHLSL::Metal::TypeNamer::mangledNameForType):
>+        (WebCore::WHLSL::Metal::TypeNamer::createNameNode):
>+        (WebCore::WHLSL::Metal::TypeNamer::insert):
>+        (WebCore::WHLSL::Metal::MetalTypeDeclarationWriter::MetalTypeDeclarationWriter):
>+        (WebCore::WHLSL::Metal::MetalTypeDeclarationWriter::toString):
>+        (WebCore::WHLSL::Metal::MetalTypeDeclarationWriter::visit):
>+        (WebCore::WHLSL::Metal::TypeNamer::metalTypeDeclarations):
>+        (WebCore::WHLSL::Metal::toString):
>+        (WebCore::WHLSL::Metal::TypeNamer::emitUnnamedTypeDefinition):
>+        (WebCore::WHLSL::Metal::TypeNamer::emitNamedTypeDefinition):
>+        (WebCore::WHLSL::Metal::TypeNamer::metalTypeDefinitions):
>+        (WebCore::WHLSL::Metal::TypeNamer::mangledNameForEnumerationMember):
>+        (WebCore::WHLSL::Metal::TypeNamer::mangledNameForStructureElement):
>+        (WebCore::WHLSL::Metal::TypeNamer::metalTypes):
>+        * Modules/webgpu/WHLSL/Metal/WHLSLTypeNamer.h: Added.
>+        (WebCore::WHLSL::Metal::TypeNamer::generateNextTypeName):
>+        (WebCore::WHLSL::Metal::TypeNamer::generateNextStructureElementName):
>+        (WebCore::WHLSL::Metal::TypeNamer::generateNextEnumerationMemberName):
>+        * Modules/webgpu/WHLSL/WHLSLChecker.cpp:
>+        (WebCore::WHLSL::checkSemantics):
>+        (WebCore::WHLSL::Checker::visit):
>+        * Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.cpp:
>+        (WebCore::WHLSL::Gatherer::visit):
>+        * Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.h:
>+        (WebCore::WHLSL::EntryPointItem::EntryPointItem):
>+        * Modules/webgpu/WHLSL/WHLSLLoopChecker.cpp: Removed.
>+        * Modules/webgpu/WHLSL/WHLSLStatementBehaviorChecker.cpp: Added.
>+        (WebCore::WHLSL::StatementBehaviorChecker::takeFunctionBehavior):
>+        (WebCore::WHLSL::checkStatementBehavior):
>+        * Modules/webgpu/WHLSL/WHLSLStatementBehaviorChecker.h: Renamed from Source/WebCore/Modules/webgpu/WHLSL/WHLSLLoopChecker.h.
>+        * Sources.txt:
>+        * WebCore.xcodeproj/project.pbxproj:
>+
> 2019-01-18  Eric Carlson  <eric.carlson@apple.com>
> 
>         Revert r238815, it broke WK1 video fullscreen on Mac
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLBuiltInSemantic.h b/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLBuiltInSemantic.h
>index 43b6986af7d4b71c28677ec65cecd8ea7c47304a..7b36caabc6bfca5d445e22b8d62db9ea2c78a082 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLBuiltInSemantic.h
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLBuiltInSemantic.h
>@@ -69,6 +69,7 @@ public:
>     BuiltInSemantic(BuiltInSemantic&&) = default;
> 
>     Variable variable() const { return m_variable; }
>+    Optional<unsigned>& targetIndex() { return m_targetIndex; }
> 
>     bool operator==(const BuiltInSemantic& other) const
>     {
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLExpression.h b/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLExpression.h
>index 8fef62959d7363e1f1bbb23d73398f8a2d48bde1..01f2ffb806a0576bc6056e061ab172be264616ee 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLExpression.h
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLExpression.h
>@@ -57,7 +57,7 @@ public:
> 
>     const Lexer::Token& origin() const { return m_origin; }
> 
>-    UnnamedType* type() { return m_type ? &*m_type : nullptr; }
>+    UnnamedType* resolvedType() { return m_type ? &*m_type : nullptr; }
> 
>     void setType(UniqueRef<UnnamedType>&& type)
>     {
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLFunctionDeclaration.h b/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLFunctionDeclaration.h
>index 508285846070577d894c534a3974bce0abb21f11..c15466a70bbbaaa75f170e126eb52dc88b84816d 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLFunctionDeclaration.h
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLFunctionDeclaration.h
>@@ -70,6 +70,7 @@ public:
>     const UnnamedType& type() const { return m_type; }
>     UnnamedType& type() { return m_type; }
>     const String& name() const { return m_name; }
>+    String& name() { return m_name; }
>     bool isCast() const { return m_name == "operator cast"; }
>     const VariableDeclarations& parameters() const { return m_parameters; }
>     VariableDeclarations& parameters() { return m_parameters; }
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLStructureDefinition.h b/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLStructureDefinition.h
>index 456c6e1238bcb57e7cfe1e017869f0f41440a769..0260d4506a51cab53819331492e5d7becbdc995b 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLStructureDefinition.h
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLStructureDefinition.h
>@@ -55,6 +55,15 @@ public:
>     bool isStructureDefinition() const override { return true; }
> 
>     StructureElements& structureElements() { return m_structureElements; }
>+    StructureElement* find(String& name)
>+    {
>+        auto iterator = std::find_if(m_structureElements.begin(), m_structureElements.end(), [&](StructureElement& structureElement) -> bool {
>+            return structureElement.name() == name;
>+        });
>+        if (iterator == m_structureElements.end())
>+            return nullptr;
>+        return &*iterator;
>+    }
> 
> private:
>     StructureElements m_structureElements;
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp
>new file mode 100644
>index 0000000000000000000000000000000000000000..0a9cff46a7937eaaf3cc53917ede941843f8105f
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp
>@@ -0,0 +1,256 @@
>+/*
>+ * 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 "WHLSLEntryPointScaffolding.h"
>+
>+#if ENABLE(WEBGPU)
>+
>+#include "WHLSLGatherEntryPointItems.h"
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+namespace Metal {
>+
>+EntryPointScaffolding::EntryPointScaffolding(AST::FunctionDefinition& functionDefinition, Intrinsics& intrinsics)
>+    : m_functionDefinition(&functionDefinition)
>+    , m_intrinsics(&intrinsics)
>+{
>+    // FIXME: Implement this.
>+    gatherEntryPointItems(*m_intrinsics, *m_functionDefinition);
>+}
>+
>+String EntryPointScaffolding::helperTypes()
>+{
>+    // FIXME: Implement this.
>+    return String();
>+}
>+
>+String EntryPointScaffolding::signature()
>+{
>+    // FIXME: Implement this.
>+    return String();
>+}
>+
>+String EntryPointScaffolding::unpack()
>+{
>+    // FIXME: Implement this.
>+    return String();
>+}
>+
>+String EntryPointScaffolding::pack(const String&)
>+{
>+    // FIXME: Implement this.
>+    return String();
>+}
>+
>+/*
>+{
>+            auto items = gatherEntryPointItems(m_intrinsics, functionDefinition);
>+            m_entryPointItems = items;
>+            m_entryPointReturnType = &functionDefinition.type();
>+            ASSERT(items);
>+            unsigned textureCount = 0;
>+            unsigned bufferCount = 0;
>+            unsigned samplerCount = 0;
>+
>+            auto semanticVisitor = WTF::makeVisitor([&](AST::BuiltInSemantic& builtInSemantic) -> String {
>+                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;
>+                }
>+            }, [&](AST::ResourceSemantic& resourceSemantic) -> String {
>+                // FIXME: Provide a better mapping between HLSL semantics and Metal attributes.
>+                // FIXME: Report this information to the API.
>+                switch (resourceSemantic.mode()) {
>+                case AST::ResourceSemantic::Mode::UnorderedAccessView:
>+                    return makeString("[[buffer(", bufferCount++, ")]]");
>+                case AST::ResourceSemantic::Mode::Texture:
>+                    return makeString("[[texture(", textureCount++, ")]]");
>+                case AST::ResourceSemantic::Mode::Buffer:
>+                    return makeString("[[buffer(", bufferCount++, ")]]");
>+                default:
>+                    ASSERT(resourceSemantic.mode() == AST::ResourceSemantic::Mode::Sampler);
>+                    return makeString("[[sampler(", samplerCount++, ")]]");
>+                }
>+            }, [&](AST::SpecializationConstantSemantic&) -> String {
>+                // FIXME: Implement this
>+                CRASH();
>+            }, [&](AST::StageInOutSemantic& stageInOutSemantic) -> String {
>+                return makeString("[[attribute(", stageInOutSemantic.index(), ")]]");
>+            });
>+
>+            m_outputStructName = m_typeNamer.generateNextTypeName();
>+            m_metalOutputNames.clear();
>+            m_stringBuilder.append(makeString("struct ", m_outputStructName, "{\n"));
>+            for (auto& output : items->outputs) {
>+                auto elementName = m_typeNamer.generateNextStructureElementName();
>+                m_metalOutputNames.append(elementName);
>+                m_stringBuilder.append(makeString("    ", m_typeNamer.mangledNameForType(*output.unnamedType), ' ', elementName, ' ', WTF::visit(semanticVisitor, *output.semantic)));
>+            }
>+            m_stringBuilder.append("};\n");
>+            auto entryPoint = ([&]() -> String {
>+                switch (*functionDefinition.entryPointType()) {
>+                case AST::EntryPointType::Vertex:
>+                    return "vertex"_str;
>+                case AST::EntryPointType::Fragment:
>+                    return "fragment"_str;
>+                default:
>+                    ASSERT(*functionDefinition.entryPointType() == AST::EntryPointType::Compute);
>+                    return "kernel"_str;
>+                }
>+            })();
>+            Vector<String> metalParameterNames;
>+            metalParameterNames.reserveCapacity(items->inputs.size());
>+            m_stringBuilder.append(makeString(entryPoint, " ", m_typeNamer.mangledNameForType(functionDefinition.type()), ' ', iterator->value, '('));
>+            for (size_t i = 0; i < items->inputs.size(); ++i) {
>+                auto& input = items->inputs[i];
>+                if (i)
>+                    m_stringBuilder.append(", ");
>+                auto parameterName = generateNextVariableName();
>+                metalParameterNames.uncheckedAppend(parameterName);
>+                m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*input.unnamedType), ' ', parameterName, ' ', WTF::visit(semanticVisitor, *input.semantic)));
>+            }
>+            m_stringBuilder.append(") {\n");
>+            HashMap<String, String> parameterNameMap;
>+            HashMap<String, AST::VariableDeclaration*> parameterObjectMap;
>+            for (auto& parameter : functionDefinition.parameters()) {
>+                auto parameterName = generateNextVariableName();
>+                {
>+                    auto addResult = m_variableMapping.add(&parameter, parameterName);
>+                    ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+                }
>+                {
>+                    auto addResult = parameterNameMap.add(parameter.name(), parameterName);
>+                    ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+                }
>+                {
>+                    auto addResult = parameterObjectMap.add(parameter.name(), &parameter);
>+                    ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+                }
>+                m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*parameter.type()), ' ', parameterName, ";\n"));
>+            }
>+            for (size_t i = 0; i < items->inputs.size(); ++i) {
>+                auto& input = items->inputs[i];
>+                AST::StructureDefinition* structureDefinition = nullptr;
>+                auto computeStructureDefinition = [](AST::UnnamedType& unnamedType) -> AST::StructureDefinition* {
>+                    auto& unifyNode = unnamedType.unifyNode();
>+                    if (!is<AST::NamedType>(unifyNode))
>+                        return nullptr;
>+                    auto& namedType = downcast<AST::NamedType>(unifyNode);
>+                    if (!is<AST::StructureDefinition>(namedType))
>+                        return nullptr;
>+                    return &downcast<AST::StructureDefinition>(namedType);
>+                };
>+                for (size_t j = 0; j < input.path.size(); ++j) {
>+                    auto& component = input.path[j];
>+                    if (!j) {
>+                        {
>+                            auto iterator = parameterNameMap.find(component);
>+                            ASSERT(iterator != parameterNameMap.end());
>+                            m_stringBuilder.append(iterator->value);
>+                        }
>+                        auto iterator = parameterObjectMap.find(component);
>+                        ASSERT(iterator != parameterObjectMap.end());
>+                        structureDefinition = computeStructureDefinition(*iterator->value->type());
>+                    } else {
>+                        auto* structureElement = structureDefinition->find(component);
>+                        ASSERT(structureElement);
>+                        m_stringBuilder.append(makeString('.', m_typeNamer.mangledNameForStructureElement(*structureElement)));
>+                        structureDefinition = computeStructureDefinition(structureElement->type());
>+                    }
>+                }
>+                m_stringBuilder.append(makeString(" = ", metalParameterNames[i], ";\n"));
>+            }
>+            checkErrorAndVisit(functionDefinition.block());
>+            m_stringBuilder.append("}\n");
>+        }
>+        */
>+
>+        /*{
>+                auto variableName = generateNextVariableName();
>+                m_stringBuilder.append(makeString(m_outputStructName, ' ', variableName, ";\n"));
>+                for (size_t i = 0; i < m_entryPointItems->outputs.size(); ++i) {
>+                    auto computeStructureDefinition = [](AST::UnnamedType& unnamedType) -> AST::StructureDefinition* {
>+                        auto& unifyNode = unnamedType.unifyNode();
>+                        if (!is<AST::NamedType>(unifyNode))
>+                            return nullptr;
>+                        auto& namedType = downcast<AST::NamedType>(unifyNode);
>+                        if (!is<AST::StructureDefinition>(namedType))
>+                            return nullptr;
>+                        return &downcast<AST::StructureDefinition>(namedType);
>+                    };
>+                    m_stringBuilder.append(makeString(variableName, '.', m_metalOutputNames[i], '='));
>+                    AST::StructureDefinition* structureDefinition = computeStructureDefinition(*m_entryPointReturnType);
>+                    ASSERT(structureDefinition);
>+                    for (size_t j = 0; j < m_entryPointItems->outputs[i].path.size(); ++j) {
>+                        auto* structureElement = structureDefinition->find(m_entryPointItems->outputs[i].path[j]);
>+                        ASSERT(structureElement);
>+                        if (!j)
>+                            m_stringBuilder.append('.');
>+                        m_stringBuilder.append(m_typeNamer.mangledNameForStructureElement(*structureElement));
>+                        structureDefinition = computeStructureDefinition(structureElement->type());
>+                    }
>+                    m_stringBuilder.append(";\n");
>+                }
>+                m_stringBuilder.append(makeString("return ", variableName, ";\n"));
>+            }*/
>+
>+}
>+
>+}
>+
>+}
>+
>+#endif
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h
>new file mode 100644
>index 0000000000000000000000000000000000000000..1aa2da3bc4eca480c252e63ee6a1bca86295101c
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h
>@@ -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)
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+namespace AST {
>+
>+class FunctionDefinition;
>+
>+}
>+
>+class Intrinsics;
>+
>+namespace Metal {
>+
>+// FIXME: This needs to know about the pipeline state object to emit function prologues and epilogues.
>+class EntryPointScaffolding {
>+public:
>+    EntryPointScaffolding(AST::FunctionDefinition&, Intrinsics&);
>+
>+    String helperTypes();
>+    String signature();
>+    String unpack();
>+    String pack(const String& variableName);
>+
>+private:
>+    AST::FunctionDefinition* m_functionDefinition;
>+    Intrinsics* m_intrinsics;
>+};
>+
>+}
>+
>+}
>+
>+}
>+
>+#endif
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp
>new file mode 100644
>index 0000000000000000000000000000000000000000..7a0916020db9a9b1b0bf272cc0ebac962205d89d
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp
>@@ -0,0 +1,625 @@
>+/*
>+ * 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 "WHLSLFunctionWriter.h"
>+
>+#if ENABLE(WEBGPU)
>+
>+#include "WHLSLArrayReferenceType.h"
>+#include "WHLSLArrayType.h"
>+#include "WHLSLAssignmentExpression.h"
>+#include "WHLSLBooleanLiteral.h"
>+#include "WHLSLBuiltInSemantic.h"
>+#include "WHLSLCallExpression.h"
>+#include "WHLSLCommaExpression.h"
>+#include "WHLSLDereferenceExpression.h"
>+#include "WHLSLDoWhileLoop.h"
>+#include "WHLSLEffectfulExpressionStatement.h"
>+#include "WHLSLEntryPointScaffolding.h"
>+#include "WHLSLEntryPointType.h"
>+#include "WHLSLFloatLiteral.h"
>+#include "WHLSLForLoop.h"
>+#include "WHLSLFunctionDeclaration.h"
>+#include "WHLSLFunctionDefinition.h"
>+#include "WHLSLIfStatement.h"
>+#include "WHLSLIntegerLiteral.h"
>+#include "WHLSLLogicalExpression.h"
>+#include "WHLSLLogicalNotExpression.h"
>+#include "WHLSLMakeArrayReferenceExpression.h"
>+#include "WHLSLMakePointerExpression.h"
>+#include "WHLSLNativeFunctionDeclaration.h"
>+#include "WHLSLNativeFunctionWriter.h"
>+#include "WHLSLNativeTypeDeclaration.h"
>+#include "WHLSLPointerType.h"
>+#include "WHLSLProgram.h"
>+#include "WHLSLReturn.h"
>+#include "WHLSLSwitchCase.h"
>+#include "WHLSLSwitchStatement.h"
>+#include "WHLSLTernaryExpression.h"
>+#include "WHLSLTypeNamer.h"
>+#include "WHLSLUnsignedIntegerLiteral.h"
>+#include "WHLSLVariableDeclaration.h"
>+#include "WHLSLVariableDeclarationsStatement.h"
>+#include "WHLSLVariableReference.h"
>+#include "WHLSLVisitor.h"
>+#include "WHLSLWhileLoop.h"
>+#include <wtf/HashMap.h>
>+#include <wtf/text/StringBuilder.h>
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+namespace Metal {
>+
>+class FunctionDeclarationWriter : public Visitor {
>+public:
>+    FunctionDeclarationWriter(TypeNamer& typeNamer, HashMap<AST::FunctionDeclaration*, String>& functionMapping)
>+        : m_typeNamer(typeNamer)
>+        , m_functionMapping(functionMapping)
>+    {
>+    }
>+
>+    virtual ~FunctionDeclarationWriter() = default;
>+
>+    String toString() { return m_stringBuilder.toString(); }
>+    void visit(AST::FunctionDeclaration& functionDeclaration) override
>+    {
>+        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");
>+    }
>+
>+private:
>+    TypeNamer& m_typeNamer;
>+    HashMap<AST::FunctionDeclaration*, String>& m_functionMapping;
>+    StringBuilder m_stringBuilder;
>+};
>+
>+class FunctionDefinitionWriter : public Visitor {
>+public:
>+    FunctionDefinitionWriter(Intrinsics& intrinsics, TypeNamer& typeNamer, HashMap<AST::FunctionDeclaration*, String>& functionMapping)
>+        : m_intrinsics(intrinsics)
>+        , m_typeNamer(typeNamer)
>+        , m_functionMapping(functionMapping)
>+    {
>+    }
>+
>+    virtual ~FunctionDefinitionWriter() = default;
>+
>+    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::FunctionDefinition& functionDefinition) override
>+    {
>+        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());
>+            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());
>+            m_stringBuilder.append("}\n");
>+        }
>+    }
>+
>+private:
>+    void visit(AST::FunctionDeclaration&) override
>+    {
>+        ASSERT_NOT_REACHED();
>+    }
>+
>+    void visit(AST::Statement& statement) override
>+    {
>+        Visitor::visit(statement);
>+    }
>+
>+    void visit(AST::Block& block) override
>+    {
>+        m_stringBuilder.append("{\n");
>+        for (auto& statement : block.statements())
>+            checkErrorAndVisit(statement);
>+        m_stringBuilder.append("}\n");
>+    }
>+
>+    void visit(AST::Break&) override
>+    {
>+        m_stringBuilder.append("break;\n");
>+    }
>+
>+    void visit(AST::Continue&) override
>+    {
>+        // FIXME: Figure out which loop we're in, and run the increment code
>+        CRASH();
>+    }
>+
>+    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 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 visit(AST::Fallthrough&) override
>+    {
>+        m_stringBuilder.append("[[clang::fallthrough]];\n"); // FIXME: Make sure this is okay. Alternatively, we could do thing and just return here instead.
>+    }
>+
>+    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_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 visit(AST::Return& returnStatement) override
>+    {
>+        if (returnStatement.value()) {
>+            checkErrorAndVisit(*returnStatement.value());
>+            if (m_entryPointScaffolding) {
>+                auto variableName = generateNextVariableName();
>+                m_stringBuilder.append(m_entryPointScaffolding->pack(variableName));
>+                m_stringBuilder.append(makeString("return ", variableName, ";\n"));
>+            } else
>+                m_stringBuilder.append(makeString("return ", m_stack.takeLast(), ";\n"));
>+        } else
>+            m_stringBuilder.append("return;\n");
>+    }
>+
>+    void visit(AST::SwitchStatement& switchStatement) override
>+    {
>+        checkErrorAndVisit(switchStatement.value());
>+
>+        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");
>+        // FIXME: Figure out whether we need to break or fallthrough.
>+        CRASH();
>+        checkErrorAndVisit(switchCase.block());
>+    }
>+
>+    void visit(AST::Trap&) override
>+    {
>+        // FIXME: Implement this
>+        CRASH();
>+    }
>+
>+    void visit(AST::VariableDeclarationsStatement& variableDeclarationsStatement) override
>+    {
>+        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 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 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 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 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 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 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 visit(AST::Expression& expression) override
>+    {
>+        Visitor::visit(expression);
>+    }
>+
>+    void visit(AST::DotExpression&) override
>+    {
>+        // This should be lowered already.
>+        ASSERT_NOT_REACHED();
>+    }
>+
>+    void visit(AST::IndexExpression&) override
>+    {
>+        // This should be lowered already.
>+        ASSERT_NOT_REACHED();
>+    }
>+
>+    void visit(AST::PropertyAccessExpression&) override
>+    {
>+        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 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 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 visit(AST::CommaExpression& commaExpression) override
>+    {
>+        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 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 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 visit(AST::MakeArrayReferenceExpression& makeArrayReferenceExpression) override
>+    {
>+        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
>+    {
>+        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 visit(AST::ReadModifyWriteExpression&) override
>+    {
>+        // This should be lowered already.
>+        ASSERT_NOT_REACHED();
>+    }
>+
>+    void visit(AST::TernaryExpression& ternaryExpression) override
>+    {
>+        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 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);
>+    }
>+
>+    String 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;
>+    }
>+
>+    String generateNextVariableName()
>+    {
>+        return makeString("variable", m_variableCount++);
>+    }
>+
>+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 };
>+};
>+
>+String metalFunctions(Program& program, TypeNamer& typeNamer)
>+{
>+    StringBuilder stringBuilder;
>+
>+    unsigned numFunctions = 0;
>+    HashMap<AST::FunctionDeclaration*, String> functionMapping;
>+    for (auto& nativeFunctionDeclaration : program.nativeFunctionDeclarations()) {
>+        auto addResult = functionMapping.add(&nativeFunctionDeclaration, makeString("function", numFunctions++));
>+        ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+    }
>+    for (auto& functionDefinition : program.functionDefinitions()) {
>+        auto addResult = functionMapping.add(&functionDefinition, makeString("function", numFunctions++));
>+        ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+    }
>+
>+    {
>+        FunctionDeclarationWriter functionDeclarationWriter(typeNamer, functionMapping);
>+        for (auto& nativeFunctionDeclaration : program.nativeFunctionDeclarations())
>+            functionDeclarationWriter.visit(nativeFunctionDeclaration);
>+        for (auto& functionDefinition : program.functionDefinitions()) {
>+            if (!functionDefinition->entryPointType())
>+                functionDeclarationWriter.visit(functionDefinition);
>+        }
>+        stringBuilder.append(functionDeclarationWriter.toString());
>+    }
>+
>+    stringBuilder.append('\n');
>+
>+    {
>+        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());
>+    }
>+
>+    return stringBuilder.toString();
>+}
>+
>+} // namespace Metal
>+
>+} // namespace WHLSL
>+
>+} // namespace WebCore
>+
>+#endif
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.h b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.h
>new file mode 100644
>index 0000000000000000000000000000000000000000..9960e5e48d58c24e35e528382456472b7d09b6d7
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.h
>@@ -0,0 +1,48 @@
>+/*
>+ * 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)
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+class Program;
>+
>+namespace Metal {
>+
>+class TypeNamer;
>+
>+String metalFunctions(Program&, TypeNamer&);
>+
>+}
>+
>+}
>+
>+}
>+
>+#endif
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.cpp b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.cpp
>new file mode 100644
>index 0000000000000000000000000000000000000000..0277be5ddae183942fb626fd58513dbecd7ff503
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.cpp
>@@ -0,0 +1,68 @@
>+/*
>+ * 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 "WHLSLMetalCodeGenerator.h"
>+
>+#if ENABLE(WEBGPU)
>+
>+#include "WHLSLFunctionWriter.h"
>+#include "WHLSLTypeNamer.h"
>+#include <wtf/text/StringBuilder.h>
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+namespace Metal {
>+
>+String generateMetalCode(Program& program)
>+{
>+    StringBuilder stringBuilder;
>+    stringBuilder.append("// This was generated by WebKit's WHLSL -> MSL code generator. Do not edit!\n");
>+    stringBuilder.append("\n");
>+    stringBuilder.append("#include <metal_stdlib>\n");
>+    stringBuilder.append("#include <metal_atomic>\n");
>+    stringBuilder.append("#include <metal_math>\n");
>+    stringBuilder.append("#include <metal_relational>\n");
>+    stringBuilder.append("#include <metal_compute>\n");
>+    stringBuilder.append("#include <metal_texture>\n");
>+    stringBuilder.append("\n");
>+    stringBuilder.append("using namespace metal;\n");
>+    stringBuilder.append("\n");
>+
>+    TypeNamer typeNamer(program);
>+    stringBuilder.append(typeNamer.metalTypes());
>+    stringBuilder.append(metalFunctions(program, typeNamer));
>+    return stringBuilder.toString();
>+}
>+
>+}
>+
>+}
>+
>+}
>+
>+#endif
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.h b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.h
>new file mode 100644
>index 0000000000000000000000000000000000000000..e088e010449948411075674347491f73d36a234f
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.h
>@@ -0,0 +1,49 @@
>+/*
>+ * 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/text/WTFString.h>
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+class Program;
>+
>+namespace Metal {
>+
>+// FIXME: This needs to know about the pipeline state object to emit function prologues and epilogues.
>+String generateMetalCode(Program&);
>+
>+}
>+
>+}
>+
>+}
>+
>+#endif
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.cpp b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.cpp
>new file mode 100644
>index 0000000000000000000000000000000000000000..bffec631b7af73508e789368fcb4c07e9ab853f8
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.cpp
>@@ -0,0 +1,394 @@
>+/*
>+ * 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 "WHLSLNativeFunctionWriter.h"
>+
>+#if ENABLE(WEBGPU)
>+
>+#include "WHLSLAddressSpace.h"
>+#include "WHLSLNamedType.h"
>+#include "WHLSLNativeFunctionDeclaration.h"
>+#include "WHLSLNativeTypeDeclaration.h"
>+#include "WHLSLTypeNamer.h"
>+#include "WHLSLUnnamedType.h"
>+#include "WHLSLVariableDeclaration.h"
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+namespace Metal {
>+
>+static String getNativeName(AST::UnnamedType& unnamedType, TypeNamer& typeNamer)
>+{
>+    ASSERT(is<AST::NamedType>(unnamedType.unifyNode()));
>+    auto& namedType = downcast<AST::NamedType>(unnamedType.unifyNode());
>+    ASSERT(is<AST::NativeTypeDeclaration>(namedType));
>+    auto& nativeTypeDeclaration = downcast<AST::NativeTypeDeclaration>(namedType);
>+    return typeNamer.mangledNameForType(nativeTypeDeclaration);
>+}
>+
>+static String mapFunctionName(String& functionName)
>+{
>+    if (functionName == "ddx")
>+        return "dfdx"_str;
>+    if (functionName == "ddy")
>+        return "dfdy"_str;
>+    if (functionName == "asint")
>+        return "as_type<int32_t>"_str;
>+    if (functionName == "asuint")
>+        return "as_type<uint32_t>"_str;
>+    if (functionName == "asfloat")
>+        return "as_type<float>"_str;
>+    return functionName;
>+}
>+
>+static String convertAddressSpace(AST::AddressSpace addressSpace)
>+{
>+    switch (addressSpace) {
>+    case AST::AddressSpace::Constant:
>+        return "constant"_str;
>+    case AST::AddressSpace::Device:
>+        return "device"_str;
>+    case AST::AddressSpace::Threadgroup:
>+        return "threadgroup"_str;
>+    default:
>+        ASSERT(addressSpace == AST::AddressSpace::Thread);
>+        return "thread"_str;
>+    }
>+}
>+
>+String writeNativeFunction(AST::NativeFunctionDeclaration& nativeFunctionDeclaration, String& outputFunctionName, TypeNamer& typeNamer)
>+{
>+    StringBuilder stringBuilder;
>+    if (nativeFunctionDeclaration.isCast()) {
>+        ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
>+        auto metalParameterName = getNativeName(*nativeFunctionDeclaration.parameters()[0].type(), typeNamer);
>+        auto metalReturnName = getNativeName(nativeFunctionDeclaration.type(), typeNamer);
>+        if (metalParameterName != "atomic_int"_str && metalParameterName != "atomic_uint"_str) {
>+            stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, "x) {\n"));
>+            stringBuilder.append(makeString("    return static_cast<", metalReturnName, ">(x);\n"));
>+            stringBuilder.append("}\n");
>+            return stringBuilder.toString();
>+        }
>+
>+        stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, "x) {\n"));
>+        stringBuilder.append("    return atomic_load_explicit(&x, memory_order_relaxed);\n");
>+        stringBuilder.append("}\n");
>+        return stringBuilder.toString();
>+    }
>+
>+    if (nativeFunctionDeclaration.name().startsWith("operator."_str)) {
>+        if (nativeFunctionDeclaration.name().endsWith("="_str)) {
>+            ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
>+            auto metalParameter1Name = getNativeName(*nativeFunctionDeclaration.parameters()[0].type(), typeNamer);
>+            auto metalParameter2Name = getNativeName(*nativeFunctionDeclaration.parameters()[1].type(), typeNamer);
>+            auto metalReturnName = getNativeName(nativeFunctionDeclaration.type(), typeNamer);
>+            auto fieldName = nativeFunctionDeclaration.name().substring("operator."_str.length());
>+            fieldName = fieldName.substring(0, fieldName.length() - 1);
>+            stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, "v, ", metalParameter2Name, " n) {\n"));
>+            stringBuilder.append(makeString("    v.", fieldName, " = n;\n"));
>+            stringBuilder.append(makeString("    return v;\n"));
>+            stringBuilder.append("}\n");
>+            return stringBuilder.toString();
>+        }
>+
>+        ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
>+        auto metalParameterName = getNativeName(*nativeFunctionDeclaration.parameters()[0].type(), typeNamer);
>+        auto metalReturnName = getNativeName(nativeFunctionDeclaration.type(), typeNamer);
>+        stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, "v) {\n"));
>+        stringBuilder.append(makeString("    return v.", nativeFunctionDeclaration.name().substring("operator."_str.length()), ";\n"));
>+        stringBuilder.append("}\n");
>+        return stringBuilder.toString();
>+
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "operator[]") {
>+        ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
>+        auto metalParameter1Name = getNativeName(*nativeFunctionDeclaration.parameters()[0].type(), typeNamer);
>+        auto metalParameter2Name = getNativeName(*nativeFunctionDeclaration.parameters()[1].type(), typeNamer);
>+        auto metalReturnName = getNativeName(nativeFunctionDeclaration.type(), typeNamer);
>+        stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, "m, ", metalParameter2Name, " i) {\n"));
>+        stringBuilder.append(makeString("    return m[i];\n"));
>+        stringBuilder.append("}\n");
>+        return stringBuilder.toString();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "operator[]=") {
>+        ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
>+        auto metalParameter1Name = getNativeName(*nativeFunctionDeclaration.parameters()[0].type(), typeNamer);
>+        auto metalParameter2Name = getNativeName(*nativeFunctionDeclaration.parameters()[1].type(), typeNamer);
>+        auto metalParameter3Name = getNativeName(*nativeFunctionDeclaration.parameters()[2].type(), typeNamer);
>+        auto metalReturnName = getNativeName(nativeFunctionDeclaration.type(), typeNamer);
>+        stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, "m, ", metalParameter2Name, " i, ", metalParameter3Name, " v) {\n"));
>+        stringBuilder.append(makeString("    m[i] = v;\n"));
>+        stringBuilder.append(makeString("    return m;\n"));
>+        stringBuilder.append("}\n");
>+        return stringBuilder.toString();
>+    }
>+
>+    if (nativeFunctionDeclaration.isOperator()) {
>+        if (nativeFunctionDeclaration.parameters().size() == 1) {
>+            auto operatorName = nativeFunctionDeclaration.name().substring("operator"_str.length());
>+            auto metalParameterName = getNativeName(*nativeFunctionDeclaration.parameters()[0].type(), typeNamer);
>+            auto metalReturnName = getNativeName(nativeFunctionDeclaration.type(), typeNamer);
>+            stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, "x) {\n"));
>+            stringBuilder.append(makeString("    return ", operatorName, "x;\n"));
>+            stringBuilder.append("}\n");
>+            return stringBuilder.toString();
>+        }
>+
>+        ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
>+        auto operatorName = nativeFunctionDeclaration.name().substring("operator"_str.length());
>+        auto metalParameter1Name = getNativeName(*nativeFunctionDeclaration.parameters()[0].type(), typeNamer);
>+        auto metalParameter2Name = getNativeName(*nativeFunctionDeclaration.parameters()[1].type(), typeNamer);
>+        auto metalReturnName = getNativeName(nativeFunctionDeclaration.type(), typeNamer);
>+        stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, "x, ", metalParameter2Name, " y) {\n"));
>+        stringBuilder.append(makeString("    return x ", operatorName, " y;\n"));
>+        stringBuilder.append("}\n");
>+        return stringBuilder.toString();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "cos"
>+        || nativeFunctionDeclaration.name() == "sin"
>+        || nativeFunctionDeclaration.name() == "tan"
>+        || nativeFunctionDeclaration.name() == "acos"
>+        || nativeFunctionDeclaration.name() == "asin"
>+        || nativeFunctionDeclaration.name() == "atan"
>+        || nativeFunctionDeclaration.name() == "cosh"
>+        || nativeFunctionDeclaration.name() == "sinh"
>+        || nativeFunctionDeclaration.name() == "tanh"
>+        || nativeFunctionDeclaration.name() == "ceil"
>+        || nativeFunctionDeclaration.name() == "exp"
>+        || nativeFunctionDeclaration.name() == "floor"
>+        || nativeFunctionDeclaration.name() == "log"
>+        || nativeFunctionDeclaration.name() == "round"
>+        || nativeFunctionDeclaration.name() == "trunc"
>+        || nativeFunctionDeclaration.name() == "ddx"
>+        || nativeFunctionDeclaration.name() == "ddy"
>+        || nativeFunctionDeclaration.name() == "isnormal"
>+        || nativeFunctionDeclaration.name() == "isfinite"
>+        || nativeFunctionDeclaration.name() == "isinf"
>+        || nativeFunctionDeclaration.name() == "isnan"
>+        || nativeFunctionDeclaration.name() == "asint"
>+        || nativeFunctionDeclaration.name() == "asuint"
>+        || nativeFunctionDeclaration.name() == "asfloat") {
>+        ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
>+        auto metalParameterName = getNativeName(*nativeFunctionDeclaration.parameters()[0].type(), typeNamer);
>+        auto metalReturnName = getNativeName(nativeFunctionDeclaration.type(), typeNamer);
>+        stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, "x) {\n"));
>+        stringBuilder.append(makeString("    return ", mapFunctionName(nativeFunctionDeclaration.name()), "(x);\n"));
>+        stringBuilder.append("}\n");
>+        return stringBuilder.toString();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "pow" || nativeFunctionDeclaration.name() == "atan2") {
>+        ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
>+        auto metalParameter1Name = getNativeName(*nativeFunctionDeclaration.parameters()[0].type(), typeNamer);
>+        auto metalParameter2Name = getNativeName(*nativeFunctionDeclaration.parameters()[1].type(), typeNamer);
>+        auto metalReturnName = getNativeName(nativeFunctionDeclaration.type(), typeNamer);
>+        stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, "x, ", metalParameter2Name, " y) {\n"));
>+        stringBuilder.append(makeString("    return ", nativeFunctionDeclaration.name(), "(x, y);\n"));
>+        stringBuilder.append("}\n");
>+        return stringBuilder.toString();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "f16tof32" || nativeFunctionDeclaration.name() == "f32tof16") {
>+        // FIXME: Implement this
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "AllMemoryBarrierWithGroupSync") {
>+        ASSERT(!nativeFunctionDeclaration.parameters().size());
>+        stringBuilder.append(makeString("void ", outputFunctionName, "() {\n"));
>+        stringBuilder.append("    threadgroup_barrier(mem_flags::mem_device);\n");
>+        stringBuilder.append("    threadgroup_barrier(mem_flags::mem_threadgroup);\n");
>+        stringBuilder.append("    threadgroup_barrier(mem_flags::mem_texture);\n");
>+        stringBuilder.append("}\n");
>+        return stringBuilder.toString();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "DeviceMemoryBarrierWithGroupSync") {
>+        ASSERT(!nativeFunctionDeclaration.parameters().size());
>+        stringBuilder.append(makeString("void ", outputFunctionName, "() {\n"));
>+        stringBuilder.append("    threadgroup_barrier(mem_flags::mem_device);\n");
>+        stringBuilder.append("}\n");
>+        return stringBuilder.toString();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "GroupMemoryBarrierWithGroupSync") {
>+        ASSERT(!nativeFunctionDeclaration.parameters().size());
>+        stringBuilder.append(makeString("void ", outputFunctionName, "() {\n"));
>+        stringBuilder.append("    threadgroup_barrier(mem_flags::mem_threadgroup);\n");
>+        stringBuilder.append("}\n");
>+        return stringBuilder.toString();
>+    }
>+
>+    if (nativeFunctionDeclaration.name().startsWith("Interlocked"_str)) {
>+        if (nativeFunctionDeclaration.name() == "InterlockedCompareExchange") {
>+            ASSERT(nativeFunctionDeclaration.parameters().size() == 4);
>+            ASSERT(is<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0].type()));
>+            auto& firstArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0].type());
>+            auto firstArgumentAddressSpace = firstArgumentPointer.addressSpace();
>+            auto firstArgumentPointee = getNativeName(firstArgumentPointer.elementType(), typeNamer);
>+            auto secondArgument = getNativeName(*nativeFunctionDeclaration.parameters()[1].type(), typeNamer);
>+            auto thirdArgument = getNativeName(*nativeFunctionDeclaration.parameters()[2].type(), typeNamer);
>+            ASSERT(is<AST::PointerType>(*nativeFunctionDeclaration.parameters()[3].type()));
>+            auto& fourthArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[3].type());
>+            auto fourthArgumentAddressSpace = fourthArgumentPointer.addressSpace();
>+            auto fourthArgumentPointee = getNativeName(fourthArgumentPointer.elementType(), typeNamer);
>+            stringBuilder.append(makeString("void ", outputFunctionName, '(', convertAddressSpace(firstArgumentAddressSpace), ' ', firstArgumentPointee, "* object, ", secondArgument, " compare, ", thirdArgument, " desired, ", convertAddressSpace(fourthArgumentAddressSpace), ' ', fourthArgumentPointee, "* out) {\n"));
>+            stringBuilder.append("    atomic_compare_exchange_weak_explicit(object, &compare, desired, memory_order_relaxed);\n");
>+            stringBuilder.append("    *out = compare;\n");
>+            stringBuilder.append("}\n");
>+            return stringBuilder.toString();
>+        }
>+
>+        ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
>+        ASSERT(is<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0].type()));
>+        auto& firstArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0].type());
>+        auto firstArgumentAddressSpace = firstArgumentPointer.addressSpace();
>+        auto firstArgumentPointee = getNativeName(firstArgumentPointer.elementType(), typeNamer);
>+        auto secondArgument = getNativeName(*nativeFunctionDeclaration.parameters()[1].type(), typeNamer);
>+        ASSERT(is<AST::PointerType>(*nativeFunctionDeclaration.parameters()[2].type()));
>+        auto& thirdArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[2].type());
>+        auto thirdArgumentAddressSpace = thirdArgumentPointer.addressSpace();
>+        auto thirdArgumentPointee = getNativeName(thirdArgumentPointer.elementType(), typeNamer);
>+        auto name = ([](String input) -> String {
>+            if (input == "Add")
>+                return "fetch_add"_str;
>+            if (input == "And")
>+                return "fetch_and"_str;
>+            if (input == "Exchange")
>+                return "exchange"_str;
>+            if (input == "Max")
>+                return "fetch_max"_str;
>+            if (input == "Min")
>+                return "fetch_min"_str;
>+            if (input == "Or")
>+                return "fetch_or"_str;
>+            ASSERT(input == "Xor");
>+            return "fetch_xor"_str;
>+        })(nativeFunctionDeclaration.name().substring("Interlocked"_str.length()));
>+        stringBuilder.append(makeString("void ", outputFunctionName, '(', convertAddressSpace(firstArgumentAddressSpace), ' ', firstArgumentPointee, "* object, ", secondArgument, " operand, ", convertAddressSpace(thirdArgumentAddressSpace), ' ', thirdArgumentPointee, "* out) {\n"));
>+        stringBuilder.append(makeString("    *out = atomic_fetch_", name, "_explicit(object, operand, memory_order_relaxed);\n"));
>+        stringBuilder.append("}\n");
>+        return stringBuilder.toString();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "Sample") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "Load") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "GetDimensions") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "SampleBias") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "SampleGrad") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "SampleLevel") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "Gather") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "GatherRed") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "SampleCmp") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "SampleCmpLevelZero") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "Store") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "GatherAlpha") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "GatherBlue") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "GatherCmp") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "GatherCmpRed") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "GatherGreen") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    // FIXME: Add all the functions that the compiler generated.
>+
>+    ASSERT_NOT_REACHED();
>+    return String();
>+}
>+
>+} // namespace Metal
>+
>+} // namespace WHLSL
>+
>+} // namespace WebCore
>+
>+#endif
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.h b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.h
>new file mode 100644
>index 0000000000000000000000000000000000000000..2e9bd4e3cbd6c626101cd08a0b82a5efb6e655e1
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.h
>@@ -0,0 +1,54 @@
>+/*
>+ * 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/text/WTFString.h>
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+namespace AST {
>+
>+class NativeFunctionDeclaration;
>+
>+}
>+
>+namespace Metal {
>+
>+class TypeNamer;
>+
>+String writeNativeFunction(AST::NativeFunctionDeclaration&, String& outputFunctionName, TypeNamer&);
>+
>+}
>+
>+}
>+
>+}
>+
>+#endif
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLTypeNamer.cpp b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLTypeNamer.cpp
>new file mode 100644
>index 0000000000000000000000000000000000000000..2318eb8676abf3840ed13237fd7b48dd2517e00c
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLTypeNamer.cpp
>@@ -0,0 +1,706 @@
>+/*
>+ * 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 "WHLSLTypeNamer.h"
>+
>+#if ENABLE(WEBGPU)
>+
>+#include "WHLSLAddressSpace.h"
>+#include "WHLSLArrayReferenceType.h"
>+#include "WHLSLArrayType.h"
>+#include "WHLSLEnumerationDefinition.h"
>+#include "WHLSLEnumerationMember.h"
>+#include "WHLSLNativeTypeDeclaration.h"
>+#include "WHLSLPointerType.h"
>+#include "WHLSLStructureDefinition.h"
>+#include "WHLSLTypeDefinition.h"
>+#include "WHLSLTypeReference.h"
>+#include "WHLSLVisitor.h"
>+#include <algorithm>
>+#include <functional>
>+#include <wtf/HashMap.h>
>+#include <wtf/HashSet.h>
>+#include <wtf/Optional.h>
>+#include <wtf/UniqueRef.h>
>+#include <wtf/Vector.h>
>+#include <wtf/text/StringBuilder.h>
>+#include <wtf/text/StringConcatenateNumbers.h>
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+namespace Metal {
>+
>+class BaseTypeNameNode {
>+public:
>+    BaseTypeNameNode(BaseTypeNameNode* parent, String&& mangledName)
>+        : m_parent(parent)
>+        , m_mangledName(mangledName)
>+    {
>+    }
>+    virtual ~BaseTypeNameNode() = default;
>+    virtual bool isArrayTypeNameNode() const { return false; }
>+    virtual bool isArrayReferenceTypeNameNode() const { return false; }
>+    virtual bool isPointerTypeNameNode() const { return false; }
>+    virtual bool isReferenceTypeNameNode() const { return false; }
>+    Vector<UniqueRef<BaseTypeNameNode>>& children() { return m_children; }
>+    void append(UniqueRef<BaseTypeNameNode>&& child)
>+    {
>+        m_children.append(WTFMove(child));
>+    }
>+    BaseTypeNameNode* parent() { return m_parent; }
>+    const String& mangledName() const { return m_mangledName; }
>+
>+private:
>+    Vector<UniqueRef<BaseTypeNameNode>> m_children;
>+    BaseTypeNameNode* m_parent;
>+    String m_mangledName;
>+};
>+
>+class ArrayTypeNameNode : public BaseTypeNameNode {
>+public:
>+    ArrayTypeNameNode(BaseTypeNameNode* parent, String&& mangledName, unsigned numElements)
>+        : BaseTypeNameNode(parent, WTFMove(mangledName))
>+        , m_numElements(numElements)
>+    {
>+    }
>+    virtual ~ArrayTypeNameNode() = default;
>+    bool isArrayTypeNameNode() const override { return true; }
>+    unsigned numElements() const { return m_numElements; }
>+
>+private:
>+    unsigned m_numElements;
>+};
>+
>+class ArrayReferenceTypeNameNode : public BaseTypeNameNode {
>+public:
>+    ArrayReferenceTypeNameNode(BaseTypeNameNode* parent, String&& mangledName, AST::AddressSpace addressSpace)
>+        : BaseTypeNameNode(parent, WTFMove(mangledName))
>+        , m_addressSpace(addressSpace)
>+    {
>+    }
>+    virtual ~ArrayReferenceTypeNameNode() = default;
>+    bool isArrayReferenceTypeNameNode() const override { return true; }
>+    AST::AddressSpace addressSpace() const { return m_addressSpace; }
>+
>+private:
>+    AST::AddressSpace m_addressSpace;
>+};
>+
>+class PointerTypeNameNode : public BaseTypeNameNode {
>+public:
>+    PointerTypeNameNode(BaseTypeNameNode* parent, String&& mangledName, AST::AddressSpace addressSpace)
>+        : BaseTypeNameNode(parent, WTFMove(mangledName))
>+        , m_addressSpace(addressSpace)
>+    {
>+    }
>+    virtual ~PointerTypeNameNode() = default;
>+    bool isPointerTypeNameNode() const override { return true; }
>+    AST::AddressSpace addressSpace() const { return m_addressSpace; }
>+
>+private:
>+    AST::AddressSpace m_addressSpace;
>+};
>+
>+class ReferenceTypeNameNode : public BaseTypeNameNode {
>+public:
>+    ReferenceTypeNameNode(BaseTypeNameNode* parent, String&& mangledName, AST::NamedType& namedType)
>+        : BaseTypeNameNode(parent, WTFMove(mangledName))
>+        , m_namedType(namedType)
>+    {
>+    }
>+    virtual ~ReferenceTypeNameNode() = default;
>+    bool isReferenceTypeNameNode() const override { return true; }
>+    AST::NamedType& namedType() { return m_namedType; }
>+
>+private:
>+    AST::NamedType& m_namedType;
>+};
>+
>+}
>+
>+}
>+
>+}
>+
>+#define SPECIALIZE_TYPE_TRAITS_WHLSL_BASE_TYPE_NAMED_NODE(ToValueTypeName, predicate) \
>+SPECIALIZE_TYPE_TRAITS_BEGIN(WebCore::WHLSL::Metal::ToValueTypeName) \
>+    static bool isType(const WebCore::WHLSL::Metal::BaseTypeNameNode& type) { return type.predicate; } \
>+SPECIALIZE_TYPE_TRAITS_END()
>+
>+SPECIALIZE_TYPE_TRAITS_WHLSL_BASE_TYPE_NAMED_NODE(ArrayTypeNameNode, isArrayTypeNameNode())
>+
>+SPECIALIZE_TYPE_TRAITS_WHLSL_BASE_TYPE_NAMED_NODE(ArrayReferenceTypeNameNode, isArrayReferenceTypeNameNode())
>+
>+SPECIALIZE_TYPE_TRAITS_WHLSL_BASE_TYPE_NAMED_NODE(PointerTypeNameNode, isPointerTypeNameNode())
>+
>+SPECIALIZE_TYPE_TRAITS_WHLSL_BASE_TYPE_NAMED_NODE(ReferenceTypeNameNode, isReferenceTypeNameNode())
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+namespace Metal {
>+
>+TypeNamer::TypeNamer(Program& program)
>+    : m_program(program)
>+{
>+}
>+
>+TypeNamer::~TypeNamer() = default;
>+
>+void TypeNamer::visit(AST::UnnamedType& unnamedType)
>+{
>+    insert(unnamedType, m_trie);
>+}
>+
>+void TypeNamer::visit(AST::EnumerationDefinition& enumerationDefinition)
>+{
>+    auto addResult = m_namedTypeMapping.add(&enumerationDefinition, generateNextTypeName());
>+    ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+    for (auto& enumerationMember : enumerationDefinition.enumerationMembers()) {
>+        auto addResult = m_enumerationMemberMapping.add(&static_cast<AST::EnumerationMember&>(enumerationMember), generateNextEnumerationMemberName());
>+        ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+    }
>+    Visitor::visit(enumerationDefinition);
>+}
>+
>+void TypeNamer::visit(AST::NativeTypeDeclaration&)
>+{
>+    // Native type declarations already have names, and are already declared in Metal.
>+}
>+
>+static Vector<UniqueRef<BaseTypeNameNode>>::iterator findInVector(AST::UnnamedType& unnamedType, Vector<UniqueRef<BaseTypeNameNode>>& types)
>+{
>+    return std::find_if(types.begin(), types.end(), [&](BaseTypeNameNode& baseTypeNameNode) -> bool {
>+        if (is<AST::TypeReference>(unnamedType) && is<ReferenceTypeNameNode>(baseTypeNameNode)) {
>+            auto* resolvedType = downcast<AST::TypeReference>(unnamedType).resolvedType();
>+            ASSERT(resolvedType);
>+            return resolvedType == &downcast<ReferenceTypeNameNode>(baseTypeNameNode).namedType();
>+        }
>+        if (is<AST::PointerType>(unnamedType) && is<PointerTypeNameNode>(baseTypeNameNode))
>+            return downcast<AST::PointerType>(unnamedType).addressSpace() == downcast<PointerTypeNameNode>(baseTypeNameNode).addressSpace();
>+        if (is<AST::ArrayReferenceType>(unnamedType) && is<ArrayReferenceTypeNameNode>(baseTypeNameNode))
>+            return downcast<AST::ArrayReferenceType>(unnamedType).addressSpace() == downcast<ArrayReferenceTypeNameNode>(baseTypeNameNode).addressSpace();
>+        if (is<AST::ArrayType>(unnamedType) && is<ArrayTypeNameNode>(baseTypeNameNode))
>+            return downcast<AST::ArrayType>(unnamedType).numElements() == downcast<ArrayTypeNameNode>(baseTypeNameNode).numElements();
>+        return false;
>+    });
>+}
>+
>+static BaseTypeNameNode& find(AST::UnnamedType& unnamedType, Vector<UniqueRef<BaseTypeNameNode>>& types)
>+{
>+    auto& vectorToSearch = ([&]() -> Vector<UniqueRef<BaseTypeNameNode>>& {
>+        if (is<AST::TypeReference>(unnamedType))
>+            return types;
>+        if (is<AST::PointerType>(unnamedType))
>+            return find(downcast<AST::PointerType>(unnamedType).elementType(), types).children();
>+        if (is<AST::ArrayReferenceType>(unnamedType))
>+            return find(downcast<AST::ArrayReferenceType>(unnamedType).elementType(), types).children();
>+        ASSERT(is<AST::ArrayType>(unnamedType));
>+        return find(downcast<AST::ArrayType>(unnamedType).type(), types).children();
>+    })();
>+    auto iterator = findInVector(unnamedType, vectorToSearch);
>+    ASSERT(iterator != types.end());
>+    return *iterator;
>+}
>+
>+void TypeNamer::visit(AST::StructureDefinition& structureDefinition)
>+{
>+    {
>+        auto addResult = m_namedTypeMapping.add(&structureDefinition, generateNextTypeName());
>+        ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+    }
>+    Visitor::visit(structureDefinition);
>+    {
>+        Vector<std::reference_wrapper<BaseTypeNameNode>> neighbors;
>+        for (auto& structureElement : structureDefinition.structureElements()) {
>+            auto addResult = m_structureElementMapping.add(&structureElement, generateNextStructureElementName());
>+            ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+            neighbors.append(find(structureElement.type(), m_trie));
>+        }
>+        auto addResult = m_dependencyGraph.add(&structureDefinition, WTFMove(neighbors));
>+        ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+    }
>+}
>+
>+void TypeNamer::visit(AST::TypeDefinition& typeDefinition)
>+{
>+    {
>+        auto addResult = m_namedTypeMapping.add(&typeDefinition, generateNextTypeName());
>+        ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+    }
>+    Visitor::visit(typeDefinition);
>+    {
>+        Vector<std::reference_wrapper<BaseTypeNameNode>> neighbors = { find(typeDefinition.type(), m_trie) };
>+        auto addResult = m_dependencyGraph.add(&typeDefinition, WTFMove(neighbors));
>+        ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+    }
>+}
>+
>+String TypeNamer::mangledNameForType(AST::NativeTypeDeclaration& nativeTypeDeclaration)
>+{
>+    if (nativeTypeDeclaration.name() == "void")
>+        return "void"_str;
>+    if (nativeTypeDeclaration.name() == "bool")
>+        return "bool"_str;
>+    if (nativeTypeDeclaration.name() == "uchar")
>+        return "uint8_t"_str;
>+    if (nativeTypeDeclaration.name() == "ushort")
>+        return "uint16_t"_str;
>+    if (nativeTypeDeclaration.name() == "uint")
>+        return "uint32_t"_str;
>+    if (nativeTypeDeclaration.name() == "char")
>+        return "int8_t"_str;
>+    if (nativeTypeDeclaration.name() == "short")
>+        return "int16_t"_str;
>+    if (nativeTypeDeclaration.name() == "int")
>+        return "int32_t"_str;
>+    if (nativeTypeDeclaration.name() == "half")
>+        return "half"_str;
>+    if (nativeTypeDeclaration.name() == "float")
>+        return "float"_str;
>+    if (nativeTypeDeclaration.name() == "atomic_int")
>+        return "atomic_int"_str;
>+    if (nativeTypeDeclaration.name() == "atomic_uint")
>+        return "atomic_uint"_str;
>+    if (nativeTypeDeclaration.name() == "sampler")
>+        return "sampler"_str;
>+    if (nativeTypeDeclaration.name() == "vector") {
>+        ASSERT(nativeTypeDeclaration.typeArguments().size() == 2);
>+        ASSERT(WTF::holds_alternative<UniqueRef<AST::TypeReference>>(nativeTypeDeclaration.typeArguments()[0]));
>+        auto& typeReference = WTF::get<UniqueRef<AST::TypeReference>>(nativeTypeDeclaration.typeArguments()[0]);
>+        auto& unifyNode = typeReference->unifyNode();
>+        ASSERT(is<AST::NamedType>(unifyNode));
>+        auto& namedType = downcast<AST::NamedType>(unifyNode);
>+        ASSERT(is<AST::NativeTypeDeclaration>(namedType));
>+        auto& parameterType = downcast<AST::NativeTypeDeclaration>(namedType);
>+        auto prefix = ([&]() -> String {
>+            if (parameterType.name() == "bool")
>+                return "bool";
>+            if (parameterType.name() == "uchar")
>+                return "uchar";
>+            if (parameterType.name() == "ushort")
>+                return "ushort";
>+            if (parameterType.name() == "uint")
>+                return "uint";
>+            if (parameterType.name() == "char")
>+                return "char";
>+            if (parameterType.name() == "short")
>+                return "short";
>+            if (parameterType.name() == "int")
>+                return "int";
>+            if (parameterType.name() == "half")
>+                return "half";
>+            ASSERT(parameterType.name() == "float");
>+            return "float";
>+        })();
>+        ASSERT(WTF::holds_alternative<AST::ConstantExpression>(nativeTypeDeclaration.typeArguments()[1]));
>+        auto& constantExpression = WTF::get<AST::ConstantExpression>(nativeTypeDeclaration.typeArguments()[0]);
>+        auto& integerLiteral = constantExpression.integerLiteral();
>+        auto suffix = ([&]() -> String {
>+            switch (integerLiteral.value()) {
>+            case 2:
>+                return "2"_str;
>+            case 3:
>+                return "3"_str;
>+            default:
>+                ASSERT(integerLiteral.value() == 4);
>+                return "4"_str;
>+            }
>+        })();
>+        return makeString(prefix, suffix);
>+    }
>+    if (nativeTypeDeclaration.name() == "matrix") {
>+        ASSERT(nativeTypeDeclaration.typeArguments().size() == 3);
>+        ASSERT(WTF::holds_alternative<UniqueRef<AST::TypeReference>>(nativeTypeDeclaration.typeArguments()[0]));
>+        auto& typeReference = WTF::get<UniqueRef<AST::TypeReference>>(nativeTypeDeclaration.typeArguments()[0]);
>+        auto& unifyNode = typeReference->unifyNode();
>+        ASSERT(is<AST::NamedType>(unifyNode));
>+        auto& namedType = downcast<AST::NamedType>(unifyNode);
>+        ASSERT(is<AST::NativeTypeDeclaration>(namedType));
>+        auto& parameterType = downcast<AST::NativeTypeDeclaration>(namedType);
>+        auto prefix = ([&]() -> String {
>+            if (parameterType.name() == "half")
>+                return "half";
>+            ASSERT(parameterType.name() == "float");
>+            return "float";
>+        })();
>+        ASSERT(WTF::holds_alternative<AST::ConstantExpression>(nativeTypeDeclaration.typeArguments()[1]));
>+        auto& constantExpression1 = WTF::get<AST::ConstantExpression>(nativeTypeDeclaration.typeArguments()[0]);
>+        auto& integerLiteral1 = constantExpression1.integerLiteral();
>+        auto middle = ([&]() -> String {
>+            switch (integerLiteral1.value()) {
>+            case 2:
>+                return "2"_str;
>+            case 3:
>+                return "3"_str;
>+            default:
>+                ASSERT(integerLiteral1.value() == 4);
>+                return "4"_str;
>+            }
>+        })();
>+        ASSERT(WTF::holds_alternative<AST::ConstantExpression>(nativeTypeDeclaration.typeArguments()[2]));
>+        auto& constantExpression2 = WTF::get<AST::ConstantExpression>(nativeTypeDeclaration.typeArguments()[0]);
>+        auto& integerLiteral2 = constantExpression2.integerLiteral();
>+        auto suffix = ([&]() -> String {
>+            switch (integerLiteral2.value()) {
>+            case 2:
>+                return "2"_str;
>+            case 3:
>+                return "3"_str;
>+            default:
>+                ASSERT(integerLiteral2.value() == 4);
>+                return "4"_str;
>+            }
>+        })();
>+        return makeString(prefix, middle, 'x', suffix);
>+    }
>+    ASSERT(nativeTypeDeclaration.typeArguments().size() == 1);
>+    ASSERT(WTF::holds_alternative<UniqueRef<AST::TypeReference>>(nativeTypeDeclaration.typeArguments()[0]));
>+    auto& typeReference = WTF::get<UniqueRef<AST::TypeReference>>(nativeTypeDeclaration.typeArguments()[0]);
>+    auto prefix = ([&]() -> String {
>+        if (nativeTypeDeclaration.name() == "Texture1D")
>+            return "texture1d"_str;
>+        if (nativeTypeDeclaration.name() == "RWTexture1D")
>+            return "texture1d"_str;
>+        if (nativeTypeDeclaration.name() == "Texture1DArray")
>+            return "texture1d_array"_str;
>+        if (nativeTypeDeclaration.name() == "RWTexture1DArray")
>+            return "texture1d_array"_str;
>+        if (nativeTypeDeclaration.name() == "Texture2D")
>+            return "texture2d"_str;
>+        if (nativeTypeDeclaration.name() == "RWTexture2D")
>+            return "texture2d"_str;
>+        if (nativeTypeDeclaration.name() == "Texture2DArray")
>+            return "texture2d_array"_str;
>+        if (nativeTypeDeclaration.name() == "RWTexture2DArray")
>+            return "texture2d_array"_str;
>+        if (nativeTypeDeclaration.name() == "Texture3D")
>+            return "texture3d"_str;
>+        if (nativeTypeDeclaration.name() == "RWTexture3D")
>+            return "texture3d"_str;
>+        if (nativeTypeDeclaration.name() == "TextureCube")
>+            return "texturecube"_str;
>+        if (nativeTypeDeclaration.name() == "TextureDepth2D")
>+            return "depth2d"_str;
>+        if (nativeTypeDeclaration.name() == "RWTextureDepth2D")
>+            return "depth2d"_str;
>+        if (nativeTypeDeclaration.name() == "TextureDepth2DArray")
>+            return "depth2d_array"_str;
>+        if (nativeTypeDeclaration.name() == "RWTextureDepth2DArray")
>+            return "depth2d_array"_str;
>+        ASSERT(nativeTypeDeclaration.name() == "TextureDepthCube");
>+        return "depthcube"_str;
>+    })();
>+    auto innerType = ([&]() -> String {
>+        if (typeReference->name() == "ushort")
>+            return "ushort"_str;
>+        if (typeReference->name() == "ushort2")
>+            return "ushort"_str;
>+        if (typeReference->name() == "ushort3")
>+            return "ushort"_str;
>+        if (typeReference->name() == "ushort4")
>+            return "ushort"_str;
>+        if (typeReference->name() == "uint")
>+            return "uint"_str;
>+        if (typeReference->name() == "uint2")
>+            return "uint"_str;
>+        if (typeReference->name() == "uint3")
>+            return "uint"_str;
>+        if (typeReference->name() == "uint4")
>+            return "uint"_str;
>+        if (typeReference->name() == "short")
>+            return "short"_str;
>+        if (typeReference->name() == "short2")
>+            return "short"_str;
>+        if (typeReference->name() == "short3")
>+            return "short"_str;
>+        if (typeReference->name() == "short4")
>+            return "short"_str;
>+        if (typeReference->name() == "int")
>+            return "int"_str;
>+        if (typeReference->name() == "int2")
>+            return "int"_str;
>+        if (typeReference->name() == "int3")
>+            return "int"_str;
>+        if (typeReference->name() == "int4")
>+            return "int"_str;
>+        if (typeReference->name() == "half")
>+            return "half"_str;
>+        if (typeReference->name() == "half2")
>+            return "half"_str;
>+        if (typeReference->name() == "half3")
>+            return "half"_str;
>+        if (typeReference->name() == "half4")
>+            return "half"_str;
>+        if (typeReference->name() == "float")
>+            return "float"_str;
>+        if (typeReference->name() == "float2")
>+            return "float"_str;
>+        if (typeReference->name() == "float3")
>+            return "float"_str;
>+        ASSERT(typeReference->name() == "float4");
>+        return "float"_str;
>+    })();
>+    // FIXME: Specify the second template argument to Metal texture types.
>+    return makeString(prefix, '<', innerType, '>');
>+}
>+
>+UniqueRef<BaseTypeNameNode> TypeNamer::createNameNode(AST::UnnamedType& unnamedType, BaseTypeNameNode* parent)
>+{
>+    if (is<AST::TypeReference>(unnamedType)) {
>+        auto& typeReference = downcast<AST::TypeReference>(unnamedType);
>+        ASSERT(typeReference.resolvedType());
>+        return makeUniqueRef<ReferenceTypeNameNode>(parent, generateNextTypeName(), *typeReference.resolvedType());
>+    }
>+    if (is<AST::PointerType>(unnamedType)) {
>+        auto& pointerType = downcast<AST::PointerType>(unnamedType);
>+        return makeUniqueRef<PointerTypeNameNode>(parent, generateNextTypeName(), pointerType.addressSpace());
>+    }
>+    if (is<AST::ArrayReferenceType>(unnamedType)) {
>+        auto& arrayReferenceType = downcast<AST::ArrayReferenceType>(unnamedType);
>+        return makeUniqueRef<ArrayReferenceTypeNameNode>(parent, generateNextTypeName(), arrayReferenceType.addressSpace());
>+    }
>+    ASSERT(is<AST::ArrayType>(unnamedType));
>+    auto& arrayType = downcast<AST::ArrayType>(unnamedType);
>+    return makeUniqueRef<ArrayTypeNameNode>(parent, generateNextTypeName(), arrayType.numElements());
>+}
>+
>+size_t TypeNamer::insert(AST::UnnamedType& unnamedType, Vector<UniqueRef<BaseTypeNameNode>>& types)
>+{
>+    Vector<UniqueRef<BaseTypeNameNode>>* vectorToInsertInto { nullptr };
>+    BaseTypeNameNode* parent { nullptr };
>+    if (is<AST::TypeReference>(unnamedType)) {
>+        vectorToInsertInto = &types;
>+        parent = nullptr;
>+    } else if (is<AST::PointerType>(unnamedType)) {
>+        auto& item = types[insert(downcast<AST::PointerType>(unnamedType).elementType(), types)];
>+        vectorToInsertInto = &item->children();
>+        parent = &item;
>+    } else if (is<AST::ArrayReferenceType>(unnamedType)) {
>+        auto& item = types[insert(downcast<AST::ArrayReferenceType>(unnamedType).elementType(), types)];
>+        vectorToInsertInto = &item->children();
>+        parent = &item;
>+    } else {
>+        ASSERT(is<AST::ArrayType>(unnamedType));
>+        auto& item = types[insert(downcast<AST::ArrayType>(unnamedType).type(), types)];
>+        vectorToInsertInto = &item->children();
>+        parent = &item;
>+    }
>+    ASSERT(vectorToInsertInto);
>+
>+    auto iterator = findInVector(unnamedType, *vectorToInsertInto);
>+    if (iterator == vectorToInsertInto->end()) {
>+        auto result = createNameNode(unnamedType, parent);
>+        {
>+            auto addResult = m_unnamedTypeMapping.add(&unnamedType, &result);
>+            ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+        }
>+        vectorToInsertInto->append(WTFMove(result));
>+        return vectorToInsertInto->size() - 1;
>+    }
>+    auto addResult = m_unnamedTypeMapping.add(&unnamedType, &*iterator);
>+    ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+    return iterator - vectorToInsertInto->begin();
>+}
>+
>+class MetalTypeDeclarationWriter : public Visitor {
>+public:
>+    MetalTypeDeclarationWriter(std::function<String(AST::NamedType&)>&& mangledNameForNamedType, std::function<String(AST::UnnamedType&)>&& mangledNameForUnnamedType, std::function<String(AST::EnumerationMember&)>&& mangledNameForEnumerationMember)
>+        : m_mangledNameForNamedType(WTFMove(mangledNameForNamedType))
>+        , m_mangledNameForUnnamedType(WTFMove(mangledNameForUnnamedType))
>+        , m_mangledNameForEnumerationMember(WTFMove(mangledNameForEnumerationMember))
>+    {
>+    }
>+
>+    String toString() { return m_stringBuilder.toString(); }
>+
>+private:
>+    void visit(AST::EnumerationDefinition& enumerationDefinition)
>+    {
>+        auto& baseType = enumerationDefinition.type().unifyNode();
>+        ASSERT(is<AST::NamedType>(baseType));
>+        m_stringBuilder.append(makeString("enum class ", m_mangledNameForNamedType(enumerationDefinition), " : ", m_mangledNameForNamedType(downcast<AST::NamedType>(baseType)), " {\n"));
>+        for (auto& enumerationMember : enumerationDefinition.enumerationMembers())
>+            m_stringBuilder.append(makeString("    ", m_mangledNameForEnumerationMember(enumerationMember), ",\n"));
>+        m_stringBuilder.append("};\n");
>+    }
>+
>+    void visit(AST::StructureDefinition& structureDefinition)
>+    {
>+        m_stringBuilder.append(makeString("struct ", m_mangledNameForNamedType(structureDefinition), ";\n"));
>+    }
>+
>+    std::function<String(AST::NamedType&)> m_mangledNameForNamedType;
>+    std::function<String(AST::UnnamedType&)> m_mangledNameForUnnamedType;
>+    std::function<String(AST::EnumerationMember&)>&& m_mangledNameForEnumerationMember;
>+    StringBuilder m_stringBuilder;
>+};
>+
>+String TypeNamer::metalTypeDeclarations()
>+{
>+    MetalTypeDeclarationWriter metalTypeDeclarationWriter([&](AST::NamedType& namedType) -> String {
>+        return mangledNameForType(namedType);
>+    }, [&](AST::UnnamedType& unnamedType) -> String {
>+        return mangledNameForType(unnamedType);
>+    }, [&](AST::EnumerationMember& enumerationMember) -> String {
>+        return mangledNameForEnumerationMember(enumerationMember);
>+    });
>+    metalTypeDeclarationWriter.Visitor::visit(m_program);
>+    return metalTypeDeclarationWriter.toString();
>+}
>+
>+static String toString(AST::AddressSpace addressSpace)
>+{
>+    switch (addressSpace) {
>+    case AST::AddressSpace::Constant:
>+        return "constant"_str;
>+    case AST::AddressSpace::Device:
>+        return "device"_str;
>+    case AST::AddressSpace::Threadgroup:
>+        return "threadgroup"_str;
>+    default:
>+        ASSERT(addressSpace == AST::AddressSpace::Thread);
>+        return "thread"_str;
>+    }
>+}
>+
>+void TypeNamer::emitUnnamedTypeDefinition(BaseTypeNameNode& baseTypeNameNode, HashSet<AST::NamedType*>& emittedNamedTypes, HashSet<BaseTypeNameNode*>& emittedUnnamedTypes, StringBuilder& stringBuilder)
>+{
>+    if (emittedUnnamedTypes.contains(&baseTypeNameNode))
>+        return;
>+    if (baseTypeNameNode.parent())
>+        emitUnnamedTypeDefinition(*baseTypeNameNode.parent(), emittedNamedTypes, emittedUnnamedTypes, stringBuilder);
>+    if (is<ReferenceTypeNameNode>(baseTypeNameNode)) {
>+        auto& namedType = downcast<ReferenceTypeNameNode>(baseTypeNameNode).namedType();
>+        emitNamedTypeDefinition(namedType, emittedNamedTypes, emittedUnnamedTypes, stringBuilder);
>+        stringBuilder.append(makeString("typedef ", mangledNameForType(namedType), ' ', baseTypeNameNode.mangledName(), ";\n"));
>+    } else if (is<PointerTypeNameNode>(baseTypeNameNode)) {
>+        auto& pointerType = downcast<PointerTypeNameNode>(baseTypeNameNode);
>+        ASSERT(baseTypeNameNode.parent());
>+        stringBuilder.append(makeString("typedef ", toString(pointerType.addressSpace()), " ", pointerType.parent()->mangledName(), "* ", pointerType.mangledName(), ";\n"));
>+    } else if (is<ArrayReferenceTypeNameNode>(baseTypeNameNode)) {
>+        auto& arrayReferenceType = downcast<ArrayReferenceTypeNameNode>(baseTypeNameNode);
>+        ASSERT(baseTypeNameNode.parent());
>+        stringBuilder.append(makeString("struct ", arrayReferenceType.mangledName(), "{ \n"));
>+        stringBuilder.append(makeString("    ", toString(arrayReferenceType.addressSpace()), " ", arrayReferenceType.parent()->mangledName(), "* pointer;\n"));
>+        stringBuilder.append("    unsigned length;\n");
>+        stringBuilder.append("};\n");
>+    } else {
>+        ASSERT(is<ArrayTypeNameNode>(baseTypeNameNode));
>+        auto& arrayType = downcast<ArrayTypeNameNode>(baseTypeNameNode);
>+        ASSERT(baseTypeNameNode.parent());
>+        stringBuilder.append(makeString("typedef Array<", arrayType.parent()->mangledName(), ", ", arrayType.numElements(), "> ", arrayType.mangledName(), ";\n"));
>+    }
>+    emittedUnnamedTypes.add(&baseTypeNameNode);
>+}
>+
>+void TypeNamer::emitNamedTypeDefinition(AST::NamedType& namedType, HashSet<AST::NamedType*>& emittedNamedTypes, HashSet<BaseTypeNameNode*>& emittedUnnamedTypes, StringBuilder& stringBuilder)
>+{
>+    if (emittedNamedTypes.contains(&namedType))
>+        return;
>+    auto iterator = m_dependencyGraph.find(&namedType);
>+    ASSERT(iterator != m_dependencyGraph.end());
>+    for (auto& baseTypeNameNode : iterator->value)
>+        emitUnnamedTypeDefinition(baseTypeNameNode, emittedNamedTypes, emittedUnnamedTypes, stringBuilder);
>+    if (is<AST::EnumerationDefinition>(namedType)) {
>+        // We already emitted this in the type declaration section. There's nothing to do.
>+    } else if (is<AST::NativeTypeDeclaration>(namedType)) {
>+        // Native types already have definitions. There's nothing to do.
>+    } else if (is<AST::StructureDefinition>(namedType)) {
>+        auto& structureDefinition = downcast<AST::StructureDefinition>(namedType);
>+        stringBuilder.append(makeString("struct ", mangledNameForType(structureDefinition), " {\n"));
>+        for (auto& structureElement : structureDefinition.structureElements())
>+            stringBuilder.append(makeString("    ", mangledNameForType(structureElement.type()), ' ', mangledNameForStructureElement(structureElement), ";\n"));
>+        stringBuilder.append("};\n");
>+    } else {
>+        ASSERT(is<AST::TypeDefinition>(namedType));
>+        auto& typeDefinition = downcast<AST::TypeDefinition>(namedType);
>+        stringBuilder.append(makeString("typedef ", mangledNameForType(typeDefinition.type()), ' ', mangledNameForType(typeDefinition), ";\n"));
>+    }
>+    emittedNamedTypes.add(&namedType);
>+}
>+
>+String TypeNamer::metalTypeDefinitions()
>+{
>+    HashSet<AST::NamedType*> emittedNamedTypes;
>+    HashSet<BaseTypeNameNode*> emittedUnnamedTypes;
>+    StringBuilder stringBuilder;
>+    for (auto& keyValuePair : m_dependencyGraph)
>+        emitNamedTypeDefinition(*keyValuePair.key, emittedNamedTypes, emittedUnnamedTypes, stringBuilder);
>+    for (auto& baseTypeNameNode : m_trie)
>+        emitUnnamedTypeDefinition(baseTypeNameNode, emittedNamedTypes, emittedUnnamedTypes, stringBuilder);
>+    return stringBuilder.toString();
>+}
>+
>+String TypeNamer::mangledNameForType(AST::UnnamedType& unnamedType)
>+{
>+    return find(unnamedType, m_trie).mangledName();
>+}
>+
>+String TypeNamer::mangledNameForType(AST::NamedType& namedType)
>+{
>+    if (is<AST::NativeTypeDeclaration>(namedType))
>+        return mangledNameForType(downcast<AST::NativeTypeDeclaration>(namedType));
>+    auto iterator = m_namedTypeMapping.find(&namedType);
>+    ASSERT(iterator != m_namedTypeMapping.end());
>+    return iterator->value;
>+}
>+
>+
>+String TypeNamer::mangledNameForEnumerationMember(AST::EnumerationMember& enumerationMember)
>+{
>+    auto iterator = m_enumerationMemberMapping.find(&enumerationMember);
>+    ASSERT(iterator != m_enumerationMemberMapping.end());
>+    return iterator->value;
>+}
>+
>+String TypeNamer::mangledNameForStructureElement(AST::StructureElement& structureElement)
>+{
>+    auto iterator = m_structureElementMapping.find(&structureElement);
>+    ASSERT(iterator != m_structureElementMapping.end());
>+    return iterator->value;
>+}
>+
>+String TypeNamer::metalTypes()
>+{
>+    Visitor::visit(m_program);
>+    StringBuilder stringBuilder;
>+    stringBuilder.append(metalTypeDeclarations());
>+    stringBuilder.append('\n');
>+    stringBuilder.append(metalTypeDefinitions());
>+    return stringBuilder.toString();
>+}
>+
>+} // namespace Metal
>+
>+} // namespace WHLSL
>+
>+} // namespace WebCore
>+
>+#endif
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLTypeNamer.h b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLTypeNamer.h
>new file mode 100644
>index 0000000000000000000000000000000000000000..63cd5fe2bc4aa715b6bd3b270b36374306ff19d6
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLTypeNamer.h
>@@ -0,0 +1,120 @@
>+/*
>+ * 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 "WHLSLVisitor.h"
>+#include <wtf/HashMap.h>
>+#include <wtf/text/StringConcatenate.h>
>+#include <wtf/text/StringConcatenateNumbers.h>
>+#include <wtf/text/WTFString.h>
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+namespace AST {
>+
>+class EnumerationDefinition;
>+class EnumerationMember;
>+class NamedType;
>+class NativeTypeDeclaration;
>+class StructureDefinition;
>+class TypeDefinition;
>+class UnnamedType;
>+
>+}
>+
>+class Program;
>+
>+namespace Metal {
>+
>+class BaseTypeNameNode;
>+
>+class TypeNamer : private Visitor {
>+public:
>+    TypeNamer(Program&);
>+    virtual ~TypeNamer();
>+
>+    String metalTypes();
>+
>+    // Must be called after calling metalTypes().
>+    String mangledNameForType(AST::NativeTypeDeclaration&);
>+    String mangledNameForType(AST::UnnamedType&);
>+    String mangledNameForType(AST::NamedType&);
>+    String mangledNameForEnumerationMember(AST::EnumerationMember&);
>+    String mangledNameForStructureElement(AST::StructureElement&);
>+
>+    String generateNextTypeName()
>+    {
>+        return makeString("type", m_typeCount++);
>+    }
>+
>+    String generateNextStructureElementName()
>+    {
>+        return makeString("structureElement", m_structureElementCount++);
>+    }
>+
>+private:
>+    void visit(AST::UnnamedType&) override;
>+    void visit(AST::EnumerationDefinition&) override;
>+    void visit(AST::NativeTypeDeclaration&) override;
>+    void visit(AST::StructureDefinition&) override;
>+    void visit(AST::TypeDefinition&) override;
>+
>+    String generateNextEnumerationMemberName()
>+    {
>+        return makeString("enumerationMember", m_enumerationMemberCount++);
>+    }
>+
>+    void emitNamedTypeDefinition(AST::NamedType&, HashSet<AST::NamedType*>& emittedNamedTypes, HashSet<BaseTypeNameNode*>& emittedUnnamedTypes, StringBuilder&);
>+    void emitUnnamedTypeDefinition(BaseTypeNameNode&, HashSet<AST::NamedType*>& emittedNamedTypes, HashSet<BaseTypeNameNode*>& emittedUnnamedTypes, StringBuilder&);
>+    String metalTypeDeclarations();
>+    String metalTypeDefinitions();
>+
>+    UniqueRef<BaseTypeNameNode> createNameNode(AST::UnnamedType&, BaseTypeNameNode* parent);
>+    size_t insert(AST::UnnamedType&, Vector<UniqueRef<BaseTypeNameNode>>&);
>+
>+    Program& m_program;
>+    Vector<UniqueRef<BaseTypeNameNode>> m_trie;
>+    HashMap<AST::UnnamedType*, BaseTypeNameNode*> m_unnamedTypeMapping;
>+    HashMap<AST::NamedType*, String> m_namedTypeMapping;
>+    HashMap<AST::NamedType*, Vector<std::reference_wrapper<BaseTypeNameNode>>> m_dependencyGraph;
>+    HashMap<AST::EnumerationMember*, String> m_enumerationMemberMapping;
>+    HashMap<AST::StructureElement*, String> m_structureElementMapping;
>+    unsigned m_typeCount { 0 };
>+    unsigned m_enumerationMemberCount { 0 };
>+    unsigned m_structureElementCount { 0 };
>+};
>+
>+} // namespace Metal
>+
>+} // namespace WHLSL
>+
>+} // namespace WebCore
>+
>+#endif
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLChecker.cpp b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLChecker.cpp
>index 60b0efeecac2463e51a8ec123dccf7a20a46f9e6..294dc98be3037a078136850d85e32650cd4146fd 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLChecker.cpp
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLChecker.cpp
>@@ -221,8 +221,8 @@ static bool checkSemantics(Vector<EntryPointItem>& inputItems, Vector<EntryPoint
>         auto checkSemanticTypes = [&](const Vector<EntryPointItem>& items) -> bool {
>             for (auto& item : items) {
>                 auto acceptable = WTF::visit(WTF::makeVisitor([&](const AST::BaseSemantic& semantic) -> bool {
>-                    return semantic.isAcceptableType(item.unnamedType, intrinsics);
>-                }), item.semantic);
>+                    return semantic.isAcceptableType(*item.unnamedType, intrinsics);
>+                }), *item.semantic);
>                 if (!acceptable)
>                     return false;
>             }
>@@ -239,7 +239,7 @@ static bool checkSemantics(Vector<EntryPointItem>& inputItems, Vector<EntryPoint
>             for (auto& item : items) {
>                 auto acceptable = WTF::visit(WTF::makeVisitor([&](const AST::BaseSemantic& semantic) -> bool {
>                     return semantic.isAcceptableForShaderItemDirection(direction, entryPointType);
>-                }), item.semantic);
>+                }), *item.semantic);
>                 if (!acceptable)
>                     return false;
>             }
>@@ -256,11 +256,11 @@ static bool checkSemantics(Vector<EntryPointItem>& inputItems, Vector<EntryPoint
>             for (auto& item : items) {
>                 PODChecker podChecker;
>                 if (is<AST::PointerType>(item.unnamedType))
>-                    podChecker.checkErrorAndVisit(downcast<AST::PointerType>(item.unnamedType).elementType());
>+                    podChecker.checkErrorAndVisit(downcast<AST::PointerType>(*item.unnamedType).elementType());
>                 else if (is<AST::ArrayReferenceType>(item.unnamedType))
>-                    podChecker.checkErrorAndVisit(downcast<AST::ArrayReferenceType>(item.unnamedType).elementType());
>+                    podChecker.checkErrorAndVisit(downcast<AST::ArrayReferenceType>(*item.unnamedType).elementType());
>                 else if (is<AST::ArrayType>(item.unnamedType))
>-                    podChecker.checkErrorAndVisit(downcast<AST::ArrayType>(item.unnamedType).type());
>+                    podChecker.checkErrorAndVisit(downcast<AST::ArrayType>(*item.unnamedType).type());
>                 else
>                     continue;
>                 if (podChecker.error())
>@@ -612,16 +612,12 @@ void Checker::visit(AST::EnumerationDefinition& enumerationDefinition)
>     auto* baseType = ([&]() -> AST::NativeTypeDeclaration* {
>         checkErrorAndVisit(enumerationDefinition.type());
>         auto& baseType = enumerationDefinition.type().unifyNode();
>-        if (!is<AST::UnnamedType>(baseType))
>+        if (!is<AST::NamedType>(baseType))
>             return nullptr;
>-        auto& unnamedBase = downcast<AST::UnnamedType>(baseType);
>-        if (!is<AST::TypeReference>(unnamedBase))
>+        auto& namedType = downcast<AST::NamedType>(baseType);
>+        if (!is<AST::NativeTypeDeclaration>(namedType))
>             return nullptr;
>-        auto& typeReferenceBase = downcast<AST::TypeReference>(unnamedBase);
>-        ASSERT(typeReferenceBase.resolvedType());
>-        if (!is<AST::NativeTypeDeclaration>(*typeReferenceBase.resolvedType()))
>-            return nullptr;
>-        auto& nativeTypeDeclaration = downcast<AST::NativeTypeDeclaration>(*typeReferenceBase.resolvedType());
>+        auto& nativeTypeDeclaration = downcast<AST::NativeTypeDeclaration>(namedType);
>         if (!nativeTypeDeclaration.isInt())
>             return nullptr;
>         return &nativeTypeDeclaration;
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.cpp b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.cpp
>index a18eac38ea3b676ae789be595a3daa813c6f29b6..45a62a6d288e3be7d2fc33508551787c20d3197d 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.cpp
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.cpp
>@@ -71,7 +71,7 @@ public:
>             setError();
>             return;
>         }
>-        m_entryPointItems.append(EntryPointItem(m_typeReferences.last().get(), *m_currentSemantic));
>+        m_entryPointItems.append(EntryPointItem(m_typeReferences.last().get(), *m_currentSemantic, m_path));
>     }
> 
>     void visit(AST::NativeTypeDeclaration& nativeTypeDeclaration)
>@@ -85,7 +85,7 @@ public:
>             return;
>         }
> 
>-        m_entryPointItems.append(EntryPointItem(m_typeReferences.last().get(), *m_currentSemantic));
>+        m_entryPointItems.append(EntryPointItem(m_typeReferences.last().get(), *m_currentSemantic, m_path));
>     }
> 
>     void visit(AST::StructureDefinition& structureDefinition)
>@@ -98,7 +98,9 @@ public:
>         for (auto& structureElement : structureDefinition.structureElements()) {
>             if (structureElement.semantic())
>                 m_currentSemantic = &*structureElement.semantic();
>+            m_path.append(structureElement.name());
>             checkErrorAndVisit(structureElement);
>+            m_path.takeLast();
>         }
>     }
> 
>@@ -122,7 +124,7 @@ public:
>             setError();
>             return;
>         }
>-        m_entryPointItems.append(EntryPointItem(pointerType, *m_currentSemantic));
>+        m_entryPointItems.append(EntryPointItem(pointerType, *m_currentSemantic, m_path));
>     }
> 
>     void visit(AST::ArrayReferenceType& arrayReferenceType)
>@@ -131,7 +133,7 @@ public:
>             setError();
>             return;
>         }
>-        m_entryPointItems.append(EntryPointItem(arrayReferenceType, *m_currentSemantic));
>+        m_entryPointItems.append(EntryPointItem(arrayReferenceType, *m_currentSemantic, m_path));
>     }
> 
>     void visit(AST::ArrayType& arrayType)
>@@ -140,7 +142,7 @@ public:
>             setError();
>             return;
>         }
>-        m_entryPointItems.append(EntryPointItem(arrayType, *m_currentSemantic));
>+        m_entryPointItems.append(EntryPointItem(arrayType, *m_currentSemantic, m_path));
>     }
> 
>     void visit(AST::VariableDeclaration& variableDeclaration)
>@@ -148,11 +150,14 @@ public:
>         ASSERT(!m_currentSemantic);
>         if (variableDeclaration.semantic())
>             m_currentSemantic = &*variableDeclaration.semantic();
>-        if (variableDeclaration.type())
>-            checkErrorAndVisit(*variableDeclaration.type());
>+        ASSERT(variableDeclaration.type());
>+        m_path.append(variableDeclaration.name());
>+        checkErrorAndVisit(*variableDeclaration.type());
>+        m_path.takeLast();
>     }
> 
> private:
>+    Vector<String> m_path;
>     const Intrinsics& m_intrinsics;
>     AST::Semantic* m_currentSemantic { nullptr };
>     Vector<std::reference_wrapper<AST::TypeReference>> m_typeReferences;
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.h b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.h
>index 20367be3e2953c6bed0e1b7ff4e7849448ebbdff..6664e2e342efd931bc6b04b0e7323fb8b8f2e0cb 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.h
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.h
>@@ -30,6 +30,7 @@
> #include "WHLSLSemantic.h"
> #include <wtf/Optional.h>
> #include <wtf/Vector.h>
>+#include <wtf/text/WTFString.h>
> 
> namespace WebCore {
> 
>@@ -45,14 +46,16 @@ class Intrinsics;
> class UnnamedType;
> 
> struct EntryPointItem {
>-    EntryPointItem(AST::UnnamedType& unnamedType, AST::Semantic& semantic)
>-        : unnamedType(unnamedType)
>-        , semantic(semantic)
>+    EntryPointItem(AST::UnnamedType& unnamedType, AST::Semantic& semantic, Vector<String>& path)
>+        : unnamedType(&unnamedType)
>+        , semantic(&semantic)
>+        , path(path)
>     {
>     }
> 
>-    AST::UnnamedType& unnamedType;
>-    AST::Semantic& semantic;
>+    AST::UnnamedType* unnamedType;
>+    AST::Semantic* semantic;
>+    Vector<String> path;
> };
> 
> struct EntryPointItems {
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLLoopChecker.cpp b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLLoopChecker.cpp
>deleted file mode 100644
>index ae6850b9b2297f59a566912306cc877951c32a86..0000000000000000000000000000000000000000
>--- a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLLoopChecker.cpp
>+++ /dev/null
>@@ -1,149 +0,0 @@
>-/*
>- * 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 "WHLSLLoopChecker.h"
>-
>-#if ENABLE(WEBGPU)
>-
>-#include "WHLSLBreak.h"
>-#include "WHLSLContinue.h"
>-#include "WHLSLDoWhileLoop.h"
>-#include "WHLSLFallthrough.h"
>-#include "WHLSLForLoop.h"
>-#include "WHLSLSwitchStatement.h"
>-#include "WHLSLVisitor.h"
>-#include "WHLSLWhileLoop.h"
>-#include <wtf/SetForScope.h>
>-
>-namespace WebCore {
>-
>-namespace WHLSL {
>-
>-// This makes sure that every occurance of "continue," "break," and "fallthrough" appear within the relevant control flow structure.
>-class LoopChecker : public Visitor {
>-public:
>-    LoopChecker() = default;
>-
>-    virtual ~LoopChecker() = default;
>-
>-private:
>-    void visit(AST::FunctionDefinition& functionDefinition) override
>-    {
>-        ASSERT(!m_loopDepth);
>-        ASSERT(!m_switchDepth);
>-        Visitor::visit(functionDefinition);
>-    }
>-
>-    void visit(AST::WhileLoop& whileLoop) override
>-    {
>-        checkErrorAndVisit(whileLoop.conditional());
>-
>-        SetForScope<unsigned> change(m_loopDepth, m_loopDepth + 1);
>-        checkErrorAndVisit(whileLoop.body());
>-        ASSERT(m_loopDepth);
>-    }
>-
>-    void visit(AST::DoWhileLoop& doWhileLoop) override
>-    {
>-        {
>-            SetForScope<unsigned> change(m_loopDepth, m_loopDepth + 1);
>-            checkErrorAndVisit(doWhileLoop.body());
>-            ASSERT(m_loopDepth);
>-        }
>-
>-        checkErrorAndVisit(doWhileLoop.conditional());
>-    }
>-
>-    void visit(AST::ForLoop& forLoop) override
>-    {
>-        WTF::visit(WTF::makeVisitor([&](AST::VariableDeclarationsStatement& variableDeclarationsStatement) {
>-            checkErrorAndVisit(variableDeclarationsStatement);
>-        }, [&](UniqueRef<AST::Expression>& expression) {
>-            checkErrorAndVisit(expression);
>-        }), forLoop.initialization());
>-        if (forLoop.condition())
>-            checkErrorAndVisit(*forLoop.condition());
>-        if (forLoop.increment())
>-            checkErrorAndVisit(*forLoop.increment());
>-
>-        SetForScope<unsigned> change(m_loopDepth, m_loopDepth + 1);
>-        checkErrorAndVisit(forLoop.body());
>-        ASSERT(m_loopDepth);
>-    }
>-
>-    void visit(AST::SwitchStatement& switchStatement) override
>-    {
>-        checkErrorAndVisit(switchStatement.value());
>-
>-        SetForScope<unsigned> change(m_switchDepth, m_switchDepth + 1);
>-        for (auto& switchCase : switchStatement.switchCases())
>-            checkErrorAndVisit(switchCase);
>-        ASSERT(m_switchDepth);
>-    }
>-
>-    void visit(AST::Break& breakStatement) override
>-    {
>-        if (!m_loopDepth && !m_switchDepth) {
>-            setError();
>-            return;
>-        }
>-        Visitor::visit(breakStatement);
>-    }
>-
>-    void visit(AST::Continue& continueStatement) override
>-    {
>-        if (!m_loopDepth) {
>-            setError();
>-            return;
>-        }
>-        Visitor::visit(continueStatement);
>-    }
>-
>-    void visit(AST::Fallthrough& fallthroughStatement) override
>-    {
>-        if (!m_switchDepth) {
>-            setError();
>-            return;
>-        }
>-        Visitor::visit(fallthroughStatement);
>-    }
>-
>-    unsigned m_loopDepth { 0 };
>-    unsigned m_switchDepth { 0 };
>-};
>-
>-bool checkLoops(Program& program)
>-{
>-    LoopChecker loopChecker;
>-    loopChecker.Visitor::visit(program);
>-    return !loopChecker.error();
>-}
>-
>-} // namespace WHLSL
>-
>-} // namespace WebCore
>-
>-#endif // ENABLE(WEBGPU)
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLLoopChecker.h b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLLoopChecker.h
>deleted file mode 100644
>index 1fcac67eb969700554d0a1e7033eae2966e12f96..0000000000000000000000000000000000000000
>--- a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLLoopChecker.h
>+++ /dev/null
>@@ -1,42 +0,0 @@
>-/*
>- * 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)
>-
>-namespace WebCore {
>-
>-namespace WHLSL {
>-
>-class Program;
>-
>-bool checkLoops(Program&);
>-
>-}
>-
>-}
>-
>-#endif
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLStatementBehaviorChecker.cpp b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLStatementBehaviorChecker.cpp
>new file mode 100644
>index 0000000000000000000000000000000000000000..864143012be5888c8cf558012301ed5ce6c19598
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLStatementBehaviorChecker.cpp
>@@ -0,0 +1,243 @@
>+/*
>+ * 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 "WHLSLStatementBehaviorChecker.h"
>+
>+#if ENABLE(WEBGPU)
>+
>+#include "WHLSLDoWhileLoop.h"
>+#include "WHLSLForLoop.h"
>+#include "WHLSLIfStatement.h"
>+#include "WHLSLInferTypes.h"
>+#include "WHLSLProgram.h"
>+#include "WHLSLSwitchStatement.h"
>+#include "WHLSLWhileLoop.h"
>+#include <cstdint>
>+#include <wtf/OptionSet.h>
>+#include <wtf/Vector.h>
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+class StatementBehaviorChecker : public Visitor {
>+public:
>+    enum class Behavior : uint8_t {
>+        Return = 1 << 0,
>+        Break = 1 << 1,
>+        Continue = 1 << 2,
>+        Fallthrough = 1 << 3,
>+        Nothing = 1 << 4
>+    };
>+
>+    OptionSet<Behavior> takeFunctionBehavior()
>+    {
>+        ASSERT(m_stack.size() == 1);
>+        return m_stack.takeLast();
>+    }
>+
>+private:
>+    void visit(AST::Break&) override
>+    {
>+        m_stack.append({ Behavior::Break });
>+    }
>+
>+    void visit(AST::Fallthrough&) override
>+    {
>+        m_stack.append({ Behavior::Fallthrough });
>+    }
>+
>+    void visit(AST::Continue&) override
>+    {
>+        m_stack.append({ Behavior::Continue });
>+    }
>+
>+    void visit(AST::Return&) override
>+    {
>+        m_stack.append({ Behavior::Return });
>+    }
>+
>+    void visit(AST::Trap&) override
>+    {
>+        m_stack.append({ Behavior::Return });
>+    }
>+
>+    void visit(AST::DoWhileLoop& doWhileLoop) override
>+    {
>+        checkErrorAndVisit(doWhileLoop.body());
>+        if (error())
>+            return;
>+        auto b = m_stack.takeLast();
>+        b.remove(Behavior::Break);
>+        b.remove(Behavior::Continue);
>+        b.add(Behavior::Nothing);
>+        m_stack.append(b);
>+    }
>+
>+    void visit(AST::ForLoop& forLoop) override
>+    {
>+        // The initialization always has a behavior of Nothing, which we already add unconditionally below,
>+        // so we can just ignore the initialization.
>+
>+        checkErrorAndVisit(forLoop.body());
>+        if (error())
>+            return;
>+        auto b = m_stack.takeLast();
>+        b.remove(Behavior::Break);
>+        b.remove(Behavior::Continue);
>+        b.add(Behavior::Nothing);
>+        m_stack.append(b);
>+    }
>+
>+    void visit(AST::WhileLoop& whileLoop) override
>+    {
>+        checkErrorAndVisit(whileLoop.body());
>+        if (error())
>+            return;
>+        auto b = m_stack.takeLast();
>+        b.remove(Behavior::Break);
>+        b.remove(Behavior::Continue);
>+        b.add(Behavior::Nothing);
>+        m_stack.append(b);
>+    }
>+
>+    void visit(AST::SwitchCase& switchCase) override
>+    {
>+        Visitor::visit(switchCase);
>+    }
>+
>+    void visit(AST::SwitchStatement& switchStatement) override
>+    {
>+        if (switchStatement.switchCases().isEmpty()) {
>+            m_stack.append({ Behavior::Nothing });
>+            return;
>+        }
>+
>+        OptionSet<Behavior> reduction = { };
>+        for (auto& switchCase : switchStatement.switchCases()) {
>+            checkErrorAndVisit(switchCase);
>+            if (error())
>+                return;
>+            auto b = m_stack.takeLast();
>+            reduction = OptionSet<Behavior>::fromRaw(reduction.toRaw() | b.toRaw());
>+        }
>+        if (reduction.contains(Behavior::Nothing)) {
>+            setError();
>+            return;
>+        }
>+        reduction.remove(Behavior::Break);
>+        reduction.remove(Behavior::Fallthrough);
>+        reduction.add(Behavior::Nothing);
>+        m_stack.append(reduction);
>+    }
>+
>+    void visit(AST::IfStatement& ifStatement) override
>+    {
>+        checkErrorAndVisit(ifStatement.body());
>+        if (error())
>+            return;
>+        auto b = m_stack.takeLast();
>+        OptionSet<Behavior> bPrime;
>+        if (ifStatement.elseBody()) {
>+            checkErrorAndVisit(*ifStatement.elseBody());
>+            if (error())
>+                return;
>+            bPrime = m_stack.takeLast();
>+        } else
>+            bPrime = { Behavior::Nothing };
>+        m_stack.append(OptionSet<Behavior>::fromRaw(b.toRaw() | bPrime.toRaw()));
>+    }
>+
>+    void visit(AST::EffectfulExpressionStatement&) override
>+    {
>+        m_stack.append({ Behavior::Nothing });
>+    }
>+
>+    void visit(AST::Block& block) override
>+    {
>+        if (block.statements().isEmpty()) {
>+            m_stack.append({ Behavior::Nothing });
>+            return;
>+        }
>+
>+        OptionSet<Behavior> reduction = { };
>+        for (size_t i = 0; i < block.statements().size() - 1; ++i) {
>+            checkErrorAndVisit(block.statements()[i]);
>+            if (error())
>+                return;
>+            auto b = m_stack.takeLast();
>+            if (!b.contains(Behavior::Nothing)) {
>+                setError();
>+                return;
>+            }
>+            b.remove(Behavior::Nothing);
>+            if (b.contains(Behavior::Fallthrough)) {
>+                setError();
>+                return;
>+            }
>+            reduction = OptionSet<Behavior>::fromRaw(reduction.toRaw() | b.toRaw());
>+        }
>+        checkErrorAndVisit(block.statements()[block.statements().size() - 1]);
>+        if (error())
>+            return;
>+        auto b = m_stack.takeLast();
>+        m_stack.append(OptionSet<Behavior>::fromRaw(reduction.toRaw() | b.toRaw()));
>+    }
>+
>+    void visit(AST::VariableDeclarationsStatement&) override
>+    {
>+        m_stack.append({ Behavior::Nothing });
>+    }
>+
>+    Vector<OptionSet<Behavior>> m_stack;
>+};
>+
>+bool checkStatementBehavior(Program& program)
>+{
>+    StatementBehaviorChecker statementBehaviorChecker;
>+    for (auto& functionDefinition : program.functionDefinitions()) {
>+        statementBehaviorChecker.Visitor::visit(functionDefinition);
>+        if (statementBehaviorChecker.error())
>+            return false;
>+        auto behavior = statementBehaviorChecker.takeFunctionBehavior();
>+        if (matches(functionDefinition->type(), program.intrinsics().voidType())) {
>+            behavior.remove(StatementBehaviorChecker::Behavior::Return);
>+            behavior.remove(StatementBehaviorChecker::Behavior::Nothing);
>+            if (behavior != OptionSet<StatementBehaviorChecker::Behavior>())
>+                return false;
>+        } else {
>+            if (behavior != StatementBehaviorChecker::Behavior::Return)
>+                return false;
>+        }
>+    }
>+    return true;
>+}
>+
>+} // namespace WHLSL
>+
>+} // namespace WebCore
>+
>+#endif // ENABLE(WEBGPU)
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLStatementBehaviorChecker.h b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLStatementBehaviorChecker.h
>new file mode 100644
>index 0000000000000000000000000000000000000000..aecdfe61fa802a2b44543a016b1d04e26e74b3ec
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLStatementBehaviorChecker.h
>@@ -0,0 +1,42 @@
>+/*
>+ * 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)
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+class Program;
>+
>+bool checkStatementBehavior(Program&);
>+
>+}
>+
>+}
>+
>+#endif
>diff --git a/Source/WebCore/Sources.txt b/Source/WebCore/Sources.txt
>index b6124192ce94452f6d1fe4bbf309fcdac41b472a..649dd7e0d6144b5f4b5f933545316d685d056287 100644
>--- a/Source/WebCore/Sources.txt
>+++ b/Source/WebCore/Sources.txt
>@@ -318,6 +318,7 @@ Modules/webgpu/WHLSL/WHLSLSynthesizeConstructors.cpp
> Modules/webgpu/WHLSL/WHLSLSynthesizeEnumerationFunctions.cpp
> Modules/webgpu/WHLSL/WHLSLSynthesizeStructureAccessors.cpp
> Modules/webgpu/WHLSL/WHLSLIntrinsics.cpp
>+Modules/webgpu/WHLSL/WHLSLStatementBehaviorChecker.cpp
> Modules/webgpu/WHLSL/WHLSLNameContext.cpp
> Modules/webgpu/WHLSL/WHLSLNameResolver.cpp
> Modules/webgpu/WHLSL/WHLSLResolveOverloadImpl.cpp
>@@ -325,7 +326,11 @@ Modules/webgpu/WHLSL/WHLSLRecursionChecker.cpp
> Modules/webgpu/WHLSL/WHLSLVisitor.cpp
> Modules/webgpu/WHLSL/WHLSLLiteralTypeChecker.cpp
> Modules/webgpu/WHLSL/WHLSLHighZombieFinder.cpp
>-Modules/webgpu/WHLSL/WHLSLLoopChecker.cpp
>+Modules/webgpu/WHLSL/Metal/WHLSLTypeNamer.cpp
>+Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp
>+Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.cpp
>+Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.cpp
>+Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp
> Modules/webgpu/WHLSL/WHLSLFunctionStageChecker.cpp
> Modules/webgpu/WHLSL/AST/WHLSLTypeArgument.cpp
> Modules/webgpu/WHLSL/AST/WHLSLBuiltInSemantic.cpp
>diff --git a/Source/WebCore/WebCore.xcodeproj/project.pbxproj b/Source/WebCore/WebCore.xcodeproj/project.pbxproj
>index 6fedc2224f8b15f25d01338710a06961f63564d8..0ec5af5dfb503f4ce2aa505e91417240f97bc9a2 100644
>--- a/Source/WebCore/WebCore.xcodeproj/project.pbxproj
>+++ b/Source/WebCore/WebCore.xcodeproj/project.pbxproj
>@@ -6437,8 +6437,6 @@
> 		1CA0C2E521EED12A00A11860 /* WHLSLFunctionStageChecker.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLFunctionStageChecker.h; sourceTree = "<group>"; };
> 		1CA0C2EA21EED6F500A11860 /* WHLSLLiteralTypeChecker.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLLiteralTypeChecker.h; sourceTree = "<group>"; };
> 		1CA0C2EC21EED6F600A11860 /* WHLSLLiteralTypeChecker.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLLiteralTypeChecker.cpp; sourceTree = "<group>"; };
>-		1CA0C2F421EEDAD000A11860 /* WHLSLLoopChecker.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLLoopChecker.cpp; sourceTree = "<group>"; };
>-		1CA0C2F521EEDAD100A11860 /* WHLSLLoopChecker.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLLoopChecker.h; sourceTree = "<group>"; };
> 		1CA19E030DC255950065A994 /* EventLoopMac.mm */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.objcpp; path = EventLoopMac.mm; sourceTree = "<group>"; };
> 		1CA19E150DC255CA0065A994 /* EventLoop.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = EventLoop.h; sourceTree = "<group>"; };
> 		1CAF347E0A6C405200ABE06E /* WebScriptObject.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = WebScriptObject.h; sourceTree = "<group>"; };
>@@ -6460,6 +6458,18 @@
> 		1CDD45E40BA9C84600F90147 /* DebugRelease.xcconfig */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = text.xcconfig; path = DebugRelease.xcconfig; sourceTree = "<group>"; };
> 		1CDD45E50BA9C84600F90147 /* WebCore.xcconfig */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = text.xcconfig; path = WebCore.xcconfig; sourceTree = "<group>"; };
> 		1CDD45E60BA9C84600F90147 /* Base.xcconfig */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = text.xcconfig; path = Base.xcconfig; sourceTree = "<group>"; };
>+		1CECB3A821F2B67300F44542 /* WHLSLStatementBehaviorChecker.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLStatementBehaviorChecker.h; sourceTree = "<group>"; };
>+		1CECB3A921F2B67300F44542 /* WHLSLStatementBehaviorChecker.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLStatementBehaviorChecker.cpp; sourceTree = "<group>"; };
>+		1CECB3AF21F2B98400F44542 /* WHLSLFunctionWriter.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLFunctionWriter.cpp; sourceTree = "<group>"; };
>+		1CECB3B021F2B98500F44542 /* WHLSLTypeNamer.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLTypeNamer.h; sourceTree = "<group>"; };
>+		1CECB3B121F2B98500F44542 /* WHLSLTypeNamer.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLTypeNamer.cpp; sourceTree = "<group>"; };
>+		1CECB3B221F2B98600F44542 /* WHLSLFunctionWriter.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLFunctionWriter.h; sourceTree = "<group>"; };
>+		1CECB3B521F50AC700F44542 /* WHLSLMetalCodeGenerator.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLMetalCodeGenerator.cpp; sourceTree = "<group>"; };
>+		1CECB3B621F50AC700F44542 /* WHLSLMetalCodeGenerator.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLMetalCodeGenerator.h; sourceTree = "<group>"; };
>+		1CECB3B821F50D1000F44542 /* WHLSLNativeFunctionWriter.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLNativeFunctionWriter.cpp; sourceTree = "<group>"; };
>+		1CECB3B921F50D1000F44542 /* WHLSLNativeFunctionWriter.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLNativeFunctionWriter.h; sourceTree = "<group>"; };
>+		1CECB3BA21F511AA00F44542 /* WHLSLEntryPointScaffolding.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLEntryPointScaffolding.cpp; sourceTree = "<group>"; };
>+		1CECB3BB21F511AA00F44542 /* WHLSLEntryPointScaffolding.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLEntryPointScaffolding.h; sourceTree = "<group>"; };
> 		1CFAE3220A6D6A3F0032593D /* libobjc.dylib */ = {isa = PBXFileReference; lastKnownFileType = "compiled.mach-o.dylib"; name = libobjc.dylib; path = /usr/lib/libobjc.dylib; sourceTree = "<absolute>"; };
> 		1DC553FD211BA12A004B780E /* NavigatorShare.idl */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = text; path = NavigatorShare.idl; sourceTree = "<group>"; };
> 		1DC553FF211BA841004B780E /* ShareData.idl */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = text; path = ShareData.idl; sourceTree = "<group>"; };
>@@ -17150,6 +17160,23 @@
> 			tabWidth = 4;
> 			usesTabs = 0;
> 		};
>+		1CECB3AD21F2B96400F44542 /* Metal */ = {
>+			isa = PBXGroup;
>+			children = (
>+				1CECB3BA21F511AA00F44542 /* WHLSLEntryPointScaffolding.cpp */,
>+				1CECB3BB21F511AA00F44542 /* WHLSLEntryPointScaffolding.h */,
>+				1CECB3AF21F2B98400F44542 /* WHLSLFunctionWriter.cpp */,
>+				1CECB3B221F2B98600F44542 /* WHLSLFunctionWriter.h */,
>+				1CECB3B521F50AC700F44542 /* WHLSLMetalCodeGenerator.cpp */,
>+				1CECB3B621F50AC700F44542 /* WHLSLMetalCodeGenerator.h */,
>+				1CECB3B821F50D1000F44542 /* WHLSLNativeFunctionWriter.cpp */,
>+				1CECB3B921F50D1000F44542 /* WHLSLNativeFunctionWriter.h */,
>+				1CECB3B121F2B98500F44542 /* WHLSLTypeNamer.cpp */,
>+				1CECB3B021F2B98500F44542 /* WHLSLTypeNamer.h */,
>+			);
>+			path = Metal;
>+			sourceTree = "<group>";
>+		};
> 		26B9998D1803ADFA00D01121 /* cssjit */ = {
> 			isa = PBXGroup;
> 			children = (
>@@ -25562,6 +25589,7 @@
> 			isa = PBXGroup;
> 			children = (
> 				1CA0C2F621EEDAD200A11860 /* AST */,
>+				1CECB3AD21F2B96400F44542 /* Metal */,
> 				C234A9B221E92C1F003C984D /* WHLSLCheckDuplicateFunctions.cpp */,
> 				C234A9AE21E92C1A003C984D /* WHLSLCheckDuplicateFunctions.h */,
> 				1C840B9B21EC400900D0500D /* WHLSLChecker.cpp */,
>@@ -25580,8 +25608,6 @@
> 				C210E91221B4BD1000B7F83D /* WHLSLLexer.h */,
> 				1CA0C2EC21EED6F600A11860 /* WHLSLLiteralTypeChecker.cpp */,
> 				1CA0C2EA21EED6F500A11860 /* WHLSLLiteralTypeChecker.h */,
>-				1CA0C2F421EEDAD000A11860 /* WHLSLLoopChecker.cpp */,
>-				1CA0C2F521EEDAD100A11860 /* WHLSLLoopChecker.h */,
> 				C234A98D21E88884003C984D /* WHLSLNameContext.cpp */,
> 				C234A98E21E88885003C984D /* WHLSLNameContext.h */,
> 				C234A98A21E8883E003C984D /* WHLSLNameResolver.cpp */,
>@@ -25597,6 +25623,8 @@
> 				C234A99721E90F28003C984D /* WHLSLResolveOverloadImpl.h */,
> 				C234A99D21E910BD003C984D /* WHLSLResolvingType.h */,
> 				C21BF74521CD969800227979 /* WHLSLStandardLibrary.txt */,
>+				1CECB3A921F2B67300F44542 /* WHLSLStatementBehaviorChecker.cpp */,
>+				1CECB3A821F2B67300F44542 /* WHLSLStatementBehaviorChecker.h */,
> 				C234A9A921E92C17003C984D /* WHLSLSynthesizeArrayOperatorLength.cpp */,
> 				C234A9AF21E92C1B003C984D /* WHLSLSynthesizeArrayOperatorLength.h */,
> 				C234A9B021E92C1C003C984D /* WHLSLSynthesizeConstructors.cpp */,

Actions: View | Formatted Diff | Diff

Attachments on bug 193531: 359362 | 359401 | 359459 | 359463 | 359473 | 359643 | 359644 | 359646 | 359647