1a392d50332a95ccb191e87c0dc579dea245d889
[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.append(makeString("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.append(makeString("    ", 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.append(makeString("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.append(makeString(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.append(makeString('.', 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.append(makeString('.', 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.append(makeString(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.append(makeString("size_t ", lengthTemporaryName, " = ", variableName, '.', lengthElementName, ".y;\n"));
334                 stringBuilder.append(makeString(lengthTemporaryName, " = ", lengthTemporaryName, " << 32;\n"));
335                 stringBuilder.append(makeString(lengthTemporaryName, " = ", lengthTemporaryName, " | ", variableName, '.', lengthElementName, ".x;\n"));
336                 stringBuilder.append(makeString(lengthTemporaryName, " = ", lengthTemporaryName, " / sizeof(", mangledTypeName, ");\n"));
337                 stringBuilder.append(makeString("if (", lengthTemporaryName, " > 0xFFFFFFFF) ", lengthTemporaryName, " = 0xFFFFFFFF;\n"));
338                 stringBuilder.append(makeString(mangledInputPath(path), " = { ", variableName, '.', elementName, ", static_cast<uint32_t>(", lengthTemporaryName, ") };\n"));
339             } else {
340                 auto& path = m_entryPointItems.inputs[iterator->value].path;
341                 auto elementName = m_namedBindGroups[i].namedBindings[j].elementName;
342                 stringBuilder.append(makeString(mangledInputPath(path), " = ", variableName, '.', elementName, ";\n"));
343             }
344         }
345     }
346
347     for (auto& namedBuiltIn : m_namedBuiltIns) {
348         auto& item = m_entryPointItems.inputs[namedBuiltIn.indexInEntryPointItems];
349         auto& path = item.path;
350         auto& variableName = namedBuiltIn.variableName;
351         auto mangledTypeName = m_typeNamer.mangledNameForType(*item.unnamedType);
352         stringBuilder.append(makeString(mangledInputPath(path), " = ", mangledTypeName, '(', variableName, ");\n"));
353     }
354     return stringBuilder.toString();
355 }
356
357 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)
358     : EntryPointScaffolding(functionDefinition, intrinsics, typeNamer, entryPointItems, resourceMap, layout, WTFMove(generateNextVariableName))
359     , m_matchedVertexAttributes(matchedVertexAttributes)
360     , m_stageInStructName(typeNamer.generateNextTypeName())
361     , m_returnStructName(typeNamer.generateNextTypeName())
362     , m_stageInParameterName(m_generateNextVariableName())
363 {
364     m_namedStageIns.reserveInitialCapacity(m_matchedVertexAttributes.size());
365     for (auto& keyValuePair : m_matchedVertexAttributes) {
366         NamedStageIn namedStageIn;
367         namedStageIn.indexInEntryPointItems = keyValuePair.value;
368         namedStageIn.elementName = m_typeNamer.generateNextStructureElementName();
369         namedStageIn.attributeIndex = keyValuePair.key->metalLocation;
370         m_namedStageIns.uncheckedAppend(WTFMove(namedStageIn));
371     }
372
373     m_namedOutputs.reserveInitialCapacity(m_entryPointItems.outputs.size());
374     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
375         auto& outputItem = m_entryPointItems.outputs[i];
376         NamedOutput namedOutput;
377         namedOutput.elementName = m_typeNamer.generateNextStructureElementName();
378         if (WTF::holds_alternative<AST::BuiltInSemantic>(*outputItem.semantic))
379             namedOutput.internalTypeName = internalTypeForSemantic(WTF::get<AST::BuiltInSemantic>(*outputItem.semantic));
380         if (namedOutput.internalTypeName.isNull())
381             namedOutput.internalTypeName = m_typeNamer.mangledNameForType(*outputItem.unnamedType);
382         m_namedOutputs.uncheckedAppend(WTFMove(namedOutput));
383     }
384 }
385
386 String VertexEntryPointScaffolding::helperTypes()
387 {
388     StringBuilder stringBuilder;
389
390     stringBuilder.append(makeString("struct ", m_stageInStructName, " {\n"));
391     for (auto& namedStageIn : m_namedStageIns) {
392         auto mangledTypeName = m_typeNamer.mangledNameForType(*m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].unnamedType);
393         auto elementName = namedStageIn.elementName;
394         auto attributeIndex = namedStageIn.attributeIndex;
395         stringBuilder.append(makeString("    ", mangledTypeName, ' ', elementName, " [[attribute(", attributeIndex, ")]];\n"));
396     }
397     stringBuilder.append("};\n\n");
398
399     stringBuilder.append(makeString("struct ", m_returnStructName, " {\n"));
400     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
401         auto& outputItem = m_entryPointItems.outputs[i];
402         auto& internalTypeName = m_namedOutputs[i].internalTypeName;
403         auto elementName = m_namedOutputs[i].elementName;
404         auto attribute = attributeForSemantic(*outputItem.semantic);
405         stringBuilder.append(makeString("    ", internalTypeName, ' ', elementName, ' ', attribute, ";\n"));
406     }
407     stringBuilder.append("};\n\n");
408
409     stringBuilder.append(resourceHelperTypes());
410
411     return stringBuilder.toString();
412 }
413
414 String VertexEntryPointScaffolding::signature(String& functionName)
415 {
416     StringBuilder stringBuilder;
417
418     stringBuilder.append(makeString("vertex ", m_returnStructName, ' ', functionName, '(', m_stageInStructName, ' ', m_stageInParameterName, " [[stage_in]]"));
419     if (auto resourceSignature = this->resourceSignature())
420         stringBuilder.append(makeString(", ", *resourceSignature));
421     if (auto builtInsSignature = this->builtInsSignature())
422         stringBuilder.append(makeString(", ", *builtInsSignature));
423     stringBuilder.append(")");
424
425     return stringBuilder.toString();
426 }
427
428 String VertexEntryPointScaffolding::unpack()
429 {
430     StringBuilder stringBuilder;
431
432     stringBuilder.append(unpackResourcesAndNamedBuiltIns());
433
434     for (auto& namedStageIn : m_namedStageIns) {
435         auto& path = m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].path;
436         auto& elementName = namedStageIn.elementName;
437         stringBuilder.append(makeString(mangledInputPath(path), " = ", m_stageInParameterName, '.', elementName, ";\n"));
438     }
439
440     return stringBuilder.toString();
441 }
442
443 String VertexEntryPointScaffolding::pack(const String& inputVariableName, const String& outputVariableName)
444 {
445     StringBuilder stringBuilder;
446     stringBuilder.append(makeString(m_returnStructName, ' ', outputVariableName, ";\n"));
447     if (m_entryPointItems.outputs.size() == 1 && !m_entryPointItems.outputs[0].path.size()) {
448         auto& elementName = m_namedOutputs[0].elementName;
449         stringBuilder.append(makeString(outputVariableName, '.', elementName, " = ", inputVariableName, ";\n"));
450         return stringBuilder.toString();
451     }
452     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
453         auto& elementName = m_namedOutputs[i].elementName;
454         auto& internalTypeName = m_namedOutputs[i].internalTypeName;
455         auto& path = m_entryPointItems.outputs[i].path;
456         stringBuilder.append(makeString(outputVariableName, '.', elementName, " = ", internalTypeName, '(', inputVariableName, mangledOutputPath(path), ");\n"));
457     }
458     return stringBuilder.toString();
459 }
460
461 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>&)
462     : EntryPointScaffolding(functionDefinition, intrinsics, typeNamer, entryPointItems, resourceMap, layout, WTFMove(generateNextVariableName))
463     , m_stageInStructName(typeNamer.generateNextTypeName())
464     , m_returnStructName(typeNamer.generateNextTypeName())
465     , m_stageInParameterName(m_generateNextVariableName())
466 {
467     for (size_t i = 0; i < m_entryPointItems.inputs.size(); ++i) {
468         auto& inputItem = m_entryPointItems.inputs[i];
469         if (!WTF::holds_alternative<AST::StageInOutSemantic>(*inputItem.semantic))
470             continue;
471         auto& stageInOutSemantic = WTF::get<AST::StageInOutSemantic>(*inputItem.semantic);
472         NamedStageIn namedStageIn;
473         namedStageIn.indexInEntryPointItems = i;
474         namedStageIn.elementName = m_typeNamer.generateNextStructureElementName();
475         namedStageIn.attributeIndex = stageInOutSemantic.index();
476         m_namedStageIns.append(WTFMove(namedStageIn));
477     }
478
479     m_namedOutputs.reserveInitialCapacity(m_entryPointItems.outputs.size());
480     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
481         auto& outputItem = m_entryPointItems.outputs[i];
482         NamedOutput namedOutput;
483         namedOutput.elementName = m_typeNamer.generateNextStructureElementName();
484         if (WTF::holds_alternative<AST::BuiltInSemantic>(*outputItem.semantic))
485             namedOutput.internalTypeName = internalTypeForSemantic(WTF::get<AST::BuiltInSemantic>(*outputItem.semantic));
486         if (namedOutput.internalTypeName.isNull())
487             namedOutput.internalTypeName = m_typeNamer.mangledNameForType(*outputItem.unnamedType);
488         m_namedOutputs.uncheckedAppend(WTFMove(namedOutput));
489     }
490 }
491
492 String FragmentEntryPointScaffolding::helperTypes()
493 {
494     StringBuilder stringBuilder;
495
496     stringBuilder.append(makeString("struct ", m_stageInStructName, " {\n"));
497     for (auto& namedStageIn : m_namedStageIns) {
498         auto mangledTypeName = m_typeNamer.mangledNameForType(*m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].unnamedType);
499         auto elementName = namedStageIn.elementName;
500         auto attributeIndex = namedStageIn.attributeIndex;
501         stringBuilder.append(makeString("    ", mangledTypeName, ' ', elementName, " [[user(user", attributeIndex, ")]];\n"));
502     }
503     stringBuilder.append("};\n\n");
504
505     stringBuilder.append(makeString("struct ", m_returnStructName, " {\n"));
506     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
507         auto& outputItem = m_entryPointItems.outputs[i];
508         auto& internalTypeName = m_namedOutputs[i].internalTypeName;
509         auto elementName = m_namedOutputs[i].elementName;
510         auto attribute = attributeForSemantic(*outputItem.semantic);
511         stringBuilder.append(makeString("    ", internalTypeName, ' ', elementName, ' ', attribute, ";\n"));
512     }
513     stringBuilder.append("};\n\n");
514
515     stringBuilder.append(resourceHelperTypes());
516
517     return stringBuilder.toString();
518 }
519
520 String FragmentEntryPointScaffolding::signature(String& functionName)
521 {
522     StringBuilder stringBuilder;
523
524     stringBuilder.append(makeString("fragment ", m_returnStructName, ' ', functionName, '(', m_stageInStructName, ' ', m_stageInParameterName, " [[stage_in]]"));
525     if (auto resourceSignature = this->resourceSignature())
526         stringBuilder.append(makeString(", ", *resourceSignature));
527     if (auto builtInsSignature = this->builtInsSignature())
528         stringBuilder.append(makeString(", ", *builtInsSignature));
529     stringBuilder.append(")");
530
531     return stringBuilder.toString();
532 }
533
534 String FragmentEntryPointScaffolding::unpack()
535 {
536     StringBuilder stringBuilder;
537
538     stringBuilder.append(unpackResourcesAndNamedBuiltIns());
539
540     for (auto& namedStageIn : m_namedStageIns) {
541         auto& path = m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].path;
542         auto& elementName = namedStageIn.elementName;
543         stringBuilder.append(makeString(mangledInputPath(path), " = ", m_stageInParameterName, '.', elementName, ";\n"));
544     }
545
546     return stringBuilder.toString();
547 }
548
549 String FragmentEntryPointScaffolding::pack(const String& inputVariableName, const String& outputVariableName)
550 {
551     StringBuilder stringBuilder;
552     stringBuilder.append(makeString(m_returnStructName, ' ', outputVariableName, ";\n"));
553     if (m_entryPointItems.outputs.size() == 1 && !m_entryPointItems.outputs[0].path.size()) {
554         auto& elementName = m_namedOutputs[0].elementName;
555         stringBuilder.append(makeString(outputVariableName, '.', elementName, " = ", inputVariableName, ";\n"));
556         return stringBuilder.toString();
557     }
558     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
559         auto& elementName = m_namedOutputs[i].elementName;
560         auto& internalTypeName = m_namedOutputs[i].internalTypeName;
561         auto& path = m_entryPointItems.outputs[i].path;
562         stringBuilder.append(makeString(outputVariableName, '.', elementName, " = ", internalTypeName, '(', inputVariableName, mangledOutputPath(path), ");\n"));
563     }
564     return stringBuilder.toString();
565 }
566
567 ComputeEntryPointScaffolding::ComputeEntryPointScaffolding(AST::FunctionDefinition& functionDefinition, Intrinsics& intrinsics, TypeNamer& typeNamer, EntryPointItems& entryPointItems, HashMap<Binding*, size_t>& resourceMap, Layout& layout, std::function<String()>&& generateNextVariableName)
568     : EntryPointScaffolding(functionDefinition, intrinsics, typeNamer, entryPointItems, resourceMap, layout, WTFMove(generateNextVariableName))
569 {
570 }
571
572 String ComputeEntryPointScaffolding::helperTypes()
573 {
574     return resourceHelperTypes();
575 }
576
577 String ComputeEntryPointScaffolding::signature(String& functionName)
578 {
579     StringBuilder stringBuilder;
580
581     stringBuilder.append(makeString("kernel void ", functionName, '('));
582     bool empty = true;
583     if (auto resourceSignature = this->resourceSignature()) {
584         empty = false;
585         stringBuilder.append(makeString(*resourceSignature));
586     }
587     if (auto builtInsSignature = this->builtInsSignature()) {
588         if (!empty)
589             stringBuilder.append(", ");
590         stringBuilder.append(*builtInsSignature);
591     }
592     stringBuilder.append(")");
593
594     return stringBuilder.toString();
595 }
596
597 String ComputeEntryPointScaffolding::unpack()
598 {
599     return unpackResourcesAndNamedBuiltIns();
600 }
601
602 String ComputeEntryPointScaffolding::pack(const String&, const String&)
603 {
604     ASSERT_NOT_REACHED();
605     return String();
606 }
607
608 }
609
610 }
611
612 }
613
614 #endif