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