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