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