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