1 //
2 // Copyright 2020 The ANGLE Project Authors. All rights reserved.
3 // Use of this source code is governed by a BSD-style license that can be
4 // found in the LICENSE file.
5 //
6
7 #include <algorithm>
8 #include <limits>
9
10 #include "compiler/translator/ImmutableStringBuilder.h"
11 #include "compiler/translator/TranslatorMetalDirect/AstHelpers.h"
12 #include "compiler/translator/TranslatorMetalDirect/SymbolEnv.h"
13 #include "compiler/translator/tree_util/IntermRebuild.h"
14
15 using namespace sh;
16
17 ////////////////////////////////////////////////////////////////////////////////
18
19 constexpr AddressSpace kAddressSpaces[] = {
20 AddressSpace::Constant,
21 AddressSpace::Device,
22 AddressSpace::Thread,
23 };
24
toString(AddressSpace space)25 char const *sh::toString(AddressSpace space)
26 {
27 switch (space)
28 {
29 case AddressSpace::Constant:
30 return "constant";
31 case AddressSpace::Device:
32 return "device";
33 case AddressSpace::Thread:
34 return "thread";
35 }
36 }
37
38 ////////////////////////////////////////////////////////////////////////////////
39
40 using NameToStruct = std::map<Name, const TStructure *>;
41
42 class StructFinder : TIntermRebuild
43 {
44 NameToStruct nameToStruct;
45
StructFinder(TCompiler & compiler)46 StructFinder(TCompiler &compiler) : TIntermRebuild(compiler, true, false) {}
47
visitDeclarationPre(TIntermDeclaration & node)48 PreResult visitDeclarationPre(TIntermDeclaration &node) override
49 {
50 Declaration decl = ViewDeclaration(node);
51 const TVariable &var = decl.symbol.variable();
52 const TType &type = var.getType();
53
54 if (var.symbolType() == SymbolType::Empty && type.isStructSpecifier())
55 {
56 const TStructure *s = type.getStruct();
57 ASSERT(s);
58 const Name name(*s);
59 const TStructure *&z = nameToStruct[name];
60 ASSERT(!z);
61 z = s;
62 }
63
64 return node;
65 }
66
visitFunctionDefinitionPre(TIntermFunctionDefinition & node)67 PreResult visitFunctionDefinitionPre(TIntermFunctionDefinition &node) override
68 {
69 return {node, VisitBits::Neither};
70 }
71
72 public:
FindStructs(TCompiler & compiler,TIntermBlock & root)73 static NameToStruct FindStructs(TCompiler &compiler, TIntermBlock &root)
74 {
75 StructFinder finder(compiler);
76 if (!finder.rebuildRoot(root))
77 {
78 UNREACHABLE();
79 }
80 return std::move(finder.nameToStruct);
81 }
82 };
83
84 ////////////////////////////////////////////////////////////////////////////////
85
TemplateArg(bool value)86 TemplateArg::TemplateArg(bool value) : mKind(Kind::Bool), mValue(value) {}
87
TemplateArg(int value)88 TemplateArg::TemplateArg(int value) : mKind(Kind::Int), mValue(value) {}
89
TemplateArg(unsigned value)90 TemplateArg::TemplateArg(unsigned value) : mKind(Kind::UInt), mValue(value) {}
91
TemplateArg(const TType & value)92 TemplateArg::TemplateArg(const TType &value) : mKind(Kind::Type), mValue(value) {}
93
operator ==(const TemplateArg & other) const94 bool TemplateArg::operator==(const TemplateArg &other) const
95 {
96 if (mKind != other.mKind)
97 {
98 return false;
99 }
100
101 switch (mKind)
102 {
103 case Kind::Bool:
104 return mValue.b == other.mValue.b;
105 case Kind::Int:
106 return mValue.i == other.mValue.i;
107 case Kind::UInt:
108 return mValue.u == other.mValue.u;
109 case Kind::Type:
110 return *mValue.t == *other.mValue.t;
111 }
112 }
113
operator <(const TemplateArg & other) const114 bool TemplateArg::operator<(const TemplateArg &other) const
115 {
116 if (mKind < other.mKind)
117 {
118 return true;
119 }
120
121 if (mKind > other.mKind)
122 {
123 return false;
124 }
125
126 switch (mKind)
127 {
128 case Kind::Bool:
129 return mValue.b < other.mValue.b;
130 case Kind::Int:
131 return mValue.i < other.mValue.i;
132 case Kind::UInt:
133 return mValue.u < other.mValue.u;
134 case Kind::Type:
135 return *mValue.t < *other.mValue.t;
136 }
137 }
138
139 ////////////////////////////////////////////////////////////////////////////////
140
operator ==(const TemplateName & other) const141 bool SymbolEnv::TemplateName::operator==(const TemplateName &other) const
142 {
143 return baseName == other.baseName && templateArgs == other.templateArgs;
144 }
145
operator <(const TemplateName & other) const146 bool SymbolEnv::TemplateName::operator<(const TemplateName &other) const
147 {
148 if (baseName < other.baseName)
149 {
150 return true;
151 }
152 if (other.baseName < baseName)
153 {
154 return false;
155 }
156 return templateArgs < other.templateArgs;
157 }
158
empty() const159 bool SymbolEnv::TemplateName::empty() const
160 {
161 return baseName.empty() && templateArgs.empty();
162 }
163
clear()164 void SymbolEnv::TemplateName::clear()
165 {
166 baseName = Name();
167 templateArgs.clear();
168 }
169
fullName(std::string & buffer) const170 Name SymbolEnv::TemplateName::fullName(std::string &buffer) const
171 {
172 ASSERT(buffer.empty());
173
174 if (templateArgs.empty())
175 {
176 return baseName;
177 }
178
179 static constexpr size_t n = std::max({
180 std::numeric_limits<unsigned>::digits10, //
181 std::numeric_limits<int>::digits10, //
182 5, // max_length("true", "false")
183 });
184
185 buffer.reserve(baseName.rawName().length() + (n + 2) * templateArgs.size() + 1);
186 buffer += baseName.rawName().data();
187
188 if (!templateArgs.empty())
189 {
190 buffer += "<";
191
192 bool first = true;
193 char argBuffer[n + 1];
194 for (const TemplateArg &arg : templateArgs)
195 {
196 if (first)
197 {
198 first = false;
199 }
200 else
201 {
202 buffer += ", ";
203 }
204
205 const TemplateArg::Value value = arg.value();
206 const TemplateArg::Kind kind = arg.kind();
207 switch (kind)
208 {
209 case TemplateArg::Kind::Bool:
210 if (value.b)
211 {
212 buffer += "true";
213 }
214 else
215 {
216 buffer += "false";
217 }
218 break;
219
220 case TemplateArg::Kind::Int:
221 sprintf(argBuffer, "%i", value.i);
222 buffer += argBuffer;
223 break;
224
225 case TemplateArg::Kind::UInt:
226 sprintf(argBuffer, "%u", value.u);
227 buffer += argBuffer;
228 break;
229
230 case TemplateArg::Kind::Type:
231 {
232 const TType &type = *value.t;
233 if (const TStructure *s = type.getStruct())
234 {
235 buffer += s->name().data();
236 }
237 else if (HasScalarBasicType(type))
238 {
239 ASSERT(!type.isArray()); // TODO
240 buffer += type.getBasicString();
241 if (type.isVector())
242 {
243 sprintf(argBuffer, "%i", type.getNominalSize());
244 buffer += argBuffer;
245 }
246 else if (type.isMatrix())
247 {
248 sprintf(argBuffer, "%i", type.getCols());
249 buffer += argBuffer;
250 buffer += "x";
251 sprintf(argBuffer, "%i", type.getRows());
252 buffer += argBuffer;
253 }
254 }
255 }
256 break;
257 }
258 }
259
260 buffer += ">";
261 }
262
263 const ImmutableString name(buffer);
264 buffer.clear();
265
266 return Name(name, baseName.symbolType());
267 }
268
assign(const Name & name,size_t argCount,const TemplateArg * args)269 void SymbolEnv::TemplateName::assign(const Name &name, size_t argCount, const TemplateArg *args)
270 {
271 baseName = name;
272 templateArgs.clear();
273 for (size_t i = 0; i < argCount; ++i)
274 {
275 templateArgs.push_back(args[i]);
276 }
277 }
278
279 ////////////////////////////////////////////////////////////////////////////////
280
SymbolEnv(TCompiler & compiler,TIntermBlock & root)281 SymbolEnv::SymbolEnv(TCompiler &compiler, TIntermBlock &root)
282 : mSymbolTable(compiler.getSymbolTable()),
283 mNameToStruct(StructFinder::FindStructs(compiler, root))
284 {}
285
remap(const TStructure & s) const286 const TStructure &SymbolEnv::remap(const TStructure &s) const
287 {
288 const Name name(s);
289 auto iter = mNameToStruct.find(name);
290 if (iter == mNameToStruct.end())
291 {
292 return s;
293 }
294 const TStructure &z = *iter->second;
295 return z;
296 }
297
remap(const TStructure * s) const298 const TStructure *SymbolEnv::remap(const TStructure *s) const
299 {
300 if (s)
301 {
302 return &remap(*s);
303 }
304 return nullptr;
305 }
306
getFunctionOverloadImpl()307 const TFunction &SymbolEnv::getFunctionOverloadImpl()
308 {
309 ASSERT(!mReusableSigBuffer.empty());
310
311 SigToFunc &sigToFunc = mOverloads[mReusableTemplateNameBuffer];
312 TFunction *&func = sigToFunc[mReusableSigBuffer];
313
314 if (!func)
315 {
316 const TType &returnType = mReusableSigBuffer.back();
317 mReusableSigBuffer.pop_back();
318
319 const Name name = mReusableTemplateNameBuffer.fullName(mReusableStringBuffer);
320
321 func = new TFunction(&mSymbolTable, name.rawName(), name.symbolType(), &returnType, false);
322 for (const TType ¶mType : mReusableSigBuffer)
323 {
324 func->addParameter(
325 new TVariable(&mSymbolTable, kEmptyImmutableString, ¶mType, SymbolType::Empty));
326 }
327 }
328
329 mReusableSigBuffer.clear();
330 mReusableTemplateNameBuffer.clear();
331
332 return *func;
333 }
334
getFunctionOverload(const Name & name,const TType & returnType,size_t paramCount,const TType ** paramTypes,size_t templateArgCount,const TemplateArg * templateArgs)335 const TFunction &SymbolEnv::getFunctionOverload(const Name &name,
336 const TType &returnType,
337 size_t paramCount,
338 const TType **paramTypes,
339 size_t templateArgCount,
340 const TemplateArg *templateArgs)
341 {
342 ASSERT(mReusableSigBuffer.empty());
343 ASSERT(mReusableTemplateNameBuffer.empty());
344
345 for (size_t i = 0; i < paramCount; ++i)
346 {
347 mReusableSigBuffer.push_back(*paramTypes[i]);
348 }
349 mReusableSigBuffer.push_back(returnType);
350 mReusableTemplateNameBuffer.assign(name, templateArgCount, templateArgs);
351 return getFunctionOverloadImpl();
352 }
353
callFunctionOverload(const Name & name,const TType & returnType,TIntermSequence & args,size_t templateArgCount,const TemplateArg * templateArgs)354 TIntermAggregate &SymbolEnv::callFunctionOverload(const Name &name,
355 const TType &returnType,
356 TIntermSequence &args,
357 size_t templateArgCount,
358 const TemplateArg *templateArgs)
359 {
360 ASSERT(mReusableSigBuffer.empty());
361 ASSERT(mReusableTemplateNameBuffer.empty());
362
363 for (TIntermNode *arg : args)
364 {
365 TIntermTyped *targ = arg->getAsTyped();
366 ASSERT(targ);
367 mReusableSigBuffer.push_back(targ->getType());
368 }
369 mReusableSigBuffer.push_back(returnType);
370 mReusableTemplateNameBuffer.assign(name, templateArgCount, templateArgs);
371 const TFunction &func = getFunctionOverloadImpl();
372 return *TIntermAggregate::CreateRawFunctionCall(func, &args);
373 }
374
newStructure(const Name & name,TFieldList & fields)375 const TStructure &SymbolEnv::newStructure(const Name &name, TFieldList &fields)
376 {
377 ASSERT(name.symbolType() == SymbolType::AngleInternal);
378
379 TStructure *&s = mAngleStructs[name.rawName()];
380 ASSERT(!s);
381 s = new TStructure(&mSymbolTable, name.rawName(), &fields, name.symbolType());
382 return *s;
383 }
384
getTextureEnv(TBasicType samplerType)385 const TStructure &SymbolEnv::getTextureEnv(TBasicType samplerType)
386 {
387 ASSERT(IsSampler(samplerType));
388 const TStructure *&env = mTextureEnvs[samplerType];
389 if (env == nullptr)
390 {
391 auto *textureType = new TType(samplerType);
392 auto *texture =
393 new TField(textureType, ImmutableString("texture"), kNoSourceLoc, SymbolType::BuiltIn);
394 markAsPointer(*texture, AddressSpace::Thread);
395
396 auto *sampler = new TField(new TType(&getSamplerStruct(), false),
397 ImmutableString("sampler"), kNoSourceLoc, SymbolType::BuiltIn);
398 markAsPointer(*sampler, AddressSpace::Thread);
399
400 std::string envName;
401 envName += "TextureEnv<";
402 envName += GetTextureTypeName(samplerType).rawName().data();
403 envName += ">";
404
405 env = &newStructure(Name(envName, SymbolType::AngleInternal),
406 *new TFieldList{texture, sampler});
407 }
408 return *env;
409 }
410
getSamplerStruct()411 const TStructure &SymbolEnv::getSamplerStruct()
412 {
413 if (!mSampler)
414 {
415 mSampler = new TStructure(&mSymbolTable, ImmutableString("metal::sampler"),
416 new TFieldList(), SymbolType::BuiltIn);
417 }
418 return *mSampler;
419 }
420
markSpace(VarField x,AddressSpace space,std::unordered_map<VarField,AddressSpace> & map)421 void SymbolEnv::markSpace(VarField x,
422 AddressSpace space,
423 std::unordered_map<VarField, AddressSpace> &map)
424 {
425 // It is in principle permissible to have references to pointers or multiple pointers, but this
426 // is not required for now and would require code changes to get right.
427 ASSERT(!isPointer(x));
428 ASSERT(!isReference(x));
429
430 map[x] = space;
431 }
432
removeSpace(VarField x,std::unordered_map<VarField,AddressSpace> & map)433 void SymbolEnv::removeSpace(VarField x, std::unordered_map<VarField, AddressSpace> &map)
434 {
435 // It is in principle permissible to have references to pointers or multiple pointers, but this
436 // is not required for now and would require code changes to get right.
437 map.erase(x);
438 }
439
isSpace(VarField x,const std::unordered_map<VarField,AddressSpace> & map) const440 const AddressSpace *SymbolEnv::isSpace(VarField x,
441 const std::unordered_map<VarField, AddressSpace> &map) const
442 {
443 const auto iter = map.find(x);
444 if (iter == map.end())
445 {
446 return nullptr;
447 }
448 const AddressSpace space = iter->second;
449 const auto index = static_cast<std::underlying_type_t<AddressSpace>>(space);
450 return &kAddressSpaces[index];
451 }
452
markAsPointer(VarField x,AddressSpace space)453 void SymbolEnv::markAsPointer(VarField x, AddressSpace space)
454 {
455 return markSpace(x, space, mPointers);
456 }
457
removePointer(VarField x)458 void SymbolEnv::removePointer(VarField x)
459 {
460 return removeSpace(x, mPointers);
461 }
462
markAsReference(VarField x,AddressSpace space)463 void SymbolEnv::markAsReference(VarField x, AddressSpace space)
464 {
465 return markSpace(x, space, mReferences);
466 }
467
isPointer(VarField x) const468 const AddressSpace *SymbolEnv::isPointer(VarField x) const
469 {
470 return isSpace(x, mPointers);
471 }
472
isReference(VarField x) const473 const AddressSpace *SymbolEnv::isReference(VarField x) const
474 {
475 return isSpace(x, mReferences);
476 }
477
markAsPacked(const TField & field)478 void SymbolEnv::markAsPacked(const TField &field)
479 {
480 mPackedFields.insert(&field);
481 }
482
isPacked(const TField & field) const483 bool SymbolEnv::isPacked(const TField &field) const
484 {
485 return mPackedFields.find(&field) != mPackedFields.end();
486 }
487
markAsUBO(VarField x)488 void SymbolEnv::markAsUBO(VarField x)
489 {
490 mUboFields.insert(x);
491 }
492
isUBO(VarField x) const493 bool SymbolEnv::isUBO(VarField x) const
494 {
495 return mUboFields.find(x) != mUboFields.end();
496 }
497
GetTextureBasicType(TBasicType basicType)498 static TBasicType GetTextureBasicType(TBasicType basicType)
499 {
500 ASSERT(IsSampler(basicType));
501
502 switch (basicType)
503 {
504 case EbtSampler2D:
505 case EbtSampler3D:
506 case EbtSamplerCube:
507 case EbtSampler2DArray:
508 case EbtSamplerExternalOES:
509 case EbtSamplerExternal2DY2YEXT:
510 case EbtSampler2DRect:
511 case EbtSampler2DMS:
512 case EbtSampler2DMSArray:
513 case EbtSamplerVideoWEBGL:
514 case EbtSampler2DShadow:
515 case EbtSamplerCubeShadow:
516 case EbtSampler2DArrayShadow:
517 case EbtSampler1D:
518 case EbtSampler1DArray:
519 case EbtSampler1DArrayShadow:
520 case EbtSamplerBuffer:
521 case EbtSamplerCubeArray:
522 case EbtSamplerCubeArrayShadow:
523 case EbtSampler1DShadow:
524 case EbtSampler2DRectShadow:
525 return TBasicType::EbtFloat;
526
527 case EbtISampler2D:
528 case EbtISampler3D:
529 case EbtISamplerCube:
530 case EbtISampler2DArray:
531 case EbtISampler2DMS:
532 case EbtISampler2DMSArray:
533 case EbtISampler1D:
534 case EbtISampler1DArray:
535 case EbtISampler2DRect:
536 case EbtISamplerBuffer:
537 case EbtISamplerCubeArray:
538 return TBasicType::EbtInt;
539
540 case EbtUSampler2D:
541 case EbtUSampler3D:
542 case EbtUSamplerCube:
543 case EbtUSampler2DArray:
544 case EbtUSampler2DMS:
545 case EbtUSampler2DMSArray:
546 case EbtUSampler1D:
547 case EbtUSampler1DArray:
548 case EbtUSampler2DRect:
549 case EbtUSamplerBuffer:
550 case EbtUSamplerCubeArray:
551 return TBasicType::EbtUInt;
552
553 default:
554 UNREACHABLE();
555 return TBasicType::EbtVoid;
556 }
557 }
558
GetTextureTypeName(TBasicType samplerType)559 Name sh::GetTextureTypeName(TBasicType samplerType)
560 {
561 ASSERT(IsSampler(samplerType));
562
563 const TBasicType textureType = GetTextureBasicType(samplerType);
564 const char *name;
565
566 #define HANDLE_TEXTURE_NAME(baseName) \
567 do \
568 { \
569 switch (textureType) \
570 { \
571 case TBasicType::EbtFloat: \
572 name = "metal::" baseName "<float>"; \
573 break; \
574 case TBasicType::EbtInt: \
575 name = "metal::" baseName "<int>"; \
576 break; \
577 case TBasicType::EbtUInt: \
578 name = "metal::" baseName "<uint>"; \
579 break; \
580 default: \
581 UNREACHABLE(); \
582 name = nullptr; \
583 break; \
584 } \
585 } while (false)
586
587 switch (samplerType)
588 {
589 // 1d
590 case EbtSampler1D: // Desktop GLSL sampler type:
591 case EbtISampler1D:
592 case EbtUSampler1D:
593 HANDLE_TEXTURE_NAME("texture1d");
594 break;
595
596 // 1d array
597 case EbtSampler1DArray:
598 case EbtISampler1DArray:
599 case EbtUSampler1DArray:
600 HANDLE_TEXTURE_NAME("texture1d_array");
601 break;
602
603 // Buffer textures
604 case EbtSamplerBuffer:
605 case EbtISamplerBuffer:
606 case EbtUSamplerBuffer:
607 HANDLE_TEXTURE_NAME("texture_buffer");
608 break;
609
610 // 2d textures
611 case EbtSampler2D:
612 case EbtISampler2D:
613 case EbtUSampler2D:
614 case EbtSampler2DRect:
615 case EbtUSampler2DRect:
616 case EbtISampler2DRect:
617 HANDLE_TEXTURE_NAME("texture2d");
618 break;
619
620 // 3d textures
621 case EbtSampler3D:
622 case EbtISampler3D:
623 case EbtUSampler3D:
624 HANDLE_TEXTURE_NAME("texture3d");
625 break;
626
627 // Cube textures
628 case EbtSamplerCube:
629 case EbtISamplerCube:
630 case EbtUSamplerCube:
631 HANDLE_TEXTURE_NAME("texturecube");
632 break;
633
634 // 2d array textures
635 case EbtSampler2DArray:
636 case EbtUSampler2DArray:
637 case EbtISampler2DArray:
638 HANDLE_TEXTURE_NAME("texture2d_array");
639 break;
640
641 case EbtSampler2DMS:
642 case EbtISampler2DMS:
643 case EbtUSampler2DMS:
644 HANDLE_TEXTURE_NAME("texture2d_ms");
645 break;
646
647 case EbtSampler2DMSArray:
648 case EbtISampler2DMSArray:
649 case EbtUSampler2DMSArray:
650 HANDLE_TEXTURE_NAME("texture2d_ms_array");
651 break;
652
653 // cube array
654 case EbtSamplerCubeArray:
655 case EbtISamplerCubeArray:
656 case EbtUSamplerCubeArray:
657 HANDLE_TEXTURE_NAME("texturecube_array");
658 break;
659
660 // Shadow
661 case EbtSampler1DShadow:
662 case EbtSampler1DArrayShadow:
663 UNIMPLEMENTED();
664 HANDLE_TEXTURE_NAME("TODO");
665 break;
666
667 case EbtSampler2DRectShadow:
668 case EbtSampler2DShadow:
669 HANDLE_TEXTURE_NAME("depth2d");
670 break;
671
672 case EbtSamplerCubeShadow:
673 HANDLE_TEXTURE_NAME("depthcube");
674 break;
675
676 case EbtSampler2DArrayShadow:
677 HANDLE_TEXTURE_NAME("depth2d_array");
678 break;
679
680 case EbtSamplerCubeArrayShadow:
681 HANDLE_TEXTURE_NAME("depthcube_array");
682 break;
683
684 // Extentions
685 case EbtSamplerExternalOES: // Only valid if OES_EGL_image_external exists:
686 case EbtSamplerExternal2DY2YEXT: // Only valid if GL_EXT_YUV_target exists:
687 case EbtSamplerVideoWEBGL:
688 UNIMPLEMENTED();
689 HANDLE_TEXTURE_NAME("TODO");
690 break;
691
692 default:
693 UNREACHABLE();
694 name = nullptr;
695 break;
696 }
697
698 #undef HANDLE_TEXTURE_NAME
699
700 return Name(name, SymbolType::BuiltIn);
701 }
702