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