StringBuilder::append(makeString(...)) is inefficient
[WebKit-https.git] / Source / WebCore / Modules / webgpu / WHLSL / Metal / WHLSLEntryPointScaffolding.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 "WHLSLEntryPointScaffolding.h"
28
29 #if ENABLE(WEBGPU)
30
31 #include "WHLSLBuiltInSemantic.h"
32 #include "WHLSLFunctionDefinition.h"
33 #include "WHLSLGatherEntryPointItems.h"
34 #include "WHLSLPipelineDescriptor.h"
35 #include "WHLSLReferenceType.h"
36 #include "WHLSLResourceSemantic.h"
37 #include "WHLSLStageInOutSemantic.h"
38 #include "WHLSLStructureDefinition.h"
39 #include "WHLSLTypeNamer.h"
40 #include <algorithm>
41 #include <wtf/Optional.h>
42 #include <wtf/text/StringBuilder.h>
43 #include <wtf/text/StringConcatenateNumbers.h>
44
45 namespace WebCore {
46
47 namespace WHLSL {
48
49 namespace Metal {
50
51 static String attributeForSemantic(AST::BuiltInSemantic& builtInSemantic)
52 {
53     switch (builtInSemantic.variable()) {
54     case AST::BuiltInSemantic::Variable::SVInstanceID:
55         return "[[instance_id]]"_str;
56     case AST::BuiltInSemantic::Variable::SVVertexID:
57         return "[[vertex_id]]"_str;
58     case AST::BuiltInSemantic::Variable::PSize:
59         return "[[point_size]]"_str;
60     case AST::BuiltInSemantic::Variable::SVPosition:
61         return "[[position]]"_str;
62     case AST::BuiltInSemantic::Variable::SVIsFrontFace:
63         return "[[front_facing]]"_str;
64     case AST::BuiltInSemantic::Variable::SVSampleIndex:
65         return "[[sample_id]]"_str;
66     case AST::BuiltInSemantic::Variable::SVInnerCoverage:
67         return "[[sample_mask]]"_str;
68     case AST::BuiltInSemantic::Variable::SVTarget:
69         return makeString("[[color(", *builtInSemantic.targetIndex(), ")]]");
70     case AST::BuiltInSemantic::Variable::SVDepth:
71         return "[[depth(any)]]"_str;
72     case AST::BuiltInSemantic::Variable::SVCoverage:
73         return "[[sample_mask]]"_str;
74     case AST::BuiltInSemantic::Variable::SVDispatchThreadID:
75         return "[[thread_position_in_grid]]"_str;
76     case AST::BuiltInSemantic::Variable::SVGroupID:
77         return "[[threadgroup_position_in_grid]]"_str;
78     case AST::BuiltInSemantic::Variable::SVGroupIndex:
79         return "[[thread_index_in_threadgroup]]"_str;
80     default:
81         ASSERT(builtInSemantic.variable() == AST::BuiltInSemantic::Variable::SVGroupThreadID);
82         return "[[thread_position_in_threadgroup]]"_str;
83     }
84 }
85
86 static String attributeForSemantic(AST::Semantic& semantic)
87 {
88     if (WTF::holds_alternative<AST::BuiltInSemantic>(semantic))
89         return attributeForSemantic(WTF::get<AST::BuiltInSemantic>(semantic));
90     auto& stageInOutSemantic = WTF::get<AST::StageInOutSemantic>(semantic);
91     return makeString("[[user(user", stageInOutSemantic.index(), ")]]");
92 }
93
94 EntryPointScaffolding::EntryPointScaffolding(AST::FunctionDefinition& functionDefinition, Intrinsics& intrinsics, TypeNamer& typeNamer, EntryPointItems& entryPointItems, HashMap<Binding*, size_t>& resourceMap, Layout& layout, std::function<String()>&& generateNextVariableName)
95     : m_functionDefinition(functionDefinition)
96     , m_intrinsics(intrinsics)
97     , m_typeNamer(typeNamer)
98     , m_entryPointItems(entryPointItems)
99     , m_resourceMap(resourceMap)
100     , m_layout(layout)
101     , m_generateNextVariableName(generateNextVariableName)
102 {
103     m_namedBindGroups.reserveInitialCapacity(m_layout.size());
104     for (size_t i = 0; i < m_layout.size(); ++i) {
105         NamedBindGroup namedBindGroup;
106         namedBindGroup.structName = m_typeNamer.generateNextTypeName();
107         namedBindGroup.variableName = m_generateNextVariableName();
108         namedBindGroup.argumentBufferIndex = m_layout[i].name; // convertLayout() in GPURenderPipelineMetal.mm makes sure these don't collide.
109         namedBindGroup.namedBindings.reserveInitialCapacity(m_layout[i].bindings.size());
110         for (size_t j = 0; j < m_layout[i].bindings.size(); ++j) {
111             NamedBinding namedBinding;
112             namedBinding.elementName = m_typeNamer.generateNextStructureElementName();
113             namedBinding.index = m_layout[i].bindings[j].internalName;
114             WTF::visit(WTF::makeVisitor([&](UniformBufferBinding& uniformBufferBinding) {
115                 LengthInformation lengthInformation { m_typeNamer.generateNextStructureElementName(), m_generateNextVariableName(), uniformBufferBinding.lengthName };
116                 namedBinding.lengthInformation = lengthInformation;
117             }, [&](SamplerBinding&) {
118             }, [&](TextureBinding&) {
119             }, [&](StorageBufferBinding& storageBufferBinding) {
120                 LengthInformation lengthInformation { m_typeNamer.generateNextStructureElementName(), m_generateNextVariableName(), storageBufferBinding.lengthName };
121                 namedBinding.lengthInformation = lengthInformation;
122             }), m_layout[i].bindings[j].binding);
123             namedBindGroup.namedBindings.uncheckedAppend(WTFMove(namedBinding));
124         }
125         m_namedBindGroups.uncheckedAppend(WTFMove(namedBindGroup));
126     }
127
128     for (size_t i = 0; i < m_entryPointItems.inputs.size(); ++i) {
129         if (!WTF::holds_alternative<AST::BuiltInSemantic>(*m_entryPointItems.inputs[i].semantic))
130             continue;
131         NamedBuiltIn namedBuiltIn;
132         namedBuiltIn.indexInEntryPointItems = i;
133         namedBuiltIn.variableName = m_generateNextVariableName();
134         m_namedBuiltIns.append(WTFMove(namedBuiltIn));
135     }
136
137     m_parameterVariables.reserveInitialCapacity(m_functionDefinition.parameters().size());
138     for (size_t i = 0; i < m_functionDefinition.parameters().size(); ++i)
139         m_parameterVariables.uncheckedAppend(m_generateNextVariableName());
140 }
141
142 String EntryPointScaffolding::resourceHelperTypes()
143 {
144     StringBuilder stringBuilder;
145     for (size_t i = 0; i < m_layout.size(); ++i) {
146         stringBuilder.flexibleAppend("struct ", m_namedBindGroups[i].structName, " {\n");
147         Vector<std::pair<unsigned, String>> structItems;
148         for (size_t j = 0; j < m_layout[i].bindings.size(); ++j) {
149             auto iterator = m_resourceMap.find(&m_layout[i].bindings[j]);
150             if (iterator == m_resourceMap.end())
151                 continue;
152             auto& type = m_entryPointItems.inputs[iterator->value].unnamedType->unifyNode();
153             if (is<AST::UnnamedType>(type) && is<AST::ReferenceType>(downcast<AST::UnnamedType>(type))) {
154                 auto& referenceType = downcast<AST::ReferenceType>(downcast<AST::UnnamedType>(type));
155                 auto mangledTypeName = m_typeNamer.mangledNameForType(referenceType.elementType());
156                 auto addressSpace = toString(referenceType.addressSpace());
157                 auto elementName = m_namedBindGroups[i].namedBindings[j].elementName;
158                 auto index = m_namedBindGroups[i].namedBindings[j].index;
159                 structItems.append(std::make_pair(index, makeString(addressSpace, " ", mangledTypeName, "* ", elementName, " [[id(", index, ")]];")));
160                 if (auto lengthInformation = m_namedBindGroups[i].namedBindings[j].lengthInformation)
161                     structItems.append(std::make_pair(lengthInformation->index, makeString("uint2 ", lengthInformation->elementName, " [[id(", lengthInformation->index, ")]];")));
162             } else if (is<AST::NamedType>(type) && is<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(type))) {
163                 auto& namedType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(type));
164                 auto mangledTypeName = m_typeNamer.mangledNameForType(namedType);
165                 auto elementName = m_namedBindGroups[i].namedBindings[j].elementName;
166                 auto index = m_namedBindGroups[i].namedBindings[j].index;
167                 structItems.append(std::make_pair(index, makeString(mangledTypeName, ' ', elementName, " [[id(", index, ")]];")));
168             }
169         }
170         std::sort(structItems.begin(), structItems.end(), [](const std::pair<unsigned, String>& left, const std::pair<unsigned, String>& right) {
171             return left.first < right.first;
172         });
173         for (const auto& structItem : structItems)
174             stringBuilder.flexibleAppend("    ", structItem.second, '\n');
175         stringBuilder.append("};\n\n");
176     }
177     return stringBuilder.toString();
178 }
179
180 Optional<String> EntryPointScaffolding::resourceSignature()
181 {
182     if (!m_layout.size())
183         return WTF::nullopt;
184
185     StringBuilder stringBuilder;
186     for (size_t i = 0; i < m_layout.size(); ++i) {
187         if (i)
188             stringBuilder.append(", ");
189         auto& namedBindGroup = m_namedBindGroups[i];
190         stringBuilder.flexibleAppend("device ", namedBindGroup.structName, "& ", namedBindGroup.variableName, " [[buffer(", namedBindGroup.argumentBufferIndex, ")]]");
191     }
192     return stringBuilder.toString();
193 }
194
195 static String internalTypeForSemantic(const AST::BuiltInSemantic& builtInSemantic)
196 {
197     switch (builtInSemantic.variable()) {
198     case AST::BuiltInSemantic::Variable::SVInstanceID:
199         return "uint"_str;
200     case AST::BuiltInSemantic::Variable::SVVertexID:
201         return "uint"_str;
202     case AST::BuiltInSemantic::Variable::PSize:
203         return "float"_str;
204     case AST::BuiltInSemantic::Variable::SVPosition:
205         return "float4"_str;
206     case AST::BuiltInSemantic::Variable::SVIsFrontFace:
207         return "bool"_str;
208     case AST::BuiltInSemantic::Variable::SVSampleIndex:
209         return "uint"_str;
210     case AST::BuiltInSemantic::Variable::SVInnerCoverage:
211         return "uint"_str;
212     case AST::BuiltInSemantic::Variable::SVTarget:
213         return String();
214     case AST::BuiltInSemantic::Variable::SVDepth:
215         return "float"_str;
216     case AST::BuiltInSemantic::Variable::SVCoverage:
217         return "uint"_str;
218     case AST::BuiltInSemantic::Variable::SVDispatchThreadID:
219         return "uint3"_str;
220     case AST::BuiltInSemantic::Variable::SVGroupID:
221         return "uint3"_str;
222     case AST::BuiltInSemantic::Variable::SVGroupIndex:
223         return "uint"_str;
224     default:
225         ASSERT(builtInSemantic.variable() == AST::BuiltInSemantic::Variable::SVGroupThreadID);
226         return "uint3"_str;
227     }
228 }
229
230 Optional<String> EntryPointScaffolding::builtInsSignature()
231 {
232     if (!m_namedBuiltIns.size())
233         return WTF::nullopt;
234
235     StringBuilder stringBuilder;
236     for (size_t i = 0; i < m_namedBuiltIns.size(); ++i) {
237         if (i)
238             stringBuilder.append(", ");
239         auto& namedBuiltIn = m_namedBuiltIns[i];
240         auto& item = m_entryPointItems.inputs[namedBuiltIn.indexInEntryPointItems];
241         auto& builtInSemantic = WTF::get<AST::BuiltInSemantic>(*item.semantic);
242         auto internalType = internalTypeForSemantic(builtInSemantic);
243         if (internalType.isNull())
244             internalType = m_typeNamer.mangledNameForType(*item.unnamedType);
245         auto variableName = namedBuiltIn.variableName;
246         stringBuilder.flexibleAppend(internalType, ' ', variableName, ' ', attributeForSemantic(builtInSemantic));
247     }
248     return stringBuilder.toString();
249 }
250
251 String EntryPointScaffolding::mangledInputPath(Vector<String>& path)
252 {
253     ASSERT(!path.isEmpty());
254     StringBuilder stringBuilder;
255     bool found = false;
256     AST::StructureDefinition* structureDefinition = nullptr;
257     for (size_t i = 0; i < m_functionDefinition.parameters().size(); ++i) {
258         if (m_functionDefinition.parameters()[i]->name() == path[0]) {
259             stringBuilder.append(m_parameterVariables[i]);
260             auto& unifyNode = m_functionDefinition.parameters()[i]->type()->unifyNode();
261             if (is<AST::NamedType>(unifyNode)) {
262                 auto& namedType = downcast<AST::NamedType>(unifyNode);
263                 if (is<AST::StructureDefinition>(namedType))
264                     structureDefinition = &downcast<AST::StructureDefinition>(namedType);
265             }
266             found = true;
267             break;
268         }
269     }
270     ASSERT(found);
271     for (size_t i = 1; i < path.size(); ++i) {
272         ASSERT(structureDefinition);
273         auto* next = structureDefinition->find(path[i]);
274         ASSERT(next);
275         stringBuilder.flexibleAppend('.', m_typeNamer.mangledNameForStructureElement(*next));
276         structureDefinition = nullptr;
277         auto& unifyNode = next->type().unifyNode();
278         if (is<AST::NamedType>(unifyNode)) {
279             auto& namedType = downcast<AST::NamedType>(unifyNode);
280             if (is<AST::StructureDefinition>(namedType))
281                 structureDefinition = &downcast<AST::StructureDefinition>(namedType);
282         }
283     }
284
285     return stringBuilder.toString();
286 }
287
288 String EntryPointScaffolding::mangledOutputPath(Vector<String>& path)
289 {
290     StringBuilder stringBuilder;
291
292     AST::StructureDefinition* structureDefinition = nullptr;
293     auto& unifyNode = m_functionDefinition.type().unifyNode();
294     structureDefinition = &downcast<AST::StructureDefinition>(downcast<AST::NamedType>(unifyNode));
295     for (auto& component : path) {
296         ASSERT(structureDefinition);
297         auto* next = structureDefinition->find(component);
298         ASSERT(next);
299         stringBuilder.flexibleAppend('.', m_typeNamer.mangledNameForStructureElement(*next));
300         structureDefinition = nullptr;
301         auto& unifyNode = next->type().unifyNode();
302         if (is<AST::NamedType>(unifyNode)) {
303             auto& namedType = downcast<AST::NamedType>(unifyNode);
304             if (is<AST::StructureDefinition>(namedType))
305                 structureDefinition = &downcast<AST::StructureDefinition>(namedType);
306         }
307     }
308
309     return stringBuilder.toString();
310 }
311
312 String EntryPointScaffolding::unpackResourcesAndNamedBuiltIns()
313 {
314     StringBuilder stringBuilder;
315     for (size_t i = 0; i < m_functionDefinition.parameters().size(); ++i)
316         stringBuilder.flexibleAppend(m_typeNamer.mangledNameForType(*m_functionDefinition.parameters()[i]->type()), ' ', m_parameterVariables[i], ";\n");
317
318     for (size_t i = 0; i < m_layout.size(); ++i) {
319         auto variableName = m_namedBindGroups[i].variableName;
320         for (size_t j = 0; j < m_layout[i].bindings.size(); ++j) {
321             auto iterator = m_resourceMap.find(&m_layout[i].bindings[j]);
322             if (iterator == m_resourceMap.end())
323                 continue;
324             if (m_namedBindGroups[i].namedBindings[j].lengthInformation) {
325                 auto& path = m_entryPointItems.inputs[iterator->value].path;
326                 auto elementName = m_namedBindGroups[i].namedBindings[j].elementName;
327                 auto lengthElementName = m_namedBindGroups[i].namedBindings[j].lengthInformation->elementName;
328                 auto lengthTemporaryName = m_namedBindGroups[i].namedBindings[j].lengthInformation->temporaryName;
329
330                 auto& unnamedType = *m_entryPointItems.inputs[iterator->value].unnamedType;
331                 auto mangledTypeName = m_typeNamer.mangledNameForType(downcast<AST::ReferenceType>(unnamedType).elementType());
332
333                 stringBuilder.flexibleAppend(
334                     "size_t ", lengthTemporaryName, " = ", variableName, '.', lengthElementName, ".y;\n",
335                     lengthTemporaryName, " = ", lengthTemporaryName, " << 32;\n",
336                     lengthTemporaryName, " = ", lengthTemporaryName, " | ", variableName, '.', lengthElementName, ".x;\n",
337                     lengthTemporaryName, " = ", lengthTemporaryName, " / sizeof(", mangledTypeName, ");\n",
338                     "if (", lengthTemporaryName, " > 0xFFFFFFFF) ", lengthTemporaryName, " = 0xFFFFFFFF;\n",
339                     mangledInputPath(path), " = { ", variableName, '.', elementName, ", static_cast<uint32_t>(", lengthTemporaryName, ") };\n"
340                 );
341             } else {
342                 auto& path = m_entryPointItems.inputs[iterator->value].path;
343                 auto elementName = m_namedBindGroups[i].namedBindings[j].elementName;
344                 stringBuilder.flexibleAppend(mangledInputPath(path), " = ", variableName, '.', elementName, ";\n");
345             }
346         }
347     }
348
349     for (auto& namedBuiltIn : m_namedBuiltIns) {
350         auto& item = m_entryPointItems.inputs[namedBuiltIn.indexInEntryPointItems];
351         auto& path = item.path;
352         auto& variableName = namedBuiltIn.variableName;
353         auto mangledTypeName = m_typeNamer.mangledNameForType(*item.unnamedType);
354         stringBuilder.flexibleAppend(mangledInputPath(path), " = ", mangledTypeName, '(', variableName, ");\n");
355     }
356     return stringBuilder.toString();
357 }
358
359 VertexEntryPointScaffolding::VertexEntryPointScaffolding(AST::FunctionDefinition& functionDefinition, Intrinsics& intrinsics, TypeNamer& typeNamer, EntryPointItems& entryPointItems, HashMap<Binding*, size_t>& resourceMap, Layout& layout, std::function<String()>&& generateNextVariableName, HashMap<VertexAttribute*, size_t>& matchedVertexAttributes)
360     : EntryPointScaffolding(functionDefinition, intrinsics, typeNamer, entryPointItems, resourceMap, layout, WTFMove(generateNextVariableName))
361     , m_matchedVertexAttributes(matchedVertexAttributes)
362     , m_stageInStructName(typeNamer.generateNextTypeName())
363     , m_returnStructName(typeNamer.generateNextTypeName())
364     , m_stageInParameterName(m_generateNextVariableName())
365 {
366     m_namedStageIns.reserveInitialCapacity(m_matchedVertexAttributes.size());
367     for (auto& keyValuePair : m_matchedVertexAttributes) {
368         NamedStageIn namedStageIn;
369         namedStageIn.indexInEntryPointItems = keyValuePair.value;
370         namedStageIn.elementName = m_typeNamer.generateNextStructureElementName();
371         namedStageIn.attributeIndex = keyValuePair.key->metalLocation;
372         m_namedStageIns.uncheckedAppend(WTFMove(namedStageIn));
373     }
374
375     m_namedOutputs.reserveInitialCapacity(m_entryPointItems.outputs.size());
376     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
377         auto& outputItem = m_entryPointItems.outputs[i];
378         NamedOutput namedOutput;
379         namedOutput.elementName = m_typeNamer.generateNextStructureElementName();
380         if (WTF::holds_alternative<AST::BuiltInSemantic>(*outputItem.semantic))
381             namedOutput.internalTypeName = internalTypeForSemantic(WTF::get<AST::BuiltInSemantic>(*outputItem.semantic));
382         if (namedOutput.internalTypeName.isNull())
383             namedOutput.internalTypeName = m_typeNamer.mangledNameForType(*outputItem.unnamedType);
384         m_namedOutputs.uncheckedAppend(WTFMove(namedOutput));
385     }
386 }
387
388 String VertexEntryPointScaffolding::helperTypes()
389 {
390     StringBuilder stringBuilder;
391
392     stringBuilder.flexibleAppend("struct ", m_stageInStructName, " {\n");
393     for (auto& namedStageIn : m_namedStageIns) {
394         auto mangledTypeName = m_typeNamer.mangledNameForType(*m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].unnamedType);
395         auto elementName = namedStageIn.elementName;
396         auto attributeIndex = namedStageIn.attributeIndex;
397         stringBuilder.flexibleAppend("    ", mangledTypeName, ' ', elementName, " [[attribute(", attributeIndex, ")]];\n");
398     }
399     stringBuilder.flexibleAppend(
400         "};\n\n"
401         "struct ", m_returnStructName, " {\n"
402     );
403     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
404         auto& outputItem = m_entryPointItems.outputs[i];
405         auto& internalTypeName = m_namedOutputs[i].internalTypeName;
406         auto elementName = m_namedOutputs[i].elementName;
407         auto attribute = attributeForSemantic(*outputItem.semantic);
408         stringBuilder.flexibleAppend("    ", internalTypeName, ' ', elementName, ' ', attribute, ";\n");
409     }
410     stringBuilder.flexibleAppend(
411         "};\n\n",
412         resourceHelperTypes()
413     );
414
415     return stringBuilder.toString();
416 }
417
418 String VertexEntryPointScaffolding::signature(String& functionName)
419 {
420     StringBuilder stringBuilder;
421
422     stringBuilder.flexibleAppend("vertex ", m_returnStructName, ' ', functionName, '(', m_stageInStructName, ' ', m_stageInParameterName, " [[stage_in]]");
423     if (auto resourceSignature = this->resourceSignature())
424         stringBuilder.flexibleAppend(", ", *resourceSignature);
425     if (auto builtInsSignature = this->builtInsSignature())
426         stringBuilder.flexibleAppend(", ", *builtInsSignature);
427     stringBuilder.append(')');
428
429     return stringBuilder.toString();
430 }
431
432 String VertexEntryPointScaffolding::unpack()
433 {
434     StringBuilder stringBuilder;
435
436     stringBuilder.append(unpackResourcesAndNamedBuiltIns());
437
438     for (auto& namedStageIn : m_namedStageIns) {
439         auto& path = m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].path;
440         auto& elementName = namedStageIn.elementName;
441         stringBuilder.flexibleAppend(mangledInputPath(path), " = ", m_stageInParameterName, '.', elementName, ";\n");
442     }
443
444     return stringBuilder.toString();
445 }
446
447 String VertexEntryPointScaffolding::pack(const String& inputVariableName, const String& outputVariableName)
448 {
449     StringBuilder stringBuilder;
450
451     stringBuilder.flexibleAppend(m_returnStructName, ' ', outputVariableName, ";\n");
452     if (m_entryPointItems.outputs.size() == 1 && !m_entryPointItems.outputs[0].path.size()) {
453         auto& elementName = m_namedOutputs[0].elementName;
454         stringBuilder.flexibleAppend(outputVariableName, '.', elementName, " = ", inputVariableName, ";\n");
455         return stringBuilder.toString();
456     }
457     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
458         auto& elementName = m_namedOutputs[i].elementName;
459         auto& internalTypeName = m_namedOutputs[i].internalTypeName;
460         auto& path = m_entryPointItems.outputs[i].path;
461         stringBuilder.flexibleAppend(outputVariableName, '.', elementName, " = ", internalTypeName, '(', inputVariableName, mangledOutputPath(path), ");\n");
462     }
463     return stringBuilder.toString();
464 }
465
466 FragmentEntryPointScaffolding::FragmentEntryPointScaffolding(AST::FunctionDefinition& functionDefinition, Intrinsics& intrinsics, TypeNamer& typeNamer, EntryPointItems& entryPointItems, HashMap<Binding*, size_t>& resourceMap, Layout& layout, std::function<String()>&& generateNextVariableName, HashMap<AttachmentDescriptor*, size_t>&)
467     : EntryPointScaffolding(functionDefinition, intrinsics, typeNamer, entryPointItems, resourceMap, layout, WTFMove(generateNextVariableName))
468     , m_stageInStructName(typeNamer.generateNextTypeName())
469     , m_returnStructName(typeNamer.generateNextTypeName())
470     , m_stageInParameterName(m_generateNextVariableName())
471 {
472     for (size_t i = 0; i < m_entryPointItems.inputs.size(); ++i) {
473         auto& inputItem = m_entryPointItems.inputs[i];
474         if (!WTF::holds_alternative<AST::StageInOutSemantic>(*inputItem.semantic))
475             continue;
476         auto& stageInOutSemantic = WTF::get<AST::StageInOutSemantic>(*inputItem.semantic);
477         NamedStageIn namedStageIn;
478         namedStageIn.indexInEntryPointItems = i;
479         namedStageIn.elementName = m_typeNamer.generateNextStructureElementName();
480         namedStageIn.attributeIndex = stageInOutSemantic.index();
481         m_namedStageIns.append(WTFMove(namedStageIn));
482     }
483
484     m_namedOutputs.reserveInitialCapacity(m_entryPointItems.outputs.size());
485     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
486         auto& outputItem = m_entryPointItems.outputs[i];
487         NamedOutput namedOutput;
488         namedOutput.elementName = m_typeNamer.generateNextStructureElementName();
489         if (WTF::holds_alternative<AST::BuiltInSemantic>(*outputItem.semantic))
490             namedOutput.internalTypeName = internalTypeForSemantic(WTF::get<AST::BuiltInSemantic>(*outputItem.semantic));
491         if (namedOutput.internalTypeName.isNull())
492             namedOutput.internalTypeName = m_typeNamer.mangledNameForType(*outputItem.unnamedType);
493         m_namedOutputs.uncheckedAppend(WTFMove(namedOutput));
494     }
495 }
496
497 String FragmentEntryPointScaffolding::helperTypes()
498 {
499     StringBuilder stringBuilder;
500
501     stringBuilder.flexibleAppend("struct ", m_stageInStructName, " {\n");
502     for (auto& namedStageIn : m_namedStageIns) {
503         auto mangledTypeName = m_typeNamer.mangledNameForType(*m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].unnamedType);
504         auto elementName = namedStageIn.elementName;
505         auto attributeIndex = namedStageIn.attributeIndex;
506         stringBuilder.flexibleAppend("    ", mangledTypeName, ' ', elementName, " [[user(user", attributeIndex, ")]];\n");
507     }
508     stringBuilder.flexibleAppend(
509         "};\n\n"
510         "struct ", m_returnStructName, " {\n"
511     );
512     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
513         auto& outputItem = m_entryPointItems.outputs[i];
514         auto& internalTypeName = m_namedOutputs[i].internalTypeName;
515         auto elementName = m_namedOutputs[i].elementName;
516         auto attribute = attributeForSemantic(*outputItem.semantic);
517         stringBuilder.flexibleAppend("    ", internalTypeName, ' ', elementName, ' ', attribute, ";\n");
518     }
519     stringBuilder.flexibleAppend(
520         "};\n\n",
521         resourceHelperTypes()
522     );
523
524     return stringBuilder.toString();
525 }
526
527 String FragmentEntryPointScaffolding::signature(String& functionName)
528 {
529     StringBuilder stringBuilder;
530
531     stringBuilder.flexibleAppend("fragment ", m_returnStructName, ' ', functionName, '(', m_stageInStructName, ' ', m_stageInParameterName, " [[stage_in]]");
532     if (auto resourceSignature = this->resourceSignature())
533         stringBuilder.flexibleAppend(", ", *resourceSignature);
534     if (auto builtInsSignature = this->builtInsSignature())
535         stringBuilder.flexibleAppend(", ", *builtInsSignature);
536     stringBuilder.append(')');
537
538     return stringBuilder.toString();
539 }
540
541 String FragmentEntryPointScaffolding::unpack()
542 {
543     StringBuilder stringBuilder;
544
545     stringBuilder.append(unpackResourcesAndNamedBuiltIns());
546
547     for (auto& namedStageIn : m_namedStageIns) {
548         auto& path = m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].path;
549         auto& elementName = namedStageIn.elementName;
550         stringBuilder.flexibleAppend(mangledInputPath(path), " = ", m_stageInParameterName, '.', elementName, ";\n");
551     }
552
553     return stringBuilder.toString();
554 }
555
556 String FragmentEntryPointScaffolding::pack(const String& inputVariableName, const String& outputVariableName)
557 {
558     StringBuilder stringBuilder;
559
560     stringBuilder.flexibleAppend(m_returnStructName, ' ', outputVariableName, ";\n");
561     if (m_entryPointItems.outputs.size() == 1 && !m_entryPointItems.outputs[0].path.size()) {
562         auto& elementName = m_namedOutputs[0].elementName;
563         stringBuilder.flexibleAppend(outputVariableName, '.', elementName, " = ", inputVariableName, ";\n");
564         return stringBuilder.toString();
565     }
566     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
567         auto& elementName = m_namedOutputs[i].elementName;
568         auto& internalTypeName = m_namedOutputs[i].internalTypeName;
569         auto& path = m_entryPointItems.outputs[i].path;
570         stringBuilder.flexibleAppend(outputVariableName, '.', elementName, " = ", internalTypeName, '(', inputVariableName, mangledOutputPath(path), ");\n");
571     }
572     return stringBuilder.toString();
573 }
574
575 ComputeEntryPointScaffolding::ComputeEntryPointScaffolding(AST::FunctionDefinition& functionDefinition, Intrinsics& intrinsics, TypeNamer& typeNamer, EntryPointItems& entryPointItems, HashMap<Binding*, size_t>& resourceMap, Layout& layout, std::function<String()>&& generateNextVariableName)
576     : EntryPointScaffolding(functionDefinition, intrinsics, typeNamer, entryPointItems, resourceMap, layout, WTFMove(generateNextVariableName))
577 {
578 }
579
580 String ComputeEntryPointScaffolding::helperTypes()
581 {
582     return resourceHelperTypes();
583 }
584
585 String ComputeEntryPointScaffolding::signature(String& functionName)
586 {
587     StringBuilder stringBuilder;
588
589     stringBuilder.flexibleAppend("kernel void ", functionName, '(');
590     bool empty = true;
591     if (auto resourceSignature = this->resourceSignature()) {
592         empty = false;
593         stringBuilder.append(*resourceSignature);
594     }
595     if (auto builtInsSignature = this->builtInsSignature()) {
596         if (!empty)
597             stringBuilder.append(", ");
598         stringBuilder.append(*builtInsSignature);
599     }
600     stringBuilder.append(')');
601
602     return stringBuilder.toString();
603 }
604
605 String ComputeEntryPointScaffolding::unpack()
606 {
607     return unpackResourcesAndNamedBuiltIns();
608 }
609
610 String ComputeEntryPointScaffolding::pack(const String&, const String&)
611 {
612     ASSERT_NOT_REACHED();
613     return String();
614 }
615
616 }
617
618 }
619
620 }
621
622 #endif