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