Rename AtomicString to AtomString
[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     if (nativeFunctionDeclaration.isOperator()) {
229         if (nativeFunctionDeclaration.parameters().size() == 1) {
230             auto operatorName = nativeFunctionDeclaration.name().substring("operator"_str.length());
231             auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
232             auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
233             stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"));
234             stringBuilder.append(makeString("    return ", operatorName, "x;\n"));
235             stringBuilder.append("}\n");
236             return stringBuilder.toString();
237         }
238
239         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
240         auto operatorName = nativeFunctionDeclaration.name().substring("operator"_str.length());
241         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
242         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
243         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
244         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " x, ", metalParameter2Name, " y) {\n"));
245         stringBuilder.append(makeString("    return x ", operatorName, " y;\n"));
246         stringBuilder.append("}\n");
247         return stringBuilder.toString();
248     }
249
250     if (nativeFunctionDeclaration.name() == "cos"
251         || nativeFunctionDeclaration.name() == "sin"
252         || nativeFunctionDeclaration.name() == "tan"
253         || nativeFunctionDeclaration.name() == "acos"
254         || nativeFunctionDeclaration.name() == "asin"
255         || nativeFunctionDeclaration.name() == "atan"
256         || nativeFunctionDeclaration.name() == "cosh"
257         || nativeFunctionDeclaration.name() == "sinh"
258         || nativeFunctionDeclaration.name() == "tanh"
259         || nativeFunctionDeclaration.name() == "ceil"
260         || nativeFunctionDeclaration.name() == "exp"
261         || nativeFunctionDeclaration.name() == "floor"
262         || nativeFunctionDeclaration.name() == "log"
263         || nativeFunctionDeclaration.name() == "round"
264         || nativeFunctionDeclaration.name() == "trunc"
265         || nativeFunctionDeclaration.name() == "ddx"
266         || nativeFunctionDeclaration.name() == "ddy"
267         || nativeFunctionDeclaration.name() == "isnormal"
268         || nativeFunctionDeclaration.name() == "isfinite"
269         || nativeFunctionDeclaration.name() == "isinf"
270         || nativeFunctionDeclaration.name() == "isnan"
271         || nativeFunctionDeclaration.name() == "asint"
272         || nativeFunctionDeclaration.name() == "asuint"
273         || nativeFunctionDeclaration.name() == "asfloat") {
274         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
275         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
276         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
277         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"));
278         stringBuilder.append(makeString("    return ", mapFunctionName(nativeFunctionDeclaration.name()), "(x);\n"));
279         stringBuilder.append("}\n");
280         return stringBuilder.toString();
281     }
282
283     if (nativeFunctionDeclaration.name() == "pow" || nativeFunctionDeclaration.name() == "atan2") {
284         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
285         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
286         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
287         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
288         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " x, ", metalParameter2Name, " y) {\n"));
289         stringBuilder.append(makeString("    return ", nativeFunctionDeclaration.name(), "(x, y);\n"));
290         stringBuilder.append("}\n");
291         return stringBuilder.toString();
292     }
293
294     if (nativeFunctionDeclaration.name() == "f16tof32" || nativeFunctionDeclaration.name() == "f32tof16") {
295         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
296         notImplemented();
297     }
298
299     if (nativeFunctionDeclaration.name() == "AllMemoryBarrierWithGroupSync") {
300         ASSERT(!nativeFunctionDeclaration.parameters().size());
301         stringBuilder.append(makeString("void ", outputFunctionName, "() {\n"));
302         stringBuilder.append("    threadgroup_barrier(mem_flags::mem_device);\n");
303         stringBuilder.append("    threadgroup_barrier(mem_flags::mem_threadgroup);\n");
304         stringBuilder.append("    threadgroup_barrier(mem_flags::mem_texture);\n");
305         stringBuilder.append("}\n");
306         return stringBuilder.toString();
307     }
308
309     if (nativeFunctionDeclaration.name() == "DeviceMemoryBarrierWithGroupSync") {
310         ASSERT(!nativeFunctionDeclaration.parameters().size());
311         stringBuilder.append(makeString("void ", outputFunctionName, "() {\n"));
312         stringBuilder.append("    threadgroup_barrier(mem_flags::mem_device);\n");
313         stringBuilder.append("}\n");
314         return stringBuilder.toString();
315     }
316
317     if (nativeFunctionDeclaration.name() == "GroupMemoryBarrierWithGroupSync") {
318         ASSERT(!nativeFunctionDeclaration.parameters().size());
319         stringBuilder.append(makeString("void ", outputFunctionName, "() {\n"));
320         stringBuilder.append("    threadgroup_barrier(mem_flags::mem_threadgroup);\n");
321         stringBuilder.append("}\n");
322         return stringBuilder.toString();
323     }
324
325     if (nativeFunctionDeclaration.name().startsWith("Interlocked"_str)) {
326         if (nativeFunctionDeclaration.name() == "InterlockedCompareExchange") {
327             ASSERT(nativeFunctionDeclaration.parameters().size() == 4);
328             auto& firstArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0]->type());
329             auto firstArgumentAddressSpace = firstArgumentPointer.addressSpace();
330             auto firstArgumentPointee = typeNamer.mangledNameForType(firstArgumentPointer.elementType());
331             auto secondArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
332             auto thirdArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[2]->type());
333             auto& fourthArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[3]->type());
334             auto fourthArgumentAddressSpace = fourthArgumentPointer.addressSpace();
335             auto fourthArgumentPointee = typeNamer.mangledNameForType(fourthArgumentPointer.elementType());
336             stringBuilder.append(makeString("void ", outputFunctionName, '(', toString(firstArgumentAddressSpace), ' ', firstArgumentPointee, "* object, ", secondArgument, " compare, ", thirdArgument, " desired, ", toString(fourthArgumentAddressSpace), ' ', fourthArgumentPointee, "* out) {\n"));
337             stringBuilder.append("    atomic_compare_exchange_weak_explicit(object, &compare, desired, memory_order_relaxed);\n");
338             stringBuilder.append("    *out = compare;\n");
339             stringBuilder.append("}\n");
340             return stringBuilder.toString();
341         }
342
343         ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
344         auto& firstArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0]->type());
345         auto firstArgumentAddressSpace = firstArgumentPointer.addressSpace();
346         auto firstArgumentPointee = typeNamer.mangledNameForType(firstArgumentPointer.elementType());
347         auto secondArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
348         auto& thirdArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[2]->type());
349         auto thirdArgumentAddressSpace = thirdArgumentPointer.addressSpace();
350         auto thirdArgumentPointee = typeNamer.mangledNameForType(thirdArgumentPointer.elementType());
351         auto name = atomicName(nativeFunctionDeclaration.name().substring("Interlocked"_str.length()));
352         stringBuilder.append(makeString("void ", outputFunctionName, '(', toString(firstArgumentAddressSpace), ' ', firstArgumentPointee, "* object, ", secondArgument, " operand, ", toString(thirdArgumentAddressSpace), ' ', thirdArgumentPointee, "* out) {\n"));
353         stringBuilder.append(makeString("    *out = atomic_fetch_", name, "_explicit(object, operand, memory_order_relaxed);\n"));
354         stringBuilder.append("}\n");
355         return stringBuilder.toString();
356     }
357
358     if (nativeFunctionDeclaration.name() == "Sample") {
359         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
360         notImplemented();
361     }
362
363     if (nativeFunctionDeclaration.name() == "Load") {
364         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
365         notImplemented();
366     }
367
368     if (nativeFunctionDeclaration.name() == "GetDimensions") {
369         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
370         notImplemented();
371     }
372
373     if (nativeFunctionDeclaration.name() == "SampleBias") {
374         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
375         notImplemented();
376     }
377
378     if (nativeFunctionDeclaration.name() == "SampleGrad") {
379         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
380         notImplemented();
381     }
382
383     if (nativeFunctionDeclaration.name() == "SampleLevel") {
384         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
385         notImplemented();
386     }
387
388     if (nativeFunctionDeclaration.name() == "Gather") {
389         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
390         notImplemented();
391     }
392
393     if (nativeFunctionDeclaration.name() == "GatherRed") {
394         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
395         notImplemented();
396     }
397
398     if (nativeFunctionDeclaration.name() == "SampleCmp") {
399         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
400         notImplemented();
401     }
402
403     if (nativeFunctionDeclaration.name() == "SampleCmpLevelZero") {
404         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
405         notImplemented();
406     }
407
408     if (nativeFunctionDeclaration.name() == "Store") {
409         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
410         notImplemented();
411     }
412
413     if (nativeFunctionDeclaration.name() == "GatherAlpha") {
414         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
415         notImplemented();
416     }
417
418     if (nativeFunctionDeclaration.name() == "GatherBlue") {
419         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
420         notImplemented();
421     }
422
423     if (nativeFunctionDeclaration.name() == "GatherCmp") {
424         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
425         notImplemented();
426     }
427
428     if (nativeFunctionDeclaration.name() == "GatherCmpRed") {
429         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
430         notImplemented();
431     }
432
433     if (nativeFunctionDeclaration.name() == "GatherGreen") {
434         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
435         notImplemented();
436     }
437
438     ASSERT_NOT_REACHED();
439     return String();
440 }
441
442 } // namespace Metal
443
444 } // namespace WHLSL
445
446 } // namespace WebCore
447
448 #endif