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