[ir][msl] Emit struct types
Add emission of struct types to the MSL IR generator.
Bug: tint:1967
Change-Id: I18eb138e6da812721908d95c07514a2b99eb98a6
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/138884
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Dan Sinclair <dsinclair@chromium.org>
diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index 8371b8f..e2dd096 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -1130,6 +1130,8 @@
"writer/msl/generator.h",
"writer/msl/generator_impl.cc",
"writer/msl/generator_impl.h",
+ "writer/msl/generator_support.cc",
+ "writer/msl/generator_support.h",
]
deps = [
@@ -2204,6 +2206,7 @@
"writer/msl/generator_impl_type_test.cc",
"writer/msl/generator_impl_unary_op_test.cc",
"writer/msl/generator_impl_variable_decl_statement_test.cc",
+ "writer/msl/generator_support_test.cc",
"writer/msl/test_helper.h",
]
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index 1b64b13..0ccef44 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -694,6 +694,8 @@
writer/msl/generator.h
writer/msl/generator_impl.cc
writer/msl/generator_impl.h
+ writer/msl/generator_support.cc
+ writer/msl/generator_support.h
)
if(${TINT_BUILD_IR})
@@ -1449,6 +1451,7 @@
writer/msl/generator_impl_type_test.cc
writer/msl/generator_impl_unary_op_test.cc
writer/msl/generator_impl_variable_decl_statement_test.cc
+ writer/msl/generator_support_test.cc
writer/msl/test_helper.h
)
diff --git a/src/tint/writer/ast_text_generator.cc b/src/tint/writer/ast_text_generator.cc
index dc0530b..65bf618 100644
--- a/src/tint/writer/ast_text_generator.cc
+++ b/src/tint/writer/ast_text_generator.cc
@@ -30,13 +30,4 @@
return builder_.Symbols().New(prefix).Name();
}
-std::string ASTTextGenerator::StructName(const type::Struct* s) {
- auto name = s->Name().Name();
- if (name.size() > 1 && name[0] == '_' && name[1] == '_') {
- name = utils::GetOrCreate(builtin_struct_names_, s,
- [&] { return UniqueIdentifier(name.substr(2)); });
- }
- return name;
-}
-
} // namespace tint::writer
diff --git a/src/tint/writer/ast_text_generator.h b/src/tint/writer/ast_text_generator.h
index 74c42ae..d385e33 100644
--- a/src/tint/writer/ast_text_generator.h
+++ b/src/tint/writer/ast_text_generator.h
@@ -36,14 +36,7 @@
/// @return a new, unique identifier with the given prefix.
/// @param prefix optional prefix to apply to the generated identifier. If
/// empty "tint_symbol" will be used.
- std::string UniqueIdentifier(const std::string& prefix = "");
-
- /// @param s the semantic structure
- /// @returns the name of the structure, taking special care of builtin
- /// structures that start with double underscores. If the structure is a
- /// builtin, then the returned name will be a unique name without the leading
- /// underscores.
- std::string StructName(const type::Struct* s);
+ std::string UniqueIdentifier(const std::string& prefix = "") override;
protected:
/// @returns the resolved type of the ast::Expression `expr`
diff --git a/src/tint/writer/ir_text_generator.cc b/src/tint/writer/ir_text_generator.cc
index c9bbaee..d467dc7 100644
--- a/src/tint/writer/ir_text_generator.cc
+++ b/src/tint/writer/ir_text_generator.cc
@@ -14,6 +14,8 @@
#include "src/tint/writer/ir_text_generator.h"
+#include "src/tint/utils/map.h"
+
namespace tint::writer {
IRTextGenerator::IRTextGenerator(ir::Module* mod) : ir_(mod) {}
diff --git a/src/tint/writer/ir_text_generator.h b/src/tint/writer/ir_text_generator.h
index 5a6e42b..60ccfbc 100644
--- a/src/tint/writer/ir_text_generator.h
+++ b/src/tint/writer/ir_text_generator.h
@@ -33,7 +33,7 @@
/// @return a new, unique identifier with the given prefix.
/// @param prefix optional prefix to apply to the generated identifier. If
/// empty "tint_symbol" will be used.
- std::string UniqueIdentifier(const std::string& prefix = "");
+ std::string UniqueIdentifier(const std::string& prefix = "") override;
/// @returns the generated shader string
std::string Result() const override {
diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc
index ae7591f..83e769a 100644
--- a/src/tint/writer/msl/generator_impl.cc
+++ b/src/tint/writer/msl/generator_impl.cc
@@ -86,6 +86,7 @@
#include "src/tint/utils/string_stream.h"
#include "src/tint/writer/check_supported_extensions.h"
#include "src/tint/writer/float_to_string.h"
+#include "src/tint/writer/msl/generator_support.h"
namespace tint::writer::msl {
namespace {
@@ -1910,73 +1911,6 @@
return true;
}
-std::string GeneratorImpl::builtin_to_attribute(builtin::BuiltinValue builtin) const {
- switch (builtin) {
- case builtin::BuiltinValue::kPosition:
- return "position";
- case builtin::BuiltinValue::kVertexIndex:
- return "vertex_id";
- case builtin::BuiltinValue::kInstanceIndex:
- return "instance_id";
- case builtin::BuiltinValue::kFrontFacing:
- return "front_facing";
- case builtin::BuiltinValue::kFragDepth:
- return "depth(any)";
- case builtin::BuiltinValue::kLocalInvocationId:
- return "thread_position_in_threadgroup";
- case builtin::BuiltinValue::kLocalInvocationIndex:
- return "thread_index_in_threadgroup";
- case builtin::BuiltinValue::kGlobalInvocationId:
- return "thread_position_in_grid";
- case builtin::BuiltinValue::kWorkgroupId:
- return "threadgroup_position_in_grid";
- case builtin::BuiltinValue::kNumWorkgroups:
- return "threadgroups_per_grid";
- case builtin::BuiltinValue::kSampleIndex:
- return "sample_id";
- case builtin::BuiltinValue::kSampleMask:
- return "sample_mask";
- case builtin::BuiltinValue::kPointSize:
- return "point_size";
- default:
- break;
- }
- return "";
-}
-
-std::string GeneratorImpl::interpolation_to_attribute(
- builtin::InterpolationType type,
- builtin::InterpolationSampling sampling) const {
- std::string attr;
- switch (sampling) {
- case builtin::InterpolationSampling::kCenter:
- attr = "center_";
- break;
- case builtin::InterpolationSampling::kCentroid:
- attr = "centroid_";
- break;
- case builtin::InterpolationSampling::kSample:
- attr = "sample_";
- break;
- case builtin::InterpolationSampling::kUndefined:
- break;
- }
- switch (type) {
- case builtin::InterpolationType::kPerspective:
- attr += "perspective";
- break;
- case builtin::InterpolationType::kLinear:
- attr += "no_perspective";
- break;
- case builtin::InterpolationType::kFlat:
- attr += "flat";
- break;
- case builtin::InterpolationType::kUndefined:
- break;
- }
- return attr;
-}
-
bool GeneratorImpl::EmitEntryPointFunction(const ast::Function* func) {
auto* func_sem = builder_.Sem().Get(func);
@@ -2086,7 +2020,7 @@
builtin_found = true;
- auto name = builtin_to_attribute(builtin);
+ auto name = BuiltinToAttribute(builtin);
if (name.empty()) {
diagnostics_.add_error(diag::System::Writer, "unknown builtin");
return false;
@@ -2835,7 +2769,7 @@
auto& attributes = mem->Attributes();
if (auto builtin = attributes.builtin) {
- auto name = builtin_to_attribute(builtin.value());
+ auto name = BuiltinToAttribute(builtin.value());
if (name.empty()) {
diagnostics_.add_error(diag::System::Writer, "unknown builtin");
return false;
@@ -2866,7 +2800,7 @@
}
if (auto interpolation = attributes.interpolation) {
- auto name = interpolation_to_attribute(interpolation->type, interpolation->sampling);
+ auto name = InterpolationToAttribute(interpolation->type, interpolation->sampling);
if (name.empty()) {
diagnostics_.add_error(diag::System::Writer, "unknown interpolation attribute");
return false;
@@ -2883,7 +2817,7 @@
if (is_host_shareable) {
// Calculate new MSL offset
- auto size_align = MslPackedTypeSizeAndAlign(ty);
+ auto size_align = MslPackedTypeSizeAndAlign(diagnostics_, ty);
if (TINT_UNLIKELY(msl_offset % size_align.align)) {
TINT_ICE(Writer, diagnostics_)
<< "Misaligned MSL structure member " << ty->FriendlyName() << " " << mem_name;
@@ -3056,122 +2990,6 @@
return true;
}
-GeneratorImpl::SizeAndAlign GeneratorImpl::MslPackedTypeSizeAndAlign(const type::Type* ty) {
- return Switch(
- ty,
-
- // https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf
- // 2.1 Scalar Data Types
- [&](const type::U32*) {
- return SizeAndAlign{4, 4};
- },
- [&](const type::I32*) {
- return SizeAndAlign{4, 4};
- },
- [&](const type::F32*) {
- return SizeAndAlign{4, 4};
- },
- [&](const type::F16*) {
- return SizeAndAlign{2, 2};
- },
-
- [&](const type::Vector* vec) {
- auto num_els = vec->Width();
- auto* el_ty = vec->type();
- SizeAndAlign el_size_align = MslPackedTypeSizeAndAlign(el_ty);
- if (el_ty->IsAnyOf<type::U32, type::I32, type::F32, type::F16>()) {
- // Use a packed_vec type for 3-element vectors only.
- if (num_els == 3) {
- // https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf
- // 2.2.3 Packed Vector Types
- return SizeAndAlign{num_els * el_size_align.size, el_size_align.align};
- } else {
- // https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf
- // 2.2 Vector Data Types
- // Vector data types are aligned to their size.
- return SizeAndAlign{num_els * el_size_align.size, num_els * el_size_align.size};
- }
- }
- TINT_UNREACHABLE(Writer, diagnostics_)
- << "Unhandled vector element type " << el_ty->TypeInfo().name;
- return SizeAndAlign{};
- },
-
- [&](const type::Matrix* mat) {
- // https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf
- // 2.3 Matrix Data Types
- auto cols = mat->columns();
- auto rows = mat->rows();
- auto* el_ty = mat->type();
- // Metal only support half and float matrix.
- if (el_ty->IsAnyOf<type::F32, type::F16>()) {
- static constexpr SizeAndAlign table_f32[] = {
- /* float2x2 */ {16, 8},
- /* float2x3 */ {32, 16},
- /* float2x4 */ {32, 16},
- /* float3x2 */ {24, 8},
- /* float3x3 */ {48, 16},
- /* float3x4 */ {48, 16},
- /* float4x2 */ {32, 8},
- /* float4x3 */ {64, 16},
- /* float4x4 */ {64, 16},
- };
- static constexpr SizeAndAlign table_f16[] = {
- /* half2x2 */ {8, 4},
- /* half2x3 */ {16, 8},
- /* half2x4 */ {16, 8},
- /* half3x2 */ {12, 4},
- /* half3x3 */ {24, 8},
- /* half3x4 */ {24, 8},
- /* half4x2 */ {16, 4},
- /* half4x3 */ {32, 8},
- /* half4x4 */ {32, 8},
- };
- if (cols >= 2 && cols <= 4 && rows >= 2 && rows <= 4) {
- if (el_ty->Is<type::F32>()) {
- return table_f32[(3 * (cols - 2)) + (rows - 2)];
- } else {
- return table_f16[(3 * (cols - 2)) + (rows - 2)];
- }
- }
- }
-
- TINT_UNREACHABLE(Writer, diagnostics_)
- << "Unhandled matrix element type " << el_ty->TypeInfo().name;
- return SizeAndAlign{};
- },
-
- [&](const type::Array* arr) {
- if (TINT_UNLIKELY(!arr->IsStrideImplicit())) {
- TINT_ICE(Writer, diagnostics_)
- << "arrays with explicit strides should not exist past the SPIR-V reader";
- return SizeAndAlign{};
- }
- if (arr->Count()->Is<type::RuntimeArrayCount>()) {
- return SizeAndAlign{arr->Stride(), arr->Align()};
- }
- if (auto count = arr->ConstantCount()) {
- return SizeAndAlign{arr->Stride() * count.value(), arr->Align()};
- }
- diagnostics_.add_error(diag::System::Writer, type::Array::kErrExpectedConstantCount);
- return SizeAndAlign{};
- },
-
- [&](const type::Struct* str) {
- // TODO(crbug.com/tint/650): There's an assumption here that MSL's
- // default structure size and alignment matches WGSL's. We need to
- // confirm this.
- return SizeAndAlign{str->Size(), str->Align()};
- },
-
- [&](const type::Atomic* atomic) { return MslPackedTypeSizeAndAlign(atomic->Type()); },
-
- [&](Default) {
- TINT_UNREACHABLE(Writer, diagnostics_) << "Unhandled type " << ty->TypeInfo().name;
- return SizeAndAlign{};
- });
-}
-
template <typename F>
bool GeneratorImpl::CallBuiltinHelper(utils::StringStream& out,
const ast::CallExpression* call,
diff --git a/src/tint/writer/msl/generator_impl.h b/src/tint/writer/msl/generator_impl.h
index ebc3905..42ff42b 100644
--- a/src/tint/writer/msl/generator_impl.h
+++ b/src/tint/writer/msl/generator_impl.h
@@ -360,25 +360,7 @@
/// @returns the name or "" if not valid
std::string generate_builtin_name(const sem::Builtin* builtin);
- /// Converts a builtin to an attribute name
- /// @param builtin the builtin to convert
- /// @returns the string name of the builtin or blank on error
- std::string builtin_to_attribute(builtin::BuiltinValue builtin) const;
-
- /// Converts interpolation attributes to an MSL attribute
- /// @param type the interpolation type
- /// @param sampling the interpolation sampling
- /// @returns the string name of the attribute or blank on error
- std::string interpolation_to_attribute(builtin::InterpolationType type,
- builtin::InterpolationSampling sampling) const;
-
private:
- // A pair of byte size and alignment `uint32_t`s.
- struct SizeAndAlign {
- uint32_t size;
- uint32_t align;
- };
-
/// CallBuiltinHelper will call the builtin helper function, creating it
/// if it hasn't been built already. If the builtin needs to be built then
/// CallBuiltinHelper will generate the function signature and will call
@@ -404,10 +386,6 @@
TextBuffer helpers_; // Helper functions emitted at the top of the output
- /// @returns the MSL packed type size and alignment in bytes for the given
- /// type.
- SizeAndAlign MslPackedTypeSizeAndAlign(const type::Type* ty);
-
std::function<bool()> emit_continuing_;
/// Name of atomicCompareExchangeWeak() helper for the given pointer storage
diff --git a/src/tint/writer/msl/generator_impl_test.cc b/src/tint/writer/msl/generator_impl_test.cc
index 2a627ec..fb5bd6b 100644
--- a/src/tint/writer/msl/generator_impl_test.cc
+++ b/src/tint/writer/msl/generator_impl_test.cc
@@ -61,42 +61,6 @@
)");
}
-struct MslBuiltinData {
- builtin::BuiltinValue builtin;
- const char* attribute_name;
-};
-inline std::ostream& operator<<(std::ostream& out, MslBuiltinData data) {
- utils::StringStream str;
- str << data.builtin;
- out << str.str();
- return out;
-}
-using MslBuiltinConversionTest = TestParamHelper<MslBuiltinData>;
-TEST_P(MslBuiltinConversionTest, Emit) {
- auto params = GetParam();
-
- GeneratorImpl& gen = Build();
-
- EXPECT_EQ(gen.builtin_to_attribute(params.builtin), std::string(params.attribute_name));
-}
-INSTANTIATE_TEST_SUITE_P(
- MslGeneratorImplTest,
- MslBuiltinConversionTest,
- testing::Values(
- MslBuiltinData{builtin::BuiltinValue::kPosition, "position"},
- MslBuiltinData{builtin::BuiltinValue::kVertexIndex, "vertex_id"},
- MslBuiltinData{builtin::BuiltinValue::kInstanceIndex, "instance_id"},
- MslBuiltinData{builtin::BuiltinValue::kFrontFacing, "front_facing"},
- MslBuiltinData{builtin::BuiltinValue::kFragDepth, "depth(any)"},
- MslBuiltinData{builtin::BuiltinValue::kLocalInvocationId, "thread_position_in_threadgroup"},
- MslBuiltinData{builtin::BuiltinValue::kLocalInvocationIndex, "thread_index_in_threadgroup"},
- MslBuiltinData{builtin::BuiltinValue::kGlobalInvocationId, "thread_position_in_grid"},
- MslBuiltinData{builtin::BuiltinValue::kWorkgroupId, "threadgroup_position_in_grid"},
- MslBuiltinData{builtin::BuiltinValue::kNumWorkgroups, "threadgroups_per_grid"},
- MslBuiltinData{builtin::BuiltinValue::kSampleIndex, "sample_id"},
- MslBuiltinData{builtin::BuiltinValue::kSampleMask, "sample_mask"},
- MslBuiltinData{builtin::BuiltinValue::kPointSize, "point_size"}));
-
TEST_F(MslGeneratorImplTest, HasInvariantAttribute_True) {
auto* out = Structure("Out", utils::Vector{
Member("pos", ty.vec4<f32>(),
diff --git a/src/tint/writer/msl/generator_support.cc b/src/tint/writer/msl/generator_support.cc
new file mode 100644
index 0000000..3739cdd
--- /dev/null
+++ b/src/tint/writer/msl/generator_support.cc
@@ -0,0 +1,214 @@
+// Copyright 2023 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "src/tint/writer/msl/generator_support.h"
+
+#include "src/tint/debug.h"
+#include "src/tint/switch.h"
+#include "src/tint/type/array.h"
+#include "src/tint/type/atomic.h"
+#include "src/tint/type/f16.h"
+#include "src/tint/type/f32.h"
+#include "src/tint/type/i32.h"
+#include "src/tint/type/matrix.h"
+#include "src/tint/type/struct.h"
+#include "src/tint/type/u32.h"
+#include "src/tint/type/vector.h"
+
+namespace tint::writer::msl {
+
+std::string BuiltinToAttribute(builtin::BuiltinValue builtin) {
+ switch (builtin) {
+ case builtin::BuiltinValue::kPosition:
+ return "position";
+ case builtin::BuiltinValue::kVertexIndex:
+ return "vertex_id";
+ case builtin::BuiltinValue::kInstanceIndex:
+ return "instance_id";
+ case builtin::BuiltinValue::kFrontFacing:
+ return "front_facing";
+ case builtin::BuiltinValue::kFragDepth:
+ return "depth(any)";
+ case builtin::BuiltinValue::kLocalInvocationId:
+ return "thread_position_in_threadgroup";
+ case builtin::BuiltinValue::kLocalInvocationIndex:
+ return "thread_index_in_threadgroup";
+ case builtin::BuiltinValue::kGlobalInvocationId:
+ return "thread_position_in_grid";
+ case builtin::BuiltinValue::kWorkgroupId:
+ return "threadgroup_position_in_grid";
+ case builtin::BuiltinValue::kNumWorkgroups:
+ return "threadgroups_per_grid";
+ case builtin::BuiltinValue::kSampleIndex:
+ return "sample_id";
+ case builtin::BuiltinValue::kSampleMask:
+ return "sample_mask";
+ case builtin::BuiltinValue::kPointSize:
+ return "point_size";
+ default:
+ break;
+ }
+ return "";
+}
+
+std::string InterpolationToAttribute(builtin::InterpolationType type,
+ builtin::InterpolationSampling sampling) {
+ std::string attr;
+ switch (sampling) {
+ case builtin::InterpolationSampling::kCenter:
+ attr = "center_";
+ break;
+ case builtin::InterpolationSampling::kCentroid:
+ attr = "centroid_";
+ break;
+ case builtin::InterpolationSampling::kSample:
+ attr = "sample_";
+ break;
+ case builtin::InterpolationSampling::kUndefined:
+ break;
+ }
+ switch (type) {
+ case builtin::InterpolationType::kPerspective:
+ attr += "perspective";
+ break;
+ case builtin::InterpolationType::kLinear:
+ attr += "no_perspective";
+ break;
+ case builtin::InterpolationType::kFlat:
+ attr += "flat";
+ break;
+ case builtin::InterpolationType::kUndefined:
+ break;
+ }
+ return attr;
+}
+
+SizeAndAlign MslPackedTypeSizeAndAlign(diag::List diagnostics, const type::Type* ty) {
+ return tint::Switch(
+ ty,
+
+ // https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf
+ // 2.1 Scalar Data Types
+ [&](const type::U32*) {
+ return SizeAndAlign{4, 4};
+ },
+ [&](const type::I32*) {
+ return SizeAndAlign{4, 4};
+ },
+ [&](const type::F32*) {
+ return SizeAndAlign{4, 4};
+ },
+ [&](const type::F16*) {
+ return SizeAndAlign{2, 2};
+ },
+
+ [&](const type::Vector* vec) {
+ auto num_els = vec->Width();
+ auto* el_ty = vec->type();
+ SizeAndAlign el_size_align = MslPackedTypeSizeAndAlign(diagnostics, el_ty);
+ if (el_ty->IsAnyOf<type::U32, type::I32, type::F32, type::F16>()) {
+ // Use a packed_vec type for 3-element vectors only.
+ if (num_els == 3) {
+ // https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf
+ // 2.2.3 Packed Vector Types
+ return SizeAndAlign{num_els * el_size_align.size, el_size_align.align};
+ } else {
+ // https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf
+ // 2.2 Vector Data Types
+ // Vector data types are aligned to their size.
+ return SizeAndAlign{num_els * el_size_align.size, num_els * el_size_align.size};
+ }
+ }
+ TINT_UNREACHABLE(Writer, diagnostics)
+ << "Unhandled vector element type " << el_ty->TypeInfo().name;
+ return SizeAndAlign{};
+ },
+
+ [&](const type::Matrix* mat) {
+ // https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf
+ // 2.3 Matrix Data Types
+ auto cols = mat->columns();
+ auto rows = mat->rows();
+ auto* el_ty = mat->type();
+ // Metal only support half and float matrix.
+ if (el_ty->IsAnyOf<type::F32, type::F16>()) {
+ static constexpr SizeAndAlign table_f32[] = {
+ /* float2x2 */ {16, 8},
+ /* float2x3 */ {32, 16},
+ /* float2x4 */ {32, 16},
+ /* float3x2 */ {24, 8},
+ /* float3x3 */ {48, 16},
+ /* float3x4 */ {48, 16},
+ /* float4x2 */ {32, 8},
+ /* float4x3 */ {64, 16},
+ /* float4x4 */ {64, 16},
+ };
+ static constexpr SizeAndAlign table_f16[] = {
+ /* half2x2 */ {8, 4},
+ /* half2x3 */ {16, 8},
+ /* half2x4 */ {16, 8},
+ /* half3x2 */ {12, 4},
+ /* half3x3 */ {24, 8},
+ /* half3x4 */ {24, 8},
+ /* half4x2 */ {16, 4},
+ /* half4x3 */ {32, 8},
+ /* half4x4 */ {32, 8},
+ };
+ if (cols >= 2 && cols <= 4 && rows >= 2 && rows <= 4) {
+ if (el_ty->Is<type::F32>()) {
+ return table_f32[(3 * (cols - 2)) + (rows - 2)];
+ } else {
+ return table_f16[(3 * (cols - 2)) + (rows - 2)];
+ }
+ }
+ }
+
+ TINT_UNREACHABLE(Writer, diagnostics)
+ << "Unhandled matrix element type " << el_ty->TypeInfo().name;
+ return SizeAndAlign{};
+ },
+
+ [&](const type::Array* arr) {
+ if (TINT_UNLIKELY(!arr->IsStrideImplicit())) {
+ TINT_ICE(Writer, diagnostics)
+ << "arrays with explicit strides should not exist past the SPIR-V reader";
+ return SizeAndAlign{};
+ }
+ if (arr->Count()->Is<type::RuntimeArrayCount>()) {
+ return SizeAndAlign{arr->Stride(), arr->Align()};
+ }
+ if (auto count = arr->ConstantCount()) {
+ return SizeAndAlign{arr->Stride() * count.value(), arr->Align()};
+ }
+ diagnostics.add_error(diag::System::Writer, type::Array::kErrExpectedConstantCount);
+ return SizeAndAlign{};
+ },
+
+ [&](const type::Struct* str) {
+ // TODO(crbug.com/tint/650): There's an assumption here that MSL's
+ // default structure size and alignment matches WGSL's. We need to
+ // confirm this.
+ return SizeAndAlign{str->Size(), str->Align()};
+ },
+
+ [&](const type::Atomic* atomic) {
+ return MslPackedTypeSizeAndAlign(diagnostics, atomic->Type());
+ },
+
+ [&](Default) {
+ TINT_UNREACHABLE(Writer, diagnostics) << "Unhandled type " << ty->TypeInfo().name;
+ return SizeAndAlign{};
+ });
+}
+} // namespace tint::writer::msl
diff --git a/src/tint/writer/msl/generator_support.h b/src/tint/writer/msl/generator_support.h
new file mode 100644
index 0000000..4f5877f
--- /dev/null
+++ b/src/tint/writer/msl/generator_support.h
@@ -0,0 +1,55 @@
+// Copyright 2023 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef SRC_TINT_WRITER_MSL_GENERATOR_SUPPORT_H_
+#define SRC_TINT_WRITER_MSL_GENERATOR_SUPPORT_H_
+
+#include <cstdint>
+#include <string>
+
+#include "src/tint/builtin/builtin_value.h"
+#include "src/tint/builtin/interpolation.h"
+#include "src/tint/diagnostic/diagnostic.h"
+#include "src/tint/type/type.h"
+
+namespace tint::writer::msl {
+
+/// A pair of byte size and alignment `uint32_t`s.
+struct SizeAndAlign {
+ /// The size
+ uint32_t size;
+ /// The alignment
+ uint32_t align;
+};
+
+/// @param diagnostics the diagnostics list
+/// @param ty the type to generate size and align for
+/// @returns the MSL packed type size and alignment in bytes for the given type.
+SizeAndAlign MslPackedTypeSizeAndAlign(diag::List diagnostics, const type::Type* ty);
+
+/// Converts a builtin to an attribute name
+/// @param builtin the builtin to convert
+/// @returns the string name of the builtin or blank on error
+std::string BuiltinToAttribute(builtin::BuiltinValue builtin);
+
+/// Converts interpolation attributes to an MSL attribute
+/// @param type the interpolation type
+/// @param sampling the interpolation sampling
+/// @returns the string name of the attribute or blank on error
+std::string InterpolationToAttribute(builtin::InterpolationType type,
+ builtin::InterpolationSampling sampling);
+
+} // namespace tint::writer::msl
+
+#endif // SRC_TINT_WRITER_MSL_GENERATOR_SUPPORT_H_
diff --git a/src/tint/writer/msl/generator_support_test.cc b/src/tint/writer/msl/generator_support_test.cc
new file mode 100644
index 0000000..24f6f92
--- /dev/null
+++ b/src/tint/writer/msl/generator_support_test.cc
@@ -0,0 +1,55 @@
+// Copyright 2023 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "src/tint/writer/msl/generator_support.h"
+#include "src/tint/writer/msl/test_helper.h"
+
+namespace tint::writer::msl {
+namespace {
+
+struct MslBuiltinData {
+ builtin::BuiltinValue builtin;
+ const char* attribute_name;
+};
+inline std::ostream& operator<<(std::ostream& out, MslBuiltinData data) {
+ utils::StringStream str;
+ str << data.builtin;
+ out << str.str();
+ return out;
+}
+using MslBuiltinConversionTest = TestParamHelper<MslBuiltinData>;
+TEST_P(MslBuiltinConversionTest, Emit) {
+ auto params = GetParam();
+ EXPECT_EQ(BuiltinToAttribute(params.builtin), std::string(params.attribute_name));
+}
+INSTANTIATE_TEST_SUITE_P(
+ MslGeneratorImplTest,
+ MslBuiltinConversionTest,
+ testing::Values(
+ MslBuiltinData{builtin::BuiltinValue::kPosition, "position"},
+ MslBuiltinData{builtin::BuiltinValue::kVertexIndex, "vertex_id"},
+ MslBuiltinData{builtin::BuiltinValue::kInstanceIndex, "instance_id"},
+ MslBuiltinData{builtin::BuiltinValue::kFrontFacing, "front_facing"},
+ MslBuiltinData{builtin::BuiltinValue::kFragDepth, "depth(any)"},
+ MslBuiltinData{builtin::BuiltinValue::kLocalInvocationId, "thread_position_in_threadgroup"},
+ MslBuiltinData{builtin::BuiltinValue::kLocalInvocationIndex, "thread_index_in_threadgroup"},
+ MslBuiltinData{builtin::BuiltinValue::kGlobalInvocationId, "thread_position_in_grid"},
+ MslBuiltinData{builtin::BuiltinValue::kWorkgroupId, "threadgroup_position_in_grid"},
+ MslBuiltinData{builtin::BuiltinValue::kNumWorkgroups, "threadgroups_per_grid"},
+ MslBuiltinData{builtin::BuiltinValue::kSampleIndex, "sample_id"},
+ MslBuiltinData{builtin::BuiltinValue::kSampleMask, "sample_mask"},
+ MslBuiltinData{builtin::BuiltinValue::kPointSize, "point_size"}));
+
+} // namespace
+} // namespace tint::writer::msl
diff --git a/src/tint/writer/msl/ir/generator_impl_ir.cc b/src/tint/writer/msl/ir/generator_impl_ir.cc
index 9b6f84f..d5e9ac9 100644
--- a/src/tint/writer/msl/ir/generator_impl_ir.cc
+++ b/src/tint/writer/msl/ir/generator_impl_ir.cc
@@ -36,6 +36,7 @@
#include "src/tint/type/vector.h"
#include "src/tint/type/void.h"
#include "src/tint/utils/scoped_assignment.h"
+#include "src/tint/writer/msl/generator_support.h"
namespace tint::writer::msl {
namespace {
@@ -283,7 +284,149 @@
diagnostics_.add_error(diag::System::Writer, "invalid texture type");
});
},
+ [&](const type::Struct* str) {
+ out << StructName(str);
+
+ TINT_SCOPED_ASSIGNMENT(current_buffer_, &preamble_buffer_);
+ EmitStructType(str);
+ },
[&](Default) { UNHANDLED_CASE(ty); });
}
+void GeneratorImplIr::EmitStructType(const type::Struct* str) {
+ auto it = emitted_structs_.emplace(str);
+ if (!it.second) {
+ return;
+ }
+
+ // This does not append directly to the preamble because a struct may require other structs, or
+ // the array template, to get emitted before it. So, the struct emits into a temporary text
+ // buffer, then anything it depends on will emit to the preamble first, and then it copies the
+ // text buffer into the preamble.
+ TextBuffer str_buf;
+ Line(&str_buf) << "struct " << StructName(str) << " {";
+
+ bool is_host_shareable = str->IsHostShareable();
+
+ // Emits a `/* 0xnnnn */` byte offset comment for a struct member.
+ auto add_byte_offset_comment = [&](utils::StringStream& out, uint32_t offset) {
+ std::ios_base::fmtflags saved_flag_state(out.flags());
+ out << "/* 0x" << std::hex << std::setfill('0') << std::setw(4) << offset << " */ ";
+ out.flags(saved_flag_state);
+ };
+
+ auto add_padding = [&](uint32_t size, uint32_t msl_offset) {
+ std::string name;
+ do {
+ name = UniqueIdentifier("tint_pad");
+ } while (str->FindMember(ir_->symbols.Get(name)));
+
+ auto out = Line(&str_buf);
+ add_byte_offset_comment(out, msl_offset);
+ out << ArrayTemplateName() << "<int8_t, " << size << "> " << name << ";";
+ };
+
+ str_buf.IncrementIndent();
+
+ uint32_t msl_offset = 0;
+ for (auto* mem : str->Members()) {
+ auto out = Line(&str_buf);
+ auto mem_name = mem->Name().Name();
+ auto ir_offset = mem->Offset();
+
+ if (is_host_shareable) {
+ if (TINT_UNLIKELY(ir_offset < msl_offset)) {
+ // Unimplementable layout
+ TINT_ICE(Writer, diagnostics_) << "Structure member offset (" << ir_offset
+ << ") is behind MSL offset (" << msl_offset << ")";
+ return;
+ }
+
+ // Generate padding if required
+ if (auto padding = ir_offset - msl_offset) {
+ add_padding(padding, msl_offset);
+ msl_offset += padding;
+ }
+
+ add_byte_offset_comment(out, msl_offset);
+ }
+
+ auto* ty = mem->Type();
+
+ EmitType(out, ty);
+ out << " " << mem_name;
+
+ // Emit attributes
+ auto& attributes = mem->Attributes();
+
+ if (auto builtin = attributes.builtin) {
+ auto name = BuiltinToAttribute(builtin.value());
+ if (name.empty()) {
+ diagnostics_.add_error(diag::System::Writer, "unknown builtin");
+ return;
+ }
+ out << " [[" << name << "]]";
+ }
+
+ if (auto location = attributes.location) {
+ auto& pipeline_stage_uses = str->PipelineStageUses();
+ if (TINT_UNLIKELY(pipeline_stage_uses.size() != 1)) {
+ TINT_ICE(Writer, diagnostics_) << "invalid entry point IO struct uses";
+ return;
+ }
+
+ if (pipeline_stage_uses.count(type::PipelineStageUsage::kVertexInput)) {
+ out << " [[attribute(" + std::to_string(location.value()) + ")]]";
+ } else if (pipeline_stage_uses.count(type::PipelineStageUsage::kVertexOutput)) {
+ out << " [[user(locn" + std::to_string(location.value()) + ")]]";
+ } else if (pipeline_stage_uses.count(type::PipelineStageUsage::kFragmentInput)) {
+ out << " [[user(locn" + std::to_string(location.value()) + ")]]";
+ } else if (TINT_LIKELY(
+ pipeline_stage_uses.count(type::PipelineStageUsage::kFragmentOutput))) {
+ out << " [[color(" + std::to_string(location.value()) + ")]]";
+ } else {
+ TINT_ICE(Writer, diagnostics_) << "invalid use of location decoration";
+ return;
+ }
+ }
+
+ if (auto interpolation = attributes.interpolation) {
+ auto name = InterpolationToAttribute(interpolation->type, interpolation->sampling);
+ if (name.empty()) {
+ diagnostics_.add_error(diag::System::Writer, "unknown interpolation attribute");
+ return;
+ }
+ out << " [[" << name << "]]";
+ }
+
+ if (attributes.invariant) {
+ invariant_define_name_ = UniqueIdentifier("TINT_INVARIANT");
+ out << " " << invariant_define_name_;
+ }
+
+ out << ";";
+
+ if (is_host_shareable) {
+ // Calculate new MSL offset
+ auto size_align = MslPackedTypeSizeAndAlign(diagnostics_, ty);
+ if (TINT_UNLIKELY(msl_offset % size_align.align)) {
+ TINT_ICE(Writer, diagnostics_)
+ << "Misaligned MSL structure member " << mem_name << " : " << ty->FriendlyName()
+ << " offset: " << msl_offset << " align: " << size_align.align;
+ return;
+ }
+ msl_offset += size_align.size;
+ }
+ }
+
+ if (is_host_shareable && str->Size() != msl_offset) {
+ add_padding(str->Size() - msl_offset, msl_offset);
+ }
+
+ str_buf.DecrementIndent();
+ Line(&str_buf) << "};";
+
+ preamble_buffer_.Append(str_buf);
+}
+
} // namespace tint::writer::msl
diff --git a/src/tint/writer/msl/ir/generator_impl_ir.h b/src/tint/writer/msl/ir/generator_impl_ir.h
index 200c24b..70e5d15 100644
--- a/src/tint/writer/msl/ir/generator_impl_ir.h
+++ b/src/tint/writer/msl/ir/generator_impl_ir.h
@@ -16,6 +16,7 @@
#define SRC_TINT_WRITER_MSL_IR_GENERATOR_IMPL_IR_H_
#include <string>
+#include <unordered_set>
#include "src/tint/diagnostic/diagnostic.h"
#include "src/tint/ir/module.h"
@@ -44,6 +45,11 @@
/// @param ty the type to emit
void EmitType(utils::StringStream& out, const type::Type* ty);
+ /// Handles generating a struct declaration. If the structure has already been emitted, then
+ /// this function will simply return without emitting anything.
+ /// @param str the struct to generate
+ void EmitStructType(const type::Struct* str);
+
/// Handles generating a address space
/// @param out the output of the type stream
/// @param sc the address space to generate
@@ -55,6 +61,13 @@
/// Unique name of the tint_array<T, N> template.
/// Non-empty only if the template has been generated.
std::string array_template_name_;
+
+ private:
+ /// Unique name of the 'TINT_INVARIANT' preprocessor define.
+ /// Non-empty only if an invariant attribute has been generated.
+ std::string invariant_define_name_;
+
+ std::unordered_set<const type::Struct*> emitted_structs_;
};
} // namespace tint::writer::msl
diff --git a/src/tint/writer/msl/ir/generator_impl_ir_type_test.cc b/src/tint/writer/msl/ir/generator_impl_ir_type_test.cc
index 237de35..3e82c1b 100644
--- a/src/tint/writer/msl/ir/generator_impl_ir_type_test.cc
+++ b/src/tint/writer/msl/ir/generator_impl_ir_type_test.cc
@@ -21,6 +21,7 @@
#include "src/tint/type/multisampled_texture.h"
#include "src/tint/type/sampled_texture.h"
#include "src/tint/type/storage_texture.h"
+#include "src/tint/type/struct.h"
#include "src/tint/utils/string.h"
#include "src/tint/writer/msl/ir/test_helper_ir.h"
@@ -30,60 +31,6 @@
using namespace tint::builtin::fluent_types; // NOLINT
using namespace tint::number_suffixes; // NOLINT
-// void FormatMSLField(utils::StringStream& out,
-// const char* addr,
-// const char* type,
-// size_t array_count,
-// const char* name) {
-// out << " /* " << std::string(addr) << " */ ";
-// if (array_count == 0) {
-// out << type << " ";
-// } else {
-// out << "tint_array<" << type << ", " << std::to_string(array_count) << "> ";
-// }
-// out << name << ";\n";
-// }
-//
-// #define CHECK_TYPE_SIZE_AND_ALIGN(TYPE, SIZE, ALIGN) \ //
-// static_assert(sizeof(TYPE) == SIZE, "Bad type size"); \ //
-// static_assert(alignof(TYPE) == ALIGN, "Bad type alignment")
-//
-// // Declare C++ types that match the size and alignment of the types of the same
-// // name in MSL.
-// #define DECLARE_TYPE(NAME, SIZE, ALIGN) \ //
-// struct alignas(ALIGN) NAME { \ //
-// uint8_t _[SIZE]; \ //
-// }; \ //
-// CHECK_TYPE_SIZE_AND_ALIGN(NAME, SIZE, ALIGN)
-//
-// // Size and alignments taken from the MSL spec:
-// // https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf
-// DECLARE_TYPE(float2, 8, 8);
-// DECLARE_TYPE(float3, 12, 4);
-// DECLARE_TYPE(float4, 16, 16);
-// DECLARE_TYPE(float2x2, 16, 8);
-// DECLARE_TYPE(float2x3, 32, 16);
-// DECLARE_TYPE(float2x4, 32, 16);
-// DECLARE_TYPE(float3x2, 24, 8);
-// DECLARE_TYPE(float3x3, 48, 16);
-// DECLARE_TYPE(float3x4, 48, 16);
-// DECLARE_TYPE(float4x2, 32, 8);
-// DECLARE_TYPE(float4x3, 64, 16);
-// DECLARE_TYPE(float4x4, 64, 16);
-// DECLARE_TYPE(half2, 4, 4);
-// DECLARE_TYPE(packed_half3, 6, 2);
-// DECLARE_TYPE(half4, 8, 8);
-// DECLARE_TYPE(half2x2, 8, 4);
-// DECLARE_TYPE(half2x3, 16, 8);
-// DECLARE_TYPE(half2x4, 16, 8);
-// DECLARE_TYPE(half3x2, 12, 4);
-// DECLARE_TYPE(half3x3, 24, 8);
-// DECLARE_TYPE(half3x4, 24, 8);
-// DECLARE_TYPE(half4x2, 16, 4);
-// DECLARE_TYPE(half4x3, 32, 8);
-// DECLARE_TYPE(half4x4, 32, 8);
-// using uint = unsigned int;
-
TEST_F(MslGeneratorImplIrTest, EmitType_Array) {
generator_.EmitType(generator_.Line(), ty.array<bool, 4>());
ASSERT_TRUE(generator_.Diagnostics().empty()) << generator_.Diagnostics().str();
@@ -266,495 +213,600 @@
MslAddressSpaceData{builtin::AddressSpace::kStorage, "device"},
MslAddressSpaceData{builtin::AddressSpace::kUniform, "constant"}));
-// TEST_F(MslGeneratorImplTest, EmitType_Struct) {
-// auto* s = Structure("S", utils::Vector{
-// Member("a", ty.i32()),
-// Member("b", ty.f32()),
-// });
-//
-// GeneratorImpl& gen = Build();
-//
-// utils::StringStream out;
-// ASSERT_TRUE(gen.EmitType(out, program->TypeOf(s), "")) << gen.Diagnostics();
-// EXPECT_EQ(out.str(), "S");
-// }
-//
-// TEST_F(MslGeneratorImplTest, EmitType_StructDecl) {
-// auto* s = Structure("S", utils::Vector{
-// Member("a", ty.i32()),
-// Member("b", ty.f32()),
-// });
-//
-// GeneratorImpl& gen = Build();
-//
-// TextGenerator::TextBuffer buf;
-// auto* str = program->TypeOf(s)->As<type::Struct>();
-// ASSERT_TRUE(gen.EmitStructType(&buf, str)) << gen.Diagnostics();
-// EXPECT_EQ(buf.String(), R"(struct S {
-// int a;
-// float b;
-// };
-// )");
-// }
-//
-// TEST_F(MslGeneratorImplTest, EmitType_Struct_Layout_NonComposites) {
-// auto* s = Structure(
-// "S", utils::Vector{
-// Member("a", ty.i32(), utils::Vector{MemberSize(32_a)}),
-// Member("b", ty.f32(), utils::Vector{MemberAlign(128_i), MemberSize(128_a)}),
-// Member("c", ty.vec2<f32>()),
-// Member("d", ty.u32()),
-// Member("e", ty.vec3<f32>()),
-// Member("f", ty.u32()),
-// Member("g", ty.vec4<f32>()),
-// Member("h", ty.u32()),
-// Member("i", ty.mat2x2<f32>()),
-// Member("j", ty.u32()),
-// Member("k", ty.mat2x3<f32>()),
-// Member("l", ty.u32()),
-// Member("m", ty.mat2x4<f32>()),
-// Member("n", ty.u32()),
-// Member("o", ty.mat3x2<f32>()),
-// Member("p", ty.u32()),
-// Member("q", ty.mat3x3<f32>()),
-// Member("r", ty.u32()),
-// Member("s", ty.mat3x4<f32>()),
-// Member("t", ty.u32()),
-// Member("u", ty.mat4x2<f32>()),
-// Member("v", ty.u32()),
-// Member("w", ty.mat4x3<f32>()),
-// Member("x", ty.u32()),
-// Member("y", ty.mat4x4<f32>()),
-// Member("z", ty.f32()),
-// });
-//
-// ast::Type type = GlobalVar("G", ty.Of(s), builtin::AddressSpace::kStorage,
-// builtin::Access::kRead, Binding(0_a), Group(0_a))
-// ->type;
-//
-// GeneratorImpl& gen = Build();
-//
-// TextGenerator::TextBuffer buf;
-// auto* str = program->TypeOf(type)->As<type::Struct>();
-// ASSERT_TRUE(gen.EmitStructType(&buf, str)) << gen.Diagnostics();
-//
-// // ALL_FIELDS() calls the macro FIELD(ADDR, TYPE, ARRAY_COUNT, NAME)
-// // for each field of the structure s.
-// #define ALL_FIELDS() \ //
-// FIELD(0x0000, int, 0, a) \ //
-// FIELD(0x0004, int8_t, 124, tint_pad) \ //
-// FIELD(0x0080, float, 0, b) \ //
-// FIELD(0x0084, int8_t, 124, tint_pad_1) \ //
-// FIELD(0x0100, float2, 0, c) \ //
-// FIELD(0x0108, uint, 0, d) \ //
-// FIELD(0x010c, int8_t, 4, tint_pad_2) \ //
-// FIELD(0x0110, float3, 0, e) \ //
-// FIELD(0x011c, uint, 0, f) \ //
-// FIELD(0x0120, float4, 0, g) \ //
-// FIELD(0x0130, uint, 0, h) \ //
-// FIELD(0x0134, int8_t, 4, tint_pad_3) \ //
-// FIELD(0x0138, float2x2, 0, i) \ //
-// FIELD(0x0148, uint, 0, j) \ //
-// FIELD(0x014c, int8_t, 4, tint_pad_4) \ //
-// FIELD(0x0150, float2x3, 0, k) \ //
-// FIELD(0x0170, uint, 0, l) \ //
-// FIELD(0x0174, int8_t, 12, tint_pad_5) \ //
-// FIELD(0x0180, float2x4, 0, m) \ //
-// FIELD(0x01a0, uint, 0, n) \ //
-// FIELD(0x01a4, int8_t, 4, tint_pad_6) \ //
-// FIELD(0x01a8, float3x2, 0, o) \ //
-// FIELD(0x01c0, uint, 0, p) \ //
-// FIELD(0x01c4, int8_t, 12, tint_pad_7) \ //
-// FIELD(0x01d0, float3x3, 0, q) \ //
-// FIELD(0x0200, uint, 0, r) \ //
-// FIELD(0x0204, int8_t, 12, tint_pad_8) \ //
-// FIELD(0x0210, float3x4, 0, s) \ //
-// FIELD(0x0240, uint, 0, t) \ //
-// FIELD(0x0244, int8_t, 4, tint_pad_9) \ //
-// FIELD(0x0248, float4x2, 0, u) \ //
-// FIELD(0x0268, uint, 0, v) \ //
-// FIELD(0x026c, int8_t, 4, tint_pad_10) \ //
-// FIELD(0x0270, float4x3, 0, w) \ //
-// FIELD(0x02b0, uint, 0, x) \ //
-// FIELD(0x02b4, int8_t, 12, tint_pad_11) \ //
-// FIELD(0x02c0, float4x4, 0, y) \ //
-// FIELD(0x0300, float, 0, z) \ //
-// FIELD(0x0304, int8_t, 124, tint_pad_12)
-//
-// // Check that the generated string is as expected.
-// utils::StringStream expect;
-// expect << "struct S {\n";
-// #define FIELD(ADDR, TYPE, ARRAY_COUNT, NAME) \ //
-// FormatMSLField(expect, #ADDR, #TYPE, ARRAY_COUNT, #NAME);
-// ALL_FIELDS()
-// #undef FIELD
-// expect << "};\n";
-// EXPECT_EQ(buf.String(), expect.str());
-//
-// // 1.4 Metal and C++14
-// // The Metal programming language is a C++14-based Specification with
-// // extensions and restrictions. Refer to the C++14 Specification (also known
-// // as the ISO/IEC JTC1/SC22/WG21 N4431 Language Specification) for a detailed
-// // description of the language grammar.
-// //
-// // Tint is written in C++14, so use the compiler to verify the generated
-// // layout is as expected for C++14 / MSL.
-// {
-// struct S {
-// #define FIELD(ADDR, TYPE, ARRAY_COUNT, NAME) std::array<TYPE, ARRAY_COUNT ? ARRAY_COUNT : 1>
-// NAME;
-// ALL_FIELDS()
-// #undef FIELD
-// };
-//
-// #define FIELD(ADDR, TYPE, ARRAY_COUNT, NAME) \ //
-// EXPECT_EQ(ADDR, static_cast<int>(offsetof(S, NAME))) << "Field " << #NAME;
-// ALL_FIELDS()
-// #undef FIELD
-// }
-// #undef ALL_FIELDS
-// }
-//
-// TEST_F(MslGeneratorImplTest, EmitType_Struct_Layout_Structures) {
-// // inner_x: size(1024), align(512)
-// auto* inner_x =
-// Structure("inner_x", utils::Vector{
-// Member("a", ty.i32()),
-// Member("b", ty.f32(), utils::Vector{MemberAlign(512_i)}),
-// });
-//
-// // inner_y: size(516), align(4)
-// auto* inner_y =
-// Structure("inner_y", utils::Vector{
-// Member("a", ty.i32(), utils::Vector{MemberSize(512_a)}),
-// Member("b", ty.f32()),
-// });
-//
-// auto* s = Structure("S", utils::Vector{
-// Member("a", ty.i32()),
-// Member("b", ty.Of(inner_x)),
-// Member("c", ty.f32()),
-// Member("d", ty.Of(inner_y)),
-// Member("e", ty.f32()),
-// });
-//
-// ast::Type type = GlobalVar("G", ty.Of(s), builtin::AddressSpace::kStorage,
-// builtin::Access::kRead, Binding(0_a), Group(0_a))
-// ->type;
-//
-// GeneratorImpl& gen = Build();
-//
-// TextGenerator::TextBuffer buf;
-// auto* str = program->TypeOf(type)->As<type::Struct>();
-// ASSERT_TRUE(gen.EmitStructType(&buf, str)) << gen.Diagnostics();
-//
-// // ALL_FIELDS() calls the macro FIELD(ADDR, TYPE, ARRAY_COUNT, NAME)
-// // for each field of the structure s.
-// #define ALL_FIELDS() \ //
-// FIELD(0x0000, int, 0, a) \ //
-// FIELD(0x0004, int8_t, 508, tint_pad) \ //
-// FIELD(0x0200, inner_x, 0, b) \ //
-// FIELD(0x0600, float, 0, c) \ //
-// FIELD(0x0604, inner_y, 0, d) \ //
-// FIELD(0x0808, float, 0, e) \ //
-// FIELD(0x080c, int8_t, 500, tint_pad_1)
-//
-// // Check that the generated string is as expected.
-// utils::StringStream expect;
-// expect << "struct S {\n";
-// #define FIELD(ADDR, TYPE, ARRAY_COUNT, NAME) \ //
-// FormatMSLField(expect, #ADDR, #TYPE, ARRAY_COUNT, #NAME);
-// ALL_FIELDS()
-// #undef FIELD
-// expect << "};\n";
-// EXPECT_EQ(buf.String(), expect.str());
-//
-// // 1.4 Metal and C++14
-// // The Metal programming language is a C++14-based Specification with
-// // extensions and restrictions. Refer to the C++14 Specification (also known
-// // as the ISO/IEC JTC1/SC22/WG21 N4431 Language Specification) for a detailed
-// // description of the language grammar.
-// //
-// // Tint is written in C++14, so use the compiler to verify the generated
-// // layout is as expected for C++14 / MSL.
-// {
-// struct inner_x {
-// uint32_t a;
-// alignas(512) float b;
-// };
-// CHECK_TYPE_SIZE_AND_ALIGN(inner_x, 1024, 512);
-//
-// struct inner_y {
-// uint32_t a[128];
-// float b;
-// };
-// CHECK_TYPE_SIZE_AND_ALIGN(inner_y, 516, 4);
-//
-// struct S {
-// #define FIELD(ADDR, TYPE, ARRAY_COUNT, NAME) std::array<TYPE, ARRAY_COUNT ? ARRAY_COUNT : 1>
-// NAME;
-// ALL_FIELDS()
-// #undef FIELD
-// };
-//
-// #define FIELD(ADDR, TYPE, ARRAY_COUNT, NAME) \ //
-// EXPECT_EQ(ADDR, static_cast<int>(offsetof(S, NAME))) << "Field " << #NAME;
-// ALL_FIELDS()
-// #undef FIELD
-// }
-//
-// #undef ALL_FIELDS
-// }
-//
-// TEST_F(MslGeneratorImplTest, EmitType_Struct_Layout_ArrayDefaultStride) {
-// // inner: size(1024), align(512)
-// auto* inner = Structure("inner", utils::Vector{
-// Member("a", ty.i32()),
-// Member("b", ty.f32(),
-// utils::Vector{MemberAlign(512_i)}),
-// });
-//
-// // array_x: size(28), align(4)
-// auto array_x = ty.array<f32, 7>();
-//
-// // array_y: size(4096), align(512)
-// auto array_y = ty.array(ty.Of(inner), 4_u);
-//
-// // array_z: size(4), align(4)
-// auto array_z = ty.array<f32>();
-//
-// auto* s = Structure("S", utils::Vector{
-// Member("a", ty.i32()),
-// Member("b", array_x),
-// Member("c", ty.f32()),
-// Member("d", array_y),
-// Member("e", ty.f32()),
-// Member("f", array_z),
-// });
-//
-// ast::Type type = GlobalVar("G", ty.Of(s), builtin::AddressSpace::kStorage,
-// builtin::Access::kRead, Binding(0_a), Group(0_a))
-// ->type;
-//
-// GeneratorImpl& gen = Build();
-//
-// TextGenerator::TextBuffer buf;
-// auto* str = program->TypeOf(type)->As<type::Struct>();
-// ASSERT_TRUE(gen.EmitStructType(&buf, str)) << gen.Diagnostics();
-//
-// // ALL_FIELDS() calls the macro FIELD(ADDR, TYPE, ARRAY_COUNT, NAME)
-// // for each field of the structure s.
-// #define ALL_FIELDS() \ //
-// FIELD(0x0000, int, 0, a) \ //
-// FIELD(0x0004, float, 7, b) \ //
-// FIELD(0x0020, float, 0, c) \ //
-// FIELD(0x0024, int8_t, 476, tint_pad) \ //
-// FIELD(0x0200, inner, 4, d) \ //
-// FIELD(0x1200, float, 0, e) \ //
-// FIELD(0x1204, float, 1, f) \ //
-// FIELD(0x1208, int8_t, 504, tint_pad_1)
-//
-// // Check that the generated string is as expected.
-// utils::StringStream expect;
-// expect << "struct S {\n";
-// #define FIELD(ADDR, TYPE, ARRAY_COUNT, NAME) \ //
-// FormatMSLField(expect, #ADDR, #TYPE, ARRAY_COUNT, #NAME);
-// ALL_FIELDS()
-// #undef FIELD
-// expect << "};\n";
-// EXPECT_EQ(buf.String(), expect.str());
-//
-// // 1.4 Metal and C++14
-// // The Metal programming language is a C++14-based Specification with
-// // extensions and restrictions. Refer to the C++14 Specification (also known
-// // as the ISO/IEC JTC1/SC22/WG21 N4431 Language Specification) for a detailed
-// // description of the language grammar.
-// //
-// // Tint is written in C++14, so use the compiler to verify the generated
-// // layout is as expected for C++14 / MSL.
-// {
-// struct inner {
-// uint32_t a;
-// alignas(512) float b;
-// };
-// CHECK_TYPE_SIZE_AND_ALIGN(inner, 1024, 512);
-//
-// // array_x: size(28), align(4)
-// using array_x = std::array<float, 7>;
-// CHECK_TYPE_SIZE_AND_ALIGN(array_x, 28, 4);
-//
-// // array_y: size(4096), align(512)
-// using array_y = std::array<inner, 4>;
-// CHECK_TYPE_SIZE_AND_ALIGN(array_y, 4096, 512);
-//
-// // array_z: size(4), align(4)
-// using array_z = std::array<float, 1>;
-// CHECK_TYPE_SIZE_AND_ALIGN(array_z, 4, 4);
-//
-// struct S {
-// #define FIELD(ADDR, TYPE, ARRAY_COUNT, NAME) std::array<TYPE, ARRAY_COUNT ? ARRAY_COUNT : 1>
-// NAME;
-// ALL_FIELDS()
-// #undef FIELD
-// };
-//
-// #define FIELD(ADDR, TYPE, ARRAY_COUNT, NAME) \ //
-// EXPECT_EQ(ADDR, static_cast<int>(offsetof(S, NAME))) << "Field " << #NAME;
-// ALL_FIELDS()
-// #undef FIELD
-// }
-//
-// #undef ALL_FIELDS
-// }
-//
-// TEST_F(MslGeneratorImplTest, EmitType_Struct_Layout_ArrayVec3DefaultStride) {
-// // array: size(64), align(16)
-// auto array = ty.array<vec3<f32>, 4>();
-//
-// auto* s = Structure("S", utils::Vector{
-// Member("a", ty.i32()),
-// Member("b", array),
-// Member("c", ty.i32()),
-// });
-//
-// ast::Type type = GlobalVar("G", ty.Of(s), builtin::AddressSpace::kStorage,
-// builtin::Access::kRead, Binding(0_a), Group(0_a))
-// ->type;
-//
-// GeneratorImpl& gen = Build();
-//
-// TextGenerator::TextBuffer buf;
-// auto* str = program->TypeOf(type)->As<type::Struct>();
-// ASSERT_TRUE(gen.EmitStructType(&buf, str)) << gen.Diagnostics();
-//
-// // ALL_FIELDS() calls the macro FIELD(ADDR, TYPE, ARRAY_COUNT, NAME)
-// // for each field of the structure s.
-// #define ALL_FIELDS() \ //
-// FIELD(0x0000, int, 0, a) \ //
-// FIELD(0x0004, int8_t, 12, tint_pad) \ //
-// FIELD(0x0010, float3, 4, b) \ //
-// FIELD(0x0050, int, 0, c) \ //
-// FIELD(0x0054, int8_t, 12, tint_pad_1)
-//
-// // Check that the generated string is as expected.
-// utils::StringStream expect;
-// expect << "struct S {\n";
-// #define FIELD(ADDR, TYPE, ARRAY_COUNT, NAME) \ //
-// FormatMSLField(expect, #ADDR, #TYPE, ARRAY_COUNT, #NAME);
-// ALL_FIELDS()
-// #undef FIELD
-// expect << "};\n";
-// EXPECT_EQ(buf.String(), expect.str());
-// }
-//
-// TEST_F(MslGeneratorImplTest, AttemptTintPadSymbolCollision) {
-// auto* s = Structure("S", utils::Vector{
-// // uses symbols tint_pad_[0..9] and tint_pad_[20..35]
-// Member("tint_pad_2", ty.i32(),
-// utils::Vector{MemberSize(32_a)}), Member("tint_pad_20",
-// ty.f32(),
-// utils::Vector{MemberAlign(128_i),
-// MemberSize(128_u)}),
-// Member("tint_pad_33", ty.vec2<f32>()),
-// Member("tint_pad_1", ty.u32()),
-// Member("tint_pad_3", ty.vec3<f32>()),
-// Member("tint_pad_7", ty.u32()),
-// Member("tint_pad_25", ty.vec4<f32>()),
-// Member("tint_pad_5", ty.u32()),
-// Member("tint_pad_27", ty.mat2x2<f32>()),
-// Member("tint_pad_24", ty.u32()),
-// Member("tint_pad_23", ty.mat2x3<f32>()),
-// Member("tint_pad", ty.u32()),
-// Member("tint_pad_8", ty.mat2x4<f32>()),
-// Member("tint_pad_26", ty.u32()),
-// Member("tint_pad_29", ty.mat3x2<f32>()),
-// Member("tint_pad_6", ty.u32()),
-// Member("tint_pad_22", ty.mat3x3<f32>()),
-// Member("tint_pad_32", ty.u32()),
-// Member("tint_pad_34", ty.mat3x4<f32>()),
-// Member("tint_pad_35", ty.u32()),
-// Member("tint_pad_30", ty.mat4x2<f32>()),
-// Member("tint_pad_9", ty.u32()),
-// Member("tint_pad_31", ty.mat4x3<f32>()),
-// Member("tint_pad_28", ty.u32()),
-// Member("tint_pad_4", ty.mat4x4<f32>()),
-// Member("tint_pad_21", ty.f32()),
-// });
-//
-// ast::Type type = GlobalVar("G", ty.Of(s), builtin::AddressSpace::kStorage,
-// builtin::Access::kRead, Binding(0_a), Group(0_a))
-// ->type;
-//
-// GeneratorImpl& gen = Build();
-//
-// TextGenerator::TextBuffer buf;
-// auto* str = program->TypeOf(type)->As<type::Struct>();
-// ASSERT_TRUE(gen.EmitStructType(&buf, str)) << gen.Diagnostics();
-// EXPECT_EQ(buf.String(), R"(struct S {
-// /* 0x0000 */ int tint_pad_2;
-// /* 0x0004 */ tint_array<int8_t, 124> tint_pad_10;
-// /* 0x0080 */ float tint_pad_20;
-// /* 0x0084 */ tint_array<int8_t, 124> tint_pad_11;
-// /* 0x0100 */ float2 tint_pad_33;
-// /* 0x0108 */ uint tint_pad_1;
-// /* 0x010c */ tint_array<int8_t, 4> tint_pad_12;
-// /* 0x0110 */ float3 tint_pad_3;
-// /* 0x011c */ uint tint_pad_7;
-// /* 0x0120 */ float4 tint_pad_25;
-// /* 0x0130 */ uint tint_pad_5;
-// /* 0x0134 */ tint_array<int8_t, 4> tint_pad_13;
-// /* 0x0138 */ float2x2 tint_pad_27;
-// /* 0x0148 */ uint tint_pad_24;
-// /* 0x014c */ tint_array<int8_t, 4> tint_pad_14;
-// /* 0x0150 */ float2x3 tint_pad_23;
-// /* 0x0170 */ uint tint_pad;
-// /* 0x0174 */ tint_array<int8_t, 12> tint_pad_15;
-// /* 0x0180 */ float2x4 tint_pad_8;
-// /* 0x01a0 */ uint tint_pad_26;
-// /* 0x01a4 */ tint_array<int8_t, 4> tint_pad_16;
-// /* 0x01a8 */ float3x2 tint_pad_29;
-// /* 0x01c0 */ uint tint_pad_6;
-// /* 0x01c4 */ tint_array<int8_t, 12> tint_pad_17;
-// /* 0x01d0 */ float3x3 tint_pad_22;
-// /* 0x0200 */ uint tint_pad_32;
-// /* 0x0204 */ tint_array<int8_t, 12> tint_pad_18;
-// /* 0x0210 */ float3x4 tint_pad_34;
-// /* 0x0240 */ uint tint_pad_35;
-// /* 0x0244 */ tint_array<int8_t, 4> tint_pad_19;
-// /* 0x0248 */ float4x2 tint_pad_30;
-// /* 0x0268 */ uint tint_pad_9;
-// /* 0x026c */ tint_array<int8_t, 4> tint_pad_36;
-// /* 0x0270 */ float4x3 tint_pad_31;
-// /* 0x02b0 */ uint tint_pad_28;
-// /* 0x02b4 */ tint_array<int8_t, 12> tint_pad_37;
-// /* 0x02c0 */ float4x4 tint_pad_4;
-// /* 0x0300 */ float tint_pad_21;
-// /* 0x0304 */ tint_array<int8_t, 124> tint_pad_38;
-// };
-// )");
-// }
-//
-// TEST_F(MslGeneratorImplTest, EmitType_Struct_WithAttribute) {
-// auto* s = Structure("S", utils::Vector{
-// Member("a", ty.i32()),
-// Member("b", ty.f32()),
-// });
-//
-// ast::Type type = GlobalVar("G", ty.Of(s), builtin::AddressSpace::kStorage,
-// builtin::Access::kRead, Binding(0_a), Group(0_a))
-// ->type;
-//
-// GeneratorImpl& gen = Build();
-//
-// TextGenerator::TextBuffer buf;
-// auto* str = program->TypeOf(type)->As<type::Struct>();
-// ASSERT_TRUE(gen.EmitStructType(&buf, str)) << gen.Diagnostics();
-// EXPECT_EQ(buf.String(), R"(struct S {
-// /* 0x0000 */ int a;
-// /* 0x0004 */ float b;
-// };
-// )");
-// }
+TEST_F(MslGeneratorImplIrTest, EmitType_Struct) {
+ auto* s = ty.Struct(mod.symbols.New("S"), {
+ {mod.symbols.Register("a"), ty.i32()},
+ {mod.symbols.Register("b"), ty.f32()},
+ });
+ generator_.EmitType(generator_.Line(), s);
+ ASSERT_TRUE(generator_.Diagnostics().empty()) << generator_.Diagnostics().str();
+ EXPECT_STREQ(std::string(utils::TrimSpace(generator_.Result())).c_str(), R"(struct S {
+ int a;
+ float b;
+};
+
+S)");
+}
+
+TEST_F(MslGeneratorImplIrTest, EmitType_Struct_Dedup) {
+ auto* s = ty.Struct(mod.symbols.New("S"), {
+ {mod.symbols.Register("a"), ty.i32()},
+ {mod.symbols.Register("b"), ty.f32()},
+ });
+ generator_.EmitType(generator_.Line(), s);
+ generator_.EmitType(generator_.Line(), s);
+ ASSERT_TRUE(generator_.Diagnostics().empty()) << generator_.Diagnostics().str();
+ EXPECT_STREQ(std::string(utils::TrimSpace(generator_.Result())).c_str(), R"(struct S {
+ int a;
+ float b;
+};
+
+S
+S)");
+}
+
+void FormatMSLField(utils::StringStream& out,
+ const char* addr,
+ const char* type,
+ size_t array_count,
+ const char* name) {
+ out << " /* " << std::string(addr) << " */ ";
+ if (array_count == 0) {
+ out << type << " ";
+ } else {
+ out << "tint_array<" << type << ", " << std::to_string(array_count) << "> ";
+ }
+ out << name << ";\n";
+}
+
+#define CHECK_TYPE_SIZE_AND_ALIGN(TYPE, SIZE, ALIGN) \
+ static_assert(sizeof(TYPE) == SIZE, "Bad type size"); \
+ static_assert(alignof(TYPE) == ALIGN, "Bad type alignment")
+
+// Declare C++ types that match the size and alignment of the types of the same
+// name in MSL.
+#define DECLARE_TYPE(NAME, SIZE, ALIGN) \
+ struct alignas(ALIGN) NAME { \
+ uint8_t _[SIZE]; \
+ }; \
+ CHECK_TYPE_SIZE_AND_ALIGN(NAME, SIZE, ALIGN)
+
+// Size and alignments taken from the MSL spec:
+// https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf
+DECLARE_TYPE(float2, 8, 8);
+DECLARE_TYPE(float3, 12, 4);
+DECLARE_TYPE(float4, 16, 16);
+DECLARE_TYPE(float2x2, 16, 8);
+DECLARE_TYPE(float2x3, 32, 16);
+DECLARE_TYPE(float2x4, 32, 16);
+DECLARE_TYPE(float3x2, 24, 8);
+DECLARE_TYPE(float3x3, 48, 16);
+DECLARE_TYPE(float3x4, 48, 16);
+DECLARE_TYPE(float4x2, 32, 8);
+DECLARE_TYPE(float4x3, 64, 16);
+DECLARE_TYPE(float4x4, 64, 16);
+DECLARE_TYPE(half2, 4, 4);
+DECLARE_TYPE(packed_half3, 6, 2);
+DECLARE_TYPE(half4, 8, 8);
+DECLARE_TYPE(half2x2, 8, 4);
+DECLARE_TYPE(half2x3, 16, 8);
+DECLARE_TYPE(half2x4, 16, 8);
+DECLARE_TYPE(half3x2, 12, 4);
+DECLARE_TYPE(half3x3, 24, 8);
+DECLARE_TYPE(half3x4, 24, 8);
+DECLARE_TYPE(half4x2, 16, 4);
+DECLARE_TYPE(half4x3, 32, 8);
+DECLARE_TYPE(half4x4, 32, 8);
+using uint = unsigned int;
+
+struct MemberData {
+ Symbol name;
+ const type::Type* type;
+ uint32_t size = 0;
+ uint32_t align = 0;
+};
+type::Struct* MkStruct(ir::Module& mod,
+ type::Manager& ty,
+ std::string_view name,
+ utils::VectorRef<MemberData> data) {
+ utils::Vector<const type::StructMember*, 26> members;
+ uint32_t align = 0;
+ uint32_t size = 0;
+ for (uint32_t i = 0; i < data.Length(); ++i) {
+ auto& d = data[i];
+
+ uint32_t mem_align = d.align == 0 ? d.type->Align() : d.align;
+ uint32_t mem_size = d.size == 0 ? d.type->Size() : d.size;
+
+ uint32_t offset = utils::RoundUp(mem_align, size);
+ members.Push(ty.Get<type::StructMember>(d.name, d.type, i, offset, mem_align, mem_size,
+ type::StructMemberAttributes{}));
+
+ align = std::max(align, mem_align);
+ size = offset + mem_size;
+ }
+
+ return ty.Get<type::Struct>(mod.symbols.New(name), std::move(members), align,
+ utils::RoundUp(align, size), size);
+}
+
+TEST_F(MslGeneratorImplIrTest, EmitType_Struct_Layout_NonComposites) {
+ utils::Vector<MemberData, 26> data = {{mod.symbols.Register("a"), ty.i32(), 32}, //
+ {mod.symbols.Register("b"), ty.f32(), 128, 128}, //
+ {mod.symbols.Register("c"), ty.vec2<f32>()}, //
+ {mod.symbols.Register("d"), ty.u32()}, //
+ {mod.symbols.Register("e"), ty.vec3<f32>()}, //
+ {mod.symbols.Register("f"), ty.u32()}, //
+ {mod.symbols.Register("g"), ty.vec4<f32>()}, //
+ {mod.symbols.Register("h"), ty.u32()}, //
+ {mod.symbols.Register("i"), ty.mat2x2<f32>()}, //
+ {mod.symbols.Register("j"), ty.u32()}, //
+ {mod.symbols.Register("k"), ty.mat2x3<f32>()}, //
+ {mod.symbols.Register("l"), ty.u32()}, //
+ {mod.symbols.Register("m"), ty.mat2x4<f32>()}, //
+ {mod.symbols.Register("n"), ty.u32()}, //
+ {mod.symbols.Register("o"), ty.mat3x2<f32>()}, //
+ {mod.symbols.Register("p"), ty.u32()}, //
+ {mod.symbols.Register("q"), ty.mat3x3<f32>()}, //
+ {mod.symbols.Register("r"), ty.u32()}, //
+ {mod.symbols.Register("s"), ty.mat3x4<f32>()}, //
+ {mod.symbols.Register("t"), ty.u32()}, //
+ {mod.symbols.Register("u"), ty.mat4x2<f32>()}, //
+ {mod.symbols.Register("v"), ty.u32()}, //
+ {mod.symbols.Register("w"), ty.mat4x3<f32>()}, //
+ {mod.symbols.Register("x"), ty.u32()}, //
+ {mod.symbols.Register("y"), ty.mat4x4<f32>()}, //
+ {mod.symbols.Register("z"), ty.f32()}};
+
+ auto* s = MkStruct(mod, ty, "S", data);
+ s->AddUsage(builtin::AddressSpace::kStorage);
+
+ // ALL_FIELDS() calls the macro FIELD(ADDR, TYPE, ARRAY_COUNT, NAME)
+ // for each field of the structure s.
+#define ALL_FIELDS() \
+ FIELD(0x0000, int, 0, a) \
+ FIELD(0x0004, int8_t, 124, tint_pad) \
+ FIELD(0x0080, float, 0, b) \
+ FIELD(0x0084, int8_t, 124, tint_pad_1) \
+ FIELD(0x0100, float2, 0, c) \
+ FIELD(0x0108, uint, 0, d) \
+ FIELD(0x010c, int8_t, 4, tint_pad_2) \
+ FIELD(0x0110, float3, 0, e) \
+ FIELD(0x011c, uint, 0, f) \
+ FIELD(0x0120, float4, 0, g) \
+ FIELD(0x0130, uint, 0, h) \
+ FIELD(0x0134, int8_t, 4, tint_pad_3) \
+ FIELD(0x0138, float2x2, 0, i) \
+ FIELD(0x0148, uint, 0, j) \
+ FIELD(0x014c, int8_t, 4, tint_pad_4) \
+ FIELD(0x0150, float2x3, 0, k) \
+ FIELD(0x0170, uint, 0, l) \
+ FIELD(0x0174, int8_t, 12, tint_pad_5) \
+ FIELD(0x0180, float2x4, 0, m) \
+ FIELD(0x01a0, uint, 0, n) \
+ FIELD(0x01a4, int8_t, 4, tint_pad_6) \
+ FIELD(0x01a8, float3x2, 0, o) \
+ FIELD(0x01c0, uint, 0, p) \
+ FIELD(0x01c4, int8_t, 12, tint_pad_7) \
+ FIELD(0x01d0, float3x3, 0, q) \
+ FIELD(0x0200, uint, 0, r) \
+ FIELD(0x0204, int8_t, 12, tint_pad_8) \
+ FIELD(0x0210, float3x4, 0, s) \
+ FIELD(0x0240, uint, 0, t) \
+ FIELD(0x0244, int8_t, 4, tint_pad_9) \
+ FIELD(0x0248, float4x2, 0, u) \
+ FIELD(0x0268, uint, 0, v) \
+ FIELD(0x026c, int8_t, 4, tint_pad_10) \
+ FIELD(0x0270, float4x3, 0, w) \
+ FIELD(0x02b0, uint, 0, x) \
+ FIELD(0x02b4, int8_t, 12, tint_pad_11) \
+ FIELD(0x02c0, float4x4, 0, y) \
+ FIELD(0x0300, float, 0, z) \
+ FIELD(0x0304, int8_t, 124, tint_pad_12)
+
+ // Check that the generated string is as expected.
+ utils::StringStream expect;
+ expect << R"(template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+)";
+
+ expect << "struct S {\n";
+#define FIELD(ADDR, TYPE, ARRAY_COUNT, NAME) \
+ FormatMSLField(expect, #ADDR, #TYPE, ARRAY_COUNT, #NAME);
+ ALL_FIELDS()
+#undef FIELD
+ expect << "};\n\nS";
+
+ generator_.EmitType(generator_.Line(), s);
+ ASSERT_TRUE(generator_.Diagnostics().empty()) << generator_.Diagnostics().str();
+ EXPECT_EQ(utils::TrimSpace(generator_.Result()), expect.str());
+
+ // 1.4 Metal and C++14
+ // The Metal programming language is a C++14-based Specification with
+ // extensions and restrictions. Refer to the C++14 Specification (also
+ // known as the ISO/IEC JTC1/SC22/WG21 N4431 Language Specification) for a
+ // detailed description of the language grammar.
+ //
+ // Tint is written in C++14, so use the compiler to verify the generated
+ // layout is as expected for C++14 / MSL.
+ {
+ struct S {
+#define FIELD(ADDR, TYPE, ARRAY_COUNT, NAME) std::array<TYPE, ARRAY_COUNT ? ARRAY_COUNT : 1> NAME;
+ ALL_FIELDS()
+#undef FIELD
+ };
+
+#define FIELD(ADDR, TYPE, ARRAY_COUNT, NAME) \
+ EXPECT_EQ(ADDR, static_cast<int>(offsetof(S, NAME))) << "Field " << #NAME;
+ ALL_FIELDS()
+#undef FIELD
+ }
+#undef ALL_FIELDS
+}
+
+TEST_F(MslGeneratorImplIrTest, EmitType_Struct_Layout_Structures) {
+ // inner_x: size(1024), align(512)
+ utils::Vector<MemberData, 2> inner_x_data = {{{mod.symbols.Register("a"), ty.i32()}, //
+ {mod.symbols.Register("b"), ty.f32(), 0, 512}}};
+ auto* inner_x = MkStruct(mod, ty, "inner_x", inner_x_data);
+
+ // inner_y: size(516), align(4)
+ utils::Vector<MemberData, 2> inner_y_data = {{mod.symbols.Register("a"), ty.i32(), 512},
+ {mod.symbols.Register("b"), ty.f32()}};
+
+ auto* inner_y = MkStruct(mod, ty, "inner_y", inner_y_data);
+
+ auto* s = ty.Struct(mod.symbols.New("S"), {{mod.symbols.Register("a"), ty.i32()},
+ {mod.symbols.Register("b"), inner_x},
+ {mod.symbols.Register("c"), ty.f32()},
+ {mod.symbols.Register("d"), inner_y},
+ {mod.symbols.Register("e"), ty.f32()}});
+ const_cast<type::Struct*>(s)->AddUsage(builtin::AddressSpace::kStorage);
+
+// ALL_FIELDS() calls the macro FIELD(ADDR, TYPE, ARRAY_COUNT, NAME)
+// for each field of the structure s.
+#define ALL_FIELDS() \
+ FIELD(0x0000, int, 0, a) \
+ FIELD(0x0004, int8_t, 508, tint_pad) \
+ FIELD(0x0200, inner_x, 0, b) \
+ FIELD(0x0600, float, 0, c) \
+ FIELD(0x0604, inner_y, 0, d) \
+ FIELD(0x0808, float, 0, e) \
+ FIELD(0x080c, int8_t, 500, tint_pad_1)
+
+ // Check that the generated string is as expected.
+ utils::StringStream expect;
+ expect << R"(template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+struct inner_x {
+ int a;
+ float b;
+};
+struct inner_y {
+ int a;
+ float b;
+};
+)";
+
+ expect << "struct S {\n";
+#define FIELD(ADDR, TYPE, ARRAY_COUNT, NAME) \
+ FormatMSLField(expect, #ADDR, #TYPE, ARRAY_COUNT, #NAME);
+ ALL_FIELDS()
+#undef FIELD
+ expect << "};\n\nS";
+
+ generator_.EmitType(generator_.Line(), s);
+ ASSERT_TRUE(generator_.Diagnostics().empty()) << generator_.Diagnostics().str();
+ EXPECT_EQ(utils::TrimSpace(generator_.Result()), expect.str());
+
+ // 1.4 Metal and C++14
+ // The Metal programming language is a C++14-based Specification with
+ // extensions and restrictions. Refer to the C++14 Specification (also
+ // known as the ISO/IEC JTC1/SC22/WG21 N4431 Language Specification) for a
+ // detailed description of the language grammar.
+ //
+ // Tint is written in C++14, so use the compiler to verify the generated
+ // layout is as expected for C++14 / MSL.
+ {
+ struct inner_x {
+ uint32_t a;
+ alignas(512) float b;
+ };
+ CHECK_TYPE_SIZE_AND_ALIGN(inner_x, 1024, 512);
+
+ struct inner_y {
+ uint32_t a[128];
+ float b;
+ };
+ CHECK_TYPE_SIZE_AND_ALIGN(inner_y, 516, 4);
+
+ struct S {
+#define FIELD(ADDR, TYPE, ARRAY_COUNT, NAME) std::array<TYPE, ARRAY_COUNT ? ARRAY_COUNT : 1> NAME;
+ ALL_FIELDS()
+#undef FIELD
+ };
+
+#define FIELD(ADDR, TYPE, ARRAY_COUNT, NAME) \
+ EXPECT_EQ(ADDR, static_cast<int>(offsetof(S, NAME))) << "Field " << #NAME;
+ ALL_FIELDS()
+#undef FIELD
+ }
+
+#undef ALL_FIELDS
+}
+
+TEST_F(MslGeneratorImplIrTest, EmitType_Struct_Layout_ArrayDefaultStride) {
+ // inner: size(1024), align(512)
+ utils::Vector<MemberData, 2> inner_data = {{mod.symbols.Register("a"), ty.i32()},
+ {mod.symbols.Register("b"), ty.f32(), 0, 512}};
+
+ auto* inner = MkStruct(mod, ty, "inner", inner_data);
+
+ // array_x: size(28), align(4)
+ auto array_x = ty.array<f32, 7>();
+
+ // array_y: size(4096), align(512)
+ auto array_y = ty.array(inner, 4_u);
+
+ // array_z: size(4), align(4)
+ auto array_z = ty.array<f32>();
+
+ auto* s = ty.Struct(mod.symbols.New("S"), {{mod.symbols.Register("a"), ty.i32()},
+ {mod.symbols.Register("b"), array_x},
+ {mod.symbols.Register("c"), ty.f32()},
+ {mod.symbols.Register("d"), array_y},
+ {mod.symbols.Register("e"), ty.f32()},
+ {mod.symbols.Register("f"), array_z}});
+ const_cast<type::Struct*>(s)->AddUsage(builtin::AddressSpace::kStorage);
+
+ // ALL_FIELDS() calls the macro FIELD(ADDR, TYPE, ARRAY_COUNT, NAME)
+ // for each field of the structure s.
+#define ALL_FIELDS() \
+ FIELD(0x0000, int, 0, a) \
+ FIELD(0x0004, float, 7, b) \
+ FIELD(0x0020, float, 0, c) \
+ FIELD(0x0024, int8_t, 476, tint_pad) \
+ FIELD(0x0200, inner, 4, d) \
+ FIELD(0x1200, float, 0, e) \
+ FIELD(0x1204, float, 1, f) \
+ FIELD(0x1208, int8_t, 504, tint_pad_1)
+
+ // Check that the generated string is as expected.
+ utils::StringStream expect;
+
+ expect << R"(template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+struct inner {
+ int a;
+ float b;
+};
+)";
+
+ expect << "struct S {\n";
+#define FIELD(ADDR, TYPE, ARRAY_COUNT, NAME) \
+ FormatMSLField(expect, #ADDR, #TYPE, ARRAY_COUNT, #NAME);
+ ALL_FIELDS()
+#undef FIELD
+ expect << "};\n\nS";
+
+ generator_.EmitType(generator_.Line(), s);
+ ASSERT_TRUE(generator_.Diagnostics().empty()) << generator_.Diagnostics().str();
+ EXPECT_EQ(utils::TrimSpace(generator_.Result()), expect.str());
+
+ // 1.4 Metal and C++14
+ // The Metal programming language is a C++14-based Specification with
+ // extensions and restrictions. Refer to the C++14 Specification (also
+ // known as the ISO/IEC JTC1/SC22/WG21 N4431 Language Specification) for a
+ // detailed description of the language grammar.
+ //
+ // Tint is written in C++14, so use the compiler to verify the generated
+ // layout is as expected for C++14 / MSL.
+ {
+ struct inner {
+ uint32_t a;
+ alignas(512) float b;
+ };
+ CHECK_TYPE_SIZE_AND_ALIGN(inner, 1024, 512);
+
+ // array_x: size(28), align(4)
+ using array_x = std::array<float, 7>;
+ CHECK_TYPE_SIZE_AND_ALIGN(array_x, 28, 4);
+
+ // array_y: size(4096), align(512)
+ using array_y = std::array<inner, 4>;
+ CHECK_TYPE_SIZE_AND_ALIGN(array_y, 4096, 512);
+
+ // array_z: size(4), align(4)
+ using array_z = std::array<float, 1>;
+ CHECK_TYPE_SIZE_AND_ALIGN(array_z, 4, 4);
+
+ struct S {
+#define FIELD(ADDR, TYPE, ARRAY_COUNT, NAME) std::array<TYPE, ARRAY_COUNT ? ARRAY_COUNT : 1> NAME;
+ ALL_FIELDS()
+#undef FIELD
+ };
+
+#define FIELD(ADDR, TYPE, ARRAY_COUNT, NAME) \
+ EXPECT_EQ(ADDR, static_cast<int>(offsetof(S, NAME))) << "Field " << #NAME;
+ ALL_FIELDS()
+#undef FIELD
+ }
+
+#undef ALL_FIELDS
+}
+
+TEST_F(MslGeneratorImplIrTest, EmitType_Struct_Layout_ArrayVec3DefaultStride) {
+ // array: size(64), align(16)
+ auto array = ty.array<vec3<f32>, 4>();
+
+ auto* s = ty.Struct(mod.symbols.New("S"), {
+ {mod.symbols.Register("a"), ty.i32()},
+ {mod.symbols.Register("b"), array},
+ {mod.symbols.Register("c"), ty.i32()},
+ });
+ const_cast<type::Struct*>(s)->AddUsage(builtin::AddressSpace::kStorage);
+
+ // ALL_FIELDS() calls the macro FIELD(ADDR, TYPE, ARRAY_COUNT, NAME)
+ // for each field of the structure s.
+#define ALL_FIELDS() \
+ FIELD(0x0000, int, 0, a) \
+ FIELD(0x0004, int8_t, 12, tint_pad) \
+ FIELD(0x0010, float3, 4, b) \
+ FIELD(0x0050, int, 0, c) \
+ FIELD(0x0054, int8_t, 12, tint_pad_1)
+
+ // Check that the generated string is as expected.
+ utils::StringStream expect;
+
+ expect << R"(template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+)";
+
+ expect << "struct S {\n";
+#define FIELD(ADDR, TYPE, ARRAY_COUNT, NAME) \
+ FormatMSLField(expect, #ADDR, #TYPE, ARRAY_COUNT, #NAME);
+ ALL_FIELDS()
+#undef FIELD
+ expect << "};\n\nS";
+
+ generator_.EmitType(generator_.Line(), s);
+ ASSERT_TRUE(generator_.Diagnostics().empty()) << generator_.Diagnostics().str();
+ EXPECT_EQ(utils::TrimSpace(generator_.Result()), expect.str());
+}
+
+TEST_F(MslGeneratorImplIrTest, AttemptTintPadSymbolCollision) {
+ utils::Vector<MemberData, 26> data = {
+ // uses symbols tint_pad_[0..9] and tint_pad_[20..35]
+ {mod.symbols.Register("tint_pad_2"), ty.i32(), 32}, //
+ {mod.symbols.Register("tint_pad_20"), ty.f32(), 128, 128}, //
+ {mod.symbols.Register("tint_pad_33"), ty.vec2<f32>()}, //
+ {mod.symbols.Register("tint_pad_1"), ty.u32()}, //
+ {mod.symbols.Register("tint_pad_3"), ty.vec3<f32>()}, //
+ {mod.symbols.Register("tint_pad_7"), ty.u32()}, //
+ {mod.symbols.Register("tint_pad_25"), ty.vec4<f32>()}, //
+ {mod.symbols.Register("tint_pad_5"), ty.u32()}, //
+ {mod.symbols.Register("tint_pad_27"), ty.mat2x2<f32>()}, //
+ {mod.symbols.Register("tint_pad_24"), ty.u32()}, //
+ {mod.symbols.Register("tint_pad_23"), ty.mat2x3<f32>()}, //
+ {mod.symbols.Register("tint_pad"), ty.u32()}, //
+ {mod.symbols.Register("tint_pad_8"), ty.mat2x4<f32>()}, //
+ {mod.symbols.Register("tint_pad_26"), ty.u32()}, //
+ {mod.symbols.Register("tint_pad_29"), ty.mat3x2<f32>()}, //
+ {mod.symbols.Register("tint_pad_6"), ty.u32()}, //
+ {mod.symbols.Register("tint_pad_22"), ty.mat3x3<f32>()}, //
+ {mod.symbols.Register("tint_pad_32"), ty.u32()}, //
+ {mod.symbols.Register("tint_pad_34"), ty.mat3x4<f32>()}, //
+ {mod.symbols.Register("tint_pad_35"), ty.u32()}, //
+ {mod.symbols.Register("tint_pad_30"), ty.mat4x2<f32>()}, //
+ {mod.symbols.Register("tint_pad_9"), ty.u32()}, //
+ {mod.symbols.Register("tint_pad_31"), ty.mat4x3<f32>()}, //
+ {mod.symbols.Register("tint_pad_28"), ty.u32()}, //
+ {mod.symbols.Register("tint_pad_4"), ty.mat4x4<f32>()}, //
+ {mod.symbols.Register("tint_pad_21"), ty.f32()}};
+
+ auto* s = MkStruct(mod, ty, "S", data);
+ s->AddUsage(builtin::AddressSpace::kStorage);
+
+ auto expect = R"(template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+struct S {
+ /* 0x0000 */ int tint_pad_2;
+ /* 0x0004 */ tint_array<int8_t, 124> tint_pad_10;
+ /* 0x0080 */ float tint_pad_20;
+ /* 0x0084 */ tint_array<int8_t, 124> tint_pad_11;
+ /* 0x0100 */ float2 tint_pad_33;
+ /* 0x0108 */ uint tint_pad_1;
+ /* 0x010c */ tint_array<int8_t, 4> tint_pad_12;
+ /* 0x0110 */ float3 tint_pad_3;
+ /* 0x011c */ uint tint_pad_7;
+ /* 0x0120 */ float4 tint_pad_25;
+ /* 0x0130 */ uint tint_pad_5;
+ /* 0x0134 */ tint_array<int8_t, 4> tint_pad_13;
+ /* 0x0138 */ float2x2 tint_pad_27;
+ /* 0x0148 */ uint tint_pad_24;
+ /* 0x014c */ tint_array<int8_t, 4> tint_pad_14;
+ /* 0x0150 */ float2x3 tint_pad_23;
+ /* 0x0170 */ uint tint_pad;
+ /* 0x0174 */ tint_array<int8_t, 12> tint_pad_15;
+ /* 0x0180 */ float2x4 tint_pad_8;
+ /* 0x01a0 */ uint tint_pad_26;
+ /* 0x01a4 */ tint_array<int8_t, 4> tint_pad_16;
+ /* 0x01a8 */ float3x2 tint_pad_29;
+ /* 0x01c0 */ uint tint_pad_6;
+ /* 0x01c4 */ tint_array<int8_t, 12> tint_pad_17;
+ /* 0x01d0 */ float3x3 tint_pad_22;
+ /* 0x0200 */ uint tint_pad_32;
+ /* 0x0204 */ tint_array<int8_t, 12> tint_pad_18;
+ /* 0x0210 */ float3x4 tint_pad_34;
+ /* 0x0240 */ uint tint_pad_35;
+ /* 0x0244 */ tint_array<int8_t, 4> tint_pad_19;
+ /* 0x0248 */ float4x2 tint_pad_30;
+ /* 0x0268 */ uint tint_pad_9;
+ /* 0x026c */ tint_array<int8_t, 4> tint_pad_36;
+ /* 0x0270 */ float4x3 tint_pad_31;
+ /* 0x02b0 */ uint tint_pad_28;
+ /* 0x02b4 */ tint_array<int8_t, 12> tint_pad_37;
+ /* 0x02c0 */ float4x4 tint_pad_4;
+ /* 0x0300 */ float tint_pad_21;
+ /* 0x0304 */ tint_array<int8_t, 124> tint_pad_38;
+};
+
+S)";
+
+ generator_.EmitType(generator_.Line(), s);
+ ASSERT_TRUE(generator_.Diagnostics().empty()) << generator_.Diagnostics().str();
+ EXPECT_STREQ(std::string(utils::TrimSpace(generator_.Result())).c_str(), expect);
+}
TEST_F(MslGeneratorImplIrTest, EmitType_Sampler) {
generator_.EmitType(generator_.Line(), ty.sampler());
diff --git a/src/tint/writer/text_generator.cc b/src/tint/writer/text_generator.cc
index 417f7bd..1585e79 100644
--- a/src/tint/writer/text_generator.cc
+++ b/src/tint/writer/text_generator.cc
@@ -18,6 +18,7 @@
#include <limits>
#include "src/tint/debug.h"
+#include "src/tint/utils/map.h"
namespace tint::writer {
@@ -25,6 +26,15 @@
TextGenerator::~TextGenerator() = default;
+std::string TextGenerator::StructName(const type::Struct* s) {
+ auto name = s->Name().Name();
+ if (name.size() > 1 && name[0] == '_' && name[1] == '_') {
+ name = utils::GetOrCreate(builtin_struct_names_, s,
+ [&] { return UniqueIdentifier(name.substr(2)); });
+ }
+ return name;
+}
+
TextGenerator::LineWriter::LineWriter(TextBuffer* buf) : buffer(buf) {}
TextGenerator::LineWriter::LineWriter(LineWriter&& other) {
diff --git a/src/tint/writer/text_generator.h b/src/tint/writer/text_generator.h
index bb9867e..348a31b 100644
--- a/src/tint/writer/text_generator.h
+++ b/src/tint/writer/text_generator.h
@@ -21,6 +21,7 @@
#include <vector>
#include "src/tint/diagnostic/diagnostic.h"
+#include "src/tint/type/struct.h"
#include "src/tint/utils/string_stream.h"
namespace tint::writer {
@@ -130,6 +131,17 @@
/// the end of `buffer`.
static LineWriter Line(TextBuffer* buffer) { return LineWriter(buffer); }
+ /// @return a new, unique identifier with the given prefix.
+ /// @param prefix optional prefix to apply to the generated identifier. If
+ /// empty "tint_symbol" will be used.
+ virtual std::string UniqueIdentifier(const std::string& prefix = "") = 0;
+
+ /// @param s the structure
+ /// @returns the name of the structure, taking special care of builtin structures that start
+ /// with double underscores. If the structure is a builtin, then the returned name will be a
+ /// unique name without the leading underscores.
+ std::string StructName(const type::Struct* s);
+
/// @returns the result data
virtual std::string Result() const { return main_buffer_.String(); }
@@ -183,6 +195,10 @@
/// The primary text buffer that the generator will emit
TextBuffer main_buffer_;
+
+ private:
+ /// Map of builtin structure to unique generated name
+ std::unordered_map<const type::Struct*, std::string> builtin_struct_names_;
};
} // namespace tint::writer
