Modified: trunk/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.cpp (248279 => 248280)
--- trunk/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.cpp 2019-08-05 23:40:40 UTC (rev 248279)
+++ trunk/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.cpp 2019-08-06 00:03:34 UTC (rev 248280)
@@ -120,26 +120,20 @@
}
}
-String writeNativeFunction(AST::NativeFunctionDeclaration& nativeFunctionDeclaration, MangledFunctionName outputFunctionName, Intrinsics& intrinsics, TypeNamer& typeNamer)
+void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDeclaration& nativeFunctionDeclaration, MangledVariableName returnName, const Vector<MangledVariableName>& args, Intrinsics& intrinsics, TypeNamer& typeNamer)
{
- StringBuilder stringBuilder;
if (nativeFunctionDeclaration.isCast()) {
auto& returnType = nativeFunctionDeclaration.type();
auto metalReturnName = typeNamer.mangledNameForType(returnType);
if (!nativeFunctionDeclaration.parameters().size()) {
- stringBuilder.flexibleAppend(
- metalReturnName, ' ', outputFunctionName, "() {\n"
- " ", metalReturnName, " x = { };\n"
- " return x;\n"
- "}\n"
- );
- return stringBuilder.toString();
+ stringBuilder.flexibleAppend(returnName, " = { };\n");
+ return;
}
ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
auto& parameterType = *nativeFunctionDeclaration.parameters()[0]->type();
auto metalParameterName = typeNamer.mangledNameForType(parameterType);
- stringBuilder.flexibleAppend(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n");
+ stringBuilder.flexibleAppend("{\n", metalParameterName, " x = ", args[0], ";\n");
{
auto isEnumerationDefinition = [] (auto& type) {
@@ -160,22 +154,9 @@
}
stringBuilder.flexibleAppend(
- " return static_cast<", metalReturnName, ">(x);\n"
- "}\n"
- );
- return stringBuilder.toString();
- }
+ returnName, " = static_cast<", metalReturnName, ">(x);\n}\n");
- if (nativeFunctionDeclaration.name() == "operator.value") {
- ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
- auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
- auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
- stringBuilder.flexibleAppend(
- metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"
- " return static_cast<", metalReturnName, ">(x);\n"
- "}\n"
- );
- return stringBuilder.toString();
+ return;
}
// 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.
@@ -182,30 +163,23 @@
if (nativeFunctionDeclaration.name() == "operator.length") {
ASSERT_UNUSED(intrinsics, matches(nativeFunctionDeclaration.type(), intrinsics.uintType()));
ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
- auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
auto& parameterType = nativeFunctionDeclaration.parameters()[0]->type()->unifyNode();
auto& unnamedParameterType = downcast<AST::UnnamedType>(parameterType);
if (is<AST::ArrayType>(unnamedParameterType)) {
auto& arrayParameterType = downcast<AST::ArrayType>(unnamedParameterType);
stringBuilder.flexibleAppend(
- "uint ", outputFunctionName, '(', metalParameterName, ") {\n"
- " return ", arrayParameterType.numElements(), ";\n"
- "}\n"
- );
- return stringBuilder.toString();
+ returnName, " = ", arrayParameterType.numElements(), ";\n");
+ return;
}
ASSERT(is<AST::ArrayReferenceType>(unnamedParameterType));
stringBuilder.flexibleAppend(
- "uint ", outputFunctionName, '(', metalParameterName, " v) {\n"
- " return v.length;\n"
- "}\n"
- );
- return stringBuilder.toString();
+ returnName, " = ", args[0], ".length;\n");
+ return;
}
if (nativeFunctionDeclaration.name().startsWith("operator."_str)) {
- auto appendMangledFieldName = [&](StringBuilder& stringBuilder, const String& fieldName) {
+ auto appendMangledFieldName = [&] (const String& fieldName) {
auto& unifyNode = nativeFunctionDeclaration.parameters()[0]->type()->unifyNode();
auto& namedType = downcast<AST::NamedType>(unifyNode);
if (is<AST::StructureDefinition>(namedType)) {
@@ -221,53 +195,34 @@
if (nativeFunctionDeclaration.name().endsWith("=")) {
ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
- auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
- auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
- auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
- stringBuilder.flexibleAppend(
- metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " v, ", metalParameter2Name, " n) {\n"
- " v."
- );
-
auto fieldName = nativeFunctionDeclaration.name().substring("operator."_str.length());
fieldName = fieldName.substring(0, fieldName.length() - 1);
- appendMangledFieldName(stringBuilder, fieldName);
- stringBuilder.append(" = n;\n"
- " return v;\n"
- "}\n"
- );
- return stringBuilder.toString();
+ stringBuilder.flexibleAppend(
+ returnName, " = ", args[0], ";\n",
+ returnName, '.');
+ appendMangledFieldName(fieldName);
+ stringBuilder.flexibleAppend(" = ", args[1], ";\n");
+
+ return;
}
ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
- auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
- auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
+ auto fieldName = nativeFunctionDeclaration.name().substring("operator."_str.length());
stringBuilder.flexibleAppend(
- metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " v) {\n"
- " return v."
- );
-
- auto fieldName = nativeFunctionDeclaration.name().substring("operator."_str.length());
- appendMangledFieldName(stringBuilder, fieldName);
-
- stringBuilder.append(";\n"
- "}\n"
- );
- return stringBuilder.toString();
+ returnName, " = ", args[0], '.');
+ appendMangledFieldName(fieldName);
+ stringBuilder.append(";\n");
+ return;
}
if (nativeFunctionDeclaration.name().startsWith("operator&."_str)) {
ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
- auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
- auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
+ auto fieldName = nativeFunctionDeclaration.name().substring("operator&."_str.length());
stringBuilder.flexibleAppend(
- metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " v) {\n"
- " return &(v->"
- );
+ returnName, " = &(", args[0], "->");
- auto fieldName = nativeFunctionDeclaration.name().substring("operator&."_str.length());
auto& unnamedType = *nativeFunctionDeclaration.parameters()[0]->type();
auto& unifyNode = downcast<AST::PointerType>(unnamedType).elementType().unifyNode();
auto& namedType = downcast<AST::NamedType>(unifyNode);
@@ -279,26 +234,19 @@
} else
stringBuilder.flexibleAppend(fieldName);
- stringBuilder.append(
- ");\n"
- "}\n"
- );
- return stringBuilder.toString();
+ stringBuilder.append(");\n");
+
+ return;
}
if (nativeFunctionDeclaration.name() == "operator&[]") {
ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
- auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
- auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
- auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
ASSERT(is<AST::ArrayReferenceType>(*nativeFunctionDeclaration.parameters()[0]->type()));
+
stringBuilder.flexibleAppend(
- metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " v, ", metalParameter2Name, " n) {\n"
- " if (n < v.length) return &(v.pointer[n]);\n"
- " return nullptr;\n"
- "}\n"
- );
- return stringBuilder.toString();
+ returnName, " = (", args[1], " < ", args[0], ".length) ? ", " &(", args[0], ".pointer[", args[1], "]) : nullptr;\n");
+
+ return;
}
auto matrixDimension = [&] (unsigned typeArgumentIndex) -> unsigned {
@@ -317,82 +265,68 @@
if (nativeFunctionDeclaration.name() == "operator[]") {
ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
- auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
- auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
unsigned numberOfRows = numberOfMatrixRows();
unsigned numberOfColumns = numberOfMatrixColumns();
+ stringBuilder.flexibleAppend("do {\n", metalReturnName, " result;\n");
stringBuilder.flexibleAppend(
- metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " m, ", metalParameter2Name, " i) {\n"
- " if (i >= ", numberOfRows, ") return ", metalReturnName, "(0);\n"
- " ", metalReturnName, " result;\n"
- " result[0] = m[i];\n"
- " result[1] = m[i + ", numberOfRows, "];\n"
- );
+ " if (", args[1], " >= ", numberOfRows, ") {", returnName, " = ", metalReturnName, "(0); break;}\n",
+ " result[0] = ", args[0], '[', args[1], "];\n",
+ " result[1] = ", args[0], '[', args[1], " + ", numberOfRows, "];\n");
+
if (numberOfColumns >= 3)
- stringBuilder.flexibleAppend(" result[2] = m[i + ", numberOfRows * 2, "];\n");
+ stringBuilder.flexibleAppend(" result[2] = ", args[0], '[', args[1], " + ", numberOfRows * 2, "];\n");
if (numberOfColumns >= 4)
- stringBuilder.flexibleAppend(" result[3] = m[i + ", numberOfRows * 3, "];\n");
- stringBuilder.append(
- " return result;\n"
- "}\n"
- );
- return stringBuilder.toString();
+ stringBuilder.flexibleAppend(" result[3] = ", args[0], '[', args[1], " + ", numberOfRows * 3, "];\n");
+
+ stringBuilder.flexibleAppend(
+ " ", returnName, " = result;\n",
+ "} while (0);\n");
+
+ return;
}
if (nativeFunctionDeclaration.name() == "operator[]=") {
ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
- auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
- auto metalParameter3Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[2]->type());
auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
unsigned numberOfRows = numberOfMatrixRows();
unsigned numberOfColumns = numberOfMatrixColumns();
+ stringBuilder.flexibleAppend("do {\n", metalReturnName, " m = ", args[0], ";\n",
+ metalParameter2Name, " i = ", args[1], ";\n");
+
stringBuilder.flexibleAppend(
- metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " m, ", metalParameter2Name, " i, ", metalParameter3Name, " v) {\n"
- " if (i >= ", numberOfRows, ") return m;\n"
- " m[i] = v[0];\n"
- " m[i + ", numberOfRows, "] = v[1];\n"
- );
+ " if (i >= ", numberOfRows, ") {", returnName, " = m;\nbreak;}\n",
+ " m[i] = ", args[2], "[0];\n",
+ " m[i + ", numberOfRows, "] = ", args[2], "[1];\n");
if (numberOfColumns >= 3)
- stringBuilder.flexibleAppend(" m[i + ", numberOfRows * 2, "] = v[2];\n");
+ stringBuilder.flexibleAppend(" m[i + ", numberOfRows * 2, "] = ", args[2], "[2];\n");
if (numberOfColumns >= 4)
- stringBuilder.flexibleAppend(" m[i + ", numberOfRows * 3, "] = v[3];\n");
- stringBuilder.append(
- " return m;"
- "}\n"
- );
- return stringBuilder.toString();
+ stringBuilder.flexibleAppend(" m[i + ", numberOfRows * 3, "] = ", args[2], "[3];\n");
+ stringBuilder.flexibleAppend(
+ " ", returnName, " = m;\n",
+ "} while(0);\n");
+ return;
}
if (nativeFunctionDeclaration.isOperator()) {
+ auto operatorName = nativeFunctionDeclaration.name().substring("operator"_str.length());
if (nativeFunctionDeclaration.parameters().size() == 1) {
- auto operatorName = nativeFunctionDeclaration.name().substring("operator"_str.length());
auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
- auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
stringBuilder.flexibleAppend(
- metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"
- " return ", operatorName, "x;\n"
- "}\n"
- );
- return stringBuilder.toString();
+ "{\n", metalParameterName, " x = ", args[0], ";\n",
+ returnName, " = ", operatorName, "x;\n}\n");
+ return;
}
ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
- auto operatorName = nativeFunctionDeclaration.name().substring("operator"_str.length());
- auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
- auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
- auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
stringBuilder.flexibleAppend(
- metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " x, ", metalParameter2Name, " y) {\n"
- " return x ", operatorName, " y;\n"
- "}\n"
- );
- return stringBuilder.toString();
+ returnName, " = ", args[0], ' ', operatorName, ' ', args[1], ";\n");
+ return;
}
if (nativeFunctionDeclaration.name() == "cos"
@@ -420,96 +354,55 @@
|| nativeFunctionDeclaration.name() == "asuint"
|| nativeFunctionDeclaration.name() == "asfloat") {
ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
- auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
- auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
stringBuilder.flexibleAppend(
- metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"
- " return ", mapFunctionName(nativeFunctionDeclaration.name()), "(x);\n"
- "}\n"
- );
- return stringBuilder.toString();
+ returnName, " = ", mapFunctionName(nativeFunctionDeclaration.name()), '(', args[0], ");\n");
+ return;
}
if (nativeFunctionDeclaration.name() == "pow" || nativeFunctionDeclaration.name() == "atan2") {
ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
- auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
- auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
- auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
stringBuilder.flexibleAppend(
- metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " x, ", metalParameter2Name, " y) {\n"
- " return ", nativeFunctionDeclaration.name(), "(x, y);\n"
- "}\n"
- );
- return stringBuilder.toString();
+ returnName, " = ", nativeFunctionDeclaration.name(), "(", args[0], ", ", args[1], ");\n");
+ return;
}
if (nativeFunctionDeclaration.name() == "AllMemoryBarrierWithGroupSync") {
ASSERT(!nativeFunctionDeclaration.parameters().size());
- stringBuilder.flexibleAppend(
- "void ", outputFunctionName, "() {\n"
- " threadgroup_barrier(mem_flags::mem_device);\n"
- " threadgroup_barrier(mem_flags::mem_threadgroup);\n"
- " threadgroup_barrier(mem_flags::mem_texture);\n"
- "}\n"
- );
- return stringBuilder.toString();
+ stringBuilder.append(
+ "threadgroup_barrier(mem_flags::mem_device);\n"
+ "threadgroup_barrier(mem_flags::mem_threadgroup);\n"
+ "threadgroup_barrier(mem_flags::mem_texture);\n");
+ return;
}
if (nativeFunctionDeclaration.name() == "DeviceMemoryBarrierWithGroupSync") {
ASSERT(!nativeFunctionDeclaration.parameters().size());
- stringBuilder.flexibleAppend(
- "void ", outputFunctionName, "() {\n"
- " threadgroup_barrier(mem_flags::mem_device);\n"
- "}\n"
- );
- return stringBuilder.toString();
+ stringBuilder.append(
+ "threadgroup_barrier(mem_flags::mem_device);\n");
+ return;
}
if (nativeFunctionDeclaration.name() == "GroupMemoryBarrierWithGroupSync") {
ASSERT(!nativeFunctionDeclaration.parameters().size());
- stringBuilder.flexibleAppend(
- "void ", outputFunctionName, "() {\n"
- " threadgroup_barrier(mem_flags::mem_threadgroup);\n"
- "}\n"
- );
- return stringBuilder.toString();
+ stringBuilder.append(
+ "threadgroup_barrier(mem_flags::mem_threadgroup);\n");
+ return;
}
if (nativeFunctionDeclaration.name().startsWith("Interlocked"_str)) {
if (nativeFunctionDeclaration.name() == "InterlockedCompareExchange") {
ASSERT(nativeFunctionDeclaration.parameters().size() == 4);
- auto& firstArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0]->type());
- auto firstArgumentAddressSpace = firstArgumentPointer.addressSpace();
- auto firstArgumentPointee = typeNamer.mangledNameForType(firstArgumentPointer.elementType());
- auto secondArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
- auto thirdArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[2]->type());
- auto& fourthArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[3]->type());
- auto fourthArgumentAddressSpace = fourthArgumentPointer.addressSpace();
- auto fourthArgumentPointee = typeNamer.mangledNameForType(fourthArgumentPointer.elementType());
stringBuilder.flexibleAppend(
- "void ", outputFunctionName, '(', toString(firstArgumentAddressSpace), ' ', firstArgumentPointee, "* object, ", secondArgument, " compare, ", thirdArgument, " desired, ", toString(fourthArgumentAddressSpace), ' ', fourthArgumentPointee, "* out) {\n"
- " atomic_compare_exchange_weak_explicit(object, &compare, desired, memory_order_relaxed, memory_order_relaxed);\n"
- " *out = compare;\n"
- "}\n"
- );
- return stringBuilder.toString();
+ "atomic_compare_exchange_weak_explicit(", args[0], ", &", args[1], ", ", args[2], ", memory_order_relaxed, memory_order_relaxed);\n",
+ '*', args[3], " = ", args[1], ";\n");
+ return;
}
ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
- auto& firstArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0]->type());
- auto firstArgumentAddressSpace = firstArgumentPointer.addressSpace();
- auto firstArgumentPointee = typeNamer.mangledNameForType(firstArgumentPointer.elementType());
- auto secondArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
- auto& thirdArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[2]->type());
- auto thirdArgumentAddressSpace = thirdArgumentPointer.addressSpace();
- auto thirdArgumentPointee = typeNamer.mangledNameForType(thirdArgumentPointer.elementType());
auto name = atomicName(nativeFunctionDeclaration.name().substring("Interlocked"_str.length()));
stringBuilder.flexibleAppend(
- "void ", outputFunctionName, '(', toString(firstArgumentAddressSpace), ' ', firstArgumentPointee, "* object, ", secondArgument, " operand, ", toString(thirdArgumentAddressSpace), ' ', thirdArgumentPointee, "* out) {\n"
- " *out = atomic_", name, "_explicit(object, operand, memory_order_relaxed);\n"
- "}\n"
- );
- return stringBuilder.toString();
+ '*', args[2], " = atomic_", name, "_explicit(", args[0], ", ", args[1], ", memory_order_relaxed);\n");
+ return;
}
if (nativeFunctionDeclaration.name() == "Sample") {
@@ -521,35 +414,22 @@
auto& returnType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.type().unifyNode()));
auto returnVectorLength = vectorLength(returnType);
- auto metalParameter1Name = typeNamer.mangledNameForType(textureType);
- auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
- auto metalParameter3Name = typeNamer.mangledNameForType(locationType);
- Optional<MangledTypeName> metalParameter4Name;
- if (nativeFunctionDeclaration.parameters().size() == 4)
- metalParameter4Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[3]->type());
- auto metalReturnName = typeNamer.mangledNameForType(returnType);
- stringBuilder.flexibleAppend(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " theTexture, ", metalParameter2Name, " theSampler, ", metalParameter3Name, " location");
- if (metalParameter4Name)
- stringBuilder.flexibleAppend(", ", *metalParameter4Name, " offset");
- stringBuilder.append(
- ") {\n"
- " return theTexture.sample(theSampler, "
- );
+ stringBuilder.flexibleAppend(
+ returnName, " = ", args[0], ".sample(", args[1], ", ");
+
if (textureType.isTextureArray()) {
ASSERT(locationVectorLength > 1);
- stringBuilder.flexibleAppend("location.", "xyzw"_str.substring(0, locationVectorLength - 1), ", location.", "xyzw"_str.substring(locationVectorLength - 1, 1));
+ stringBuilder.flexibleAppend(args[2], '.', "xyzw"_str.substring(0, locationVectorLength - 1), ", ", args[2], '.', "xyzw"_str.substring(locationVectorLength - 1, 1));
} else
- stringBuilder.append("location");
- if (metalParameter4Name)
- stringBuilder.append(", offset");
+ stringBuilder.flexibleAppend(args[2]);
+ if (nativeFunctionDeclaration.parameters().size() == 4)
+ stringBuilder.flexibleAppend(", ", args[3]);
stringBuilder.append(")");
if (!textureType.isDepthTexture())
stringBuilder.flexibleAppend(".", "xyzw"_str.substring(0, returnVectorLength));
- stringBuilder.append(
- ";\n"
- "}\n"
- );
- return stringBuilder.toString();
+ stringBuilder.append(";\n");
+
+ return;
}
if (nativeFunctionDeclaration.name() == "Load") {
@@ -561,67 +441,58 @@
auto& returnType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.type().unifyNode()));
auto returnVectorLength = vectorLength(returnType);
- auto metalParameter1Name = typeNamer.mangledNameForType(textureType);
- auto metalParameter2Name = typeNamer.mangledNameForType(locationType);
auto metalReturnName = typeNamer.mangledNameForType(returnType);
- stringBuilder.flexibleAppend(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " theTexture, ", metalParameter2Name, " location) {\n");
+ stringBuilder.append("do {\n");
+
if (textureType.isTextureArray()) {
ASSERT(locationVectorLength > 1);
String dimensions[] = { "width"_str, "height"_str, "depth"_str };
for (int i = 0; i < locationVectorLength - 1; ++i) {
auto suffix = "xyzw"_str.substring(i, 1);
- stringBuilder.flexibleAppend(" if (location.", suffix, " < 0 || static_cast<uint32_t>(location.", suffix, ") >= theTexture.get_", dimensions[i], "()) return ", metalReturnName, "(0);\n");
+ stringBuilder.flexibleAppend(" if (", args[1], '.', suffix, " < 0 || static_cast<uint32_t>(", args[1], '.', suffix, ") >= ", args[0], ".get_", dimensions[i], "()) {", returnName, " = ", metalReturnName, "(0); break;}\n");
}
auto suffix = "xyzw"_str.substring(locationVectorLength - 1, 1);
- stringBuilder.flexibleAppend(" if (location.", suffix, " < 0 || static_cast<uint32_t>(location.", suffix, ") >= theTexture.get_array_size()) return ", metalReturnName, "(0);\n");
+ stringBuilder.flexibleAppend(" if (", args[1], '.', suffix, " < 0 || static_cast<uint32_t>(", args[1], '.', suffix, ") >= ", args[0], ".get_array_size()) {", returnName, " = ", metalReturnName, "(0); break;}\n");
} else {
if (locationVectorLength == 1)
- stringBuilder.flexibleAppend(" if (location < 0 || static_cast<uint32_t>(location) >= theTexture.get_width()) return ", metalReturnName, "(0);\n");
+ stringBuilder.flexibleAppend(" if (", args[1], " < 0 || static_cast<uint32_t>(", args[1], ") >= ", args[0], ".get_width()) { ", returnName, " = ", metalReturnName, "(0); break;}\n");
else {
stringBuilder.flexibleAppend(
- " if (location.x < 0 || static_cast<uint32_t>(location.x) >= theTexture.get_width()) return ", metalReturnName, "(0);\n"
- " if (location.y < 0 || static_cast<uint32_t>(location.y) >= theTexture.get_height()) return ", metalReturnName, "(0);\n"
- );
+ " if (", args[1], ".x < 0 || static_cast<uint32_t>(", args[1], ".x) >= ", args[0], ".get_width()) {", returnName, " = ", metalReturnName, "(0); break;}\n"
+ " if (", args[1], ".y < 0 || static_cast<uint32_t>(", args[1], ".y) >= ", args[0], ".get_height()) {", returnName, " = ", metalReturnName, "(0); break;}\n");
+
if (locationVectorLength >= 3)
- stringBuilder.flexibleAppend(" if (location.z < 0 || static_cast<uint32_t>(location.z) >= theTexture.get_depth()) return ", metalReturnName, "(0);\n");
+ stringBuilder.flexibleAppend(" if (", args[1], ".z < 0 || static_cast<uint32_t>(", args[1], ".z) >= ", args[0], ".get_depth()) {", returnName, " = ", metalReturnName, "(0); break;}\n");
}
}
- stringBuilder.append(" return theTexture.read(");
+ stringBuilder.flexibleAppend(" ", returnName, " = ", args[0], ".read(");
if (textureType.isTextureArray()) {
ASSERT(locationVectorLength > 1);
- stringBuilder.flexibleAppend("uint", vectorSuffix(locationVectorLength - 1), "(location.", "xyzw"_str.substring(0, locationVectorLength - 1), "), uint(location.", "xyzw"_str.substring(locationVectorLength - 1, 1), ')');
+ stringBuilder.flexibleAppend("uint", vectorSuffix(locationVectorLength - 1), '(', args[1], '.', "xyzw"_str.substring(0, locationVectorLength - 1), "), uint(", args[1], '.', "xyzw"_str.substring(locationVectorLength - 1, 1), ')');
} else
- stringBuilder.flexibleAppend("uint", vectorSuffix(locationVectorLength), "(location)");
+ stringBuilder.flexibleAppend("uint", vectorSuffix(locationVectorLength), '(', args[1], ')');
stringBuilder.append(')');
if (!textureType.isDepthTexture())
stringBuilder.flexibleAppend('.', "xyzw"_str.substring(0, returnVectorLength));
stringBuilder.append(
";\n"
- "}\n"
- );
- return stringBuilder.toString();
+ "} while(0);\n");
+
+ return;
}
if (nativeFunctionDeclaration.name() == "load") {
ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
- auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
- auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
stringBuilder.flexibleAppend(
- metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"
- " return atomic_load_explicit(x, memory_order_relaxed);\n"
- "}\n"
- );
- return stringBuilder.toString();
+ returnName, " = atomic_load_explicit(", args[0], ", memory_order_relaxed);\n");
+ return;
}
if (nativeFunctionDeclaration.name() == "store") {
ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
- auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
- auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
- stringBuilder.flexibleAppend("void ", outputFunctionName, '(', metalParameter1Name, " x, ", metalParameter2Name, " y) {\n"
- " atomic_store_explicit(x, y, memory_order_relaxed);\n"
- "}\n");
- return stringBuilder.toString();
+ stringBuilder.flexibleAppend(
+ "atomic_store_explicit(", args[0], ", ", args[1], ", memory_order_relaxed);\n");
+ return;
}
if (nativeFunctionDeclaration.name() == "GetDimensions") {
@@ -628,85 +499,67 @@
auto& textureType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[0]->type()->unifyNode()));
size_t index = 1;
- if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
+ bool hasMipLevel = !textureType.isWritableTexture() && textureType.textureDimension() != 1;
+ if (hasMipLevel)
++index;
- auto widthTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
+ const MangledVariableName& widthName = args[index];
++index;
- Optional<MangledTypeName> heightTypeName;
+ Optional<MangledVariableName> heightName;
if (textureType.textureDimension() >= 2) {
- heightTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
+ heightName = args[index];
++index;
}
- Optional<MangledTypeName> depthTypeName;
+ Optional<MangledVariableName> depthName;
if (textureType.textureDimension() >= 3) {
- depthTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
+ depthName = args[index];
++index;
}
- Optional<MangledTypeName> elementsTypeName;
+ Optional<MangledVariableName> elementsName;
if (textureType.isTextureArray()) {
- elementsTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
+ elementsName = args[index];
++index;
}
- Optional<MangledTypeName> numberOfLevelsTypeName;
+ Optional<MangledVariableName> numberOfLevelsName;
if (!textureType.isWritableTexture() && textureType.textureDimension() != 1) {
- numberOfLevelsTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
+ numberOfLevelsName = args[index];
++index;
}
ASSERT(index == nativeFunctionDeclaration.parameters().size());
- auto metalParameter1Name = typeNamer.mangledNameForType(textureType);
- stringBuilder.flexibleAppend("void ", outputFunctionName, '(', metalParameter1Name, " theTexture");
- if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
- stringBuilder.append(", uint mipLevel");
- stringBuilder.flexibleAppend(", ", widthTypeName, " width");
- if (heightTypeName)
- stringBuilder.flexibleAppend(", ", *heightTypeName, " height");
- if (depthTypeName)
- stringBuilder.flexibleAppend(", ", *depthTypeName, " depth");
- if (elementsTypeName)
- stringBuilder.flexibleAppend(", ", *elementsTypeName, " elements");
- if (numberOfLevelsTypeName)
- stringBuilder.flexibleAppend(", ", *numberOfLevelsTypeName, " numberOfLevels");
- stringBuilder.append(
- ") {\n"
- " if (width)\n"
- " *width = theTexture.get_width("
- );
- if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
- stringBuilder.append("mipLevel");
+ stringBuilder.flexibleAppend(
+ "if (", widthName, ")\n"
+ " *", widthName, " = ", args[0], ".get_width(");
+
+ if (hasMipLevel)
+ stringBuilder.flexibleAppend(args[1]);
stringBuilder.append(");\n");
- if (heightTypeName) {
- stringBuilder.append(
- " if (height)\n"
- " *height = theTexture.get_height("
- );
- if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
- stringBuilder.append("mipLevel");
+ if (heightName) {
+ stringBuilder.flexibleAppend(
+ " if (", *heightName, ")\n"
+ " *", *heightName, " = ", args[0], ".get_height(");
+ if (hasMipLevel)
+ stringBuilder.flexibleAppend(args[1]);
stringBuilder.append(");\n");
}
- if (depthTypeName) {
- stringBuilder.append(
- " if (depth)\n"
- " *depth = theTexture.get_depth("
- );
- if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
- stringBuilder.append("mipLevel");
+ if (depthName) {
+ stringBuilder.flexibleAppend(
+ " if (", *depthName, ")\n"
+ " *", *depthName, " = ", args[0], ".get_depth(");
+ if (hasMipLevel)
+ stringBuilder.flexibleAppend(args[1]);
stringBuilder.append(");\n");
}
- if (elementsTypeName) {
- stringBuilder.append(
- " if (elements)\n"
- " *elements = theTexture.get_array_size();\n"
- );
+ if (elementsName) {
+ stringBuilder.flexibleAppend(
+ " if (", *elementsName, ")\n"
+ " *", *elementsName, " = ", args[0], ".get_array_size();\n");
}
- if (numberOfLevelsTypeName) {
- stringBuilder.append(
- " if (numberOfLevels)\n"
- " *numberOfLevels = theTexture.get_num_mip_levels();\n"
- );
+ if (numberOfLevelsName) {
+ stringBuilder.flexibleAppend(
+ " if (", *numberOfLevelsName, ")\n"
+ " *", *numberOfLevelsName, " = ", args[0], ".get_num_mip_levels();\n");
}
- stringBuilder.append("}\n");
- return stringBuilder.toString();
+ return;
}
if (nativeFunctionDeclaration.name() == "SampleBias") {
@@ -754,46 +607,43 @@
auto& locationType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(nativeFunctionDeclaration.parameters()[2]->type()->unifyNode()));
auto locationVectorLength = vectorLength(locationType);
- auto metalParameter1Name = typeNamer.mangledNameForType(textureType);
- auto metalParameter2Name = typeNamer.mangledNameForType(itemType);
- auto metalParameter3Name = typeNamer.mangledNameForType(locationType);
auto metalInnerTypeName = typeNamer.mangledNameForType(itemVectorInnerType);
- stringBuilder.flexibleAppend("void ", outputFunctionName, '(', metalParameter1Name, " theTexture, ", metalParameter2Name, " item, ", metalParameter3Name, " location) {\n");
+
+ stringBuilder.append("do {\n");
if (textureType.isTextureArray()) {
ASSERT(locationVectorLength > 1);
String dimensions[] = { "width"_str, "height"_str, "depth"_str };
for (int i = 0; i < locationVectorLength - 1; ++i) {
auto suffix = "xyzw"_str.substring(i, 1);
- stringBuilder.flexibleAppend(" if (location.", suffix, " >= theTexture.get_", dimensions[i], "()) return;\n");
+ stringBuilder.flexibleAppend(" if (", args[2], ".", suffix, " >= ", args[0], ".get_", dimensions[i], "()) break;\n");
}
auto suffix = "xyzw"_str.substring(locationVectorLength - 1, 1);
- stringBuilder.flexibleAppend(" if (location.", suffix, " >= theTexture.get_array_size()) return;\n");
+ stringBuilder.flexibleAppend(" if (", args[2], '.', suffix, " >= ", args[0], ".get_array_size()) break;\n");
} else {
if (locationVectorLength == 1)
- stringBuilder.append(" if (location >= theTexture.get_width()) return;\n");
+ stringBuilder.flexibleAppend(" if (", args[2], " >= ", args[0], ".get_width()) break;\n");
else {
- stringBuilder.append(
- " if (location.x >= theTexture.get_width()) return;\n"
- " if (location.y >= theTexture.get_height()) return;\n"
- );
+ stringBuilder.flexibleAppend(
+ " if (", args[2], ".x >= ", args[0], ".get_width()) break;\n"
+ " if (", args[2], ".y >= ", args[0], ".get_height()) break;\n");
if (locationVectorLength >= 3)
- stringBuilder.append(" if (location.z >= theTexture.get_depth()) return;\n");
+ stringBuilder.flexibleAppend(" if (", args[2], ".z >= ", args[0], ".get_depth()) break;\n");
}
}
- stringBuilder.flexibleAppend(" theTexture.write(vec<", metalInnerTypeName, ", 4>(item");
+ stringBuilder.flexibleAppend(" ", args[0], ".write(vec<", metalInnerTypeName, ", 4>(", args[1]);
for (int i = 0; i < 4 - itemVectorLength; ++i)
stringBuilder.append(", 0");
stringBuilder.append("), ");
if (textureType.isTextureArray()) {
ASSERT(locationVectorLength > 1);
- stringBuilder.flexibleAppend("uint", vectorSuffix(locationVectorLength - 1), "(location.", "xyzw"_str.substring(0, locationVectorLength - 1), "), uint(location.", "xyzw"_str.substring(locationVectorLength - 1, 1), ')');
+ stringBuilder.flexibleAppend("uint", vectorSuffix(locationVectorLength - 1), '(', args[2], '.', "xyzw"_str.substring(0, locationVectorLength - 1), "), uint(", args[2], ".", "xyzw"_str.substring(locationVectorLength - 1, 1), ')');
} else
- stringBuilder.flexibleAppend("uint", vectorSuffix(locationVectorLength), "(location)");
+ stringBuilder.flexibleAppend("uint", vectorSuffix(locationVectorLength), '(', args[2], ')');
stringBuilder.append(
");\n"
- "}\n"
- );
- return stringBuilder.toString();
+ "} while(0);\n");
+
+ return;
}
if (nativeFunctionDeclaration.name() == "GatherAlpha") {
@@ -822,7 +672,6 @@
}
ASSERT_NOT_REACHED();
- return String();
}
} // namespace Metal
