StringBuilder::append(makeString(...)) is inefficient
[WebKit-https.git] / Source / WebCore / Modules / webgpu / WHLSL / Metal / WHLSLNativeFunctionWriter.cpp
1 /*
2  * Copyright (C) 2019 Apple Inc. All rights reserved.
3  *
4  * Redistribution and use in source and binary forms, with or without
5  * modification, are permitted provided that the following conditions
6  * are met:
7  * 1. Redistributions of source code must retain the above copyright
8  *    notice, this list of conditions and the following disclaimer.
9  * 2. Redistributions in binary form must reproduce the above copyright
10  *    notice, this list of conditions and the following disclaimer in the
11  *    documentation and/or other materials provided with the distribution.
12  *
13  * THIS SOFTWARE IS PROVIDED BY APPLE INC. AND ITS CONTRIBUTORS ``AS IS''
14  * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO,
15  * THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
16  * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL APPLE INC. OR ITS CONTRIBUTORS
17  * BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
18  * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
19  * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
20  * INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
21  * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
22  * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
23  * THE POSSIBILITY OF SUCH DAMAGE.
24  */
25
26 #include "config.h"
27 #include "WHLSLNativeFunctionWriter.h"
28
29 #if ENABLE(WEBGPU)
30
31 #include "NotImplemented.h"
32 #include "WHLSLAddressSpace.h"
33 #include "WHLSLArrayType.h"
34 #include "WHLSLEnumerationDefinition.h"
35 #include "WHLSLInferTypes.h"
36 #include "WHLSLIntrinsics.h"
37 #include "WHLSLNamedType.h"
38 #include "WHLSLNativeFunctionDeclaration.h"
39 #include "WHLSLNativeTypeDeclaration.h"
40 #include "WHLSLPointerType.h"
41 #include "WHLSLStructureDefinition.h"
42 #include "WHLSLTypeDefinition.h"
43 #include "WHLSLTypeNamer.h"
44 #include "WHLSLUnnamedType.h"
45 #include "WHLSLVariableDeclaration.h"
46 #include <wtf/text/StringBuilder.h>
47
48 namespace WebCore {
49
50 namespace WHLSL {
51
52 namespace Metal {
53
54 static String mapFunctionName(String& functionName)
55 {
56     if (functionName == "ddx")
57         return "dfdx"_str;
58     if (functionName == "ddy")
59         return "dfdy"_str;
60     if (functionName == "asint")
61         return "as_type<int32_t>"_str;
62     if (functionName == "asuint")
63         return "as_type<uint32_t>"_str;
64     if (functionName == "asfloat")
65         return "as_type<float>"_str;
66     return functionName;
67 }
68
69 static String atomicName(String input)
70 {
71     if (input == "Add")
72         return "fetch_add"_str;
73     if (input == "And")
74         return "fetch_and"_str;
75     if (input == "Exchange")
76         return "exchange"_str;
77     if (input == "Max")
78         return "fetch_max"_str;
79     if (input == "Min")
80         return "fetch_min"_str;
81     if (input == "Or")
82         return "fetch_or"_str;
83     ASSERT(input == "Xor");
84         return "fetch_xor"_str;
85 }
86
87 static int vectorLength(AST::NativeTypeDeclaration& nativeTypeDeclaration)
88 {
89     int vectorLength = 1;
90     if (!nativeTypeDeclaration.typeArguments().isEmpty()) {
91         ASSERT(nativeTypeDeclaration.typeArguments().size() == 2);
92         ASSERT(WTF::holds_alternative<AST::ConstantExpression>(nativeTypeDeclaration.typeArguments()[1]));
93         vectorLength = WTF::get<AST::ConstantExpression>(nativeTypeDeclaration.typeArguments()[1]).integerLiteral().value();
94     }
95     return vectorLength;
96 }
97
98 static AST::NamedType& vectorInnerType(AST::NativeTypeDeclaration& nativeTypeDeclaration)
99 {
100     if (nativeTypeDeclaration.typeArguments().isEmpty())
101         return nativeTypeDeclaration;
102
103     ASSERT(nativeTypeDeclaration.typeArguments().size() == 2);
104     ASSERT(WTF::holds_alternative<Ref<AST::TypeReference>>(nativeTypeDeclaration.typeArguments()[0]));
105     return WTF::get<Ref<AST::TypeReference>>(nativeTypeDeclaration.typeArguments()[0])->resolvedType();
106 }
107
108 static const char* vectorSuffix(int vectorLength)
109 {
110     switch (vectorLength) {
111     case 1:
112         return "";
113     case 2:
114         return "2";
115     case 3:
116         return "3";
117     default:
118         ASSERT(vectorLength == 4);
119         return "4";
120     }
121 }
122
123 String writeNativeFunction(AST::NativeFunctionDeclaration& nativeFunctionDeclaration, String& outputFunctionName, Intrinsics& intrinsics, TypeNamer& typeNamer, const char* memsetZeroFunctionName)
124 {
125     StringBuilder stringBuilder;
126     if (nativeFunctionDeclaration.isCast()) {
127         auto& returnType = nativeFunctionDeclaration.type();
128         auto metalReturnName = typeNamer.mangledNameForType(returnType);
129         if (!nativeFunctionDeclaration.parameters().size()) {
130             stringBuilder.flexibleAppend(
131                 metalReturnName, ' ', outputFunctionName, "() {\n"
132                 "    ", metalReturnName, " x;\n"
133                 "    ", memsetZeroFunctionName, "(x);\n"
134                 "    return x;\n"
135                 "}\n"
136             );
137             return stringBuilder.toString();
138         }
139
140         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
141         auto& parameterType = *nativeFunctionDeclaration.parameters()[0]->type();
142         auto metalParameterName = typeNamer.mangledNameForType(parameterType);
143         stringBuilder.flexibleAppend(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n");
144
145         {
146             auto isEnumerationDefinition = [] (auto& type) {
147                 return is<AST::NamedType>(type) && is<AST::EnumerationDefinition>(downcast<AST::NamedType>(type));
148             };
149             auto& unifiedReturnType = returnType.unifyNode();
150             if (isEnumerationDefinition(unifiedReturnType) && !isEnumerationDefinition(parameterType.unifyNode())) { 
151                 auto& enumerationDefinition = downcast<AST::EnumerationDefinition>(downcast<AST::NamedType>(unifiedReturnType));
152                 stringBuilder.append("    switch (x) {\n");
153                 bool hasZeroCase = false;
154                 for (auto& member : enumerationDefinition.enumerationMembers()) {
155                     hasZeroCase |= !member.get().value();
156                     stringBuilder.flexibleAppend("        case ", member.get().value(), ": break;\n");
157                 }
158                 ASSERT_UNUSED(hasZeroCase, hasZeroCase);
159                 stringBuilder.append("        default: x = 0; break; }\n");
160             }
161         }
162
163         stringBuilder.flexibleAppend(
164             "    return static_cast<", metalReturnName, ">(x);\n"
165             "}\n"
166         );
167         return stringBuilder.toString();
168     }
169
170     if (nativeFunctionDeclaration.name() == "operator.value") {
171         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
172         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
173         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
174         stringBuilder.flexibleAppend(
175             metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"
176             "    return static_cast<", metalReturnName, ">(x);\n"
177             "}\n"
178         );
179         return stringBuilder.toString();
180     }
181
182     // FIXME: https://bugs.webkit.org/show_bug.cgi?id=198077 Authors can make a struct field named "length" too. Autogenerated getters for those shouldn't take this codepath.
183     if (nativeFunctionDeclaration.name() == "operator.length") {
184         ASSERT_UNUSED(intrinsics, matches(nativeFunctionDeclaration.type(), intrinsics.uintType()));
185         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
186         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
187         auto& parameterType = nativeFunctionDeclaration.parameters()[0]->type()->unifyNode();
188         auto& unnamedParameterType = downcast<AST::UnnamedType>(parameterType);
189         if (is<AST::ArrayType>(unnamedParameterType)) {
190             auto& arrayParameterType = downcast<AST::ArrayType>(unnamedParameterType);
191             stringBuilder.flexibleAppend(
192                 "uint ", outputFunctionName, '(', metalParameterName, ") {\n"
193                 "    return ", arrayParameterType.numElements(), ";\n"
194                 "}\n"
195             );
196             return stringBuilder.toString();
197         }
198
199         ASSERT(is<AST::ArrayReferenceType>(unnamedParameterType));
200         stringBuilder.flexibleAppend(
201             "uint ", outputFunctionName, '(', metalParameterName, " v) {\n"
202             "    return v.length;\n"
203             "}\n"
204         );
205         return stringBuilder.toString();
206     }
207
208     if (nativeFunctionDeclaration.name().startsWith("operator."_str)) {
209         auto mangledFieldName = [&](const String& fieldName) -> String {
210             auto& unifyNode = nativeFunctionDeclaration.parameters()[0]->type()->unifyNode();
211             auto& namedType = downcast<AST::NamedType>(unifyNode);
212             if (is<AST::StructureDefinition>(namedType)) {
213                 auto& structureDefinition = downcast<AST::StructureDefinition>(namedType);
214                 auto* structureElement = structureDefinition.find(fieldName);
215                 ASSERT(structureElement);
216                 return typeNamer.mangledNameForStructureElement(*structureElement);
217             }
218             ASSERT(is<AST::NativeTypeDeclaration>(namedType));
219             return fieldName;
220         };
221
222         if (nativeFunctionDeclaration.name().endsWith("=")) {
223             ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
224             auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
225             auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
226             auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
227             auto fieldName = nativeFunctionDeclaration.name().substring("operator."_str.length());
228             fieldName = fieldName.substring(0, fieldName.length() - 1);
229             auto metalFieldName = mangledFieldName(fieldName);
230             stringBuilder.flexibleAppend(
231                 metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " v, ", metalParameter2Name, " n) {\n"
232                 "    v.", metalFieldName, " = n;\n"
233                 "    return v;\n"
234                 "}\n"
235             );
236             return stringBuilder.toString();
237         }
238
239         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
240         auto fieldName = nativeFunctionDeclaration.name().substring("operator."_str.length());
241         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
242         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
243         auto metalFieldName = mangledFieldName(fieldName);
244         stringBuilder.flexibleAppend(
245             metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " v) {\n"
246             "    return v.", metalFieldName, ";\n"
247             "}\n"
248         );
249         return stringBuilder.toString();
250     }
251
252     if (nativeFunctionDeclaration.name().startsWith("operator&."_str)) {
253         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
254         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
255         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
256         auto fieldName = nativeFunctionDeclaration.name().substring("operator&."_str.length());
257
258         String metalFieldName;
259         auto& unnamedType = *nativeFunctionDeclaration.parameters()[0]->type();
260         auto& unifyNode = downcast<AST::PointerType>(unnamedType).elementType().unifyNode();
261         auto& namedType = downcast<AST::NamedType>(unifyNode);
262         if (is<AST::StructureDefinition>(namedType)) {
263             auto& structureDefinition = downcast<AST::StructureDefinition>(namedType);
264             auto* structureElement = structureDefinition.find(fieldName);
265             ASSERT(structureElement);
266             metalFieldName = typeNamer.mangledNameForStructureElement(*structureElement);
267         } else
268             metalFieldName = fieldName;
269
270         stringBuilder.flexibleAppend(
271             metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " v) {\n"
272             "    return &(v->", metalFieldName, ");\n"
273             "}\n"
274         );
275         return stringBuilder.toString();
276     }
277
278     if (nativeFunctionDeclaration.name() == "operator&[]") {
279         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
280         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
281         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
282         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
283         ASSERT(is<AST::ArrayReferenceType>(*nativeFunctionDeclaration.parameters()[0]->type()));
284         stringBuilder.flexibleAppend(
285             metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " v, ", metalParameter2Name, " n) {\n"
286             "    if (n < v.length) return &(v.pointer[n]);\n"
287             "    return nullptr;\n"
288             "}\n"
289         );
290         return stringBuilder.toString();
291     }
292
293     auto matrixDimension = [&] (unsigned typeArgumentIndex) -> unsigned {
294         auto& typeReference = downcast<AST::TypeReference>(*nativeFunctionDeclaration.parameters()[0]->type());
295         auto& matrixType = downcast<AST::NativeTypeDeclaration>(downcast<AST::TypeReference>(downcast<AST::TypeDefinition>(typeReference.resolvedType()).type()).resolvedType());
296         ASSERT(matrixType.name() == "matrix");
297         ASSERT(matrixType.typeArguments().size() == 3);
298         return WTF::get<AST::ConstantExpression>(matrixType.typeArguments()[typeArgumentIndex]).integerLiteral().value();
299     };
300     auto numberOfMatrixRows = [&] {
301         return matrixDimension(1);
302     };
303     auto numberOfMatrixColumns = [&] {
304         return matrixDimension(2);
305     };
306
307     if (nativeFunctionDeclaration.name() == "operator[]") {
308         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
309         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
310         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
311         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
312
313         unsigned numberOfRows = numberOfMatrixRows();
314         unsigned numberOfColumns = numberOfMatrixColumns();
315
316         stringBuilder.flexibleAppend(
317             metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " m, ", metalParameter2Name, " i) {\n"
318             "    if (i >= ", numberOfRows, ") return ", metalReturnName, "(0);\n"
319             "    ", metalReturnName, " result;\n"
320             "    result[0] = m[i];\n"
321             "    result[1] = m[i + ", numberOfRows, "];\n"
322         );
323         if (numberOfColumns >= 3)
324             stringBuilder.flexibleAppend("    result[2] = m[i + ", numberOfRows * 2, "];\n");
325         if (numberOfColumns >= 4)
326             stringBuilder.flexibleAppend("    result[3] = m[i + ", numberOfRows * 3, "];\n");
327         stringBuilder.append(
328             "    return result;\n"
329             "}\n"
330         );
331         return stringBuilder.toString();
332     }
333
334     if (nativeFunctionDeclaration.name() == "operator[]=") {
335         ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
336         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
337         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
338         auto metalParameter3Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[2]->type());
339         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
340
341         unsigned numberOfRows = numberOfMatrixRows();
342         unsigned numberOfColumns = numberOfMatrixColumns();
343
344         stringBuilder.flexibleAppend(
345             metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " m, ", metalParameter2Name, " i, ", metalParameter3Name, " v) {\n"
346             "    if (i >= ", numberOfRows, ") return m;\n"
347             "    m[i] = v[0];\n"
348             "    m[i + ", numberOfRows, "] = v[1];\n"
349         );
350         if (numberOfColumns >= 3)
351             stringBuilder.flexibleAppend("    m[i + ", numberOfRows * 2, "] = v[2];\n");
352         if (numberOfColumns >= 4)
353             stringBuilder.flexibleAppend("    m[i + ", numberOfRows * 3, "] = v[3];\n");
354         stringBuilder.append(
355             "    return m;"
356             "}\n"
357         );
358         return stringBuilder.toString();
359     }
360
361     if (nativeFunctionDeclaration.isOperator()) {
362         if (nativeFunctionDeclaration.parameters().size() == 1) {
363             auto operatorName = nativeFunctionDeclaration.name().substring("operator"_str.length());
364             auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
365             auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
366             stringBuilder.flexibleAppend(
367                 metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"
368                 "    return ", operatorName, "x;\n"
369                 "}\n"
370             );
371             return stringBuilder.toString();
372         }
373
374         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
375         auto operatorName = nativeFunctionDeclaration.name().substring("operator"_str.length());
376         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
377         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
378         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
379         stringBuilder.flexibleAppend(
380             metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " x, ", metalParameter2Name, " y) {\n"
381             "    return x ", operatorName, " y;\n"
382             "}\n"
383         );
384         return stringBuilder.toString();
385     }
386
387     if (nativeFunctionDeclaration.name() == "cos"
388         || nativeFunctionDeclaration.name() == "sin"
389         || nativeFunctionDeclaration.name() == "tan"
390         || nativeFunctionDeclaration.name() == "acos"
391         || nativeFunctionDeclaration.name() == "asin"
392         || nativeFunctionDeclaration.name() == "atan"
393         || nativeFunctionDeclaration.name() == "cosh"
394         || nativeFunctionDeclaration.name() == "sinh"
395         || nativeFunctionDeclaration.name() == "tanh"
396         || nativeFunctionDeclaration.name() == "ceil"
397         || nativeFunctionDeclaration.name() == "exp"
398         || nativeFunctionDeclaration.name() == "floor"
399         || nativeFunctionDeclaration.name() == "log"
400         || nativeFunctionDeclaration.name() == "round"
401         || nativeFunctionDeclaration.name() == "trunc"
402         || nativeFunctionDeclaration.name() == "ddx"
403         || nativeFunctionDeclaration.name() == "ddy"
404         || nativeFunctionDeclaration.name() == "isnormal"
405         || nativeFunctionDeclaration.name() == "isfinite"
406         || nativeFunctionDeclaration.name() == "isinf"
407         || nativeFunctionDeclaration.name() == "isnan"
408         || nativeFunctionDeclaration.name() == "asint"
409         || nativeFunctionDeclaration.name() == "asuint"
410         || nativeFunctionDeclaration.name() == "asfloat") {
411         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
412         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
413         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
414         stringBuilder.flexibleAppend(
415             metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"
416             "    return ", mapFunctionName(nativeFunctionDeclaration.name()), "(x);\n"
417             "}\n"
418         );
419         return stringBuilder.toString();
420     }
421
422     if (nativeFunctionDeclaration.name() == "pow" || nativeFunctionDeclaration.name() == "atan2") {
423         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
424         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
425         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
426         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
427         stringBuilder.flexibleAppend(
428             metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " x, ", metalParameter2Name, " y) {\n"
429             "    return ", nativeFunctionDeclaration.name(), "(x, y);\n"
430             "}\n"
431         );
432         return stringBuilder.toString();
433     }
434
435     if (nativeFunctionDeclaration.name() == "AllMemoryBarrierWithGroupSync") {
436         ASSERT(!nativeFunctionDeclaration.parameters().size());
437         stringBuilder.flexibleAppend(
438             "void ", outputFunctionName, "() {\n"
439             "    threadgroup_barrier(mem_flags::mem_device);\n"
440             "    threadgroup_barrier(mem_flags::mem_threadgroup);\n"
441             "    threadgroup_barrier(mem_flags::mem_texture);\n"
442             "}\n"
443         );
444         return stringBuilder.toString();
445     }
446
447     if (nativeFunctionDeclaration.name() == "DeviceMemoryBarrierWithGroupSync") {
448         ASSERT(!nativeFunctionDeclaration.parameters().size());
449         stringBuilder.flexibleAppend(
450             "void ", outputFunctionName, "() {\n"
451             "    threadgroup_barrier(mem_flags::mem_device);\n"
452             "}\n"
453         );
454         return stringBuilder.toString();
455     }
456
457     if (nativeFunctionDeclaration.name() == "GroupMemoryBarrierWithGroupSync") {
458         ASSERT(!nativeFunctionDeclaration.parameters().size());
459         stringBuilder.flexibleAppend(
460             "void ", outputFunctionName, "() {\n"
461             "    threadgroup_barrier(mem_flags::mem_threadgroup);\n"
462             "}\n"
463         );
464         return stringBuilder.toString();
465     }
466
467     if (nativeFunctionDeclaration.name().startsWith("Interlocked"_str)) {
468         if (nativeFunctionDeclaration.name() == "InterlockedCompareExchange") {
469             ASSERT(nativeFunctionDeclaration.parameters().size() == 4);
470             auto& firstArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0]->type());
471             auto firstArgumentAddressSpace = firstArgumentPointer.addressSpace();
472             auto firstArgumentPointee = typeNamer.mangledNameForType(firstArgumentPointer.elementType());
473             auto secondArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
474             auto thirdArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[2]->type());
475             auto& fourthArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[3]->type());
476             auto fourthArgumentAddressSpace = fourthArgumentPointer.addressSpace();
477             auto fourthArgumentPointee = typeNamer.mangledNameForType(fourthArgumentPointer.elementType());
478             stringBuilder.flexibleAppend(
479                 "void ", outputFunctionName, '(', toString(firstArgumentAddressSpace), ' ', firstArgumentPointee, "* object, ", secondArgument, " compare, ", thirdArgument, " desired, ", toString(fourthArgumentAddressSpace), ' ', fourthArgumentPointee, "* out) {\n"
480                 "    atomic_compare_exchange_weak_explicit(object, &compare, desired, memory_order_relaxed, memory_order_relaxed);\n"
481                 "    *out = compare;\n"
482                 "}\n"
483             );
484             return stringBuilder.toString();
485         }
486
487         ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
488         auto& firstArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0]->type());
489         auto firstArgumentAddressSpace = firstArgumentPointer.addressSpace();
490         auto firstArgumentPointee = typeNamer.mangledNameForType(firstArgumentPointer.elementType());
491         auto secondArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
492         auto& thirdArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[2]->type());
493         auto thirdArgumentAddressSpace = thirdArgumentPointer.addressSpace();
494         auto thirdArgumentPointee = typeNamer.mangledNameForType(thirdArgumentPointer.elementType());
495         auto name = atomicName(nativeFunctionDeclaration.name().substring("Interlocked"_str.length()));
496         stringBuilder.flexibleAppend(
497             "void ", outputFunctionName, '(', toString(firstArgumentAddressSpace), ' ', firstArgumentPointee, "* object, ", secondArgument, " operand, ", toString(thirdArgumentAddressSpace), ' ', thirdArgumentPointee, "* out) {\n"
498             "    *out = atomic_", name, "_explicit(object, operand, memory_order_relaxed);\n"
499             "}\n"
500         );
501         return stringBuilder.toString();
502     }
503
504     if (nativeFunctionDeclaration.name() == "Sample") {
505         ASSERT(nativeFunctionDeclaration.parameters().size() == 3 || nativeFunctionDeclaration.parameters().size() == 4);
506         
507         auto& textureType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[0]->type()->unifyNode()));
508         auto& locationType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[2]->type()->unifyNode()));
509         auto locationVectorLength = vectorLength(locationType);
510         auto& returnType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.type().unifyNode()));
511         auto returnVectorLength = vectorLength(returnType);
512
513         auto metalParameter1Name = typeNamer.mangledNameForType(textureType);
514         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
515         auto metalParameter3Name = typeNamer.mangledNameForType(locationType);
516         String metalParameter4Name;
517         if (nativeFunctionDeclaration.parameters().size() == 4)
518             metalParameter4Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[3]->type());
519         auto metalReturnName = typeNamer.mangledNameForType(returnType);
520         stringBuilder.flexibleAppend(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " theTexture, ", metalParameter2Name, " theSampler, ", metalParameter3Name, " location");
521         if (!metalParameter4Name.isNull())
522             stringBuilder.flexibleAppend(", ", metalParameter4Name, " offset");
523         stringBuilder.append(
524             ") {\n"
525             "    return theTexture.sample(theSampler, "
526         );
527         if (textureType.isTextureArray()) {
528             ASSERT(locationVectorLength > 1);
529             stringBuilder.flexibleAppend("location.", "xyzw"_str.substring(0, locationVectorLength - 1), ", location.", "xyzw"_str.substring(locationVectorLength - 1, 1));
530         } else
531             stringBuilder.append("location");
532         if (!metalParameter4Name.isNull())
533             stringBuilder.append(", offset");
534         stringBuilder.append(")");
535         if (!textureType.isDepthTexture())
536             stringBuilder.flexibleAppend(".", "xyzw"_str.substring(0, returnVectorLength));
537         stringBuilder.append(
538             ";\n"
539             "}\n"
540         );
541         return stringBuilder.toString();
542     }
543
544     if (nativeFunctionDeclaration.name() == "Load") {
545         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
546         
547         auto& textureType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[0]->type()->unifyNode()));
548         auto& locationType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[1]->type()->unifyNode()));
549         auto locationVectorLength = vectorLength(locationType);
550         auto& returnType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.type().unifyNode()));
551         auto returnVectorLength = vectorLength(returnType);
552
553         auto metalParameter1Name = typeNamer.mangledNameForType(textureType);
554         auto metalParameter2Name = typeNamer.mangledNameForType(locationType);
555         auto metalReturnName = typeNamer.mangledNameForType(returnType);
556         stringBuilder.flexibleAppend(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " theTexture, ", metalParameter2Name, " location) {\n");
557         if (textureType.isTextureArray()) {
558             ASSERT(locationVectorLength > 1);
559             String dimensions[] = { "width"_str, "height"_str, "depth"_str };
560             for (int i = 0; i < locationVectorLength - 1; ++i) {
561                 auto suffix = "xyzw"_str.substring(i, 1);
562                 stringBuilder.flexibleAppend("    if (location.", suffix, " < 0 || static_cast<uint32_t>(location.", suffix, ") >= theTexture.get_", dimensions[i], "()) return ", metalReturnName, "(0);\n");
563             }
564             auto suffix = "xyzw"_str.substring(locationVectorLength - 1, 1);
565             stringBuilder.flexibleAppend("    if (location.", suffix, " < 0 || static_cast<uint32_t>(location.", suffix, ") >= theTexture.get_array_size()) return ", metalReturnName, "(0);\n");
566         } else {
567             if (locationVectorLength == 1)
568                 stringBuilder.flexibleAppend("    if (location < 0 || static_cast<uint32_t>(location) >= theTexture.get_width()) return ", metalReturnName, "(0);\n");
569             else {
570                 stringBuilder.flexibleAppend(
571                     "    if (location.x < 0 || static_cast<uint32_t>(location.x) >= theTexture.get_width()) return ", metalReturnName, "(0);\n"
572                     "    if (location.y < 0 || static_cast<uint32_t>(location.y) >= theTexture.get_height()) return ", metalReturnName, "(0);\n"
573                 );
574                 if (locationVectorLength >= 3)
575                     stringBuilder.flexibleAppend("    if (location.z < 0 || static_cast<uint32_t>(location.z) >= theTexture.get_depth()) return ", metalReturnName, "(0);\n");
576             }
577         }
578         stringBuilder.append("    return theTexture.read(");
579         if (textureType.isTextureArray()) {
580             ASSERT(locationVectorLength > 1);
581             stringBuilder.flexibleAppend("uint", vectorSuffix(locationVectorLength - 1), "(location.", "xyzw"_str.substring(0, locationVectorLength - 1), "), uint(location.", "xyzw"_str.substring(locationVectorLength - 1, 1), ')');
582         } else
583             stringBuilder.flexibleAppend("uint", vectorSuffix(locationVectorLength), "(location)");
584         stringBuilder.append(')');
585         if (!textureType.isDepthTexture())
586             stringBuilder.flexibleAppend('.', "xyzw"_str.substring(0, returnVectorLength));
587         stringBuilder.append(
588             ";\n"
589             "}\n"
590         );
591         return stringBuilder.toString();
592     }
593
594     if (nativeFunctionDeclaration.name() == "load") {
595         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
596         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
597         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
598         stringBuilder.flexibleAppend(
599             metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"
600             "    return atomic_load_explicit(x, memory_order_relaxed);\n"
601             "}\n"
602         );
603         return stringBuilder.toString();
604     }
605
606     if (nativeFunctionDeclaration.name() == "store") {
607         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
608         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
609         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
610         stringBuilder.flexibleAppend("void ", outputFunctionName, '(', metalParameter1Name, " x, ", metalParameter2Name, " y) {\n"
611             "    atomic_store_explicit(x, y, memory_order_relaxed);\n"
612             "}\n");
613         return stringBuilder.toString();
614     }
615
616     if (nativeFunctionDeclaration.name() == "GetDimensions") {
617         auto& textureType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[0]->type()->unifyNode()));
618
619         size_t index = 1;
620         if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
621             ++index;
622         auto widthTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
623         ++index;
624         String heightTypeName;
625         if (textureType.textureDimension() >= 2) {
626             heightTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
627             ++index;
628         }
629         String depthTypeName;
630         if (textureType.textureDimension() >= 3) {
631             depthTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
632             ++index;
633         }
634         String elementsTypeName;
635         if (textureType.isTextureArray()) {
636             elementsTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
637             ++index;
638         }
639         String numberOfLevelsTypeName;
640         if (!textureType.isWritableTexture() && textureType.textureDimension() != 1) {
641             numberOfLevelsTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
642             ++index;
643         }
644         ASSERT(index == nativeFunctionDeclaration.parameters().size());
645
646         auto metalParameter1Name = typeNamer.mangledNameForType(textureType);
647         stringBuilder.flexibleAppend("void ", outputFunctionName, '(', metalParameter1Name, " theTexture");
648         if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
649             stringBuilder.append(", uint mipLevel");
650         stringBuilder.flexibleAppend(", ", widthTypeName, " width");
651         if (!heightTypeName.isNull())
652             stringBuilder.flexibleAppend(", ", heightTypeName, " height");
653         if (!depthTypeName.isNull())
654             stringBuilder.flexibleAppend(", ", depthTypeName, " depth");
655         if (!elementsTypeName.isNull())
656             stringBuilder.flexibleAppend(", ", elementsTypeName, " elements");
657         if (!numberOfLevelsTypeName.isNull())
658             stringBuilder.flexibleAppend(", ", numberOfLevelsTypeName, " numberOfLevels");
659         stringBuilder.append(
660             ") {\n"
661             "    if (width)\n"
662             "        *width = theTexture.get_width("
663         );
664         if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
665             stringBuilder.append("mipLevel");
666         stringBuilder.append(");\n");
667         if (!heightTypeName.isNull()) {
668             stringBuilder.append(
669                 "    if (height)\n"
670                 "        *height = theTexture.get_height("
671             );
672             if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
673                 stringBuilder.append("mipLevel");
674             stringBuilder.append(");\n");
675         }
676         if (!depthTypeName.isNull()) {
677             stringBuilder.append(
678                 "    if (depth)\n"
679                 "        *depth = theTexture.get_depth("
680             );
681             if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
682                 stringBuilder.append("mipLevel");
683             stringBuilder.append(");\n");
684         }
685         if (!elementsTypeName.isNull()) {
686             stringBuilder.append(
687                 "    if (elements)\n"
688                 "        *elements = theTexture.get_array_size();\n"
689             );
690         }
691         if (!numberOfLevelsTypeName.isNull()) {
692             stringBuilder.append(
693                 "    if (numberOfLevels)\n"
694                 "        *numberOfLevels = theTexture.get_num_mip_levels();\n"
695             );
696         }
697         stringBuilder.append("}\n");
698         return stringBuilder.toString();
699     }
700
701     if (nativeFunctionDeclaration.name() == "SampleBias") {
702         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
703         notImplemented();
704     }
705
706     if (nativeFunctionDeclaration.name() == "SampleGrad") {
707         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
708         notImplemented();
709     }
710
711     if (nativeFunctionDeclaration.name() == "SampleLevel") {
712         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
713         notImplemented();
714     }
715
716     if (nativeFunctionDeclaration.name() == "Gather") {
717         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
718         notImplemented();
719     }
720
721     if (nativeFunctionDeclaration.name() == "GatherRed") {
722         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
723         notImplemented();
724     }
725
726     if (nativeFunctionDeclaration.name() == "SampleCmp") {
727         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
728         notImplemented();
729     }
730
731     if (nativeFunctionDeclaration.name() == "SampleCmpLevelZero") {
732         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
733         notImplemented();
734     }
735
736     if (nativeFunctionDeclaration.name() == "Store") {
737         ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
738         
739         auto& textureType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[0]->type()->unifyNode()));
740         auto& itemType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[1]->type()->unifyNode()));
741         auto& itemVectorInnerType = vectorInnerType(itemType);
742         auto itemVectorLength = vectorLength(itemType);
743         auto& locationType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[2]->type()->unifyNode()));
744         auto locationVectorLength = vectorLength(locationType);
745
746         auto metalParameter1Name = typeNamer.mangledNameForType(textureType);
747         auto metalParameter2Name = typeNamer.mangledNameForType(itemType);
748         auto metalParameter3Name = typeNamer.mangledNameForType(locationType);
749         auto metalInnerTypeName = typeNamer.mangledNameForType(itemVectorInnerType);
750         stringBuilder.flexibleAppend("void ", outputFunctionName, '(', metalParameter1Name, " theTexture, ", metalParameter2Name, " item, ", metalParameter3Name, " location) {\n");
751         if (textureType.isTextureArray()) {
752             ASSERT(locationVectorLength > 1);
753             String dimensions[] = { "width"_str, "height"_str, "depth"_str };
754             for (int i = 0; i < locationVectorLength - 1; ++i) {
755                 auto suffix = "xyzw"_str.substring(i, 1);
756                 stringBuilder.flexibleAppend("    if (location.", suffix, " >= theTexture.get_", dimensions[i], "()) return;\n");
757             }
758             auto suffix = "xyzw"_str.substring(locationVectorLength - 1, 1);
759             stringBuilder.flexibleAppend("    if (location.", suffix, " >= theTexture.get_array_size()) return;\n");
760         } else {
761             if (locationVectorLength == 1)
762                 stringBuilder.append("    if (location >= theTexture.get_width()) return;\n");
763             else {
764                 stringBuilder.append(
765                     "    if (location.x >= theTexture.get_width()) return;\n"
766                     "    if (location.y >= theTexture.get_height()) return;\n"
767                 );
768                 if (locationVectorLength >= 3)
769                     stringBuilder.append("    if (location.z >= theTexture.get_depth()) return;\n");
770             }
771         }
772         stringBuilder.flexibleAppend("    theTexture.write(vec<", metalInnerTypeName, ", 4>(item");
773         for (int i = 0; i < 4 - itemVectorLength; ++i)
774             stringBuilder.append(", 0");
775         stringBuilder.append("), ");
776         if (textureType.isTextureArray()) {
777             ASSERT(locationVectorLength > 1);
778             stringBuilder.flexibleAppend("uint", vectorSuffix(locationVectorLength - 1), "(location.", "xyzw"_str.substring(0, locationVectorLength - 1), "), uint(location.", "xyzw"_str.substring(locationVectorLength - 1, 1), ')');
779         } else
780             stringBuilder.flexibleAppend("uint", vectorSuffix(locationVectorLength), "(location)");
781         stringBuilder.append(
782             ");\n"
783             "}\n"
784         );
785         return stringBuilder.toString();
786     }
787
788     if (nativeFunctionDeclaration.name() == "GatherAlpha") {
789         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
790         notImplemented();
791     }
792
793     if (nativeFunctionDeclaration.name() == "GatherBlue") {
794         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
795         notImplemented();
796     }
797
798     if (nativeFunctionDeclaration.name() == "GatherCmp") {
799         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
800         notImplemented();
801     }
802
803     if (nativeFunctionDeclaration.name() == "GatherCmpRed") {
804         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
805         notImplemented();
806     }
807
808     if (nativeFunctionDeclaration.name() == "GatherGreen") {
809         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
810         notImplemented();
811     }
812
813     ASSERT_NOT_REACHED();
814     return String();
815 }
816
817 } // namespace Metal
818
819 } // namespace WHLSL
820
821 } // namespace WebCore
822
823 #endif