1f3d3cd23a55eff3a0fe32a3428d0c8e38a7470a
[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 "WHLSLTypeNamer.h"
41 #include "WHLSLUnnamedType.h"
42 #include "WHLSLVariableDeclaration.h"
43 #include <wtf/text/StringBuilder.h>
44
45 namespace WebCore {
46
47 namespace WHLSL {
48
49 namespace Metal {
50
51 static String mapFunctionName(String& functionName)
52 {
53     if (functionName == "ddx")
54         return "dfdx"_str;
55     if (functionName == "ddy")
56         return "dfdy"_str;
57     if (functionName == "asint")
58         return "as_type<int32_t>"_str;
59     if (functionName == "asuint")
60         return "as_type<uint32_t>"_str;
61     if (functionName == "asfloat")
62         return "as_type<float>"_str;
63     return functionName;
64 }
65
66 static String atomicName(String input)
67 {
68     if (input == "Add")
69         return "fetch_add"_str;
70     if (input == "And")
71         return "fetch_and"_str;
72     if (input == "Exchange")
73         return "exchange"_str;
74     if (input == "Max")
75         return "fetch_max"_str;
76     if (input == "Min")
77         return "fetch_min"_str;
78     if (input == "Or")
79         return "fetch_or"_str;
80     ASSERT(input == "Xor");
81         return "fetch_xor"_str;
82 }
83
84 String writeNativeFunction(AST::NativeFunctionDeclaration& nativeFunctionDeclaration, String& outputFunctionName, Intrinsics& intrinsics, TypeNamer& typeNamer, const char* memsetZeroFunctionName)
85 {
86     StringBuilder stringBuilder;
87     if (nativeFunctionDeclaration.isCast()) {
88         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
89         if (!nativeFunctionDeclaration.parameters().size()) {
90             stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, "() {\n"));
91             stringBuilder.append(makeString("    ", metalReturnName, " x;\n"));
92             stringBuilder.append(makeString("    ", memsetZeroFunctionName, "(x);\n"));
93             stringBuilder.append("    return x;\n");
94             stringBuilder.append("}\n");
95             return stringBuilder.toString();
96         }
97
98         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
99         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
100         auto& parameterType = nativeFunctionDeclaration.parameters()[0]->type()->unifyNode();
101         if (is<AST::NamedType>(parameterType)) {
102             auto& parameterNamedType = downcast<AST::NamedType>(parameterType);
103             if (is<AST::NativeTypeDeclaration>(parameterNamedType)) {
104                 auto& parameterNativeTypeDeclaration = downcast<AST::NativeTypeDeclaration>(parameterNamedType);
105                 if (parameterNativeTypeDeclaration.isAtom()) {
106                     stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"));
107                     stringBuilder.append("    return atomic_load_explicit(&x, memory_order_relaxed);\n");
108                     stringBuilder.append("}\n");
109                     return stringBuilder.toString();
110                 }
111             }
112         }
113
114         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"));
115         stringBuilder.append(makeString("    return static_cast<", metalReturnName, ">(x);\n"));
116         stringBuilder.append("}\n");
117         return stringBuilder.toString();
118     }
119
120     if (nativeFunctionDeclaration.name() == "operator.value") {
121         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
122         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
123         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
124         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"));
125         stringBuilder.append(makeString("    return static_cast<", metalReturnName, ">(x);\n"));
126         stringBuilder.append("}\n");
127         return stringBuilder.toString();
128     }
129
130     // 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.
131     if (nativeFunctionDeclaration.name() == "operator.length") {
132         ASSERT_UNUSED(intrinsics, matches(nativeFunctionDeclaration.type(), intrinsics.uintType()));
133         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
134         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
135         auto& parameterType = nativeFunctionDeclaration.parameters()[0]->type()->unifyNode();
136         auto& unnamedParameterType = downcast<AST::UnnamedType>(parameterType);
137         if (is<AST::ArrayType>(unnamedParameterType)) {
138             auto& arrayParameterType = downcast<AST::ArrayType>(unnamedParameterType);
139             stringBuilder.append(makeString("uint ", outputFunctionName, '(', metalParameterName, " v) {\n"));
140             stringBuilder.append(makeString("    return ", arrayParameterType.numElements(), "u;\n"));
141             stringBuilder.append("}\n");
142             return stringBuilder.toString();
143         }
144
145         ASSERT(is<AST::ArrayReferenceType>(unnamedParameterType));
146         stringBuilder.append(makeString("uint ", outputFunctionName, '(', metalParameterName, " v) {\n"));
147         stringBuilder.append(makeString("    return v.length;\n"));
148         stringBuilder.append("}\n");
149         return stringBuilder.toString();
150     }
151
152     if (nativeFunctionDeclaration.name().startsWith("operator."_str)) {
153         auto mangledFieldName = [&](const String& fieldName) -> String {
154             auto& unifyNode = nativeFunctionDeclaration.parameters()[0]->type()->unifyNode();
155             auto& namedType = downcast<AST::NamedType>(unifyNode);
156             if (is<AST::StructureDefinition>(namedType)) {
157                 auto& structureDefinition = downcast<AST::StructureDefinition>(namedType);
158                 auto* structureElement = structureDefinition.find(fieldName);
159                 ASSERT(structureElement);
160                 return typeNamer.mangledNameForStructureElement(*structureElement);
161             }
162             ASSERT(is<AST::NativeTypeDeclaration>(namedType));
163             return fieldName;
164         };
165
166         if (nativeFunctionDeclaration.name().endsWith("=")) {
167             ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
168             auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
169             auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
170             auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
171             auto fieldName = nativeFunctionDeclaration.name().substring("operator."_str.length());
172             fieldName = fieldName.substring(0, fieldName.length() - 1);
173             auto metalFieldName = mangledFieldName(fieldName);
174             stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " v, ", metalParameter2Name, " n) {\n"));
175             stringBuilder.append(makeString("    v.", metalFieldName, " = n;\n"));
176             stringBuilder.append(makeString("    return v;\n"));
177             stringBuilder.append("}\n");
178             return stringBuilder.toString();
179         }
180
181         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
182         auto fieldName = nativeFunctionDeclaration.name().substring("operator."_str.length());
183         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
184         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
185         auto metalFieldName = mangledFieldName(fieldName);
186         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " v) {\n"));
187         stringBuilder.append(makeString("    return v.", metalFieldName, ";\n"));
188         stringBuilder.append("}\n");
189         return stringBuilder.toString();
190     }
191
192     if (nativeFunctionDeclaration.name().startsWith("operator&."_str)) {
193         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
194         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
195         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
196         auto fieldName = nativeFunctionDeclaration.name().substring("operator&."_str.length());
197
198         String metalFieldName;
199         auto& unnamedType = *nativeFunctionDeclaration.parameters()[0]->type();
200         auto& unifyNode = downcast<AST::PointerType>(unnamedType).elementType().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             metalFieldName = typeNamer.mangledNameForStructureElement(*structureElement);
207         } else
208             metalFieldName = fieldName;
209
210         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " v) {\n"));
211         stringBuilder.append(makeString("    return &(v->", metalFieldName, ");\n"));
212         stringBuilder.append("}\n");
213         return stringBuilder.toString();
214     }
215
216     if (nativeFunctionDeclaration.name() == "operator&[]") {
217         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
218         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
219         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
220         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
221         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " v, ", metalParameter2Name, " n) {\n"));
222         stringBuilder.append("    if (n < v.length) return &(v.pointer[n]);\n");
223         stringBuilder.append("    return nullptr;\n");
224         stringBuilder.append("}\n");
225         return stringBuilder.toString();
226     }
227
228     auto numberOfMatrixRows = [&] {
229         auto& typeReference = downcast<AST::TypeReference>(*nativeFunctionDeclaration.parameters()[0]->type());
230         auto& matrixType = downcast<AST::NativeTypeDeclaration>(downcast<AST::TypeReference>(downcast<AST::TypeDefinition>(typeReference.resolvedType()).type()).resolvedType());
231         ASSERT(matrixType.name() == "matrix");
232         ASSERT(matrixType.typeArguments().size() == 3);
233         return String::number(WTF::get<AST::ConstantExpression>(matrixType.typeArguments()[1]).integerLiteral().value());
234     };
235
236     if (nativeFunctionDeclaration.name() == "operator[]") {
237         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
238         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
239         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
240         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
241         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " m, ", metalParameter2Name, " i) {\n"));
242         stringBuilder.append(makeString("    if (i < ", numberOfMatrixRows(), ") return m[i];\n"));
243         stringBuilder.append(makeString("    return ", metalReturnName, "(0);\n"));
244         stringBuilder.append("}\n");
245         return stringBuilder.toString();
246     }
247
248     if (nativeFunctionDeclaration.name() == "operator[]=") {
249         ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
250         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
251         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
252         auto metalParameter3Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[2]->type());
253         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
254         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " m, ", metalParameter2Name, " i, ", metalParameter3Name, " v) {\n"));
255         stringBuilder.append(makeString("    if (i < ", numberOfMatrixRows(), ") m[i] = v;\n"));
256         stringBuilder.append("    return m;\n");
257         stringBuilder.append("}\n");
258         return stringBuilder.toString();
259     }
260
261     if (nativeFunctionDeclaration.isOperator()) {
262         if (nativeFunctionDeclaration.parameters().size() == 1) {
263             auto operatorName = nativeFunctionDeclaration.name().substring("operator"_str.length());
264             auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
265             auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
266             stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"));
267             stringBuilder.append(makeString("    return ", operatorName, "x;\n"));
268             stringBuilder.append("}\n");
269             return stringBuilder.toString();
270         }
271
272         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
273         auto operatorName = nativeFunctionDeclaration.name().substring("operator"_str.length());
274         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
275         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
276         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
277         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " x, ", metalParameter2Name, " y) {\n"));
278         stringBuilder.append(makeString("    return x ", operatorName, " y;\n"));
279         stringBuilder.append("}\n");
280         return stringBuilder.toString();
281     }
282
283     if (nativeFunctionDeclaration.name() == "cos"
284         || nativeFunctionDeclaration.name() == "sin"
285         || nativeFunctionDeclaration.name() == "tan"
286         || nativeFunctionDeclaration.name() == "acos"
287         || nativeFunctionDeclaration.name() == "asin"
288         || nativeFunctionDeclaration.name() == "atan"
289         || nativeFunctionDeclaration.name() == "cosh"
290         || nativeFunctionDeclaration.name() == "sinh"
291         || nativeFunctionDeclaration.name() == "tanh"
292         || nativeFunctionDeclaration.name() == "ceil"
293         || nativeFunctionDeclaration.name() == "exp"
294         || nativeFunctionDeclaration.name() == "floor"
295         || nativeFunctionDeclaration.name() == "log"
296         || nativeFunctionDeclaration.name() == "round"
297         || nativeFunctionDeclaration.name() == "trunc"
298         || nativeFunctionDeclaration.name() == "ddx"
299         || nativeFunctionDeclaration.name() == "ddy"
300         || nativeFunctionDeclaration.name() == "isnormal"
301         || nativeFunctionDeclaration.name() == "isfinite"
302         || nativeFunctionDeclaration.name() == "isinf"
303         || nativeFunctionDeclaration.name() == "isnan"
304         || nativeFunctionDeclaration.name() == "asint"
305         || nativeFunctionDeclaration.name() == "asuint"
306         || nativeFunctionDeclaration.name() == "asfloat") {
307         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
308         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
309         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
310         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"));
311         stringBuilder.append(makeString("    return ", mapFunctionName(nativeFunctionDeclaration.name()), "(x);\n"));
312         stringBuilder.append("}\n");
313         return stringBuilder.toString();
314     }
315
316     if (nativeFunctionDeclaration.name() == "pow" || nativeFunctionDeclaration.name() == "atan2") {
317         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
318         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
319         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
320         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
321         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " x, ", metalParameter2Name, " y) {\n"));
322         stringBuilder.append(makeString("    return ", nativeFunctionDeclaration.name(), "(x, y);\n"));
323         stringBuilder.append("}\n");
324         return stringBuilder.toString();
325     }
326
327     if (nativeFunctionDeclaration.name() == "f16tof32" || nativeFunctionDeclaration.name() == "f32tof16") {
328         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
329         notImplemented();
330     }
331
332     if (nativeFunctionDeclaration.name() == "AllMemoryBarrierWithGroupSync") {
333         ASSERT(!nativeFunctionDeclaration.parameters().size());
334         stringBuilder.append(makeString("void ", outputFunctionName, "() {\n"));
335         stringBuilder.append("    threadgroup_barrier(mem_flags::mem_device);\n");
336         stringBuilder.append("    threadgroup_barrier(mem_flags::mem_threadgroup);\n");
337         stringBuilder.append("    threadgroup_barrier(mem_flags::mem_texture);\n");
338         stringBuilder.append("}\n");
339         return stringBuilder.toString();
340     }
341
342     if (nativeFunctionDeclaration.name() == "DeviceMemoryBarrierWithGroupSync") {
343         ASSERT(!nativeFunctionDeclaration.parameters().size());
344         stringBuilder.append(makeString("void ", outputFunctionName, "() {\n"));
345         stringBuilder.append("    threadgroup_barrier(mem_flags::mem_device);\n");
346         stringBuilder.append("}\n");
347         return stringBuilder.toString();
348     }
349
350     if (nativeFunctionDeclaration.name() == "GroupMemoryBarrierWithGroupSync") {
351         ASSERT(!nativeFunctionDeclaration.parameters().size());
352         stringBuilder.append(makeString("void ", outputFunctionName, "() {\n"));
353         stringBuilder.append("    threadgroup_barrier(mem_flags::mem_threadgroup);\n");
354         stringBuilder.append("}\n");
355         return stringBuilder.toString();
356     }
357
358     if (nativeFunctionDeclaration.name().startsWith("Interlocked"_str)) {
359         if (nativeFunctionDeclaration.name() == "InterlockedCompareExchange") {
360             ASSERT(nativeFunctionDeclaration.parameters().size() == 4);
361             auto& firstArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0]->type());
362             auto firstArgumentAddressSpace = firstArgumentPointer.addressSpace();
363             auto firstArgumentPointee = typeNamer.mangledNameForType(firstArgumentPointer.elementType());
364             auto secondArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
365             auto thirdArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[2]->type());
366             auto& fourthArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[3]->type());
367             auto fourthArgumentAddressSpace = fourthArgumentPointer.addressSpace();
368             auto fourthArgumentPointee = typeNamer.mangledNameForType(fourthArgumentPointer.elementType());
369             stringBuilder.append(makeString("void ", outputFunctionName, '(', toString(firstArgumentAddressSpace), ' ', firstArgumentPointee, "* object, ", secondArgument, " compare, ", thirdArgument, " desired, ", toString(fourthArgumentAddressSpace), ' ', fourthArgumentPointee, "* out) {\n"));
370             stringBuilder.append("    atomic_compare_exchange_weak_explicit(object, &compare, desired, memory_order_relaxed);\n");
371             stringBuilder.append("    *out = compare;\n");
372             stringBuilder.append("}\n");
373             return stringBuilder.toString();
374         }
375
376         ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
377         auto& firstArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0]->type());
378         auto firstArgumentAddressSpace = firstArgumentPointer.addressSpace();
379         auto firstArgumentPointee = typeNamer.mangledNameForType(firstArgumentPointer.elementType());
380         auto secondArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
381         auto& thirdArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[2]->type());
382         auto thirdArgumentAddressSpace = thirdArgumentPointer.addressSpace();
383         auto thirdArgumentPointee = typeNamer.mangledNameForType(thirdArgumentPointer.elementType());
384         auto name = atomicName(nativeFunctionDeclaration.name().substring("Interlocked"_str.length()));
385         stringBuilder.append(makeString("void ", outputFunctionName, '(', toString(firstArgumentAddressSpace), ' ', firstArgumentPointee, "* object, ", secondArgument, " operand, ", toString(thirdArgumentAddressSpace), ' ', thirdArgumentPointee, "* out) {\n"));
386         stringBuilder.append(makeString("    *out = atomic_fetch_", name, "_explicit(object, operand, memory_order_relaxed);\n"));
387         stringBuilder.append("}\n");
388         return stringBuilder.toString();
389     }
390
391     if (nativeFunctionDeclaration.name() == "Sample") {
392         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
393         notImplemented();
394     }
395
396     if (nativeFunctionDeclaration.name() == "Load") {
397         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
398         notImplemented();
399     }
400
401     if (nativeFunctionDeclaration.name() == "GetDimensions") {
402         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
403         notImplemented();
404     }
405
406     if (nativeFunctionDeclaration.name() == "SampleBias") {
407         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
408         notImplemented();
409     }
410
411     if (nativeFunctionDeclaration.name() == "SampleGrad") {
412         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
413         notImplemented();
414     }
415
416     if (nativeFunctionDeclaration.name() == "SampleLevel") {
417         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
418         notImplemented();
419     }
420
421     if (nativeFunctionDeclaration.name() == "Gather") {
422         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
423         notImplemented();
424     }
425
426     if (nativeFunctionDeclaration.name() == "GatherRed") {
427         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
428         notImplemented();
429     }
430
431     if (nativeFunctionDeclaration.name() == "SampleCmp") {
432         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
433         notImplemented();
434     }
435
436     if (nativeFunctionDeclaration.name() == "SampleCmpLevelZero") {
437         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
438         notImplemented();
439     }
440
441     if (nativeFunctionDeclaration.name() == "Store") {
442         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
443         notImplemented();
444     }
445
446     if (nativeFunctionDeclaration.name() == "GatherAlpha") {
447         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
448         notImplemented();
449     }
450
451     if (nativeFunctionDeclaration.name() == "GatherBlue") {
452         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
453         notImplemented();
454     }
455
456     if (nativeFunctionDeclaration.name() == "GatherCmp") {
457         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
458         notImplemented();
459     }
460
461     if (nativeFunctionDeclaration.name() == "GatherCmpRed") {
462         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
463         notImplemented();
464     }
465
466     if (nativeFunctionDeclaration.name() == "GatherGreen") {
467         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
468         notImplemented();
469     }
470
471     ASSERT_NOT_REACHED();
472     return String();
473 }
474
475 } // namespace Metal
476
477 } // namespace WHLSL
478
479 } // namespace WebCore
480
481 #endif