// // Copyright 2020 The ANGLE Project Authors. All rights reserved. // Use of this source code is governed by a BSD-style license that can be // found in the LICENSE file. // #include #include #include #include #include #include #include "compiler/translator/Symbol.h" #include "compiler/translator/msl/Layout.h" using namespace sh; //////////////////////////////////////////////////////////////////////////////// enum class Language { Metal, GLSL, }; static size_t RoundUpToMultipleOf(size_t x, size_t multiple) { const size_t rem = x % multiple; return rem == 0 ? x : x + (multiple - rem); } void Layout::requireAlignment(size_t align, bool pad) { alignOf = std::max(alignOf, align); if (pad) { sizeOf = RoundUpToMultipleOf(sizeOf, align); } } bool Layout::operator==(const Layout &other) const { return sizeOf == other.sizeOf && alignOf == other.alignOf; } void Layout::operator+=(const Layout &other) { requireAlignment(other.alignOf, true); sizeOf += other.sizeOf; } void Layout::operator*=(size_t scale) { sizeOf *= scale; } Layout Layout::operator+(const Layout &other) const { auto self = *this; self += other; return self; } Layout Layout::operator*(size_t scale) const { auto self = *this; self *= scale; return self; } static Layout ScalarLayoutOf(const TType &type, Language language) { const TBasicType basicType = type.getBasicType(); switch (basicType) { case TBasicType::EbtBool: return {1, 1}; case TBasicType::EbtInt: case TBasicType::EbtUInt: case TBasicType::EbtFloat: return {4, 4}; default: if (IsSampler(basicType)) { switch (language) { case Language::Metal: return {8, 8}; case Language::GLSL: return {4, 4}; } } UNIMPLEMENTED(); return Layout::Invalid(); } } static const size_t innerScalesPacked[] = {0, 1, 2, 3, 4}; static const size_t innerScalesUnpacked[] = {0, 1, 2, 4, 4}; Layout sh::MetalLayoutOf(const TType &type, MetalLayoutOfConfig config) { ASSERT(type.getNominalSize() <= 4); ASSERT(type.getSecondarySize() <= 4); const TLayoutBlockStorage storage = type.getLayoutQualifier().blockStorage; const bool isPacked = !config.disablePacking && (storage == TLayoutBlockStorage::EbsPacked || storage == TLayoutBlockStorage::EbsShared); if (type.isArray() && !config.maskArray) { config.maskArray = true; const Layout layout = MetalLayoutOf(type, config); config.maskArray = false; const size_t vol = type.getArraySizeProduct(); return layout * vol; } if (const TStructure *structure = type.getStruct()) { ASSERT(type.getNominalSize() == 1); ASSERT(type.getSecondarySize() == 1); auto config2 = config; config2.maskArray = false; auto layout = Layout::Identity(); const TFieldList &fields = structure->fields(); for (const TField *field : fields) { layout += MetalLayoutOf(*field->type(), config2); } if (config.assumeStructsAreTailPadded) { size_t pad = (kDefaultStructAlignmentSize - layout.sizeOf) % kDefaultStructAlignmentSize; layout.sizeOf += pad; } layout.sizeOf = RoundUpToMultipleOf(layout.sizeOf, layout.alignOf); return layout; } if (config.treatSamplersAsTextureEnv && IsSampler(type.getBasicType())) { return {16, 8}; // pointer{8,8} and metal::sampler{8,8} } const Layout scalarLayout = ScalarLayoutOf(type, Language::Metal); if (type.isRank0()) { return scalarLayout; } else if (type.isVector()) { if (isPacked) { const size_t innerScale = innerScalesPacked[type.getNominalSize()]; auto layout = Layout{scalarLayout.sizeOf * innerScale, scalarLayout.alignOf}; return layout; } else { const size_t innerScale = innerScalesUnpacked[type.getNominalSize()]; auto layout = Layout::Both(scalarLayout.sizeOf * innerScale); return layout; } } else { ASSERT(type.isMatrix()); ASSERT(type.getBasicType() == TBasicType::EbtFloat); // typeCxR <=> typeR[C] const size_t innerScale = innerScalesUnpacked[type.getRows()]; const size_t outerScale = static_cast(type.getCols()); const size_t n = scalarLayout.sizeOf * innerScale; return {n * outerScale, n}; } } TLayoutBlockStorage sh::Overlay(TLayoutBlockStorage oldStorage, const TType &type) { const TLayoutBlockStorage newStorage = type.getLayoutQualifier().blockStorage; switch (newStorage) { case TLayoutBlockStorage::EbsUnspecified: return oldStorage == TLayoutBlockStorage::EbsUnspecified ? kDefaultLayoutBlockStorage : oldStorage; default: return newStorage; } } TLayoutMatrixPacking sh::Overlay(TLayoutMatrixPacking oldPacking, const TType &type) { const TLayoutMatrixPacking newPacking = type.getLayoutQualifier().matrixPacking; switch (newPacking) { case TLayoutMatrixPacking::EmpUnspecified: return oldPacking == TLayoutMatrixPacking::EmpUnspecified ? kDefaultLayoutMatrixPacking : oldPacking; default: return newPacking; } } bool sh::CanBePacked(TLayoutBlockStorage storage) { switch (storage) { case TLayoutBlockStorage::EbsPacked: case TLayoutBlockStorage::EbsShared: return true; case TLayoutBlockStorage::EbsStd140: case TLayoutBlockStorage::EbsStd430: return false; case TLayoutBlockStorage::EbsUnspecified: UNREACHABLE(); return false; } } bool sh::CanBePacked(TLayoutQualifier layoutQualifier) { return CanBePacked(layoutQualifier.blockStorage); } bool sh::CanBePacked(const TType &type) { if (!type.isVector()) { return false; } return CanBePacked(type.getLayoutQualifier()); } void sh::SetBlockStorage(TType &type, TLayoutBlockStorage storage) { auto qual = type.getLayoutQualifier(); qual.blockStorage = storage; type.setLayoutQualifier(qual); } static Layout CommonGlslStructLayoutOf(TField const *const *begin, TField const *const *end, const TLayoutBlockStorage storage, const TLayoutMatrixPacking matrixPacking, const bool maskArray, const size_t baseAlignment) { const bool isPacked = storage == TLayoutBlockStorage::EbsPacked || storage == TLayoutBlockStorage::EbsShared; auto layout = Layout::Identity(); for (auto iter = begin; iter != end; ++iter) { layout += GlslLayoutOf(*(*iter)->type(), storage, matrixPacking, false); } if (!isPacked) // XXX: Correct? { layout.sizeOf = RoundUpToMultipleOf(layout.sizeOf, layout.alignOf); } layout.requireAlignment(baseAlignment, true); return layout; } static Layout CommonGlslLayoutOf(const TType &type, const TLayoutBlockStorage storage, const TLayoutMatrixPacking matrixPacking, const bool maskArray, const size_t baseAlignment) { ASSERT(storage != TLayoutBlockStorage::EbsUnspecified); const bool isPacked = storage == TLayoutBlockStorage::EbsPacked || storage == TLayoutBlockStorage::EbsShared; if (type.isArray() && !type.isMatrix() && !maskArray) { auto layout = GlslLayoutOf(type, storage, matrixPacking, true); layout *= type.getArraySizeProduct(); layout.requireAlignment(baseAlignment, true); return layout; } if (const TStructure *structure = type.getStruct()) { ASSERT(type.getNominalSize() == 1); ASSERT(type.getSecondarySize() == 1); const TFieldList &fields = structure->fields(); return CommonGlslStructLayoutOf(fields.data(), fields.data() + fields.size(), storage, matrixPacking, maskArray, baseAlignment); } const auto scalarLayout = ScalarLayoutOf(type, Language::GLSL); if (type.isRank0()) { return scalarLayout; } else if (type.isVector()) { if (isPacked) { const size_t sizeScale = innerScalesPacked[type.getNominalSize()]; const size_t alignScale = innerScalesUnpacked[type.getNominalSize()]; auto layout = Layout{scalarLayout.sizeOf * sizeScale, scalarLayout.alignOf * alignScale}; return layout; } else { const size_t innerScale = innerScalesUnpacked[type.getNominalSize()]; auto layout = Layout::Both(scalarLayout.sizeOf * innerScale); return layout; } } else { ASSERT(type.isMatrix()); size_t innerDim; size_t outerDim; switch (matrixPacking) { case TLayoutMatrixPacking::EmpColumnMajor: innerDim = static_cast(type.getRows()); outerDim = static_cast(type.getCols()); break; case TLayoutMatrixPacking::EmpRowMajor: innerDim = static_cast(type.getCols()); outerDim = static_cast(type.getRows()); break; case TLayoutMatrixPacking::EmpUnspecified: UNREACHABLE(); innerDim = 0; outerDim = 0; } outerDim *= type.getArraySizeProduct(); const size_t innerScale = innerScalesUnpacked[innerDim]; const size_t n = innerScale * scalarLayout.sizeOf; Layout layout = {outerDim * n, n}; layout.requireAlignment(baseAlignment, true); return layout; } } Layout sh::GlslLayoutOf(const TType &type, TLayoutBlockStorage storage, TLayoutMatrixPacking matrixPacking, bool maskArray) { ASSERT(type.getNominalSize() <= 4); ASSERT(type.getSecondarySize() <= 4); storage = Overlay(storage, type); matrixPacking = Overlay(matrixPacking, type); switch (storage) { case TLayoutBlockStorage::EbsPacked: return CommonGlslLayoutOf(type, storage, matrixPacking, maskArray, 1); case TLayoutBlockStorage::EbsShared: return CommonGlslLayoutOf(type, storage, matrixPacking, maskArray, 16); case TLayoutBlockStorage::EbsStd140: return CommonGlslLayoutOf(type, storage, matrixPacking, maskArray, 16); case TLayoutBlockStorage::EbsStd430: return CommonGlslLayoutOf(type, storage, matrixPacking, maskArray, 1); case TLayoutBlockStorage::EbsUnspecified: UNREACHABLE(); return Layout::Invalid(); } } [[nodiscard]] Layout sh::GlslStructLayoutOf(TField const *const *begin, TField const *const *end, TLayoutBlockStorage storage, TLayoutMatrixPacking matrixPacking, bool maskArray) { ASSERT(storage != TLayoutBlockStorage::EbsUnspecified); ASSERT(matrixPacking != TLayoutMatrixPacking::EmpUnspecified); switch (storage) { case TLayoutBlockStorage::EbsPacked: return CommonGlslStructLayoutOf(begin, end, storage, matrixPacking, maskArray, 1); case TLayoutBlockStorage::EbsShared: return CommonGlslStructLayoutOf(begin, end, storage, matrixPacking, maskArray, 16); case TLayoutBlockStorage::EbsStd140: return CommonGlslStructLayoutOf(begin, end, storage, matrixPacking, maskArray, 16); case TLayoutBlockStorage::EbsStd430: return CommonGlslStructLayoutOf(begin, end, storage, matrixPacking, maskArray, 1); case TLayoutBlockStorage::EbsUnspecified: UNREACHABLE(); return Layout::Invalid(); } }