Add Optional to Forward.h.
[WebKit.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 "WHLSLResourceSemantic.h"
36 #include "WHLSLStageInOutSemantic.h"
37 #include "WHLSLStructureDefinition.h"
38 #include "WHLSLTypeNamer.h"
39 #include <wtf/Optional.h>
40 #include <wtf/text/StringBuilder.h>
41 #include <wtf/text/StringConcatenateNumbers.h>
42
43 namespace WebCore {
44
45 namespace WHLSL {
46
47 namespace Metal {
48
49 static String attributeForSemantic(AST::BuiltInSemantic& builtInSemantic)
50 {
51     switch (builtInSemantic.variable()) {
52     case AST::BuiltInSemantic::Variable::SVInstanceID:
53         return "[[instance_id]]"_str;
54     case AST::BuiltInSemantic::Variable::SVVertexID:
55         return "[[vertex_id]]"_str;
56     case AST::BuiltInSemantic::Variable::PSize:
57         return "[[point_size]]"_str;
58     case AST::BuiltInSemantic::Variable::SVPosition:
59         return "[[position]]"_str;
60     case AST::BuiltInSemantic::Variable::SVIsFrontFace:
61         return "[[front_facing]]"_str;
62     case AST::BuiltInSemantic::Variable::SVSampleIndex:
63         return "[[sample_id]]"_str;
64     case AST::BuiltInSemantic::Variable::SVInnerCoverage:
65         return "[[sample_mask]]"_str;
66     case AST::BuiltInSemantic::Variable::SVTarget:
67         return makeString("[[color(", *builtInSemantic.targetIndex(), ")]]");
68     case AST::BuiltInSemantic::Variable::SVDepth:
69         return "[[depth(any)]]"_str;
70     case AST::BuiltInSemantic::Variable::SVCoverage:
71         return "[[sample_mask]]"_str;
72     case AST::BuiltInSemantic::Variable::SVDispatchThreadID:
73         return "[[thread_position_in_grid]]"_str;
74     case AST::BuiltInSemantic::Variable::SVGroupID:
75         return "[[threadgroup_position_in_grid]]"_str;
76     case AST::BuiltInSemantic::Variable::SVGroupIndex:
77         return "[[thread_index_in_threadgroup]]"_str;
78     default:
79         ASSERT(builtInSemantic.variable() == AST::BuiltInSemantic::Variable::SVGroupThreadID);
80         return "[[thread_position_in_threadgroup]]"_str;
81     }
82 }
83
84 static String attributeForSemantic(AST::Semantic& semantic)
85 {
86     if (WTF::holds_alternative<AST::BuiltInSemantic>(semantic))
87         return attributeForSemantic(WTF::get<AST::BuiltInSemantic>(semantic));
88     auto& stageInOutSemantic = WTF::get<AST::StageInOutSemantic>(semantic);
89     return makeString("[[user(", stageInOutSemantic.index(), ")]]");
90 }
91
92 EntryPointScaffolding::EntryPointScaffolding(AST::FunctionDefinition& functionDefinition, Intrinsics& intrinsics, TypeNamer& typeNamer, EntryPointItems& entryPointItems, HashMap<Binding*, size_t>& resourceMap, Layout& layout, std::function<String()>&& generateNextVariableName)
93     : m_functionDefinition(functionDefinition)
94     , m_intrinsics(intrinsics)
95     , m_typeNamer(typeNamer)
96     , m_entryPointItems(entryPointItems)
97     , m_resourceMap(resourceMap)
98     , m_layout(layout)
99     , m_generateNextVariableName(generateNextVariableName)
100 {
101     unsigned argumentBufferIndex = 0;
102     m_namedBindGroups.reserveInitialCapacity(m_layout.size());
103     for (size_t i = 0; i < m_layout.size(); ++i) {
104         NamedBindGroup namedBindGroup;
105         namedBindGroup.structName = m_typeNamer.generateNextTypeName();
106         namedBindGroup.variableName = m_generateNextVariableName();
107         namedBindGroup.argumentBufferIndex = argumentBufferIndex++;
108         namedBindGroup.namedBindings.reserveInitialCapacity(m_layout[i].bindings.size());
109         unsigned index = 0;
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 = index++;
114             namedBindGroup.namedBindings.uncheckedAppend(WTFMove(namedBinding));
115         }
116         m_namedBindGroups.uncheckedAppend(WTFMove(namedBindGroup));
117     }
118
119     for (size_t i = 0; i < m_entryPointItems.inputs.size(); ++i) {
120         if (!WTF::holds_alternative<AST::BuiltInSemantic>(*m_entryPointItems.inputs[i].semantic))
121             continue;
122         NamedBuiltIn namedBuiltIn;
123         namedBuiltIn.indexInEntryPointItems = i;
124         namedBuiltIn.variableName = m_generateNextVariableName();
125         m_namedBuiltIns.append(WTFMove(namedBuiltIn));
126     }
127
128     m_parameterVariables.reserveInitialCapacity(m_functionDefinition.parameters().size());
129     for (size_t i = 0; i < m_functionDefinition.parameters().size(); ++i)
130         m_parameterVariables.uncheckedAppend(m_generateNextVariableName());
131 }
132
133 MappedBindGroups EntryPointScaffolding::mappedBindGroups() const
134 {
135     MappedBindGroups result;
136     result.reserveInitialCapacity(m_layout.size());
137     for (auto& namedBindGroup : m_namedBindGroups) {
138         MappedBindGroup mappedBindGroup;
139         mappedBindGroup.argumentBufferIndex = namedBindGroup.argumentBufferIndex;
140         mappedBindGroup.bindingIndices.reserveInitialCapacity(namedBindGroup.namedBindings.size());
141         for (auto& namedBinding : namedBindGroup.namedBindings)
142             mappedBindGroup.bindingIndices.uncheckedAppend(namedBinding.index);
143         result.uncheckedAppend(WTFMove(mappedBindGroup));
144     }
145     return result;
146 }
147
148 String EntryPointScaffolding::resourceHelperTypes()
149 {
150     StringBuilder stringBuilder;
151     for (size_t i = 0; i < m_layout.size(); ++i) {
152         stringBuilder.append(makeString("struct ", m_namedBindGroups[i].structName, " {\n"));
153         for (size_t j = 0; j < m_layout[i].bindings.size(); ++j) {
154             auto iterator = m_resourceMap.find(&m_layout[i].bindings[j]);
155             if (iterator == m_resourceMap.end())
156                 continue;
157             auto mangledTypeName = m_typeNamer.mangledNameForType(*m_entryPointItems.inputs[iterator->value].unnamedType);
158             auto elementName = m_namedBindGroups[i].namedBindings[j].elementName;
159             auto index = m_namedBindGroups[i].namedBindings[j].index;
160             stringBuilder.append(makeString("    ", mangledTypeName, ' ', elementName, " [[id(", index, ")]];\n"));
161         }
162         stringBuilder.append("}\n\n");
163     }
164     return stringBuilder.toString();
165 }
166
167 Optional<String> EntryPointScaffolding::resourceSignature()
168 {
169     if (!m_layout.size())
170         return WTF::nullopt;
171
172     StringBuilder stringBuilder;
173     for (size_t i = 0; i < m_layout.size(); ++i) {
174         if (i)
175             stringBuilder.append(", ");
176         auto& namedBindGroup = m_namedBindGroups[i];
177         stringBuilder.append(makeString(namedBindGroup.structName, "& ", namedBindGroup.variableName, " [[buffer(", namedBindGroup.argumentBufferIndex, ")]]"));
178     }
179     return stringBuilder.toString();
180 }
181
182 Optional<String> EntryPointScaffolding::builtInsSignature()
183 {
184     if (!m_namedBuiltIns.size())
185         return WTF::nullopt;
186
187     StringBuilder stringBuilder;
188     for (size_t i = 0; i < m_namedBuiltIns.size(); ++i) {
189         if (i)
190             stringBuilder.append(", ");
191         auto& namedBuiltIn = m_namedBuiltIns[i];
192         auto& item = m_entryPointItems.inputs[namedBuiltIn.indexInEntryPointItems];
193         auto& builtInSemantic = WTF::get<AST::BuiltInSemantic>(*item.semantic);
194         auto mangledTypeName = m_typeNamer.mangledNameForType(*item.unnamedType);
195         auto variableName = namedBuiltIn.variableName;
196         stringBuilder.append(makeString(mangledTypeName, ' ', variableName, ' ', attributeForSemantic(builtInSemantic)));
197     }
198     return stringBuilder.toString();
199 }
200
201 String EntryPointScaffolding::mangledInputPath(Vector<String>& path)
202 {
203     ASSERT(!path.isEmpty());
204     StringBuilder stringBuilder;
205     bool found = false;
206     AST::StructureDefinition* structureDefinition = nullptr;
207     for (size_t i = 0; i < m_functionDefinition.parameters().size(); ++i) {
208         if (m_functionDefinition.parameters()[i].name() == path[0]) {
209             stringBuilder.append(m_parameterVariables[i]);
210             auto& unifyNode = m_functionDefinition.parameters()[i].type()->unifyNode();
211             if (is<AST::NamedType>(unifyNode)) {
212                 auto& namedType = downcast<AST::NamedType>(unifyNode);
213                 if (is<AST::StructureDefinition>(namedType))
214                     structureDefinition = &downcast<AST::StructureDefinition>(namedType);
215             }
216             found = true;
217             break;
218         }
219     }
220     ASSERT(found);
221     for (size_t i = 1; i < path.size(); ++i) {
222         ASSERT(structureDefinition);
223         auto* next = structureDefinition->find(path[i]);
224         ASSERT(next);
225         stringBuilder.append(makeString('.', m_typeNamer.mangledNameForStructureElement(*next)));
226         structureDefinition = nullptr;
227         auto& unifyNode = next->type().unifyNode();
228         if (is<AST::NamedType>(unifyNode)) {
229             auto& namedType = downcast<AST::NamedType>(unifyNode);
230             if (is<AST::StructureDefinition>(namedType))
231                 structureDefinition = &downcast<AST::StructureDefinition>(namedType);
232         }
233     }
234
235     return stringBuilder.toString();
236 }
237
238 String EntryPointScaffolding::mangledOutputPath(Vector<String>& path)
239 {
240     StringBuilder stringBuilder;
241
242     AST::StructureDefinition* structureDefinition = nullptr;
243     auto& unifyNode = m_functionDefinition.type().unifyNode();
244     ASSERT(is<AST::NamedType>(unifyNode));
245     auto& namedType = downcast<AST::NamedType>(unifyNode);
246     ASSERT(is<AST::StructureDefinition>(namedType));
247     structureDefinition = &downcast<AST::StructureDefinition>(namedType);
248     for (auto& component : path) {
249         ASSERT(structureDefinition);
250         auto* next = structureDefinition->find(component);
251         ASSERT(next);
252         stringBuilder.append(makeString('.', m_typeNamer.mangledNameForStructureElement(*next)));
253         structureDefinition = nullptr;
254         auto& unifyNode = next->type().unifyNode();
255         if (is<AST::NamedType>(unifyNode)) {
256             auto& namedType = downcast<AST::NamedType>(unifyNode);
257             if (is<AST::StructureDefinition>(namedType))
258                 structureDefinition = &downcast<AST::StructureDefinition>(namedType);
259         }
260     }
261
262     return stringBuilder.toString();
263 }
264
265 String EntryPointScaffolding::unpackResourcesAndNamedBuiltIns()
266 {
267     StringBuilder stringBuilder;
268     for (size_t i = 0; i < m_functionDefinition.parameters().size(); ++i)
269         stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*m_functionDefinition.parameters()[i].type()), ' ', m_parameterVariables[i], ";\n"));
270
271     for (size_t i = 0; i < m_layout.size(); ++i) {
272         auto variableName = m_namedBindGroups[i].variableName;
273         for (size_t j = 0; j < m_layout[i].bindings.size(); ++j) {
274             auto iterator = m_resourceMap.find(&m_layout[i].bindings[j]);
275             if (iterator == m_resourceMap.end())
276                 continue;
277             auto& path = m_entryPointItems.inputs[iterator->value].path;
278             auto elementName = m_namedBindGroups[i].namedBindings[j].elementName;
279             stringBuilder.append(makeString(mangledInputPath(path), " = ", variableName, '.', elementName, ";\n"));
280         }
281     }
282
283     for (auto& namedBuiltIn : m_namedBuiltIns) {
284         auto& path = m_entryPointItems.inputs[namedBuiltIn.indexInEntryPointItems].path;
285         auto& variableName = namedBuiltIn.variableName;
286         stringBuilder.append(makeString(mangledInputPath(path), " = ", variableName, ";\n"));
287     }
288     return stringBuilder.toString();
289 }
290
291 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)
292     : EntryPointScaffolding(functionDefinition, intrinsics, typeNamer, entryPointItems, resourceMap, layout, WTFMove(generateNextVariableName))
293     , m_matchedVertexAttributes(matchedVertexAttributes)
294     , m_stageInStructName(typeNamer.generateNextTypeName())
295     , m_returnStructName(typeNamer.generateNextTypeName())
296     , m_stageInParameterName(m_generateNextVariableName())
297 {
298     m_namedStageIns.reserveInitialCapacity(m_matchedVertexAttributes.size());
299     for (auto& keyValuePair : m_matchedVertexAttributes) {
300         NamedStageIn namedStageIn;
301         namedStageIn.indexInEntryPointItems = keyValuePair.value;
302         namedStageIn.elementName = m_typeNamer.generateNextStructureElementName();
303         namedStageIn.attributeIndex = keyValuePair.key->name;
304         m_namedStageIns.uncheckedAppend(WTFMove(namedStageIn));
305     }
306
307     m_namedOutputs.reserveInitialCapacity(m_entryPointItems.outputs.size());
308     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
309         NamedOutput namedOutput;
310         namedOutput.elementName = m_typeNamer.generateNextStructureElementName();
311         m_namedOutputs.uncheckedAppend(WTFMove(namedOutput));
312     }
313 }
314
315 String VertexEntryPointScaffolding::helperTypes()
316 {
317     StringBuilder stringBuilder;
318
319     stringBuilder.append(makeString("struct ", m_stageInStructName, " {\n"));
320     for (auto& namedStageIn : m_namedStageIns) {
321         auto mangledTypeName = m_typeNamer.mangledNameForType(*m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].unnamedType);
322         auto elementName = namedStageIn.elementName;
323         auto attributeIndex = namedStageIn.elementName;
324         stringBuilder.append(makeString("    ", mangledTypeName, ' ', elementName, " [[attribute(", attributeIndex, ")]];\n"));
325     }
326     stringBuilder.append("}\n\n");
327
328     stringBuilder.append(makeString("struct ", m_returnStructName, " {\n"));
329     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
330         auto& outputItem = m_entryPointItems.outputs[i];
331         auto mangledTypeName = m_typeNamer.mangledNameForType(*outputItem.unnamedType);
332         auto elementName = m_namedOutputs[i].elementName;
333         auto attribute = attributeForSemantic(*outputItem.semantic);
334         stringBuilder.append(makeString("    ", mangledTypeName, ' ', elementName, ' ', attribute, ";\n"));
335     }
336     stringBuilder.append("}\n\n");
337
338     stringBuilder.append(resourceHelperTypes());
339
340     return stringBuilder.toString();
341 }
342
343 String VertexEntryPointScaffolding::signature(String& functionName)
344 {
345     StringBuilder stringBuilder;
346
347     stringBuilder.append(makeString("vertex ", m_returnStructName, ' ', functionName, '(', m_stageInStructName, ' ', m_stageInParameterName, " [[stage_in]]"));
348     if (auto resourceSignature = this->resourceSignature())
349         stringBuilder.append(makeString(", ", *resourceSignature));
350     if (auto builtInsSignature = this->builtInsSignature())
351         stringBuilder.append(makeString(", ", *builtInsSignature));
352     stringBuilder.append(")");
353
354     return stringBuilder.toString();
355 }
356
357 String VertexEntryPointScaffolding::unpack()
358 {
359     StringBuilder stringBuilder;
360
361     stringBuilder.append(unpackResourcesAndNamedBuiltIns());
362
363     for (auto& namedStageIn : m_namedStageIns) {
364         auto& path = m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].path;
365         auto& elementName = namedStageIn.elementName;
366         stringBuilder.append(makeString(mangledInputPath(path), " = ", m_stageInStructName, '.', elementName, ";\n"));
367     }
368
369     return stringBuilder.toString();
370 }
371
372 String VertexEntryPointScaffolding::pack(const String& inputVariableName, const String& outputVariableName)
373 {
374     StringBuilder stringBuilder;
375     stringBuilder.append(makeString(m_returnStructName, ' ', outputVariableName));
376     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
377         auto& elementName = m_namedOutputs[i].elementName;
378         auto& path = m_entryPointItems.outputs[i].path;
379         stringBuilder.append(makeString(outputVariableName, '.', elementName, " = ", inputVariableName, mangledOutputPath(path), ";\n"));
380     }
381     return stringBuilder.toString();
382 }
383
384 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>&)
385     : EntryPointScaffolding(functionDefinition, intrinsics, typeNamer, entryPointItems, resourceMap, layout, WTFMove(generateNextVariableName))
386     , m_stageInStructName(typeNamer.generateNextTypeName())
387     , m_returnStructName(typeNamer.generateNextTypeName())
388     , m_stageInParameterName(m_generateNextVariableName())
389 {
390     for (size_t i = 0; i < m_entryPointItems.inputs.size(); ++i) {
391         auto& inputItem = m_entryPointItems.inputs[i];
392         if (!WTF::holds_alternative<AST::StageInOutSemantic>(*inputItem.semantic))
393             continue;
394         auto& stageInOutSemantic = WTF::get<AST::StageInOutSemantic>(*inputItem.semantic);
395         NamedStageIn namedStageIn;
396         namedStageIn.indexInEntryPointItems = i;
397         namedStageIn.elementName = m_typeNamer.generateNextStructureElementName();
398         namedStageIn.attributeIndex = stageInOutSemantic.index();
399         m_namedStageIns.append(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         NamedOutput namedOutput;
405         namedOutput.elementName = m_typeNamer.generateNextStructureElementName();
406         m_namedOutputs.uncheckedAppend(WTFMove(namedOutput));
407     }
408 }
409
410 String FragmentEntryPointScaffolding::helperTypes()
411 {
412     StringBuilder stringBuilder;
413
414     stringBuilder.append(makeString("struct ", m_stageInStructName, " {\n"));
415     for (auto& namedStageIn : m_namedStageIns) {
416         auto mangledTypeName = m_typeNamer.mangledNameForType(*m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].unnamedType);
417         auto elementName = namedStageIn.elementName;
418         auto attributeIndex = namedStageIn.elementName;
419         stringBuilder.append(makeString("    ", mangledTypeName, ' ', elementName, " [[user(", attributeIndex, ")]];\n"));
420     }
421     stringBuilder.append("}\n\n");
422
423     stringBuilder.append(makeString("struct ", m_returnStructName, " {\n"));
424     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
425         auto& outputItem = m_entryPointItems.outputs[i];
426         auto mangledTypeName = m_typeNamer.mangledNameForType(*outputItem.unnamedType);
427         auto elementName = m_namedOutputs[i].elementName;
428         auto attribute = attributeForSemantic(*outputItem.semantic);
429         stringBuilder.append(makeString("    ", mangledTypeName, ' ', elementName, ' ', attribute, ";\n"));
430     }
431     stringBuilder.append("}\n\n");
432
433     stringBuilder.append(resourceHelperTypes());
434
435     return stringBuilder.toString();
436 }
437
438 String FragmentEntryPointScaffolding::signature(String& functionName)
439 {
440     StringBuilder stringBuilder;
441
442     stringBuilder.append(makeString("fragment ", m_returnStructName, ' ', functionName, '(', m_stageInStructName, ' ', m_stageInParameterName, " [[stage_in]]"));
443     if (auto resourceSignature = this->resourceSignature())
444         stringBuilder.append(makeString(", ", *resourceSignature));
445     if (auto builtInsSignature = this->builtInsSignature())
446         stringBuilder.append(makeString(", ", *builtInsSignature));
447     stringBuilder.append(")");
448
449     return stringBuilder.toString();
450 }
451
452 String FragmentEntryPointScaffolding::unpack()
453 {
454     StringBuilder stringBuilder;
455
456     stringBuilder.append(unpackResourcesAndNamedBuiltIns());
457
458     for (auto& namedStageIn : m_namedStageIns) {
459         auto& path = m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].path;
460         auto& elementName = namedStageIn.elementName;
461         stringBuilder.append(makeString(mangledInputPath(path), " = ", m_stageInStructName, '.', elementName, ";\n"));
462     }
463
464     return stringBuilder.toString();
465 }
466
467 String FragmentEntryPointScaffolding::pack(const String& inputVariableName, const String& outputVariableName)
468 {
469     StringBuilder stringBuilder;
470     stringBuilder.append(makeString(m_returnStructName, ' ', outputVariableName));
471     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
472         auto& elementName = m_namedOutputs[i].elementName;
473         auto& path = m_entryPointItems.outputs[i].path;
474         stringBuilder.append(makeString(outputVariableName, '.', elementName, " = ", inputVariableName, mangledOutputPath(path), ";\n"));
475     }
476     return stringBuilder.toString();
477 }
478
479 ComputeEntryPointScaffolding::ComputeEntryPointScaffolding(AST::FunctionDefinition& functionDefinition, Intrinsics& intrinsics, TypeNamer& typeNamer, EntryPointItems& entryPointItems, HashMap<Binding*, size_t>& resourceMap, Layout& layout, std::function<String()>&& generateNextVariableName)
480     : EntryPointScaffolding(functionDefinition, intrinsics, typeNamer, entryPointItems, resourceMap, layout, WTFMove(generateNextVariableName))
481 {
482 }
483
484 String ComputeEntryPointScaffolding::helperTypes()
485 {
486     return resourceHelperTypes();
487 }
488
489 String ComputeEntryPointScaffolding::signature(String& functionName)
490 {
491     StringBuilder stringBuilder;
492
493     stringBuilder.append(makeString("compute void ", functionName, '('));
494     bool empty = true;
495     if (auto resourceSignature = this->resourceSignature()) {
496         empty = false;
497         stringBuilder.append(makeString(*resourceSignature));
498     }
499     if (auto builtInsSignature = this->builtInsSignature()) {
500         if (!empty)
501             stringBuilder.append(", ");
502         stringBuilder.append(*builtInsSignature);
503     }
504     stringBuilder.append(")");
505
506     return stringBuilder.toString();
507 }
508
509 String ComputeEntryPointScaffolding::unpack()
510 {
511     return unpackResourcesAndNamedBuiltIns();
512 }
513
514 String ComputeEntryPointScaffolding::pack(const String&, const String&)
515 {
516     ASSERT_NOT_REACHED();
517     return String();
518 }
519
520 }
521
522 }
523
524 }
525
526 #endif