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
45namespace WebCore {
46
47namespace WHLSL {
48
49namespace Metal {
50
51static 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
66static 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
81static 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
99String 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
466