Mangled WHLSL names don't need to allocate Strings
[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, MangledFunctionName outputFunctionName, Intrinsics& intrinsics, TypeNamer& typeNamer)
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                 "    return x;\n"
134                 "}\n"
135             );
136             return stringBuilder.toString();
137         }
138
139         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
140         auto& parameterType = *nativeFunctionDeclaration.parameters()[0]->type();
141         auto metalParameterName = typeNamer.mangledNameForType(parameterType);
142         stringBuilder.flexibleAppend(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n");
143
144         {
145             auto isEnumerationDefinition = [] (auto& type) {
146                 return is<AST::NamedType>(type) && is<AST::EnumerationDefinition>(downcast<AST::NamedType>(type));
147             };
148             auto& unifiedReturnType = returnType.unifyNode();
149             if (isEnumerationDefinition(unifiedReturnType) && !isEnumerationDefinition(parameterType.unifyNode())) { 
150                 auto& enumerationDefinition = downcast<AST::EnumerationDefinition>(downcast<AST::NamedType>(unifiedReturnType));
151                 stringBuilder.append("    switch (x) {\n");
152                 bool hasZeroCase = false;
153                 for (auto& member : enumerationDefinition.enumerationMembers()) {
154                     hasZeroCase |= !member.get().value();
155                     stringBuilder.flexibleAppend("        case ", member.get().value(), ": break;\n");
156                 }
157                 ASSERT_UNUSED(hasZeroCase, hasZeroCase);
158                 stringBuilder.append("        default: x = 0; break; }\n");
159             }
160         }
161
162         stringBuilder.flexibleAppend(
163             "    return static_cast<", metalReturnName, ">(x);\n"
164             "}\n"
165         );
166         return stringBuilder.toString();
167     }
168
169     if (nativeFunctionDeclaration.name() == "operator.value") {
170         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
171         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
172         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
173         stringBuilder.flexibleAppend(
174             metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"
175             "    return static_cast<", metalReturnName, ">(x);\n"
176             "}\n"
177         );
178         return stringBuilder.toString();
179     }
180
181     // 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.
182     if (nativeFunctionDeclaration.name() == "operator.length") {
183         ASSERT_UNUSED(intrinsics, matches(nativeFunctionDeclaration.type(), intrinsics.uintType()));
184         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
185         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
186         auto& parameterType = nativeFunctionDeclaration.parameters()[0]->type()->unifyNode();
187         auto& unnamedParameterType = downcast<AST::UnnamedType>(parameterType);
188         if (is<AST::ArrayType>(unnamedParameterType)) {
189             auto& arrayParameterType = downcast<AST::ArrayType>(unnamedParameterType);
190             stringBuilder.flexibleAppend(
191                 "uint ", outputFunctionName, '(', metalParameterName, ") {\n"
192                 "    return ", arrayParameterType.numElements(), ";\n"
193                 "}\n"
194             );
195             return stringBuilder.toString();
196         }
197
198         ASSERT(is<AST::ArrayReferenceType>(unnamedParameterType));
199         stringBuilder.flexibleAppend(
200             "uint ", outputFunctionName, '(', metalParameterName, " v) {\n"
201             "    return v.length;\n"
202             "}\n"
203         );
204         return stringBuilder.toString();
205     }
206
207     if (nativeFunctionDeclaration.name().startsWith("operator."_str)) {
208         auto appendMangledFieldName = [&](StringBuilder& stringBuilder, const String& fieldName) {
209             auto& unifyNode = nativeFunctionDeclaration.parameters()[0]->type()->unifyNode();
210             auto& namedType = downcast<AST::NamedType>(unifyNode);
211             if (is<AST::StructureDefinition>(namedType)) {
212                 auto& structureDefinition = downcast<AST::StructureDefinition>(namedType);
213                 auto* structureElement = structureDefinition.find(fieldName);
214                 ASSERT(structureElement);
215                 stringBuilder.flexibleAppend(typeNamer.mangledNameForStructureElement(*structureElement));
216                 return;
217             }
218             ASSERT(is<AST::NativeTypeDeclaration>(namedType));
219             stringBuilder.append(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             stringBuilder.flexibleAppend(
228                 metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " v, ", metalParameter2Name, " n) {\n"
229                 "    v."
230             );
231
232             auto fieldName = nativeFunctionDeclaration.name().substring("operator."_str.length());
233             fieldName = fieldName.substring(0, fieldName.length() - 1);
234             appendMangledFieldName(stringBuilder, fieldName);
235
236             stringBuilder.append(" = n;\n"
237                 "    return v;\n"
238                 "}\n"
239             );
240             return stringBuilder.toString();
241         }
242
243         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
244         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
245         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
246         stringBuilder.flexibleAppend(
247             metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " v) {\n"
248             "    return v."
249         );
250
251         auto fieldName = nativeFunctionDeclaration.name().substring("operator."_str.length());
252         appendMangledFieldName(stringBuilder, fieldName);
253             
254         stringBuilder.append(";\n"
255             "}\n"
256         );
257         return stringBuilder.toString();
258     }
259
260     if (nativeFunctionDeclaration.name().startsWith("operator&."_str)) {
261         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
262         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
263         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
264
265         stringBuilder.flexibleAppend(
266             metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " v) {\n"
267             "    return &(v->"
268         );
269
270         auto fieldName = nativeFunctionDeclaration.name().substring("operator&."_str.length());
271         auto& unnamedType = *nativeFunctionDeclaration.parameters()[0]->type();
272         auto& unifyNode = downcast<AST::PointerType>(unnamedType).elementType().unifyNode();
273         auto& namedType = downcast<AST::NamedType>(unifyNode);
274         if (is<AST::StructureDefinition>(namedType)) {
275             auto& structureDefinition = downcast<AST::StructureDefinition>(namedType);
276             auto* structureElement = structureDefinition.find(fieldName);
277             ASSERT(structureElement);
278             stringBuilder.flexibleAppend(typeNamer.mangledNameForStructureElement(*structureElement));
279         } else
280             stringBuilder.flexibleAppend(fieldName);
281
282         stringBuilder.append(
283             ");\n"
284             "}\n"
285         );
286         return stringBuilder.toString();
287     }
288
289     if (nativeFunctionDeclaration.name() == "operator&[]") {
290         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
291         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
292         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
293         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
294         ASSERT(is<AST::ArrayReferenceType>(*nativeFunctionDeclaration.parameters()[0]->type()));
295         stringBuilder.flexibleAppend(
296             metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " v, ", metalParameter2Name, " n) {\n"
297             "    if (n < v.length) return &(v.pointer[n]);\n"
298             "    return nullptr;\n"
299             "}\n"
300         );
301         return stringBuilder.toString();
302     }
303
304     auto matrixDimension = [&] (unsigned typeArgumentIndex) -> unsigned {
305         auto& typeReference = downcast<AST::TypeReference>(*nativeFunctionDeclaration.parameters()[0]->type());
306         auto& matrixType = downcast<AST::NativeTypeDeclaration>(downcast<AST::TypeReference>(downcast<AST::TypeDefinition>(typeReference.resolvedType()).type()).resolvedType());
307         ASSERT(matrixType.name() == "matrix");
308         ASSERT(matrixType.typeArguments().size() == 3);
309         return WTF::get<AST::ConstantExpression>(matrixType.typeArguments()[typeArgumentIndex]).integerLiteral().value();
310     };
311     auto numberOfMatrixRows = [&] {
312         return matrixDimension(1);
313     };
314     auto numberOfMatrixColumns = [&] {
315         return matrixDimension(2);
316     };
317
318     if (nativeFunctionDeclaration.name() == "operator[]") {
319         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
320         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
321         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
322         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
323
324         unsigned numberOfRows = numberOfMatrixRows();
325         unsigned numberOfColumns = numberOfMatrixColumns();
326
327         stringBuilder.flexibleAppend(
328             metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " m, ", metalParameter2Name, " i) {\n"
329             "    if (i >= ", numberOfRows, ") return ", metalReturnName, "(0);\n"
330             "    ", metalReturnName, " result;\n"
331             "    result[0] = m[i];\n"
332             "    result[1] = m[i + ", numberOfRows, "];\n"
333         );
334         if (numberOfColumns >= 3)
335             stringBuilder.flexibleAppend("    result[2] = m[i + ", numberOfRows * 2, "];\n");
336         if (numberOfColumns >= 4)
337             stringBuilder.flexibleAppend("    result[3] = m[i + ", numberOfRows * 3, "];\n");
338         stringBuilder.append(
339             "    return result;\n"
340             "}\n"
341         );
342         return stringBuilder.toString();
343     }
344
345     if (nativeFunctionDeclaration.name() == "operator[]=") {
346         ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
347         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
348         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
349         auto metalParameter3Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[2]->type());
350         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
351
352         unsigned numberOfRows = numberOfMatrixRows();
353         unsigned numberOfColumns = numberOfMatrixColumns();
354
355         stringBuilder.flexibleAppend(
356             metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " m, ", metalParameter2Name, " i, ", metalParameter3Name, " v) {\n"
357             "    if (i >= ", numberOfRows, ") return m;\n"
358             "    m[i] = v[0];\n"
359             "    m[i + ", numberOfRows, "] = v[1];\n"
360         );
361         if (numberOfColumns >= 3)
362             stringBuilder.flexibleAppend("    m[i + ", numberOfRows * 2, "] = v[2];\n");
363         if (numberOfColumns >= 4)
364             stringBuilder.flexibleAppend("    m[i + ", numberOfRows * 3, "] = v[3];\n");
365         stringBuilder.append(
366             "    return m;"
367             "}\n"
368         );
369         return stringBuilder.toString();
370     }
371
372     if (nativeFunctionDeclaration.isOperator()) {
373         if (nativeFunctionDeclaration.parameters().size() == 1) {
374             auto operatorName = nativeFunctionDeclaration.name().substring("operator"_str.length());
375             auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
376             auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
377             stringBuilder.flexibleAppend(
378                 metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"
379                 "    return ", operatorName, "x;\n"
380                 "}\n"
381             );
382             return stringBuilder.toString();
383         }
384
385         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
386         auto operatorName = nativeFunctionDeclaration.name().substring("operator"_str.length());
387         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
388         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
389         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
390         stringBuilder.flexibleAppend(
391             metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " x, ", metalParameter2Name, " y) {\n"
392             "    return x ", operatorName, " y;\n"
393             "}\n"
394         );
395         return stringBuilder.toString();
396     }
397
398     if (nativeFunctionDeclaration.name() == "cos"
399         || nativeFunctionDeclaration.name() == "sin"
400         || nativeFunctionDeclaration.name() == "tan"
401         || nativeFunctionDeclaration.name() == "acos"
402         || nativeFunctionDeclaration.name() == "asin"
403         || nativeFunctionDeclaration.name() == "atan"
404         || nativeFunctionDeclaration.name() == "cosh"
405         || nativeFunctionDeclaration.name() == "sinh"
406         || nativeFunctionDeclaration.name() == "tanh"
407         || nativeFunctionDeclaration.name() == "ceil"
408         || nativeFunctionDeclaration.name() == "exp"
409         || nativeFunctionDeclaration.name() == "floor"
410         || nativeFunctionDeclaration.name() == "log"
411         || nativeFunctionDeclaration.name() == "round"
412         || nativeFunctionDeclaration.name() == "trunc"
413         || nativeFunctionDeclaration.name() == "ddx"
414         || nativeFunctionDeclaration.name() == "ddy"
415         || nativeFunctionDeclaration.name() == "isnormal"
416         || nativeFunctionDeclaration.name() == "isfinite"
417         || nativeFunctionDeclaration.name() == "isinf"
418         || nativeFunctionDeclaration.name() == "isnan"
419         || nativeFunctionDeclaration.name() == "asint"
420         || nativeFunctionDeclaration.name() == "asuint"
421         || nativeFunctionDeclaration.name() == "asfloat") {
422         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
423         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
424         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
425         stringBuilder.flexibleAppend(
426             metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"
427             "    return ", mapFunctionName(nativeFunctionDeclaration.name()), "(x);\n"
428             "}\n"
429         );
430         return stringBuilder.toString();
431     }
432
433     if (nativeFunctionDeclaration.name() == "pow" || nativeFunctionDeclaration.name() == "atan2") {
434         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
435         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
436         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
437         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
438         stringBuilder.flexibleAppend(
439             metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " x, ", metalParameter2Name, " y) {\n"
440             "    return ", nativeFunctionDeclaration.name(), "(x, y);\n"
441             "}\n"
442         );
443         return stringBuilder.toString();
444     }
445
446     if (nativeFunctionDeclaration.name() == "AllMemoryBarrierWithGroupSync") {
447         ASSERT(!nativeFunctionDeclaration.parameters().size());
448         stringBuilder.flexibleAppend(
449             "void ", outputFunctionName, "() {\n"
450             "    threadgroup_barrier(mem_flags::mem_device);\n"
451             "    threadgroup_barrier(mem_flags::mem_threadgroup);\n"
452             "    threadgroup_barrier(mem_flags::mem_texture);\n"
453             "}\n"
454         );
455         return stringBuilder.toString();
456     }
457
458     if (nativeFunctionDeclaration.name() == "DeviceMemoryBarrierWithGroupSync") {
459         ASSERT(!nativeFunctionDeclaration.parameters().size());
460         stringBuilder.flexibleAppend(
461             "void ", outputFunctionName, "() {\n"
462             "    threadgroup_barrier(mem_flags::mem_device);\n"
463             "}\n"
464         );
465         return stringBuilder.toString();
466     }
467
468     if (nativeFunctionDeclaration.name() == "GroupMemoryBarrierWithGroupSync") {
469         ASSERT(!nativeFunctionDeclaration.parameters().size());
470         stringBuilder.flexibleAppend(
471             "void ", outputFunctionName, "() {\n"
472             "    threadgroup_barrier(mem_flags::mem_threadgroup);\n"
473             "}\n"
474         );
475         return stringBuilder.toString();
476     }
477
478     if (nativeFunctionDeclaration.name().startsWith("Interlocked"_str)) {
479         if (nativeFunctionDeclaration.name() == "InterlockedCompareExchange") {
480             ASSERT(nativeFunctionDeclaration.parameters().size() == 4);
481             auto& firstArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0]->type());
482             auto firstArgumentAddressSpace = firstArgumentPointer.addressSpace();
483             auto firstArgumentPointee = typeNamer.mangledNameForType(firstArgumentPointer.elementType());
484             auto secondArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
485             auto thirdArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[2]->type());
486             auto& fourthArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[3]->type());
487             auto fourthArgumentAddressSpace = fourthArgumentPointer.addressSpace();
488             auto fourthArgumentPointee = typeNamer.mangledNameForType(fourthArgumentPointer.elementType());
489             stringBuilder.flexibleAppend(
490                 "void ", outputFunctionName, '(', toString(firstArgumentAddressSpace), ' ', firstArgumentPointee, "* object, ", secondArgument, " compare, ", thirdArgument, " desired, ", toString(fourthArgumentAddressSpace), ' ', fourthArgumentPointee, "* out) {\n"
491                 "    atomic_compare_exchange_weak_explicit(object, &compare, desired, memory_order_relaxed, memory_order_relaxed);\n"
492                 "    *out = compare;\n"
493                 "}\n"
494             );
495             return stringBuilder.toString();
496         }
497
498         ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
499         auto& firstArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0]->type());
500         auto firstArgumentAddressSpace = firstArgumentPointer.addressSpace();
501         auto firstArgumentPointee = typeNamer.mangledNameForType(firstArgumentPointer.elementType());
502         auto secondArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
503         auto& thirdArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[2]->type());
504         auto thirdArgumentAddressSpace = thirdArgumentPointer.addressSpace();
505         auto thirdArgumentPointee = typeNamer.mangledNameForType(thirdArgumentPointer.elementType());
506         auto name = atomicName(nativeFunctionDeclaration.name().substring("Interlocked"_str.length()));
507         stringBuilder.flexibleAppend(
508             "void ", outputFunctionName, '(', toString(firstArgumentAddressSpace), ' ', firstArgumentPointee, "* object, ", secondArgument, " operand, ", toString(thirdArgumentAddressSpace), ' ', thirdArgumentPointee, "* out) {\n"
509             "    *out = atomic_", name, "_explicit(object, operand, memory_order_relaxed);\n"
510             "}\n"
511         );
512         return stringBuilder.toString();
513     }
514
515     if (nativeFunctionDeclaration.name() == "Sample") {
516         ASSERT(nativeFunctionDeclaration.parameters().size() == 3 || nativeFunctionDeclaration.parameters().size() == 4);
517         
518         auto& textureType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[0]->type()->unifyNode()));
519         auto& locationType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[2]->type()->unifyNode()));
520         auto locationVectorLength = vectorLength(locationType);
521         auto& returnType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.type().unifyNode()));
522         auto returnVectorLength = vectorLength(returnType);
523
524         auto metalParameter1Name = typeNamer.mangledNameForType(textureType);
525         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
526         auto metalParameter3Name = typeNamer.mangledNameForType(locationType);
527         Optional<MangledTypeName> metalParameter4Name;
528         if (nativeFunctionDeclaration.parameters().size() == 4)
529             metalParameter4Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[3]->type());
530         auto metalReturnName = typeNamer.mangledNameForType(returnType);
531         stringBuilder.flexibleAppend(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " theTexture, ", metalParameter2Name, " theSampler, ", metalParameter3Name, " location");
532         if (metalParameter4Name)
533             stringBuilder.flexibleAppend(", ", *metalParameter4Name, " offset");
534         stringBuilder.append(
535             ") {\n"
536             "    return theTexture.sample(theSampler, "
537         );
538         if (textureType.isTextureArray()) {
539             ASSERT(locationVectorLength > 1);
540             stringBuilder.flexibleAppend("location.", "xyzw"_str.substring(0, locationVectorLength - 1), ", location.", "xyzw"_str.substring(locationVectorLength - 1, 1));
541         } else
542             stringBuilder.append("location");
543         if (metalParameter4Name)
544             stringBuilder.append(", offset");
545         stringBuilder.append(")");
546         if (!textureType.isDepthTexture())
547             stringBuilder.flexibleAppend(".", "xyzw"_str.substring(0, returnVectorLength));
548         stringBuilder.append(
549             ";\n"
550             "}\n"
551         );
552         return stringBuilder.toString();
553     }
554
555     if (nativeFunctionDeclaration.name() == "Load") {
556         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
557         
558         auto& textureType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[0]->type()->unifyNode()));
559         auto& locationType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[1]->type()->unifyNode()));
560         auto locationVectorLength = vectorLength(locationType);
561         auto& returnType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.type().unifyNode()));
562         auto returnVectorLength = vectorLength(returnType);
563
564         auto metalParameter1Name = typeNamer.mangledNameForType(textureType);
565         auto metalParameter2Name = typeNamer.mangledNameForType(locationType);
566         auto metalReturnName = typeNamer.mangledNameForType(returnType);
567         stringBuilder.flexibleAppend(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " theTexture, ", metalParameter2Name, " location) {\n");
568         if (textureType.isTextureArray()) {
569             ASSERT(locationVectorLength > 1);
570             String dimensions[] = { "width"_str, "height"_str, "depth"_str };
571             for (int i = 0; i < locationVectorLength - 1; ++i) {
572                 auto suffix = "xyzw"_str.substring(i, 1);
573                 stringBuilder.flexibleAppend("    if (location.", suffix, " < 0 || static_cast<uint32_t>(location.", suffix, ") >= theTexture.get_", dimensions[i], "()) return ", metalReturnName, "(0);\n");
574             }
575             auto suffix = "xyzw"_str.substring(locationVectorLength - 1, 1);
576             stringBuilder.flexibleAppend("    if (location.", suffix, " < 0 || static_cast<uint32_t>(location.", suffix, ") >= theTexture.get_array_size()) return ", metalReturnName, "(0);\n");
577         } else {
578             if (locationVectorLength == 1)
579                 stringBuilder.flexibleAppend("    if (location < 0 || static_cast<uint32_t>(location) >= theTexture.get_width()) return ", metalReturnName, "(0);\n");
580             else {
581                 stringBuilder.flexibleAppend(
582                     "    if (location.x < 0 || static_cast<uint32_t>(location.x) >= theTexture.get_width()) return ", metalReturnName, "(0);\n"
583                     "    if (location.y < 0 || static_cast<uint32_t>(location.y) >= theTexture.get_height()) return ", metalReturnName, "(0);\n"
584                 );
585                 if (locationVectorLength >= 3)
586                     stringBuilder.flexibleAppend("    if (location.z < 0 || static_cast<uint32_t>(location.z) >= theTexture.get_depth()) return ", metalReturnName, "(0);\n");
587             }
588         }
589         stringBuilder.append("    return theTexture.read(");
590         if (textureType.isTextureArray()) {
591             ASSERT(locationVectorLength > 1);
592             stringBuilder.flexibleAppend("uint", vectorSuffix(locationVectorLength - 1), "(location.", "xyzw"_str.substring(0, locationVectorLength - 1), "), uint(location.", "xyzw"_str.substring(locationVectorLength - 1, 1), ')');
593         } else
594             stringBuilder.flexibleAppend("uint", vectorSuffix(locationVectorLength), "(location)");
595         stringBuilder.append(')');
596         if (!textureType.isDepthTexture())
597             stringBuilder.flexibleAppend('.', "xyzw"_str.substring(0, returnVectorLength));
598         stringBuilder.append(
599             ";\n"
600             "}\n"
601         );
602         return stringBuilder.toString();
603     }
604
605     if (nativeFunctionDeclaration.name() == "load") {
606         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
607         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
608         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
609         stringBuilder.flexibleAppend(
610             metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"
611             "    return atomic_load_explicit(x, memory_order_relaxed);\n"
612             "}\n"
613         );
614         return stringBuilder.toString();
615     }
616
617     if (nativeFunctionDeclaration.name() == "store") {
618         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
619         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
620         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
621         stringBuilder.flexibleAppend("void ", outputFunctionName, '(', metalParameter1Name, " x, ", metalParameter2Name, " y) {\n"
622             "    atomic_store_explicit(x, y, memory_order_relaxed);\n"
623             "}\n");
624         return stringBuilder.toString();
625     }
626
627     if (nativeFunctionDeclaration.name() == "GetDimensions") {
628         auto& textureType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[0]->type()->unifyNode()));
629
630         size_t index = 1;
631         if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
632             ++index;
633         auto widthTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
634         ++index;
635         Optional<MangledTypeName> heightTypeName;
636         if (textureType.textureDimension() >= 2) {
637             heightTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
638             ++index;
639         }
640         Optional<MangledTypeName> depthTypeName;
641         if (textureType.textureDimension() >= 3) {
642             depthTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
643             ++index;
644         }
645         Optional<MangledTypeName> elementsTypeName;
646         if (textureType.isTextureArray()) {
647             elementsTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
648             ++index;
649         }
650         Optional<MangledTypeName> numberOfLevelsTypeName;
651         if (!textureType.isWritableTexture() && textureType.textureDimension() != 1) {
652             numberOfLevelsTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
653             ++index;
654         }
655         ASSERT(index == nativeFunctionDeclaration.parameters().size());
656
657         auto metalParameter1Name = typeNamer.mangledNameForType(textureType);
658         stringBuilder.flexibleAppend("void ", outputFunctionName, '(', metalParameter1Name, " theTexture");
659         if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
660             stringBuilder.append(", uint mipLevel");
661         stringBuilder.flexibleAppend(", ", widthTypeName, " width");
662         if (heightTypeName)
663             stringBuilder.flexibleAppend(", ", *heightTypeName, " height");
664         if (depthTypeName)
665             stringBuilder.flexibleAppend(", ", *depthTypeName, " depth");
666         if (elementsTypeName)
667             stringBuilder.flexibleAppend(", ", *elementsTypeName, " elements");
668         if (numberOfLevelsTypeName)
669             stringBuilder.flexibleAppend(", ", *numberOfLevelsTypeName, " numberOfLevels");
670         stringBuilder.append(
671             ") {\n"
672             "    if (width)\n"
673             "        *width = theTexture.get_width("
674         );
675         if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
676             stringBuilder.append("mipLevel");
677         stringBuilder.append(");\n");
678         if (heightTypeName) {
679             stringBuilder.append(
680                 "    if (height)\n"
681                 "        *height = theTexture.get_height("
682             );
683             if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
684                 stringBuilder.append("mipLevel");
685             stringBuilder.append(");\n");
686         }
687         if (depthTypeName) {
688             stringBuilder.append(
689                 "    if (depth)\n"
690                 "        *depth = theTexture.get_depth("
691             );
692             if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
693                 stringBuilder.append("mipLevel");
694             stringBuilder.append(");\n");
695         }
696         if (elementsTypeName) {
697             stringBuilder.append(
698                 "    if (elements)\n"
699                 "        *elements = theTexture.get_array_size();\n"
700             );
701         }
702         if (numberOfLevelsTypeName) {
703             stringBuilder.append(
704                 "    if (numberOfLevels)\n"
705                 "        *numberOfLevels = theTexture.get_num_mip_levels();\n"
706             );
707         }
708         stringBuilder.append("}\n");
709         return stringBuilder.toString();
710     }
711
712     if (nativeFunctionDeclaration.name() == "SampleBias") {
713         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
714         notImplemented();
715     }
716
717     if (nativeFunctionDeclaration.name() == "SampleGrad") {
718         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
719         notImplemented();
720     }
721
722     if (nativeFunctionDeclaration.name() == "SampleLevel") {
723         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
724         notImplemented();
725     }
726
727     if (nativeFunctionDeclaration.name() == "Gather") {
728         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
729         notImplemented();
730     }
731
732     if (nativeFunctionDeclaration.name() == "GatherRed") {
733         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
734         notImplemented();
735     }
736
737     if (nativeFunctionDeclaration.name() == "SampleCmp") {
738         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
739         notImplemented();
740     }
741
742     if (nativeFunctionDeclaration.name() == "SampleCmpLevelZero") {
743         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
744         notImplemented();
745     }
746
747     if (nativeFunctionDeclaration.name() == "Store") {
748         ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
749         
750         auto& textureType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[0]->type()->unifyNode()));
751         auto& itemType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[1]->type()->unifyNode()));
752         auto& itemVectorInnerType = vectorInnerType(itemType);
753         auto itemVectorLength = vectorLength(itemType);
754         auto& locationType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[2]->type()->unifyNode()));
755         auto locationVectorLength = vectorLength(locationType);
756
757         auto metalParameter1Name = typeNamer.mangledNameForType(textureType);
758         auto metalParameter2Name = typeNamer.mangledNameForType(itemType);
759         auto metalParameter3Name = typeNamer.mangledNameForType(locationType);
760         auto metalInnerTypeName = typeNamer.mangledNameForType(itemVectorInnerType);
761         stringBuilder.flexibleAppend("void ", outputFunctionName, '(', metalParameter1Name, " theTexture, ", metalParameter2Name, " item, ", metalParameter3Name, " location) {\n");
762         if (textureType.isTextureArray()) {
763             ASSERT(locationVectorLength > 1);
764             String dimensions[] = { "width"_str, "height"_str, "depth"_str };
765             for (int i = 0; i < locationVectorLength - 1; ++i) {
766                 auto suffix = "xyzw"_str.substring(i, 1);
767                 stringBuilder.flexibleAppend("    if (location.", suffix, " >= theTexture.get_", dimensions[i], "()) return;\n");
768             }
769             auto suffix = "xyzw"_str.substring(locationVectorLength - 1, 1);
770             stringBuilder.flexibleAppend("    if (location.", suffix, " >= theTexture.get_array_size()) return;\n");
771         } else {
772             if (locationVectorLength == 1)
773                 stringBuilder.append("    if (location >= theTexture.get_width()) return;\n");
774             else {
775                 stringBuilder.append(
776                     "    if (location.x >= theTexture.get_width()) return;\n"
777                     "    if (location.y >= theTexture.get_height()) return;\n"
778                 );
779                 if (locationVectorLength >= 3)
780                     stringBuilder.append("    if (location.z >= theTexture.get_depth()) return;\n");
781             }
782         }
783         stringBuilder.flexibleAppend("    theTexture.write(vec<", metalInnerTypeName, ", 4>(item");
784         for (int i = 0; i < 4 - itemVectorLength; ++i)
785             stringBuilder.append(", 0");
786         stringBuilder.append("), ");
787         if (textureType.isTextureArray()) {
788             ASSERT(locationVectorLength > 1);
789             stringBuilder.flexibleAppend("uint", vectorSuffix(locationVectorLength - 1), "(location.", "xyzw"_str.substring(0, locationVectorLength - 1), "), uint(location.", "xyzw"_str.substring(locationVectorLength - 1, 1), ')');
790         } else
791             stringBuilder.flexibleAppend("uint", vectorSuffix(locationVectorLength), "(location)");
792         stringBuilder.append(
793             ");\n"
794             "}\n"
795         );
796         return stringBuilder.toString();
797     }
798
799     if (nativeFunctionDeclaration.name() == "GatherAlpha") {
800         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
801         notImplemented();
802     }
803
804     if (nativeFunctionDeclaration.name() == "GatherBlue") {
805         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
806         notImplemented();
807     }
808
809     if (nativeFunctionDeclaration.name() == "GatherCmp") {
810         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
811         notImplemented();
812     }
813
814     if (nativeFunctionDeclaration.name() == "GatherCmpRed") {
815         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
816         notImplemented();
817     }
818
819     if (nativeFunctionDeclaration.name() == "GatherGreen") {
820         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
821         notImplemented();
822     }
823
824     ASSERT_NOT_REACHED();
825     return String();
826 }
827
828 } // namespace Metal
829
830 } // namespace WHLSL
831
832 } // namespace WebCore
833
834 #endif