Rename StringBuilder::flexibleAppend(...) to StringBuilder::append(...)
[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 void EntryPointScaffolding::emitResourceHelperTypes(StringBuilder& stringBuilder)
143 {
144     for (size_t i = 0; i < m_layout.size(); ++i) {
145         stringBuilder.append("struct ", m_namedBindGroups[i].structName, " {\n");
146         Vector<std::pair<unsigned, String>> structItems;
147         for (size_t j = 0; j < m_layout[i].bindings.size(); ++j) {
148             auto iterator = m_resourceMap.find(&m_layout[i].bindings[j]);
149             if (iterator == m_resourceMap.end())
150                 continue;
151             auto& type = m_entryPointItems.inputs[iterator->value].unnamedType->unifyNode();
152             if (is<AST::UnnamedType>(type) && is<AST::ReferenceType>(downcast<AST::UnnamedType>(type))) {
153                 auto& referenceType = downcast<AST::ReferenceType>(downcast<AST::UnnamedType>(type));
154                 auto mangledTypeName = m_typeNamer.mangledNameForType(referenceType.elementType());
155                 auto addressSpace = toString(referenceType.addressSpace());
156                 auto elementName = m_namedBindGroups[i].namedBindings[j].elementName;
157                 auto index = m_namedBindGroups[i].namedBindings[j].index;
158                 structItems.append(std::make_pair(index, makeString(addressSpace, " ", mangledTypeName, "* ", elementName, " [[id(", index, ")]];")));
159                 if (auto lengthInformation = m_namedBindGroups[i].namedBindings[j].lengthInformation)
160                     structItems.append(std::make_pair(lengthInformation->index, makeString("uint2 ", lengthInformation->elementName, " [[id(", lengthInformation->index, ")]];")));
161             } else if (is<AST::NamedType>(type) && is<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(type))) {
162                 auto& namedType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(type));
163                 auto mangledTypeName = m_typeNamer.mangledNameForType(namedType);
164                 auto elementName = m_namedBindGroups[i].namedBindings[j].elementName;
165                 auto index = m_namedBindGroups[i].namedBindings[j].index;
166                 structItems.append(std::make_pair(index, makeString(mangledTypeName, ' ', elementName, " [[id(", index, ")]];")));
167             }
168         }
169         std::sort(structItems.begin(), structItems.end(), [](const std::pair<unsigned, String>& left, const std::pair<unsigned, String>& right) {
170             return left.first < right.first;
171         });
172         for (const auto& structItem : structItems)
173             stringBuilder.append("    ", structItem.second, '\n');
174         stringBuilder.append("};\n\n");
175     }
176 }
177
178 bool EntryPointScaffolding::emitResourceSignature(StringBuilder& stringBuilder, IncludePrecedingComma includePrecedingComma)
179 {
180     if (!m_layout.size())
181         return false;
182
183     if (includePrecedingComma == IncludePrecedingComma::Yes)
184         stringBuilder.append(", ");
185
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.append("device ", namedBindGroup.structName, "& ", namedBindGroup.variableName, " [[buffer(", namedBindGroup.argumentBufferIndex, ")]]");
191     }
192     return true;
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 bool EntryPointScaffolding::emitBuiltInsSignature(StringBuilder& stringBuilder, IncludePrecedingComma includePrecedingComma)
231 {
232     if (!m_namedBuiltIns.size())
233         return false;
234
235     if (includePrecedingComma == IncludePrecedingComma::Yes)
236         stringBuilder.append(", ");
237
238     for (size_t i = 0; i < m_namedBuiltIns.size(); ++i) {
239         if (i)
240             stringBuilder.append(", ");
241         auto& namedBuiltIn = m_namedBuiltIns[i];
242         auto& item = m_entryPointItems.inputs[namedBuiltIn.indexInEntryPointItems];
243         auto& builtInSemantic = WTF::get<AST::BuiltInSemantic>(*item.semantic);
244         auto internalType = internalTypeForSemantic(builtInSemantic);
245         if (!internalType.isNull())
246             stringBuilder.append(internalType);
247         else
248             stringBuilder.append(m_typeNamer.mangledNameForType(*item.unnamedType));
249         stringBuilder.append(' ', namedBuiltIn.variableName, ' ', attributeForSemantic(builtInSemantic));
250     }
251     return true;
252 }
253
254 void EntryPointScaffolding::emitMangledInputPath(StringBuilder& stringBuilder, Vector<String>& path)
255 {
256     ASSERT(!path.isEmpty());
257     bool found = false;
258     AST::StructureDefinition* structureDefinition = nullptr;
259     for (size_t i = 0; i < m_functionDefinition.parameters().size(); ++i) {
260         if (m_functionDefinition.parameters()[i]->name() == path[0]) {
261             stringBuilder.append(m_parameterVariables[i]);
262             auto& unifyNode = m_functionDefinition.parameters()[i]->type()->unifyNode();
263             if (is<AST::NamedType>(unifyNode)) {
264                 auto& namedType = downcast<AST::NamedType>(unifyNode);
265                 if (is<AST::StructureDefinition>(namedType))
266                     structureDefinition = &downcast<AST::StructureDefinition>(namedType);
267             }
268             found = true;
269             break;
270         }
271     }
272     ASSERT(found);
273     for (size_t i = 1; i < path.size(); ++i) {
274         ASSERT(structureDefinition);
275         auto* next = structureDefinition->find(path[i]);
276         ASSERT(next);
277         stringBuilder.append('.', m_typeNamer.mangledNameForStructureElement(*next));
278         structureDefinition = nullptr;
279         auto& unifyNode = next->type().unifyNode();
280         if (is<AST::NamedType>(unifyNode)) {
281             auto& namedType = downcast<AST::NamedType>(unifyNode);
282             if (is<AST::StructureDefinition>(namedType))
283                 structureDefinition = &downcast<AST::StructureDefinition>(namedType);
284         }
285     }
286 }
287
288 void EntryPointScaffolding::emitMangledOutputPath(StringBuilder& stringBuilder, Vector<String>& path)
289 {
290     AST::StructureDefinition* structureDefinition = nullptr;
291     auto& unifyNode = m_functionDefinition.type().unifyNode();
292     structureDefinition = &downcast<AST::StructureDefinition>(downcast<AST::NamedType>(unifyNode));
293     for (auto& component : path) {
294         ASSERT(structureDefinition);
295         auto* next = structureDefinition->find(component);
296         ASSERT(next);
297         stringBuilder.append('.', m_typeNamer.mangledNameForStructureElement(*next));
298         structureDefinition = nullptr;
299         auto& unifyNode = next->type().unifyNode();
300         if (is<AST::NamedType>(unifyNode)) {
301             auto& namedType = downcast<AST::NamedType>(unifyNode);
302             if (is<AST::StructureDefinition>(namedType))
303                 structureDefinition = &downcast<AST::StructureDefinition>(namedType);
304         }
305     }
306 }
307
308 void EntryPointScaffolding::emitUnpackResourcesAndNamedBuiltIns(StringBuilder& stringBuilder)
309 {
310     for (size_t i = 0; i < m_functionDefinition.parameters().size(); ++i)
311         stringBuilder.append(m_typeNamer.mangledNameForType(*m_functionDefinition.parameters()[i]->type()), ' ', m_parameterVariables[i], ";\n");
312
313     for (size_t i = 0; i < m_layout.size(); ++i) {
314         auto variableName = m_namedBindGroups[i].variableName;
315         for (size_t j = 0; j < m_layout[i].bindings.size(); ++j) {
316             auto iterator = m_resourceMap.find(&m_layout[i].bindings[j]);
317             if (iterator == m_resourceMap.end())
318                 continue;
319             if (m_namedBindGroups[i].namedBindings[j].lengthInformation) {
320                 auto& path = m_entryPointItems.inputs[iterator->value].path;
321                 auto elementName = m_namedBindGroups[i].namedBindings[j].elementName;
322                 auto lengthElementName = m_namedBindGroups[i].namedBindings[j].lengthInformation->elementName;
323                 auto lengthTemporaryName = m_namedBindGroups[i].namedBindings[j].lengthInformation->temporaryName;
324
325                 auto& unnamedType = *m_entryPointItems.inputs[iterator->value].unnamedType;
326                 auto mangledTypeName = m_typeNamer.mangledNameForType(downcast<AST::ReferenceType>(unnamedType).elementType());
327
328                 stringBuilder.append(
329                     "size_t ", lengthTemporaryName, " = ", variableName, '.', lengthElementName, ".y;\n",
330                     lengthTemporaryName, " = ", lengthTemporaryName, " << 32;\n",
331                     lengthTemporaryName, " = ", lengthTemporaryName, " | ", variableName, '.', lengthElementName, ".x;\n",
332                     lengthTemporaryName, " = ", lengthTemporaryName, " / sizeof(", mangledTypeName, ");\n",
333                     "if (", lengthTemporaryName, " > 0xFFFFFFFF) ", lengthTemporaryName, " = 0xFFFFFFFF;\n"
334                 );
335                     
336                 emitMangledInputPath(stringBuilder, path);
337                 stringBuilder.append(
338                     " = { ", variableName, '.', elementName, ", static_cast<uint32_t>(", lengthTemporaryName, ") };\n"
339                 );
340             } else {
341                 auto& path = m_entryPointItems.inputs[iterator->value].path;
342                 auto elementName = m_namedBindGroups[i].namedBindings[j].elementName;
343                 
344                 emitMangledInputPath(stringBuilder, path);
345                 stringBuilder.append(" = ", 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
356         emitMangledInputPath(stringBuilder, path);
357         stringBuilder.append(" = ", mangledTypeName, '(', variableName, ");\n");
358     }
359 }
360
361 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)
362     : EntryPointScaffolding(functionDefinition, intrinsics, typeNamer, entryPointItems, resourceMap, layout, WTFMove(generateNextVariableName))
363     , m_matchedVertexAttributes(matchedVertexAttributes)
364     , m_stageInStructName(typeNamer.generateNextTypeName())
365     , m_returnStructName(typeNamer.generateNextTypeName())
366     , m_stageInParameterName(m_generateNextVariableName())
367 {
368     m_namedStageIns.reserveInitialCapacity(m_matchedVertexAttributes.size());
369     for (auto& keyValuePair : m_matchedVertexAttributes) {
370         NamedStageIn namedStageIn;
371         namedStageIn.indexInEntryPointItems = keyValuePair.value;
372         namedStageIn.elementName = m_typeNamer.generateNextStructureElementName();
373         namedStageIn.attributeIndex = keyValuePair.key->metalLocation;
374         m_namedStageIns.uncheckedAppend(WTFMove(namedStageIn));
375     }
376
377     m_namedOutputs.reserveInitialCapacity(m_entryPointItems.outputs.size());
378     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
379         auto& outputItem = m_entryPointItems.outputs[i];
380         NamedOutput namedOutput;
381         namedOutput.elementName = m_typeNamer.generateNextStructureElementName();
382         StringView internalType;
383         if (WTF::holds_alternative<AST::BuiltInSemantic>(*outputItem.semantic))
384             internalType = internalTypeForSemantic(WTF::get<AST::BuiltInSemantic>(*outputItem.semantic));
385         if (!internalType.isNull())
386             namedOutput.internalTypeName = internalType.toString();
387         else
388             namedOutput.internalTypeName = m_typeNamer.mangledNameForType(*outputItem.unnamedType);
389         m_namedOutputs.uncheckedAppend(WTFMove(namedOutput));
390     }
391 }
392
393 void VertexEntryPointScaffolding::emitHelperTypes(StringBuilder& stringBuilder)
394 {
395     stringBuilder.append("struct ", m_stageInStructName, " {\n");
396     for (auto& namedStageIn : m_namedStageIns) {
397         auto mangledTypeName = m_typeNamer.mangledNameForType(*m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].unnamedType);
398         auto elementName = namedStageIn.elementName;
399         auto attributeIndex = namedStageIn.attributeIndex;
400         stringBuilder.append("    ", mangledTypeName, ' ', elementName, " [[attribute(", attributeIndex, ")]];\n");
401     }
402     stringBuilder.append(
403         "};\n\n"
404         "struct ", m_returnStructName, " {\n"
405     );
406     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
407         auto& outputItem = m_entryPointItems.outputs[i];
408         auto& internalTypeName = m_namedOutputs[i].internalTypeName;
409         auto elementName = m_namedOutputs[i].elementName;
410         auto attribute = attributeForSemantic(*outputItem.semantic);
411         stringBuilder.append("    ", internalTypeName, ' ', elementName, ' ', attribute, ";\n");
412     }
413     stringBuilder.append(
414         "};\n\n"
415     );
416     
417     emitResourceHelperTypes(stringBuilder);
418 }
419
420 void VertexEntryPointScaffolding::emitSignature(StringBuilder& stringBuilder, MangledFunctionName functionName)
421 {
422     stringBuilder.append("vertex ", m_returnStructName, ' ', functionName, '(', m_stageInStructName, ' ', m_stageInParameterName, " [[stage_in]]");
423     emitResourceSignature(stringBuilder, IncludePrecedingComma::Yes);
424     emitBuiltInsSignature(stringBuilder, IncludePrecedingComma::Yes);
425     stringBuilder.append(')');
426 }
427
428 void VertexEntryPointScaffolding::emitUnpack(StringBuilder& stringBuilder)
429 {
430     emitUnpackResourcesAndNamedBuiltIns(stringBuilder);
431
432     for (auto& namedStageIn : m_namedStageIns) {
433         auto& path = m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].path;
434         auto& elementName = namedStageIn.elementName;
435         emitMangledInputPath(stringBuilder, path);
436         stringBuilder.append(" = ", m_stageInParameterName, '.', elementName, ";\n");
437     }
438 }
439
440 void VertexEntryPointScaffolding::emitPack(StringBuilder& stringBuilder, MangledVariableName inputVariableName, MangledVariableName outputVariableName)
441 {
442     stringBuilder.append(m_returnStructName, ' ', outputVariableName, ";\n");
443     if (m_entryPointItems.outputs.size() == 1 && !m_entryPointItems.outputs[0].path.size()) {
444         auto& elementName = m_namedOutputs[0].elementName;
445         stringBuilder.append(outputVariableName, '.', elementName, " = ", inputVariableName, ";\n");
446         return;
447     }
448     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
449         auto& elementName = m_namedOutputs[i].elementName;
450         auto& internalTypeName = m_namedOutputs[i].internalTypeName;
451         auto& path = m_entryPointItems.outputs[i].path;
452         stringBuilder.append(outputVariableName, '.', elementName, " = ", internalTypeName, '(', inputVariableName);
453         emitMangledOutputPath(stringBuilder, path);
454         stringBuilder.append(");\n");
455     }
456 }
457
458 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>&)
459     : EntryPointScaffolding(functionDefinition, intrinsics, typeNamer, entryPointItems, resourceMap, layout, WTFMove(generateNextVariableName))
460     , m_stageInStructName(typeNamer.generateNextTypeName())
461     , m_returnStructName(typeNamer.generateNextTypeName())
462     , m_stageInParameterName(m_generateNextVariableName())
463 {
464     for (size_t i = 0; i < m_entryPointItems.inputs.size(); ++i) {
465         auto& inputItem = m_entryPointItems.inputs[i];
466         if (!WTF::holds_alternative<AST::StageInOutSemantic>(*inputItem.semantic))
467             continue;
468         auto& stageInOutSemantic = WTF::get<AST::StageInOutSemantic>(*inputItem.semantic);
469         NamedStageIn namedStageIn;
470         namedStageIn.indexInEntryPointItems = i;
471         namedStageIn.elementName = m_typeNamer.generateNextStructureElementName();
472         namedStageIn.attributeIndex = stageInOutSemantic.index();
473         m_namedStageIns.append(WTFMove(namedStageIn));
474     }
475
476     m_namedOutputs.reserveInitialCapacity(m_entryPointItems.outputs.size());
477     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
478         auto& outputItem = m_entryPointItems.outputs[i];
479         NamedOutput namedOutput;
480         namedOutput.elementName = m_typeNamer.generateNextStructureElementName();
481         StringView internalType;
482         if (WTF::holds_alternative<AST::BuiltInSemantic>(*outputItem.semantic))
483             internalType = internalTypeForSemantic(WTF::get<AST::BuiltInSemantic>(*outputItem.semantic));
484         if (!internalType.isNull())
485             namedOutput.internalTypeName = internalType.toString();
486         else
487             namedOutput.internalTypeName = m_typeNamer.mangledNameForType(*outputItem.unnamedType);
488         m_namedOutputs.uncheckedAppend(WTFMove(namedOutput));
489     }
490 }
491
492 void FragmentEntryPointScaffolding::emitHelperTypes(StringBuilder& stringBuilder)
493 {
494     stringBuilder.append("struct ", m_stageInStructName, " {\n");
495     for (auto& namedStageIn : m_namedStageIns) {
496         auto mangledTypeName = m_typeNamer.mangledNameForType(*m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].unnamedType);
497         auto elementName = namedStageIn.elementName;
498         auto attributeIndex = namedStageIn.attributeIndex;
499         stringBuilder.append("    ", mangledTypeName, ' ', elementName, " [[user(user", attributeIndex, ")]];\n");
500     }
501     stringBuilder.append(
502         "};\n\n"
503         "struct ", m_returnStructName, " {\n"
504     );
505     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
506         auto& outputItem = m_entryPointItems.outputs[i];
507         auto& internalTypeName = m_namedOutputs[i].internalTypeName;
508         auto elementName = m_namedOutputs[i].elementName;
509         auto attribute = attributeForSemantic(*outputItem.semantic);
510         stringBuilder.append("    ", internalTypeName, ' ', elementName, ' ', attribute, ";\n");
511     }
512     stringBuilder.append(
513         "};\n\n"
514     );
515
516     emitResourceHelperTypes(stringBuilder);
517 }
518
519 void FragmentEntryPointScaffolding::emitSignature(StringBuilder& stringBuilder, MangledFunctionName functionName)
520 {
521     stringBuilder.append("fragment ", m_returnStructName, ' ', functionName, '(', m_stageInStructName, ' ', m_stageInParameterName, " [[stage_in]]");
522     emitResourceSignature(stringBuilder, IncludePrecedingComma::Yes);
523     emitBuiltInsSignature(stringBuilder, IncludePrecedingComma::Yes);
524     stringBuilder.append(')');
525 }
526
527 void FragmentEntryPointScaffolding::emitUnpack(StringBuilder& stringBuilder)
528 {
529     emitUnpackResourcesAndNamedBuiltIns(stringBuilder);
530
531     for (auto& namedStageIn : m_namedStageIns) {
532         auto& path = m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].path;
533         auto& elementName = namedStageIn.elementName;
534         emitMangledInputPath(stringBuilder, path);
535         stringBuilder.append(" = ", m_stageInParameterName, '.', elementName, ";\n");
536     }
537 }
538
539 void FragmentEntryPointScaffolding::emitPack(StringBuilder& stringBuilder, MangledVariableName inputVariableName, MangledVariableName outputVariableName)
540 {
541     stringBuilder.append(m_returnStructName, ' ', outputVariableName, ";\n");
542     if (m_entryPointItems.outputs.size() == 1 && !m_entryPointItems.outputs[0].path.size()) {
543         auto& elementName = m_namedOutputs[0].elementName;
544         stringBuilder.append(outputVariableName, '.', elementName, " = ", inputVariableName, ";\n");
545         return;
546     }
547     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
548         auto& elementName = m_namedOutputs[i].elementName;
549         auto& internalTypeName = m_namedOutputs[i].internalTypeName;
550         auto& path = m_entryPointItems.outputs[i].path;
551         stringBuilder.append(outputVariableName, '.', elementName, " = ", internalTypeName, '(', inputVariableName);
552         emitMangledOutputPath(stringBuilder, path);
553         stringBuilder.append(");\n");
554     }
555 }
556
557 ComputeEntryPointScaffolding::ComputeEntryPointScaffolding(AST::FunctionDefinition& functionDefinition, Intrinsics& intrinsics, TypeNamer& typeNamer, EntryPointItems& entryPointItems, HashMap<Binding*, size_t>& resourceMap, Layout& layout, std::function<MangledVariableName()>&& generateNextVariableName)
558     : EntryPointScaffolding(functionDefinition, intrinsics, typeNamer, entryPointItems, resourceMap, layout, WTFMove(generateNextVariableName))
559 {
560 }
561
562 void ComputeEntryPointScaffolding::emitHelperTypes(StringBuilder& stringBuilder)
563 {
564     emitResourceHelperTypes(stringBuilder);
565 }
566
567 void ComputeEntryPointScaffolding::emitSignature(StringBuilder& stringBuilder, MangledFunctionName functionName)
568 {
569     stringBuilder.append("kernel void ", functionName, '(');
570     bool addedToSignature = emitResourceSignature(stringBuilder, IncludePrecedingComma::No);
571     emitBuiltInsSignature(stringBuilder, addedToSignature ? IncludePrecedingComma::Yes : IncludePrecedingComma::No);
572     stringBuilder.append(')');
573 }
574
575 void ComputeEntryPointScaffolding::emitUnpack(StringBuilder& stringBuilder)
576 {
577     emitUnpackResourcesAndNamedBuiltIns(stringBuilder);
578 }
579
580 void ComputeEntryPointScaffolding::emitPack(StringBuilder&, MangledVariableName, MangledVariableName)
581 {
582     ASSERT_NOT_REACHED();
583 }
584
585 }
586
587 }
588
589 }
590
591 #endif