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