3aec9ff80e4d654fa8e52704759cae764d809e82
[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 "WHLSLInferTypes.h"
35 #include "WHLSLIntrinsics.h"
36 #include "WHLSLNamedType.h"
37 #include "WHLSLNativeFunctionDeclaration.h"
38 #include "WHLSLNativeTypeDeclaration.h"
39 #include "WHLSLPointerType.h"
40 #include "WHLSLStructureDefinition.h"
41 #include "WHLSLTypeDefinition.h"
42 #include "WHLSLTypeNamer.h"
43 #include "WHLSLUnnamedType.h"
44 #include "WHLSLVariableDeclaration.h"
45 #include <wtf/text/StringBuilder.h>
46
47 namespace WebCore {
48
49 namespace WHLSL {
50
51 namespace Metal {
52
53 static String mapFunctionName(String& functionName)
54 {
55     if (functionName == "ddx")
56         return "dfdx"_str;
57     if (functionName == "ddy")
58         return "dfdy"_str;
59     if (functionName == "asint")
60         return "as_type<int32_t>"_str;
61     if (functionName == "asuint")
62         return "as_type<uint32_t>"_str;
63     if (functionName == "asfloat")
64         return "as_type<float>"_str;
65     return functionName;
66 }
67
68 static String atomicName(String input)
69 {
70     if (input == "Add")
71         return "fetch_add"_str;
72     if (input == "And")
73         return "fetch_and"_str;
74     if (input == "Exchange")
75         return "exchange"_str;
76     if (input == "Max")
77         return "fetch_max"_str;
78     if (input == "Min")
79         return "fetch_min"_str;
80     if (input == "Or")
81         return "fetch_or"_str;
82     ASSERT(input == "Xor");
83         return "fetch_xor"_str;
84 }
85
86 static int vectorLength(AST::NativeTypeDeclaration& nativeTypeDeclaration)
87 {
88     int vectorLength = 1;
89     if (!nativeTypeDeclaration.typeArguments().isEmpty()) {
90         ASSERT(nativeTypeDeclaration.typeArguments().size() == 2);
91         ASSERT(WTF::holds_alternative<AST::ConstantExpression>(nativeTypeDeclaration.typeArguments()[1]));
92         vectorLength = WTF::get<AST::ConstantExpression>(nativeTypeDeclaration.typeArguments()[1]).integerLiteral().value();
93     }
94     return vectorLength;
95 }
96
97 static AST::NamedType& vectorInnerType(AST::NativeTypeDeclaration& nativeTypeDeclaration)
98 {
99     if (nativeTypeDeclaration.typeArguments().isEmpty())
100         return nativeTypeDeclaration;
101
102     ASSERT(nativeTypeDeclaration.typeArguments().size() == 2);
103     ASSERT(WTF::holds_alternative<UniqueRef<AST::TypeReference>>(nativeTypeDeclaration.typeArguments()[0]));
104     return WTF::get<UniqueRef<AST::TypeReference>>(nativeTypeDeclaration.typeArguments()[0])->resolvedType();
105 }
106
107 static const char* vectorSuffix(int vectorLength)
108 {
109     switch (vectorLength) {
110     case 1:
111         return "";
112     case 2:
113         return "2";
114     case 3:
115         return "3";
116     default:
117         ASSERT(vectorLength == 4);
118         return "4";
119     }
120 }
121
122 String writeNativeFunction(AST::NativeFunctionDeclaration& nativeFunctionDeclaration, String& outputFunctionName, Intrinsics& intrinsics, TypeNamer& typeNamer, const char* memsetZeroFunctionName)
123 {
124     StringBuilder stringBuilder;
125     if (nativeFunctionDeclaration.isCast()) {
126         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
127         if (!nativeFunctionDeclaration.parameters().size()) {
128             stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, "() {\n"));
129             stringBuilder.append(makeString("    ", metalReturnName, " x;\n"));
130             stringBuilder.append(makeString("    ", memsetZeroFunctionName, "(x);\n"));
131             stringBuilder.append("    return x;\n");
132             stringBuilder.append("}\n");
133             return stringBuilder.toString();
134         }
135
136         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
137         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
138         auto& parameterType = nativeFunctionDeclaration.parameters()[0]->type()->unifyNode();
139         if (is<AST::NamedType>(parameterType)) {
140             auto& parameterNamedType = downcast<AST::NamedType>(parameterType);
141             if (is<AST::NativeTypeDeclaration>(parameterNamedType)) {
142                 auto& parameterNativeTypeDeclaration = downcast<AST::NativeTypeDeclaration>(parameterNamedType);
143                 if (parameterNativeTypeDeclaration.isAtomic()) {
144                     stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"));
145                     stringBuilder.append("    return atomic_load_explicit(&x, memory_order_relaxed);\n");
146                     stringBuilder.append("}\n");
147                     return stringBuilder.toString();
148                 }
149             }
150         }
151
152         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"));
153         stringBuilder.append(makeString("    return static_cast<", metalReturnName, ">(x);\n"));
154         stringBuilder.append("}\n");
155         return stringBuilder.toString();
156     }
157
158     if (nativeFunctionDeclaration.name() == "operator.value") {
159         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
160         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
161         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
162         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"));
163         stringBuilder.append(makeString("    return static_cast<", metalReturnName, ">(x);\n"));
164         stringBuilder.append("}\n");
165         return stringBuilder.toString();
166     }
167
168     // 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.
169     if (nativeFunctionDeclaration.name() == "operator.length") {
170         ASSERT_UNUSED(intrinsics, matches(nativeFunctionDeclaration.type(), intrinsics.uintType()));
171         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
172         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
173         auto& parameterType = nativeFunctionDeclaration.parameters()[0]->type()->unifyNode();
174         auto& unnamedParameterType = downcast<AST::UnnamedType>(parameterType);
175         if (is<AST::ArrayType>(unnamedParameterType)) {
176             auto& arrayParameterType = downcast<AST::ArrayType>(unnamedParameterType);
177             stringBuilder.append(makeString("uint ", outputFunctionName, '(', metalParameterName, ") {\n"));
178             stringBuilder.append(makeString("    return ", arrayParameterType.numElements(), ";\n"));
179             stringBuilder.append("}\n");
180             return stringBuilder.toString();
181         }
182
183         ASSERT(is<AST::ArrayReferenceType>(unnamedParameterType));
184         stringBuilder.append(makeString("uint ", outputFunctionName, '(', metalParameterName, " v) {\n"));
185         stringBuilder.append(makeString("    return v.length;\n"));
186         stringBuilder.append("}\n");
187         return stringBuilder.toString();
188     }
189
190     if (nativeFunctionDeclaration.name().startsWith("operator."_str)) {
191         auto mangledFieldName = [&](const String& fieldName) -> String {
192             auto& unifyNode = nativeFunctionDeclaration.parameters()[0]->type()->unifyNode();
193             auto& namedType = downcast<AST::NamedType>(unifyNode);
194             if (is<AST::StructureDefinition>(namedType)) {
195                 auto& structureDefinition = downcast<AST::StructureDefinition>(namedType);
196                 auto* structureElement = structureDefinition.find(fieldName);
197                 ASSERT(structureElement);
198                 return typeNamer.mangledNameForStructureElement(*structureElement);
199             }
200             ASSERT(is<AST::NativeTypeDeclaration>(namedType));
201             return fieldName;
202         };
203
204         if (nativeFunctionDeclaration.name().endsWith("=")) {
205             ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
206             auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
207             auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
208             auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
209             auto fieldName = nativeFunctionDeclaration.name().substring("operator."_str.length());
210             fieldName = fieldName.substring(0, fieldName.length() - 1);
211             auto metalFieldName = mangledFieldName(fieldName);
212             stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " v, ", metalParameter2Name, " n) {\n"));
213             stringBuilder.append(makeString("    v.", metalFieldName, " = n;\n"));
214             stringBuilder.append(makeString("    return v;\n"));
215             stringBuilder.append("}\n");
216             return stringBuilder.toString();
217         }
218
219         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
220         auto fieldName = nativeFunctionDeclaration.name().substring("operator."_str.length());
221         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
222         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
223         auto metalFieldName = mangledFieldName(fieldName);
224         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " v) {\n"));
225         stringBuilder.append(makeString("    return v.", metalFieldName, ";\n"));
226         stringBuilder.append("}\n");
227         return stringBuilder.toString();
228     }
229
230     if (nativeFunctionDeclaration.name().startsWith("operator&."_str)) {
231         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
232         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
233         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
234         auto fieldName = nativeFunctionDeclaration.name().substring("operator&."_str.length());
235
236         String metalFieldName;
237         auto& unnamedType = *nativeFunctionDeclaration.parameters()[0]->type();
238         auto& unifyNode = downcast<AST::PointerType>(unnamedType).elementType().unifyNode();
239         auto& namedType = downcast<AST::NamedType>(unifyNode);
240         if (is<AST::StructureDefinition>(namedType)) {
241             auto& structureDefinition = downcast<AST::StructureDefinition>(namedType);
242             auto* structureElement = structureDefinition.find(fieldName);
243             ASSERT(structureElement);
244             metalFieldName = typeNamer.mangledNameForStructureElement(*structureElement);
245         } else
246             metalFieldName = fieldName;
247
248         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " v) {\n"));
249         stringBuilder.append(makeString("    return &(v->", metalFieldName, ");\n"));
250         stringBuilder.append("}\n");
251         return stringBuilder.toString();
252     }
253
254     if (nativeFunctionDeclaration.name() == "operator&[]") {
255         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
256         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
257         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
258         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
259         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " v, ", metalParameter2Name, " n) {\n"));
260         ASSERT(is<AST::ArrayReferenceType>(*nativeFunctionDeclaration.parameters()[0]->type()));
261         stringBuilder.append("    if (n < v.length) return &(v.pointer[n]);\n");
262         stringBuilder.append("    return nullptr;\n");
263         stringBuilder.append("}\n");
264         return stringBuilder.toString();
265     }
266
267     auto numberOfMatrixRows = [&] {
268         auto& typeReference = downcast<AST::TypeReference>(*nativeFunctionDeclaration.parameters()[0]->type());
269         auto& matrixType = downcast<AST::NativeTypeDeclaration>(downcast<AST::TypeReference>(downcast<AST::TypeDefinition>(typeReference.resolvedType()).type()).resolvedType());
270         ASSERT(matrixType.name() == "matrix");
271         ASSERT(matrixType.typeArguments().size() == 3);
272         return String::number(WTF::get<AST::ConstantExpression>(matrixType.typeArguments()[1]).integerLiteral().value());
273     };
274
275     if (nativeFunctionDeclaration.name() == "operator[]") {
276         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
277         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
278         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
279         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
280         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " m, ", metalParameter2Name, " i) {\n"));
281         stringBuilder.append(makeString("    if (i < ", numberOfMatrixRows(), ") return m[i];\n"));
282         stringBuilder.append(makeString("    return ", metalReturnName, "(0);\n"));
283         stringBuilder.append("}\n");
284         return stringBuilder.toString();
285     }
286
287     if (nativeFunctionDeclaration.name() == "operator[]=") {
288         ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
289         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
290         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
291         auto metalParameter3Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[2]->type());
292         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
293         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " m, ", metalParameter2Name, " i, ", metalParameter3Name, " v) {\n"));
294         stringBuilder.append(makeString("    if (i < ", numberOfMatrixRows(), ") m[i] = v;\n"));
295         stringBuilder.append("    return m;\n");
296         stringBuilder.append("}\n");
297         return stringBuilder.toString();
298     }
299
300     if (nativeFunctionDeclaration.isOperator()) {
301         if (nativeFunctionDeclaration.parameters().size() == 1) {
302             auto operatorName = nativeFunctionDeclaration.name().substring("operator"_str.length());
303             auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
304             auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
305             stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"));
306             stringBuilder.append(makeString("    return ", operatorName, "x;\n"));
307             stringBuilder.append("}\n");
308             return stringBuilder.toString();
309         }
310
311         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
312         auto operatorName = nativeFunctionDeclaration.name().substring("operator"_str.length());
313         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
314         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
315         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
316         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " x, ", metalParameter2Name, " y) {\n"));
317         stringBuilder.append(makeString("    return x ", operatorName, " y;\n"));
318         stringBuilder.append("}\n");
319         return stringBuilder.toString();
320     }
321
322     if (nativeFunctionDeclaration.name() == "cos"
323         || nativeFunctionDeclaration.name() == "sin"
324         || nativeFunctionDeclaration.name() == "tan"
325         || nativeFunctionDeclaration.name() == "acos"
326         || nativeFunctionDeclaration.name() == "asin"
327         || nativeFunctionDeclaration.name() == "atan"
328         || nativeFunctionDeclaration.name() == "cosh"
329         || nativeFunctionDeclaration.name() == "sinh"
330         || nativeFunctionDeclaration.name() == "tanh"
331         || nativeFunctionDeclaration.name() == "ceil"
332         || nativeFunctionDeclaration.name() == "exp"
333         || nativeFunctionDeclaration.name() == "floor"
334         || nativeFunctionDeclaration.name() == "log"
335         || nativeFunctionDeclaration.name() == "round"
336         || nativeFunctionDeclaration.name() == "trunc"
337         || nativeFunctionDeclaration.name() == "ddx"
338         || nativeFunctionDeclaration.name() == "ddy"
339         || nativeFunctionDeclaration.name() == "isnormal"
340         || nativeFunctionDeclaration.name() == "isfinite"
341         || nativeFunctionDeclaration.name() == "isinf"
342         || nativeFunctionDeclaration.name() == "isnan"
343         || nativeFunctionDeclaration.name() == "asint"
344         || nativeFunctionDeclaration.name() == "asuint"
345         || nativeFunctionDeclaration.name() == "asfloat") {
346         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
347         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
348         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
349         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"));
350         stringBuilder.append(makeString("    return ", mapFunctionName(nativeFunctionDeclaration.name()), "(x);\n"));
351         stringBuilder.append("}\n");
352         return stringBuilder.toString();
353     }
354
355     if (nativeFunctionDeclaration.name() == "pow" || nativeFunctionDeclaration.name() == "atan2") {
356         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
357         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
358         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
359         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
360         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " x, ", metalParameter2Name, " y) {\n"));
361         stringBuilder.append(makeString("    return ", nativeFunctionDeclaration.name(), "(x, y);\n"));
362         stringBuilder.append("}\n");
363         return stringBuilder.toString();
364     }
365
366     if (nativeFunctionDeclaration.name() == "f16tof32" || nativeFunctionDeclaration.name() == "f32tof16") {
367         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
368         notImplemented();
369     }
370
371     if (nativeFunctionDeclaration.name() == "AllMemoryBarrierWithGroupSync") {
372         ASSERT(!nativeFunctionDeclaration.parameters().size());
373         stringBuilder.append(makeString("void ", outputFunctionName, "() {\n"));
374         stringBuilder.append("    threadgroup_barrier(mem_flags::mem_device);\n");
375         stringBuilder.append("    threadgroup_barrier(mem_flags::mem_threadgroup);\n");
376         stringBuilder.append("    threadgroup_barrier(mem_flags::mem_texture);\n");
377         stringBuilder.append("}\n");
378         return stringBuilder.toString();
379     }
380
381     if (nativeFunctionDeclaration.name() == "DeviceMemoryBarrierWithGroupSync") {
382         ASSERT(!nativeFunctionDeclaration.parameters().size());
383         stringBuilder.append(makeString("void ", outputFunctionName, "() {\n"));
384         stringBuilder.append("    threadgroup_barrier(mem_flags::mem_device);\n");
385         stringBuilder.append("}\n");
386         return stringBuilder.toString();
387     }
388
389     if (nativeFunctionDeclaration.name() == "GroupMemoryBarrierWithGroupSync") {
390         ASSERT(!nativeFunctionDeclaration.parameters().size());
391         stringBuilder.append(makeString("void ", outputFunctionName, "() {\n"));
392         stringBuilder.append("    threadgroup_barrier(mem_flags::mem_threadgroup);\n");
393         stringBuilder.append("}\n");
394         return stringBuilder.toString();
395     }
396
397     if (nativeFunctionDeclaration.name().startsWith("Interlocked"_str)) {
398         if (nativeFunctionDeclaration.name() == "InterlockedCompareExchange") {
399             ASSERT(nativeFunctionDeclaration.parameters().size() == 4);
400             auto& firstArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0]->type());
401             auto firstArgumentAddressSpace = firstArgumentPointer.addressSpace();
402             auto firstArgumentPointee = typeNamer.mangledNameForType(firstArgumentPointer.elementType());
403             auto secondArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
404             auto thirdArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[2]->type());
405             auto& fourthArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[3]->type());
406             auto fourthArgumentAddressSpace = fourthArgumentPointer.addressSpace();
407             auto fourthArgumentPointee = typeNamer.mangledNameForType(fourthArgumentPointer.elementType());
408             stringBuilder.append(makeString("void ", outputFunctionName, '(', toString(firstArgumentAddressSpace), ' ', firstArgumentPointee, "* object, ", secondArgument, " compare, ", thirdArgument, " desired, ", toString(fourthArgumentAddressSpace), ' ', fourthArgumentPointee, "* out) {\n"));
409             stringBuilder.append("    atomic_compare_exchange_weak_explicit(object, &compare, desired, memory_order_relaxed);\n");
410             stringBuilder.append("    *out = compare;\n");
411             stringBuilder.append("}\n");
412             return stringBuilder.toString();
413         }
414
415         ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
416         auto& firstArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0]->type());
417         auto firstArgumentAddressSpace = firstArgumentPointer.addressSpace();
418         auto firstArgumentPointee = typeNamer.mangledNameForType(firstArgumentPointer.elementType());
419         auto secondArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
420         auto& thirdArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[2]->type());
421         auto thirdArgumentAddressSpace = thirdArgumentPointer.addressSpace();
422         auto thirdArgumentPointee = typeNamer.mangledNameForType(thirdArgumentPointer.elementType());
423         auto name = atomicName(nativeFunctionDeclaration.name().substring("Interlocked"_str.length()));
424         stringBuilder.append(makeString("void ", outputFunctionName, '(', toString(firstArgumentAddressSpace), ' ', firstArgumentPointee, "* object, ", secondArgument, " operand, ", toString(thirdArgumentAddressSpace), ' ', thirdArgumentPointee, "* out) {\n"));
425         stringBuilder.append(makeString("    *out = atomic_fetch_", name, "_explicit(object, operand, memory_order_relaxed);\n"));
426         stringBuilder.append("}\n");
427         return stringBuilder.toString();
428     }
429
430     if (nativeFunctionDeclaration.name() == "Sample") {
431         ASSERT(nativeFunctionDeclaration.parameters().size() == 3 || nativeFunctionDeclaration.parameters().size() == 4);
432         
433         auto& textureType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[0]->type()->unifyNode()));
434         auto& locationType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[2]->type()->unifyNode()));
435         auto locationVectorLength = vectorLength(locationType);
436         auto& returnType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.type().unifyNode()));
437         auto returnVectorLength = vectorLength(returnType);
438
439         auto metalParameter1Name = typeNamer.mangledNameForType(textureType);
440         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
441         auto metalParameter3Name = typeNamer.mangledNameForType(locationType);
442         String metalParameter4Name;
443         if (nativeFunctionDeclaration.parameters().size() == 4)
444             metalParameter4Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[3]->type());
445         auto metalReturnName = typeNamer.mangledNameForType(returnType);
446         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " theTexture, ", metalParameter2Name, " theSampler, ", metalParameter3Name, " location"));
447         if (!metalParameter4Name.isNull())
448             stringBuilder.append(makeString(", ", metalParameter4Name, " offset"));
449         stringBuilder.append(") {\n");
450         stringBuilder.append("    return theTexture.sample(theSampler, ");
451         if (textureType.isTextureArray()) {
452             ASSERT(locationVectorLength > 1);
453             stringBuilder.append(makeString("location.", "xyzw"_str.substring(0, locationVectorLength - 1), ", location.", "xyzw"_str.substring(locationVectorLength - 1, 1)));
454         } else
455             stringBuilder.append("location");
456         if (!metalParameter4Name.isNull())
457             stringBuilder.append(", offset");
458         stringBuilder.append(")");
459         if (!textureType.isDepthTexture())
460             stringBuilder.append(makeString(".", "xyzw"_str.substring(0, returnVectorLength)));
461         stringBuilder.append(";\n");
462         stringBuilder.append("}\n");
463         return stringBuilder.toString();
464     }
465
466     if (nativeFunctionDeclaration.name() == "Load") {
467         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
468         
469         auto& textureType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[0]->type()->unifyNode()));
470         auto& locationType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[1]->type()->unifyNode()));
471         auto locationVectorLength = vectorLength(locationType);
472         auto& returnType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.type().unifyNode()));
473         auto returnVectorLength = vectorLength(returnType);
474
475         auto metalParameter1Name = typeNamer.mangledNameForType(textureType);
476         auto metalParameter2Name = typeNamer.mangledNameForType(locationType);
477         auto metalReturnName = typeNamer.mangledNameForType(returnType);
478         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " theTexture, ", metalParameter2Name, " location) {\n"));
479         if (textureType.isTextureArray()) {
480             ASSERT(locationVectorLength > 1);
481             String dimensions[] = { "width"_str, "height"_str, "depth"_str };
482             for (int i = 0; i < locationVectorLength - 1; ++i) {
483                 auto suffix = "xyzw"_str.substring(i, 1);
484                 stringBuilder.append(makeString("    if (location.", suffix, " < 0 || static_cast<uint32_t>(location.", suffix, ") >= theTexture.get_", dimensions[i], "()) return ", metalReturnName, "(0);\n"));
485             }
486             auto suffix = "xyzw"_str.substring(locationVectorLength - 1, 1);
487             stringBuilder.append(makeString("    if (location.", suffix, " < 0 || static_cast<uint32_t>(location.", suffix, ") >= theTexture.get_array_size()) return ", metalReturnName, "(0);\n"));
488         } else {
489             if (locationVectorLength == 1)
490                 stringBuilder.append(makeString("    if (location < 0 || static_cast<uint32_t>(location) >= theTexture.get_width()) return ", metalReturnName, "(0);\n"));
491             else {
492                 stringBuilder.append(makeString("    if (location.x < 0 || static_cast<uint32_t>(location.x) >= theTexture.get_width()) return ", metalReturnName, "(0);\n"));
493                 stringBuilder.append(makeString("    if (location.y < 0 || static_cast<uint32_t>(location.y) >= theTexture.get_height()) return ", metalReturnName, "(0);\n"));
494                 if (locationVectorLength >= 3)
495                     stringBuilder.append(makeString("    if (location.z < 0 || static_cast<uint32_t>(location.z) >= theTexture.get_depth()) return ", metalReturnName, "(0);\n"));
496             }
497         }
498         stringBuilder.append("    return theTexture.read(");
499         if (textureType.isTextureArray()) {
500             ASSERT(locationVectorLength > 1);
501             stringBuilder.append(makeString("uint", vectorSuffix(locationVectorLength - 1), "(location.", "xyzw"_str.substring(0, locationVectorLength - 1), "), uint(location.", "xyzw"_str.substring(locationVectorLength - 1, 1), ")"));
502         } else
503             stringBuilder.append(makeString("uint", vectorSuffix(locationVectorLength), "(location)"));
504         stringBuilder.append(")");
505         if (!textureType.isDepthTexture())
506             stringBuilder.append(makeString(".", "xyzw"_str.substring(0, returnVectorLength)));
507         stringBuilder.append(";\n");
508         stringBuilder.append("}\n");
509         return stringBuilder.toString();
510     }
511
512     if (nativeFunctionDeclaration.name() == "GetDimensions") {
513         auto& textureType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[0]->type()->unifyNode()));
514
515         size_t index = 1;
516         if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
517             ++index;
518         auto widthTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
519         ++index;
520         String heightTypeName;
521         if (textureType.textureDimension() >= 2) {
522             heightTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
523             ++index;
524         }
525         String depthTypeName;
526         if (textureType.textureDimension() >= 3) {
527             depthTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
528             ++index;
529         }
530         String elementsTypeName;
531         if (textureType.isTextureArray()) {
532             elementsTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
533             ++index;
534         }
535         String numberOfLevelsTypeName;
536         if (!textureType.isWritableTexture()) {
537             numberOfLevelsTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
538             ++index;
539         }
540         ASSERT(index == nativeFunctionDeclaration.parameters().size());
541
542         auto metalParameter1Name = typeNamer.mangledNameForType(textureType);
543         stringBuilder.append(makeString("void ", outputFunctionName, '(', metalParameter1Name, " theTexture"));
544         if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
545             stringBuilder.append(", uint mipLevel");
546         stringBuilder.append(makeString(", ", widthTypeName, " width"));
547         if (!heightTypeName.isNull())
548             stringBuilder.append(makeString(", ", heightTypeName, " height"));
549         if (!depthTypeName.isNull())
550             stringBuilder.append(makeString(", ", depthTypeName, " depth"));
551         if (!elementsTypeName.isNull())
552             stringBuilder.append(makeString(", ", elementsTypeName, " elements"));
553         if (!numberOfLevelsTypeName.isNull())
554             stringBuilder.append(makeString(", ", numberOfLevelsTypeName, " numberOfLevels"));
555         stringBuilder.append(") {\n");
556         stringBuilder.append("    if (width)\n");
557         stringBuilder.append("        *width = theTexture.get_width(");
558         if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
559             stringBuilder.append("mipLevel");
560         stringBuilder.append(");\n");
561         if (!heightTypeName.isNull()) {
562             stringBuilder.append("    if (height)\n");
563             stringBuilder.append("        *height = theTexture.get_height(");
564             if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
565                 stringBuilder.append("mipLevel");
566             stringBuilder.append(");\n");
567         }
568         if (!depthTypeName.isNull()) {
569             stringBuilder.append("    if (depth)\n");
570             stringBuilder.append("        *depth = theTexture.get_depth(");
571             if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
572                 stringBuilder.append("mipLevel");
573             stringBuilder.append(");\n");
574         }
575         if (!elementsTypeName.isNull()) {
576             stringBuilder.append("    if (elements)\n");
577             stringBuilder.append("        *elements = theTexture.get_array_size();\n");
578         }
579         if (!numberOfLevelsTypeName.isNull()) {
580             stringBuilder.append("    if (numberOfLevels)\n");
581             stringBuilder.append("        *numberOfLevels = theTexture.get_num_mip_levels();\n");
582         }
583         stringBuilder.append("}\n");
584         return stringBuilder.toString();
585     }
586
587     if (nativeFunctionDeclaration.name() == "SampleBias") {
588         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
589         notImplemented();
590     }
591
592     if (nativeFunctionDeclaration.name() == "SampleGrad") {
593         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
594         notImplemented();
595     }
596
597     if (nativeFunctionDeclaration.name() == "SampleLevel") {
598         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
599         notImplemented();
600     }
601
602     if (nativeFunctionDeclaration.name() == "Gather") {
603         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
604         notImplemented();
605     }
606
607     if (nativeFunctionDeclaration.name() == "GatherRed") {
608         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
609         notImplemented();
610     }
611
612     if (nativeFunctionDeclaration.name() == "SampleCmp") {
613         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
614         notImplemented();
615     }
616
617     if (nativeFunctionDeclaration.name() == "SampleCmpLevelZero") {
618         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
619         notImplemented();
620     }
621
622     if (nativeFunctionDeclaration.name() == "Store") {
623         ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
624         
625         auto& textureType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[0]->type()->unifyNode()));
626         auto& itemType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[1]->type()->unifyNode()));
627         auto& itemVectorInnerType = vectorInnerType(itemType);
628         auto itemVectorLength = vectorLength(itemType);
629         auto& locationType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[2]->type()->unifyNode()));
630         auto locationVectorLength = vectorLength(locationType);
631
632         auto metalParameter1Name = typeNamer.mangledNameForType(textureType);
633         auto metalParameter2Name = typeNamer.mangledNameForType(itemType);
634         auto metalParameter3Name = typeNamer.mangledNameForType(locationType);
635         auto metalInnerTypeName = typeNamer.mangledNameForType(itemVectorInnerType);
636         stringBuilder.append(makeString("void ", outputFunctionName, '(', metalParameter1Name, " theTexture, ", metalParameter2Name, " item, ", metalParameter3Name, " location) {\n"));
637         if (textureType.isTextureArray()) {
638             ASSERT(locationVectorLength > 1);
639             String dimensions[] = { "width"_str, "height"_str, "depth"_str };
640             for (int i = 0; i < locationVectorLength - 1; ++i) {
641                 auto suffix = "xyzw"_str.substring(i, 1);
642                 stringBuilder.append(makeString("    if (location.", suffix, " >= theTexture.get_", dimensions[i], "()) return;\n"));
643             }
644             auto suffix = "xyzw"_str.substring(locationVectorLength - 1, 1);
645             stringBuilder.append(makeString("    if (location.", suffix, " >= theTexture.get_array_size()) return;\n"));
646         } else {
647             if (locationVectorLength == 1)
648                 stringBuilder.append(makeString("    if (location >= theTexture.get_width()) return;\n"));
649             else {
650                 stringBuilder.append(makeString("    if (location.x >= theTexture.get_width()) return;\n"));
651                 stringBuilder.append(makeString("    if (location.y >= theTexture.get_height()) return;\n"));
652                 if (locationVectorLength >= 3)
653                     stringBuilder.append(makeString("    if (location.z >= theTexture.get_depth()) return;\n"));
654             }
655         }
656         stringBuilder.append(makeString("    theTexture.write(vec<", metalInnerTypeName, ", 4>(item"));
657         for (int i = 0; i < 4 - itemVectorLength; ++i)
658             stringBuilder.append(", 0");
659         stringBuilder.append("), ");
660         if (textureType.isTextureArray()) {
661             ASSERT(locationVectorLength > 1);
662             stringBuilder.append(makeString("uint", vectorSuffix(locationVectorLength - 1), "(location.", "xyzw"_str.substring(0, locationVectorLength - 1), "), uint(location.", "xyzw"_str.substring(locationVectorLength - 1, 1), ")"));
663         } else
664             stringBuilder.append(makeString("uint", vectorSuffix(locationVectorLength), "(location)"));
665         stringBuilder.append(");\n");
666         stringBuilder.append("}\n");
667         return stringBuilder.toString();
668     }
669
670     if (nativeFunctionDeclaration.name() == "GatherAlpha") {
671         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
672         notImplemented();
673     }
674
675     if (nativeFunctionDeclaration.name() == "GatherBlue") {
676         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
677         notImplemented();
678     }
679
680     if (nativeFunctionDeclaration.name() == "GatherCmp") {
681         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
682         notImplemented();
683     }
684
685     if (nativeFunctionDeclaration.name() == "GatherCmpRed") {
686         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
687         notImplemented();
688     }
689
690     if (nativeFunctionDeclaration.name() == "GatherGreen") {
691         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
692         notImplemented();
693     }
694
695     ASSERT_NOT_REACHED();
696     return String();
697 }
698
699 } // namespace Metal
700
701 } // namespace WHLSL
702
703 } // namespace WebCore
704
705 #endif