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