[WHLSL] Hook up the compiler to our WebGPU implementation
[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     if (nativeFunctionDeclaration.name() == "operator.length") {
146         ASSERT_UNUSED(intrinsics, matches(nativeFunctionDeclaration.type(), intrinsics.uintType()));
147         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
148         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0].type());
149         auto& parameterType = nativeFunctionDeclaration.parameters()[0].type()->unifyNode();
150         ASSERT(is<AST::UnnamedType>(parameterType));
151         auto& unnamedParameterType = downcast<AST::UnnamedType>(parameterType);
152         if (is<AST::ArrayType>(unnamedParameterType)) {
153             auto& arrayParameterType = downcast<AST::ArrayType>(unnamedParameterType);
154             stringBuilder.append(makeString("uint ", outputFunctionName, '(', metalParameterName, " v) {\n"));
155             stringBuilder.append(makeString("    return ", arrayParameterType.numElements(), "u;\n"));
156             stringBuilder.append("}\n");
157             return stringBuilder.toString();
158         }
159
160         ASSERT(is<AST::ArrayReferenceType>(unnamedParameterType));
161         stringBuilder.append(makeString("uint ", outputFunctionName, '(', metalParameterName, " v) {\n"));
162         stringBuilder.append(makeString("    return v.length;\n"));
163         stringBuilder.append("}\n");
164         return stringBuilder.toString();
165     }
166
167     if (nativeFunctionDeclaration.name().startsWith("operator."_str)) {
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             stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " v, ", metalParameter2Name, " n) {\n"));
176             stringBuilder.append(makeString("    v.", fieldName, " = n;\n"));
177             stringBuilder.append(makeString("    return v;\n"));
178             stringBuilder.append("}\n");
179             return stringBuilder.toString();
180         }
181
182         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
183         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0].type());
184         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
185         auto fieldName = nativeFunctionDeclaration.name().substring("operator."_str.length());
186         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " v) {\n"));
187         stringBuilder.append(makeString("    return v.", fieldName, ";\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         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " v) {\n"));
198         stringBuilder.append(makeString("    return &(v->", fieldName, ");\n"));
199         stringBuilder.append("}\n");
200         return stringBuilder.toString();
201     }
202
203     if (nativeFunctionDeclaration.name() == "operator[]") {
204         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
205         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0].type());
206         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1].type());
207         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
208         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " m, ", metalParameter2Name, " i) {\n"));
209         stringBuilder.append(makeString("    return m[i];\n"));
210         stringBuilder.append("}\n");
211         return stringBuilder.toString();
212     }
213
214     if (nativeFunctionDeclaration.name() == "operator&[]") {
215         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
216         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0].type());
217         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0].type());
218         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
219         auto fieldName = nativeFunctionDeclaration.name().substring("operator&[]."_str.length());
220         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " v, ", metalParameter2Name, " n) {\n"));
221         stringBuilder.append(makeString("    return &(v.pointer[n]);\n"));
222         stringBuilder.append("}\n");
223         return stringBuilder.toString();
224     }
225
226     if (nativeFunctionDeclaration.name() == "operator[]=") {
227         ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
228         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0].type());
229         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1].type());
230         auto metalParameter3Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[2].type());
231         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
232         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " m, ", metalParameter2Name, " i, ", metalParameter3Name, " v) {\n"));
233         stringBuilder.append(makeString("    m[i] = v;\n"));
234         stringBuilder.append(makeString("    return m;\n"));
235         stringBuilder.append("}\n");
236         return stringBuilder.toString();
237     }
238
239     if (nativeFunctionDeclaration.isOperator()) {
240         if (nativeFunctionDeclaration.parameters().size() == 1) {
241             auto operatorName = nativeFunctionDeclaration.name().substring("operator"_str.length());
242             auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0].type());
243             auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
244             stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"));
245             stringBuilder.append(makeString("    return ", operatorName, "x;\n"));
246             stringBuilder.append("}\n");
247             return stringBuilder.toString();
248         }
249
250         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
251         auto operatorName = nativeFunctionDeclaration.name().substring("operator"_str.length());
252         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0].type());
253         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1].type());
254         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
255         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " x, ", metalParameter2Name, " y) {\n"));
256         stringBuilder.append(makeString("    return x ", operatorName, " y;\n"));
257         stringBuilder.append("}\n");
258         return stringBuilder.toString();
259     }
260
261     if (nativeFunctionDeclaration.name() == "cos"
262         || nativeFunctionDeclaration.name() == "sin"
263         || nativeFunctionDeclaration.name() == "tan"
264         || nativeFunctionDeclaration.name() == "acos"
265         || nativeFunctionDeclaration.name() == "asin"
266         || nativeFunctionDeclaration.name() == "atan"
267         || nativeFunctionDeclaration.name() == "cosh"
268         || nativeFunctionDeclaration.name() == "sinh"
269         || nativeFunctionDeclaration.name() == "tanh"
270         || nativeFunctionDeclaration.name() == "ceil"
271         || nativeFunctionDeclaration.name() == "exp"
272         || nativeFunctionDeclaration.name() == "floor"
273         || nativeFunctionDeclaration.name() == "log"
274         || nativeFunctionDeclaration.name() == "round"
275         || nativeFunctionDeclaration.name() == "trunc"
276         || nativeFunctionDeclaration.name() == "ddx"
277         || nativeFunctionDeclaration.name() == "ddy"
278         || nativeFunctionDeclaration.name() == "isnormal"
279         || nativeFunctionDeclaration.name() == "isfinite"
280         || nativeFunctionDeclaration.name() == "isinf"
281         || nativeFunctionDeclaration.name() == "isnan"
282         || nativeFunctionDeclaration.name() == "asint"
283         || nativeFunctionDeclaration.name() == "asuint"
284         || nativeFunctionDeclaration.name() == "asfloat") {
285         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
286         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0].type());
287         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
288         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"));
289         stringBuilder.append(makeString("    return ", mapFunctionName(nativeFunctionDeclaration.name()), "(x);\n"));
290         stringBuilder.append("}\n");
291         return stringBuilder.toString();
292     }
293
294     if (nativeFunctionDeclaration.name() == "pow" || nativeFunctionDeclaration.name() == "atan2") {
295         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
296         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0].type());
297         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1].type());
298         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
299         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " x, ", metalParameter2Name, " y) {\n"));
300         stringBuilder.append(makeString("    return ", nativeFunctionDeclaration.name(), "(x, y);\n"));
301         stringBuilder.append("}\n");
302         return stringBuilder.toString();
303     }
304
305     if (nativeFunctionDeclaration.name() == "f16tof32" || nativeFunctionDeclaration.name() == "f32tof16") {
306         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
307         notImplemented();
308     }
309
310     if (nativeFunctionDeclaration.name() == "AllMemoryBarrierWithGroupSync") {
311         ASSERT(!nativeFunctionDeclaration.parameters().size());
312         stringBuilder.append(makeString("void ", outputFunctionName, "() {\n"));
313         stringBuilder.append("    threadgroup_barrier(mem_flags::mem_device);\n");
314         stringBuilder.append("    threadgroup_barrier(mem_flags::mem_threadgroup);\n");
315         stringBuilder.append("    threadgroup_barrier(mem_flags::mem_texture);\n");
316         stringBuilder.append("}\n");
317         return stringBuilder.toString();
318     }
319
320     if (nativeFunctionDeclaration.name() == "DeviceMemoryBarrierWithGroupSync") {
321         ASSERT(!nativeFunctionDeclaration.parameters().size());
322         stringBuilder.append(makeString("void ", outputFunctionName, "() {\n"));
323         stringBuilder.append("    threadgroup_barrier(mem_flags::mem_device);\n");
324         stringBuilder.append("}\n");
325         return stringBuilder.toString();
326     }
327
328     if (nativeFunctionDeclaration.name() == "GroupMemoryBarrierWithGroupSync") {
329         ASSERT(!nativeFunctionDeclaration.parameters().size());
330         stringBuilder.append(makeString("void ", outputFunctionName, "() {\n"));
331         stringBuilder.append("    threadgroup_barrier(mem_flags::mem_threadgroup);\n");
332         stringBuilder.append("}\n");
333         return stringBuilder.toString();
334     }
335
336     if (nativeFunctionDeclaration.name().startsWith("Interlocked"_str)) {
337         if (nativeFunctionDeclaration.name() == "InterlockedCompareExchange") {
338             ASSERT(nativeFunctionDeclaration.parameters().size() == 4);
339             ASSERT(is<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0].type()));
340             auto& firstArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0].type());
341             auto firstArgumentAddressSpace = firstArgumentPointer.addressSpace();
342             auto firstArgumentPointee = typeNamer.mangledNameForType(firstArgumentPointer.elementType());
343             auto secondArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1].type());
344             auto thirdArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[2].type());
345             ASSERT(is<AST::PointerType>(*nativeFunctionDeclaration.parameters()[3].type()));
346             auto& fourthArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[3].type());
347             auto fourthArgumentAddressSpace = fourthArgumentPointer.addressSpace();
348             auto fourthArgumentPointee = typeNamer.mangledNameForType(fourthArgumentPointer.elementType());
349             stringBuilder.append(makeString("void ", outputFunctionName, '(', convertAddressSpace(firstArgumentAddressSpace), ' ', firstArgumentPointee, "* object, ", secondArgument, " compare, ", thirdArgument, " desired, ", convertAddressSpace(fourthArgumentAddressSpace), ' ', fourthArgumentPointee, "* out) {\n"));
350             stringBuilder.append("    atomic_compare_exchange_weak_explicit(object, &compare, desired, memory_order_relaxed);\n");
351             stringBuilder.append("    *out = compare;\n");
352             stringBuilder.append("}\n");
353             return stringBuilder.toString();
354         }
355
356         ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
357         ASSERT(is<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0].type()));
358         auto& firstArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0].type());
359         auto firstArgumentAddressSpace = firstArgumentPointer.addressSpace();
360         auto firstArgumentPointee = typeNamer.mangledNameForType(firstArgumentPointer.elementType());
361         auto secondArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1].type());
362         ASSERT(is<AST::PointerType>(*nativeFunctionDeclaration.parameters()[2].type()));
363         auto& thirdArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[2].type());
364         auto thirdArgumentAddressSpace = thirdArgumentPointer.addressSpace();
365         auto thirdArgumentPointee = typeNamer.mangledNameForType(thirdArgumentPointer.elementType());
366         auto name = atomicName(nativeFunctionDeclaration.name().substring("Interlocked"_str.length()));
367         stringBuilder.append(makeString("void ", outputFunctionName, '(', convertAddressSpace(firstArgumentAddressSpace), ' ', firstArgumentPointee, "* object, ", secondArgument, " operand, ", convertAddressSpace(thirdArgumentAddressSpace), ' ', thirdArgumentPointee, "* out) {\n"));
368         stringBuilder.append(makeString("    *out = atomic_fetch_", name, "_explicit(object, operand, memory_order_relaxed);\n"));
369         stringBuilder.append("}\n");
370         return stringBuilder.toString();
371     }
372
373     if (nativeFunctionDeclaration.name() == "Sample") {
374         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
375         notImplemented();
376     }
377
378     if (nativeFunctionDeclaration.name() == "Load") {
379         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
380         notImplemented();
381     }
382
383     if (nativeFunctionDeclaration.name() == "GetDimensions") {
384         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
385         notImplemented();
386     }
387
388     if (nativeFunctionDeclaration.name() == "SampleBias") {
389         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
390         notImplemented();
391     }
392
393     if (nativeFunctionDeclaration.name() == "SampleGrad") {
394         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
395         notImplemented();
396     }
397
398     if (nativeFunctionDeclaration.name() == "SampleLevel") {
399         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
400         notImplemented();
401     }
402
403     if (nativeFunctionDeclaration.name() == "Gather") {
404         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
405         notImplemented();
406     }
407
408     if (nativeFunctionDeclaration.name() == "GatherRed") {
409         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
410         notImplemented();
411     }
412
413     if (nativeFunctionDeclaration.name() == "SampleCmp") {
414         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
415         notImplemented();
416     }
417
418     if (nativeFunctionDeclaration.name() == "SampleCmpLevelZero") {
419         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
420         notImplemented();
421     }
422
423     if (nativeFunctionDeclaration.name() == "Store") {
424         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
425         notImplemented();
426     }
427
428     if (nativeFunctionDeclaration.name() == "GatherAlpha") {
429         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
430         notImplemented();
431     }
432
433     if (nativeFunctionDeclaration.name() == "GatherBlue") {
434         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
435         notImplemented();
436     }
437
438     if (nativeFunctionDeclaration.name() == "GatherCmp") {
439         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
440         notImplemented();
441     }
442
443     if (nativeFunctionDeclaration.name() == "GatherCmpRed") {
444         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
445         notImplemented();
446     }
447
448     if (nativeFunctionDeclaration.name() == "GatherGreen") {
449         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
450         notImplemented();
451     }
452
453     // FIXME: Add all the functions that the compiler generated.
454
455     ASSERT_NOT_REACHED();
456     return String();
457 }
458
459 } // namespace Metal
460
461 } // namespace WHLSL
462
463 } // namespace WebCore
464
465 #endif