Mangled WHLSL names don't need to allocate Strings
[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<MangledVariableName()>&& 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 StringView internalTypeForSemantic(const AST::BuiltInSemantic& builtInSemantic)
196 {
197     switch (builtInSemantic.variable()) {
198     case AST::BuiltInSemantic::Variable::SVInstanceID:
199         return "uint";
200     case AST::BuiltInSemantic::Variable::SVVertexID:
201         return "uint";
202     case AST::BuiltInSemantic::Variable::PSize:
203         return "float";
204     case AST::BuiltInSemantic::Variable::SVPosition:
205         return "float4";
206     case AST::BuiltInSemantic::Variable::SVIsFrontFace:
207         return "bool";
208     case AST::BuiltInSemantic::Variable::SVSampleIndex:
209         return "uint";
210     case AST::BuiltInSemantic::Variable::SVInnerCoverage:
211         return "uint";
212     case AST::BuiltInSemantic::Variable::SVTarget:
213         return { };
214     case AST::BuiltInSemantic::Variable::SVDepth:
215         return "float";
216     case AST::BuiltInSemantic::Variable::SVCoverage:
217         return "uint";
218     case AST::BuiltInSemantic::Variable::SVDispatchThreadID:
219         return "uint3";
220     case AST::BuiltInSemantic::Variable::SVGroupID:
221         return "uint3";
222     case AST::BuiltInSemantic::Variable::SVGroupIndex:
223         return "uint";
224     default:
225         ASSERT(builtInSemantic.variable() == AST::BuiltInSemantic::Variable::SVGroupThreadID);
226         return "uint3";
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             stringBuilder.flexibleAppend(internalType);
245         else
246             stringBuilder.flexibleAppend(m_typeNamer.mangledNameForType(*item.unnamedType));
247         stringBuilder.flexibleAppend(' ', namedBuiltIn.variableName, ' ', attributeForSemantic(builtInSemantic));
248     }
249     return stringBuilder.toString();
250 }
251
252 String EntryPointScaffolding::mangledInputPath(Vector<String>& path)
253 {
254     ASSERT(!path.isEmpty());
255     StringBuilder stringBuilder;
256     bool found = false;
257     AST::StructureDefinition* structureDefinition = nullptr;
258     for (size_t i = 0; i < m_functionDefinition.parameters().size(); ++i) {
259         if (m_functionDefinition.parameters()[i]->name() == path[0]) {
260             stringBuilder.flexibleAppend(m_parameterVariables[i]);
261             auto& unifyNode = m_functionDefinition.parameters()[i]->type()->unifyNode();
262             if (is<AST::NamedType>(unifyNode)) {
263                 auto& namedType = downcast<AST::NamedType>(unifyNode);
264                 if (is<AST::StructureDefinition>(namedType))
265                     structureDefinition = &downcast<AST::StructureDefinition>(namedType);
266             }
267             found = true;
268             break;
269         }
270     }
271     ASSERT(found);
272     for (size_t i = 1; i < path.size(); ++i) {
273         ASSERT(structureDefinition);
274         auto* next = structureDefinition->find(path[i]);
275         ASSERT(next);
276         stringBuilder.flexibleAppend('.', m_typeNamer.mangledNameForStructureElement(*next));
277         structureDefinition = nullptr;
278         auto& unifyNode = next->type().unifyNode();
279         if (is<AST::NamedType>(unifyNode)) {
280             auto& namedType = downcast<AST::NamedType>(unifyNode);
281             if (is<AST::StructureDefinition>(namedType))
282                 structureDefinition = &downcast<AST::StructureDefinition>(namedType);
283         }
284     }
285
286     return stringBuilder.toString();
287 }
288
289 String EntryPointScaffolding::mangledOutputPath(Vector<String>& path)
290 {
291     StringBuilder stringBuilder;
292
293     AST::StructureDefinition* structureDefinition = nullptr;
294     auto& unifyNode = m_functionDefinition.type().unifyNode();
295     structureDefinition = &downcast<AST::StructureDefinition>(downcast<AST::NamedType>(unifyNode));
296     for (auto& component : path) {
297         ASSERT(structureDefinition);
298         auto* next = structureDefinition->find(component);
299         ASSERT(next);
300         stringBuilder.flexibleAppend('.', m_typeNamer.mangledNameForStructureElement(*next));
301         structureDefinition = nullptr;
302         auto& unifyNode = next->type().unifyNode();
303         if (is<AST::NamedType>(unifyNode)) {
304             auto& namedType = downcast<AST::NamedType>(unifyNode);
305             if (is<AST::StructureDefinition>(namedType))
306                 structureDefinition = &downcast<AST::StructureDefinition>(namedType);
307         }
308     }
309
310     return stringBuilder.toString();
311 }
312
313 String EntryPointScaffolding::unpackResourcesAndNamedBuiltIns()
314 {
315     StringBuilder stringBuilder;
316     for (size_t i = 0; i < m_functionDefinition.parameters().size(); ++i)
317         stringBuilder.flexibleAppend(m_typeNamer.mangledNameForType(*m_functionDefinition.parameters()[i]->type()), ' ', m_parameterVariables[i], ";\n");
318
319     for (size_t i = 0; i < m_layout.size(); ++i) {
320         auto variableName = m_namedBindGroups[i].variableName;
321         for (size_t j = 0; j < m_layout[i].bindings.size(); ++j) {
322             auto iterator = m_resourceMap.find(&m_layout[i].bindings[j]);
323             if (iterator == m_resourceMap.end())
324                 continue;
325             if (m_namedBindGroups[i].namedBindings[j].lengthInformation) {
326                 auto& path = m_entryPointItems.inputs[iterator->value].path;
327                 auto elementName = m_namedBindGroups[i].namedBindings[j].elementName;
328                 auto lengthElementName = m_namedBindGroups[i].namedBindings[j].lengthInformation->elementName;
329                 auto lengthTemporaryName = m_namedBindGroups[i].namedBindings[j].lengthInformation->temporaryName;
330
331                 auto& unnamedType = *m_entryPointItems.inputs[iterator->value].unnamedType;
332                 auto mangledTypeName = m_typeNamer.mangledNameForType(downcast<AST::ReferenceType>(unnamedType).elementType());
333
334                 stringBuilder.flexibleAppend(
335                     "size_t ", lengthTemporaryName, " = ", variableName, '.', lengthElementName, ".y;\n",
336                     lengthTemporaryName, " = ", lengthTemporaryName, " << 32;\n",
337                     lengthTemporaryName, " = ", lengthTemporaryName, " | ", variableName, '.', lengthElementName, ".x;\n",
338                     lengthTemporaryName, " = ", lengthTemporaryName, " / sizeof(", mangledTypeName, ");\n",
339                     "if (", lengthTemporaryName, " > 0xFFFFFFFF) ", lengthTemporaryName, " = 0xFFFFFFFF;\n",
340                     mangledInputPath(path), " = { ", variableName, '.', elementName, ", static_cast<uint32_t>(", lengthTemporaryName, ") };\n"
341                 );
342             } else {
343                 auto& path = m_entryPointItems.inputs[iterator->value].path;
344                 auto elementName = m_namedBindGroups[i].namedBindings[j].elementName;
345                 stringBuilder.flexibleAppend(mangledInputPath(path), " = ", variableName, '.', elementName, ";\n");
346             }
347         }
348     }
349
350     for (auto& namedBuiltIn : m_namedBuiltIns) {
351         auto& item = m_entryPointItems.inputs[namedBuiltIn.indexInEntryPointItems];
352         auto& path = item.path;
353         auto& variableName = namedBuiltIn.variableName;
354         auto mangledTypeName = m_typeNamer.mangledNameForType(*item.unnamedType);
355         stringBuilder.flexibleAppend(mangledInputPath(path), " = ", mangledTypeName, '(', variableName, ");\n");
356     }
357     return stringBuilder.toString();
358 }
359
360 VertexEntryPointScaffolding::VertexEntryPointScaffolding(AST::FunctionDefinition& functionDefinition, Intrinsics& intrinsics, TypeNamer& typeNamer, EntryPointItems& entryPointItems, HashMap<Binding*, size_t>& resourceMap, Layout& layout, std::function<MangledVariableName()>&& generateNextVariableName, HashMap<VertexAttribute*, size_t>& matchedVertexAttributes)
361     : EntryPointScaffolding(functionDefinition, intrinsics, typeNamer, entryPointItems, resourceMap, layout, WTFMove(generateNextVariableName))
362     , m_matchedVertexAttributes(matchedVertexAttributes)
363     , m_stageInStructName(typeNamer.generateNextTypeName())
364     , m_returnStructName(typeNamer.generateNextTypeName())
365     , m_stageInParameterName(m_generateNextVariableName())
366 {
367     m_namedStageIns.reserveInitialCapacity(m_matchedVertexAttributes.size());
368     for (auto& keyValuePair : m_matchedVertexAttributes) {
369         NamedStageIn namedStageIn;
370         namedStageIn.indexInEntryPointItems = keyValuePair.value;
371         namedStageIn.elementName = m_typeNamer.generateNextStructureElementName();
372         namedStageIn.attributeIndex = keyValuePair.key->metalLocation;
373         m_namedStageIns.uncheckedAppend(WTFMove(namedStageIn));
374     }
375
376     m_namedOutputs.reserveInitialCapacity(m_entryPointItems.outputs.size());
377     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
378         auto& outputItem = m_entryPointItems.outputs[i];
379         NamedOutput namedOutput;
380         namedOutput.elementName = m_typeNamer.generateNextStructureElementName();
381         StringView internalType;
382         if (WTF::holds_alternative<AST::BuiltInSemantic>(*outputItem.semantic))
383             internalType = internalTypeForSemantic(WTF::get<AST::BuiltInSemantic>(*outputItem.semantic));
384         if (!internalType.isNull())
385             namedOutput.internalTypeName = internalType.toString();
386         else
387             namedOutput.internalTypeName = m_typeNamer.mangledNameForType(*outputItem.unnamedType);
388         m_namedOutputs.uncheckedAppend(WTFMove(namedOutput));
389     }
390 }
391
392 String VertexEntryPointScaffolding::helperTypes()
393 {
394     StringBuilder stringBuilder;
395
396     stringBuilder.flexibleAppend("struct ", m_stageInStructName, " {\n");
397     for (auto& namedStageIn : m_namedStageIns) {
398         auto mangledTypeName = m_typeNamer.mangledNameForType(*m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].unnamedType);
399         auto elementName = namedStageIn.elementName;
400         auto attributeIndex = namedStageIn.attributeIndex;
401         stringBuilder.flexibleAppend("    ", mangledTypeName, ' ', elementName, " [[attribute(", attributeIndex, ")]];\n");
402     }
403     stringBuilder.flexibleAppend(
404         "};\n\n"
405         "struct ", m_returnStructName, " {\n"
406     );
407     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
408         auto& outputItem = m_entryPointItems.outputs[i];
409         auto& internalTypeName = m_namedOutputs[i].internalTypeName;
410         auto elementName = m_namedOutputs[i].elementName;
411         auto attribute = attributeForSemantic(*outputItem.semantic);
412         stringBuilder.flexibleAppend("    ", internalTypeName, ' ', elementName, ' ', attribute, ";\n");
413     }
414     stringBuilder.flexibleAppend(
415         "};\n\n",
416         resourceHelperTypes()
417     );
418
419     return stringBuilder.toString();
420 }
421
422 String VertexEntryPointScaffolding::signature(MangledFunctionName functionName)
423 {
424     StringBuilder stringBuilder;
425
426     stringBuilder.flexibleAppend("vertex ", m_returnStructName, ' ', functionName, '(', m_stageInStructName, ' ', m_stageInParameterName, " [[stage_in]]");
427     if (auto resourceSignature = this->resourceSignature())
428         stringBuilder.flexibleAppend(", ", *resourceSignature);
429     if (auto builtInsSignature = this->builtInsSignature())
430         stringBuilder.flexibleAppend(", ", *builtInsSignature);
431     stringBuilder.append(')');
432
433     return stringBuilder.toString();
434 }
435
436 String VertexEntryPointScaffolding::unpack()
437 {
438     StringBuilder stringBuilder;
439
440     stringBuilder.append(unpackResourcesAndNamedBuiltIns());
441
442     for (auto& namedStageIn : m_namedStageIns) {
443         auto& path = m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].path;
444         auto& elementName = namedStageIn.elementName;
445         stringBuilder.flexibleAppend(mangledInputPath(path), " = ", m_stageInParameterName, '.', elementName, ";\n");
446     }
447
448     return stringBuilder.toString();
449 }
450
451 String VertexEntryPointScaffolding::pack(MangledVariableName inputVariableName, MangledVariableName outputVariableName)
452 {
453     StringBuilder stringBuilder;
454
455     stringBuilder.flexibleAppend(m_returnStructName, ' ', outputVariableName, ";\n");
456     if (m_entryPointItems.outputs.size() == 1 && !m_entryPointItems.outputs[0].path.size()) {
457         auto& elementName = m_namedOutputs[0].elementName;
458         stringBuilder.flexibleAppend(outputVariableName, '.', elementName, " = ", inputVariableName, ";\n");
459         return stringBuilder.toString();
460     }
461     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
462         auto& elementName = m_namedOutputs[i].elementName;
463         auto& internalTypeName = m_namedOutputs[i].internalTypeName;
464         auto& path = m_entryPointItems.outputs[i].path;
465         stringBuilder.flexibleAppend(outputVariableName, '.', elementName, " = ", internalTypeName, '(', inputVariableName, mangledOutputPath(path), ");\n");
466     }
467     return stringBuilder.toString();
468 }
469
470 FragmentEntryPointScaffolding::FragmentEntryPointScaffolding(AST::FunctionDefinition& functionDefinition, Intrinsics& intrinsics, TypeNamer& typeNamer, EntryPointItems& entryPointItems, HashMap<Binding*, size_t>& resourceMap, Layout& layout, std::function<MangledVariableName()>&& generateNextVariableName, HashMap<AttachmentDescriptor*, size_t>&)
471     : EntryPointScaffolding(functionDefinition, intrinsics, typeNamer, entryPointItems, resourceMap, layout, WTFMove(generateNextVariableName))
472     , m_stageInStructName(typeNamer.generateNextTypeName())
473     , m_returnStructName(typeNamer.generateNextTypeName())
474     , m_stageInParameterName(m_generateNextVariableName())
475 {
476     for (size_t i = 0; i < m_entryPointItems.inputs.size(); ++i) {
477         auto& inputItem = m_entryPointItems.inputs[i];
478         if (!WTF::holds_alternative<AST::StageInOutSemantic>(*inputItem.semantic))
479             continue;
480         auto& stageInOutSemantic = WTF::get<AST::StageInOutSemantic>(*inputItem.semantic);
481         NamedStageIn namedStageIn;
482         namedStageIn.indexInEntryPointItems = i;
483         namedStageIn.elementName = m_typeNamer.generateNextStructureElementName();
484         namedStageIn.attributeIndex = stageInOutSemantic.index();
485         m_namedStageIns.append(WTFMove(namedStageIn));
486     }
487
488     m_namedOutputs.reserveInitialCapacity(m_entryPointItems.outputs.size());
489     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
490         auto& outputItem = m_entryPointItems.outputs[i];
491         NamedOutput namedOutput;
492         namedOutput.elementName = m_typeNamer.generateNextStructureElementName();
493         StringView internalType;
494         if (WTF::holds_alternative<AST::BuiltInSemantic>(*outputItem.semantic))
495             internalType = internalTypeForSemantic(WTF::get<AST::BuiltInSemantic>(*outputItem.semantic));
496         if (!internalType.isNull())
497             namedOutput.internalTypeName = internalType.toString();
498         else
499             namedOutput.internalTypeName = m_typeNamer.mangledNameForType(*outputItem.unnamedType);
500         m_namedOutputs.uncheckedAppend(WTFMove(namedOutput));
501     }
502 }
503
504 String FragmentEntryPointScaffolding::helperTypes()
505 {
506     StringBuilder stringBuilder;
507
508     stringBuilder.flexibleAppend("struct ", m_stageInStructName, " {\n");
509     for (auto& namedStageIn : m_namedStageIns) {
510         auto mangledTypeName = m_typeNamer.mangledNameForType(*m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].unnamedType);
511         auto elementName = namedStageIn.elementName;
512         auto attributeIndex = namedStageIn.attributeIndex;
513         stringBuilder.flexibleAppend("    ", mangledTypeName, ' ', elementName, " [[user(user", attributeIndex, ")]];\n");
514     }
515     stringBuilder.flexibleAppend(
516         "};\n\n"
517         "struct ", m_returnStructName, " {\n"
518     );
519     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
520         auto& outputItem = m_entryPointItems.outputs[i];
521         auto& internalTypeName = m_namedOutputs[i].internalTypeName;
522         auto elementName = m_namedOutputs[i].elementName;
523         auto attribute = attributeForSemantic(*outputItem.semantic);
524         stringBuilder.flexibleAppend("    ", internalTypeName, ' ', elementName, ' ', attribute, ";\n");
525     }
526     stringBuilder.flexibleAppend(
527         "};\n\n",
528         resourceHelperTypes()
529     );
530
531     return stringBuilder.toString();
532 }
533
534 String FragmentEntryPointScaffolding::signature(MangledFunctionName functionName)
535 {
536     StringBuilder stringBuilder;
537
538     stringBuilder.flexibleAppend("fragment ", m_returnStructName, ' ', functionName, '(', m_stageInStructName, ' ', m_stageInParameterName, " [[stage_in]]");
539     if (auto resourceSignature = this->resourceSignature())
540         stringBuilder.flexibleAppend(", ", *resourceSignature);
541     if (auto builtInsSignature = this->builtInsSignature())
542         stringBuilder.flexibleAppend(", ", *builtInsSignature);
543     stringBuilder.append(')');
544
545     return stringBuilder.toString();
546 }
547
548 String FragmentEntryPointScaffolding::unpack()
549 {
550     StringBuilder stringBuilder;
551
552     stringBuilder.append(unpackResourcesAndNamedBuiltIns());
553
554     for (auto& namedStageIn : m_namedStageIns) {
555         auto& path = m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].path;
556         auto& elementName = namedStageIn.elementName;
557         stringBuilder.flexibleAppend(mangledInputPath(path), " = ", m_stageInParameterName, '.', elementName, ";\n");
558     }
559
560     return stringBuilder.toString();
561 }
562
563 String FragmentEntryPointScaffolding::pack(MangledVariableName inputVariableName, MangledVariableName outputVariableName)
564 {
565     StringBuilder stringBuilder;
566
567     stringBuilder.flexibleAppend(m_returnStructName, ' ', outputVariableName, ";\n");
568     if (m_entryPointItems.outputs.size() == 1 && !m_entryPointItems.outputs[0].path.size()) {
569         auto& elementName = m_namedOutputs[0].elementName;
570         stringBuilder.flexibleAppend(outputVariableName, '.', elementName, " = ", inputVariableName, ";\n");
571         return stringBuilder.toString();
572     }
573     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
574         auto& elementName = m_namedOutputs[i].elementName;
575         auto& internalTypeName = m_namedOutputs[i].internalTypeName;
576         auto& path = m_entryPointItems.outputs[i].path;
577         stringBuilder.flexibleAppend(outputVariableName, '.', elementName, " = ", internalTypeName, '(', inputVariableName, mangledOutputPath(path), ");\n");
578     }
579     return stringBuilder.toString();
580 }
581
582 ComputeEntryPointScaffolding::ComputeEntryPointScaffolding(AST::FunctionDefinition& functionDefinition, Intrinsics& intrinsics, TypeNamer& typeNamer, EntryPointItems& entryPointItems, HashMap<Binding*, size_t>& resourceMap, Layout& layout, std::function<MangledVariableName()>&& generateNextVariableName)
583     : EntryPointScaffolding(functionDefinition, intrinsics, typeNamer, entryPointItems, resourceMap, layout, WTFMove(generateNextVariableName))
584 {
585 }
586
587 String ComputeEntryPointScaffolding::helperTypes()
588 {
589     return resourceHelperTypes();
590 }
591
592 String ComputeEntryPointScaffolding::signature(MangledFunctionName functionName)
593 {
594     StringBuilder stringBuilder;
595
596     stringBuilder.flexibleAppend("kernel void ", functionName, '(');
597     bool empty = true;
598     if (auto resourceSignature = this->resourceSignature()) {
599         empty = false;
600         stringBuilder.append(*resourceSignature);
601     }
602     if (auto builtInsSignature = this->builtInsSignature()) {
603         if (!empty)
604             stringBuilder.append(", ");
605         stringBuilder.append(*builtInsSignature);
606     }
607     stringBuilder.append(')');
608
609     return stringBuilder.toString();
610 }
611
612 String ComputeEntryPointScaffolding::unpack()
613 {
614     return unpackResourcesAndNamedBuiltIns();
615 }
616
617 String ComputeEntryPointScaffolding::pack(MangledVariableName, MangledVariableName)
618 {
619     ASSERT_NOT_REACHED();
620     return String();
621 }
622
623 }
624
625 }
626
627 }
628
629 #endif