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 <cctype>
9 #include <cstring>
10 #include <limits>
11 #include <unordered_map>
12 #include <unordered_set>
13 
14 #include "compiler/translator/Symbol.h"
15 #include "compiler/translator/TranslatorMetalDirect/Layout.h"
16 
17 using namespace sh;
18 
19 ////////////////////////////////////////////////////////////////////////////////
20 
21 enum class Language
22 {
23     Metal,
24     GLSL,
25 };
26 
RoundUpToMultipleOf(size_t x,size_t multiple)27 static size_t RoundUpToMultipleOf(size_t x, size_t multiple)
28 {
29     const size_t rem = x % multiple;
30     return rem == 0 ? x : x + (multiple - rem);
31 }
32 
requireAlignment(size_t align,bool pad)33 void Layout::requireAlignment(size_t align, bool pad)
34 {
35     alignOf = std::max(alignOf, align);
36     if (pad)
37     {
38         sizeOf = RoundUpToMultipleOf(sizeOf, align);
39     }
40 }
41 
operator ==(const Layout & other) const42 bool Layout::operator==(const Layout &other) const
43 {
44     return sizeOf == other.sizeOf && alignOf == other.alignOf;
45 }
46 
operator +=(const Layout & other)47 void Layout::operator+=(const Layout &other)
48 {
49     requireAlignment(other.alignOf, true);
50     sizeOf += other.sizeOf;
51 }
52 
operator *=(size_t scale)53 void Layout::operator*=(size_t scale)
54 {
55     sizeOf *= scale;
56 }
57 
operator +(const Layout & other) const58 Layout Layout::operator+(const Layout &other) const
59 {
60     auto self = *this;
61     self += other;
62     return self;
63 }
64 
operator *(size_t scale) const65 Layout Layout::operator*(size_t scale) const
66 {
67     auto self = *this;
68     self *= scale;
69     return self;
70 }
71 
ScalarLayoutOf(const TType & type,Language language)72 static Layout ScalarLayoutOf(const TType &type, Language language)
73 {
74     const TBasicType basicType = type.getBasicType();
75     switch (basicType)
76     {
77         case TBasicType::EbtBool:
78             return {1, 1};
79         case TBasicType::EbtInt:
80         case TBasicType::EbtUInt:
81         case TBasicType::EbtFloat:
82             return {4, 4};
83         default:
84             if (IsSampler(basicType))
85             {
86                 switch (language)
87                 {
88                     case Language::Metal:
89                         return {8, 8};
90                     case Language::GLSL:
91                         return {4, 4};
92                 }
93             }
94             UNIMPLEMENTED();
95             return Layout::Invalid();
96     }
97 }
98 
99 static const size_t innerScalesPacked[]   = {0, 1, 2, 3, 4};
100 static const size_t innerScalesUnpacked[] = {0, 1, 2, 4, 4};
101 
MetalLayoutOf(const TType & type,MetalLayoutOfConfig config)102 Layout sh::MetalLayoutOf(const TType &type, MetalLayoutOfConfig config)
103 {
104     ASSERT(type.getNominalSize() <= 4);
105     ASSERT(type.getSecondarySize() <= 4);
106 
107     const TLayoutBlockStorage storage = type.getLayoutQualifier().blockStorage;
108 
109     const bool isPacked = !config.disablePacking && (storage == TLayoutBlockStorage::EbsPacked ||
110                                                      storage == TLayoutBlockStorage::EbsShared);
111 
112     if (type.isArray() && !config.maskArray)
113     {
114         config.maskArray    = true;
115         const Layout layout = MetalLayoutOf(type, config);
116         config.maskArray    = false;
117         const size_t vol    = type.getArraySizeProduct();
118         return layout * vol;
119     }
120 
121     if (const TStructure *structure = type.getStruct())
122     {
123         ASSERT(type.getNominalSize() == 1);
124         ASSERT(type.getSecondarySize() == 1);
125         auto config2             = config;
126         config2.maskArray        = false;
127         auto layout              = Layout::Identity();
128         const TFieldList &fields = structure->fields();
129         for (const TField *field : fields)
130         {
131             layout += MetalLayoutOf(*field->type(), config2);
132         }
133         if (config.assumeStructsAreTailPadded)
134         {
135             size_t pad =
136                 (kDefaultStructAlignmentSize - layout.sizeOf) % kDefaultStructAlignmentSize;
137             layout.sizeOf += pad;
138         }
139         layout.sizeOf = RoundUpToMultipleOf(layout.sizeOf, layout.alignOf);
140         return layout;
141     }
142 
143     if (config.treatSamplersAsTextureEnv && IsSampler(type.getBasicType()))
144     {
145         return {16, 8};  // pointer{8,8} and metal::sampler{8,8}
146     }
147 
148     const Layout scalarLayout = ScalarLayoutOf(type, Language::Metal);
149 
150     if (type.isRank0())
151     {
152         return scalarLayout;
153     }
154     else if (type.isVector())
155     {
156         if (isPacked)
157         {
158             const size_t innerScale = innerScalesPacked[type.getNominalSize()];
159             auto layout = Layout{scalarLayout.sizeOf * innerScale, scalarLayout.alignOf};
160             return layout;
161         }
162         else
163         {
164             const size_t innerScale = innerScalesUnpacked[type.getNominalSize()];
165             auto layout             = Layout::Both(scalarLayout.sizeOf * innerScale);
166             return layout;
167         }
168     }
169     else
170     {
171         ASSERT(type.isMatrix());
172         ASSERT(type.getBasicType() == TBasicType::EbtFloat);
173         // typeCxR <=> typeR[C]
174         const size_t innerScale = innerScalesUnpacked[type.getRows()];
175         const size_t outerScale = static_cast<size_t>(type.getCols());
176         const size_t n          = scalarLayout.sizeOf * innerScale;
177         return {n * outerScale, n};
178     }
179 }
180 
Overlay(TLayoutBlockStorage oldStorage,const TType & type)181 TLayoutBlockStorage sh::Overlay(TLayoutBlockStorage oldStorage, const TType &type)
182 {
183     const TLayoutBlockStorage newStorage = type.getLayoutQualifier().blockStorage;
184     switch (newStorage)
185     {
186         case TLayoutBlockStorage::EbsUnspecified:
187             return oldStorage == TLayoutBlockStorage::EbsUnspecified ? kDefaultLayoutBlockStorage
188                                                                      : oldStorage;
189         default:
190             return newStorage;
191     }
192 }
193 
Overlay(TLayoutMatrixPacking oldPacking,const TType & type)194 TLayoutMatrixPacking sh::Overlay(TLayoutMatrixPacking oldPacking, const TType &type)
195 {
196     const TLayoutMatrixPacking newPacking = type.getLayoutQualifier().matrixPacking;
197     switch (newPacking)
198     {
199         case TLayoutMatrixPacking::EmpUnspecified:
200             return oldPacking == TLayoutMatrixPacking::EmpUnspecified ? kDefaultLayoutMatrixPacking
201                                                                       : oldPacking;
202         default:
203             return newPacking;
204     }
205 }
206 
CanBePacked(TLayoutBlockStorage storage)207 bool sh::CanBePacked(TLayoutBlockStorage storage)
208 {
209     switch (storage)
210     {
211         case TLayoutBlockStorage::EbsPacked:
212         case TLayoutBlockStorage::EbsShared:
213             return true;
214         case TLayoutBlockStorage::EbsStd140:
215         case TLayoutBlockStorage::EbsStd430:
216             return false;
217         case TLayoutBlockStorage::EbsUnspecified:
218             UNREACHABLE();
219             return false;
220     }
221 }
222 
CanBePacked(TLayoutQualifier layoutQualifier)223 bool sh::CanBePacked(TLayoutQualifier layoutQualifier)
224 {
225     return CanBePacked(layoutQualifier.blockStorage);
226 }
227 
CanBePacked(const TType & type)228 bool sh::CanBePacked(const TType &type)
229 {
230     if (!type.isVector())
231     {
232         return false;
233     }
234     return CanBePacked(type.getLayoutQualifier());
235 }
236 
SetBlockStorage(TType & type,TLayoutBlockStorage storage)237 void sh::SetBlockStorage(TType &type, TLayoutBlockStorage storage)
238 {
239     auto qual         = type.getLayoutQualifier();
240     qual.blockStorage = storage;
241     type.setLayoutQualifier(qual);
242 }
243 
CommonGlslStructLayoutOf(TField const * const * begin,TField const * const * end,const TLayoutBlockStorage storage,const TLayoutMatrixPacking matrixPacking,const bool maskArray,const size_t baseAlignment)244 static Layout CommonGlslStructLayoutOf(TField const *const *begin,
245                                        TField const *const *end,
246                                        const TLayoutBlockStorage storage,
247                                        const TLayoutMatrixPacking matrixPacking,
248                                        const bool maskArray,
249                                        const size_t baseAlignment)
250 {
251     const bool isPacked =
252         storage == TLayoutBlockStorage::EbsPacked || storage == TLayoutBlockStorage::EbsShared;
253 
254     auto layout = Layout::Identity();
255     for (auto iter = begin; iter != end; ++iter)
256     {
257         layout += GlslLayoutOf(*(*iter)->type(), storage, matrixPacking, false);
258     }
259     if (!isPacked)  // XXX: Correct?
260     {
261         layout.sizeOf = RoundUpToMultipleOf(layout.sizeOf, layout.alignOf);
262     }
263     layout.requireAlignment(baseAlignment, true);
264     return layout;
265 }
266 
CommonGlslLayoutOf(const TType & type,const TLayoutBlockStorage storage,const TLayoutMatrixPacking matrixPacking,const bool maskArray,const size_t baseAlignment)267 static Layout CommonGlslLayoutOf(const TType &type,
268                                  const TLayoutBlockStorage storage,
269                                  const TLayoutMatrixPacking matrixPacking,
270                                  const bool maskArray,
271                                  const size_t baseAlignment)
272 {
273     ASSERT(storage != TLayoutBlockStorage::EbsUnspecified);
274 
275     const bool isPacked =
276         storage == TLayoutBlockStorage::EbsPacked || storage == TLayoutBlockStorage::EbsShared;
277 
278     if (type.isArray() && !type.isMatrix() && !maskArray)
279     {
280         auto layout = GlslLayoutOf(type, storage, matrixPacking, true);
281         layout *= type.getArraySizeProduct();
282         layout.requireAlignment(baseAlignment, true);
283         return layout;
284     }
285 
286     if (const TStructure *structure = type.getStruct())
287     {
288         ASSERT(type.getNominalSize() == 1);
289         ASSERT(type.getSecondarySize() == 1);
290         const TFieldList &fields = structure->fields();
291         return CommonGlslStructLayoutOf(fields.data(), fields.data() + fields.size(), storage,
292                                         matrixPacking, maskArray, baseAlignment);
293     }
294 
295     const auto scalarLayout = ScalarLayoutOf(type, Language::GLSL);
296 
297     if (type.isRank0())
298     {
299         return scalarLayout;
300     }
301     else if (type.isVector())
302     {
303         if (isPacked)
304         {
305             const size_t sizeScale  = innerScalesPacked[type.getNominalSize()];
306             const size_t alignScale = innerScalesUnpacked[type.getNominalSize()];
307             auto layout =
308                 Layout{scalarLayout.sizeOf * sizeScale, scalarLayout.alignOf * alignScale};
309             return layout;
310         }
311         else
312         {
313             const size_t innerScale = innerScalesUnpacked[type.getNominalSize()];
314             auto layout             = Layout::Both(scalarLayout.sizeOf * innerScale);
315             return layout;
316         }
317     }
318     else
319     {
320         ASSERT(type.isMatrix());
321 
322         size_t innerDim;
323         size_t outerDim;
324 
325         switch (matrixPacking)
326         {
327             case TLayoutMatrixPacking::EmpColumnMajor:
328                 innerDim = static_cast<size_t>(type.getRows());
329                 outerDim = static_cast<size_t>(type.getCols());
330                 break;
331             case TLayoutMatrixPacking::EmpRowMajor:
332                 innerDim = static_cast<size_t>(type.getCols());
333                 outerDim = static_cast<size_t>(type.getRows());
334                 break;
335             case TLayoutMatrixPacking::EmpUnspecified:
336                 UNREACHABLE();
337                 innerDim = 0;
338                 outerDim = 0;
339         }
340 
341         outerDim *= type.getArraySizeProduct();
342 
343         const size_t innerScale = innerScalesUnpacked[innerDim];
344         const size_t n          = innerScale * scalarLayout.sizeOf;
345         Layout layout           = {outerDim * n, n};
346         layout.requireAlignment(baseAlignment, true);
347         return layout;
348     }
349 }
350 
GlslLayoutOf(const TType & type,TLayoutBlockStorage storage,TLayoutMatrixPacking matrixPacking,bool maskArray)351 Layout sh::GlslLayoutOf(const TType &type,
352                         TLayoutBlockStorage storage,
353                         TLayoutMatrixPacking matrixPacking,
354                         bool maskArray)
355 {
356     ASSERT(type.getNominalSize() <= 4);
357     ASSERT(type.getSecondarySize() <= 4);
358 
359     storage       = Overlay(storage, type);
360     matrixPacking = Overlay(matrixPacking, type);
361 
362     switch (storage)
363     {
364         case TLayoutBlockStorage::EbsPacked:
365             return CommonGlslLayoutOf(type, storage, matrixPacking, maskArray, 1);
366         case TLayoutBlockStorage::EbsShared:
367             return CommonGlslLayoutOf(type, storage, matrixPacking, maskArray, 16);
368         case TLayoutBlockStorage::EbsStd140:
369             return CommonGlslLayoutOf(type, storage, matrixPacking, maskArray, 16);
370         case TLayoutBlockStorage::EbsStd430:
371             return CommonGlslLayoutOf(type, storage, matrixPacking, maskArray, 1);
372         case TLayoutBlockStorage::EbsUnspecified:
373             UNREACHABLE();
374             return Layout::Invalid();
375     }
376 }
377 
GlslStructLayoutOf(TField const * const * begin,TField const * const * end,TLayoutBlockStorage storage,TLayoutMatrixPacking matrixPacking,bool maskArray)378 ANGLE_NO_DISCARD Layout sh::GlslStructLayoutOf(TField const *const *begin,
379                                                TField const *const *end,
380                                                TLayoutBlockStorage storage,
381                                                TLayoutMatrixPacking matrixPacking,
382                                                bool maskArray)
383 {
384     ASSERT(storage != TLayoutBlockStorage::EbsUnspecified);
385     ASSERT(matrixPacking != TLayoutMatrixPacking::EmpUnspecified);
386 
387     switch (storage)
388     {
389         case TLayoutBlockStorage::EbsPacked:
390             return CommonGlslStructLayoutOf(begin, end, storage, matrixPacking, maskArray, 1);
391         case TLayoutBlockStorage::EbsShared:
392             return CommonGlslStructLayoutOf(begin, end, storage, matrixPacking, maskArray, 16);
393         case TLayoutBlockStorage::EbsStd140:
394             return CommonGlslStructLayoutOf(begin, end, storage, matrixPacking, maskArray, 16);
395         case TLayoutBlockStorage::EbsStd430:
396             return CommonGlslStructLayoutOf(begin, end, storage, matrixPacking, maskArray, 1);
397         case TLayoutBlockStorage::EbsUnspecified:
398             UNREACHABLE();
399             return Layout::Invalid();
400     }
401 }
402