⚠ Archived content — this site is no longer maintained.   Current WebKit documentation is at docs.webkit.org.

Changeset 248280 in webkit


Ignore:
Timestamp:
Aug 5, 2019, 5:03:34 PM (7 years ago)
Author:
sbarati@apple.com
Message:

[WHLSL] Inline all native function calls
​https://bugs.webkit.org/show_bug.cgi?id=200350

Reviewed by Robin Morisset.

Native functions calls tend to be really small. If we inline in the generated
Metal code, we end up with faster Metal compile times. On compute_boids, this
provides a ~10ms improvement.

  • Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp:

(WebCore::WHLSL::Metal::FunctionDefinitionWriter::visit):
(WebCore::WHLSL::Metal::sharedMetalFunctions):

  • Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.cpp:

(WebCore::WHLSL::Metal::inlineNativeFunction):
(WebCore::WHLSL::Metal::writeNativeFunction): Deleted.

  • Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.h:
Location:
trunk/Source/WebCore
Files:
4 edited

Legend:

Unmodified
Added
Removed

  • trunk/Source/WebCore/ChangeLog

    r248278 r248280  
     12019-08-05  Saam Barati  <sbarati@apple.com>
     2
     3        [WHLSL] Inline all native function calls
     4        https://bugs.webkit.org/show_bug.cgi?id=200350
     5
     6        Reviewed by Robin Morisset.
     7
     8        Native functions calls tend to be really small. If we inline in the generated
     9        Metal code, we end up with faster Metal compile times. On compute_boids, this
     10        provides a ~10ms improvement.
     11
     12        * Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp:
     13        (WebCore::WHLSL::Metal::FunctionDefinitionWriter::visit):
     14        (WebCore::WHLSL::Metal::sharedMetalFunctions):
     15        * Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.cpp:
     16        (WebCore::WHLSL::Metal::inlineNativeFunction):
     17        (WebCore::WHLSL::Metal::writeNativeFunction): Deleted.
     18        * Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.h:
     19
    1202019-08-05  Youenn Fablet  <youenn@apple.com>
    221

  • trunk/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp

    r248266 r248280  
    229229};
    230230
    231 void FunctionDefinitionWriter::visit(AST::NativeFunctionDeclaration& nativeFunctionDeclaration)
    232 {
    233     auto iterator = m_functionMapping.find(&nativeFunctionDeclaration);
    234     ASSERT(iterator != m_functionMapping.end());
    235     m_stringBuilder.append(writeNativeFunction(nativeFunctionDeclaration, iterator->value, m_intrinsics, m_typeNamer));
     231void FunctionDefinitionWriter::visit(AST::NativeFunctionDeclaration&)
     232{
     233    // We inline native function calls.
    236234}
    237235
    … …  
    578576        argumentNames.append(takeLastValue());
    579577    }
    580     auto iterator = m_functionMapping.find(&callExpression.function());
    581     ASSERT(iterator != m_functionMapping.end());
    582     auto variableName = generateNextVariableName();
    583     if (!matches(callExpression.resolvedType(), m_intrinsics.voidType()))
    584         m_stringBuilder.flexibleAppend(m_typeNamer.mangledNameForType(callExpression.resolvedType()), ' ', variableName, " = ");
    585     m_stringBuilder.flexibleAppend(iterator->value, '(');
    586     for (size_t i = 0; i < argumentNames.size(); ++i) {
    587         if (i)
    588             m_stringBuilder.append(", ");
    589         m_stringBuilder.flexibleAppend(argumentNames[i]);
    590     }
    591     m_stringBuilder.append(");\n");
    592     appendRightValue(callExpression, variableName);
     578
     579    bool isVoid = matches(callExpression.resolvedType(), m_intrinsics.voidType());
     580    MangledVariableName returnName;
     581    if (!isVoid) {
     582        returnName = generateNextVariableName();
     583        m_stringBuilder.flexibleAppend(m_typeNamer.mangledNameForType(callExpression.resolvedType()), ' ', returnName, ";\n");
     584    }
     585
     586    if (is<AST::NativeFunctionDeclaration>(callExpression.function()))
     587        inlineNativeFunction(m_stringBuilder, downcast<AST::NativeFunctionDeclaration>(callExpression.function()), returnName, argumentNames, m_intrinsics, m_typeNamer);
     588    else {
     589        auto iterator = m_functionMapping.find(&callExpression.function());
     590        ASSERT(iterator != m_functionMapping.end());
     591        if (!isVoid)
     592            m_stringBuilder.flexibleAppend(returnName, " = ");
     593        m_stringBuilder.flexibleAppend(iterator->value, '(');
     594        for (size_t i = 0; i < argumentNames.size(); ++i) {
     595            if (i)
     596                m_stringBuilder.append(", ");
     597            m_stringBuilder.flexibleAppend(argumentNames[i]);
     598        }
     599        m_stringBuilder.append(");\n");
     600    }
     601
     602    appendRightValue(callExpression, returnName);
    593603}
    594604
    … …  
    805815    unsigned numFunctions = 0;
    806816    HashMap<AST::FunctionDeclaration*, MangledFunctionName> functionMapping;
    807     for (auto& nativeFunctionDeclaration : program.nativeFunctionDeclarations()) {
    808         auto addResult = functionMapping.add(&nativeFunctionDeclaration, MangledFunctionName { numFunctions++ });
    809         ASSERT_UNUSED(addResult, addResult.isNewEntry);
    810     }
    811817    for (auto& functionDefinition : program.functionDefinitions()) {
    812818        auto addResult = functionMapping.add(&functionDefinition, MangledFunctionName { numFunctions++ });
    … …  
    816822    {
    817823        FunctionDeclarationWriter functionDeclarationWriter(typeNamer, functionMapping);
    818         for (auto& nativeFunctionDeclaration : program.nativeFunctionDeclarations()) {
    819             if (reachableFunctions.contains(&nativeFunctionDeclaration))
    820                 functionDeclarationWriter.visit(nativeFunctionDeclaration);
    821         }
    822824        for (auto& functionDefinition : program.functionDefinitions()) {
    823825            if (!functionDefinition->entryPointType() && reachableFunctions.contains(&functionDefinition))
    … …  
    870872
    871873    RenderFunctionDefinitionWriter functionDefinitionWriter(program.intrinsics(), typeNamer, sharedMetalFunctions.functionMapping, WTFMove(matchedSemantics), layout);
    872     for (auto& nativeFunctionDeclaration : program.nativeFunctionDeclarations()) {
    873         if (reachableFunctions.contains(&nativeFunctionDeclaration))
    874             functionDefinitionWriter.visit(nativeFunctionDeclaration);
    875     }
    876874    for (auto& functionDefinition : program.functionDefinitions()) {
    877875        if (reachableFunctions.contains(&functionDefinition))
    … …  
    901899
    902900    ComputeFunctionDefinitionWriter functionDefinitionWriter(program.intrinsics(), typeNamer, sharedMetalFunctions.functionMapping, WTFMove(matchedSemantics), layout);
    903     for (auto& nativeFunctionDeclaration : program.nativeFunctionDeclarations()) {
    904         if (reachableFunctions.contains(&nativeFunctionDeclaration))
    905             functionDefinitionWriter.visit(nativeFunctionDeclaration);
    906     }
    907901    for (auto& functionDefinition : program.functionDefinitions()) {
    908902        if (reachableFunctions.contains(&functionDefinition))

  • trunk/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.cpp

    r248266 r248280  
    121121}
    122122
    123 String writeNativeFunction(AST::NativeFunctionDeclaration& nativeFunctionDeclaration, MangledFunctionName outputFunctionName, Intrinsics& intrinsics, TypeNamer& typeNamer)
     123void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDeclaration& nativeFunctionDeclaration, MangledVariableName returnName, const Vector<MangledVariableName>& args, Intrinsics& intrinsics, TypeNamer& typeNamer)
    124124{
    125     StringBuilder stringBuilder;
    126125    if (nativeFunctionDeclaration.isCast()) {
    127126        auto& returnType = nativeFunctionDeclaration.type();
    128127        auto metalReturnName = typeNamer.mangledNameForType(returnType);
    129128        if (!nativeFunctionDeclaration.parameters().size()) {
    130             stringBuilder.flexibleAppend(
    131                 metalReturnName, ' ', outputFunctionName, "() {\n"
    132                 "    ", metalReturnName, " x = { };\n"
    133                 "    return x;\n"
    134                 "}\n"
    135             );
    136             return stringBuilder.toString();
     129            stringBuilder.flexibleAppend(returnName, " = { };\n");
     130            return;
    137131        }
    138132
    … …  
    140134        auto& parameterType = *nativeFunctionDeclaration.parameters()[0]->type();
    141135        auto metalParameterName = typeNamer.mangledNameForType(parameterType);
    142         stringBuilder.flexibleAppend(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n");
     136        stringBuilder.flexibleAppend("{\n", metalParameterName, " x = ", args[0], ";\n");
    143137
    144138        {
    … …  
    161155
    162156        stringBuilder.flexibleAppend(
    163             "    return static_cast<", metalReturnName, ">(x);\n"
    164             "}\n"
    165         );
    166         return stringBuilder.toString();
    167     }
    168 
    169     if (nativeFunctionDeclaration.name() == "operator.value") {
    170         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
    171         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
    172         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
    173         stringBuilder.flexibleAppend(
    174             metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"
    175             "    return static_cast<", metalReturnName, ">(x);\n"
    176             "}\n"
    177         );
    178         return stringBuilder.toString();
     157            returnName, " = static_cast<", metalReturnName, ">(x);\n}\n");
     158
     159        return;
    179160    }
    180161
    … …  
    183164        ASSERT_UNUSED(intrinsics, matches(nativeFunctionDeclaration.type(), intrinsics.uintType()));
    184165        ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
    185         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
    186166        auto& parameterType = nativeFunctionDeclaration.parameters()[0]->type()->unifyNode();
    187167        auto& unnamedParameterType = downcast<AST::UnnamedType>(parameterType);
    … …  
    189169            auto& arrayParameterType = downcast<AST::ArrayType>(unnamedParameterType);
    190170            stringBuilder.flexibleAppend(
    191                 "uint ", outputFunctionName, '(', metalParameterName, ") {\n"
    192                 "    return ", arrayParameterType.numElements(), ";\n"
    193                 "}\n"
    194             );
    195             return stringBuilder.toString();
     171                returnName, " = ", arrayParameterType.numElements(), ";\n");
     172            return;
    196173        }
    197174
    198175        ASSERT(is<AST::ArrayReferenceType>(unnamedParameterType));
    199176        stringBuilder.flexibleAppend(
    200             "uint ", outputFunctionName, '(', metalParameterName, " v) {\n"
    201             "    return v.length;\n"
    202             "}\n"
    203         );
    204         return stringBuilder.toString();
     177            returnName, " = ", args[0], ".length;\n");
     178        return;
    205179    }
    206180
    207181    if (nativeFunctionDeclaration.name().startsWith("operator."_str)) {
    208         auto appendMangledFieldName = [&](StringBuilder& stringBuilder, const String& fieldName) {
     182        auto appendMangledFieldName = [&] (const String& fieldName) {
    209183            auto& unifyNode = nativeFunctionDeclaration.parameters()[0]->type()->unifyNode();
    210184            auto& namedType = downcast<AST::NamedType>(unifyNode);
    … …  
    222196        if (nativeFunctionDeclaration.name().endsWith("=")) {
    223197            ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
    224             auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
    225             auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
    226             auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
    227             stringBuilder.flexibleAppend(
    228                 metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " v, ", metalParameter2Name, " n) {\n"
    229                 "    v."
    230             );
    231 
    232198            auto fieldName = nativeFunctionDeclaration.name().substring("operator."_str.length());
    233199            fieldName = fieldName.substring(0, fieldName.length() - 1);
    234             appendMangledFieldName(stringBuilder, fieldName);
    235 
    236             stringBuilder.append(" = n;\n"
    237                 "    return v;\n"
    238                 "}\n"
    239             );
    240             return stringBuilder.toString();
     200
     201            stringBuilder.flexibleAppend(
     202                returnName, " = ", args[0], ";\n",
     203                returnName, '.');
     204            appendMangledFieldName(fieldName);
     205            stringBuilder.flexibleAppend(" = ", args[1], ";\n");
     206
     207            return;
    241208        }
    242209
    243210        ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
    244         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
    245         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
    246         stringBuilder.flexibleAppend(
    247             metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " v) {\n"
    248             "    return v."
    249         );
    250 
    251211        auto fieldName = nativeFunctionDeclaration.name().substring("operator."_str.length());
    252         appendMangledFieldName(stringBuilder, fieldName);
    253            
    254         stringBuilder.append(";\n"
    255             "}\n"
    256         );
    257         return stringBuilder.toString();
     212        stringBuilder.flexibleAppend(
     213            returnName, " = ", args[0], '.');
     214        appendMangledFieldName(fieldName);
     215        stringBuilder.append(";\n");
     216        return;
    258217    }
    259218
    260219    if (nativeFunctionDeclaration.name().startsWith("operator&."_str)) {
    261220        ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
    262         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
    263         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
    264 
    265         stringBuilder.flexibleAppend(
    266             metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " v) {\n"
    267             "    return &(v->"
    268         );
    269 
    270221        auto fieldName = nativeFunctionDeclaration.name().substring("operator&."_str.length());
     222
     223        stringBuilder.flexibleAppend(
     224            returnName, " = &(", args[0], "->");
     225
    271226        auto& unnamedType = *nativeFunctionDeclaration.parameters()[0]->type();
    272227        auto& unifyNode = downcast<AST::PointerType>(unnamedType).elementType().unifyNode();
    … …  
    280235            stringBuilder.flexibleAppend(fieldName);
    281236
    282         stringBuilder.append(
    283             ");\n"
    284             "}\n"
    285         );
    286         return stringBuilder.toString();
     237        stringBuilder.append(");\n");
     238
     239        return;
    287240    }
    288241
    289242    if (nativeFunctionDeclaration.name() == "operator&[]") {
    290243        ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
    291         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
    292         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
    293         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
    294244        ASSERT(is<AST::ArrayReferenceType>(*nativeFunctionDeclaration.parameters()[0]->type()));
    295         stringBuilder.flexibleAppend(
    296             metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " v, ", metalParameter2Name, " n) {\n"
    297             "    if (n < v.length) return &(v.pointer[n]);\n"
    298             "    return nullptr;\n"
    299             "}\n"
    300         );
    301         return stringBuilder.toString();
     245
     246        stringBuilder.flexibleAppend(
     247            returnName, " = (", args[1], " < ", args[0], ".length) ? ", " &(", args[0], ".pointer[", args[1], "]) : nullptr;\n");
     248           
     249        return;
    302250    }
    303251
    … …  
    318266    if (nativeFunctionDeclaration.name() == "operator[]") {
    319267        ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
    320         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
     268        auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
     269
     270        unsigned numberOfRows = numberOfMatrixRows();
     271        unsigned numberOfColumns = numberOfMatrixColumns();
     272        stringBuilder.flexibleAppend("do {\n", metalReturnName, " result;\n");
     273
     274        stringBuilder.flexibleAppend(
     275            "    if (", args[1], " >= ", numberOfRows, ") {", returnName, " = ", metalReturnName, "(0); break;}\n",
     276            "    result[0] = ", args[0], '[', args[1], "];\n",
     277            "    result[1] = ", args[0], '[', args[1], " + ", numberOfRows, "];\n");
     278
     279        if (numberOfColumns >= 3)
     280            stringBuilder.flexibleAppend("    result[2] = ", args[0], '[', args[1], " + ", numberOfRows * 2, "];\n");
     281        if (numberOfColumns >= 4)
     282            stringBuilder.flexibleAppend("    result[3] = ", args[0], '[', args[1], " + ", numberOfRows * 3, "];\n");
     283
     284        stringBuilder.flexibleAppend(
     285            "    ", returnName, " = result;\n",
     286            "} while (0);\n");
     287
     288        return;
     289    }
     290
     291    if (nativeFunctionDeclaration.name() == "operator[]=") {
     292        ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
    321293        auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
    322294        auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
    … …  
    325297        unsigned numberOfColumns = numberOfMatrixColumns();
    326298
    327         stringBuilder.flexibleAppend(
    328             metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " m, ", metalParameter2Name, " i) {\n"
    329             "    if (i >= ", numberOfRows, ") return ", metalReturnName, "(0);\n"
    330             "    ", metalReturnName, " result;\n"
    331             "    result[0] = m[i];\n"
    332             "    result[1] = m[i + ", numberOfRows, "];\n"
    333         );
     299        stringBuilder.flexibleAppend("do {\n", metalReturnName, " m = ", args[0], ";\n",
     300            metalParameter2Name, " i = ", args[1], ";\n");
     301
     302        stringBuilder.flexibleAppend(
     303            "    if (i >= ", numberOfRows, ") {", returnName, " = m;\nbreak;}\n",
     304            "    m[i] = ", args[2], "[0];\n",
     305            "    m[i + ", numberOfRows, "] = ", args[2], "[1];\n");
    334306        if (numberOfColumns >= 3)
    335             stringBuilder.flexibleAppend("    result[2] = m[i + ", numberOfRows * 2, "];\n");
     307            stringBuilder.flexibleAppend("    m[i + ", numberOfRows * 2, "] = ", args[2], "[2];\n");
    336308        if (numberOfColumns >= 4)
    337             stringBuilder.flexibleAppend("    result[3] = m[i + ", numberOfRows * 3, "];\n");
    338         stringBuilder.append(
    339             "    return result;\n"
    340             "}\n"
    341         );
    342         return stringBuilder.toString();
    343     }
    344 
    345     if (nativeFunctionDeclaration.name() == "operator[]=") {
    346         ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
    347         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
    348         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
    349         auto metalParameter3Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[2]->type());
    350         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
    351 
    352         unsigned numberOfRows = numberOfMatrixRows();
    353         unsigned numberOfColumns = numberOfMatrixColumns();
    354 
    355         stringBuilder.flexibleAppend(
    356             metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " m, ", metalParameter2Name, " i, ", metalParameter3Name, " v) {\n"
    357             "    if (i >= ", numberOfRows, ") return m;\n"
    358             "    m[i] = v[0];\n"
    359             "    m[i + ", numberOfRows, "] = v[1];\n"
    360         );
    361         if (numberOfColumns >= 3)
    362             stringBuilder.flexibleAppend("    m[i + ", numberOfRows * 2, "] = v[2];\n");
    363         if (numberOfColumns >= 4)
    364             stringBuilder.flexibleAppend("    m[i + ", numberOfRows * 3, "] = v[3];\n");
    365         stringBuilder.append(
    366             "    return m;"
    367             "}\n"
    368         );
    369         return stringBuilder.toString();
     309            stringBuilder.flexibleAppend("    m[i + ", numberOfRows * 3, "] = ", args[2], "[3];\n");
     310        stringBuilder.flexibleAppend(
     311            "    ", returnName, " = m;\n",
     312            "} while(0);\n");
     313        return;
    370314    }
    371315
    372316    if (nativeFunctionDeclaration.isOperator()) {
     317        auto operatorName = nativeFunctionDeclaration.name().substring("operator"_str.length());
    373318        if (nativeFunctionDeclaration.parameters().size() == 1) {
    374             auto operatorName = nativeFunctionDeclaration.name().substring("operator"_str.length());
    375319            auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
    376             auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
    377             stringBuilder.flexibleAppend(
    378                 metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"
    379                 "    return ", operatorName, "x;\n"
    380                 "}\n"
    381             );
    382             return stringBuilder.toString();
     320            stringBuilder.flexibleAppend(
     321                "{\n", metalParameterName, " x = ", args[0], ";\n",
     322                returnName, " = ", operatorName, "x;\n}\n");
     323            return;
    383324        }
    384325
    385326        ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
    386         auto operatorName = nativeFunctionDeclaration.name().substring("operator"_str.length());
    387         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
    388         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
    389         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
    390         stringBuilder.flexibleAppend(
    391             metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " x, ", metalParameter2Name, " y) {\n"
    392             "    return x ", operatorName, " y;\n"
    393             "}\n"
    394         );
    395         return stringBuilder.toString();
     327        stringBuilder.flexibleAppend(
     328            returnName, " = ", args[0], ' ', operatorName, ' ', args[1], ";\n");
     329        return;
    396330    }
    397331
    … …  
    421355        || nativeFunctionDeclaration.name() == "asfloat") {
    422356        ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
    423         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
    424         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
    425         stringBuilder.flexibleAppend(
    426             metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"
    427             "    return ", mapFunctionName(nativeFunctionDeclaration.name()), "(x);\n"
    428             "}\n"
    429         );
    430         return stringBuilder.toString();
     357        stringBuilder.flexibleAppend(
     358            returnName, " = ", mapFunctionName(nativeFunctionDeclaration.name()), '(', args[0], ");\n");
     359        return;
    431360    }
    432361
    433362    if (nativeFunctionDeclaration.name() == "pow" || nativeFunctionDeclaration.name() == "atan2") {
    434363        ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
    435         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
    436         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
    437         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
    438         stringBuilder.flexibleAppend(
    439             metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " x, ", metalParameter2Name, " y) {\n"
    440             "    return ", nativeFunctionDeclaration.name(), "(x, y);\n"
    441             "}\n"
    442         );
    443         return stringBuilder.toString();
     364        stringBuilder.flexibleAppend(
     365            returnName, " = ", nativeFunctionDeclaration.name(), "(", args[0], ", ", args[1], ");\n");
     366        return;
    444367    }
    445368
    446369    if (nativeFunctionDeclaration.name() == "AllMemoryBarrierWithGroupSync") {
    447370        ASSERT(!nativeFunctionDeclaration.parameters().size());
    448         stringBuilder.flexibleAppend(
    449             "void ", outputFunctionName, "() {\n"
    450             "    threadgroup_barrier(mem_flags::mem_device);\n"
    451             "    threadgroup_barrier(mem_flags::mem_threadgroup);\n"
    452             "    threadgroup_barrier(mem_flags::mem_texture);\n"
    453             "}\n"
    454         );
    455         return stringBuilder.toString();
     371        stringBuilder.append(
     372            "threadgroup_barrier(mem_flags::mem_device);\n"
     373            "threadgroup_barrier(mem_flags::mem_threadgroup);\n"
     374            "threadgroup_barrier(mem_flags::mem_texture);\n");
     375        return;
    456376    }
    457377
    458378    if (nativeFunctionDeclaration.name() == "DeviceMemoryBarrierWithGroupSync") {
    459379        ASSERT(!nativeFunctionDeclaration.parameters().size());
    460         stringBuilder.flexibleAppend(
    461             "void ", outputFunctionName, "() {\n"
    462             "    threadgroup_barrier(mem_flags::mem_device);\n"
    463             "}\n"
    464         );
    465         return stringBuilder.toString();
     380        stringBuilder.append(
     381            "threadgroup_barrier(mem_flags::mem_device);\n");
     382        return;
    466383    }
    467384
    468385    if (nativeFunctionDeclaration.name() == "GroupMemoryBarrierWithGroupSync") {
    469386        ASSERT(!nativeFunctionDeclaration.parameters().size());
    470         stringBuilder.flexibleAppend(
    471             "void ", outputFunctionName, "() {\n"
    472             "    threadgroup_barrier(mem_flags::mem_threadgroup);\n"
    473             "}\n"
    474         );
    475         return stringBuilder.toString();
     387        stringBuilder.append(
     388            "threadgroup_barrier(mem_flags::mem_threadgroup);\n");
     389        return;
    476390    }
    477391
    … …  
    479393        if (nativeFunctionDeclaration.name() == "InterlockedCompareExchange") {
    480394            ASSERT(nativeFunctionDeclaration.parameters().size() == 4);
    481             auto& firstArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0]->type());
    482             auto firstArgumentAddressSpace = firstArgumentPointer.addressSpace();
    483             auto firstArgumentPointee = typeNamer.mangledNameForType(firstArgumentPointer.elementType());
    484             auto secondArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
    485             auto thirdArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[2]->type());
    486             auto& fourthArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[3]->type());
    487             auto fourthArgumentAddressSpace = fourthArgumentPointer.addressSpace();
    488             auto fourthArgumentPointee = typeNamer.mangledNameForType(fourthArgumentPointer.elementType());
    489             stringBuilder.flexibleAppend(
    490                 "void ", outputFunctionName, '(', toString(firstArgumentAddressSpace), ' ', firstArgumentPointee, "* object, ", secondArgument, " compare, ", thirdArgument, " desired, ", toString(fourthArgumentAddressSpace), ' ', fourthArgumentPointee, "* out) {\n"
    491                 "    atomic_compare_exchange_weak_explicit(object, &compare, desired, memory_order_relaxed, memory_order_relaxed);\n"
    492                 "    *out = compare;\n"
    493                 "}\n"
    494             );
    495             return stringBuilder.toString();
     395            stringBuilder.flexibleAppend(
     396                "atomic_compare_exchange_weak_explicit(", args[0], ", &", args[1], ", ", args[2], ", memory_order_relaxed, memory_order_relaxed);\n",
     397                '*', args[3], " = ", args[1], ";\n");
     398            return;
    496399        }
    497400
    498401        ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
    499         auto& firstArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0]->type());
    500         auto firstArgumentAddressSpace = firstArgumentPointer.addressSpace();
    501         auto firstArgumentPointee = typeNamer.mangledNameForType(firstArgumentPointer.elementType());
    502         auto secondArgument = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
    503         auto& thirdArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[2]->type());
    504         auto thirdArgumentAddressSpace = thirdArgumentPointer.addressSpace();
    505         auto thirdArgumentPointee = typeNamer.mangledNameForType(thirdArgumentPointer.elementType());
    506402        auto name = atomicName(nativeFunctionDeclaration.name().substring("Interlocked"_str.length()));
    507403        stringBuilder.flexibleAppend(
    508             "void ", outputFunctionName, '(', toString(firstArgumentAddressSpace), ' ', firstArgumentPointee, "* object, ", secondArgument, " operand, ", toString(thirdArgumentAddressSpace), ' ', thirdArgumentPointee, "* out) {\n"
    509             "    *out = atomic_", name, "_explicit(object, operand, memory_order_relaxed);\n"
    510             "}\n"
    511         );
    512         return stringBuilder.toString();
     404            '*', args[2], " = atomic_", name, "_explicit(", args[0], ", ", args[1], ", memory_order_relaxed);\n");
     405        return;
    513406    }
    514407
    … …  
    522415        auto returnVectorLength = vectorLength(returnType);
    523416
    524         auto metalParameter1Name = typeNamer.mangledNameForType(textureType);
    525         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
    526         auto metalParameter3Name = typeNamer.mangledNameForType(locationType);
    527         Optional<MangledTypeName> metalParameter4Name;
    528         if (nativeFunctionDeclaration.parameters().size() == 4)
    529             metalParameter4Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[3]->type());
    530         auto metalReturnName = typeNamer.mangledNameForType(returnType);
    531         stringBuilder.flexibleAppend(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " theTexture, ", metalParameter2Name, " theSampler, ", metalParameter3Name, " location");
    532         if (metalParameter4Name)
    533             stringBuilder.flexibleAppend(", ", *metalParameter4Name, " offset");
    534         stringBuilder.append(
    535             ") {\n"
    536             "    return theTexture.sample(theSampler, "
    537         );
     417        stringBuilder.flexibleAppend(
     418            returnName, " = ", args[0], ".sample(", args[1], ", ");
     419
    538420        if (textureType.isTextureArray()) {
    539421            ASSERT(locationVectorLength > 1);
    540             stringBuilder.flexibleAppend("location.", "xyzw"_str.substring(0, locationVectorLength - 1), ", location.", "xyzw"_str.substring(locationVectorLength - 1, 1));
     422            stringBuilder.flexibleAppend(args[2], '.', "xyzw"_str.substring(0, locationVectorLength - 1), ", ", args[2], '.', "xyzw"_str.substring(locationVectorLength - 1, 1));
    541423        } else
    542             stringBuilder.append("location");
    543         if (metalParameter4Name)
    544             stringBuilder.append(", offset");
     424            stringBuilder.flexibleAppend(args[2]);
     425        if (nativeFunctionDeclaration.parameters().size() == 4)
     426            stringBuilder.flexibleAppend(", ", args[3]);
    545427        stringBuilder.append(")");
    546428        if (!textureType.isDepthTexture())
    547429            stringBuilder.flexibleAppend(".", "xyzw"_str.substring(0, returnVectorLength));
    548         stringBuilder.append(
    549             ";\n"
    550             "}\n"
    551         );
    552         return stringBuilder.toString();
     430        stringBuilder.append(";\n");
     431
     432        return;
    553433    }
    554434
    … …  
    562442        auto returnVectorLength = vectorLength(returnType);
    563443
    564         auto metalParameter1Name = typeNamer.mangledNameForType(textureType);
    565         auto metalParameter2Name = typeNamer.mangledNameForType(locationType);
    566444        auto metalReturnName = typeNamer.mangledNameForType(returnType);
    567         stringBuilder.flexibleAppend(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, " theTexture, ", metalParameter2Name, " location) {\n");
     445        stringBuilder.append("do {\n");
     446
    568447        if (textureType.isTextureArray()) {
    569448            ASSERT(locationVectorLength > 1);
    … …  
    571450            for (int i = 0; i < locationVectorLength - 1; ++i) {
    572451                auto suffix = "xyzw"_str.substring(i, 1);
    573                 stringBuilder.flexibleAppend("    if (location.", suffix, " < 0 || static_cast<uint32_t>(location.", suffix, ") >= theTexture.get_", dimensions[i], "()) return ", metalReturnName, "(0);\n");
     452                stringBuilder.flexibleAppend("    if (", args[1], '.', suffix, " < 0 || static_cast<uint32_t>(", args[1], '.', suffix, ") >= ", args[0], ".get_", dimensions[i], "()) {", returnName, " = ", metalReturnName, "(0); break;}\n");
    574453            }
    575454            auto suffix = "xyzw"_str.substring(locationVectorLength - 1, 1);
    576             stringBuilder.flexibleAppend("    if (location.", suffix, " < 0 || static_cast<uint32_t>(location.", suffix, ") >= theTexture.get_array_size()) return ", metalReturnName, "(0);\n");
     455            stringBuilder.flexibleAppend("    if (", args[1], '.', suffix, " < 0 || static_cast<uint32_t>(", args[1], '.', suffix, ") >= ", args[0], ".get_array_size()) {", returnName, " = ", metalReturnName, "(0); break;}\n");
    577456        } else {
    578457            if (locationVectorLength == 1)
    579                 stringBuilder.flexibleAppend("    if (location < 0 || static_cast<uint32_t>(location) >= theTexture.get_width()) return ", metalReturnName, "(0);\n");
     458                stringBuilder.flexibleAppend("    if (", args[1], " < 0 || static_cast<uint32_t>(", args[1], ") >= ", args[0], ".get_width()) { ", returnName, " = ", metalReturnName, "(0); break;}\n");
    580459            else {
    581460                stringBuilder.flexibleAppend(
    582                     "    if (location.x < 0 || static_cast<uint32_t>(location.x) >= theTexture.get_width()) return ", metalReturnName, "(0);\n"
    583                     "    if (location.y < 0 || static_cast<uint32_t>(location.y) >= theTexture.get_height()) return ", metalReturnName, "(0);\n"
    584                 );
     461                    "    if (", args[1], ".x < 0 || static_cast<uint32_t>(", args[1], ".x) >= ", args[0], ".get_width()) {", returnName, " = ", metalReturnName, "(0); break;}\n"
     462                    "    if (", args[1], ".y < 0 || static_cast<uint32_t>(", args[1], ".y) >= ", args[0], ".get_height()) {", returnName, " = ", metalReturnName, "(0); break;}\n");
     463
    585464                if (locationVectorLength >= 3)
    586                     stringBuilder.flexibleAppend("    if (location.z < 0 || static_cast<uint32_t>(location.z) >= theTexture.get_depth()) return ", metalReturnName, "(0);\n");
     465                    stringBuilder.flexibleAppend("    if (", args[1], ".z < 0 || static_cast<uint32_t>(", args[1], ".z) >= ", args[0], ".get_depth()) {", returnName, " = ", metalReturnName, "(0); break;}\n");
    587466            }
    588467        }
    589         stringBuilder.append("    return theTexture.read(");
     468        stringBuilder.flexibleAppend("    ", returnName, " = ", args[0], ".read(");
    590469        if (textureType.isTextureArray()) {
    591470            ASSERT(locationVectorLength > 1);
    592             stringBuilder.flexibleAppend("uint", vectorSuffix(locationVectorLength - 1), "(location.", "xyzw"_str.substring(0, locationVectorLength - 1), "), uint(location.", "xyzw"_str.substring(locationVectorLength - 1, 1), ')');
     471            stringBuilder.flexibleAppend("uint", vectorSuffix(locationVectorLength - 1), '(', args[1], '.', "xyzw"_str.substring(0, locationVectorLength - 1), "), uint(", args[1], '.', "xyzw"_str.substring(locationVectorLength - 1, 1), ')');
    593472        } else
    594             stringBuilder.flexibleAppend("uint", vectorSuffix(locationVectorLength), "(location)");
     473            stringBuilder.flexibleAppend("uint", vectorSuffix(locationVectorLength), '(', args[1], ')');
    595474        stringBuilder.append(')');
    596475        if (!textureType.isDepthTexture())
    … …  
    598477        stringBuilder.append(
    599478            ";\n"
    600             "}\n"
    601         );
    602         return stringBuilder.toString();
     479            "} while(0);\n");
     480
     481        return;
    603482    }
    604483
    605484    if (nativeFunctionDeclaration.name() == "load") {
    606485        ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
    607         auto metalParameterName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
    608         auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
    609         stringBuilder.flexibleAppend(
    610             metalReturnName, ' ', outputFunctionName, '(', metalParameterName, " x) {\n"
    611             "    return atomic_load_explicit(x, memory_order_relaxed);\n"
    612             "}\n"
    613         );
    614         return stringBuilder.toString();
     486        stringBuilder.flexibleAppend(
     487            returnName, " = atomic_load_explicit(", args[0], ", memory_order_relaxed);\n");
     488        return;
    615489    }
    616490
    617491    if (nativeFunctionDeclaration.name() == "store") {
    618492        ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
    619         auto metalParameter1Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[0]->type());
    620         auto metalParameter2Name = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[1]->type());
    621         stringBuilder.flexibleAppend("void ", outputFunctionName, '(', metalParameter1Name, " x, ", metalParameter2Name, " y) {\n"
    622             "    atomic_store_explicit(x, y, memory_order_relaxed);\n"
    623             "}\n");
    624         return stringBuilder.toString();
     493        stringBuilder.flexibleAppend(
     494            "atomic_store_explicit(", args[0], ", ", args[1], ", memory_order_relaxed);\n");
     495        return;
    625496    }
    626497
    … …  
    629500
    630501        size_t index = 1;
    631         if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
     502        bool hasMipLevel = !textureType.isWritableTexture() && textureType.textureDimension() != 1;
     503        if (hasMipLevel)
    632504            ++index;
    633         auto widthTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
     505        const MangledVariableName& widthName = args[index];
    634506        ++index;
    635         Optional<MangledTypeName> heightTypeName;
     507        Optional<MangledVariableName> heightName;
    636508        if (textureType.textureDimension() >= 2) {
    637             heightTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
     509            heightName = args[index];
    638510            ++index;
    639511        }
    640         Optional<MangledTypeName> depthTypeName;
     512        Optional<MangledVariableName> depthName;
    641513        if (textureType.textureDimension() >= 3) {
    642             depthTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
     514            depthName = args[index];
    643515            ++index;
    644516        }
    645         Optional<MangledTypeName> elementsTypeName;
     517        Optional<MangledVariableName> elementsName;
    646518        if (textureType.isTextureArray()) {
    647             elementsTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
     519            elementsName = args[index];
    648520            ++index;
    649521        }
    650         Optional<MangledTypeName> numberOfLevelsTypeName;
     522        Optional<MangledVariableName> numberOfLevelsName;
    651523        if (!textureType.isWritableTexture() && textureType.textureDimension() != 1) {
    652             numberOfLevelsTypeName = typeNamer.mangledNameForType(*nativeFunctionDeclaration.parameters()[index]->type());
     524            numberOfLevelsName = args[index];
    653525            ++index;
    654526        }
    655527        ASSERT(index == nativeFunctionDeclaration.parameters().size());
    656528
    657         auto metalParameter1Name = typeNamer.mangledNameForType(textureType);
    658         stringBuilder.flexibleAppend("void ", outputFunctionName, '(', metalParameter1Name, " theTexture");
    659         if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
    660             stringBuilder.append(", uint mipLevel");
    661         stringBuilder.flexibleAppend(", ", widthTypeName, " width");
    662         if (heightTypeName)
    663             stringBuilder.flexibleAppend(", ", *heightTypeName, " height");
    664         if (depthTypeName)
    665             stringBuilder.flexibleAppend(", ", *depthTypeName, " depth");
    666         if (elementsTypeName)
    667             stringBuilder.flexibleAppend(", ", *elementsTypeName, " elements");
    668         if (numberOfLevelsTypeName)
    669             stringBuilder.flexibleAppend(", ", *numberOfLevelsTypeName, " numberOfLevels");
    670         stringBuilder.append(
    671             ") {\n"
    672             "    if (width)\n"
    673             "        *width = theTexture.get_width("
    674         );
    675         if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
    676             stringBuilder.append("mipLevel");
     529        stringBuilder.flexibleAppend(
     530            "if (", widthName, ")\n"
     531            "    *", widthName, " = ", args[0], ".get_width(");
     532
     533        if (hasMipLevel)
     534            stringBuilder.flexibleAppend(args[1]);
    677535        stringBuilder.append(");\n");
    678         if (heightTypeName) {
    679             stringBuilder.append(
    680                 "    if (height)\n"
    681                 "        *height = theTexture.get_height("
    682             );
    683             if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
    684                 stringBuilder.append("mipLevel");
     536        if (heightName) {
     537            stringBuilder.flexibleAppend(
     538                "    if (", *heightName, ")\n"
     539                "        *", *heightName, " = ", args[0], ".get_height(");
     540            if (hasMipLevel)
     541                stringBuilder.flexibleAppend(args[1]);
    685542            stringBuilder.append(");\n");
    686543        }
    687         if (depthTypeName) {
    688             stringBuilder.append(
    689                 "    if (depth)\n"
    690                 "        *depth = theTexture.get_depth("
    691             );
    692             if (!textureType.isWritableTexture() && textureType.textureDimension() != 1)
    693                 stringBuilder.append("mipLevel");
     544        if (depthName) {
     545            stringBuilder.flexibleAppend(
     546                "    if (", *depthName, ")\n"
     547                "        *", *depthName, " = ", args[0], ".get_depth(");
     548            if (hasMipLevel)
     549                stringBuilder.flexibleAppend(args[1]);
    694550            stringBuilder.append(");\n");
    695551        }
    696         if (elementsTypeName) {
    697             stringBuilder.append(
    698                 "    if (elements)\n"
    699                 "        *elements = theTexture.get_array_size();\n"
    700             );
    701         }
    702         if (numberOfLevelsTypeName) {
    703             stringBuilder.append(
    704                 "    if (numberOfLevels)\n"
    705                 "        *numberOfLevels = theTexture.get_num_mip_levels();\n"
    706             );
    707         }
    708         stringBuilder.append("}\n");
    709         return stringBuilder.toString();
     552        if (elementsName) {
     553            stringBuilder.flexibleAppend(
     554                "    if (", *elementsName, ")\n"
     555                "        *", *elementsName, " = ", args[0], ".get_array_size();\n");
     556        }
     557        if (numberOfLevelsName) {
     558            stringBuilder.flexibleAppend(
     559                "    if (", *numberOfLevelsName, ")\n"
     560                "        *", *numberOfLevelsName, " = ", args[0], ".get_num_mip_levels();\n");
     561        }
     562        return;
    710563    }
    711564
    … …  
    755608        auto locationVectorLength = vectorLength(locationType);
    756609
    757         auto metalParameter1Name = typeNamer.mangledNameForType(textureType);
    758         auto metalParameter2Name = typeNamer.mangledNameForType(itemType);
    759         auto metalParameter3Name = typeNamer.mangledNameForType(locationType);
    760610        auto metalInnerTypeName = typeNamer.mangledNameForType(itemVectorInnerType);
    761         stringBuilder.flexibleAppend("void ", outputFunctionName, '(', metalParameter1Name, " theTexture, ", metalParameter2Name, " item, ", metalParameter3Name, " location) {\n");
     611
     612        stringBuilder.append("do {\n");
    762613        if (textureType.isTextureArray()) {
    763614            ASSERT(locationVectorLength > 1);
    … …  
    765616            for (int i = 0; i < locationVectorLength - 1; ++i) {
    766617                auto suffix = "xyzw"_str.substring(i, 1);
    767                 stringBuilder.flexibleAppend("    if (location.", suffix, " >= theTexture.get_", dimensions[i], "()) return;\n");
     618                stringBuilder.flexibleAppend("    if (", args[2], ".", suffix, " >= ", args[0], ".get_", dimensions[i], "()) break;\n");
    768619            }
    769620            auto suffix = "xyzw"_str.substring(locationVectorLength - 1, 1);
    770             stringBuilder.flexibleAppend("    if (location.", suffix, " >= theTexture.get_array_size()) return;\n");
     621            stringBuilder.flexibleAppend("    if (", args[2], '.', suffix, " >= ", args[0], ".get_array_size()) break;\n");
    771622        } else {
    772623            if (locationVectorLength == 1)
    773                 stringBuilder.append("    if (location >= theTexture.get_width()) return;\n");
     624                stringBuilder.flexibleAppend("    if (", args[2], " >= ", args[0], ".get_width()) break;\n");
    774625            else {
    775                 stringBuilder.append(
    776                     "    if (location.x >= theTexture.get_width()) return;\n"
    777                     "    if (location.y >= theTexture.get_height()) return;\n"
    778                 );
     626                stringBuilder.flexibleAppend(
     627                    "    if (", args[2], ".x >= ", args[0], ".get_width()) break;\n"
     628                    "    if (", args[2], ".y >= ", args[0], ".get_height()) break;\n");
    779629                if (locationVectorLength >= 3)
    780                     stringBuilder.append("    if (location.z >= theTexture.get_depth()) return;\n");
     630                    stringBuilder.flexibleAppend("    if (", args[2], ".z >= ", args[0], ".get_depth()) break;\n");
    781631            }
    782632        }
    783         stringBuilder.flexibleAppend("    theTexture.write(vec<", metalInnerTypeName, ", 4>(item");
     633        stringBuilder.flexibleAppend("    ", args[0], ".write(vec<", metalInnerTypeName, ", 4>(", args[1]);
    784634        for (int i = 0; i < 4 - itemVectorLength; ++i)
    785635            stringBuilder.append(", 0");
    … …  
    787637        if (textureType.isTextureArray()) {
    788638            ASSERT(locationVectorLength > 1);
    789             stringBuilder.flexibleAppend("uint", vectorSuffix(locationVectorLength - 1), "(location.", "xyzw"_str.substring(0, locationVectorLength - 1), "), uint(location.", "xyzw"_str.substring(locationVectorLength - 1, 1), ')');
     639            stringBuilder.flexibleAppend("uint", vectorSuffix(locationVectorLength - 1), '(', args[2], '.', "xyzw"_str.substring(0, locationVectorLength - 1), "), uint(", args[2], ".", "xyzw"_str.substring(locationVectorLength - 1, 1), ')');
    790640        } else
    791             stringBuilder.flexibleAppend("uint", vectorSuffix(locationVectorLength), "(location)");
     641            stringBuilder.flexibleAppend("uint", vectorSuffix(locationVectorLength), '(', args[2], ')');
    792642        stringBuilder.append(
    793643            ");\n"
    794             "}\n"
    795         );
    796         return stringBuilder.toString();
     644            "} while(0);\n");
     645
     646        return;
    797647    }
    798648
    … …  
    823673
    824674    ASSERT_NOT_REACHED();
    825     return String();
    826675}
    827676

  • trunk/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.h

    r248266 r248280  
    4646class TypeNamer;
    4747
    48 String writeNativeFunction(AST::NativeFunctionDeclaration&, MangledFunctionName outputFunctionName, Intrinsics&, TypeNamer&);
     48void inlineNativeFunction(StringBuilder&, AST::NativeFunctionDeclaration&, MangledVariableName returnName, const Vector<MangledVariableName>& argumentNames, Intrinsics&, TypeNamer&);
    4949
    5050}
Note: See TracChangeset for help on using the changeset viewer.