Attachment 359463 Details for Bug 193531 – WIP

[patch] WIP

bug-193531-20190118020011.patch (text/plain), 111.96 KB, created by Myles C. Maxfield on 2019-01-18 02:00:12 PST

(hide)

Description:

Filename:

MIME Type:

Creator: Myles C. Maxfield

Created: 2019-01-18 02:00:12 PST

Size: 111.96 KB

patch

obsolete

>Subversion Revision: 240132
>diff --git a/Source/WebCore/ChangeLog b/Source/WebCore/ChangeLog
>index 7c219c8368e2608a3751b873d1e92c04de5cc681..241de990d8dea4c09dfb4d590a1bd743d73e0a2d 100644
>--- a/Source/WebCore/ChangeLog
>+++ b/Source/WebCore/ChangeLog
>@@ -1,3 +1,40 @@
>+2019-01-17  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/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::ArrayTypeNameNode::ArrayTypeNameNode):
>+        (WebCore::WHLSL::Metal::ArrayTypeNameNode::numElements const):
>+        (WebCore::WHLSL::Metal::ArrayTypeNameNode::children):
>+        (WebCore::WHLSL::Metal::ArrayTypeNameNode::append):
>+        (WebCore::WHLSL::Metal::ArrayReferenceTypeNameNode::ArrayReferenceTypeNameNode):
>+        (WebCore::WHLSL::Metal::ArrayReferenceTypeNameNode::addressSpace const):
>+        (WebCore::WHLSL::Metal::ArrayReferenceTypeNameNode::children):
>+        (WebCore::WHLSL::Metal::ArrayReferenceTypeNameNode::append):
>+        (WebCore::WHLSL::Metal::PointerTypeNameNode::PointerTypeNameNode):
>+        (WebCore::WHLSL::Metal::PointerTypeNameNode::addressSpace const):
>+        (WebCore::WHLSL::Metal::PointerTypeNameNode::children):
>+        (WebCore::WHLSL::Metal::PointerTypeNameNode::append):
>+        (WebCore::WHLSL::Metal::ReferenceTypeNameNode::ReferenceTypeNameNode):
>+        (WebCore::WHLSL::Metal::ReferenceTypeNameNode::namedType const):
>+        (WebCore::WHLSL::Metal::TypeNamer::generateNextName):
>+        (WebCore::WHLSL::Metal::TypeNamer::build):
>+        (WebCore::WHLSL::Metal::TypeNamer::insert):
>+        (WebCore::WHLSL::Metal::typeNames):
>+        * Modules/webgpu/WHLSL/Metal/WHLSLTypeNamer.h: Added.
>+        * Sources.txt:
>+        * WebCore.xcodeproj/project.pbxproj:
>+
> 2019-01-17  Ross Kirsling  <ross.kirsling@sony.com>
> 
>         Unreviewed WinCairo fix -- hundreds of tests crash after r240031.
>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/Metal/WHLSLFunctionWriter.cpp b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp
>new file mode 100644
>index 0000000000000000000000000000000000000000..0983d4f6160e329e27a7b80c663afb10fe316df9
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp
>@@ -0,0 +1,988 @@
>+/*
>+ * 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 "WHLSLAssignmentExpression.h"
>+#include "WHLSLBooleanLiteral.h"
>+#include "WHLSLBuiltInSemantic.h"
>+#include "WHLSLCallExpression.h"
>+#include "WHLSLCommaExpression.h"
>+#include "WHLSLDereferenceExpression.h"
>+#include "WHLSLDoWhileLoop.h"
>+#include "WHLSLEffectfulExpressionStatement.h"
>+#include "WHLSLEntryPointType.h"
>+#include "WHLSLFloatLiteral.h"
>+#include "WHLSLForLoop.h"
>+#include "WHLSLFunctionDeclaration.h"
>+#include "WHLSLFunctionDefinition.h"
>+#include "WHLSLGatherEntryPointItems.h"
>+#include "WHLSLIfStatement.h"
>+#include "WHLSLIntegerLiteral.h"
>+#include "WHLSLLogicalExpression.h"
>+#include "WHLSLLogicalNotExpression.h"
>+#include "WHLSLMakeArrayReferenceExpression.h"
>+#include "WHLSLMakePointerExpression.h"
>+#include "WHLSLNativeFunctionDeclaration.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 getNativeName = [&](AST::UnnamedType& unnamedType) -> String {
>+            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 m_typeNamer.mangledNameForType(nativeTypeDeclaration);
>+        };
>+        auto mapFunctionName = [](String& functionName) -> String {
>+            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;
>+        };
>+        auto iterator = m_functionMapping.find(&nativeFunctionDeclaration);
>+        ASSERT(iterator != m_functionMapping.end());
>+        if (nativeFunctionDeclaration.isCast()) {
>+            ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
>+            auto metalParameterName = getNativeName(*nativeFunctionDeclaration.parameters()[0].type());
>+            auto metalReturnName = getNativeName(nativeFunctionDeclaration.type());
>+            if (metalParameterName != "atomic_int"_str && metalParameterName != "atomic_uint"_str) {
>+                m_stringBuilder.append(makeString(metalReturnName, ' ', iterator->value, '(', metalParameterName, "x) {\n"));
>+                m_stringBuilder.append(makeString("    return static_cast<", metalReturnName, ">(x);\n"));
>+                m_stringBuilder.append("}\n");
>+            } else {
>+                m_stringBuilder.append(makeString(metalReturnName, ' ', iterator->value, '(', metalParameterName, "x) {\n"));
>+                m_stringBuilder.append("    return atomic_load_explicit(&x, memory_order_relaxed);\n");
>+                m_stringBuilder.append("}\n");
>+            }
>+        } else if (nativeFunctionDeclaration.name().startsWith("operator."_str)) {
>+            if (nativeFunctionDeclaration.name().endsWith("="_str)) {
>+                ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
>+                auto metalParameter1Name = getNativeName(*nativeFunctionDeclaration.parameters()[0].type());
>+                auto metalParameter2Name = getNativeName(*nativeFunctionDeclaration.parameters()[1].type());
>+                auto metalReturnName = getNativeName(nativeFunctionDeclaration.type());
>+                auto fieldName = nativeFunctionDeclaration.name().substring("operator."_str.length());
>+                fieldName = fieldName.substring(0, fieldName.length() - 1);
>+                m_stringBuilder.append(makeString(metalReturnName, ' ', iterator->value, '(', metalParameter1Name, "v, ", metalParameter2Name, " n) {\n"));
>+                m_stringBuilder.append(makeString("    v.", fieldName, " = n;\n"));
>+                m_stringBuilder.append(makeString("    return v;\n"));
>+                m_stringBuilder.append("}\n");
>+            } else {
>+                ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
>+                auto metalParameterName = getNativeName(*nativeFunctionDeclaration.parameters()[0].type());
>+                auto metalReturnName = getNativeName(nativeFunctionDeclaration.type());
>+                m_stringBuilder.append(makeString(metalReturnName, ' ', iterator->value, '(', metalParameterName, "v) {\n"));
>+                m_stringBuilder.append(makeString("    return v.", nativeFunctionDeclaration.name().substring("operator."_str.length()), ";\n"));
>+                m_stringBuilder.append("}\n");
>+            }
>+        } else if (nativeFunctionDeclaration.name() == "operator[]") {
>+            ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
>+            auto metalParameter1Name = getNativeName(*nativeFunctionDeclaration.parameters()[0].type());
>+            auto metalParameter2Name = getNativeName(*nativeFunctionDeclaration.parameters()[1].type());
>+            auto metalReturnName = getNativeName(nativeFunctionDeclaration.type());
>+            m_stringBuilder.append(makeString(metalReturnName, ' ', iterator->value, '(', metalParameter1Name, "m, ", metalParameter2Name, " i) {\n"));
>+            m_stringBuilder.append(makeString("    return m[i];\n"));
>+            m_stringBuilder.append("}\n");
>+        } else if (nativeFunctionDeclaration.name() == "operator[]=") {
>+            ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
>+            auto metalParameter1Name = getNativeName(*nativeFunctionDeclaration.parameters()[0].type());
>+            auto metalParameter2Name = getNativeName(*nativeFunctionDeclaration.parameters()[1].type());
>+            auto metalParameter3Name = getNativeName(*nativeFunctionDeclaration.parameters()[2].type());
>+            auto metalReturnName = getNativeName(nativeFunctionDeclaration.type());
>+            m_stringBuilder.append(makeString(metalReturnName, ' ', iterator->value, '(', metalParameter1Name, "m, ", metalParameter2Name, " i, ", metalParameter3Name, " v) {\n"));
>+            m_stringBuilder.append(makeString("    m[i] = v;\n"));
>+            m_stringBuilder.append(makeString("    return m;\n"));
>+            m_stringBuilder.append("}\n");
>+        } else if (nativeFunctionDeclaration.isOperator()) {
>+            if (nativeFunctionDeclaration.parameters().size() == 1) {
>+                auto operatorName = nativeFunctionDeclaration.name().substring("operator"_str.length());
>+                auto metalParameterName = getNativeName(*nativeFunctionDeclaration.parameters()[0].type());
>+                auto metalReturnName = getNativeName(nativeFunctionDeclaration.type());
>+                m_stringBuilder.append(makeString(metalReturnName, ' ', iterator->value, '(', metalParameterName, "x) {\n"));
>+                m_stringBuilder.append(makeString("    return ", operatorName, "x;\n"));
>+                m_stringBuilder.append("}\n");
>+            } else {
>+                ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
>+                auto operatorName = nativeFunctionDeclaration.name().substring("operator"_str.length());
>+                auto metalParameter1Name = getNativeName(*nativeFunctionDeclaration.parameters()[0].type());
>+                auto metalParameter2Name = getNativeName(*nativeFunctionDeclaration.parameters()[1].type());
>+                auto metalReturnName = getNativeName(nativeFunctionDeclaration.type());
>+                m_stringBuilder.append(makeString(metalReturnName, ' ', iterator->value, '(', metalParameter1Name, "x, ", metalParameter2Name, " y) {\n"));
>+                m_stringBuilder.append(makeString("    return x ", operatorName, " y;\n"));
>+                m_stringBuilder.append("}\n");
>+            }
>+        } else 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());
>+            auto metalReturnName = getNativeName(nativeFunctionDeclaration.type());
>+            m_stringBuilder.append(makeString(metalReturnName, ' ', iterator->value, '(', metalParameterName, "x) {\n"));
>+            m_stringBuilder.append(makeString("    return ", mapFunctionName(nativeFunctionDeclaration.name()), "(x);\n"));
>+            m_stringBuilder.append("}\n");
>+        } else if (nativeFunctionDeclaration.name() == "pow" || nativeFunctionDeclaration.name() == "atan2") {
>+            ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
>+            auto metalParameter1Name = getNativeName(*nativeFunctionDeclaration.parameters()[0].type());
>+            auto metalParameter2Name = getNativeName(*nativeFunctionDeclaration.parameters()[1].type());
>+            auto metalReturnName = getNativeName(nativeFunctionDeclaration.type());
>+            m_stringBuilder.append(makeString(metalReturnName, ' ', iterator->value, '(', metalParameter1Name, "x, ", metalParameter2Name, " y) {\n"));
>+            m_stringBuilder.append(makeString("    return ", nativeFunctionDeclaration.name(), "(x, y);\n"));
>+            m_stringBuilder.append("}\n");
>+        } else if (nativeFunctionDeclaration.name() == "f16tof32" || nativeFunctionDeclaration.name() == "f32tof16") {
>+            // FIXME: Implement this
>+            CRASH();
>+        } else if (nativeFunctionDeclaration.name() == "AllMemoryBarrierWithGroupSync") {
>+            ASSERT(!nativeFunctionDeclaration.parameters().size());
>+            m_stringBuilder.append(makeString("void ", iterator->value, "() {\n"));
>+            m_stringBuilder.append("    threadgroup_barrier(mem_flags::mem_device);\n");
>+            m_stringBuilder.append("    threadgroup_barrier(mem_flags::mem_threadgroup);\n");
>+            m_stringBuilder.append("    threadgroup_barrier(mem_flags::mem_texture);\n");
>+            m_stringBuilder.append("}\n");
>+        } else if (nativeFunctionDeclaration.name() == "DeviceMemoryBarrierWithGroupSync") {
>+            ASSERT(!nativeFunctionDeclaration.parameters().size());
>+            m_stringBuilder.append(makeString("void ", iterator->value, "() {\n"));
>+            m_stringBuilder.append("    threadgroup_barrier(mem_flags::mem_device);\n");
>+            m_stringBuilder.append("}\n");
>+        } else if (nativeFunctionDeclaration.name() == "GroupMemoryBarrierWithGroupSync") {
>+            ASSERT(!nativeFunctionDeclaration.parameters().size());
>+            m_stringBuilder.append(makeString("void ", iterator->value, "() {\n"));
>+            m_stringBuilder.append("    threadgroup_barrier(mem_flags::mem_threadgroup);\n");
>+            m_stringBuilder.append("}\n");
>+        } else if (nativeFunctionDeclaration.name().startsWith("Interlocked"_str)) {
>+            auto convertAddressSpace = [](AST::AddressSpace addressSpace) -> String {
>+                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;
>+                }
>+            };
>+            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());
>+                auto secondArgument = getNativeName(*nativeFunctionDeclaration.parameters()[1].type());
>+                auto thirdArgument = getNativeName(*nativeFunctionDeclaration.parameters()[2].type());
>+                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());
>+                m_stringBuilder.append(makeString("void ", iterator->value, '(', convertAddressSpace(firstArgumentAddressSpace), ' ', firstArgumentPointee, "* object, ", secondArgument, " compare, ", thirdArgument, " desired, ", convertAddressSpace(fourthArgumentAddressSpace), ' ', fourthArgumentPointee, "* out) {\n"));
>+                m_stringBuilder.append("    atomic_compare_exchange_weak_explicit(object, &compare, desired, memory_order_relaxed);\n");
>+                m_stringBuilder.append("    *out = compare;\n");
>+                m_stringBuilder.append("}\n");
>+            } else {
>+                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());
>+                auto secondArgument = getNativeName(*nativeFunctionDeclaration.parameters()[1].type());
>+                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());
>+                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()));
>+                m_stringBuilder.append(makeString("void ", iterator->value, '(', convertAddressSpace(firstArgumentAddressSpace), ' ', firstArgumentPointee, "* object, ", secondArgument, " operand, ", convertAddressSpace(thirdArgumentAddressSpace), ' ', thirdArgumentPointee, "* out) {\n"));
>+                m_stringBuilder.append(makeString("    *out = atomic_fetch_", name, "_explicit(object, operand, memory_order_relaxed);\n"));
>+                m_stringBuilder.append("}\n");
>+            }
>+        } else if (nativeFunctionDeclaration.name() == "Sample") {
>+            // FIXME: Implement this.
>+            CRASH();
>+        } else if (nativeFunctionDeclaration.name() == "Load") {
>+            // FIXME: Implement this.
>+            CRASH();
>+        } else if (nativeFunctionDeclaration.name() == "GetDimensions") {
>+            // FIXME: Implement this.
>+            CRASH();
>+        } else if (nativeFunctionDeclaration.name() == "SampleBias") {
>+            // FIXME: Implement this.
>+            CRASH();
>+        } else if (nativeFunctionDeclaration.name() == "SampleGrad") {
>+            // FIXME: Implement this.
>+            CRASH();
>+        } else if (nativeFunctionDeclaration.name() == "SampleLevel") {
>+            // FIXME: Implement this.
>+            CRASH();
>+        } else if (nativeFunctionDeclaration.name() == "Gather") {
>+            // FIXME: Implement this.
>+            CRASH();
>+        } else if (nativeFunctionDeclaration.name() == "GatherRed") {
>+            // FIXME: Implement this.
>+            CRASH();
>+        } else if (nativeFunctionDeclaration.name() == "SampleCmp") {
>+            // FIXME: Implement this.
>+            CRASH();
>+        } else if (nativeFunctionDeclaration.name() == "SampleCmpLevelZero") {
>+            // FIXME: Implement this.
>+            CRASH();
>+        } else if (nativeFunctionDeclaration.name() == "Store") {
>+            // FIXME: Implement this.
>+            CRASH();
>+        } else if (nativeFunctionDeclaration.name() == "GatherAlpha") {
>+            // FIXME: Implement this.
>+            CRASH();
>+        } else if (nativeFunctionDeclaration.name() == "GatherBlue") {
>+            // FIXME: Implement this.
>+            CRASH();
>+        } else if (nativeFunctionDeclaration.name() == "GatherCmp") {
>+            // FIXME: Implement this.
>+            CRASH();
>+        } else if (nativeFunctionDeclaration.name() == "GatherCmpRed") {
>+            // FIXME: Implement this.
>+            CRASH();
>+        } else if (nativeFunctionDeclaration.name() == "GatherGreen") {
>+            // FIXME: Implement this.
>+            CRASH();
>+        }
>+    }
>+
>+    void visit(AST::FunctionDefinition& functionDefinition) override
>+    {
>+        auto iterator = m_functionMapping.find(&functionDefinition);
>+        ASSERT(iterator != m_functionMapping.end());
>+        if (functionDefinition.entryPointType()) {
>+            auto items = gatherEntryPointItems(m_intrinsics, functionDefinition);
>+            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(), ")]]");
>+            });
>+
>+            auto outputStructName = m_typeNamer.generateNextTypeName();
>+            Vector<String> outputNames;
>+            outputNames.reserveCapacity(items->outputs.size());
>+            m_stringBuilder.append(makeString("struct ", outputStructName, "{\n"));
>+            for (auto& output : items->outputs) {
>+                auto elementName = m_typeNamer.generateNextStructureElementName();
>+                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;
>+                }
>+            })();
>+            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();
>+                m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(input.unnamedType), ' ', parameterName, ' ', WTF::visit(semanticVisitor, input.semantic)));
>+            }
>+            m_stringBuilder.append(") {\n");
>+            #error Pack the data
>+            checkErrorAndVisit(functionDefinition.block());
>+            m_stringBuilder.append("}\n");
>+        } else {
>+            m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(functionDefinition.type()), ' ', iterator->value, '('));
>+            for (size_t i = 0; i < functionDefinition.parameters().size(); ++i) {
>+                if (i)
>+                    m_stringBuilder.append(", ");
>+                m_stringBuilder.append(m_typeNamer.mangledNameForType(*functionDefinition.parameters()[i].type()));
>+                #error emit variable names
>+            }
>+            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
>+    {
>+        m_stringBuilder.append("continue;\n");
>+    }
>+
>+    void visit(AST::DoWhileLoop& doWhileLoop) override
>+    {
>+        #error lambda
>+        auto lambdaName = generateNextVariableName();
>+        m_stringBuilder.append(makeString("auto ", lambdaName, " = [&]() {\n"));
>+        checkErrorAndVisit(doWhileLoop.conditional());
>+        m_stringBuilder.append(makeString("return ", m_stack.takeLast(), ";\n};\n"));
>+
>+        m_stringBuilder.append("do {\n");
>+        checkErrorAndVisit(doWhileLoop.body());
>+        m_stringBuilder.append(makeString("} while(", lambdaName, "());\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());
>+
>+        #error lambda
>+        auto conditionLambdaName = generateNextVariableName();
>+        if (forLoop.condition()) {
>+            m_stringBuilder.append(makeString("auto ", conditionLambdaName, " = [&]() {\n"));
>+            checkErrorAndVisit(*forLoop.condition());
>+            m_stringBuilder.append("return ");
>+            m_stringBuilder.append(makeString(m_stack.takeLast(), ";\n};\n"));
>+        }
>+
>+        auto incrementLambdaName = generateNextVariableName();
>+        if (forLoop.increment()) {
>+            m_stringBuilder.append(makeString("auto ", conditionLambdaName, " = [&]() {\n"));
>+            checkErrorAndVisit(*forLoop.increment());
>+            m_stack.takeLast(); // We don't need to do anything with the result.
>+            m_stringBuilder.append("};\n");
>+        }
>+
>+        m_stringBuilder.append("for ( ; ");
>+        if (forLoop.condition())
>+            m_stringBuilder.append(makeString(conditionLambdaName, "() "));
>+        m_stringBuilder.append("; ");
>+        if (forLoop.increment())
>+            m_stringBuilder.append(makeString(incrementLambdaName, "() "));
>+        m_stringBuilder.append(") {\n");
>+        checkErrorAndVisit(forLoop.body());
>+        m_stringBuilder.append("}\n");
>+    }
>+
>+    void visit(AST::IfStatement& ifStatement) override
>+    {
>+        #error lambda
>+        auto conditionalLambdaName = generateNextVariableName();
>+        m_stringBuilder.append(makeString("auto ", conditionalLambdaName, " = [&]() {\n"));
>+        checkErrorAndVisit(ifStatement.conditional());
>+        m_stringBuilder.append(makeString("return ", m_stack.takeLast(), ";\n};\n"));
>+
>+        m_stringBuilder.append(makeString("if (", conditionalLambdaName, "() {\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
>+    {
>+        #error detect if we're returning from an entry point, and fixup the result
>+        if (returnStatement.value()) {
>+            checkErrorAndVisit(*returnStatement.value());
>+            auto value = m_stack.takeLast();
>+            m_stringBuilder.append(makeString("return ", value, ";\n"));
>+        } else
>+            m_stringBuilder.append("return;\n");
>+    }
>+
>+    void visit(AST::SwitchStatement& switchStatement) override
>+    {
>+        #error lambda
>+        auto valueLambdaName = generateNextVariableName();
>+        m_stringBuilder.append(makeString("auto ", valueLambdaName, " = [&]() {\n"));
>+        checkErrorAndVisit(switchStatement.value());
>+        m_stringBuilder.append(makeString("return ", m_stack.takeLast(), ";\n};\n"));
>+
>+        m_stringBuilder.append(makeString("switch (", valueLambdaName, "()) {"));
>+        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
>+    {
>+        #error lambda
>+        auto lambdaName = generateNextVariableName();
>+        m_stringBuilder.append(makeString("auto ", lambdaName, " = [&]() {\n"));
>+        checkErrorAndVisit(whileLoop.conditional());
>+        m_stringBuilder.append(makeString("return ", m_stack.takeLast(), ";\n};\n"));
>+
>+        m_stringBuilder.append(makeString("while (", lambdaName, "()) {\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;
>+    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();
>+}
>+
>+}
>+
>+}
>+
>+}
>+
>+#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..0d0dfc8cc7fd0c350741c0d36f4bed18395be025
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.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)
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+class Program;
>+
>+namespace Metal {
>+
>+class TypeNamer;
>+
>+// metal_stdlib, metal_atomic, metal_math, metal_relational, metal_compute
>+String metalFunctions(Program&, 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..1cb673f941f543b95036ad5c70c33eae214ce510
>--- /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<TypeReference>>(nativeTypeDeclaration.typeArguments()[0]));
>+        auto& typeReference = WTF::get<UniqueRef<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<ConstantExpression>(nativeTypeDeclaration.typeArguments()[1]));
>+        auto& constantExpression = WTF::get<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<TypeReference>>(nativeTypeDeclaration.typeArguments()[0]));
>+        auto& typeReference = WTF::get<UniqueRef<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<ConstantExpression>(nativeTypeDeclaration.typeArguments()[1]));
>+        auto& constantExpression1 = WTF::get<ConstantExpression>(nativeTypeDeclaration.typeArguments()[0]);
>+        auto& integerLiteral1 = constantExpression.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<ConstantExpression>(nativeTypeDeclaration.typeArguments()[2]));
>+        auto& constantExpression2 = WTF::get<ConstantExpression>(nativeTypeDeclaration.typeArguments()[0]);
>+        auto& integerLiteral2 = constantExpression.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<TypeReference>>(nativeTypeDeclaration.typeArguments()[0]));
>+    auto& typeReference = WTF::get<UniqueRef<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();
>+}
>+
>+}
>+
>+}
>+
>+}
>+
>+#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..c70a0e71a9338d6780a7e74b74356fd427395a09
>--- /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 };
>+};
>+
>+}
>+
>+}
>+
>+}
>+
>+#endif
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLChecker.cpp b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLChecker.cpp
>index 60b0efeecac2463e51a8ec123dccf7a20a46f9e6..3c7e89302335867bfc71a59f4c8cc5f7e5ef549d 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLChecker.cpp
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLChecker.cpp
>@@ -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/Sources.txt b/Source/WebCore/Sources.txt
>index b6124192ce94452f6d1fe4bbf309fcdac41b472a..ea7684e836d2e01adca8a67b57c12c7352e58bea 100644
>--- a/Source/WebCore/Sources.txt
>+++ b/Source/WebCore/Sources.txt
>@@ -326,6 +326,8 @@ 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/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 5ea09b0bc7b1ce11584f2d689834db2754864ba4..86e769db896d014b1b49725668b8fe6f0de2f842 100644
>--- a/Source/WebCore/WebCore.xcodeproj/project.pbxproj
>+++ b/Source/WebCore/WebCore.xcodeproj/project.pbxproj
>@@ -6430,17 +6430,16 @@
> 		1C840B9A21EC400900D0500D /* WHLSLGatherEntryPointItems.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLGatherEntryPointItems.h; sourceTree = "<group>"; };
> 		1C840B9B21EC400900D0500D /* WHLSLChecker.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLChecker.cpp; sourceTree = "<group>"; };
> 		1C904DF90BA9D2C80081E9D0 /* Version.xcconfig */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = text.xcconfig; path = Version.xcconfig; sourceTree = "<group>"; };
>-		1CA0C2E421EED12A00A11860 /* WHLSLFunctionStageChecker.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLFunctionStageChecker.cpp; sourceTree = "<group>"; };
>-		1CA0C2E521EED12A00A11860 /* WHLSLFunctionStageChecker.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLFunctionStageChecker.h; sourceTree = "<group>"; };
>+		1C9AE5CA21ED9DF50069D5F2 /* WHLSLHighZombieFinder.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLHighZombieFinder.cpp; sourceTree = "<group>"; };
>+		1C9AE5CB21ED9DF50069D5F2 /* WHLSLHighZombieFinder.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLHighZombieFinder.h; sourceTree = "<group>"; };
> 		1CA0C2DE21EEB5F400A11860 /* WHLSLRecursionChecker.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLRecursionChecker.h; sourceTree = "<group>"; };
> 		1CA0C2E021EEB5F500A11860 /* WHLSLRecursionChecker.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLRecursionChecker.cpp; sourceTree = "<group>"; };
>+		1CA0C2E421EED12A00A11860 /* WHLSLFunctionStageChecker.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLFunctionStageChecker.cpp; sourceTree = "<group>"; };
>+		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>"; };
>-		1C9AE5CA21ED9DF50069D5F2 /* WHLSLHighZombieFinder.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLHighZombieFinder.cpp; sourceTree = "<group>"; };
>-		1C9AE5CB21ED9DF50069D5F2 /* WHLSLHighZombieFinder.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLHighZombieFinder.h; 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>"; };
>-		1CA0C2F621EEDAD200A11860 /* AST */ = {isa = PBXFileReference; lastKnownFileType = folder; path = AST; 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>"; };
>@@ -6457,6 +6456,8 @@
> 		1CB69B3B21DF041E006E846A /* WHLSLUnsignedIntegerLiteralType.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLUnsignedIntegerLiteralType.cpp; sourceTree = "<group>"; };
> 		1CB6B4F8217B83930093B9CD /* TextDecorationThickness.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = TextDecorationThickness.h; sourceTree = "<group>"; };
> 		1CB6B4FB217B83940093B9CD /* TextUnderlineOffset.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = TextUnderlineOffset.h; sourceTree = "<group>"; };
>+		1CC83DE121F15D3D0064042E /* WHLSLFunctionWriter.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLFunctionWriter.cpp; sourceTree = "<group>"; };
>+		1CC83DE221F15D3D0064042E /* WHLSLFunctionWriter.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLFunctionWriter.h; sourceTree = "<group>"; };
> 		1CCDF5BB1990332400BCEBAD /* SVGToOTFFontConversion.cpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.cpp; path = SVGToOTFFontConversion.cpp; sourceTree = "<group>"; };
> 		1CCDF5BC1990332400BCEBAD /* SVGToOTFFontConversion.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = SVGToOTFFontConversion.h; sourceTree = "<group>"; };
> 		1CDD45E40BA9C84600F90147 /* DebugRelease.xcconfig */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = text.xcconfig; path = DebugRelease.xcconfig; sourceTree = "<group>"; };
>@@ -13391,6 +13392,8 @@
> 		C2458E611FE8979E00594759 /* FontCacheCoreText.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = FontCacheCoreText.h; sourceTree = "<group>"; };
> 		C26017A11C72DC9900F74A16 /* CSSFontFaceSet.cpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.cpp; path = CSSFontFaceSet.cpp; sourceTree = "<group>"; };
> 		C26017A21C72DC9900F74A16 /* CSSFontFaceSet.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = CSSFontFaceSet.h; sourceTree = "<group>"; };
>+		C26D2C7021F05D5500497899 /* WHLSLTypeNamer.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLTypeNamer.cpp; sourceTree = "<group>"; };
>+		C26D2C7121F05D5500497899 /* WHLSLTypeNamer.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLTypeNamer.h; sourceTree = "<group>"; };
> 		C280833C1C6DB194001451B6 /* FontFace.idl */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = text; path = FontFace.idl; sourceTree = "<group>"; };
> 		C280833D1C6DC22C001451B6 /* JSFontFace.cpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.cpp; name = JSFontFace.cpp; path = DerivedSources/WebCore/JSFontFace.cpp; sourceTree = BUILT_PRODUCTS_DIR; };
> 		C280833E1C6DC22C001451B6 /* JSFontFace.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; name = JSFontFace.h; path = DerivedSources/WebCore/JSFontFace.h; sourceTree = BUILT_PRODUCTS_DIR; };
>@@ -17038,6 +17041,100 @@
> 			tabWidth = 8;
> 			usesTabs = 0;
> 		};
>+		1CA0C2F621EEDAD200A11860 /* AST */ = {
>+			isa = PBXGroup;
>+			children = (
>+				1C840B9021EC30F900D0500D /* WHLSLAddressSpace.h */,
>+				C21BF72521CD89E200227979 /* WHLSLArrayReferenceType.h */,
>+				C21BF70921CD89CA00227979 /* WHLSLArrayType.h */,
>+				C21BF73021CD89ED00227979 /* WHLSLAssignmentExpression.h */,
>+				C21BF70A21CD89CB00227979 /* WHLSLBaseFunctionAttribute.h */,
>+				C21BF6FA21CD89BE00227979 /* WHLSLBaseSemantic.h */,
>+				C21BF71E21CD89DC00227979 /* WHLSLBlock.h */,
>+				C21BF6F621CD89B700227979 /* WHLSLBooleanLiteral.h */,
>+				C21BF71A21CD89D800227979 /* WHLSLBreak.h */,
>+				C2138A1321DDECD300F516BA /* WHLSLBuiltInSemantic.cpp */,
>+				C21BF72221CD89DF00227979 /* WHLSLBuiltInSemantic.h */,
>+				C21BF71621CD89D500227979 /* WHLSLCallExpression.h */,
>+				C21BF70621CD89C700227979 /* WHLSLCommaExpression.h */,
>+				C21BF73321CD89F000227979 /* WHLSLConstantExpression.h */,
>+				C21BF6F521CD89B500227979 /* WHLSLContinue.h */,
>+				C21BF70121CD89C400227979 /* WHLSLDereferenceExpression.h */,
>+				C21BF71821CD89D700227979 /* WHLSLDotExpression.h */,
>+				C21BF6FB21CD89BE00227979 /* WHLSLDoWhileLoop.h */,
>+				C21BF72821CD89E500227979 /* WHLSLEffectfulExpressionStatement.h */,
>+				1C840B7D21EBE0B800D0500D /* WHLSLEntryPointType.h */,
>+				C21BF72021CD89DD00227979 /* WHLSLEnumerationDefinition.h */,
>+				C21BF72621CD89E300227979 /* WHLSLEnumerationMember.h */,
>+				C21BF70221CD89C400227979 /* WHLSLEnumerationMemberLiteral.h */,
>+				C21BF70C21CD89CC00227979 /* WHLSLExpression.h */,
>+				C21BF70021CD89C200227979 /* WHLSLFallthrough.h */,
>+				C21BF73521CD89F200227979 /* WHLSLFloatLiteral.h */,
>+				1CB69B3821DF03E1006E846A /* WHLSLFloatLiteralType.cpp */,
>+				1CB69B3421DED63A006E846A /* WHLSLFloatLiteralType.h */,
>+				C21BF73421CD89F100227979 /* WHLSLForLoop.h */,
>+				C21BF70521CD89C700227979 /* WHLSLFunctionAttribute.h */,
>+				C21BF6FD21CD89C000227979 /* WHLSLFunctionDeclaration.h */,
>+				C21BF6F421CD89B300227979 /* WHLSLFunctionDefinition.h */,
>+				C21BF6FF21CD89C200227979 /* WHLSLIfStatement.h */,
>+				C21BF6F721CD89B900227979 /* WHLSLIndexExpression.h */,
>+				C20F4F6621DFF2360070C45A /* WHLSLIntegerLiteral.cpp */,
>+				C21BF6F821CD89BB00227979 /* WHLSLIntegerLiteral.h */,
>+				1CB69B3921DF03F3006E846A /* WHLSLIntegerLiteralType.cpp */,
>+				1CB69B3521DED649006E846A /* WHLSLIntegerLiteralType.h */,
>+				C21BF73221CD89EF00227979 /* WHLSLLogicalExpression.h */,
>+				C21BF71521CD89D400227979 /* WHLSLLogicalNotExpression.h */,
>+				C21BF72D21CD89E900227979 /* WHLSLMakeArrayReferenceExpression.h */,
>+				C21BF72C21CD89E900227979 /* WHLSLMakePointerExpression.h */,
>+				1C33277121CF0BE1000DC9F2 /* WHLSLNamedType.h */,
>+				C21BF72321CD89E100227979 /* WHLSLNativeFunctionDeclaration.h */,
>+				C21BF72A21CD89E700227979 /* WHLSLNativeTypeDeclaration.h */,
>+				C21BF72421CD89E100227979 /* WHLSLNode.h */,
>+				C21BF70721CD89C800227979 /* WHLSLNullLiteral.h */,
>+				1CB69B3A21DF0403006E846A /* WHLSLNullLiteralType.cpp */,
>+				1CB69B3621DED657006E846A /* WHLSLNullLiteralType.h */,
>+				C21BF72121CD89DE00227979 /* WHLSLNumThreadsFunctionAttribute.h */,
>+				C21BF72F21CD89EC00227979 /* WHLSLPointerType.h */,
>+				C21BF72E21CD89EA00227979 /* WHLSLPropertyAccessExpression.h */,
>+				C21BF70B21CD89CC00227979 /* WHLSLQualifier.h */,
>+				C21BF71B21CD89D900227979 /* WHLSLReadModifyWriteExpression.h */,
>+				C21BF70D21CD89CD00227979 /* WHLSLReferenceType.h */,
>+				1CB69B3221DED40B006E846A /* WHLSLResolvableType.h */,
>+				C2138A1521DDECE900F516BA /* WHLSLResourceSemantic.cpp */,
>+				C21BF70821CD89C900227979 /* WHLSLResourceSemantic.h */,
>+				C21BF70321CD89C500227979 /* WHLSLReturn.h */,
>+				C21BF6F921CD89BD00227979 /* WHLSLSemantic.h */,
>+				C2138A1621DDECFB00F516BA /* WHLSLSpecializationConstantSemantic.cpp */,
>+				C21BF73621CD89F300227979 /* WHLSLSpecializationConstantSemantic.h */,
>+				C2138A1721DDED0D00F516BA /* WHLSLStageInOutSemantic.cpp */,
>+				C21BF70E21CD89CE00227979 /* WHLSLStageInOutSemantic.h */,
>+				C21BF71221CD89D100227979 /* WHLSLStatement.h */,
>+				C21BF72721CD89E400227979 /* WHLSLStructureDefinition.h */,
>+				C21BF6FE21CD89C100227979 /* WHLSLStructureElement.h */,
>+				C21BF71921CD89D700227979 /* WHLSLSwitchCase.h */,
>+				C21BF73121CD89EE00227979 /* WHLSLSwitchStatement.h */,
>+				C21BF71C21CD89DA00227979 /* WHLSLTernaryExpression.h */,
>+				C21BF6F321CD89AD00227979 /* WHLSLTrap.h */,
>+				C21BF71D21CD89DB00227979 /* WHLSLType.h */,
>+				C288C72D21C991DA002DF5CA /* WHLSLTypeArgument.cpp */,
>+				C21BF71121CD89D100227979 /* WHLSLTypeArgument.h */,
>+				C21BF72921CD89E600227979 /* WHLSLTypeDefinition.h */,
>+				C20F4F6421DFBE5C0070C45A /* WHLSLTypeReference.cpp */,
>+				C21BF71F21CD89DC00227979 /* WHLSLTypeReference.h */,
>+				1C33277221CF0D2E000DC9F2 /* WHLSLUnnamedType.h */,
>+				C20F4F6721DFF3A70070C45A /* WHLSLUnsignedIntegerLiteral.cpp */,
>+				C21BF72B21CD89E800227979 /* WHLSLUnsignedIntegerLiteral.h */,
>+				1CB69B3B21DF041E006E846A /* WHLSLUnsignedIntegerLiteralType.cpp */,
>+				1CB69B3721DED66B006E846A /* WHLSLUnsignedIntegerLiteralType.h */,
>+				C21BF6FC21CD89BF00227979 /* WHLSLValue.h */,
>+				C21BF71021CD89D000227979 /* WHLSLVariableDeclaration.h */,
>+				C21BF71421CD89D300227979 /* WHLSLVariableDeclarationsStatement.h */,
>+				C21BF71321CD89D200227979 /* WHLSLVariableReference.h */,
>+				C21BF70421CD89C600227979 /* WHLSLWhileLoop.h */,
>+			);
>+			path = AST;
>+			sourceTree = "<group>";
>+		};
> 		1CDD44660BA9C80000F90147 /* Configurations */ = {
> 			isa = PBXGroup;
> 			children = (
>@@ -25466,6 +25563,7 @@
> 			isa = PBXGroup;
> 			children = (
> 				1CA0C2F621EEDAD200A11860 /* AST */,
>+				C26D2C6E21F05CDE00497899 /* Metal */,
> 				C234A9B221E92C1F003C984D /* WHLSLCheckDuplicateFunctions.cpp */,
> 				C234A9AE21E92C1A003C984D /* WHLSLCheckDuplicateFunctions.h */,
> 				1C840B9B21EC400900D0500D /* WHLSLChecker.cpp */,
>@@ -25515,98 +25613,15 @@
> 			path = WHLSL;
> 			sourceTree = "<group>";
> 		};
>-		1CA0C2F621EEDAD200A11860 /* AST */ = {
>+		C26D2C6E21F05CDE00497899 /* Metal */ = {
> 			isa = PBXGroup;
> 			children = (
>-				1C840B9021EC30F900D0500D /* WHLSLAddressSpace.h */,
>-				C21BF72521CD89E200227979 /* WHLSLArrayReferenceType.h */,
>-				C21BF70921CD89CA00227979 /* WHLSLArrayType.h */,
>-				C21BF73021CD89ED00227979 /* WHLSLAssignmentExpression.h */,
>-				C21BF70A21CD89CB00227979 /* WHLSLBaseFunctionAttribute.h */,
>-				C21BF6FA21CD89BE00227979 /* WHLSLBaseSemantic.h */,
>-				C21BF71E21CD89DC00227979 /* WHLSLBlock.h */,
>-				C21BF6F621CD89B700227979 /* WHLSLBooleanLiteral.h */,
>-				C21BF71A21CD89D800227979 /* WHLSLBreak.h */,
>-				C2138A1321DDECD300F516BA /* WHLSLBuiltInSemantic.cpp */,
>-				C21BF72221CD89DF00227979 /* WHLSLBuiltInSemantic.h */,
>-				C21BF71621CD89D500227979 /* WHLSLCallExpression.h */,
>-				C21BF70621CD89C700227979 /* WHLSLCommaExpression.h */,
>-				C21BF73321CD89F000227979 /* WHLSLConstantExpression.h */,
>-				C21BF6F521CD89B500227979 /* WHLSLContinue.h */,
>-				C21BF70121CD89C400227979 /* WHLSLDereferenceExpression.h */,
>-				C21BF71821CD89D700227979 /* WHLSLDotExpression.h */,
>-				C21BF6FB21CD89BE00227979 /* WHLSLDoWhileLoop.h */,
>-				C21BF72821CD89E500227979 /* WHLSLEffectfulExpressionStatement.h */,
>-				1C840B7D21EBE0B800D0500D /* WHLSLEntryPointType.h */,
>-				C21BF72021CD89DD00227979 /* WHLSLEnumerationDefinition.h */,
>-				C21BF72621CD89E300227979 /* WHLSLEnumerationMember.h */,
>-				C21BF70221CD89C400227979 /* WHLSLEnumerationMemberLiteral.h */,
>-				C21BF70C21CD89CC00227979 /* WHLSLExpression.h */,
>-				C21BF70021CD89C200227979 /* WHLSLFallthrough.h */,
>-				C21BF73521CD89F200227979 /* WHLSLFloatLiteral.h */,
>-				1CB69B3821DF03E1006E846A /* WHLSLFloatLiteralType.cpp */,
>-				1CB69B3421DED63A006E846A /* WHLSLFloatLiteralType.h */,
>-				C21BF73421CD89F100227979 /* WHLSLForLoop.h */,
>-				C21BF70521CD89C700227979 /* WHLSLFunctionAttribute.h */,
>-				C21BF6FD21CD89C000227979 /* WHLSLFunctionDeclaration.h */,
>-				C21BF6F421CD89B300227979 /* WHLSLFunctionDefinition.h */,
>-				C21BF6FF21CD89C200227979 /* WHLSLIfStatement.h */,
>-				C21BF6F721CD89B900227979 /* WHLSLIndexExpression.h */,
>-				C20F4F6621DFF2360070C45A /* WHLSLIntegerLiteral.cpp */,
>-				C21BF6F821CD89BB00227979 /* WHLSLIntegerLiteral.h */,
>-				1CB69B3921DF03F3006E846A /* WHLSLIntegerLiteralType.cpp */,
>-				1CB69B3521DED649006E846A /* WHLSLIntegerLiteralType.h */,
>-				C21BF73221CD89EF00227979 /* WHLSLLogicalExpression.h */,
>-				C21BF71521CD89D400227979 /* WHLSLLogicalNotExpression.h */,
>-				C21BF72D21CD89E900227979 /* WHLSLMakeArrayReferenceExpression.h */,
>-				C21BF72C21CD89E900227979 /* WHLSLMakePointerExpression.h */,
>-				1C33277121CF0BE1000DC9F2 /* WHLSLNamedType.h */,
>-				C21BF72321CD89E100227979 /* WHLSLNativeFunctionDeclaration.h */,
>-				C21BF72A21CD89E700227979 /* WHLSLNativeTypeDeclaration.h */,
>-				C21BF72421CD89E100227979 /* WHLSLNode.h */,
>-				C21BF70721CD89C800227979 /* WHLSLNullLiteral.h */,
>-				1CB69B3A21DF0403006E846A /* WHLSLNullLiteralType.cpp */,
>-				1CB69B3621DED657006E846A /* WHLSLNullLiteralType.h */,
>-				C21BF72121CD89DE00227979 /* WHLSLNumThreadsFunctionAttribute.h */,
>-				C21BF72F21CD89EC00227979 /* WHLSLPointerType.h */,
>-				C21BF72E21CD89EA00227979 /* WHLSLPropertyAccessExpression.h */,
>-				C21BF70B21CD89CC00227979 /* WHLSLQualifier.h */,
>-				C21BF71B21CD89D900227979 /* WHLSLReadModifyWriteExpression.h */,
>-				C21BF70D21CD89CD00227979 /* WHLSLReferenceType.h */,
>-				1CB69B3221DED40B006E846A /* WHLSLResolvableType.h */,
>-				C2138A1521DDECE900F516BA /* WHLSLResourceSemantic.cpp */,
>-				C21BF70821CD89C900227979 /* WHLSLResourceSemantic.h */,
>-				C21BF70321CD89C500227979 /* WHLSLReturn.h */,
>-				C21BF6F921CD89BD00227979 /* WHLSLSemantic.h */,
>-				C2138A1621DDECFB00F516BA /* WHLSLSpecializationConstantSemantic.cpp */,
>-				C21BF73621CD89F300227979 /* WHLSLSpecializationConstantSemantic.h */,
>-				C2138A1721DDED0D00F516BA /* WHLSLStageInOutSemantic.cpp */,
>-				C21BF70E21CD89CE00227979 /* WHLSLStageInOutSemantic.h */,
>-				C21BF71221CD89D100227979 /* WHLSLStatement.h */,
>-				C21BF72721CD89E400227979 /* WHLSLStructureDefinition.h */,
>-				C21BF6FE21CD89C100227979 /* WHLSLStructureElement.h */,
>-				C21BF71921CD89D700227979 /* WHLSLSwitchCase.h */,
>-				C21BF73121CD89EE00227979 /* WHLSLSwitchStatement.h */,
>-				C21BF71C21CD89DA00227979 /* WHLSLTernaryExpression.h */,
>-				C21BF6F321CD89AD00227979 /* WHLSLTrap.h */,
>-				C21BF71D21CD89DB00227979 /* WHLSLType.h */,
>-				C288C72D21C991DA002DF5CA /* WHLSLTypeArgument.cpp */,
>-				C21BF71121CD89D100227979 /* WHLSLTypeArgument.h */,
>-				C21BF72921CD89E600227979 /* WHLSLTypeDefinition.h */,
>-				C20F4F6421DFBE5C0070C45A /* WHLSLTypeReference.cpp */,
>-				C21BF71F21CD89DC00227979 /* WHLSLTypeReference.h */,
>-				1C33277221CF0D2E000DC9F2 /* WHLSLUnnamedType.h */,
>-				C20F4F6721DFF3A70070C45A /* WHLSLUnsignedIntegerLiteral.cpp */,
>-				C21BF72B21CD89E800227979 /* WHLSLUnsignedIntegerLiteral.h */,
>-				1CB69B3B21DF041E006E846A /* WHLSLUnsignedIntegerLiteralType.cpp */,
>-				1CB69B3721DED66B006E846A /* WHLSLUnsignedIntegerLiteralType.h */,
>-				C21BF6FC21CD89BF00227979 /* WHLSLValue.h */,
>-				C21BF71021CD89D000227979 /* WHLSLVariableDeclaration.h */,
>-				C21BF71421CD89D300227979 /* WHLSLVariableDeclarationsStatement.h */,
>-				C21BF71321CD89D200227979 /* WHLSLVariableReference.h */,
>-				C21BF70421CD89C600227979 /* WHLSLWhileLoop.h */,
>+				1CC83DE121F15D3D0064042E /* WHLSLFunctionWriter.cpp */,
>+				1CC83DE221F15D3D0064042E /* WHLSLFunctionWriter.h */,
>+				C26D2C7021F05D5500497899 /* WHLSLTypeNamer.cpp */,
>+				C26D2C7121F05D5500497899 /* WHLSLTypeNamer.h */,
> 			);
>-			path = AST;
>+			path = Metal;
> 			sourceTree = "<group>";
> 		};
> 		C96F5EBF1B5872260091EA9D /* mediasession */ = {

Actions: View | Formatted Diff | Diff

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