Import Tint changes from Dawn
Changes:
- 6198bea2ac17ff1a085354205e11c16d2894a2bd Dawn&Tint: Implement F16 pipeline IO by Zhaoming Jiang <zhaoming.jiang@intel.com>
- 94751e1a7d2e9f7746bf1566ab5f339abe4d491f tint/utils: Vector::Clear: disable "maybe initialized" wa... by Jason Erb <jason.erb@sparist.com>
- ad9cd0a141b427e0af0d7ebe89e4744dbce1a515 Add base class for sem Struct and StructMember by dan sinclair <dsinclair@chromium.org>
- c07de73b0b56efcabc355c2a4df9211dcf1b0f30 tint/transform: Strip unused atomic builtins by Ben Clayton <bclayton@google.com>
- 49334b05cf77d272b1eb479ad8afb1e111633365 tint/utils: Remove non-const accessors on VectorRef by Ben Clayton <bclayton@google.com>
- 6016d1e5cdc913d405076b411584e1b56caef3f4 tint: Fix unused-but-set-parameter warning by Jason Erb <jason.erb@sparist.com>
- 8753796aaca097153af3ccebdd0f17d80d25bb9f tint: Add PreservePadding transform by James Price <jrprice@google.com>
- a31d89d6a330565858c1d61421015ceb7bdc3afb tint/utils: Add support for unsafe pointer downcasts by Ben Clayton <bclayton@google.com>
GitOrigin-RevId: 6198bea2ac17ff1a085354205e11c16d2894a2bd
Change-Id: Id93347fd22429517fa2c40b9eb98fb90468a039c
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/113220
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index 6a0e70d..cbcd804 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -529,6 +529,8 @@
"transform/packed_vec3.h",
"transform/pad_structs.cc",
"transform/pad_structs.h",
+ "transform/preserve_padding.cc",
+ "transform/preserve_padding.h",
"transform/promote_initializers_to_let.cc",
"transform/promote_initializers_to_let.h",
"transform/promote_side_effects_to_decl.cc",
@@ -1244,6 +1246,7 @@
"transform/num_workgroups_from_uniform_test.cc",
"transform/packed_vec3_test.cc",
"transform/pad_structs_test.cc",
+ "transform/preserve_padding_test.cc",
"transform/promote_initializers_to_let_test.cc",
"transform/promote_side_effects_to_decl_test.cc",
"transform/remove_continue_in_switch_test.cc",
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index 72cec7b..9156cb0 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -453,6 +453,8 @@
transform/packed_vec3.h
transform/pad_structs.cc
transform/pad_structs.h
+ transform/preserve_padding.cc
+ transform/preserve_padding.h
transform/promote_initializers_to_let.cc
transform/promote_initializers_to_let.h
transform/promote_side_effects_to_decl.cc
@@ -1212,6 +1214,7 @@
transform/num_workgroups_from_uniform_test.cc
transform/packed_vec3_test.cc
transform/pad_structs_test.cc
+ transform/preserve_padding_test.cc
transform/promote_initializers_to_let_test.cc
transform/promote_side_effects_to_decl_test.cc
transform/remove_continue_in_switch_test.cc
diff --git a/src/tint/castable.h b/src/tint/castable.h
index fe820e1..86f11c2 100644
--- a/src/tint/castable.h
+++ b/src/tint/castable.h
@@ -573,14 +573,12 @@
/// then the case handler is called.
/// @returns true if a case handler was found, otherwise false.
template <typename T, typename RETURN_TYPE, typename... CASES>
-inline bool NonDefaultCases(T* object,
+inline bool NonDefaultCases([[maybe_unused]] T* object,
const TypeInfo* type,
- RETURN_TYPE* result,
+ [[maybe_unused]] RETURN_TYPE* result,
std::tuple<CASES...>&& cases) {
using Cases = std::tuple<CASES...>;
- (void)result; // Not always used, avoid warning.
-
static constexpr bool kHasReturnType = !std::is_same_v<RETURN_TYPE, void>;
static constexpr size_t kNumCases = sizeof...(CASES);
diff --git a/src/tint/inspector/entry_point.h b/src/tint/inspector/entry_point.h
index 4a4706b..eabe601 100644
--- a/src/tint/inspector/entry_point.h
+++ b/src/tint/inspector/entry_point.h
@@ -30,9 +30,10 @@
/// Base component type of a stage variable.
enum class ComponentType {
kUnknown = -1,
- kFloat,
- kUInt,
- kSInt,
+ kF32,
+ kU32,
+ kI32,
+ kF16,
};
/// Composition of components of a stage variable.
diff --git a/src/tint/inspector/inspector.cc b/src/tint/inspector/inspector.cc
index 25cde92..b893ed1 100644
--- a/src/tint/inspector/inspector.cc
+++ b/src/tint/inspector/inspector.cc
@@ -69,41 +69,48 @@
}
std::tuple<ComponentType, CompositionType> CalculateComponentAndComposition(const sem::Type* type) {
- if (type->is_float_scalar()) {
- return {ComponentType::kFloat, CompositionType::kScalar};
- } else if (type->is_float_vector()) {
- auto* vec = type->As<sem::Vector>();
- if (vec->Width() == 2) {
- return {ComponentType::kFloat, CompositionType::kVec2};
- } else if (vec->Width() == 3) {
- return {ComponentType::kFloat, CompositionType::kVec3};
- } else if (vec->Width() == 4) {
- return {ComponentType::kFloat, CompositionType::kVec4};
+ // entry point in/out variables must of numeric scalar or vector types.
+ TINT_ASSERT(Inspector, type->is_numeric_scalar_or_vector());
+
+ ComponentType componentType = Switch(
+ sem::Type::DeepestElementOf(type), //
+ [&](const sem::F32*) { return ComponentType::kF32; },
+ [&](const sem::F16*) { return ComponentType::kF16; },
+ [&](const sem::I32*) { return ComponentType::kI32; },
+ [&](const sem::U32*) { return ComponentType::kU32; },
+ [&](Default) {
+ tint::diag::List diagnostics;
+ TINT_UNREACHABLE(Inspector, diagnostics) << "unhandled component type";
+ return ComponentType::kUnknown;
+ });
+
+ CompositionType compositionType;
+ if (auto* vec = type->As<sem::Vector>()) {
+ switch (vec->Width()) {
+ case 2: {
+ compositionType = CompositionType::kVec2;
+ break;
+ }
+ case 3: {
+ compositionType = CompositionType::kVec3;
+ break;
+ }
+ case 4: {
+ compositionType = CompositionType::kVec4;
+ break;
+ }
+ default: {
+ tint::diag::List diagnostics;
+ TINT_UNREACHABLE(Inspector, diagnostics) << "unhandled composition type";
+ compositionType = CompositionType::kUnknown;
+ break;
+ }
}
- } else if (type->is_unsigned_integer_scalar()) {
- return {ComponentType::kUInt, CompositionType::kScalar};
- } else if (type->is_unsigned_integer_vector()) {
- auto* vec = type->As<sem::Vector>();
- if (vec->Width() == 2) {
- return {ComponentType::kUInt, CompositionType::kVec2};
- } else if (vec->Width() == 3) {
- return {ComponentType::kUInt, CompositionType::kVec3};
- } else if (vec->Width() == 4) {
- return {ComponentType::kUInt, CompositionType::kVec4};
- }
- } else if (type->is_signed_integer_scalar()) {
- return {ComponentType::kSInt, CompositionType::kScalar};
- } else if (type->is_signed_integer_vector()) {
- auto* vec = type->As<sem::Vector>();
- if (vec->Width() == 2) {
- return {ComponentType::kSInt, CompositionType::kVec2};
- } else if (vec->Width() == 3) {
- return {ComponentType::kSInt, CompositionType::kVec3};
- } else if (vec->Width() == 4) {
- return {ComponentType::kSInt, CompositionType::kVec4};
- }
+ } else {
+ compositionType = CompositionType::kScalar;
}
- return {ComponentType::kUnknown, CompositionType::kUnknown};
+
+ return {componentType, compositionType};
}
std::tuple<InterpolationType, InterpolationSampling> CalculateInterpolationData(
@@ -532,7 +539,7 @@
ResourceBinding::ResourceType::kExternalTexture);
}
-utils::Vector<sem::SamplerTexturePair, 4> Inspector::GetSamplerTextureUses(
+utils::VectorRef<sem::SamplerTexturePair> Inspector::GetSamplerTextureUses(
const std::string& entry_point) {
auto* func = FindEntryPointByName(entry_point);
if (!func) {
diff --git a/src/tint/inspector/inspector.h b/src/tint/inspector/inspector.h
index 49e4bdf..62a1c73 100644
--- a/src/tint/inspector/inspector.h
+++ b/src/tint/inspector/inspector.h
@@ -126,7 +126,7 @@
/// @param entry_point name of the entry point to get information about.
/// @returns vector of all of the sampler/texture sampling pairs that are used
/// by that entry point.
- utils::Vector<sem::SamplerTexturePair, 4> GetSamplerTextureUses(const std::string& entry_point);
+ utils::VectorRef<sem::SamplerTexturePair> GetSamplerTextureUses(const std::string& entry_point);
/// @param entry_point name of the entry point to get information about.
/// @param placeholder the sampler binding point to use for texture-only
diff --git a/src/tint/inspector/inspector_test.cc b/src/tint/inspector/inspector_test.cc
index 05f3c19..78cfb5a 100644
--- a/src/tint/inspector/inspector_test.cc
+++ b/src/tint/inspector/inspector_test.cc
@@ -12,7 +12,8 @@
// See the License for the specific language governing permissions and
// limitations under the License.
-#include "gtest/gtest.h"
+#include "gmock/gmock.h"
+
#include "src/tint/ast/call_statement.h"
#include "src/tint/ast/disable_validation_attribute.h"
#include "src/tint/ast/id_attribute.h"
@@ -33,16 +34,13 @@
namespace tint::inspector {
namespace {
-// All the tests that descend from InspectorBuilder are expected to define their
-// test state via building up the AST through InspectorBuilder and then generate
-// the program with ::Build.
-// The returned Inspector from ::Build can then be used to test expecations.
+// All the tests that descend from InspectorBuilder are expected to define their test state via
+// building up the AST through InspectorBuilder and then generate the program with ::Build. The
+// returned Inspector from ::Build can then be used to test expectations.
//
-// All the tests that descend from InspectorRunner are expected to define their
-// test state via a WGSL shader, which will be parsed to generate a Program and
-// Inspector in ::Initialize.
-// The returned Inspector from ::Initialize can then be used to test
-// expecations.
+// All the tests that descend from InspectorRunner are expected to define their test state via a
+// WGSL shader, which will be parsed to generate a Program and Inspector in ::Initialize. The
+// returned Inspector from ::Initialize can then be used to test expectations.
class InspectorGetEntryPointTest : public InspectorBuilder, public testing::Test {};
@@ -289,6 +287,10 @@
std::tie(component, composition) = GetParam();
std::function<const ast::Type*()> tint_type = GetTypeFunction(component, composition);
+ if (component == ComponentType::kF16) {
+ Enable(ast::Extension::kF16);
+ }
+
auto* in_var = Param("in_var", tint_type(),
utils::Vector{
Location(0_u),
@@ -325,9 +327,10 @@
}
INSTANTIATE_TEST_SUITE_P(InspectorGetEntryPointTest,
InspectorGetEntryPointComponentAndCompositionTest,
- testing::Combine(testing::Values(ComponentType::kFloat,
- ComponentType::kSInt,
- ComponentType::kUInt),
+ testing::Combine(testing::Values(ComponentType::kF32,
+ ComponentType::kI32,
+ ComponentType::kU32,
+ ComponentType::kF16),
testing::Values(CompositionType::kScalar,
CompositionType::kVec2,
CompositionType::kVec3,
@@ -371,23 +374,23 @@
EXPECT_TRUE(result[0].input_variables[0].has_location_attribute);
EXPECT_EQ(0u, result[0].input_variables[0].location_attribute);
EXPECT_EQ(InterpolationType::kFlat, result[0].input_variables[0].interpolation_type);
- EXPECT_EQ(ComponentType::kUInt, result[0].input_variables[0].component_type);
+ EXPECT_EQ(ComponentType::kU32, result[0].input_variables[0].component_type);
EXPECT_EQ("in_var1", result[0].input_variables[1].name);
EXPECT_TRUE(result[0].input_variables[1].has_location_attribute);
EXPECT_EQ(1u, result[0].input_variables[1].location_attribute);
EXPECT_EQ(InterpolationType::kFlat, result[0].input_variables[1].interpolation_type);
- EXPECT_EQ(ComponentType::kUInt, result[0].input_variables[1].component_type);
+ EXPECT_EQ(ComponentType::kU32, result[0].input_variables[1].component_type);
EXPECT_EQ("in_var4", result[0].input_variables[2].name);
EXPECT_TRUE(result[0].input_variables[2].has_location_attribute);
EXPECT_EQ(4u, result[0].input_variables[2].location_attribute);
EXPECT_EQ(InterpolationType::kFlat, result[0].input_variables[2].interpolation_type);
- EXPECT_EQ(ComponentType::kUInt, result[0].input_variables[2].component_type);
+ EXPECT_EQ(ComponentType::kU32, result[0].input_variables[2].component_type);
ASSERT_EQ(1u, result[0].output_variables.size());
EXPECT_EQ("<retval>", result[0].output_variables[0].name);
EXPECT_TRUE(result[0].output_variables[0].has_location_attribute);
EXPECT_EQ(0u, result[0].output_variables[0].location_attribute);
- EXPECT_EQ(ComponentType::kUInt, result[0].output_variables[0].component_type);
+ EXPECT_EQ(ComponentType::kU32, result[0].output_variables[0].component_type);
}
TEST_F(InspectorGetEntryPointTest, MultipleEntryPointsInOutVariables) {
@@ -435,26 +438,26 @@
EXPECT_TRUE(result[0].input_variables[0].has_location_attribute);
EXPECT_EQ(0u, result[0].input_variables[0].location_attribute);
EXPECT_EQ(InterpolationType::kFlat, result[0].input_variables[0].interpolation_type);
- EXPECT_EQ(ComponentType::kUInt, result[0].input_variables[0].component_type);
+ EXPECT_EQ(ComponentType::kU32, result[0].input_variables[0].component_type);
ASSERT_EQ(1u, result[0].output_variables.size());
EXPECT_EQ("<retval>", result[0].output_variables[0].name);
EXPECT_TRUE(result[0].output_variables[0].has_location_attribute);
EXPECT_EQ(0u, result[0].output_variables[0].location_attribute);
- EXPECT_EQ(ComponentType::kUInt, result[0].output_variables[0].component_type);
+ EXPECT_EQ(ComponentType::kU32, result[0].output_variables[0].component_type);
ASSERT_EQ(1u, result[1].input_variables.size());
EXPECT_EQ("in_var_bar", result[1].input_variables[0].name);
EXPECT_TRUE(result[1].input_variables[0].has_location_attribute);
EXPECT_EQ(0u, result[1].input_variables[0].location_attribute);
EXPECT_EQ(InterpolationType::kFlat, result[1].input_variables[0].interpolation_type);
- EXPECT_EQ(ComponentType::kUInt, result[1].input_variables[0].component_type);
+ EXPECT_EQ(ComponentType::kU32, result[1].input_variables[0].component_type);
ASSERT_EQ(1u, result[1].output_variables.size());
EXPECT_EQ("<retval>", result[1].output_variables[0].name);
EXPECT_TRUE(result[1].output_variables[0].has_location_attribute);
EXPECT_EQ(1u, result[1].output_variables[0].location_attribute);
- EXPECT_EQ(ComponentType::kUInt, result[1].output_variables[0].component_type);
+ EXPECT_EQ(ComponentType::kU32, result[1].output_variables[0].component_type);
}
TEST_F(InspectorGetEntryPointTest, BuiltInsNotStageVariables) {
@@ -487,7 +490,7 @@
EXPECT_EQ("in_var1", result[0].input_variables[0].name);
EXPECT_TRUE(result[0].input_variables[0].has_location_attribute);
EXPECT_EQ(0u, result[0].input_variables[0].location_attribute);
- EXPECT_EQ(ComponentType::kFloat, result[0].input_variables[0].component_type);
+ EXPECT_EQ(ComponentType::kF32, result[0].input_variables[0].component_type);
ASSERT_EQ(0u, result[0].output_variables.size());
}
@@ -519,21 +522,21 @@
EXPECT_EQ("param.a", result[0].input_variables[0].name);
EXPECT_TRUE(result[0].input_variables[0].has_location_attribute);
EXPECT_EQ(0u, result[0].input_variables[0].location_attribute);
- EXPECT_EQ(ComponentType::kUInt, result[0].input_variables[0].component_type);
+ EXPECT_EQ(ComponentType::kU32, result[0].input_variables[0].component_type);
EXPECT_EQ("param.b", result[0].input_variables[1].name);
EXPECT_TRUE(result[0].input_variables[1].has_location_attribute);
EXPECT_EQ(1u, result[0].input_variables[1].location_attribute);
- EXPECT_EQ(ComponentType::kUInt, result[0].input_variables[1].component_type);
+ EXPECT_EQ(ComponentType::kU32, result[0].input_variables[1].component_type);
ASSERT_EQ(2u, result[0].output_variables.size());
EXPECT_EQ("<retval>.a", result[0].output_variables[0].name);
EXPECT_TRUE(result[0].output_variables[0].has_location_attribute);
EXPECT_EQ(0u, result[0].output_variables[0].location_attribute);
- EXPECT_EQ(ComponentType::kUInt, result[0].output_variables[0].component_type);
+ EXPECT_EQ(ComponentType::kU32, result[0].output_variables[0].component_type);
EXPECT_EQ("<retval>.b", result[0].output_variables[1].name);
EXPECT_TRUE(result[0].output_variables[1].has_location_attribute);
EXPECT_EQ(1u, result[0].output_variables[1].location_attribute);
- EXPECT_EQ(ComponentType::kUInt, result[0].output_variables[1].component_type);
+ EXPECT_EQ(ComponentType::kU32, result[0].output_variables[1].component_type);
}
TEST_F(InspectorGetEntryPointTest, MultipleEntryPointsInOutSharedStruct) {
@@ -565,21 +568,21 @@
EXPECT_EQ("<retval>.a", result[0].output_variables[0].name);
EXPECT_TRUE(result[0].output_variables[0].has_location_attribute);
EXPECT_EQ(0u, result[0].output_variables[0].location_attribute);
- EXPECT_EQ(ComponentType::kUInt, result[0].output_variables[0].component_type);
+ EXPECT_EQ(ComponentType::kU32, result[0].output_variables[0].component_type);
EXPECT_EQ("<retval>.b", result[0].output_variables[1].name);
EXPECT_TRUE(result[0].output_variables[1].has_location_attribute);
EXPECT_EQ(1u, result[0].output_variables[1].location_attribute);
- EXPECT_EQ(ComponentType::kUInt, result[0].output_variables[1].component_type);
+ EXPECT_EQ(ComponentType::kU32, result[0].output_variables[1].component_type);
ASSERT_EQ(2u, result[1].input_variables.size());
EXPECT_EQ("param.a", result[1].input_variables[0].name);
EXPECT_TRUE(result[1].input_variables[0].has_location_attribute);
EXPECT_EQ(0u, result[1].input_variables[0].location_attribute);
- EXPECT_EQ(ComponentType::kUInt, result[1].input_variables[0].component_type);
+ EXPECT_EQ(ComponentType::kU32, result[1].input_variables[0].component_type);
EXPECT_EQ("param.b", result[1].input_variables[1].name);
EXPECT_TRUE(result[1].input_variables[1].has_location_attribute);
EXPECT_EQ(1u, result[1].input_variables[1].location_attribute);
- EXPECT_EQ(ComponentType::kUInt, result[1].input_variables[1].component_type);
+ EXPECT_EQ(ComponentType::kU32, result[1].input_variables[1].component_type);
ASSERT_EQ(0u, result[1].output_variables.size());
}
@@ -617,33 +620,33 @@
EXPECT_EQ("param_a.a", result[0].input_variables[0].name);
EXPECT_TRUE(result[0].input_variables[0].has_location_attribute);
EXPECT_EQ(0u, result[0].input_variables[0].location_attribute);
- EXPECT_EQ(ComponentType::kUInt, result[0].input_variables[0].component_type);
+ EXPECT_EQ(ComponentType::kU32, result[0].input_variables[0].component_type);
EXPECT_EQ("param_a.b", result[0].input_variables[1].name);
EXPECT_TRUE(result[0].input_variables[1].has_location_attribute);
EXPECT_EQ(1u, result[0].input_variables[1].location_attribute);
- EXPECT_EQ(ComponentType::kUInt, result[0].input_variables[1].component_type);
+ EXPECT_EQ(ComponentType::kU32, result[0].input_variables[1].component_type);
EXPECT_EQ("param_b.a", result[0].input_variables[2].name);
EXPECT_TRUE(result[0].input_variables[2].has_location_attribute);
EXPECT_EQ(2u, result[0].input_variables[2].location_attribute);
- EXPECT_EQ(ComponentType::kUInt, result[0].input_variables[2].component_type);
+ EXPECT_EQ(ComponentType::kU32, result[0].input_variables[2].component_type);
EXPECT_EQ("param_c", result[0].input_variables[3].name);
EXPECT_TRUE(result[0].input_variables[3].has_location_attribute);
EXPECT_EQ(3u, result[0].input_variables[3].location_attribute);
- EXPECT_EQ(ComponentType::kFloat, result[0].input_variables[3].component_type);
+ EXPECT_EQ(ComponentType::kF32, result[0].input_variables[3].component_type);
EXPECT_EQ("param_d", result[0].input_variables[4].name);
EXPECT_TRUE(result[0].input_variables[4].has_location_attribute);
EXPECT_EQ(4u, result[0].input_variables[4].location_attribute);
- EXPECT_EQ(ComponentType::kFloat, result[0].input_variables[4].component_type);
+ EXPECT_EQ(ComponentType::kF32, result[0].input_variables[4].component_type);
ASSERT_EQ(2u, result[0].output_variables.size());
EXPECT_EQ("<retval>.a", result[0].output_variables[0].name);
EXPECT_TRUE(result[0].output_variables[0].has_location_attribute);
EXPECT_EQ(0u, result[0].output_variables[0].location_attribute);
- EXPECT_EQ(ComponentType::kUInt, result[0].output_variables[0].component_type);
+ EXPECT_EQ(ComponentType::kU32, result[0].output_variables[0].component_type);
EXPECT_EQ("<retval>.b", result[0].output_variables[1].name);
EXPECT_TRUE(result[0].output_variables[1].has_location_attribute);
EXPECT_EQ(1u, result[0].output_variables[1].location_attribute);
- EXPECT_EQ(ComponentType::kUInt, result[0].output_variables[1].component_type);
+ EXPECT_EQ(ComponentType::kU32, result[0].output_variables[1].component_type);
}
TEST_F(InspectorGetEntryPointTest, OverrideUnreferenced) {
@@ -3202,7 +3205,7 @@
})";
Inspector& inspector = Initialize(shader);
- auto result = inspector.GetSamplerTextureUses("foo");
+ inspector.GetSamplerTextureUses("foo");
ASSERT_TRUE(inspector.has_error()) << inspector.error();
}
@@ -3224,7 +3227,8 @@
auto result_1 = inspector.GetSamplerTextureUses("main");
ASSERT_FALSE(inspector.has_error()) << inspector.error();
- EXPECT_EQ(result_0, result_1);
+ EXPECT_EQ((utils::Vector<tint::sem::SamplerTexturePair, 4>(result_0)),
+ (utils::Vector<tint::sem::SamplerTexturePair, 4>(result_1)));
}
TEST_F(InspectorGetSamplerTextureUsesTest, BothIndirect) {
@@ -3641,7 +3645,7 @@
})";
Inspector& inspector = Initialize(shader);
- auto result = inspector.GetSamplerTextureUses("main");
+ inspector.GetSamplerTextureUses("main");
}
} // namespace
diff --git a/src/tint/inspector/test_inspector_builder.cc b/src/tint/inspector/test_inspector_builder.cc
index 79122dc..97ae097 100644
--- a/src/tint/inspector/test_inspector_builder.cc
+++ b/src/tint/inspector/test_inspector_builder.cc
@@ -307,15 +307,18 @@
CompositionType composition) {
std::function<const ast::Type*()> func;
switch (component) {
- case ComponentType::kFloat:
+ case ComponentType::kF32:
func = [this]() -> const ast::Type* { return ty.f32(); };
break;
- case ComponentType::kSInt:
+ case ComponentType::kI32:
func = [this]() -> const ast::Type* { return ty.i32(); };
break;
- case ComponentType::kUInt:
+ case ComponentType::kU32:
func = [this]() -> const ast::Type* { return ty.u32(); };
break;
+ case ComponentType::kF16:
+ func = [this]() -> const ast::Type* { return ty.f16(); };
+ break;
case ComponentType::kUnknown:
return []() -> const ast::Type* { return nullptr; };
}
diff --git a/src/tint/ir/builder.cc b/src/tint/ir/builder.cc
index 0f0200d..35e4f70 100644
--- a/src/tint/ir/builder.cc
+++ b/src/tint/ir/builder.cc
@@ -77,7 +77,7 @@
return ir_switch;
}
-Block* Builder::CreateCase(Switch* s, const utils::VectorRef<const ast::CaseSelector*> selectors) {
+Block* Builder::CreateCase(Switch* s, utils::VectorRef<const ast::CaseSelector*> selectors) {
s->cases.Push(Switch::Case{selectors, CreateBlock()});
Block* b = s->cases.Back().start_target;
diff --git a/src/tint/ir/builder.h b/src/tint/ir/builder.h
index 3f2e011..ceabfa6 100644
--- a/src/tint/ir/builder.h
+++ b/src/tint/ir/builder.h
@@ -78,7 +78,7 @@
/// @param s the switch to create the case into
/// @param selectors the case selectors for the case statement
/// @returns the start block for the case flow node
- Block* CreateCase(Switch* s, const utils::VectorRef<const ast::CaseSelector*> selectors);
+ Block* CreateCase(Switch* s, utils::VectorRef<const ast::CaseSelector*> selectors);
/// Branches the given block to the given flow node.
/// @param from the block to branch from
diff --git a/src/tint/ir/switch.h b/src/tint/ir/switch.h
index 3ea9c3f..7fff0a0 100644
--- a/src/tint/ir/switch.h
+++ b/src/tint/ir/switch.h
@@ -33,7 +33,7 @@
/// A case label in the struct
struct Case {
/// The case selector for this node
- const utils::VectorRef<const ast::CaseSelector*> selectors;
+ utils::Vector<const ast::CaseSelector*, 4> selectors;
/// The start block for the case block.
Block* start_target;
};
diff --git a/src/tint/resolver/atomics_test.cc b/src/tint/resolver/atomics_test.cc
index bb08e83..e3d0ca9 100644
--- a/src/tint/resolver/atomics_test.cc
+++ b/src/tint/resolver/atomics_test.cc
@@ -55,7 +55,7 @@
ASSERT_TRUE(TypeOf(g)->Is<sem::Reference>());
auto* str = TypeOf(g)->UnwrapRef()->As<sem::Struct>();
ASSERT_NE(str, nullptr);
- ASSERT_EQ(str->Members().size(), 1u);
+ ASSERT_EQ(str->Members().Length(), 1u);
auto* atomic = str->Members()[0]->Type()->As<sem::Atomic>();
ASSERT_NE(atomic, nullptr);
ASSERT_TRUE(atomic->Type()->Is<sem::I32>());
diff --git a/src/tint/resolver/builtin_test.cc b/src/tint/resolver/builtin_test.cc
index 15bf5de..4b33064 100644
--- a/src/tint/resolver/builtin_test.cc
+++ b/src/tint/resolver/builtin_test.cc
@@ -926,7 +926,7 @@
ASSERT_NE(TypeOf(call), nullptr);
auto* ty = TypeOf(call)->As<sem::Struct>();
ASSERT_NE(ty, nullptr);
- ASSERT_EQ(ty->Members().size(), 2u);
+ ASSERT_EQ(ty->Members().Length(), 2u);
auto* fract = ty->Members()[0];
EXPECT_TRUE(fract->Type()->Is<sem::F32>());
@@ -957,7 +957,7 @@
ASSERT_NE(TypeOf(call), nullptr);
auto* ty = TypeOf(call)->As<sem::Struct>();
ASSERT_NE(ty, nullptr);
- ASSERT_EQ(ty->Members().size(), 2u);
+ ASSERT_EQ(ty->Members().Length(), 2u);
auto* fract = ty->Members()[0];
EXPECT_TRUE(fract->Type()->Is<sem::F16>());
@@ -986,7 +986,7 @@
ASSERT_NE(TypeOf(call), nullptr);
auto* ty = TypeOf(call)->As<sem::Struct>();
ASSERT_NE(ty, nullptr);
- ASSERT_EQ(ty->Members().size(), 2u);
+ ASSERT_EQ(ty->Members().Length(), 2u);
auto* fract = ty->Members()[0];
ASSERT_TRUE(fract->Type()->Is<sem::Vector>());
@@ -1021,7 +1021,7 @@
ASSERT_NE(TypeOf(call), nullptr);
auto* ty = TypeOf(call)->As<sem::Struct>();
ASSERT_NE(ty, nullptr);
- ASSERT_EQ(ty->Members().size(), 2u);
+ ASSERT_EQ(ty->Members().Length(), 2u);
auto* fract = ty->Members()[0];
ASSERT_TRUE(fract->Type()->Is<sem::Vector>());
@@ -1058,7 +1058,7 @@
ASSERT_NE(TypeOf(call), nullptr);
auto* ty = TypeOf(call)->As<sem::Struct>();
ASSERT_NE(ty, nullptr);
- ASSERT_EQ(ty->Members().size(), 2u);
+ ASSERT_EQ(ty->Members().Length(), 2u);
auto* sig = ty->Members()[0];
EXPECT_TYPE(sig->Type(), TypeOf(expr));
@@ -1198,7 +1198,7 @@
ASSERT_NE(TypeOf(call), nullptr);
auto* ty = TypeOf(call)->As<sem::Struct>();
ASSERT_NE(ty, nullptr);
- ASSERT_EQ(ty->Members().size(), 2u);
+ ASSERT_EQ(ty->Members().Length(), 2u);
auto* fract = ty->Members()[0];
EXPECT_TRUE(fract->Type()->Is<sem::F32>());
@@ -1229,7 +1229,7 @@
ASSERT_NE(TypeOf(call), nullptr);
auto* ty = TypeOf(call)->As<sem::Struct>();
ASSERT_NE(ty, nullptr);
- ASSERT_EQ(ty->Members().size(), 2u);
+ ASSERT_EQ(ty->Members().Length(), 2u);
auto* fract = ty->Members()[0];
EXPECT_TRUE(fract->Type()->Is<sem::F16>());
@@ -1258,7 +1258,7 @@
ASSERT_NE(TypeOf(call), nullptr);
auto* ty = TypeOf(call)->As<sem::Struct>();
ASSERT_NE(ty, nullptr);
- ASSERT_EQ(ty->Members().size(), 2u);
+ ASSERT_EQ(ty->Members().Length(), 2u);
auto* fract = ty->Members()[0];
ASSERT_TRUE(fract->Type()->Is<sem::Vector>());
@@ -1293,7 +1293,7 @@
ASSERT_NE(TypeOf(call), nullptr);
auto* ty = TypeOf(call)->As<sem::Struct>();
ASSERT_NE(ty, nullptr);
- ASSERT_EQ(ty->Members().size(), 2u);
+ ASSERT_EQ(ty->Members().Length(), 2u);
auto* fract = ty->Members()[0];
ASSERT_TRUE(fract->Type()->Is<sem::Vector>());
diff --git a/src/tint/resolver/builtins_validation_test.cc b/src/tint/resolver/builtins_validation_test.cc
index c0452ce..d60b5b1 100644
--- a/src/tint/resolver/builtins_validation_test.cc
+++ b/src/tint/resolver/builtins_validation_test.cc
@@ -892,8 +892,8 @@
EXPECT_TRUE(r()->Resolve()) << r()->error();
auto* res_ty = TypeOf(builtin)->As<sem::Struct>();
ASSERT_TRUE(res_ty != nullptr);
- auto& members = res_ty->Members();
- ASSERT_EQ(members.size(), 2u);
+ auto members = res_ty->Members();
+ ASSERT_EQ(members.Length(), 2u);
EXPECT_TRUE(members[0]->Type()->Is<sem::F32>());
EXPECT_TRUE(members[1]->Type()->Is<sem::I32>());
}
@@ -905,8 +905,8 @@
EXPECT_TRUE(r()->Resolve()) << r()->error();
auto* res_ty = TypeOf(builtin)->As<sem::Struct>();
ASSERT_TRUE(res_ty != nullptr);
- auto& members = res_ty->Members();
- ASSERT_EQ(members.size(), 2u);
+ auto members = res_ty->Members();
+ ASSERT_EQ(members.Length(), 2u);
ASSERT_TRUE(members[0]->Type()->Is<sem::Vector>());
ASSERT_TRUE(members[1]->Type()->Is<sem::Vector>());
EXPECT_EQ(members[0]->Type()->As<sem::Vector>()->Width(), 2u);
@@ -922,8 +922,8 @@
EXPECT_TRUE(r()->Resolve()) << r()->error();
auto* res_ty = TypeOf(builtin)->As<sem::Struct>();
ASSERT_TRUE(res_ty != nullptr);
- auto& members = res_ty->Members();
- ASSERT_EQ(members.size(), 2u);
+ auto members = res_ty->Members();
+ ASSERT_EQ(members.Length(), 2u);
ASSERT_TRUE(members[0]->Type()->Is<sem::Vector>());
ASSERT_TRUE(members[1]->Type()->Is<sem::Vector>());
EXPECT_EQ(members[0]->Type()->As<sem::Vector>()->Width(), 3u);
@@ -939,8 +939,8 @@
EXPECT_TRUE(r()->Resolve()) << r()->error();
auto* res_ty = TypeOf(builtin)->As<sem::Struct>();
ASSERT_TRUE(res_ty != nullptr);
- auto& members = res_ty->Members();
- ASSERT_EQ(members.size(), 2u);
+ auto members = res_ty->Members();
+ ASSERT_EQ(members.Length(), 2u);
ASSERT_TRUE(members[0]->Type()->Is<sem::Vector>());
ASSERT_TRUE(members[1]->Type()->Is<sem::Vector>());
EXPECT_EQ(members[0]->Type()->As<sem::Vector>()->Width(), 4u);
@@ -956,8 +956,8 @@
EXPECT_TRUE(r()->Resolve()) << r()->error();
auto* res_ty = TypeOf(builtin)->As<sem::Struct>();
ASSERT_TRUE(res_ty != nullptr);
- auto& members = res_ty->Members();
- ASSERT_EQ(members.size(), 2u);
+ auto members = res_ty->Members();
+ ASSERT_EQ(members.Length(), 2u);
EXPECT_TRUE(members[0]->Type()->Is<sem::F32>());
EXPECT_TRUE(members[1]->Type()->Is<sem::F32>());
}
@@ -969,8 +969,8 @@
EXPECT_TRUE(r()->Resolve()) << r()->error();
auto* res_ty = TypeOf(builtin)->As<sem::Struct>();
ASSERT_TRUE(res_ty != nullptr);
- auto& members = res_ty->Members();
- ASSERT_EQ(members.size(), 2u);
+ auto members = res_ty->Members();
+ ASSERT_EQ(members.Length(), 2u);
ASSERT_TRUE(members[0]->Type()->Is<sem::Vector>());
ASSERT_TRUE(members[1]->Type()->Is<sem::Vector>());
EXPECT_EQ(members[0]->Type()->As<sem::Vector>()->Width(), 2u);
@@ -986,8 +986,8 @@
EXPECT_TRUE(r()->Resolve()) << r()->error();
auto* res_ty = TypeOf(builtin)->As<sem::Struct>();
ASSERT_TRUE(res_ty != nullptr);
- auto& members = res_ty->Members();
- ASSERT_EQ(members.size(), 2u);
+ auto members = res_ty->Members();
+ ASSERT_EQ(members.Length(), 2u);
ASSERT_TRUE(members[0]->Type()->Is<sem::Vector>());
ASSERT_TRUE(members[1]->Type()->Is<sem::Vector>());
EXPECT_EQ(members[0]->Type()->As<sem::Vector>()->Width(), 3u);
@@ -1003,8 +1003,8 @@
EXPECT_TRUE(r()->Resolve()) << r()->error();
auto* res_ty = TypeOf(builtin)->As<sem::Struct>();
ASSERT_TRUE(res_ty != nullptr);
- auto& members = res_ty->Members();
- ASSERT_EQ(members.size(), 2u);
+ auto members = res_ty->Members();
+ ASSERT_EQ(members.Length(), 2u);
ASSERT_TRUE(members[0]->Type()->Is<sem::Vector>());
ASSERT_TRUE(members[1]->Type()->Is<sem::Vector>());
EXPECT_EQ(members[0]->Type()->As<sem::Vector>()->Width(), 4u);
diff --git a/src/tint/resolver/const_eval.cc b/src/tint/resolver/const_eval.cc
index 45ab2de..ce38b1d 100644
--- a/src/tint/resolver/const_eval.cc
+++ b/src/tint/resolver/const_eval.cc
@@ -416,7 +416,7 @@
conv_els.Reserve(elements.Length());
std::function<const sem::Type*(size_t idx)> target_el_ty;
if (auto* str = target_ty->As<sem::Struct>()) {
- if (str->Members().size() != elements.Length()) {
+ if (str->Members().Length() != elements.Length()) {
TINT_ICE(Resolver, builder.Diagnostics())
<< "const-eval conversion of structure has mismatched element counts";
return utils::Failure;
@@ -496,7 +496,7 @@
[&](const sem::Struct* s) -> const ImplConstant* {
utils::Hashmap<const sem::Type*, const ImplConstant*, 8> zero_by_type;
utils::Vector<const sem::Constant*, 4> zeros;
- zeros.Reserve(s->Members().size());
+ zeros.Reserve(s->Members().Length());
for (auto* member : s->Members()) {
auto* zero = zero_by_type.GetOrCreate(
member->Type(), [&] { return ZeroValue(builder, member->Type()); });
@@ -507,7 +507,7 @@
}
if (zero_by_type.Count() == 1) {
// All members were of the same type, so the zero value is the same for all members.
- return builder.create<Splat>(type, zeros[0], s->Members().size());
+ return builder.create<Splat>(type, zeros[0], s->Members().Length());
}
return CreateComposite(builder, s, std::move(zeros));
},
diff --git a/src/tint/resolver/const_eval_construction_test.cc b/src/tint/resolver/const_eval_construction_test.cc
index cabc4ab..fe323c2 100644
--- a/src/tint/resolver/const_eval_construction_test.cc
+++ b/src/tint/resolver/const_eval_construction_test.cc
@@ -1625,7 +1625,7 @@
ASSERT_NE(sem, nullptr);
auto* str = sem->Type()->As<sem::Struct>();
ASSERT_NE(str, nullptr);
- EXPECT_EQ(str->Members().size(), 3u);
+ EXPECT_EQ(str->Members().Length(), 3u);
ASSERT_NE(sem->ConstantValue(), nullptr);
EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
EXPECT_TRUE(sem->ConstantValue()->AllEqual());
@@ -1670,7 +1670,7 @@
ASSERT_NE(sem, nullptr);
auto* str = sem->Type()->As<sem::Struct>();
ASSERT_NE(str, nullptr);
- EXPECT_EQ(str->Members().size(), 5u);
+ EXPECT_EQ(str->Members().Length(), 5u);
ASSERT_NE(sem->ConstantValue(), nullptr);
EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
EXPECT_FALSE(sem->ConstantValue()->AllEqual());
@@ -1723,7 +1723,7 @@
ASSERT_NE(sem, nullptr);
auto* str = sem->Type()->As<sem::Struct>();
ASSERT_NE(str, nullptr);
- EXPECT_EQ(str->Members().size(), 3u);
+ EXPECT_EQ(str->Members().Length(), 3u);
ASSERT_NE(sem->ConstantValue(), nullptr);
EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
EXPECT_TRUE(sem->ConstantValue()->AllEqual());
@@ -1777,7 +1777,7 @@
ASSERT_NE(sem, nullptr);
auto* str = sem->Type()->As<sem::Struct>();
ASSERT_NE(str, nullptr);
- EXPECT_EQ(str->Members().size(), 5u);
+ EXPECT_EQ(str->Members().Length(), 5u);
ASSERT_NE(sem->ConstantValue(), nullptr);
EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
EXPECT_FALSE(sem->ConstantValue()->AllEqual());
@@ -1849,7 +1849,7 @@
ASSERT_NE(sem, nullptr);
auto* str = sem->Type()->As<sem::Struct>();
ASSERT_NE(str, nullptr);
- EXPECT_EQ(str->Members().size(), 2u);
+ EXPECT_EQ(str->Members().Length(), 2u);
ASSERT_NE(sem->ConstantValue(), nullptr);
EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
EXPECT_TRUE(sem->ConstantValue()->AllEqual());
@@ -1892,7 +1892,7 @@
ASSERT_NE(sem, nullptr);
auto* str = sem->Type()->As<sem::Struct>();
ASSERT_NE(str, nullptr);
- EXPECT_EQ(str->Members().size(), 5u);
+ EXPECT_EQ(str->Members().Length(), 5u);
ASSERT_NE(sem->ConstantValue(), nullptr);
EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
EXPECT_FALSE(sem->ConstantValue()->AllEqual());
@@ -1950,7 +1950,7 @@
ASSERT_NE(sem, nullptr);
auto* str = sem->Type()->As<sem::Struct>();
ASSERT_NE(str, nullptr);
- EXPECT_EQ(str->Members().size(), 5u);
+ EXPECT_EQ(str->Members().Length(), 5u);
ASSERT_NE(sem->ConstantValue(), nullptr);
EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
EXPECT_FALSE(sem->ConstantValue()->AllEqual());
@@ -2024,7 +2024,7 @@
ASSERT_NE(sem, nullptr);
auto* str = sem->Type()->As<sem::Struct>();
ASSERT_NE(str, nullptr);
- EXPECT_EQ(str->Members().size(), 2u);
+ EXPECT_EQ(str->Members().Length(), 2u);
ASSERT_NE(sem->ConstantValue(), nullptr);
EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
EXPECT_FALSE(sem->ConstantValue()->AllEqual());
@@ -2064,7 +2064,7 @@
ASSERT_NE(sem, nullptr);
auto* str = sem->Type()->As<sem::Struct>();
ASSERT_NE(str, nullptr);
- EXPECT_EQ(str->Members().size(), 2u);
+ EXPECT_EQ(str->Members().Length(), 2u);
ASSERT_NE(sem->ConstantValue(), nullptr);
EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
EXPECT_FALSE(sem->ConstantValue()->AllEqual());
diff --git a/src/tint/resolver/const_eval_member_access_test.cc b/src/tint/resolver/const_eval_member_access_test.cc
index 25d1f26..2c5b15c 100644
--- a/src/tint/resolver/const_eval_member_access_test.cc
+++ b/src/tint/resolver/const_eval_member_access_test.cc
@@ -43,7 +43,7 @@
ASSERT_NE(outer, nullptr);
auto* str = outer->Type()->As<sem::Struct>();
ASSERT_NE(str, nullptr);
- EXPECT_EQ(str->Members().size(), 2u);
+ EXPECT_EQ(str->Members().Length(), 2u);
ASSERT_NE(outer->ConstantValue(), nullptr);
EXPECT_TYPE(outer->ConstantValue()->Type(), outer->Type());
EXPECT_FALSE(outer->ConstantValue()->AllEqual());
diff --git a/src/tint/resolver/entry_point_validation_test.cc b/src/tint/resolver/entry_point_validation_test.cc
index 79b41d7..cd49ded 100644
--- a/src/tint/resolver/entry_point_validation_test.cc
+++ b/src/tint/resolver/entry_point_validation_test.cc
@@ -606,17 +606,14 @@
ParamsFor<alias<i32>>(true), //
ParamsFor<alias<u32>>(true), //
ParamsFor<alias<bool>>(false), //
- // Currently entry point IO of f16 types are not implemented yet.
- // TODO(tint:1473, tint:1502): Change f16 and vecN<f16> cases to valid after f16 is supported in
- // entry point IO.
- ParamsFor<f16>(false), //
- ParamsFor<vec2<f16>>(false), //
- ParamsFor<vec3<f16>>(false), //
- ParamsFor<vec4<f16>>(false), //
+ ParamsFor<f16>(true), //
+ ParamsFor<vec2<f16>>(true), //
+ ParamsFor<vec3<f16>>(true), //
+ ParamsFor<vec4<f16>>(true), //
ParamsFor<mat2x2<f16>>(false), //
ParamsFor<mat3x3<f16>>(false), //
ParamsFor<mat4x4<f16>>(false), //
- ParamsFor<alias<f16>>(false), //
+ ParamsFor<alias<f16>>(true), //
};
TEST_P(TypeValidationTest, BareInputs) {
diff --git a/src/tint/resolver/inferred_type_test.cc b/src/tint/resolver/inferred_type_test.cc
index ddbc8f8..025ea95 100644
--- a/src/tint/resolver/inferred_type_test.cc
+++ b/src/tint/resolver/inferred_type_test.cc
@@ -150,12 +150,11 @@
auto* member = Member("x", ty.i32());
auto* str = Structure("S", utils::Vector{member});
- auto* expected_type =
- create<sem::Struct>(str, str->source, str->name,
- sem::StructMemberList{create<sem::StructMember>(
- member, member->source, member->symbol, create<sem::I32>(), 0u, 0u,
- 0u, 4u, std::nullopt)},
- 0u, 4u, 4u);
+ auto* expected_type = create<sem::Struct>(
+ str, str->source, str->name,
+ utils::Vector{create<sem::StructMember>(member, member->source, member->symbol,
+ create<sem::I32>(), 0u, 0u, 0u, 4u, std::nullopt)},
+ 0u, 4u, 4u);
auto* ctor_expr = Construct(ty.Of(str));
diff --git a/src/tint/resolver/intrinsic_table.cc b/src/tint/resolver/intrinsic_table.cc
index ebd5043..2578881 100644
--- a/src/tint/resolver/intrinsic_table.cc
+++ b/src/tint/resolver/intrinsic_table.cc
@@ -802,18 +802,18 @@
std::initializer_list<NameAndType> member_names_and_types) {
uint32_t offset = 0;
uint32_t max_align = 0;
- sem::StructMemberList members;
+ utils::Vector<const sem::StructMember*, 4> members;
for (auto& m : member_names_and_types) {
uint32_t align = std::max<uint32_t>(m.type->Align(), 1);
uint32_t size = m.type->Size();
offset = utils::RoundUp(align, offset);
max_align = std::max(max_align, align);
- members.emplace_back(b.create<sem::StructMember>(
+ members.Push(b.create<sem::StructMember>(
/* declaration */ nullptr,
/* source */ Source{},
/* name */ b.Sym(m.name),
/* type */ m.type,
- /* index */ static_cast<uint32_t>(members.size()),
+ /* index */ static_cast<uint32_t>(members.Length()),
/* offset */ offset,
/* align */ align,
/* size */ size,
@@ -826,7 +826,7 @@
/* declaration */ nullptr,
/* source */ Source{},
/* name */ b.Sym(name),
- /* members */ members,
+ /* members */ std::move(members),
/* align */ max_align,
/* size */ size_with_padding,
/* size_no_padding */ size_without_padding);
diff --git a/src/tint/resolver/resolver.cc b/src/tint/resolver/resolver.cc
index ca66ac2..e47f83c 100644
--- a/src/tint/resolver/resolver.cc
+++ b/src/tint/resolver/resolver.cc
@@ -1745,7 +1745,7 @@
return nullptr;
},
[&](const sem::Struct* s) -> const sem::Type* {
- if (auto& tys = s->ConcreteTypes(); !tys.IsEmpty()) {
+ if (auto tys = s->ConcreteTypes(); !tys.IsEmpty()) {
return target_ty ? target_ty : tys[0];
}
return nullptr;
@@ -2069,7 +2069,7 @@
StructInitializerSig{{str, args.Length(), args_stage}},
[&]() -> sem::TypeInitializer* {
utils::Vector<const sem::Parameter*, 8> params;
- params.Resize(std::min(args.Length(), str->Members().size()));
+ params.Resize(std::min(args.Length(), str->Members().Length()));
for (size_t i = 0, n = params.Length(); i < n; i++) {
params[i] = builder_->create<sem::Parameter>(
nullptr, // declaration
@@ -3097,19 +3097,15 @@
Mark(attr);
}
- sem::StructMemberList sem_members;
- sem_members.reserve(str->members.Length());
+ utils::Vector<const sem::StructMember*, 8> sem_members;
+ sem_members.Reserve(str->members.Length());
- // Calculate the effective size and alignment of each field, and the overall
- // size of the structure.
- // For size, use the size attribute if provided, otherwise use the default
- // size for the type.
- // For alignment, use the alignment attribute if provided, otherwise use the
- // default alignment for the member type.
- // Diagnostic errors are raised if a basic rule is violated.
- // Validation of storage-class rules requires analyzing the actual variable
- // usage of the structure, and so is performed as part of the variable
- // validation.
+ // Calculate the effective size and alignment of each field, and the overall size of the
+ // structure. For size, use the size attribute if provided, otherwise use the default size for
+ // the type. For alignment, use the alignment attribute if provided, otherwise use the default
+ // alignment for the member type. Diagnostic errors are raised if a basic rule is violated.
+ // Validation of storage-class rules requires analyzing the actual variable usage of the
+ // structure, and so is performed as part of the variable validation.
uint64_t struct_size = 0;
uint64_t struct_align = 1;
utils::Hashmap<Symbol, const ast::StructMember*, 8> member_map;
@@ -3274,11 +3270,11 @@
}
auto* sem_member = builder_->create<sem::StructMember>(
- member, member->source, member->symbol, type, static_cast<uint32_t>(sem_members.size()),
- static_cast<uint32_t>(offset), static_cast<uint32_t>(align),
- static_cast<uint32_t>(size), location);
+ member, member->source, member->symbol, type,
+ static_cast<uint32_t>(sem_members.Length()), static_cast<uint32_t>(offset),
+ static_cast<uint32_t>(align), static_cast<uint32_t>(size), location);
builder_->Sem().Add(member, sem_member);
- sem_members.emplace_back(sem_member);
+ sem_members.Push(sem_member);
struct_size = offset + size;
struct_align = std::max(struct_align, align);
@@ -3299,10 +3295,10 @@
}
auto* out = builder_->create<sem::Struct>(
- str, str->source, str->name, sem_members, static_cast<uint32_t>(struct_align),
+ str, str->source, str->name, std::move(sem_members), static_cast<uint32_t>(struct_align),
static_cast<uint32_t>(struct_size), static_cast<uint32_t>(size_no_padding));
- for (size_t i = 0; i < sem_members.size(); i++) {
+ for (size_t i = 0; i < sem_members.Length(); i++) {
auto* mem_type = sem_members[i]->Type();
if (mem_type->Is<sem::Atomic>()) {
atomic_composite_info_.Add(out, &sem_members[i]->Source());
diff --git a/src/tint/resolver/struct_layout_test.cc b/src/tint/resolver/struct_layout_test.cc
index acfe691..7c50a3e 100644
--- a/src/tint/resolver/struct_layout_test.cc
+++ b/src/tint/resolver/struct_layout_test.cc
@@ -39,7 +39,7 @@
EXPECT_EQ(sem->Size(), 12u);
EXPECT_EQ(sem->SizeNoPadding(), 12u);
EXPECT_EQ(sem->Align(), 4u);
- ASSERT_EQ(sem->Members().size(), 3u);
+ ASSERT_EQ(sem->Members().Length(), 3u);
EXPECT_EQ(sem->Members()[0]->Offset(), 0u);
EXPECT_EQ(sem->Members()[0]->Align(), 4u);
EXPECT_EQ(sem->Members()[0]->Size(), 4u);
@@ -74,7 +74,7 @@
EXPECT_EQ(sem->Size(), 24u);
EXPECT_EQ(sem->SizeNoPadding(), 22u);
EXPECT_EQ(sem->Align(), 4u);
- ASSERT_EQ(sem->Members().size(), 7u);
+ ASSERT_EQ(sem->Members().Length(), 7u);
// f32
EXPECT_EQ(sem->Members()[0]->Offset(), 0u);
EXPECT_EQ(sem->Members()[0]->Align(), 4u);
@@ -121,7 +121,7 @@
EXPECT_EQ(sem->Size(), 8u);
EXPECT_EQ(sem->SizeNoPadding(), 8u);
EXPECT_EQ(sem->Align(), 4u);
- ASSERT_EQ(sem->Members().size(), 2u);
+ ASSERT_EQ(sem->Members().Length(), 2u);
EXPECT_EQ(sem->Members()[0]->Offset(), 0u);
EXPECT_EQ(sem->Members()[0]->Align(), 4u);
EXPECT_EQ(sem->Members()[0]->Size(), 4u);
@@ -150,7 +150,7 @@
EXPECT_EQ(sem->Size(), 52u);
EXPECT_EQ(sem->SizeNoPadding(), 52u);
EXPECT_EQ(sem->Align(), 4u);
- ASSERT_EQ(sem->Members().size(), 4u);
+ ASSERT_EQ(sem->Members().Length(), 4u);
// array<i32, 3>
EXPECT_EQ(sem->Members()[0]->Offset(), 0u);
EXPECT_EQ(sem->Members()[0]->Align(), 4u);
@@ -190,7 +190,7 @@
EXPECT_EQ(sem->Size(), 164u);
EXPECT_EQ(sem->SizeNoPadding(), 164u);
EXPECT_EQ(sem->Align(), 4u);
- ASSERT_EQ(sem->Members().size(), 4u);
+ ASSERT_EQ(sem->Members().Length(), 4u);
// array<i32, 3>, stride = 8
EXPECT_EQ(sem->Members()[0]->Offset(), 0u);
EXPECT_EQ(sem->Members()[0]->Align(), 4u);
@@ -225,7 +225,7 @@
EXPECT_EQ(sem->Size(), 4u);
EXPECT_EQ(sem->SizeNoPadding(), 4u);
EXPECT_EQ(sem->Align(), 4u);
- ASSERT_EQ(sem->Members().size(), 1u);
+ ASSERT_EQ(sem->Members().Length(), 1u);
EXPECT_EQ(sem->Members()[0]->Offset(), 0u);
EXPECT_EQ(sem->Members()[0]->Align(), 4u);
EXPECT_EQ(sem->Members()[0]->Size(), 4u);
@@ -246,7 +246,7 @@
EXPECT_EQ(sem->Size(), 32u);
EXPECT_EQ(sem->SizeNoPadding(), 32u);
EXPECT_EQ(sem->Align(), 4u);
- ASSERT_EQ(sem->Members().size(), 1u);
+ ASSERT_EQ(sem->Members().Length(), 1u);
EXPECT_EQ(sem->Members()[0]->Offset(), 0u);
EXPECT_EQ(sem->Members()[0]->Align(), 4u);
EXPECT_EQ(sem->Members()[0]->Size(), 32u);
@@ -269,7 +269,7 @@
EXPECT_EQ(sem->Size(), 384u);
EXPECT_EQ(sem->SizeNoPadding(), 384u);
EXPECT_EQ(sem->Align(), 4u);
- ASSERT_EQ(sem->Members().size(), 1u);
+ ASSERT_EQ(sem->Members().Length(), 1u);
EXPECT_EQ(sem->Members()[0]->Offset(), 0u);
EXPECT_EQ(sem->Members()[0]->Align(), 4u);
EXPECT_EQ(sem->Members()[0]->Size(), 384u);
@@ -296,7 +296,7 @@
EXPECT_EQ(sem->Size(), 576u);
EXPECT_EQ(sem->SizeNoPadding(), 576u);
EXPECT_EQ(sem->Align(), 16u);
- ASSERT_EQ(sem->Members().size(), 1u);
+ ASSERT_EQ(sem->Members().Length(), 1u);
EXPECT_EQ(sem->Members()[0]->Offset(), 0u);
EXPECT_EQ(sem->Members()[0]->Align(), 16u);
EXPECT_EQ(sem->Members()[0]->Size(), 576u);
@@ -319,7 +319,7 @@
EXPECT_EQ(sem->Size(), 48u);
EXPECT_EQ(sem->SizeNoPadding(), 48u);
EXPECT_EQ(sem->Align(), 16u);
- ASSERT_EQ(sem->Members().size(), 3u);
+ ASSERT_EQ(sem->Members().Length(), 3u);
EXPECT_EQ(sem->Members()[0]->Offset(), 0u); // vec2
EXPECT_EQ(sem->Members()[0]->Align(), 8u);
EXPECT_EQ(sem->Members()[0]->Size(), 8u);
@@ -365,7 +365,7 @@
EXPECT_EQ(sem->Size(), 576u);
EXPECT_EQ(sem->SizeNoPadding(), 576u);
EXPECT_EQ(sem->Align(), 16u);
- ASSERT_EQ(sem->Members().size(), 18u);
+ ASSERT_EQ(sem->Members().Length(), 18u);
EXPECT_EQ(sem->Members()[0]->Offset(), 0u); // mat2x2<f32>
EXPECT_EQ(sem->Members()[0]->Align(), 8u);
EXPECT_EQ(sem->Members()[0]->Size(), 16u);
@@ -443,7 +443,7 @@
EXPECT_EQ(sem->Size(), 80u);
EXPECT_EQ(sem->SizeNoPadding(), 68u);
EXPECT_EQ(sem->Align(), 16u);
- ASSERT_EQ(sem->Members().size(), 3u);
+ ASSERT_EQ(sem->Members().Length(), 3u);
EXPECT_EQ(sem->Members()[0]->Offset(), 0u);
EXPECT_EQ(sem->Members()[0]->Align(), 4u);
EXPECT_EQ(sem->Members()[0]->Size(), 4u);
@@ -478,7 +478,7 @@
EXPECT_EQ(sem->Size(), 76u);
EXPECT_EQ(sem->SizeNoPadding(), 76u);
EXPECT_EQ(sem->Align(), 4u);
- ASSERT_EQ(sem->Members().size(), 4u);
+ ASSERT_EQ(sem->Members().Length(), 4u);
EXPECT_EQ(sem->Members()[0]->Offset(), 0u);
EXPECT_EQ(sem->Members()[0]->Align(), 4u);
EXPECT_EQ(sem->Members()[0]->Size(), 4u);
@@ -516,7 +516,7 @@
EXPECT_EQ(sem->Size(), 96u);
EXPECT_EQ(sem->SizeNoPadding(), 68u);
EXPECT_EQ(sem->Align(), 32u);
- ASSERT_EQ(sem->Members().size(), 4u);
+ ASSERT_EQ(sem->Members().Length(), 4u);
EXPECT_EQ(sem->Members()[0]->Offset(), 0u);
EXPECT_EQ(sem->Members()[0]->Align(), 4u);
EXPECT_EQ(sem->Members()[0]->Size(), 4u);
@@ -546,7 +546,7 @@
EXPECT_EQ(sem->Size(), 1024u);
EXPECT_EQ(sem->SizeNoPadding(), 4u);
EXPECT_EQ(sem->Align(), 1024u);
- ASSERT_EQ(sem->Members().size(), 1u);
+ ASSERT_EQ(sem->Members().Length(), 1u);
EXPECT_EQ(sem->Members()[0]->Offset(), 0u);
EXPECT_EQ(sem->Members()[0]->Align(), 1024u);
EXPECT_EQ(sem->Members()[0]->Size(), 4u);
@@ -576,7 +576,7 @@
EXPECT_EQ(sem->Size(), 132u);
EXPECT_EQ(sem->SizeNoPadding(), 132u);
EXPECT_EQ(sem->Align(), 4u);
- ASSERT_EQ(sem->Members().size(), 5u);
+ ASSERT_EQ(sem->Members().Length(), 5u);
EXPECT_EQ(sem->Members()[0]->Offset(), 4u);
EXPECT_EQ(sem->Members()[0]->Align(), 4u);
EXPECT_EQ(sem->Members()[0]->Size(), 4u);
diff --git a/src/tint/resolver/struct_pipeline_stage_use_test.cc b/src/tint/resolver/struct_pipeline_stage_use_test.cc
index 107b241..c4455b7 100644
--- a/src/tint/resolver/struct_pipeline_stage_use_test.cc
+++ b/src/tint/resolver/struct_pipeline_stage_use_test.cc
@@ -184,7 +184,7 @@
auto* sem = TypeOf(s)->As<sem::Struct>();
ASSERT_NE(sem, nullptr);
- ASSERT_EQ(1u, sem->Members().size());
+ ASSERT_EQ(1u, sem->Members().Length());
EXPECT_EQ(3u, sem->Members()[0]->Location());
}
@@ -214,7 +214,7 @@
auto* sem = TypeOf(s)->As<sem::Struct>();
ASSERT_NE(sem, nullptr);
- ASSERT_EQ(1u, sem->Members().size());
+ ASSERT_EQ(1u, sem->Members().Length());
EXPECT_EQ(3u, sem->Members()[0]->Location());
}
diff --git a/src/tint/resolver/validator.cc b/src/tint/resolver/validator.cc
index e5c7c1a..558217d 100644
--- a/src/tint/resolver/validator.cc
+++ b/src/tint/resolver/validator.cc
@@ -440,7 +440,7 @@
}
if (auto* str = store_ty->As<sem::Struct>()) {
- for (size_t i = 0; i < str->Members().size(); ++i) {
+ for (size_t i = 0; i < str->Members().Length(); ++i) {
auto* const m = str->Members()[i];
uint32_t required_align = required_alignment_of(m->Type());
@@ -1068,13 +1068,6 @@
ParamOrRetType param_or_ret,
bool is_struct_member,
std::optional<uint32_t> location) {
- // Temporally forbid using f16 types in entry point IO.
- // TODO(tint:1473, tint:1502): Remove this error after f16 is supported in entry point IO.
- if (Is<sem::F16>(sem::Type::DeepestElementOf(ty))) {
- AddError("entry point IO of f16 types is not implemented yet", source);
- return false;
- }
-
// Scan attributes for pipeline IO attributes.
// Check for overlap with attributes that have been seen previously.
const ast::Attribute* pipeline_io_attribute = nullptr;
@@ -1724,10 +1717,10 @@
}
if (ctor->args.Length() > 0) {
- if (ctor->args.Length() != struct_type->Members().size()) {
- std::string fm = ctor->args.Length() < struct_type->Members().size() ? "few" : "many";
+ if (ctor->args.Length() != struct_type->Members().Length()) {
+ std::string fm = ctor->args.Length() < struct_type->Members().Length() ? "few" : "many";
AddError("struct initializer has too " + fm + " inputs: expected " +
- std::to_string(struct_type->Members().size()) + ", found " +
+ std::to_string(struct_type->Members().Length()) + ", found " +
std::to_string(ctor->args.Length()),
ctor->source);
return false;
@@ -1811,7 +1804,7 @@
return true;
}
-bool Validator::PipelineStages(const utils::VectorRef<sem::Function*> entry_points) const {
+bool Validator::PipelineStages(utils::VectorRef<sem::Function*> entry_points) const {
auto backtrace = [&](const sem::Function* func, const sem::Function* entry_point) {
if (func != entry_point) {
TraverseCallChain(diagnostics_, entry_point, func, [&](const sem::Function* f) {
@@ -1906,7 +1899,7 @@
return true;
}
-bool Validator::PushConstants(const utils::VectorRef<sem::Function*> entry_points) const {
+bool Validator::PushConstants(utils::VectorRef<sem::Function*> entry_points) const {
for (auto* entry_point : entry_points) {
// State checked and modified by check_push_constant so that it remembers previously seen
// push_constant variables for an entry-point.
@@ -2019,7 +2012,7 @@
}
bool Validator::Structure(const sem::Struct* str, ast::PipelineStage stage) const {
- if (str->Members().empty()) {
+ if (str->Members().IsEmpty()) {
AddError("structures must have at least one member", str->Source());
return false;
}
@@ -2028,7 +2021,7 @@
for (auto* member : str->Members()) {
if (auto* r = member->Type()->As<sem::Array>()) {
if (r->Count()->Is<sem::RuntimeArrayCount>()) {
- if (member != str->Members().back()) {
+ if (member != str->Members().Back()) {
AddError("runtime arrays may only appear as the last member of a struct",
member->Source());
return false;
@@ -2422,7 +2415,7 @@
const sem::Type* store_ty,
ast::Access access,
ast::AddressSpace address_space,
- const utils::VectorRef<const tint::ast::Attribute*> attributes,
+ utils::VectorRef<const tint::ast::Attribute*> attributes,
const Source& source) const {
if (!AddressSpaceLayout(store_ty, address_space, source)) {
return false;
diff --git a/src/tint/resolver/validator.h b/src/tint/resolver/validator.h
index 0f183bb..9f0664c 100644
--- a/src/tint/resolver/validator.h
+++ b/src/tint/resolver/validator.h
@@ -135,12 +135,12 @@
/// Validates pipeline stages
/// @param entry_points the entry points to the module
/// @returns true on success, false otherwise.
- bool PipelineStages(const utils::VectorRef<sem::Function*> entry_points) const;
+ bool PipelineStages(utils::VectorRef<sem::Function*> entry_points) const;
/// Validates push_constant variables
/// @param entry_points the entry points to the module
/// @returns true on success, false otherwise.
- bool PushConstants(const utils::VectorRef<sem::Function*> entry_points) const;
+ bool PushConstants(utils::VectorRef<sem::Function*> entry_points) const;
/// Validates aliases
/// @param alias the alias to validate
@@ -508,7 +508,7 @@
bool CheckTypeAccessAddressSpace(const sem::Type* store_ty,
ast::Access access,
ast::AddressSpace address_space,
- const utils::VectorRef<const tint::ast::Attribute*> attributes,
+ utils::VectorRef<const tint::ast::Attribute*> attributes,
const Source& source) const;
SymbolTable& symbols_;
diag::List& diagnostics_;
diff --git a/src/tint/sem/function.h b/src/tint/sem/function.h
index 8f21f74..7b5eaff 100644
--- a/src/tint/sem/function.h
+++ b/src/tint/sem/function.h
@@ -135,9 +135,7 @@
/// @returns the list of texture/sampler pairs that this function uses
/// (directly or indirectly).
- const utils::Vector<VariablePair, 8>& TextureSamplerPairs() const {
- return texture_sampler_pairs_;
- }
+ utils::VectorRef<VariablePair> TextureSamplerPairs() const { return texture_sampler_pairs_; }
/// @returns the list of direct calls to functions / builtins made by this
/// function
diff --git a/src/tint/sem/module.h b/src/tint/sem/module.h
index b451c5b..216a2c4 100644
--- a/src/tint/sem/module.h
+++ b/src/tint/sem/module.h
@@ -39,7 +39,7 @@
~Module() override;
/// @returns the dependency-ordered global declarations for the module
- const utils::Vector<const ast::Node*, 64>& DependencyOrderedDeclarations() const {
+ utils::VectorRef<const ast::Node*> DependencyOrderedDeclarations() const {
return dep_ordered_decls_;
}
diff --git a/src/tint/sem/struct.cc b/src/tint/sem/struct.cc
index c46e310..ecb06a3 100644
--- a/src/tint/sem/struct.cc
+++ b/src/tint/sem/struct.cc
@@ -23,13 +23,15 @@
#include "src/tint/symbol_table.h"
#include "src/tint/utils/hash.h"
+TINT_INSTANTIATE_TYPEINFO(tint::sem::StructBase);
TINT_INSTANTIATE_TYPEINFO(tint::sem::Struct);
+TINT_INSTANTIATE_TYPEINFO(tint::sem::StructMemberBase);
TINT_INSTANTIATE_TYPEINFO(tint::sem::StructMember);
namespace tint::sem {
namespace {
-TypeFlags FlagsFrom(const StructMemberList& members) {
+TypeFlags FlagsFrom(utils::VectorRef<const StructMemberBase*> members) {
TypeFlags flags{
TypeFlag::kConstructable,
TypeFlag::kCreationFixedFootprint,
@@ -54,12 +56,21 @@
Struct::Struct(const ast::Struct* declaration,
tint::Source source,
Symbol name,
- StructMemberList members,
+ utils::VectorRef<const StructMember*> members,
uint32_t align,
uint32_t size,
uint32_t size_no_padding)
+ : Base(source, name, members, align, size, size_no_padding), declaration_(declaration) {}
+
+Struct::~Struct() = default;
+
+StructBase::StructBase(tint::Source source,
+ Symbol name,
+ utils::VectorRef<const StructMemberBase*> members,
+ uint32_t align,
+ uint32_t size,
+ uint32_t size_no_padding)
: Base(FlagsFrom(members)),
- declaration_(declaration),
source_(source),
name_(name),
members_(std::move(members)),
@@ -67,20 +78,20 @@
size_(size),
size_no_padding_(size_no_padding) {}
-Struct::~Struct() = default;
+StructBase::~StructBase() = default;
-size_t Struct::Hash() const {
+size_t StructBase::Hash() const {
return utils::Hash(TypeInfo::Of<Struct>().full_hashcode, name_);
}
-bool Struct::Equals(const sem::Type& other) const {
+bool StructBase::Equals(const sem::Type& other) const {
if (auto* o = other.As<Struct>()) {
return o->name_ == name_;
}
return false;
}
-const StructMember* Struct::FindMember(Symbol name) const {
+const StructMemberBase* StructBase::FindMember(Symbol name) const {
for (auto* member : members_) {
if (member->Name() == name) {
return member;
@@ -89,27 +100,29 @@
return nullptr;
}
-uint32_t Struct::Align() const {
+uint32_t StructBase::Align() const {
return align_;
}
-uint32_t Struct::Size() const {
+uint32_t StructBase::Size() const {
return size_;
}
-std::string Struct::FriendlyName(const SymbolTable& symbols) const {
+std::string StructBase::FriendlyName(const SymbolTable& symbols) const {
return symbols.NameFor(name_);
}
-std::string Struct::Layout(const tint::SymbolTable& symbols) const {
+std::string StructBase::Layout(const tint::SymbolTable& symbols) const {
std::stringstream ss;
- auto member_name_of = [&](const sem::StructMember* sm) { return symbols.NameFor(sm->Name()); };
+ auto member_name_of = [&](const sem::StructMemberBase* sm) {
+ return symbols.NameFor(sm->Name());
+ };
- if (Members().empty()) {
+ if (Members().IsEmpty()) {
return {};
}
- const auto* const last_member = Members().back();
+ const auto* const last_member = Members().Back();
const uint32_t last_member_struct_padding_offset = last_member->Offset() + last_member->Size();
// Compute max widths to align output
@@ -135,7 +148,7 @@
print_struct_begin_line(Align(), Size(), UnwrapRef()->FriendlyName(symbols));
- for (size_t i = 0; i < Members().size(); ++i) {
+ for (size_t i = 0; i < Members().Length(); ++i) {
auto* const m = Members()[i];
// Output field alignment padding, if any
@@ -176,8 +189,19 @@
uint32_t align,
uint32_t size,
std::optional<uint32_t> location)
- : declaration_(declaration),
- source_(source),
+ : Base(source, name, type, index, offset, align, size, location), declaration_(declaration) {}
+
+StructMember::~StructMember() = default;
+
+StructMemberBase::StructMemberBase(tint::Source source,
+ Symbol name,
+ const sem::Type* type,
+ uint32_t index,
+ uint32_t offset,
+ uint32_t align,
+ uint32_t size,
+ std::optional<uint32_t> location)
+ : source_(source),
name_(name),
type_(type),
index_(index),
@@ -186,6 +210,6 @@
size_(size),
location_(location) {}
-StructMember::~StructMember() = default;
+StructMemberBase::~StructMemberBase() = default;
} // namespace tint::sem
diff --git a/src/tint/sem/struct.h b/src/tint/sem/struct.h
index 831cf3e..5ae2a93 100644
--- a/src/tint/sem/struct.h
+++ b/src/tint/sem/struct.h
@@ -20,13 +20,13 @@
#include <optional>
#include <string>
#include <unordered_set>
-#include <vector>
#include "src/tint/ast/address_space.h"
#include "src/tint/ast/struct.h"
#include "src/tint/sem/node.h"
#include "src/tint/sem/type.h"
#include "src/tint/symbol.h"
+#include "src/tint/utils/vector.h"
// Forward declarations
namespace tint::ast {
@@ -34,14 +34,12 @@
} // namespace tint::ast
namespace tint::sem {
class StructMember;
+class StructMemberBase;
class Type;
} // namespace tint::sem
namespace tint::sem {
-/// A vector of StructMember pointers.
-using StructMemberList = std::vector<const StructMember*>;
-
/// Metadata to capture how a structure is used in a shader module.
enum class PipelineStageUsage {
kVertexInput,
@@ -52,11 +50,10 @@
kComputeOutput,
};
-/// Struct holds the semantic information for structures.
-class Struct final : public Castable<Struct, Type> {
+/// StructBase holds the semantic information for structures.
+class StructBase : public Castable<StructBase, Type> {
public:
/// Constructor
- /// @param declaration the AST structure declaration
/// @param source the source of the structure
/// @param name the name of the structure
/// @param members the structure members
@@ -64,16 +61,15 @@
/// @param size the byte size of the structure
/// @param size_no_padding size of the members without the end of structure
/// alignment padding
- Struct(const ast::Struct* declaration,
- tint::Source source,
- Symbol name,
- StructMemberList members,
- uint32_t align,
- uint32_t size,
- uint32_t size_no_padding);
+ StructBase(tint::Source source,
+ Symbol name,
+ utils::VectorRef<const StructMemberBase*> members,
+ uint32_t align,
+ uint32_t size,
+ uint32_t size_no_padding);
/// Destructor
- ~Struct() override;
+ ~StructBase() override;
/// @returns a hash of the type.
size_t Hash() const override;
@@ -82,9 +78,6 @@
/// @returns true if the this type is equal to the given type
bool Equals(const Type& other) const override;
- /// @returns the struct
- const ast::Struct* Declaration() const { return declaration_; }
-
/// @returns the source of the structure
tint::Source Source() const { return source_; }
@@ -92,11 +85,11 @@
Symbol Name() const { return name_; }
/// @returns the members of the structure
- const StructMemberList& Members() const { return members_; }
+ utils::VectorRef<const StructMemberBase*> Members() const { return members_; }
/// @param name the member name to look for
/// @returns the member with the given name, or nullptr if it was not found.
- const StructMember* FindMember(Symbol name) const;
+ const StructMemberBase* FindMember(Symbol name) const;
/// @returns the byte alignment of the structure
/// @note this may differ from the alignment of a structure member of this
@@ -158,28 +151,131 @@
std::string Layout(const tint::SymbolTable& symbols) const;
/// @param concrete the conversion-rank ordered concrete versions of this abstract structure.
- void SetConcreteTypes(utils::VectorRef<const Struct*> concrete) { concrete_types_ = concrete; }
+ void SetConcreteTypes(utils::VectorRef<const StructBase*> concrete) {
+ concrete_types_ = concrete;
+ }
/// @returns the conversion-rank ordered concrete versions of this abstract structure, or an
/// empty vector if this structure is not abstract.
/// @note only structures returned by builtins may be abstract (e.g. modf, frexp)
- const utils::Vector<const Struct*, 2>& ConcreteTypes() const { return concrete_types_; }
+ utils::VectorRef<const StructBase*> ConcreteTypes() const { return concrete_types_; }
private:
- ast::Struct const* const declaration_;
const tint::Source source_;
const Symbol name_;
- const StructMemberList members_;
+ const utils::Vector<const StructMemberBase*, 4> members_;
const uint32_t align_;
const uint32_t size_;
const uint32_t size_no_padding_;
std::unordered_set<ast::AddressSpace> address_space_usage_;
std::unordered_set<PipelineStageUsage> pipeline_stage_uses_;
- utils::Vector<const Struct*, 2> concrete_types_;
+ utils::Vector<const StructBase*, 2> concrete_types_;
+};
+
+/// Struct holds the semantic information for structures.
+class Struct final : public Castable<Struct, StructBase> {
+ public:
+ /// Constructor
+ /// @param declaration the AST structure declaration
+ /// @param source the source of the structure
+ /// @param name the name of the structure
+ /// @param members the structure members
+ /// @param align the byte alignment of the structure
+ /// @param size the byte size of the structure
+ /// @param size_no_padding size of the members without the end of structure
+ /// alignment padding
+ Struct(const ast::Struct* declaration,
+ tint::Source source,
+ Symbol name,
+ utils::VectorRef<const StructMember*> members,
+ uint32_t align,
+ uint32_t size,
+ uint32_t size_no_padding);
+
+ /// Destructor
+ ~Struct() override;
+
+ /// @returns the struct
+ const ast::Struct* Declaration() const { return declaration_; }
+
+ /// @returns the members of the structure
+ utils::VectorRef<const StructMember*> Members() const {
+ return Base::Members().ReinterpretCast<const StructMember*>();
+ }
+
+ private:
+ ast::Struct const* const declaration_;
+};
+
+/// StructMemberBase holds the semantic information for structure members.
+class StructMemberBase : public Castable<StructMemberBase, Node> {
+ public:
+ /// Constructor
+ /// @param source the source of the struct member
+ /// @param name the name of the structure member
+ /// @param type the type of the member
+ /// @param index the index of the member in the structure
+ /// @param offset the byte offset from the base of the structure
+ /// @param align the byte alignment of the member
+ /// @param size the byte size of the member
+ /// @param location the location attribute, if present
+ StructMemberBase(tint::Source source,
+ Symbol name,
+ const sem::Type* type,
+ uint32_t index,
+ uint32_t offset,
+ uint32_t align,
+ uint32_t size,
+ std::optional<uint32_t> location);
+
+ /// Destructor
+ ~StructMemberBase() override;
+
+ /// @returns the source the struct member
+ const tint::Source& Source() const { return source_; }
+
+ /// @returns the name of the structure member
+ Symbol Name() const { return name_; }
+
+ /// Sets the owning structure to `s`
+ /// @param s the new structure owner
+ void SetStruct(const sem::StructBase* s) { struct_ = s; }
+
+ /// @returns the structure that owns this member
+ const sem::StructBase* Struct() const { return struct_; }
+
+ /// @returns the type of the member
+ const sem::Type* Type() const { return type_; }
+
+ /// @returns the member index
+ uint32_t Index() const { return index_; }
+
+ /// @returns byte offset from base of structure
+ uint32_t Offset() const { return offset_; }
+
+ /// @returns the alignment of the member in bytes
+ uint32_t Align() const { return align_; }
+
+ /// @returns byte size
+ uint32_t Size() const { return size_; }
+
+ /// @returns the location, if set
+ std::optional<uint32_t> Location() const { return location_; }
+
+ private:
+ const tint::Source source_;
+ const Symbol name_;
+ const sem::StructBase* struct_;
+ const sem::Type* type_;
+ const uint32_t index_;
+ const uint32_t offset_;
+ const uint32_t align_;
+ const uint32_t size_;
+ const std::optional<uint32_t> location_;
};
/// StructMember holds the semantic information for structure members.
-class StructMember final : public Castable<StructMember, Node> {
+class StructMember final : public Castable<StructMember, StructMemberBase> {
public:
/// Constructor
/// @param declaration the AST declaration node
@@ -207,48 +303,11 @@
/// @returns the AST declaration node
const ast::StructMember* Declaration() const { return declaration_; }
- /// @returns the source the struct member
- const tint::Source& Source() const { return source_; }
-
- /// @returns the name of the structure member
- Symbol Name() const { return name_; }
-
- /// Sets the owning structure to `s`
- /// @param s the new structure owner
- void SetStruct(const sem::Struct* s) { struct_ = s; }
-
/// @returns the structure that owns this member
- const sem::Struct* Struct() const { return struct_; }
-
- /// @returns the type of the member
- const sem::Type* Type() const { return type_; }
-
- /// @returns the member index
- uint32_t Index() const { return index_; }
-
- /// @returns byte offset from base of structure
- uint32_t Offset() const { return offset_; }
-
- /// @returns the alignment of the member in bytes
- uint32_t Align() const { return align_; }
-
- /// @returns byte size
- uint32_t Size() const { return size_; }
-
- /// @returns the location, if set
- std::optional<uint32_t> Location() const { return location_; }
+ const sem::Struct* Struct() const { return static_cast<const sem::Struct*>(Base::Struct()); }
private:
const ast::StructMember* const declaration_;
- const tint::Source source_;
- const Symbol name_;
- const sem::Struct* struct_;
- const sem::Type* type_;
- const uint32_t index_;
- const uint32_t offset_;
- const uint32_t align_;
- const uint32_t size_;
- const std::optional<uint32_t> location_;
};
} // namespace tint::sem
diff --git a/src/tint/sem/struct_test.cc b/src/tint/sem/struct_test.cc
index 6a88307..55744d1 100644
--- a/src/tint/sem/struct_test.cc
+++ b/src/tint/sem/struct_test.cc
@@ -26,8 +26,8 @@
auto name = Sym("S");
auto* impl = create<ast::Struct>(name, utils::Empty, utils::Empty);
auto* ptr = impl;
- auto* s = create<sem::Struct>(impl, impl->source, impl->name, StructMemberList{},
- 4u /* align */, 8u /* size */, 16u /* size_no_padding */);
+ auto* s = create<sem::Struct>(impl, impl->source, impl->name, utils::Empty, 4u /* align */,
+ 8u /* size */, 16u /* size_no_padding */);
EXPECT_EQ(s->Declaration(), ptr);
EXPECT_EQ(s->Align(), 4u);
EXPECT_EQ(s->Size(), 8u);
@@ -36,10 +36,10 @@
TEST_F(StructTest, Hash) {
auto* a_impl = create<ast::Struct>(Sym("a"), utils::Empty, utils::Empty);
- auto* a = create<sem::Struct>(a_impl, a_impl->source, a_impl->name, StructMemberList{},
+ auto* a = create<sem::Struct>(a_impl, a_impl->source, a_impl->name, utils::Empty,
4u /* align */, 4u /* size */, 4u /* size_no_padding */);
auto* b_impl = create<ast::Struct>(Sym("b"), utils::Empty, utils::Empty);
- auto* b = create<sem::Struct>(b_impl, b_impl->source, b_impl->name, StructMemberList{},
+ auto* b = create<sem::Struct>(b_impl, b_impl->source, b_impl->name, utils::Empty,
4u /* align */, 4u /* size */, 4u /* size_no_padding */);
EXPECT_NE(a->Hash(), b->Hash());
@@ -47,10 +47,10 @@
TEST_F(StructTest, Equals) {
auto* a_impl = create<ast::Struct>(Sym("a"), utils::Empty, utils::Empty);
- auto* a = create<sem::Struct>(a_impl, a_impl->source, a_impl->name, StructMemberList{},
+ auto* a = create<sem::Struct>(a_impl, a_impl->source, a_impl->name, utils::Empty,
4u /* align */, 4u /* size */, 4u /* size_no_padding */);
auto* b_impl = create<ast::Struct>(Sym("b"), utils::Empty, utils::Empty);
- auto* b = create<sem::Struct>(b_impl, b_impl->source, b_impl->name, StructMemberList{},
+ auto* b = create<sem::Struct>(b_impl, b_impl->source, b_impl->name, utils::Empty,
4u /* align */, 4u /* size */, 4u /* size_no_padding */);
EXPECT_TRUE(a->Equals(*a));
@@ -61,8 +61,8 @@
TEST_F(StructTest, FriendlyName) {
auto name = Sym("my_struct");
auto* impl = create<ast::Struct>(name, utils::Empty, utils::Empty);
- auto* s = create<sem::Struct>(impl, impl->source, impl->name, StructMemberList{},
- 4u /* align */, 4u /* size */, 4u /* size_no_padding */);
+ auto* s = create<sem::Struct>(impl, impl->source, impl->name, utils::Empty, 4u /* align */,
+ 4u /* size */, 4u /* size_no_padding */);
EXPECT_EQ(s->FriendlyName(Symbols()), "my_struct");
}
@@ -116,7 +116,7 @@
ASSERT_TRUE(p.IsValid()) << p.Diagnostics().str();
auto* sem = p.Sem().Get(st);
- ASSERT_EQ(2u, sem->Members().size());
+ ASSERT_EQ(2u, sem->Members().Length());
EXPECT_TRUE(sem->Members()[0]->Location().has_value());
EXPECT_EQ(sem->Members()[0]->Location().value(), 1u);
diff --git a/src/tint/sem/type.cc b/src/tint/sem/type.cc
index 3d25e7e..51afc12 100644
--- a/src/tint/sem/type.cc
+++ b/src/tint/sem/type.cc
@@ -239,7 +239,7 @@
return kNoConversion;
},
[&](const Struct* from_str) {
- auto& concrete_tys = from_str->ConcreteTypes();
+ auto concrete_tys = from_str->ConcreteTypes();
for (size_t i = 0; i < concrete_tys.Length(); i++) {
if (concrete_tys[i] == to) {
return static_cast<uint32_t>(i + 1);
diff --git a/src/tint/sem/type_test.cc b/src/tint/sem/type_test.cc
index 9148322..90e114f 100644
--- a/src/tint/sem/type_test.cc
+++ b/src/tint/sem/type_test.cc
@@ -47,7 +47,7 @@
const sem::Struct* str_f32 = create<Struct>(nullptr,
Source{},
Sym("str_f32"),
- StructMemberList{
+ utils::Vector{
create<StructMember>(
/* declaration */ nullptr,
/* source */ Source{},
@@ -65,7 +65,7 @@
const sem::Struct* str_f16 = create<Struct>(nullptr,
Source{},
Sym("str_f16"),
- StructMemberList{
+ utils::Vector{
create<StructMember>(
/* declaration */ nullptr,
/* source */ Source{},
@@ -83,7 +83,7 @@
sem::Struct* str_af = create<Struct>(nullptr,
Source{},
Sym("str_af"),
- StructMemberList{
+ utils::Vector{
create<StructMember>(
/* declaration */ nullptr,
/* source */ Source{},
diff --git a/src/tint/transform/array_length_from_uniform.cc b/src/tint/transform/array_length_from_uniform.cc
index 70097f2..e49c215 100644
--- a/src/tint/transform/array_length_from_uniform.cc
+++ b/src/tint/transform/array_length_from_uniform.cc
@@ -145,7 +145,7 @@
if (auto* str = storage_buffer_type->As<sem::Struct>()) {
// The variable is a struct, so subtract the byte offset of the array
// member.
- auto* array_member_sem = str->Members().back();
+ auto* array_member_sem = str->Members().Back();
array_type = array_member_sem->Type()->As<sem::Array>();
total_size = b.Sub(total_storage_buffer_size, u32(array_member_sem->Offset()));
} else if (auto* arr = storage_buffer_type->As<sem::Array>()) {
diff --git a/src/tint/transform/calculate_array_length.cc b/src/tint/transform/calculate_array_length.cc
index 9dcdd7b..a87e11d 100644
--- a/src/tint/transform/calculate_array_length.cc
+++ b/src/tint/transform/calculate_array_length.cc
@@ -207,7 +207,7 @@
[&](const sem::Struct* str) {
// The variable is a struct, so subtract the byte offset of
// the array member.
- auto* array_member_sem = str->Members().back();
+ auto* array_member_sem = str->Members().Back();
total_size = b.Sub(total_size, u32(array_member_sem->Offset()));
return array_member_sem->Type()->As<sem::Array>();
},
diff --git a/src/tint/transform/decompose_memory_access.cc b/src/tint/transform/decompose_memory_access.cc
index 671ebc4..1b674ba 100644
--- a/src/tint/transform/decompose_memory_access.cc
+++ b/src/tint/transform/decompose_memory_access.cc
@@ -714,7 +714,7 @@
TINT_ASSERT(Transform, str && str->Declaration() == nullptr);
utils::Vector<const ast::StructMember*, 8> ast_members;
- ast_members.Reserve(str->Members().size());
+ ast_members.Reserve(str->Members().Length());
for (auto& m : str->Members()) {
ast_members.Push(
b.Member(ctx.Clone(m->Name()), CreateASTTypeFor(ctx, m->Type())));
diff --git a/src/tint/transform/direct_variable_access.cc b/src/tint/transform/direct_variable_access.cc
index d1c1339..26c4aba 100644
--- a/src/tint/transform/direct_variable_access.cc
+++ b/src/tint/transform/direct_variable_access.cc
@@ -228,7 +228,8 @@
// will have the pointer parameters replaced with an array of u32s, used to perform the
// pointer indexing in the variant.
// Function call pointer arguments are replaced with an array of these dynamic indices.
- for (auto* decl : utils::Reverse(sem.Module()->DependencyOrderedDeclarations())) {
+ auto decls = sem.Module()->DependencyOrderedDeclarations();
+ for (auto* decl : utils::Reverse(decls)) {
if (auto* fn = sem.Get<sem::Function>(decl)) {
auto* fn_info = FnInfoFor(fn);
ProcessFunction(fn, fn_info);
diff --git a/src/tint/transform/num_workgroups_from_uniform.cc b/src/tint/transform/num_workgroups_from_uniform.cc
index b6c93b2..c8f8322 100644
--- a/src/tint/transform/num_workgroups_from_uniform.cc
+++ b/src/tint/transform/num_workgroups_from_uniform.cc
@@ -117,7 +117,7 @@
ctx.Remove(str->Declaration()->members, member->Declaration());
// If this is the only member, remove the struct and parameter too.
- if (str->Members().size() == 1) {
+ if (str->Members().Length() == 1) {
ctx.Remove(func->params, param->Declaration());
ctx.Remove(src->AST().GlobalDeclarations(), str->Declaration());
}
diff --git a/src/tint/transform/preserve_padding.cc b/src/tint/transform/preserve_padding.cc
new file mode 100644
index 0000000..d02dacc
--- /dev/null
+++ b/src/tint/transform/preserve_padding.cc
@@ -0,0 +1,226 @@
+// Copyright 2022 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/transform/preserve_padding.h"
+
+#include <unordered_set>
+#include <utility>
+
+#include "src/tint/program_builder.h"
+#include "src/tint/sem/reference.h"
+#include "src/tint/sem/struct.h"
+#include "src/tint/utils/map.h"
+#include "src/tint/utils/vector.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::transform::PreservePadding);
+
+using namespace tint::number_suffixes; // NOLINT
+
+namespace tint::transform {
+
+PreservePadding::PreservePadding() = default;
+
+PreservePadding::~PreservePadding() = default;
+
+/// The PIMPL state for the PreservePadding transform
+struct PreservePadding::State {
+ /// Constructor
+ /// @param src the source Program
+ explicit State(const Program* src) : ctx{&b, src, /* auto_clone_symbols */ true} {}
+
+ /// The main function for the transform.
+ /// @returns the ApplyResult
+ ApplyResult Run() {
+ // Gather a list of assignments that need to be transformed.
+ std::unordered_set<const ast::AssignmentStatement*> assignments_to_transform;
+ for (auto* node : ctx.src->ASTNodes().Objects()) {
+ Switch(
+ node, //
+ [&](const ast::AssignmentStatement* assign) {
+ auto* ty = sem.Get(assign->lhs)->Type();
+ if (assign->lhs->Is<ast::PhonyExpression>()) {
+ // Ignore phony assignment.
+ return;
+ }
+ if (ty->As<sem::Reference>()->AddressSpace() != ast::AddressSpace::kStorage) {
+ // We only care about assignments that write to variables in the storage
+ // address space, as nothing else is host-visible.
+ return;
+ }
+ if (HasPadding(ty->UnwrapRef())) {
+ // The assigned type has padding bytes, so we need to decompose the writes.
+ assignments_to_transform.insert(assign);
+ }
+ },
+ [&](const ast::Enable* enable) {
+ // Check if the full pointer parameters extension is already enabled.
+ if (enable->extension ==
+ ast::Extension::kChromiumExperimentalFullPtrParameters) {
+ ext_enabled = true;
+ }
+ });
+ }
+ if (assignments_to_transform.empty()) {
+ return SkipTransform;
+ }
+
+ // Replace all assignments that include padding with decomposed versions.
+ ctx.ReplaceAll([&](const ast::AssignmentStatement* assign) -> const ast::Statement* {
+ if (!assignments_to_transform.count(assign)) {
+ return nullptr;
+ }
+ auto* ty = sem.Get(assign->lhs)->Type()->UnwrapRef();
+ return MakeAssignment(ty, ctx.Clone(assign->lhs), ctx.Clone(assign->rhs));
+ });
+
+ ctx.Clone();
+ return Program(std::move(b));
+ }
+
+ /// Create a statement that will perform the assignment `lhs = rhs`, creating and using helper
+ /// functions to decompose the assignment into element-wise copies if needed.
+ /// @param ty the type of the assignment
+ /// @param lhs the lhs expression (in the destination program)
+ /// @param rhs the rhs expression (in the destination program)
+ /// @returns the statement that performs the assignment
+ const ast::Statement* MakeAssignment(const sem::Type* ty,
+ const ast::Expression* lhs,
+ const ast::Expression* rhs) {
+ if (!HasPadding(ty)) {
+ // No padding - use a regular assignment.
+ return b.Assign(lhs, rhs);
+ }
+
+ // Call (and create if necessary) a helper function that assigns a composite using the
+ // statements in `body`. The helper will have the form:
+ // fn assign_helper_T(dest : ptr<storage, T, read_write>, value : T) {
+ // <body>
+ // }
+ // It will be called by passing a pointer to the original LHS:
+ // assign_helper_T(&lhs, rhs);
+ //
+ // Since this requires passing pointers to the storage address space, this will also enable
+ // the chromium_experimental_full_ptr_parameters extension.
+ constexpr const char* kDestParamName = "dest";
+ constexpr const char* kValueParamName = "value";
+ auto call_helper = [&](auto&& body) {
+ EnableExtension();
+ auto helper = helpers.GetOrCreate(ty, [&]() {
+ auto helper_name = b.Symbols().New("assign_and_preserve_padding");
+ utils::Vector<const ast::Parameter*, 2> params = {
+ b.Param(kDestParamName,
+ b.ty.pointer(CreateASTTypeFor(ctx, ty), ast::AddressSpace::kStorage,
+ ast::Access::kReadWrite)),
+ b.Param(kValueParamName, CreateASTTypeFor(ctx, ty)),
+ };
+ b.Func(helper_name, params, b.ty.void_(), body());
+ return helper_name;
+ });
+ return b.CallStmt(b.Call(helper, b.AddressOf(lhs), rhs));
+ };
+
+ return Switch(
+ ty, //
+ [&](const sem::Array* arr) {
+ // Call a helper function that uses a loop to assigns each element separately.
+ return call_helper([&]() {
+ utils::Vector<const ast::Statement*, 8> body;
+ auto* idx = b.Var("i", b.Expr(0_u));
+ body.Push(
+ b.For(b.Decl(idx), b.LessThan(idx, u32(arr->ConstantCount().value())),
+ b.Assign(idx, b.Add(idx, 1_u)),
+ b.Block(MakeAssignment(arr->ElemType(),
+ b.IndexAccessor(b.Deref(kDestParamName), idx),
+ b.IndexAccessor(kValueParamName, idx)))));
+ return body;
+ });
+ },
+ [&](const sem::Struct* str) {
+ // Call a helper function that assigns each member separately.
+ return call_helper([&]() {
+ utils::Vector<const ast::Statement*, 8> body;
+ for (auto member : str->Members()) {
+ auto name = sym.NameFor(member->Declaration()->symbol);
+ body.Push(MakeAssignment(member->Type(),
+ b.MemberAccessor(b.Deref(kDestParamName), name),
+ b.MemberAccessor(kValueParamName, name)));
+ }
+ return body;
+ });
+ },
+ [&](Default) {
+ TINT_ICE(Transform, b.Diagnostics()) << "unhandled type with padding";
+ return nullptr;
+ });
+ }
+
+ /// Checks if a type contains padding bytes.
+ /// @param ty the type to check
+ /// @returns true if `ty` (or any of its contained types) have padding bytes
+ bool HasPadding(const sem::Type* ty) {
+ return Switch(
+ ty, //
+ [&](const sem::Array* arr) {
+ auto* elem_ty = arr->ElemType();
+ if (elem_ty->Size() % elem_ty->Align() > 0) {
+ return true;
+ }
+ return HasPadding(elem_ty);
+ },
+ [&](const sem::Struct* str) {
+ uint32_t current_offset = 0;
+ for (auto* member : str->Members()) {
+ if (member->Offset() > current_offset) {
+ return true;
+ }
+ if (HasPadding(member->Type())) {
+ return true;
+ }
+ current_offset += member->Type()->Size();
+ }
+ return (current_offset < str->Size());
+ },
+ [&](Default) { return false; });
+ }
+
+ /// Enable the full pointer parameters extension, if we have not already done so.
+ void EnableExtension() {
+ if (!ext_enabled) {
+ b.Enable(ast::Extension::kChromiumExperimentalFullPtrParameters);
+ ext_enabled = true;
+ }
+ }
+
+ private:
+ /// The program builder
+ ProgramBuilder b;
+ /// The clone context
+ CloneContext ctx;
+ /// Alias to the semantic info in ctx.src
+ const sem::Info& sem = ctx.src->Sem();
+ /// Alias to the symbols in ctx.src
+ const SymbolTable& sym = ctx.src->Symbols();
+ /// Flag to track whether we have already enabled the full pointer parameters extension.
+ bool ext_enabled = false;
+ /// Map of semantic types to their assignment helper functions.
+ utils::Hashmap<const sem::Type*, Symbol, 8> helpers;
+};
+
+Transform::ApplyResult PreservePadding::Apply(const Program* program,
+ const DataMap&,
+ DataMap&) const {
+ return State(program).Run();
+}
+
+} // namespace tint::transform
diff --git a/src/tint/transform/preserve_padding.h b/src/tint/transform/preserve_padding.h
new file mode 100644
index 0000000..3bf0a35
--- /dev/null
+++ b/src/tint/transform/preserve_padding.h
@@ -0,0 +1,47 @@
+// Copyright 2022 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_TRANSFORM_PRESERVE_PADDING_H_
+#define SRC_TINT_TRANSFORM_PRESERVE_PADDING_H_
+
+#include "src/tint/transform/transform.h"
+
+namespace tint::transform {
+
+/// Decompose assignments of whole structure and array types to preserve padding bytes.
+///
+/// WGSL states that memory operations on structures and arrays will not access padding bytes. To
+/// avoid overwriting padding bytes when writing to buffers, this transform decomposes those
+/// assignments into element-wise assignments via helper functions.
+///
+/// @note Assumes that the DirectVariableTransform will be run afterwards for backends that need it.
+class PreservePadding final : public Castable<PreservePadding, Transform> {
+ public:
+ /// Constructor
+ PreservePadding();
+ /// Destructor
+ ~PreservePadding() override;
+
+ /// @copydoc Transform::Apply
+ ApplyResult Apply(const Program* program,
+ const DataMap& inputs,
+ DataMap& outputs) const override;
+
+ private:
+ struct State;
+};
+
+} // namespace tint::transform
+
+#endif // SRC_TINT_TRANSFORM_PRESERVE_PADDING_H_
diff --git a/src/tint/transform/preserve_padding_test.cc b/src/tint/transform/preserve_padding_test.cc
new file mode 100644
index 0000000..f47cb3e
--- /dev/null
+++ b/src/tint/transform/preserve_padding_test.cc
@@ -0,0 +1,677 @@
+// Copyright 2022 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/transform/preserve_padding.h"
+
+#include <utility>
+
+#include "src/tint/transform/test_helper.h"
+
+namespace tint::transform {
+namespace {
+
+using PreservePaddingTest = TransformTest;
+
+TEST_F(PreservePaddingTest, ShouldRun_EmptyModule) {
+ auto* src = R"()";
+
+ EXPECT_FALSE(ShouldRun<PreservePadding>(src));
+}
+
+TEST_F(PreservePaddingTest, ShouldRun_NonStructVec3) {
+ auto* src = R"(
+@group(0) @binding(0) var<storage, read_write> v : vec3<u32>;
+
+@compute @workgroup_size(1)
+fn foo() {
+ v = vec3<u32>();
+}
+ )";
+
+ EXPECT_FALSE(ShouldRun<PreservePadding>(src));
+}
+
+TEST_F(PreservePaddingTest, ShouldRun_StructWithoutPadding) {
+ auto* src = R"(
+struct S {
+ a : u32,
+ b : u32,
+ c : u32,
+ d : u32,
+ e : vec3<u32>,
+ f : u32,
+}
+
+@group(0) @binding(0) var<storage, read_write> v : S;
+
+@compute @workgroup_size(1)
+fn foo() {
+ v = S();
+}
+ )";
+
+ EXPECT_FALSE(ShouldRun<PreservePadding>(src));
+}
+
+TEST_F(PreservePaddingTest, ShouldRun_ArrayWithoutPadding) {
+ auto* src = R"(
+@group(0) @binding(0) var<storage, read_write> v : array<vec4<u32>, 4>;
+
+@compute @workgroup_size(1)
+fn foo() {
+ v = array<vec4<u32>, 4>();
+}
+ )";
+
+ EXPECT_FALSE(ShouldRun<PreservePadding>(src));
+}
+
+TEST_F(PreservePaddingTest, EmptyModule) {
+ auto* src = R"()";
+
+ auto* expect = src;
+
+ auto got = Run<PreservePadding>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PreservePaddingTest, StructTrailingPadding) {
+ auto* src = R"(
+struct S {
+ a : u32,
+ b : u32,
+ c : u32,
+ d : u32,
+ e : vec3<u32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> v : S;
+
+@compute @workgroup_size(1)
+fn foo() {
+ v = S();
+}
+)";
+
+ auto* expect = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+struct S {
+ a : u32,
+ b : u32,
+ c : u32,
+ d : u32,
+ e : vec3<u32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> v : S;
+
+fn assign_and_preserve_padding(dest : ptr<storage, S, read_write>, value : S) {
+ (*(dest)).a = value.a;
+ (*(dest)).b = value.b;
+ (*(dest)).c = value.c;
+ (*(dest)).d = value.d;
+ (*(dest)).e = value.e;
+}
+
+@compute @workgroup_size(1)
+fn foo() {
+ assign_and_preserve_padding(&(v), S());
+}
+)";
+
+ auto got = Run<PreservePadding>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PreservePaddingTest, StructInternalPadding) {
+ auto* src = R"(
+struct S {
+ a : u32,
+ b : vec4<u32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> v : S;
+
+@compute @workgroup_size(1)
+fn foo() {
+ v = S();
+}
+)";
+
+ auto* expect = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+struct S {
+ a : u32,
+ b : vec4<u32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> v : S;
+
+fn assign_and_preserve_padding(dest : ptr<storage, S, read_write>, value : S) {
+ (*(dest)).a = value.a;
+ (*(dest)).b = value.b;
+}
+
+@compute @workgroup_size(1)
+fn foo() {
+ assign_and_preserve_padding(&(v), S());
+}
+)";
+
+ auto got = Run<PreservePadding>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PreservePaddingTest, StructExplicitSize_TrailingPadding) {
+ auto* src = R"(
+struct S {
+ @size(16) a : u32,
+}
+
+@group(0) @binding(0) var<storage, read_write> v : S;
+
+@compute @workgroup_size(1)
+fn foo() {
+ v = S();
+}
+)";
+
+ auto* expect = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+struct S {
+ @size(16)
+ a : u32,
+}
+
+@group(0) @binding(0) var<storage, read_write> v : S;
+
+fn assign_and_preserve_padding(dest : ptr<storage, S, read_write>, value : S) {
+ (*(dest)).a = value.a;
+}
+
+@compute @workgroup_size(1)
+fn foo() {
+ assign_and_preserve_padding(&(v), S());
+}
+)";
+
+ auto got = Run<PreservePadding>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PreservePaddingTest, StructExplicitSize_InternalPadding) {
+ auto* src = R"(
+struct S {
+ @size(16) a : u32,
+ b : u32,
+}
+
+@group(0) @binding(0) var<storage, read_write> v : S;
+
+@compute @workgroup_size(1)
+fn foo() {
+ v = S();
+}
+)";
+
+ auto* expect = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+struct S {
+ @size(16)
+ a : u32,
+ b : u32,
+}
+
+@group(0) @binding(0) var<storage, read_write> v : S;
+
+fn assign_and_preserve_padding(dest : ptr<storage, S, read_write>, value : S) {
+ (*(dest)).a = value.a;
+ (*(dest)).b = value.b;
+}
+
+@compute @workgroup_size(1)
+fn foo() {
+ assign_and_preserve_padding(&(v), S());
+}
+)";
+
+ auto got = Run<PreservePadding>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PreservePaddingTest, NestedStructs) {
+ auto* src = R"(
+struct S1 {
+ a1 : u32,
+ b1 : vec3<u32>,
+ c1 : u32,
+}
+
+struct S2 {
+ a2 : u32,
+ b2 : S1,
+ c2 : S1,
+}
+
+struct S3 {
+ a3 : S1,
+ b3 : S2,
+ c3 : S2,
+}
+
+@group(0) @binding(0) var<storage, read_write> v : S3;
+
+@compute @workgroup_size(1)
+fn foo() {
+ v = S3();
+}
+)";
+
+ auto* expect = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+struct S1 {
+ a1 : u32,
+ b1 : vec3<u32>,
+ c1 : u32,
+}
+
+struct S2 {
+ a2 : u32,
+ b2 : S1,
+ c2 : S1,
+}
+
+struct S3 {
+ a3 : S1,
+ b3 : S2,
+ c3 : S2,
+}
+
+@group(0) @binding(0) var<storage, read_write> v : S3;
+
+fn assign_and_preserve_padding_1(dest : ptr<storage, S1, read_write>, value : S1) {
+ (*(dest)).a1 = value.a1;
+ (*(dest)).b1 = value.b1;
+ (*(dest)).c1 = value.c1;
+}
+
+fn assign_and_preserve_padding_2(dest : ptr<storage, S2, read_write>, value : S2) {
+ (*(dest)).a2 = value.a2;
+ assign_and_preserve_padding_1(&((*(dest)).b2), value.b2);
+ assign_and_preserve_padding_1(&((*(dest)).c2), value.c2);
+}
+
+fn assign_and_preserve_padding(dest : ptr<storage, S3, read_write>, value : S3) {
+ assign_and_preserve_padding_1(&((*(dest)).a3), value.a3);
+ assign_and_preserve_padding_2(&((*(dest)).b3), value.b3);
+ assign_and_preserve_padding_2(&((*(dest)).c3), value.c3);
+}
+
+@compute @workgroup_size(1)
+fn foo() {
+ assign_and_preserve_padding(&(v), S3());
+}
+)";
+
+ auto got = Run<PreservePadding>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PreservePaddingTest, ArrayOfVec3) {
+ auto* src = R"(
+@group(0) @binding(0) var<storage, read_write> v : array<vec3<u32>, 4>;
+
+@compute @workgroup_size(1)
+fn foo() {
+ v = array<vec3<u32>, 4>();
+}
+)";
+
+ auto* expect = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+@group(0) @binding(0) var<storage, read_write> v : array<vec3<u32>, 4>;
+
+fn assign_and_preserve_padding(dest : ptr<storage, array<vec3<u32>, 4u>, read_write>, value : array<vec3<u32>, 4u>) {
+ for(var i = 0u; (i < 4u); i = (i + 1u)) {
+ (*(dest))[i] = value[i];
+ }
+}
+
+@compute @workgroup_size(1)
+fn foo() {
+ assign_and_preserve_padding(&(v), array<vec3<u32>, 4>());
+}
+)";
+
+ auto got = Run<PreservePadding>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PreservePaddingTest, ArrayOfArray) {
+ auto* src = R"(
+type Array = array<array<vec3<u32>, 4>, 3>;
+
+@group(0) @binding(0) var<storage, read_write> v : Array;
+
+@compute @workgroup_size(1)
+fn foo() {
+ v = Array();
+}
+)";
+
+ auto* expect = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+type Array = array<array<vec3<u32>, 4>, 3>;
+
+@group(0) @binding(0) var<storage, read_write> v : Array;
+
+fn assign_and_preserve_padding_1(dest : ptr<storage, array<vec3<u32>, 4u>, read_write>, value : array<vec3<u32>, 4u>) {
+ for(var i = 0u; (i < 4u); i = (i + 1u)) {
+ (*(dest))[i] = value[i];
+ }
+}
+
+fn assign_and_preserve_padding(dest : ptr<storage, array<array<vec3<u32>, 4u>, 3u>, read_write>, value : array<array<vec3<u32>, 4u>, 3u>) {
+ for(var i = 0u; (i < 3u); i = (i + 1u)) {
+ assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
+ }
+}
+
+@compute @workgroup_size(1)
+fn foo() {
+ assign_and_preserve_padding(&(v), Array());
+}
+)";
+
+ auto got = Run<PreservePadding>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PreservePaddingTest, ArrayOfStructOfArray) {
+ auto* src = R"(
+struct S {
+ a : u32,
+ b : array<vec3<u32>, 4>,
+}
+
+@group(0) @binding(0) var<storage, read_write> v : array<S, 3>;
+
+@compute @workgroup_size(1)
+fn foo() {
+ v = array<S, 3>();
+}
+)";
+
+ auto* expect = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+struct S {
+ a : u32,
+ b : array<vec3<u32>, 4>,
+}
+
+@group(0) @binding(0) var<storage, read_write> v : array<S, 3>;
+
+fn assign_and_preserve_padding_2(dest : ptr<storage, array<vec3<u32>, 4u>, read_write>, value : array<vec3<u32>, 4u>) {
+ for(var i = 0u; (i < 4u); i = (i + 1u)) {
+ (*(dest))[i] = value[i];
+ }
+}
+
+fn assign_and_preserve_padding_1(dest : ptr<storage, S, read_write>, value : S) {
+ (*(dest)).a = value.a;
+ assign_and_preserve_padding_2(&((*(dest)).b), value.b);
+}
+
+fn assign_and_preserve_padding(dest : ptr<storage, array<S, 3u>, read_write>, value : array<S, 3u>) {
+ for(var i = 0u; (i < 3u); i = (i + 1u)) {
+ assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
+ }
+}
+
+@compute @workgroup_size(1)
+fn foo() {
+ assign_and_preserve_padding(&(v), array<S, 3>());
+}
+)";
+
+ auto got = Run<PreservePadding>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PreservePaddingTest, NoModify_Vec3) {
+ auto* src = R"(
+@group(0) @binding(0) var<storage, read_write> v : vec3<u32>;
+
+@compute @workgroup_size(1)
+fn foo() {
+ v = vec3<u32>();
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<PreservePadding>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PreservePaddingTest, AvoidDuplicateEnables) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+struct S {
+ @size(16) a : u32,
+}
+
+@group(0) @binding(0) var<storage, read_write> v : S;
+
+@compute @workgroup_size(1)
+fn foo() {
+ v = S();
+}
+)";
+
+ auto* expect = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+struct S {
+ @size(16)
+ a : u32,
+}
+
+@group(0) @binding(0) var<storage, read_write> v : S;
+
+fn assign_and_preserve_padding(dest : ptr<storage, S, read_write>, value : S) {
+ (*(dest)).a = value.a;
+}
+
+@compute @workgroup_size(1)
+fn foo() {
+ assign_and_preserve_padding(&(v), S());
+}
+)";
+
+ auto got = Run<PreservePadding>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PreservePaddingTest, NoModify_Mat3x3) {
+ auto* src = R"(
+@group(0) @binding(0) var<storage, read_write> v : mat3x3<f32>;
+
+@compute @workgroup_size(1)
+fn foo() {
+ v = mat3x3<f32>();
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<PreservePadding>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PreservePaddingTest, NoModify_StructNoPadding) {
+ auto* src = R"(
+struct S {
+ a : u32,
+ b : u32,
+ c : u32,
+ d : u32,
+ e : vec4<u32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> v : S;
+
+@compute @workgroup_size(1)
+fn foo() {
+ v = S();
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<PreservePadding>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PreservePaddingTest, NoModify_ArrayNoPadding) {
+ auto* src = R"(
+@group(0) @binding(0) var<storage, read_write> v : array<vec4<u32>, 4>;
+
+@compute @workgroup_size(1)
+fn foo() {
+ v = array<vec4<u32>, 4>();
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<PreservePadding>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PreservePaddingTest, NoModify_ArrayOfStructNoPadding) {
+ auto* src = R"(
+struct S {
+ a : u32,
+ b : u32,
+ c : u32,
+ d : u32,
+ e : vec4<u32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> v : array<S, 4>;
+
+@compute @workgroup_size(1)
+fn foo() {
+ v = array<S, 4>();
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<PreservePadding>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PreservePaddingTest, NoModify_Workgroup) {
+ auto* src = R"(
+struct S {
+ a : u32,
+ b : vec3<u32>,
+}
+
+var<workgroup> v : S;
+
+@compute @workgroup_size(1)
+fn foo() {
+ v = S();
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<PreservePadding>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PreservePaddingTest, NoModify_Private) {
+ auto* src = R"(
+struct S {
+ a : u32,
+ b : vec3<u32>,
+}
+
+var<private> v : S;
+
+@compute @workgroup_size(1)
+fn foo() {
+ v = S();
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<PreservePadding>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PreservePaddingTest, NoModify_Function) {
+ auto* src = R"(
+struct S {
+ a : u32,
+ b : vec3<u32>,
+}
+
+@compute @workgroup_size(1)
+fn foo() {
+ var<function> v : S;
+ v = S();
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<PreservePadding>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+} // namespace
+} // namespace tint::transform
diff --git a/src/tint/transform/spirv_atomic.cc b/src/tint/transform/spirv_atomic.cc
index de3cdab..72304e6 100644
--- a/src/tint/transform/spirv_atomic.cc
+++ b/src/tint/transform/spirv_atomic.cc
@@ -64,6 +64,8 @@
/// Runs the transform
/// @returns the new program or SkipTransform if the transform is not required
ApplyResult Run() {
+ bool made_changes = false;
+
// Look for stub functions generated by the SPIR-V reader, which are used as placeholders
// for atomic builtin calls.
for (auto* fn : ctx.src->AST().Functions()) {
@@ -104,10 +106,11 @@
// Remove the stub from the output program
ctx.Remove(ctx.src->AST().GlobalDeclarations(), fn);
+ made_changes = true;
}
}
- if (atomic_expressions.IsEmpty()) {
+ if (!made_changes) {
return SkipTransform;
}
diff --git a/src/tint/transform/spirv_atomic_test.cc b/src/tint/transform/spirv_atomic_test.cc
index d9371bf..36312f7 100644
--- a/src/tint/transform/spirv_atomic_test.cc
+++ b/src/tint/transform/spirv_atomic_test.cc
@@ -151,6 +151,19 @@
std::vector<std::unique_ptr<Source::File>> files_;
};
+TEST_F(SpirvAtomicTest, StripUnusedBuiltins) {
+ auto* src = R"(
+fn f() {
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(SpirvAtomicTest, ArrayOfU32) {
auto* src = R"(
var<workgroup> wg : array<u32, 4>;
diff --git a/src/tint/transform/transform_test.cc b/src/tint/transform/transform_test.cc
index 29a21e1..001d38d 100644
--- a/src/tint/transform/transform_test.cc
+++ b/src/tint/transform/transform_test.cc
@@ -122,8 +122,8 @@
TEST_F(CreateASTTypeForTest, Struct) {
auto* str = create([](ProgramBuilder& b) {
auto* decl = b.Structure("S", {});
- return b.create<sem::Struct>(decl, decl->source, decl->name, sem::StructMemberList{},
- 4u /* align */, 4u /* size */, 4u /* size_no_padding */);
+ return b.create<sem::Struct>(decl, decl->source, decl->name, utils::Empty, 4u /* align */,
+ 4u /* size */, 4u /* size_no_padding */);
});
ASSERT_TRUE(str->Is<ast::TypeName>());
EXPECT_EQ(ast_type_builder.Symbols().NameFor(str->As<ast::TypeName>()->name), "S");
diff --git a/src/tint/transform/vertex_pulling.cc b/src/tint/transform/vertex_pulling.cc
index 0196002..e213ac6 100644
--- a/src/tint/transform/vertex_pulling.cc
+++ b/src/tint/transform/vertex_pulling.cc
@@ -41,6 +41,7 @@
kU32,
kI32,
kF32,
+ kF16,
};
/// The data type of a vertex format.
@@ -138,6 +139,7 @@
bool IsTypeCompatible(AttributeWGSLType wgslType, VertexFormatType vertexFormatType) {
switch (wgslType.base_type) {
case BaseWGSLType::kF32:
+ case BaseWGSLType::kF16:
return (vertexFormatType.base_type == VertexDataType::kFloat);
case BaseWGSLType::kU32:
return (vertexFormatType.base_type == VertexDataType::kUInt);
@@ -149,19 +151,26 @@
}
AttributeWGSLType WGSLTypeOf(const sem::Type* ty) {
- if (ty->Is<sem::I32>()) {
- return {BaseWGSLType::kI32, 1};
- }
- if (ty->Is<sem::U32>()) {
- return {BaseWGSLType::kU32, 1};
- }
- if (ty->Is<sem::F32>()) {
- return {BaseWGSLType::kF32, 1};
- }
- if (auto* vec = ty->As<sem::Vector>()) {
- return {WGSLTypeOf(vec->type()).base_type, vec->Width()};
- }
- return {BaseWGSLType::kInvalid, 0};
+ return Switch(
+ ty,
+ [](const sem::I32*) -> AttributeWGSLType {
+ return {BaseWGSLType::kI32, 1};
+ },
+ [](const sem::U32*) -> AttributeWGSLType {
+ return {BaseWGSLType::kU32, 1};
+ },
+ [](const sem::F32*) -> AttributeWGSLType {
+ return {BaseWGSLType::kF32, 1};
+ },
+ [](const sem::F16*) -> AttributeWGSLType {
+ return {BaseWGSLType::kF16, 1};
+ },
+ [](const sem::Vector* vec) -> AttributeWGSLType {
+ return {WGSLTypeOf(vec->type()).base_type, vec->Width()};
+ },
+ [](Default) -> AttributeWGSLType {
+ return {BaseWGSLType::kInvalid, 0};
+ });
}
VertexFormatType VertexFormatTypeOf(VertexFormat format) {
@@ -378,9 +387,22 @@
// Load the attribute value according to vertex format and convert the element type
// of result to match target WGSL variable. The result of `Fetch` should be of WGSL
- // types `f32`, `i32`, `u32`, and their vectors.
+ // types `f32`, `i32`, `u32`, and their vectors, while WGSL variable can be of
+ // `f16`.
auto* fetch = Fetch(buffer_array_base, attribute_desc.offset, buffer_idx,
attribute_desc.format);
+ // Convert the fetched scalar/vector if WGSL variable is of `f16` types
+ if (var_dt.base_type == BaseWGSLType::kF16) {
+ // The type of the same element number of base type of target WGSL variable
+ const ast::Type* loaded_data_target_type;
+ if (fmt_dt.width == 1) {
+ loaded_data_target_type = b.ty.f16();
+ } else {
+ loaded_data_target_type = b.ty.vec(b.ty.f16(), fmt_dt.width);
+ }
+
+ fetch = b.Construct(loaded_data_target_type, fetch);
+ }
// The attribute value may not be of the desired vector width. If it is not, we'll
// need to either reduce the width with a swizzle, or append 0's and / or a 1.
diff --git a/src/tint/transform/vertex_pulling_test.cc b/src/tint/transform/vertex_pulling_test.cc
index 54c348e..a6dc2d2 100644
--- a/src/tint/transform/vertex_pulling_test.cc
+++ b/src/tint/transform/vertex_pulling_test.cc
@@ -736,6 +736,63 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(VertexPullingTest, FloatVectorAttributes_F16) {
+ auto* src = R"(
+enable f16;
+
+@vertex
+fn main(@location(0) var_a : vec2<f16>,
+ @location(1) var_b : vec3<f16>,
+ @location(2) var_c : vec4<f16>
+ ) -> @builtin(position) vec4<f32> {
+ return vec4<f32>();
+}
+)";
+
+ auto* expect = R"(
+enable f16;
+
+struct TintVertexData {
+ tint_vertex_data : array<u32>,
+}
+
+@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
+
+@binding(1) @group(4) var<storage, read> tint_pulling_vertex_buffer_1 : TintVertexData;
+
+@binding(2) @group(4) var<storage, read> tint_pulling_vertex_buffer_2 : TintVertexData;
+
+@vertex
+fn main(@builtin(vertex_index) tint_pulling_vertex_index : u32) -> @builtin(position) vec4<f32> {
+ var var_a : vec2<f16>;
+ var var_b : vec3<f16>;
+ var var_c : vec4<f16>;
+ {
+ let buffer_array_base_0 = (tint_pulling_vertex_index * 2u);
+ var_a = vec2<f16>(vec2<f32>(bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[buffer_array_base_0]), bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 1u)])));
+ let buffer_array_base_1 = (tint_pulling_vertex_index * 3u);
+ var_b = vec3<f16>(vec3<f32>(bitcast<f32>(tint_pulling_vertex_buffer_1.tint_vertex_data[buffer_array_base_1]), bitcast<f32>(tint_pulling_vertex_buffer_1.tint_vertex_data[(buffer_array_base_1 + 1u)]), bitcast<f32>(tint_pulling_vertex_buffer_1.tint_vertex_data[(buffer_array_base_1 + 2u)])));
+ let buffer_array_base_2 = (tint_pulling_vertex_index * 4u);
+ var_c = vec4<f16>(vec4<f32>(unpack2x16float(tint_pulling_vertex_buffer_2.tint_vertex_data[buffer_array_base_2]), unpack2x16float(tint_pulling_vertex_buffer_2.tint_vertex_data[(buffer_array_base_2 + 1u)])));
+ }
+ return vec4<f32>();
+}
+)";
+
+ VertexPulling::Config cfg;
+ cfg.vertex_state = {{
+ {8, VertexStepMode::kVertex, {{VertexFormat::kFloat32x2, 0, 0}}},
+ {12, VertexStepMode::kVertex, {{VertexFormat::kFloat32x3, 0, 1}}},
+ {16, VertexStepMode::kVertex, {{VertexFormat::kFloat16x4, 0, 2}}},
+ }};
+
+ DataMap data;
+ data.Add<VertexPulling::Config>(cfg);
+ auto got = Run<VertexPulling>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(VertexPullingTest, AttemptSymbolCollision) {
auto* src = R"(
@vertex
@@ -1019,6 +1076,104 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(VertexPullingTest, FormatsAligned_Float_F16) {
+ auto* src = R"(
+enable f16;
+
+@vertex
+fn main(
+ @location(0) unorm8x2 : vec2<f16>,
+ @location(1) unorm8x4 : vec4<f16>,
+ @location(2) snorm8x2 : vec2<f16>,
+ @location(3) snorm8x4 : vec4<f16>,
+ @location(4) unorm16x2 : vec2<f16>,
+ @location(5) unorm16x4 : vec4<f16>,
+ @location(6) snorm16x2 : vec2<f16>,
+ @location(7) snorm16x4 : vec4<f16>,
+ @location(8) float16x2 : vec2<f16>,
+ @location(9) float16x4 : vec4<f16>,
+ @location(10) float32 : f16,
+ @location(11) float32x2 : vec2<f16>,
+ @location(12) float32x3 : vec3<f16>,
+ @location(13) float32x4 : vec4<f16>,
+ ) -> @builtin(position) vec4<f32> {
+ return vec4<f32>(0.0, 0.0, 0.0, 1.0);
+}
+)";
+
+ auto* expect = R"(
+enable f16;
+
+struct TintVertexData {
+ tint_vertex_data : array<u32>,
+}
+
+@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
+
+@vertex
+fn main(@builtin(vertex_index) tint_pulling_vertex_index : u32) -> @builtin(position) vec4<f32> {
+ var unorm8x2 : vec2<f16>;
+ var unorm8x4 : vec4<f16>;
+ var snorm8x2 : vec2<f16>;
+ var snorm8x4 : vec4<f16>;
+ var unorm16x2 : vec2<f16>;
+ var unorm16x4 : vec4<f16>;
+ var snorm16x2 : vec2<f16>;
+ var snorm16x4 : vec4<f16>;
+ var float16x2 : vec2<f16>;
+ var float16x4 : vec4<f16>;
+ var float32 : f16;
+ var float32x2 : vec2<f16>;
+ var float32x3 : vec3<f16>;
+ var float32x4 : vec4<f16>;
+ {
+ let buffer_array_base_0 = (tint_pulling_vertex_index * 64u);
+ unorm8x2 = vec2<f16>(unpack4x8unorm((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] & 65535u)).xy);
+ unorm8x4 = vec4<f16>(unpack4x8unorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]));
+ snorm8x2 = vec2<f16>(unpack4x8snorm((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] & 65535u)).xy);
+ snorm8x4 = vec4<f16>(unpack4x8snorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]));
+ unorm16x2 = vec2<f16>(unpack2x16unorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]));
+ unorm16x4 = vec4<f16>(vec4<f32>(unpack2x16unorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]), unpack2x16unorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)])));
+ snorm16x2 = vec2<f16>(unpack2x16snorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]));
+ snorm16x4 = vec4<f16>(vec4<f32>(unpack2x16snorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]), unpack2x16snorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)])));
+ float16x2 = vec2<f16>(unpack2x16float(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]));
+ float16x4 = vec4<f16>(vec4<f32>(unpack2x16float(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]), unpack2x16float(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)])));
+ float32 = f16(bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]));
+ float32x2 = vec2<f16>(vec2<f32>(bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]), bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)])));
+ float32x3 = vec3<f16>(vec3<f32>(bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]), bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)]), bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 18u)])));
+ float32x4 = vec4<f16>(vec4<f32>(bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]), bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)]), bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 18u)]), bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 19u)])));
+ }
+ return vec4<f32>(0.0, 0.0, 0.0, 1.0);
+}
+)";
+
+ VertexPulling::Config cfg;
+ cfg.vertex_state = {{{256,
+ VertexStepMode::kVertex,
+ {
+ {VertexFormat::kUnorm8x2, 64, 0},
+ {VertexFormat::kUnorm8x4, 64, 1},
+ {VertexFormat::kSnorm8x2, 64, 2},
+ {VertexFormat::kSnorm8x4, 64, 3},
+ {VertexFormat::kUnorm16x2, 64, 4},
+ {VertexFormat::kUnorm16x4, 64, 5},
+ {VertexFormat::kSnorm16x2, 64, 6},
+ {VertexFormat::kSnorm16x4, 64, 7},
+ {VertexFormat::kFloat16x2, 64, 8},
+ {VertexFormat::kFloat16x4, 64, 9},
+ {VertexFormat::kFloat32, 64, 10},
+ {VertexFormat::kFloat32x2, 64, 11},
+ {VertexFormat::kFloat32x3, 64, 12},
+ {VertexFormat::kFloat32x4, 64, 13},
+ }}}};
+
+ DataMap data;
+ data.Add<VertexPulling::Config>(cfg);
+ auto got = Run<VertexPulling>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(VertexPullingTest, FormatsUnaligned_SInt) {
auto* src = R"(
@vertex
@@ -1253,6 +1408,104 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(VertexPullingTest, FormatsUnaligned_Float_F16) {
+ auto* src = R"(
+enable f16;
+
+@vertex
+fn main(
+ @location(0) unorm8x2 : vec2<f16>,
+ @location(1) unorm8x4 : vec4<f16>,
+ @location(2) snorm8x2 : vec2<f16>,
+ @location(3) snorm8x4 : vec4<f16>,
+ @location(4) unorm16x2 : vec2<f16>,
+ @location(5) unorm16x4 : vec4<f16>,
+ @location(6) snorm16x2 : vec2<f16>,
+ @location(7) snorm16x4 : vec4<f16>,
+ @location(8) float16x2 : vec2<f16>,
+ @location(9) float16x4 : vec4<f16>,
+ @location(10) float32 : f16,
+ @location(11) float32x2 : vec2<f16>,
+ @location(12) float32x3 : vec3<f16>,
+ @location(13) float32x4 : vec4<f16>,
+ ) -> @builtin(position) vec4<f32> {
+ return vec4<f32>(0.0, 0.0, 0.0, 1.0);
+}
+)";
+
+ auto* expect = R"(
+enable f16;
+
+struct TintVertexData {
+ tint_vertex_data : array<u32>,
+}
+
+@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
+
+@vertex
+fn main(@builtin(vertex_index) tint_pulling_vertex_index : u32) -> @builtin(position) vec4<f32> {
+ var unorm8x2 : vec2<f16>;
+ var unorm8x4 : vec4<f16>;
+ var snorm8x2 : vec2<f16>;
+ var snorm8x4 : vec4<f16>;
+ var unorm16x2 : vec2<f16>;
+ var unorm16x4 : vec4<f16>;
+ var snorm16x2 : vec2<f16>;
+ var snorm16x4 : vec4<f16>;
+ var float16x2 : vec2<f16>;
+ var float16x4 : vec4<f16>;
+ var float32 : f16;
+ var float32x2 : vec2<f16>;
+ var float32x3 : vec3<f16>;
+ var float32x4 : vec4<f16>;
+ {
+ let buffer_array_base_0 = (tint_pulling_vertex_index * 64u);
+ unorm8x2 = vec2<f16>(unpack4x8unorm((((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] << 8u) | (tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 15u)] >> 24u)) & 65535u)).xy);
+ unorm8x4 = vec4<f16>(unpack4x8unorm(((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 15u)] >> 24u) | (tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] << 8u))));
+ snorm8x2 = vec2<f16>(unpack4x8snorm((((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] << 8u) | (tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 15u)] >> 24u)) & 65535u)).xy);
+ snorm8x4 = vec4<f16>(unpack4x8snorm(((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 15u)] >> 24u) | (tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] << 8u))));
+ unorm16x2 = vec2<f16>(unpack2x16unorm(((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 15u)] >> 24u) | (tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] << 8u))));
+ unorm16x4 = vec4<f16>(vec4<f32>(unpack2x16unorm(((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 15u)] >> 24u) | (tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] << 8u))), unpack2x16unorm(((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] >> 24u) | (tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)] << 8u)))));
+ snorm16x2 = vec2<f16>(unpack2x16snorm(((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 15u)] >> 24u) | (tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] << 8u))));
+ snorm16x4 = vec4<f16>(vec4<f32>(unpack2x16snorm(((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 15u)] >> 24u) | (tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] << 8u))), unpack2x16snorm(((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] >> 24u) | (tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)] << 8u)))));
+ float16x2 = vec2<f16>(unpack2x16float(((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 15u)] >> 24u) | (tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] << 8u))));
+ float16x4 = vec4<f16>(vec4<f32>(unpack2x16float(((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 15u)] >> 24u) | (tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] << 8u))), unpack2x16float(((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] >> 24u) | (tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)] << 8u)))));
+ float32 = f16(bitcast<f32>(((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 15u)] >> 24u) | (tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] << 8u))));
+ float32x2 = vec2<f16>(vec2<f32>(bitcast<f32>(((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 15u)] >> 24u) | (tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] << 8u))), bitcast<f32>(((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] >> 24u) | (tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)] << 8u)))));
+ float32x3 = vec3<f16>(vec3<f32>(bitcast<f32>(((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 15u)] >> 24u) | (tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] << 8u))), bitcast<f32>(((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] >> 24u) | (tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)] << 8u))), bitcast<f32>(((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)] >> 24u) | (tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 18u)] << 8u)))));
+ float32x4 = vec4<f16>(vec4<f32>(bitcast<f32>(((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 15u)] >> 24u) | (tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] << 8u))), bitcast<f32>(((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] >> 24u) | (tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)] << 8u))), bitcast<f32>(((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)] >> 24u) | (tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 18u)] << 8u))), bitcast<f32>(((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 18u)] >> 24u) | (tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 19u)] << 8u)))));
+ }
+ return vec4<f32>(0.0, 0.0, 0.0, 1.0);
+}
+)";
+
+ VertexPulling::Config cfg;
+ cfg.vertex_state = {{{256,
+ VertexStepMode::kVertex,
+ {
+ {VertexFormat::kUnorm8x2, 63, 0},
+ {VertexFormat::kUnorm8x4, 63, 1},
+ {VertexFormat::kSnorm8x2, 63, 2},
+ {VertexFormat::kSnorm8x4, 63, 3},
+ {VertexFormat::kUnorm16x2, 63, 4},
+ {VertexFormat::kUnorm16x4, 63, 5},
+ {VertexFormat::kSnorm16x2, 63, 6},
+ {VertexFormat::kSnorm16x4, 63, 7},
+ {VertexFormat::kFloat16x2, 63, 8},
+ {VertexFormat::kFloat16x4, 63, 9},
+ {VertexFormat::kFloat32, 63, 10},
+ {VertexFormat::kFloat32x2, 63, 11},
+ {VertexFormat::kFloat32x3, 63, 12},
+ {VertexFormat::kFloat32x4, 63, 13},
+ }}}};
+
+ DataMap data;
+ data.Add<VertexPulling::Config>(cfg);
+ auto got = Run<VertexPulling>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(VertexPullingTest, FormatsWithVectorsResized_Padding_SInt) {
auto* src = R"(
@vertex
@@ -1511,6 +1764,112 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(VertexPullingTest, FormatsWithVectorsResized_Padding_Float_F16) {
+ auto* src = R"(
+enable f16;
+
+@vertex
+fn main(
+ @location(0) vec3_unorm8x2 : vec3<f16>,
+ @location(1) vec4_unorm8x2 : vec4<f16>,
+ @location(2) vec3_snorm8x2 : vec3<f16>,
+ @location(3) vec4_snorm8x2 : vec4<f16>,
+ @location(4) vec3_unorm16x2 : vec3<f16>,
+ @location(5) vec4_unorm16x2 : vec4<f16>,
+ @location(6) vec3_snorm16x2 : vec3<f16>,
+ @location(7) vec4_snorm16x2 : vec4<f16>,
+ @location(8) vec3_float16x2 : vec3<f16>,
+ @location(9) vec4_float16x2 : vec4<f16>,
+ @location(10) vec2_float32 : vec2<f16>,
+ @location(11) vec3_float32 : vec3<f16>,
+ @location(12) vec4_float32 : vec4<f16>,
+ @location(13) vec3_float32x2 : vec3<f16>,
+ @location(14) vec4_float32x2 : vec4<f16>,
+ @location(15) vec4_float32x3 : vec4<f16>,
+ ) -> @builtin(position) vec4<f32> {
+ return vec4<f32>(0.0, 0.0, 0.0, 1.0);
+}
+)";
+
+ auto* expect = R"(
+enable f16;
+
+struct TintVertexData {
+ tint_vertex_data : array<u32>,
+}
+
+@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
+
+@vertex
+fn main(@builtin(vertex_index) tint_pulling_vertex_index : u32) -> @builtin(position) vec4<f32> {
+ var vec3_unorm8x2 : vec3<f16>;
+ var vec4_unorm8x2 : vec4<f16>;
+ var vec3_snorm8x2 : vec3<f16>;
+ var vec4_snorm8x2 : vec4<f16>;
+ var vec3_unorm16x2 : vec3<f16>;
+ var vec4_unorm16x2 : vec4<f16>;
+ var vec3_snorm16x2 : vec3<f16>;
+ var vec4_snorm16x2 : vec4<f16>;
+ var vec3_float16x2 : vec3<f16>;
+ var vec4_float16x2 : vec4<f16>;
+ var vec2_float32 : vec2<f16>;
+ var vec3_float32 : vec3<f16>;
+ var vec4_float32 : vec4<f16>;
+ var vec3_float32x2 : vec3<f16>;
+ var vec4_float32x2 : vec4<f16>;
+ var vec4_float32x3 : vec4<f16>;
+ {
+ let buffer_array_base_0 = (tint_pulling_vertex_index * 64u);
+ vec3_unorm8x2 = vec3<f16>(vec2<f16>(unpack4x8unorm((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] & 65535u)).xy), 0.0);
+ vec4_unorm8x2 = vec4<f16>(vec2<f16>(unpack4x8unorm((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] & 65535u)).xy), 0.0, 1.0);
+ vec3_snorm8x2 = vec3<f16>(vec2<f16>(unpack4x8snorm((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] & 65535u)).xy), 0.0);
+ vec4_snorm8x2 = vec4<f16>(vec2<f16>(unpack4x8snorm((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] & 65535u)).xy), 0.0, 1.0);
+ vec3_unorm16x2 = vec3<f16>(vec2<f16>(unpack2x16unorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)])), 0.0);
+ vec4_unorm16x2 = vec4<f16>(vec2<f16>(unpack2x16unorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)])), 0.0, 1.0);
+ vec3_snorm16x2 = vec3<f16>(vec2<f16>(unpack2x16snorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)])), 0.0);
+ vec4_snorm16x2 = vec4<f16>(vec2<f16>(unpack2x16snorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)])), 0.0, 1.0);
+ vec3_float16x2 = vec3<f16>(vec2<f16>(unpack2x16float(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)])), 0.0);
+ vec4_float16x2 = vec4<f16>(vec2<f16>(unpack2x16float(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)])), 0.0, 1.0);
+ vec2_float32 = vec2<f16>(f16(bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)])), 0.0);
+ vec3_float32 = vec3<f16>(f16(bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)])), 0.0, 0.0);
+ vec4_float32 = vec4<f16>(f16(bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)])), 0.0, 0.0, 1.0);
+ vec3_float32x2 = vec3<f16>(vec2<f16>(vec2<f32>(bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]), bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)]))), 0.0);
+ vec4_float32x2 = vec4<f16>(vec2<f16>(vec2<f32>(bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]), bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)]))), 0.0, 1.0);
+ vec4_float32x3 = vec4<f16>(vec3<f16>(vec3<f32>(bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]), bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)]), bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 18u)]))), 1.0);
+ }
+ return vec4<f32>(0.0, 0.0, 0.0, 1.0);
+}
+)";
+
+ VertexPulling::Config cfg;
+ cfg.vertex_state = {{{256,
+ VertexStepMode::kVertex,
+ {
+ {VertexFormat::kUnorm8x2, 64, 0},
+ {VertexFormat::kUnorm8x2, 64, 1},
+ {VertexFormat::kSnorm8x2, 64, 2},
+ {VertexFormat::kSnorm8x2, 64, 3},
+ {VertexFormat::kUnorm16x2, 64, 4},
+ {VertexFormat::kUnorm16x2, 64, 5},
+ {VertexFormat::kSnorm16x2, 64, 6},
+ {VertexFormat::kSnorm16x2, 64, 7},
+ {VertexFormat::kFloat16x2, 64, 8},
+ {VertexFormat::kFloat16x2, 64, 9},
+ {VertexFormat::kFloat32, 64, 10},
+ {VertexFormat::kFloat32, 64, 11},
+ {VertexFormat::kFloat32, 64, 12},
+ {VertexFormat::kFloat32x2, 64, 13},
+ {VertexFormat::kFloat32x2, 64, 14},
+ {VertexFormat::kFloat32x3, 64, 15},
+ }}}};
+
+ DataMap data;
+ data.Add<VertexPulling::Config>(cfg);
+ auto got = Run<VertexPulling>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(VertexPullingTest, FormatsWithVectorsResized_Shrinking_SInt) {
auto* src = R"(
@vertex
@@ -1829,5 +2188,139 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(VertexPullingTest, FormatsWithVectorsResized_Shrinking_Float_F16) {
+ auto* src = R"(
+enable f16;
+
+@vertex
+fn main(
+ @location(0) sclr_unorm8x2 : f16 ,
+ @location(1) sclr_unorm8x4 : f16 ,
+ @location(2) vec2_unorm8x4 : vec2<f16>,
+ @location(3) vec3_unorm8x4 : vec3<f16>,
+ @location(4) sclr_snorm8x2 : f16 ,
+ @location(5) sclr_snorm8x4 : f16 ,
+ @location(6) vec2_snorm8x4 : vec2<f16>,
+ @location(7) vec3_snorm8x4 : vec3<f16>,
+ @location(8) sclr_unorm16x2 : f16 ,
+ @location(9) sclr_unorm16x4 : f16 ,
+ @location(10) vec2_unorm16x4 : vec2<f16>,
+ @location(11) vec3_unorm16x4 : vec3<f16>,
+ @location(12) sclr_snorm16x2 : f16 ,
+ @location(13) sclr_snorm16x4 : f16 ,
+ @location(14) vec2_snorm16x4 : vec2<f16>,
+ @location(15) vec3_snorm16x4 : vec3<f16>,
+ @location(16) sclr_float16x2 : f16 ,
+ @location(17) sclr_float16x4 : f16 ,
+ @location(18) vec2_float16x4 : vec2<f16>,
+ @location(19) vec3_float16x4 : vec3<f16>,
+ @location(20) sclr_float32x2 : f16 ,
+ @location(21) sclr_float32x3 : f16 ,
+ @location(22) vec2_float32x3 : vec2<f16>,
+ @location(23) sclr_float32x4 : f16 ,
+ @location(24) vec2_float32x4 : vec2<f16>,
+ @location(25) vec3_float32x4 : vec3<f16>,
+ ) -> @builtin(position) vec4<f32> {
+ return vec4<f32>(0.0, 0.0, 0.0, 1.0);
+}
+)";
+
+ auto* expect = R"(
+enable f16;
+
+struct TintVertexData {
+ tint_vertex_data : array<u32>,
+}
+
+@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
+
+@vertex
+fn main(@builtin(vertex_index) tint_pulling_vertex_index : u32) -> @builtin(position) vec4<f32> {
+ var sclr_unorm8x2 : f16;
+ var sclr_unorm8x4 : f16;
+ var vec2_unorm8x4 : vec2<f16>;
+ var vec3_unorm8x4 : vec3<f16>;
+ var sclr_snorm8x2 : f16;
+ var sclr_snorm8x4 : f16;
+ var vec2_snorm8x4 : vec2<f16>;
+ var vec3_snorm8x4 : vec3<f16>;
+ var sclr_unorm16x2 : f16;
+ var sclr_unorm16x4 : f16;
+ var vec2_unorm16x4 : vec2<f16>;
+ var vec3_unorm16x4 : vec3<f16>;
+ var sclr_snorm16x2 : f16;
+ var sclr_snorm16x4 : f16;
+ var vec2_snorm16x4 : vec2<f16>;
+ var vec3_snorm16x4 : vec3<f16>;
+ var sclr_float16x2 : f16;
+ var sclr_float16x4 : f16;
+ var vec2_float16x4 : vec2<f16>;
+ var vec3_float16x4 : vec3<f16>;
+ var sclr_float32x2 : f16;
+ var sclr_float32x3 : f16;
+ var vec2_float32x3 : vec2<f16>;
+ var sclr_float32x4 : f16;
+ var vec2_float32x4 : vec2<f16>;
+ var vec3_float32x4 : vec3<f16>;
+ {
+ let buffer_array_base_0 = (tint_pulling_vertex_index * 64u);
+ sclr_unorm8x2 = vec2<f16>(unpack4x8unorm((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] & 65535u)).xy).x;
+ sclr_unorm8x4 = vec4<f16>(unpack4x8unorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)])).x;
+ vec2_unorm8x4 = vec4<f16>(unpack4x8unorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)])).xy;
+ vec3_unorm8x4 = vec4<f16>(unpack4x8unorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)])).xyz;
+ sclr_snorm8x2 = vec2<f16>(unpack4x8snorm((tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)] & 65535u)).xy).x;
+ sclr_snorm8x4 = vec4<f16>(unpack4x8snorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)])).x;
+ vec2_snorm8x4 = vec4<f16>(unpack4x8snorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)])).xy;
+ vec3_snorm8x4 = vec4<f16>(unpack4x8snorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)])).xyz;
+ sclr_unorm16x2 = vec2<f16>(unpack2x16unorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)])).x;
+ sclr_unorm16x4 = vec4<f16>(vec4<f32>(unpack2x16unorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]), unpack2x16unorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)]))).x;
+ vec2_unorm16x4 = vec4<f16>(vec4<f32>(unpack2x16unorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]), unpack2x16unorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)]))).xy;
+ vec3_unorm16x4 = vec4<f16>(vec4<f32>(unpack2x16unorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]), unpack2x16unorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)]))).xyz;
+ sclr_snorm16x2 = vec2<f16>(unpack2x16snorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)])).x;
+ sclr_snorm16x4 = vec4<f16>(vec4<f32>(unpack2x16snorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]), unpack2x16snorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)]))).x;
+ vec2_snorm16x4 = vec4<f16>(vec4<f32>(unpack2x16snorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]), unpack2x16snorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)]))).xy;
+ vec3_snorm16x4 = vec4<f16>(vec4<f32>(unpack2x16snorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]), unpack2x16snorm(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)]))).xyz;
+ sclr_float16x2 = vec2<f16>(unpack2x16float(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)])).x;
+ sclr_float16x4 = vec4<f16>(vec4<f32>(unpack2x16float(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]), unpack2x16float(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)]))).x;
+ vec2_float16x4 = vec4<f16>(vec4<f32>(unpack2x16float(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]), unpack2x16float(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)]))).xy;
+ vec3_float16x4 = vec4<f16>(vec4<f32>(unpack2x16float(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]), unpack2x16float(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)]))).xyz;
+ sclr_float32x2 = vec2<f16>(vec2<f32>(bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]), bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)]))).x;
+ sclr_float32x3 = vec3<f16>(vec3<f32>(bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]), bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)]), bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 18u)]))).x;
+ vec2_float32x3 = vec3<f16>(vec3<f32>(bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]), bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)]), bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 18u)]))).xy;
+ sclr_float32x4 = vec4<f16>(vec4<f32>(bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]), bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)]), bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 18u)]), bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 19u)]))).x;
+ vec2_float32x4 = vec4<f16>(vec4<f32>(bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]), bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)]), bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 18u)]), bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 19u)]))).xy;
+ vec3_float32x4 = vec4<f16>(vec4<f32>(bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 16u)]), bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 17u)]), bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 18u)]), bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[(buffer_array_base_0 + 19u)]))).xyz;
+ }
+ return vec4<f32>(0.0, 0.0, 0.0, 1.0);
+}
+)";
+
+ VertexPulling::Config cfg;
+ cfg.vertex_state = {
+ {{256,
+ VertexStepMode::kVertex,
+ {
+ {VertexFormat::kUnorm8x2, 64, 0}, {VertexFormat::kUnorm8x4, 64, 1},
+ {VertexFormat::kUnorm8x4, 64, 2}, {VertexFormat::kUnorm8x4, 64, 3},
+ {VertexFormat::kSnorm8x2, 64, 4}, {VertexFormat::kSnorm8x4, 64, 5},
+ {VertexFormat::kSnorm8x4, 64, 6}, {VertexFormat::kSnorm8x4, 64, 7},
+ {VertexFormat::kUnorm16x2, 64, 8}, {VertexFormat::kUnorm16x4, 64, 9},
+ {VertexFormat::kUnorm16x4, 64, 10}, {VertexFormat::kUnorm16x4, 64, 11},
+ {VertexFormat::kSnorm16x2, 64, 12}, {VertexFormat::kSnorm16x4, 64, 13},
+ {VertexFormat::kSnorm16x4, 64, 14}, {VertexFormat::kSnorm16x4, 64, 15},
+ {VertexFormat::kFloat16x2, 64, 16}, {VertexFormat::kFloat16x4, 64, 17},
+ {VertexFormat::kFloat16x4, 64, 18}, {VertexFormat::kFloat16x4, 64, 19},
+ {VertexFormat::kFloat32x2, 64, 20}, {VertexFormat::kFloat32x3, 64, 21},
+ {VertexFormat::kFloat32x3, 64, 22}, {VertexFormat::kFloat32x4, 64, 23},
+ {VertexFormat::kFloat32x4, 64, 24}, {VertexFormat::kFloat32x4, 64, 25},
+ }}}};
+
+ DataMap data;
+ data.Add<VertexPulling::Config>(cfg);
+ auto got = Run<VertexPulling>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
} // namespace
} // namespace tint::transform
diff --git a/src/tint/utils/transform.h b/src/tint/utils/transform.h
index 2faca46..d96a4bb 100644
--- a/src/tint/utils/transform.h
+++ b/src/tint/utils/transform.h
@@ -91,8 +91,7 @@
/// @tparam N the small-array size of the returned Vector
/// @returns a new vector with each element of the source vector transformed by `transform`.
template <size_t N, typename IN, typename TRANSFORMER>
-auto Transform(const VectorRef<IN>& in, TRANSFORMER&& transform)
- -> Vector<decltype(transform(in[0])), N> {
+auto Transform(VectorRef<IN> in, TRANSFORMER&& transform) -> Vector<decltype(transform(in[0])), N> {
const auto count = in.Length();
Vector<decltype(transform(in[0])), N> result;
result.Reserve(count);
@@ -108,7 +107,7 @@
/// @tparam N the small-array size of the returned Vector
/// @returns a new vector with each element of the source vector transformed by `transform`.
template <size_t N, typename IN, typename TRANSFORMER>
-auto Transform(const VectorRef<IN>& in, TRANSFORMER&& transform)
+auto Transform(VectorRef<IN> in, TRANSFORMER&& transform)
-> Vector<decltype(transform(in[0], 1u)), N> {
const auto count = in.Length();
Vector<decltype(transform(in[0], 1u)), N> result;
diff --git a/src/tint/utils/unique_vector.h b/src/tint/utils/unique_vector.h
index bda090c..6cb8f88 100644
--- a/src/tint/utils/unique_vector.h
+++ b/src/tint/utils/unique_vector.h
@@ -87,8 +87,8 @@
/// @returns an iterator to the end of the reversed vector
auto rend() const { return vector.rend(); }
- /// @returns a const reference to the internal vector
- operator const Vector<T, N>&() const { return vector; }
+ /// @returns a reference to the internal vector
+ operator VectorRef<T>() const { return vector; }
/// @returns the std::move()'d vector.
/// @note The UniqueVector must not be used after calling this method
diff --git a/src/tint/utils/unique_vector_test.cc b/src/tint/utils/unique_vector_test.cc
index 9b015c2..c198615 100644
--- a/src/tint/utils/unique_vector_test.cc
+++ b/src/tint/utils/unique_vector_test.cc
@@ -94,11 +94,11 @@
unique_vec.Add(1);
unique_vec.Add(2);
- const utils::Vector<int, 4>& vec = unique_vec;
- EXPECT_EQ(vec.Length(), 3u);
+ utils::VectorRef<int> ref = unique_vec;
+ EXPECT_EQ(ref.Length(), 3u);
EXPECT_EQ(unique_vec.IsEmpty(), false);
int i = 0;
- for (auto n : vec) {
+ for (auto n : ref) {
EXPECT_EQ(n, i);
i++;
}
diff --git a/src/tint/utils/vector.h b/src/tint/utils/vector.h
index cefd536..83e1c12 100644
--- a/src/tint/utils/vector.h
+++ b/src/tint/utils/vector.h
@@ -27,6 +27,7 @@
#include "src/tint/castable.h"
#include "src/tint/traits.h"
#include "src/tint/utils/bitcast.h"
+#include "src/tint/utils/compiler_macros.h"
#include "src/tint/utils/string.h"
namespace tint::utils {
@@ -106,12 +107,20 @@
auto rend() const { return std::reverse_iterator<const T*>(begin()); }
};
+/// Mode enumerator for ReinterpretSlice
+enum class ReinterpretMode {
+ /// Only upcasts of pointers are permitted
+ kSafe,
+ /// Potentially unsafe downcasts of pointers are also permitted
+ kUnsafe,
+};
+
namespace detail {
/// Private implementation of tint::utils::CanReinterpretSlice.
/// Specialized for the case of TO equal to FROM, which is the common case, and avoids inspection of
/// the base classes, which can be troublesome if the slice is of an incomplete type.
-template <typename TO, typename FROM>
+template <ReinterpretMode MODE, typename TO, typename FROM>
struct CanReinterpretSlice {
/// True if a slice of FROM can be reinterpreted as a slice of TO
static constexpr bool value =
@@ -122,13 +131,14 @@
!std::is_const_v<std::remove_pointer_t<FROM>>)&& //
// TO and FROM are both Castable
IsCastable<std::remove_pointer_t<FROM>, std::remove_pointer_t<TO>> && //
- // FROM is of, or derives from TO
- traits::IsTypeOrDerived<std::remove_pointer_t<FROM>, std::remove_pointer_t<TO>>;
+ // MODE is kUnsafe, or FROM is of, or derives from TO
+ (MODE == ReinterpretMode::kUnsafe ||
+ traits::IsTypeOrDerived<std::remove_pointer_t<FROM>, std::remove_pointer_t<TO>>);
};
/// Specialization of 'CanReinterpretSlice' for when TO and FROM are equal types.
-template <typename T>
-struct CanReinterpretSlice<T, T> {
+template <typename T, ReinterpretMode MODE>
+struct CanReinterpretSlice<MODE, T, T> {
/// Always `true` as TO and FROM are the same type.
static constexpr bool value = true;
};
@@ -140,16 +150,16 @@
/// CastableBase, and the pointee type of `TO` is of the same type as, or is an ancestor of the
/// pointee type of `FROM`. Vectors of non-`const` Castable pointers can be converted to a vector of
/// `const` Castable pointers.
-template <typename TO, typename FROM>
-static constexpr bool CanReinterpretSlice = detail::CanReinterpretSlice<TO, FROM>::value;
+template <ReinterpretMode MODE, typename TO, typename FROM>
+static constexpr bool CanReinterpretSlice = detail::CanReinterpretSlice<MODE, TO, FROM>::value;
/// Reinterprets `const Slice<FROM>*` as `const Slice<TO>*`
/// @param slice a pointer to the slice to reinterpret
/// @returns the reinterpreted slice
/// @see CanReinterpretSlice
-template <typename TO, typename FROM>
+template <ReinterpretMode MODE, typename TO, typename FROM>
const Slice<TO>* ReinterpretSlice(const Slice<FROM>* slice) {
- static_assert(CanReinterpretSlice<TO, FROM>);
+ static_assert(CanReinterpretSlice<MODE, TO, FROM>);
return Bitcast<const Slice<TO>*>(slice);
}
@@ -157,9 +167,9 @@
/// @param slice a pointer to the slice to reinterpret
/// @returns the reinterpreted slice
/// @see CanReinterpretSlice
-template <typename TO, typename FROM>
+template <ReinterpretMode MODE, typename TO, typename FROM>
Slice<TO>* ReinterpretSlice(Slice<FROM>* slice) {
- static_assert(CanReinterpretSlice<TO, FROM>);
+ static_assert(CanReinterpretSlice<MODE, TO, FROM>);
return Bitcast<Slice<TO>*>(slice);
}
@@ -230,15 +240,21 @@
/// Copy constructor with covariance / const conversion
/// @param other the vector to copy
/// @see CanReinterpretSlice for rules about conversion
- template <typename U, size_t N2, typename = std::enable_if_t<CanReinterpretSlice<T, U>>>
+ template <typename U,
+ size_t N2,
+ ReinterpretMode MODE,
+ typename = std::enable_if_t<CanReinterpretSlice<MODE, T, U>>>
Vector(const Vector<U, N2>& other) { // NOLINT(runtime/explicit)
- Copy(*ReinterpretSlice<T>(&other.impl_.slice));
+ Copy(*ReinterpretSlice<MODE, T>(&other.impl_.slice));
}
/// Move constructor with covariance / const conversion
/// @param other the vector to move
/// @see CanReinterpretSlice for rules about conversion
- template <typename U, size_t N2, typename = std::enable_if_t<CanReinterpretSlice<T, U>>>
+ template <typename U,
+ size_t N2,
+ ReinterpretMode MODE,
+ typename = std::enable_if_t<CanReinterpretSlice<MODE, T, U>>>
Vector(Vector<U, N2>&& other) { // NOLINT(runtime/explicit)
MoveOrCopy(VectorRef<T>(std::move(other)));
}
@@ -380,10 +396,12 @@
/// Clears all elements from the vector, keeping the capacity the same.
void Clear() {
+ TINT_BEGIN_DISABLE_WARNING(MAYBE_UNINITIALIZED);
for (size_t i = 0; i < impl_.slice.len; i++) {
impl_.slice.data[i].~T();
}
impl_.slice.len = 0;
+ TINT_END_DISABLE_WARNING(MAYBE_UNINITIALIZED);
}
/// Appends a new element to the vector.
@@ -665,11 +683,11 @@
Vector(Ts...) -> Vector<VectorCommonType<Ts...>, sizeof...(Ts)>;
/// VectorRef is a weak reference to a Vector, used to pass vectors as parameters, avoiding copies
-/// between the caller and the callee. VectorRef can accept a Vector of any 'N' value, decoupling
-/// the caller's vector internal size from the callee's vector size. A VectorRef tracks the usage of
-/// moves either side of the call. If at the call site, a Vector argument is moved to a VectorRef
-/// parameter, and within the callee, the VectorRef parameter is moved to a Vector, then the Vector
-/// heap allocation will be moved. For example:
+/// between the caller and the callee, or as an non-static sized accessor on a vector. VectorRef can
+/// accept a Vector of any 'N' value, decoupling the caller's vector internal size from the callee's
+/// vector size. A VectorRef tracks the usage of moves either side of the call. If at the call site,
+/// a Vector argument is moved to a VectorRef parameter, and within the callee, the VectorRef
+/// parameter is moved to a Vector, then the Vector heap allocation will be moved. For example:
///
/// ```
/// void func_a() {
@@ -683,6 +701,8 @@
/// Vector<std::string, 2> vec(std::move(vec_ref));
/// }
/// ```
+///
+/// Aside from this move pattern, a VectorRef provides an immutable reference to the Vector.
template <typename T>
class VectorRef {
/// The slice type used by this vector reference
@@ -695,12 +715,20 @@
}
public:
+ /// Type of `T`.
+ using value_type = T;
+
/// Constructor - empty reference
VectorRef() : slice_(EmptySlice()) {}
/// Constructor
VectorRef(EmptyType) : slice_(EmptySlice()) {} // NOLINT(runtime/explicit)
+ /// Constructor from a Slice
+ /// @param slice the slice
+ VectorRef(Slice& slice) // NOLINT(runtime/explicit)
+ : slice_(slice) {}
+
/// Constructor from a Vector
/// @param vector the vector to create a reference of
template <size_t N>
@@ -729,29 +757,37 @@
/// Copy constructor with covariance / const conversion
/// @param other the other vector reference
- template <typename U, typename = std::enable_if_t<CanReinterpretSlice<T, U>>>
+ template <typename U,
+ typename = std::enable_if_t<CanReinterpretSlice<ReinterpretMode::kSafe, T, U>>>
VectorRef(const VectorRef<U>& other) // NOLINT(runtime/explicit)
- : slice_(*ReinterpretSlice<T>(&other.slice_)) {}
+ : slice_(*ReinterpretSlice<ReinterpretMode::kSafe, T>(&other.slice_)) {}
/// Move constructor with covariance / const conversion
/// @param other the vector reference
- template <typename U, typename = std::enable_if_t<CanReinterpretSlice<T, U>>>
+ template <typename U,
+ typename = std::enable_if_t<CanReinterpretSlice<ReinterpretMode::kSafe, T, U>>>
VectorRef(VectorRef<U>&& other) // NOLINT(runtime/explicit)
- : slice_(*ReinterpretSlice<T>(&other.slice_)), can_move_(other.can_move_) {}
+ : slice_(*ReinterpretSlice<ReinterpretMode::kSafe, T>(&other.slice_)),
+ can_move_(other.can_move_) {}
/// Constructor from a Vector with covariance / const conversion
/// @param vector the vector to create a reference of
/// @see CanReinterpretSlice for rules about conversion
- template <typename U, size_t N, typename = std::enable_if_t<CanReinterpretSlice<T, U>>>
+ template <typename U,
+ size_t N,
+ typename = std::enable_if_t<CanReinterpretSlice<ReinterpretMode::kSafe, T, U>>>
VectorRef(Vector<U, N>& vector) // NOLINT(runtime/explicit)
- : slice_(*ReinterpretSlice<T>(&vector.impl_.slice)) {}
+ : slice_(*ReinterpretSlice<ReinterpretMode::kSafe, T>(&vector.impl_.slice)) {}
/// Constructor from a moved Vector with covariance / const conversion
/// @param vector the vector to create a reference of
/// @see CanReinterpretSlice for rules about conversion
- template <typename U, size_t N, typename = std::enable_if_t<CanReinterpretSlice<T, U>>>
+ template <typename U,
+ size_t N,
+ typename = std::enable_if_t<CanReinterpretSlice<ReinterpretMode::kSafe, T, U>>>
VectorRef(Vector<U, N>&& vector) // NOLINT(runtime/explicit)
- : slice_(*ReinterpretSlice<T>(&vector.impl_.slice)), can_move_(vector.impl_.CanMove()) {}
+ : slice_(*ReinterpretSlice<ReinterpretMode::kSafe, T>(&vector.impl_.slice)),
+ can_move_(vector.impl_.CanMove()) {}
/// Index operator
/// @param i the element index. Must be less than `len`.
@@ -765,43 +801,33 @@
/// be made
size_t Capacity() const { return slice_.cap; }
+ /// @return a reinterpretation of this VectorRef as elements of type U.
+ /// @note this is doing a reinterpret_cast of elements. It is up to the caller to ensure that
+ /// this is a safe operation.
+ template <typename U>
+ VectorRef<U> ReinterpretCast() const {
+ return {*ReinterpretSlice<ReinterpretMode::kUnsafe, U>(&slice_)};
+ }
+
/// @returns true if the vector is empty.
bool IsEmpty() const { return slice_.len == 0; }
/// @returns a reference to the first element in the vector
- T& Front() { return slice_.Front(); }
-
- /// @returns a reference to the first element in the vector
const T& Front() const { return slice_.Front(); }
/// @returns a reference to the last element in the vector
- T& Back() { return slice_.Back(); }
-
- /// @returns a reference to the last element in the vector
const T& Back() const { return slice_.Back(); }
/// @returns a pointer to the first element in the vector
- T* begin() { return slice_.begin(); }
-
- /// @returns a pointer to the first element in the vector
const T* begin() const { return slice_.begin(); }
/// @returns a pointer to one past the last element in the vector
- T* end() { return slice_.end(); }
-
- /// @returns a pointer to one past the last element in the vector
const T* end() const { return slice_.end(); }
/// @returns a reverse iterator starting with the last element in the vector
- auto rbegin() { return slice_.rbegin(); }
-
- /// @returns a reverse iterator starting with the last element in the vector
auto rbegin() const { return slice_.rbegin(); }
/// @returns the end for a reverse iterator
- auto rend() { return slice_.rend(); }
-
- /// @returns the end for a reverse iterator
auto rend() const { return slice_.rend(); }
private:
@@ -871,7 +897,7 @@
/// @param vec the vector reference
/// @return the std::ostream so calls can be chained
template <typename T>
-inline std::ostream& operator<<(std::ostream& o, const utils::VectorRef<T>& vec) {
+inline std::ostream& operator<<(std::ostream& o, utils::VectorRef<T> vec) {
o << "[";
bool first = true;
for (auto& el : vec) {
diff --git a/src/tint/utils/vector_test.cc b/src/tint/utils/vector_test.cc
index 8e6fa7a..ed9a97a 100644
--- a/src/tint/utils/vector_test.cc
+++ b/src/tint/utils/vector_test.cc
@@ -79,22 +79,30 @@
static_assert(std::is_same_v<VectorCommonType<C2a*, const C2b*>, const C1*>);
static_assert(std::is_same_v<VectorCommonType<const C2a*, const C2b*>, const C1*>);
-static_assert(CanReinterpretSlice<const C0*, C0*>, "apply const");
-static_assert(!CanReinterpretSlice<C0*, const C0*>, "remove const");
-static_assert(CanReinterpretSlice<C0*, C1*>, "up cast");
-static_assert(CanReinterpretSlice<const C0*, const C1*>, "up cast");
-static_assert(CanReinterpretSlice<const C0*, C1*>, "up cast, apply const");
-static_assert(!CanReinterpretSlice<C0*, const C1*>, "up cast, remove const");
-static_assert(!CanReinterpretSlice<C1*, C0*>, "down cast");
-static_assert(!CanReinterpretSlice<const C1*, const C0*>, "down cast");
-static_assert(!CanReinterpretSlice<const C1*, C0*>, "down cast, apply const");
-static_assert(!CanReinterpretSlice<C1*, const C0*>, "down cast, remove const");
-static_assert(!CanReinterpretSlice<const C1*, C0*>, "down cast, apply const");
-static_assert(!CanReinterpretSlice<C1*, const C0*>, "down cast, remove const");
-static_assert(!CanReinterpretSlice<C2a*, C2b*>, "sideways cast");
-static_assert(!CanReinterpretSlice<const C2a*, const C2b*>, "sideways cast");
-static_assert(!CanReinterpretSlice<const C2a*, C2b*>, "sideways cast, apply const");
-static_assert(!CanReinterpretSlice<C2a*, const C2b*>, "sideways cast, remove const");
+static_assert(CanReinterpretSlice<ReinterpretMode::kSafe, const C0*, C0*>, "apply const");
+static_assert(!CanReinterpretSlice<ReinterpretMode::kSafe, C0*, const C0*>, "remove const");
+static_assert(CanReinterpretSlice<ReinterpretMode::kSafe, C0*, C1*>, "up cast");
+static_assert(CanReinterpretSlice<ReinterpretMode::kSafe, const C0*, const C1*>, "up cast");
+static_assert(CanReinterpretSlice<ReinterpretMode::kSafe, const C0*, C1*>, "up cast, apply const");
+static_assert(!CanReinterpretSlice<ReinterpretMode::kSafe, C0*, const C1*>,
+ "up cast, remove const");
+static_assert(!CanReinterpretSlice<ReinterpretMode::kSafe, C1*, C0*>, "down cast");
+static_assert(!CanReinterpretSlice<ReinterpretMode::kSafe, const C1*, const C0*>, "down cast");
+static_assert(!CanReinterpretSlice<ReinterpretMode::kSafe, const C1*, C0*>,
+ "down cast, apply const");
+static_assert(!CanReinterpretSlice<ReinterpretMode::kSafe, C1*, const C0*>,
+ "down cast, remove const");
+static_assert(!CanReinterpretSlice<ReinterpretMode::kSafe, const C1*, C0*>,
+ "down cast, apply const");
+static_assert(!CanReinterpretSlice<ReinterpretMode::kSafe, C1*, const C0*>,
+ "down cast, remove const");
+static_assert(!CanReinterpretSlice<ReinterpretMode::kSafe, C2a*, C2b*>, "sideways cast");
+static_assert(!CanReinterpretSlice<ReinterpretMode::kSafe, const C2a*, const C2b*>,
+ "sideways cast");
+static_assert(!CanReinterpretSlice<ReinterpretMode::kSafe, const C2a*, C2b*>,
+ "sideways cast, apply const");
+static_assert(!CanReinterpretSlice<ReinterpretMode::kSafe, C2a*, const C2b*>,
+ "sideways cast, remove const");
////////////////////////////////////////////////////////////////////////////////
// TintVectorTest
@@ -2001,6 +2009,18 @@
EXPECT_TRUE(AllExternallyHeld(vec_b)); // Moved, not copied
}
+TEST(TintVectorRefTest, MoveVector_ReinterpretCast) {
+ C2a c2a;
+ C2b c2b;
+ Vector<C0*, 1> vec_a{&c2a, &c2b};
+ VectorRef<const C0*> vec_ref(std::move(vec_a)); // Move
+ EXPECT_EQ(vec_ref[0], &c2a);
+ EXPECT_EQ(vec_ref[1], &c2b);
+ VectorRef<const C1*> reinterpret = vec_ref.ReinterpretCast<const C1*>();
+ EXPECT_EQ(reinterpret[0], &c2a);
+ EXPECT_EQ(reinterpret[1], &c2b);
+}
+
TEST(TintVectorRefTest, Index) {
Vector<std::string, 2> vec{"one", "two"};
VectorRef<std::string> vec_ref(vec);
@@ -2053,15 +2073,6 @@
TEST(TintVectorRefTest, FrontBack) {
Vector<std::string, 3> vec{"front", "mid", "back"};
- VectorRef<std::string> vec_ref(vec);
- static_assert(!std::is_const_v<std::remove_reference_t<decltype(vec_ref.Front())>>);
- static_assert(!std::is_const_v<std::remove_reference_t<decltype(vec_ref.Back())>>);
- EXPECT_EQ(vec_ref.Front(), "front");
- EXPECT_EQ(vec_ref.Back(), "back");
-}
-
-TEST(TintVectorRefTest, ConstFrontBack) {
- Vector<std::string, 3> vec{"front", "mid", "back"};
const VectorRef<std::string> vec_ref(vec);
static_assert(std::is_const_v<std::remove_reference_t<decltype(vec_ref.Front())>>);
static_assert(std::is_const_v<std::remove_reference_t<decltype(vec_ref.Back())>>);
@@ -2071,15 +2082,6 @@
TEST(TintVectorRefTest, BeginEnd) {
Vector<std::string, 3> vec{"front", "mid", "back"};
- VectorRef<std::string> vec_ref(vec);
- static_assert(!std::is_const_v<std::remove_reference_t<decltype(*vec_ref.begin())>>);
- static_assert(!std::is_const_v<std::remove_reference_t<decltype(*vec_ref.end())>>);
- EXPECT_EQ(vec_ref.begin(), &vec[0]);
- EXPECT_EQ(vec_ref.end(), &vec[0] + 3);
-}
-
-TEST(TintVectorRefTest, ConstBeginEnd) {
- Vector<std::string, 3> vec{"front", "mid", "back"};
const VectorRef<std::string> vec_ref(vec);
static_assert(std::is_const_v<std::remove_reference_t<decltype(*vec_ref.begin())>>);
static_assert(std::is_const_v<std::remove_reference_t<decltype(*vec_ref.end())>>);
diff --git a/src/tint/writer/glsl/generator_impl.cc b/src/tint/writer/glsl/generator_impl.cc
index c9ca429..4f86d35 100644
--- a/src/tint/writer/glsl/generator_impl.cc
+++ b/src/tint/writer/glsl/generator_impl.cc
@@ -59,6 +59,7 @@
#include "src/tint/transform/expand_compound_assignment.h"
#include "src/tint/transform/manager.h"
#include "src/tint/transform/pad_structs.h"
+#include "src/tint/transform/preserve_padding.h"
#include "src/tint/transform/promote_initializers_to_let.h"
#include "src/tint/transform/promote_side_effects_to_decl.h"
#include "src/tint/transform/remove_phonies.h"
@@ -210,6 +211,9 @@
manager.Add<transform::Renamer>();
data.Add<transform::Renamer::Config>(transform::Renamer::Target::kGlslKeywords,
/* preserve_unicode */ false);
+
+ manager.Add<transform::PreservePadding>(); // Must come before DirectVariableAccess
+
manager.Add<transform::Unshadow>(); // Must come before DirectVariableAccess
manager.Add<transform::DirectVariableAccess>();
@@ -291,7 +295,7 @@
} else if (auto* str = decl->As<ast::Struct>()) {
auto* sem = builder_.Sem().Get(str);
bool has_rt_arr = false;
- if (auto* arr = sem->Members().back()->Type()->As<sem::Array>()) {
+ if (auto* arr = sem->Members().Back()->Type()->As<sem::Array>()) {
has_rt_arr = arr->Count()->Is<sem::RuntimeArrayCount>();
}
bool is_block =
@@ -2381,7 +2385,7 @@
ScopedParen sp(out);
- for (size_t i = 0; i < s->Members().size(); i++) {
+ for (size_t i = 0; i < s->Members().Length(); i++) {
if (i > 0) {
out << ", ";
}
diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc
index 0c3856a..944481c 100644
--- a/src/tint/writer/hlsl/generator_impl.cc
+++ b/src/tint/writer/hlsl/generator_impl.cc
@@ -3380,7 +3380,7 @@
auto emit_member_values = [&](std::ostream& o) {
o << "{";
- for (size_t i = 0; i < s->Members().size(); i++) {
+ for (size_t i = 0; i < s->Members().Length(); i++) {
if (i > 0) {
o << ", ";
}
diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc
index c7354f0..1056026 100644
--- a/src/tint/writer/msl/generator_impl.cc
+++ b/src/tint/writer/msl/generator_impl.cc
@@ -67,6 +67,7 @@
#include "src/tint/transform/manager.h"
#include "src/tint/transform/module_scope_var_to_entry_point_param.h"
#include "src/tint/transform/packed_vec3.h"
+#include "src/tint/transform/preserve_padding.h"
#include "src/tint/transform/promote_initializers_to_let.h"
#include "src/tint/transform/promote_side_effects_to_decl.h"
#include "src/tint/transform/remove_phonies.h"
@@ -219,6 +220,8 @@
}
manager.Add<transform::MultiplanarExternalTexture>();
+ manager.Add<transform::PreservePadding>();
+
manager.Add<transform::Unshadow>();
if (!options.disable_workgroup_init) {
@@ -1758,8 +1761,8 @@
return true;
}
- auto& members = s->Members();
- for (size_t i = 0; i < members.size(); i++) {
+ auto members = s->Members();
+ for (size_t i = 0; i < members.Length(); i++) {
if (i > 0) {
out << ", ";
}
diff --git a/src/tint/writer/spirv/builder.cc b/src/tint/writer/spirv/builder.cc
index f2b8f85..b5910ff 100644
--- a/src/tint/writer/spirv/builder.cc
+++ b/src/tint/writer/spirv/builder.cc
@@ -1706,7 +1706,7 @@
}
return composite(count.value());
},
- [&](const sem::Struct* s) { return composite(s->Members().size()); },
+ [&](const sem::Struct* s) { return composite(s->Members().Length()); },
[&](Default) {
error_ = "unhandled constant type: " + builder_.FriendlyName(ty);
return 0;
@@ -3922,7 +3922,7 @@
push_annot(spv::Op::OpDecorate, {Operand(struct_id), U32Operand(SpvDecorationBlock)});
}
- for (uint32_t i = 0; i < struct_type->Members().size(); ++i) {
+ for (uint32_t i = 0; i < struct_type->Members().Length(); ++i) {
auto mem_id = GenerateStructMember(struct_id, i, struct_type->Members()[i]);
if (mem_id == 0) {
return false;
diff --git a/src/tint/writer/spirv/builder_type_test.cc b/src/tint/writer/spirv/builder_type_test.cc
index 4377c42..4ab8aa0 100644
--- a/src/tint/writer/spirv/builder_type_test.cc
+++ b/src/tint/writer/spirv/builder_type_test.cc
@@ -442,8 +442,8 @@
auto* arr_arr_mat2x3_f32 =
ty.array(ty.array(ty.mat2x3<f32>(), 1_u), 2_u); // Doubly nested array
auto* arr_arr_mat2x3_f16 =
- ty.array(ty.array(ty.mat2x3<f16>(), 1_u), 2_u); // Doubly nested array
- auto* rtarr_mat4x4 = ty.array(ty.mat4x4<f32>()); // Runtime array
+ ty.array(ty.array(ty.mat2x3<f16>(), 1_u), 2_u); // Doubly nested array
+ auto* rtarr_mat4x4 = ty.array(ty.mat4x4<f32>()); // Runtime array
auto* s = Structure(
"S", utils::Vector{
diff --git a/src/tint/writer/spirv/generator_impl.cc b/src/tint/writer/spirv/generator_impl.cc
index b6bb0ff..7bd37f2 100644
--- a/src/tint/writer/spirv/generator_impl.cc
+++ b/src/tint/writer/spirv/generator_impl.cc
@@ -28,6 +28,7 @@
#include "src/tint/transform/for_loop_to_loop.h"
#include "src/tint/transform/manager.h"
#include "src/tint/transform/merge_return.h"
+#include "src/tint/transform/preserve_padding.h"
#include "src/tint/transform/promote_side_effects_to_decl.h"
#include "src/tint/transform/remove_phonies.h"
#include "src/tint/transform/remove_unreachable_statements.h"
@@ -78,6 +79,8 @@
}
manager.Add<transform::MultiplanarExternalTexture>();
+ manager.Add<transform::PreservePadding>(); // Must come before DirectVariableAccess
+
manager.Add<transform::Unshadow>(); // Must come before DirectVariableAccess
bool disable_workgroup_init_in_sanitizer =
options.disable_workgroup_init || options.use_zero_initialize_workgroup_memory_extension;
