[WHLSL] Hook up common texture functions
[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 static int vectorLength(AST::NativeTypeDeclaration& nativeTypeDeclaration)
85 {
86     int vectorLength = 1;
87     if (!nativeTypeDeclaration.typeArguments().isEmpty()) {
88         ASSERT(nativeTypeDeclaration.typeArguments().size() == 2);
89         ASSERT(WTF::holds_alternative<AST::ConstantExpression>(nativeTypeDeclaration.typeArguments()[1]));
90         vectorLength = WTF::get<AST::ConstantExpression>(nativeTypeDeclaration.typeArguments()[1]).integerLiteral().value();
91     }
92     return vectorLength;
93 }
94
95 static AST::NamedType& vectorInnerType(AST::NativeTypeDeclaration& nativeTypeDeclaration)
96 {
97     if (nativeTypeDeclaration.typeArguments().isEmpty())
98         return nativeTypeDeclaration;
99
100     ASSERT(nativeTypeDeclaration.typeArguments().size() == 2);
101     ASSERT(WTF::holds_alternative<UniqueRef<AST::TypeReference>>(nativeTypeDeclaration.typeArguments()[0]));
102     return WTF::get<UniqueRef<AST::TypeReference>>(nativeTypeDeclaration.typeArguments()[0])->resolvedType();
103 }
104
105 static const char* vectorSuffix(int vectorLength)
106 {
107     switch (vectorLength) {
108     case 1:
109         return "";
110     case 2:
111         return "2";
112     case 3:
113         return "3";
114     default:
115         ASSERT(vectorLength == 4);
116         return "4";
117     }
118 }
119
120 String writeNativeFunction(AST::NativeFunctionDeclaration& nativeFunctionDeclaration, String& outputFunctionName, Intrinsics& intrinsics, TypeNamer& typeNamer, const char* memsetZeroFunctionName)
121 {
122     StringBuilder stringBuilder;
123     if (nativeFunctionDeclaration.isCast()) {
124         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
125         if (!nativeFunctionDeclaration.parameters().size()) {
126             stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, "() {\n"));
127             stringBuilder.append(makeString("    ", metalReturnName, " x;\n"));
128             stringBuilder.append(makeString("    ", memsetZeroFunctionName, "(x);\n"));
129             stringBuilder.append("    return x;\n");
130             stringBuilder.append("}\n");
131             return stringBuilder.toString();
132         }
133
134         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
135         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
136         auto& parameterType = nativeFunctionDeclaration.parameters()[0]->type()->unifyNode();
137         if (is<AST::NamedType>(parameterType)) {
138             auto& parameterNamedType = downcast<AST::NamedType>(parameterType);
139             if (is<AST::NativeTypeDeclaration>(parameterNamedType)) {
140                 auto& parameterNativeTypeDeclaration = downcast<AST::NativeTypeDeclaration>(parameterNamedType);
141                 if (parameterNativeTypeDeclaration.isAtomic()) {
142                     stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"));
143                     stringBuilder.append("    return atomic_load_explicit(&x, memory_order_relaxed);\n");
144                     stringBuilder.append("}\n");
145                     return stringBuilder.toString();
146                 }
147             }
148         }
149
150         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"));
151         stringBuilder.append(makeString("    return static_cast<", metalReturnName, ">(x);\n"));
152         stringBuilder.append("}\n");
153         return stringBuilder.toString();
154     }
155
156     if (nativeFunctionDeclaration.name() == "operator.value") {
157         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
158         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
159         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
160         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"));
161         stringBuilder.append(makeString("    return static_cast<", metalReturnName, ">(x);\n"));
162         stringBuilder.append("}\n");
163         return stringBuilder.toString();
164     }
165
166     // 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.
167     if (nativeFunctionDeclaration.name() == "operator.length") {
168         ASSERT_UNUSED(intrinsics, matches(nativeFunctionDeclaration.type(), intrinsics.uintType()));
169         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
170         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
171         auto& parameterType = nativeFunctionDeclaration.parameters()[0]->type()->unifyNode();
172         auto& unnamedParameterType = downcast<AST::UnnamedType>(parameterType);
173         if (is<AST::ArrayType>(unnamedParameterType)) {
174             auto& arrayParameterType = downcast<AST::ArrayType>(unnamedParameterType);
175             stringBuilder.append(makeString("uint ", outputFunctionName, '(', metalParameterName, " v) {\n"));
176             stringBuilder.append(makeString("    return ", arrayParameterType.numElements(), "u;\n"));
177             stringBuilder.append("}\n");
178             return stringBuilder.toString();
179         }
180
181         ASSERT(is<AST::ArrayReferenceType>(unnamedParameterType));
182         stringBuilder.append(makeString("uint ", outputFunctionName, '(', metalParameterName, " v) {\n"));
183         stringBuilder.append(makeString("    return v.length;\n"));
184         stringBuilder.append("}\n");
185         return stringBuilder.toString();
186     }
187
188     if (nativeFunctionDeclaration.name().startsWith("operator."_str)) {
189         auto mangledFieldName = [&](const String& fieldName) -> String {
190             auto& unifyNode = nativeFunctionDeclaration.parameters()[0]->type()->unifyNode();
191             auto& namedType = downcast<AST::NamedType>(unifyNode);
192             if (is<AST::StructureDefinition>(namedType)) {
193                 auto& structureDefinition = downcast<AST::StructureDefinition>(namedType);
194                 auto* structureElement = structureDefinition.find(fieldName);
195                 ASSERT(structureElement);
196                 return typeNamer.mangledNameForStructureElement(*structureElement);
197             }
198             ASSERT(is<AST::NativeTypeDeclaration>(namedType));
199             return fieldName;
200         };
201
202         if (nativeFunctionDeclaration.name().endsWith("=")) {
203             ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
204             auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
205             auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
206             auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
207             auto fieldName = nativeFunctionDeclaration.name().substring("operator."_str.length());
208             fieldName = fieldName.substring(0, fieldName.length() - 1);
209             auto metalFieldName = mangledFieldName(fieldName);
210             stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " v, ", metalParameter2Name, " n) {\n"));
211             stringBuilder.append(makeString("    v.", metalFieldName, " = n;\n"));
212             stringBuilder.append(makeString("    return v;\n"));
213             stringBuilder.append("}\n");
214             return stringBuilder.toString();
215         }
216
217         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
218         auto fieldName = nativeFunctionDeclaration.name().substring("operator."_str.length());
219         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
220         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
221         auto metalFieldName = mangledFieldName(fieldName);
222         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " v) {\n"));
223         stringBuilder.append(makeString("    return v.", metalFieldName, ";\n"));
224         stringBuilder.append("}\n");
225         return stringBuilder.toString();
226     }
227
228     if (nativeFunctionDeclaration.name().startsWith("operator&."_str)) {
229         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
230         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
231         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
232         auto fieldName = nativeFunctionDeclaration.name().substring("operator&."_str.length());
233
234         String metalFieldName;
235         auto& unnamedType = *nativeFunctionDeclaration.parameters()[0]->type();
236         auto& unifyNode = downcast<AST::PointerType>(unnamedType).elementType().unifyNode();
237         auto& namedType = downcast<AST::NamedType>(unifyNode);
238         if (is<AST::StructureDefinition>(namedType)) {
239             auto& structureDefinition = downcast<AST::StructureDefinition>(namedType);
240             auto* structureElement = structureDefinition.find(fieldName);
241             ASSERT(structureElement);
242             metalFieldName = typeNamer.mangledNameForStructureElement(*structureElement);
243         } else
244             metalFieldName = fieldName;
245
246         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " v) {\n"));
247         stringBuilder.append(makeString("    return &(v->", metalFieldName, ");\n"));
248         stringBuilder.append("}\n");
249         return stringBuilder.toString();
250     }
251
252     if (nativeFunctionDeclaration.name() == "operator&[]") {
253         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
254         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
255         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
256         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
257         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " v, ", metalParameter2Name, " n) {\n"));
258         stringBuilder.append("    if (n < v.length) return &(v.pointer[n]);\n");
259         stringBuilder.append("    return nullptr;\n");
260         stringBuilder.append("}\n");
261         return stringBuilder.toString();
262     }
263
264     auto numberOfMatrixRows = [&] {
265         auto& typeReference = downcast<AST::TypeReference>(*nativeFunctionDeclaration.parameters()[0]->type());
266         auto& matrixType = downcast<AST::NativeTypeDeclaration>(downcast<AST::TypeReference>(downcast<AST::TypeDefinition>(typeReference.resolvedType()).type()).resolvedType());
267         ASSERT(matrixType.name() == "matrix");
268         ASSERT(matrixType.typeArguments().size() == 3);
269         return String::number(WTF::get<AST::ConstantExpression>(matrixType.typeArguments()[1]).integerLiteral().value());
270     };
271
272     if (nativeFunctionDeclaration.name() == "operator[]") {
273         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
274         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
275         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
276         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
277         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " m, ", metalParameter2Name, " i) {\n"));
278         stringBuilder.append(makeString("    if (i < ", numberOfMatrixRows(), ") return m[i];\n"));
279         stringBuilder.append(makeString("    return ", metalReturnName, "(0);\n"));
280         stringBuilder.append("}\n");
281         return stringBuilder.toString();
282     }
283
284     if (nativeFunctionDeclaration.name() == "operator[]=") {
285         ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
286         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
287         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
288         auto metalParameter3Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[2]->type());
289         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
290         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " m, ", metalParameter2Name, " i, ", metalParameter3Name, " v) {\n"));
291         stringBuilder.append(makeString("    if (i < ", numberOfMatrixRows(), ") m[i] = v;\n"));
292         stringBuilder.append("    return m;\n");
293         stringBuilder.append("}\n");
294         return stringBuilder.toString();
295     }
296
297     if (nativeFunctionDeclaration.isOperator()) {
298         if (nativeFunctionDeclaration.parameters().size() == 1) {
299             auto operatorName = nativeFunctionDeclaration.name().substring("operator"_str.length());
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 ", operatorName, "x;\n"));
304             stringBuilder.append("}\n");
305             return stringBuilder.toString();
306         }
307
308         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
309         auto operatorName = nativeFunctionDeclaration.name().substring("operator"_str.length());
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 x ", operatorName, " y;\n"));
315         stringBuilder.append("}\n");
316         return stringBuilder.toString();
317     }
318
319     if (nativeFunctionDeclaration.name() == "cos"
320         || nativeFunctionDeclaration.name() == "sin"
321         || nativeFunctionDeclaration.name() == "tan"
322         || nativeFunctionDeclaration.name() == "acos"
323         || nativeFunctionDeclaration.name() == "asin"
324         || nativeFunctionDeclaration.name() == "atan"
325         || nativeFunctionDeclaration.name() == "cosh"
326         || nativeFunctionDeclaration.name() == "sinh"
327         || nativeFunctionDeclaration.name() == "tanh"
328         || nativeFunctionDeclaration.name() == "ceil"
329         || nativeFunctionDeclaration.name() == "exp"
330         || nativeFunctionDeclaration.name() == "floor"
331         || nativeFunctionDeclaration.name() == "log"
332         || nativeFunctionDeclaration.name() == "round"
333         || nativeFunctionDeclaration.name() == "trunc"
334         || nativeFunctionDeclaration.name() == "ddx"
335         || nativeFunctionDeclaration.name() == "ddy"
336         || nativeFunctionDeclaration.name() == "isnormal"
337         || nativeFunctionDeclaration.name() == "isfinite"
338         || nativeFunctionDeclaration.name() == "isinf"
339         || nativeFunctionDeclaration.name() == "isnan"
340         || nativeFunctionDeclaration.name() == "asint"
341         || nativeFunctionDeclaration.name() == "asuint"
342         || nativeFunctionDeclaration.name() == "asfloat") {
343         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
344         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
345         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
346         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"));
347         stringBuilder.append(makeString("    return ", mapFunctionName(nativeFunctionDeclaration.name()), "(x);\n"));
348         stringBuilder.append("}\n");
349         return stringBuilder.toString();
350     }
351
352     if (nativeFunctionDeclaration.name() == "pow" || nativeFunctionDeclaration.name() == "atan2") {
353         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
354         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
355         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
356         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
357         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " x, ", metalParameter2Name, " y) {\n"));
358         stringBuilder.append(makeString("    return ", nativeFunctionDeclaration.name(), "(x, y);\n"));
359         stringBuilder.append("}\n");
360         return stringBuilder.toString();
361     }
362
363     if (nativeFunctionDeclaration.name() == "f16tof32" || nativeFunctionDeclaration.name() == "f32tof16") {
364         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
365         notImplemented();
366     }
367
368     if (nativeFunctionDeclaration.name() == "AllMemoryBarrierWithGroupSync") {
369         ASSERT(!nativeFunctionDeclaration.parameters().size());
370         stringBuilder.append(makeString("void ", outputFunctionName, "() {\n"));
371         stringBuilder.append("    threadgroup_barrier(mem_flags::mem_device);\n");
372         stringBuilder.append("    threadgroup_barrier(mem_flags::mem_threadgroup);\n");
373         stringBuilder.append("    threadgroup_barrier(mem_flags::mem_texture);\n");
374         stringBuilder.append("}\n");
375         return stringBuilder.toString();
376     }
377
378     if (nativeFunctionDeclaration.name() == "DeviceMemoryBarrierWithGroupSync") {
379         ASSERT(!nativeFunctionDeclaration.parameters().size());
380         stringBuilder.append(makeString("void ", outputFunctionName, "() {\n"));
381         stringBuilder.append("    threadgroup_barrier(mem_flags::mem_device);\n");
382         stringBuilder.append("}\n");
383         return stringBuilder.toString();
384     }
385
386     if (nativeFunctionDeclaration.name() == "GroupMemoryBarrierWithGroupSync") {
387         ASSERT(!nativeFunctionDeclaration.parameters().size());
388         stringBuilder.append(makeString("void ", outputFunctionName, "() {\n"));
389         stringBuilder.append("    threadgroup_barrier(mem_flags::mem_threadgroup);\n");
390         stringBuilder.append("}\n");
391         return stringBuilder.toString();
392     }
393
394     if (nativeFunctionDeclaration.name().startsWith("Interlocked"_str)) {
395         if (nativeFunctionDeclaration.name() == "InterlockedCompareExchange") {
396             ASSERT(nativeFunctionDeclaration.parameters().size() == 4);
397             auto& firstArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0]->type());
398             auto firstArgumentAddressSpace = firstArgumentPointer.addressSpace();
399             auto firstArgumentPointee = typeNamer.mangledNameForType(firstArgumentPointer.elementType());
400             auto secondArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
401             auto thirdArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[2]->type());
402             auto& fourthArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[3]->type());
403             auto fourthArgumentAddressSpace = fourthArgumentPointer.addressSpace();
404             auto fourthArgumentPointee = typeNamer.mangledNameForType(fourthArgumentPointer.elementType());
405             stringBuilder.append(makeString("void ", outputFunctionName, '(', toString(firstArgumentAddressSpace), ' ', firstArgumentPointee, "* object, ", secondArgument, " compare, ", thirdArgument, " desired, ", toString(fourthArgumentAddressSpace), ' ', fourthArgumentPointee, "* out) {\n"));
406             stringBuilder.append("    atomic_compare_exchange_weak_explicit(object, &compare, desired, memory_order_relaxed);\n");
407             stringBuilder.append("    *out = compare;\n");
408             stringBuilder.append("}\n");
409             return stringBuilder.toString();
410         }
411
412         ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
413         auto& firstArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0]->type());
414         auto firstArgumentAddressSpace = firstArgumentPointer.addressSpace();
415         auto firstArgumentPointee = typeNamer.mangledNameForType(firstArgumentPointer.elementType());
416         auto secondArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
417         auto& thirdArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[2]->type());
418         auto thirdArgumentAddressSpace = thirdArgumentPointer.addressSpace();
419         auto thirdArgumentPointee = typeNamer.mangledNameForType(thirdArgumentPointer.elementType());
420         auto name = atomicName(nativeFunctionDeclaration.name().substring("Interlocked"_str.length()));
421         stringBuilder.append(makeString("void ", outputFunctionName, '(', toString(firstArgumentAddressSpace), ' ', firstArgumentPointee, "* object, ", secondArgument, " operand, ", toString(thirdArgumentAddressSpace), ' ', thirdArgumentPointee, "* out) {\n"));
422         stringBuilder.append(makeString("    *out = atomic_fetch_", name, "_explicit(object, operand, memory_order_relaxed);\n"));
423         stringBuilder.append("}\n");
424         return stringBuilder.toString();
425     }
426
427     if (nativeFunctionDeclaration.name() == "Sample") {
428         ASSERT(nativeFunctionDeclaration.parameters().size() == 3 || nativeFunctionDeclaration.parameters().size() == 4);
429         
430         auto& textureType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[0]->type()->unifyNode()));
431         auto& locationType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[2]->type()->unifyNode()));
432         auto locationVectorLength = vectorLength(locationType);
433         auto& returnType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.type().unifyNode()));
434         auto returnVectorLength = vectorLength(returnType);
435
436         auto metalParameter1Name = typeNamer.mangledNameForType(textureType);
437         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
438         auto metalParameter3Name = typeNamer.mangledNameForType(locationType);
439         String metalParameter4Name;
440         if (nativeFunctionDeclaration.parameters().size() == 4)
441             metalParameter4Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[3]->type());
442         auto metalReturnName = typeNamer.mangledNameForType(returnType);
443         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " theTexture, ", metalParameter2Name, " theSampler, ", metalParameter3Name, " location"));
444         if (!metalParameter4Name.isNull())
445             stringBuilder.append(makeString(", ", metalParameter4Name, " offset"));
446         stringBuilder.append(") {\n");
447         stringBuilder.append("    return theTexture.sample(theSampler, ");
448         if (textureType.isTextureArray()) {
449             ASSERT(locationVectorLength > 1);
450             stringBuilder.append(makeString("location.", "xyzw"_str.substring(0, locationVectorLength - 1), ", location.", "xyzw"_str.substring(locationVectorLength - 1, 1)));
451         } else
452             stringBuilder.append("location");
453         if (!metalParameter4Name.isNull())
454             stringBuilder.append(", offset");
455         stringBuilder.append(")");
456         if (!textureType.isDepthTexture())
457             stringBuilder.append(makeString(".", "xyzw"_str.substring(0, returnVectorLength)));
458         stringBuilder.append(";\n");
459         stringBuilder.append("}\n");
460         return stringBuilder.toString();
461     }
462
463     if (nativeFunctionDeclaration.name() == "Load") {
464         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
465         
466         auto& textureType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[0]->type()->unifyNode()));
467         auto& locationType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[1]->type()->unifyNode()));
468         auto locationVectorLength = vectorLength(locationType);
469         auto& returnType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.type().unifyNode()));
470         auto returnVectorLength = vectorLength(returnType);
471
472         auto metalParameter1Name = typeNamer.mangledNameForType(textureType);
473         auto metalParameter2Name = typeNamer.mangledNameForType(locationType);
474         auto metalReturnName = typeNamer.mangledNameForType(returnType);
475         stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " theTexture, ", metalParameter2Name, " location) {\n"));
476         if (textureType.isTextureArray()) {
477             ASSERT(locationVectorLength > 1);
478             String dimensions[] = { "width"_str, "height"_str, "depth"_str };
479             for (int i = 0; i < locationVectorLength - 1; ++i) {
480                 auto suffix = "xyzw"_str.substring(i, 1);
481                 stringBuilder.append(makeString("    if (location.", suffix, " < 0 || static_cast<uint32_t>(location.", suffix, ") >= theTexture.get_", dimensions[i], "()) return ", metalReturnName, "(0);\n"));
482             }
483             auto suffix = "xyzw"_str.substring(locationVectorLength - 1, 1);
484             stringBuilder.append(makeString("    if (location.", suffix, " < 0 || static_cast<uint32_t>(location.", suffix, ") >= theTexture.get_array_size()) return ", metalReturnName, "(0);\n"));
485         } else {
486             if (locationVectorLength == 1)
487                 stringBuilder.append(makeString("    if (location < 0 || static_cast<uint32_t>(location) >= theTexture.get_width()) return ", metalReturnName, "(0);\n"));
488             else {
489                 stringBuilder.append(makeString("    if (location.x < 0 || static_cast<uint32_t>(location.x) >= theTexture.get_width()) return ", metalReturnName, "(0);\n"));
490                 stringBuilder.append(makeString("    if (location.y < 0 || static_cast<uint32_t>(location.y) >= theTexture.get_height()) return ", metalReturnName, "(0);\n"));
491                 if (locationVectorLength >= 3)
492                     stringBuilder.append(makeString("    if (location.z < 0 || static_cast<uint32_t>(location.z) >= theTexture.get_depth()) return ", metalReturnName, "(0);\n"));
493             }
494         }
495         stringBuilder.append("    return theTexture.read(");
496         if (textureType.isTextureArray()) {
497             ASSERT(locationVectorLength > 1);
498             stringBuilder.append(makeString("uint", vectorSuffix(locationVectorLength - 1), "(location.", "xyzw"_str.substring(0, locationVectorLength - 1), "), uint(location.", "xyzw"_str.substring(locationVectorLength - 1, 1), ")"));
499         } else
500             stringBuilder.append(makeString("uint", vectorSuffix(locationVectorLength), "(location)"));
501         stringBuilder.append(")");
502         if (!textureType.isDepthTexture())
503             stringBuilder.append(makeString(".", "xyzw"_str.substring(0, returnVectorLength)));
504         stringBuilder.append(";\n");
505         stringBuilder.append("}\n");
506         return stringBuilder.toString();
507     }
508
509     if (nativeFunctionDeclaration.name() == "GetDimensions") {
510         auto& textureType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[0]->type()->unifyNode()));
511
512         size_t index = 1;
513         if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
514             ++index;
515         auto widthTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
516         ++index;
517         String heightTypeName;
518         if (textureType.textureDimension() >= 2) {
519             heightTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
520             ++index;
521         }
522         String depthTypeName;
523         if (textureType.textureDimension() >= 3) {
524             depthTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
525             ++index;
526         }
527         String elementsTypeName;
528         if (textureType.isTextureArray()) {
529             elementsTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
530             ++index;
531         }
532         String numberOfLevelsTypeName;
533         if (!textureType.isWritableTexture()) {
534             numberOfLevelsTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
535             ++index;
536         }
537         ASSERT(index == nativeFunctionDeclaration.parameters().size());
538
539         auto metalParameter1Name = typeNamer.mangledNameForType(textureType);
540         stringBuilder.append(makeString("void ", outputFunctionName, '(', metalParameter1Name, " theTexture"));
541         if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
542             stringBuilder.append(", uint mipLevel");
543         stringBuilder.append(makeString(", ", widthTypeName, " width"));
544         if (!heightTypeName.isNull())
545             stringBuilder.append(makeString(", ", heightTypeName, " height"));
546         if (!depthTypeName.isNull())
547             stringBuilder.append(makeString(", ", depthTypeName, " depth"));
548         if (!elementsTypeName.isNull())
549             stringBuilder.append(makeString(", ", elementsTypeName, " elements"));
550         if (!numberOfLevelsTypeName.isNull())
551             stringBuilder.append(makeString(", ", numberOfLevelsTypeName, " numberOfLevels"));
552         stringBuilder.append(") {\n");
553         stringBuilder.append("    if (width)\n");
554         stringBuilder.append("        *width = theTexture.get_width(");
555         if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
556             stringBuilder.append("mipLevel");
557         stringBuilder.append(");\n");
558         if (!heightTypeName.isNull()) {
559             stringBuilder.append("    if (height)\n");
560             stringBuilder.append("        *height = theTexture.get_height(");
561             if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
562                 stringBuilder.append("mipLevel");
563             stringBuilder.append(");\n");
564         }
565         if (!depthTypeName.isNull()) {
566             stringBuilder.append("    if (depth)\n");
567             stringBuilder.append("        *depth = theTexture.get_depth(");
568             if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
569                 stringBuilder.append("mipLevel");
570             stringBuilder.append(");\n");
571         }
572         if (!elementsTypeName.isNull()) {
573             stringBuilder.append("    if (elements)\n");
574             stringBuilder.append("        *elements = theTexture.get_array_size();\n");
575         }
576         if (!numberOfLevelsTypeName.isNull()) {
577             stringBuilder.append("    if (numberOfLevels)\n");
578             stringBuilder.append("        *numberOfLevels = theTexture.get_num_mip_levels();\n");
579         }
580         stringBuilder.append("}\n");
581         return stringBuilder.toString();
582     }
583
584     if (nativeFunctionDeclaration.name() == "SampleBias") {
585         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
586         notImplemented();
587     }
588
589     if (nativeFunctionDeclaration.name() == "SampleGrad") {
590         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
591         notImplemented();
592     }
593
594     if (nativeFunctionDeclaration.name() == "SampleLevel") {
595         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
596         notImplemented();
597     }
598
599     if (nativeFunctionDeclaration.name() == "Gather") {
600         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
601         notImplemented();
602     }
603
604     if (nativeFunctionDeclaration.name() == "GatherRed") {
605         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
606         notImplemented();
607     }
608
609     if (nativeFunctionDeclaration.name() == "SampleCmp") {
610         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
611         notImplemented();
612     }
613
614     if (nativeFunctionDeclaration.name() == "SampleCmpLevelZero") {
615         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
616         notImplemented();
617     }
618
619     if (nativeFunctionDeclaration.name() == "Store") {
620         ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
621         
622         auto& textureType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[0]->type()->unifyNode()));
623         auto& itemType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[1]->type()->unifyNode()));
624         auto& itemVectorInnerType = vectorInnerType(itemType);
625         auto itemVectorLength = vectorLength(itemType);
626         auto& locationType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[2]->type()->unifyNode()));
627         auto locationVectorLength = vectorLength(locationType);
628
629         auto metalParameter1Name = typeNamer.mangledNameForType(textureType);
630         auto metalParameter2Name = typeNamer.mangledNameForType(itemType);
631         auto metalParameter3Name = typeNamer.mangledNameForType(locationType);
632         auto metalInnerTypeName = typeNamer.mangledNameForType(itemVectorInnerType);
633         stringBuilder.append(makeString("void ", outputFunctionName, '(', metalParameter1Name, " theTexture, ", metalParameter2Name, " item, ", metalParameter3Name, " location) {\n"));
634         if (textureType.isTextureArray()) {
635             ASSERT(locationVectorLength > 1);
636             String dimensions[] = { "width"_str, "height"_str, "depth"_str };
637             for (int i = 0; i < locationVectorLength - 1; ++i) {
638                 auto suffix = "xyzw"_str.substring(i, 1);
639                 stringBuilder.append(makeString("    if (location.", suffix, " >= theTexture.get_", dimensions[i], "()) return;\n"));
640             }
641             auto suffix = "xyzw"_str.substring(locationVectorLength - 1, 1);
642             stringBuilder.append(makeString("    if (location.", suffix, " >= theTexture.get_array_size()) return;\n"));
643         } else {
644             if (locationVectorLength == 1)
645                 stringBuilder.append(makeString("    if (location >= theTexture.get_width()) return;\n"));
646             else {
647                 stringBuilder.append(makeString("    if (location.x >= theTexture.get_width()) return;\n"));
648                 stringBuilder.append(makeString("    if (location.y >= theTexture.get_height()) return;\n"));
649                 if (locationVectorLength >= 3)
650                     stringBuilder.append(makeString("    if (location.z >= theTexture.get_depth()) return;\n"));
651             }
652         }
653         stringBuilder.append(makeString("    theTexture.write(vec<", metalInnerTypeName, ", 4>(item"));
654         for (int i = 0; i < 4 - itemVectorLength; ++i)
655             stringBuilder.append(", 0");
656         stringBuilder.append("), ");
657         if (textureType.isTextureArray()) {
658             ASSERT(locationVectorLength > 1);
659             stringBuilder.append(makeString("uint", vectorSuffix(locationVectorLength - 1), "(location.", "xyzw"_str.substring(0, locationVectorLength - 1), "), uint(location.", "xyzw"_str.substring(locationVectorLength - 1, 1), ")"));
660         } else
661             stringBuilder.append(makeString("uint", vectorSuffix(locationVectorLength), "(location)"));
662         stringBuilder.append(");\n");
663         stringBuilder.append("}\n");
664         return stringBuilder.toString();
665     }
666
667     if (nativeFunctionDeclaration.name() == "GatherAlpha") {
668         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
669         notImplemented();
670     }
671
672     if (nativeFunctionDeclaration.name() == "GatherBlue") {
673         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
674         notImplemented();
675     }
676
677     if (nativeFunctionDeclaration.name() == "GatherCmp") {
678         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
679         notImplemented();
680     }
681
682     if (nativeFunctionDeclaration.name() == "GatherCmpRed") {
683         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
684         notImplemented();
685     }
686
687     if (nativeFunctionDeclaration.name() == "GatherGreen") {
688         // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195813 Implement this
689         notImplemented();
690     }
691
692     ASSERT_NOT_REACHED();
693     return String();
694 }
695
696 } // namespace Metal
697
698 } // namespace WHLSL
699
700 } // namespace WebCore
701
702 #endif