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