Import Tint changes from Dawn
Changes:
- 11e2571aedb28e66100c0528c4d86503f057f4f4 tint: Minor improvements by Antonio Maiorano <amaiorano@google.com>
- 8b4ed5093ed872cc07a874d00893da2f6bad2927 tint: Fix / remove tests that access OOB by Ben Clayton <bclayton@google.com>
- 72340d00ce890b4fe6205aa0dd3f013705326d8e Add location into `sem::Parameter`. by dan sinclair <dsinclair@chromium.org>
- 766a458f53116baf45263d64324ac6a991cf2073 Store location value into `sem::Function`. by dan sinclair <dsinclair@chromium.org>
- c8278e5cbda0f3697379ebf46376aa83948b85c5 tint/transform: Fix hasher return type by Ben Clayton <bclayton@google.com>
- 97c9e0ba9f975054cede16f0709da9182a4fc2a0 Store struct member location into the sem. by dan sinclair <dsinclair@chromium.org>
GitOrigin-RevId: 11e2571aedb28e66100c0528c4d86503f057f4f4
Change-Id: I2bbca8c226e8817015d7af41c6c37035a56c4702
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/101501
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/src/tint/number.h b/src/tint/number.h
index 5321158..29ae227 100644
--- a/src/tint/number.h
+++ b/src/tint/number.h
@@ -78,16 +78,18 @@
std::is_floating_point_v<U> || std::is_same_v<T, detail::NumberKindF16>;
/// Evaluates to true iff T or Number<T> is an integral type.
-template <typename T, typename U = std::conditional_t<IsNumber<T>, UnwrapNumber<T>, T>>
-constexpr bool IsIntegral = std::is_integral_v<U>;
+template <typename T>
+constexpr bool IsIntegral = std::is_integral_v<UnwrapNumber<T>>;
/// Evaluates to true iff T or Number<T> is a signed integer type.
-template <typename T, typename U = std::conditional_t<IsNumber<T>, UnwrapNumber<T>, T>>
-constexpr bool IsSignedIntegral = std::is_integral_v<U> && std::is_signed_v<U>;
+template <typename T>
+constexpr bool IsSignedIntegral =
+ std::is_integral_v<UnwrapNumber<T>> && std::is_signed_v<UnwrapNumber<T>>;
/// Evaluates to true iff T or Number<T> is an unsigned integer type.
-template <typename T, typename U = std::conditional_t<IsNumber<T>, UnwrapNumber<T>, T>>
-constexpr bool IsUnsignedIntegral = std::is_integral_v<U> && std::is_unsigned_v<U>;
+template <typename T>
+constexpr bool IsUnsignedIntegral =
+ std::is_integral_v<UnwrapNumber<T>> && std::is_unsigned_v<UnwrapNumber<T>>;
/// Evaluates to true iff T is an integer type, floating-point type or is NumberKindF16.
template <typename T>
diff --git a/src/tint/resolver/array_accessor_test.cc b/src/tint/resolver/array_accessor_test.cc
index a07f6df..f63bb1f 100644
--- a/src/tint/resolver/array_accessor_test.cc
+++ b/src/tint/resolver/array_accessor_test.cc
@@ -102,7 +102,7 @@
TEST_F(ResolverIndexAccessorTest, Matrix) {
GlobalVar("my_var", ty.mat2x3<f32>(), ast::StorageClass::kPrivate);
- auto* acc = IndexAccessor("my_var", 2_i);
+ auto* acc = IndexAccessor("my_var", 1_i);
WrapInFunction(acc);
EXPECT_TRUE(r()->Resolve()) << r()->error();
@@ -123,7 +123,7 @@
TEST_F(ResolverIndexAccessorTest, Matrix_BothDimensions) {
GlobalVar("my_var", ty.mat2x3<f32>(), ast::StorageClass::kPrivate);
- auto* acc = IndexAccessor(IndexAccessor("my_var", 2_i), 1_i);
+ auto* acc = IndexAccessor(IndexAccessor("my_var", 0_i), 1_i);
WrapInFunction(acc);
EXPECT_TRUE(r()->Resolve()) << r()->error();
diff --git a/src/tint/resolver/const_eval_test.cc b/src/tint/resolver/const_eval_test.cc
index d21d60a..1f745a3 100644
--- a/src/tint/resolver/const_eval_test.cc
+++ b/src/tint/resolver/const_eval_test.cc
@@ -99,6 +99,7 @@
return std::move(v1);
}
+// Concats vectors `vs` into `v1`
template <typename Vec, typename... Vecs>
void ConcatInto(Vec& v1, Vecs&&... vs) {
auto total_size = v1.size() + (vs.size() + ...);
@@ -106,6 +107,7 @@
(std::move(vs.begin(), vs.end(), std::back_inserter(v1)), ...);
}
+// Concats vectors `vs` into `v1` iff `condition` is true
template <bool condition, typename Vec, typename... Vecs>
void ConcatIntoIf([[maybe_unused]] Vec& v1, [[maybe_unused]] Vecs&&... vs) {
if constexpr (condition) {
diff --git a/src/tint/resolver/inferred_type_test.cc b/src/tint/resolver/inferred_type_test.cc
index 83ca0bd..047ac78 100644
--- a/src/tint/resolver/inferred_type_test.cc
+++ b/src/tint/resolver/inferred_type_test.cc
@@ -149,11 +149,11 @@
auto* member = Member("x", ty.i32());
auto* str = Structure("S", utils::Vector{member});
- auto* expected_type =
- create<sem::Struct>(str, str->name,
- sem::StructMemberList{create<sem::StructMember>(
- member, member->symbol, create<sem::I32>(), 0u, 0u, 0u, 4u)},
- 0u, 4u, 4u);
+ auto* expected_type = create<sem::Struct>(
+ str, str->name,
+ sem::StructMemberList{create<sem::StructMember>(member, 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 86b44a4..4db3cb4 100644
--- a/src/tint/resolver/intrinsic_table.cc
+++ b/src/tint/resolver/intrinsic_table.cc
@@ -793,7 +793,8 @@
/* index */ static_cast<uint32_t>(members.size()),
/* offset */ offset,
/* align */ align,
- /* size */ size));
+ /* size */ size,
+ /* location */ std::nullopt));
offset += size;
}
uint32_t size_without_padding = offset;
diff --git a/src/tint/resolver/materialize_test.cc b/src/tint/resolver/materialize_test.cc
index 239efd5..af22c53 100644
--- a/src/tint/resolver/materialize_test.cc
+++ b/src/tint/resolver/materialize_test.cc
@@ -954,14 +954,6 @@
Method::kBitcastF32Arg,
};
-/// Methods that support abstract-integer materialization
-/// Note: Doesn't contain kWorkgroupSize or kArrayLength as they have tighter constraints on the
-/// range of allowed integer values.
-constexpr Method kAIntMethods[] = {
- Method::kSwitch,
- Method::kIndex,
-};
-
/// Methods that support vector materialization
constexpr Method kVectorMethods[] = {
Method::kLet, Method::kVar, Method::kBuiltinArg, Method::kBitcastVec3F32Arg,
@@ -1039,11 +1031,36 @@
Types<f32M, AFloatM>(AFloat(-kSubnormalF32), -kSubnormalF32), //
})));
+INSTANTIATE_TEST_SUITE_P(MaterializeAInt,
+ MaterializeAbstractNumericToDefaultType,
+ testing::Combine(testing::Values(Expectation::kMaterialize),
+ testing::Values(Method::kWorkgroupSize,
+ Method::kArrayLength),
+ testing::ValuesIn(std::vector<Data>{
+ Types<i32, AInt>(1_a, 1.0), //
+ Types<i32, AInt>(10_a, 10.0), //
+ Types<i32, AInt>(100_a, 100.0), //
+ Types<i32, AInt>(1000_a, 1000.0), //
+ Types<i32, AInt>(10000_a, 10000.0), //
+ Types<i32, AInt>(65535_a, 65535.0), //
+ })));
+
+INSTANTIATE_TEST_SUITE_P(MaterializeAIntIndex,
+ MaterializeAbstractNumericToDefaultType,
+ testing::Combine(testing::Values(Expectation::kMaterialize),
+ testing::Values(Method::kIndex),
+ testing::ValuesIn(std::vector<Data>{
+ Types<i32, AInt>(0_a, 0.0), //
+ Types<i32, AInt>(1_a, 1.0), //
+ Types<i32, AInt>(2_a, 2.0), //
+ Types<i32, AInt>(3_a, 3.0), //
+ })));
+
INSTANTIATE_TEST_SUITE_P(
- MaterializeAInt,
+ MaterializeAIntSwitch,
MaterializeAbstractNumericToDefaultType,
testing::Combine(testing::Values(Expectation::kMaterialize),
- testing::ValuesIn(kAIntMethods),
+ testing::Values(Method::kSwitch),
testing::ValuesIn(std::vector<Data>{
Types<i32, AInt>(0_a, 0.0), //
Types<i32, AInt>(10_a, 10.0), //
@@ -1133,7 +1150,10 @@
AIntValueCannotBeRepresented,
MaterializeAbstractNumericToDefaultType,
testing::Combine(testing::Values(Expectation::kValueCannotBeRepresented),
- testing::ValuesIn(kAIntMethods),
+ testing::Values(Method::kWorkgroupSize,
+ Method::kArrayLength,
+ Method::kSwitch,
+ Method::kIndex),
testing::ValuesIn(std::vector<Data>{
Types<i32, AInt>(0_a, static_cast<double>(i32::kHighestValue) + 1), //
Types<i32, AInt>(0_a, static_cast<double>(i32::kLowestValue) - 1), //
diff --git a/src/tint/resolver/resolver.cc b/src/tint/resolver/resolver.cc
index 1599334..1eb3b35 100644
--- a/src/tint/resolver/resolver.cc
+++ b/src/tint/resolver/resolver.cc
@@ -716,9 +716,14 @@
}
}
- auto* sem = builder_->create<sem::Parameter>(param, index, ty, ast::StorageClass::kNone,
- ast::Access::kUndefined,
- sem::ParameterUsage::kNone, binding_point);
+ std::optional<uint32_t> location;
+ if (auto* l = ast::GetAttribute<ast::LocationAttribute>(param->attributes)) {
+ location = l->value;
+ }
+
+ auto* sem = builder_->create<sem::Parameter>(
+ param, index, ty, ast::StorageClass::kNone, ast::Access::kUndefined,
+ sem::ParameterUsage::kNone, binding_point, location);
builder_->Sem().Add(param, sem);
return sem;
}
@@ -906,6 +911,19 @@
return_type = builder_->create<sem::Void>();
}
+ // Determine if the return type has a location
+ std::optional<uint32_t> return_location;
+ for (auto* attr : decl->return_type_attributes) {
+ Mark(attr);
+
+ if (auto* a = attr->As<ast::LocationAttribute>()) {
+ return_location = a->value;
+ }
+ }
+ if (!validator_.NoDuplicateAttributes(decl->attributes)) {
+ return nullptr;
+ }
+
if (auto* str = return_type->As<sem::Struct>()) {
if (!ApplyStorageClassUsageToType(ast::StorageClass::kNone, str, decl->source)) {
AddNote(
@@ -929,7 +947,8 @@
}
}
- auto* func = builder_->create<sem::Function>(decl, return_type, std::move(parameters));
+ auto* func =
+ builder_->create<sem::Function>(decl, return_type, return_location, std::move(parameters));
builder_->Sem().Add(decl, func);
TINT_SCOPED_ASSIGNMENT(current_function_, func);
@@ -968,13 +987,7 @@
for (auto* attr : decl->attributes) {
Mark(attr);
}
- if (!validator_.NoDuplicateAttributes(decl->attributes)) {
- return nullptr;
- }
- for (auto* attr : decl->return_type_attributes) {
- Mark(attr);
- }
if (!validator_.NoDuplicateAttributes(decl->return_type_attributes)) {
return nullptr;
}
@@ -2747,6 +2760,7 @@
bool has_offset_attr = false;
bool has_align_attr = false;
bool has_size_attr = false;
+ std::optional<uint32_t> location;
for (auto* attr : member->attributes) {
Mark(attr);
if (auto* o = attr->As<ast::StructMemberOffsetAttribute>()) {
@@ -2760,11 +2774,7 @@
align = 1;
has_offset_attr = true;
} else if (auto* a = attr->As<ast::StructMemberAlignAttribute>()) {
- const auto* expr = Expression(a->align);
- if (!expr) {
- return nullptr;
- }
- auto* materialized = Materialize(expr);
+ auto* materialized = Materialize(Expression(a->align));
if (!materialized) {
return nullptr;
}
@@ -2790,6 +2800,8 @@
}
size = s->size;
has_size_attr = true;
+ } else if (auto* l = attr->As<ast::LocationAttribute>()) {
+ location = l->value;
}
}
@@ -2811,7 +2823,7 @@
auto* sem_member = builder_->create<sem::StructMember>(
member, 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));
+ static_cast<uint32_t>(size), location);
builder_->Sem().Add(member, sem_member);
sem_members.emplace_back(sem_member);
diff --git a/src/tint/resolver/resolver_test.cc b/src/tint/resolver/resolver_test.cc
index 72d741e..523dd89 100644
--- a/src/tint/resolver/resolver_test.cc
+++ b/src/tint/resolver/resolver_test.cc
@@ -773,6 +773,39 @@
EXPECT_TRUE(func_sem->ReturnType()->Is<sem::Void>());
}
+TEST_F(ResolverTest, Function_Parameters_Locations) {
+ auto* param_a = Param("a", ty.f32(), utils::Vector{Location(3)});
+ auto* param_b = Param("b", ty.u32(), utils::Vector{Builtin(ast::BuiltinValue::kVertexIndex)});
+ auto* param_c = Param("c", ty.u32(), utils::Vector{Location(1)});
+
+ GlobalVar("my_vec", ty.vec4<f32>(), ast::StorageClass::kPrivate);
+ auto* func = Func("my_func",
+ utils::Vector{
+ param_a,
+ param_b,
+ param_c,
+ },
+ ty.vec4<f32>(),
+ utils::Vector{
+ Return("my_vec"),
+ },
+ utils::Vector{
+ Stage(ast::PipelineStage::kVertex),
+ },
+ utils::Vector{
+ Builtin(ast::BuiltinValue::kPosition),
+ });
+
+ EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+ auto* func_sem = Sem().Get(func);
+ ASSERT_NE(func_sem, nullptr);
+ EXPECT_EQ(func_sem->Parameters().Length(), 3u);
+ EXPECT_EQ(3u, func_sem->Parameters()[0]->Location());
+ EXPECT_FALSE(func_sem->Parameters()[1]->Location().has_value());
+ EXPECT_EQ(1u, func_sem->Parameters()[2]->Location());
+}
+
TEST_F(ResolverTest, Function_RegisterInputOutputVariables) {
auto* s = Structure("S", utils::Vector{Member("m", ty.u32())});
@@ -802,6 +835,45 @@
EXPECT_EQ(vars[2]->Declaration(), priv_var);
}
+TEST_F(ResolverTest, Function_ReturnType_Location) {
+ auto* func = Func("my_func", utils::Empty, ty.f32(),
+ utils::Vector{
+ Return(1_f),
+ },
+ utils::Vector{
+ Stage(ast::PipelineStage::kFragment),
+ },
+ utils::Vector{
+ Location(2),
+ });
+
+ EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+ auto* sem = Sem().Get(func);
+ ASSERT_NE(nullptr, sem);
+ EXPECT_EQ(2u, sem->ReturnLocation());
+}
+
+TEST_F(ResolverTest, Function_ReturnType_NoLocation) {
+ GlobalVar("my_vec", ty.vec4<f32>(), ast::StorageClass::kPrivate);
+ auto* func = Func("my_func", utils::Empty, ty.vec4<f32>(),
+ utils::Vector{
+ Return("my_vec"),
+ },
+ utils::Vector{
+ Stage(ast::PipelineStage::kVertex),
+ },
+ utils::Vector{
+ Builtin(ast::BuiltinValue::kPosition),
+ });
+
+ EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+ auto* sem = Sem().Get(func);
+ ASSERT_NE(nullptr, sem);
+ EXPECT_FALSE(sem->ReturnLocation());
+}
+
TEST_F(ResolverTest, Function_RegisterInputOutputVariables_SubFunction) {
auto* s = Structure("S", utils::Vector{Member("m", ty.u32())});
diff --git a/src/tint/resolver/struct_pipeline_stage_use_test.cc b/src/tint/resolver/struct_pipeline_stage_use_test.cc
index acf30f6..c8e77ea 100644
--- a/src/tint/resolver/struct_pipeline_stage_use_test.cc
+++ b/src/tint/resolver/struct_pipeline_stage_use_test.cc
@@ -174,6 +174,20 @@
UnorderedElementsAre(sem::PipelineStageUsage::kFragmentInput));
}
+TEST_F(ResolverPipelineStageUseTest, StructUsedAsShaderParamLocationSet) {
+ auto* s = Structure("S", utils::Vector{Member("a", ty.f32(), utils::Vector{Location(3)})});
+
+ Func("main", utils::Vector{Param("param", ty.Of(s))}, ty.void_(), utils::Empty,
+ utils::Vector{Stage(ast::PipelineStage::kFragment)});
+
+ ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+ auto* sem = TypeOf(s)->As<sem::Struct>();
+ ASSERT_NE(sem, nullptr);
+ ASSERT_EQ(1u, sem->Members().size());
+ EXPECT_EQ(3u, sem->Members()[0]->Location());
+}
+
TEST_F(ResolverPipelineStageUseTest, StructUsedAsShaderReturnTypeViaAlias) {
auto* s = Structure("S", utils::Vector{Member("a", ty.f32(), utils::Vector{Location(0)})});
auto* s_alias = Alias("S_alias", ty.Of(s));
@@ -190,5 +204,19 @@
UnorderedElementsAre(sem::PipelineStageUsage::kFragmentOutput));
}
+TEST_F(ResolverPipelineStageUseTest, StructUsedAsShaderReturnTypeLocationSet) {
+ auto* s = Structure("S", utils::Vector{Member("a", ty.f32(), utils::Vector{Location(3)})});
+
+ Func("main", utils::Empty, ty.Of(s), utils::Vector{Return(Construct(ty.Of(s), Expr(0_f)))},
+ utils::Vector{Stage(ast::PipelineStage::kFragment)});
+
+ ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+ auto* sem = TypeOf(s)->As<sem::Struct>();
+ ASSERT_NE(sem, nullptr);
+ ASSERT_EQ(1u, sem->Members().size());
+ EXPECT_EQ(3u, sem->Members()[0]->Location());
+}
+
} // namespace
} // namespace tint::resolver
diff --git a/src/tint/resolver/variable_test.cc b/src/tint/resolver/variable_test.cc
index 52d240c..8842141 100644
--- a/src/tint/resolver/variable_test.cc
+++ b/src/tint/resolver/variable_test.cc
@@ -466,14 +466,14 @@
// }
// @group(0) @binding(0) var<storage, read_write> s : S;
// fn f() {
- // let p = &s.inner.arr[4];
+ // let p = &s.inner.arr[3];
// }
auto* inner = Structure("Inner", utils::Vector{Member("arr", ty.array<i32, 4>())});
auto* buf = Structure("S", utils::Vector{Member("inner", ty.Of(inner))});
auto* storage = GlobalVar("s", ty.Of(buf), ast::StorageClass::kStorage, ast::Access::kReadWrite,
Binding(0_a), Group(0_a));
- auto* expr = IndexAccessor(MemberAccessor(MemberAccessor(storage, "inner"), "arr"), 4_i);
+ auto* expr = IndexAccessor(MemberAccessor(MemberAccessor(storage, "inner"), "arr"), 3_i);
auto* ptr = Let("p", AddressOf(expr));
WrapInFunction(ptr);
diff --git a/src/tint/sem/function.cc b/src/tint/sem/function.cc
index 97171fe..6562526 100644
--- a/src/tint/sem/function.cc
+++ b/src/tint/sem/function.cc
@@ -40,10 +40,12 @@
Function::Function(const ast::Function* declaration,
Type* return_type,
+ std::optional<uint32_t> return_location,
utils::VectorRef<Parameter*> parameters)
: Base(return_type, SetOwner(std::move(parameters), this), EvaluationStage::kRuntime),
declaration_(declaration),
- workgroup_size_{WorkgroupDimension{1}, WorkgroupDimension{1}, WorkgroupDimension{1}} {}
+ workgroup_size_{WorkgroupDimension{1}, WorkgroupDimension{1}, WorkgroupDimension{1}},
+ return_location_(return_location) {}
Function::~Function() = default;
diff --git a/src/tint/sem/function.h b/src/tint/sem/function.h
index d4cb1c7..3f7256a 100644
--- a/src/tint/sem/function.h
+++ b/src/tint/sem/function.h
@@ -16,6 +16,7 @@
#define SRC_TINT_SEM_FUNCTION_H_
#include <array>
+#include <optional>
#include <utility>
#include <vector>
@@ -60,9 +61,11 @@
/// Constructor
/// @param declaration the ast::Function
/// @param return_type the return type of the function
+ /// @param return_location the location value for the return, if provided
/// @param parameters the parameters to the function
Function(const ast::Function* declaration,
Type* return_type,
+ std::optional<uint32_t> return_location,
utils::VectorRef<Parameter*> parameters);
/// Destructor
@@ -254,6 +257,9 @@
/// @return the behaviors of this function
sem::Behaviors& Behaviors() { return behaviors_; }
+ /// @return the location for the return, if provided
+ std::optional<uint32_t> ReturnLocation() const { return return_location_; }
+
private:
Function(const Function&) = delete;
Function(Function&&) = delete;
@@ -274,6 +280,8 @@
std::vector<const Function*> ancestor_entry_points_;
bool has_discard_ = false;
sem::Behaviors behaviors_{sem::Behavior::kNext};
+
+ std::optional<uint32_t> return_location_;
};
} // namespace tint::sem
diff --git a/src/tint/sem/sem_struct_test.cc b/src/tint/sem/sem_struct_test.cc
index 3ea14e1..453b8f7 100644
--- a/src/tint/sem/sem_struct_test.cc
+++ b/src/tint/sem/sem_struct_test.cc
@@ -105,5 +105,23 @@
/* */ };)");
}
+TEST_F(StructTest, Location) {
+ auto* st = Structure("st", utils::Vector{
+ Member("a", ty.i32(), utils::Vector{Location(1u)}),
+ Member("b", ty.u32()),
+ });
+
+ auto p = Build();
+ ASSERT_TRUE(p.IsValid()) << p.Diagnostics().str();
+
+ auto* sem = p.Sem().Get(st);
+ ASSERT_EQ(2u, sem->Members().size());
+
+ EXPECT_TRUE(sem->Members()[0]->Location().has_value());
+ EXPECT_EQ(sem->Members()[0]->Location().value(), 1u);
+
+ EXPECT_FALSE(sem->Members()[1]->Location().has_value());
+}
+
} // namespace
} // namespace tint::sem
diff --git a/src/tint/sem/struct.cc b/src/tint/sem/struct.cc
index 0ace9db..7d457e8 100644
--- a/src/tint/sem/struct.cc
+++ b/src/tint/sem/struct.cc
@@ -161,14 +161,16 @@
uint32_t index,
uint32_t offset,
uint32_t align,
- uint32_t size)
+ uint32_t size,
+ std::optional<uint32_t> location)
: declaration_(declaration),
name_(name),
type_(type),
index_(index),
offset_(offset),
align_(align),
- size_(size) {}
+ size_(size),
+ location_(location) {}
StructMember::~StructMember() = default;
diff --git a/src/tint/sem/struct.h b/src/tint/sem/struct.h
index 62b1008..0f3213a 100644
--- a/src/tint/sem/struct.h
+++ b/src/tint/sem/struct.h
@@ -17,6 +17,7 @@
#include <stdint.h>
+#include <optional>
#include <string>
#include <unordered_set>
#include <vector>
@@ -180,13 +181,15 @@
/// @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
StructMember(const ast::StructMember* declaration,
Symbol name,
const sem::Type* type,
uint32_t index,
uint32_t offset,
uint32_t align,
- uint32_t size);
+ uint32_t size,
+ std::optional<uint32_t> location);
/// Destructor
~StructMember() override;
@@ -219,6 +222,9 @@
/// @returns byte size
uint32_t Size() const { return size_; }
+ /// @returns the location, if set
+ std::optional<uint32_t> Location() const { return location_; }
+
private:
const ast::StructMember* const declaration_;
const Symbol name_;
@@ -228,6 +234,7 @@
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/type_test.cc b/src/tint/sem/type_test.cc
index 837579e..56dcbff 100644
--- a/src/tint/sem/type_test.cc
+++ b/src/tint/sem/type_test.cc
@@ -54,7 +54,8 @@
/* index */ 0u,
/* offset */ 0u,
/* align */ 4u,
- /* size */ 4u),
+ /* size */ 4u,
+ /* location */ std::nullopt),
},
/* align*/ 4u,
/* size*/ 4u,
diff --git a/src/tint/sem/variable.cc b/src/tint/sem/variable.cc
index f233053..67e7945 100644
--- a/src/tint/sem/variable.cc
+++ b/src/tint/sem/variable.cc
@@ -73,11 +73,13 @@
ast::StorageClass storage_class,
ast::Access access,
const ParameterUsage usage /* = ParameterUsage::kNone */,
- sem::BindingPoint binding_point /* = {} */)
+ sem::BindingPoint binding_point /* = {} */,
+ std::optional<uint32_t> location /* = std::nullopt */)
: Base(declaration, type, EvaluationStage::kRuntime, storage_class, access, nullptr),
index_(index),
usage_(usage),
- binding_point_(binding_point) {}
+ binding_point_(binding_point),
+ location_(location) {}
Parameter::~Parameter() = default;
diff --git a/src/tint/sem/variable.h b/src/tint/sem/variable.h
index b511563..a0eb75e 100644
--- a/src/tint/sem/variable.h
+++ b/src/tint/sem/variable.h
@@ -15,6 +15,7 @@
#ifndef SRC_TINT_SEM_VARIABLE_H_
#define SRC_TINT_SEM_VARIABLE_H_
+#include <optional>
#include <utility>
#include <vector>
@@ -189,13 +190,15 @@
/// @param access the variable access control type
/// @param usage the semantic usage for the parameter
/// @param binding_point the optional resource binding point of the parameter
+ /// @param location the location value, if set
Parameter(const ast::Parameter* declaration,
uint32_t index,
const sem::Type* type,
ast::StorageClass storage_class,
ast::Access access,
const ParameterUsage usage = ParameterUsage::kNone,
- sem::BindingPoint binding_point = {});
+ sem::BindingPoint binding_point = {},
+ std::optional<uint32_t> location = std::nullopt);
/// Destructor
~Parameter() override;
@@ -222,12 +225,16 @@
/// @returns the resource binding point for the parameter
sem::BindingPoint BindingPoint() const { return binding_point_; }
+ /// @returns the location value for the parameter, if set
+ std::optional<uint32_t> Location() const { return location_; }
+
private:
const uint32_t index_;
const ParameterUsage usage_;
CallTarget const* owner_ = nullptr;
const sem::Node* shadows_ = nullptr;
const sem::BindingPoint binding_point_;
+ const std::optional<uint32_t> location_;
};
/// VariableUser holds the semantic information for an identifier expression
diff --git a/src/tint/transform/std140.cc b/src/tint/transform/std140.cc
index 92d3d8b..f46fdea 100644
--- a/src/tint/transform/std140.cc
+++ b/src/tint/transform/std140.cc
@@ -54,7 +54,7 @@
/// The hash function for the DynamicIndex
/// @param d the DynamicIndex to hash
/// @return the hash for the given DynamicIndex
- uint64_t operator()(const DynamicIndex& d) const { return utils::Hash(d.slot); }
+ size_t operator()(const DynamicIndex& d) const { return utils::Hash(d.slot); }
};
} // namespace tint::utils
@@ -149,9 +149,7 @@
struct Hasher {
/// @param fn the LoadFnKey to hash
/// @return the hash for the given LoadFnKey
- uint64_t operator()(const LoadFnKey& fn) const {
- return utils::Hash(fn.var, fn.indices);
- }
+ size_t operator()(const LoadFnKey& fn) const { return utils::Hash(fn.var, fn.indices); }
};
/// Equality operator
diff --git a/src/tint/transform/std140_test.cc b/src/tint/transform/std140_test.cc
index 4681cdf..57e23e6 100644
--- a/src/tint/transform/std140_test.cc
+++ b/src/tint/transform/std140_test.cc
@@ -1584,7 +1584,7 @@
let l_a_3_a_1 : Inner = a[3].a[1];
let l_a_0_a_2_m : mat4x2<f32> = a[0].a[2].m;
let l_a_1_a_3_m_0 : vec2<f32> = a[1].a[3].m[0];
- let l_a_2_a_0_m_1_2 : f32 = a[2].a[0].m[1][2];
+ let l_a_2_a_0_m_1_0 : f32 = a[2].a[0].m[1][0];
}
)";
@@ -1647,7 +1647,7 @@
let l_a_3_a_1 : Inner = conv_Inner(a[3u].a[1u]);
let l_a_0_a_2_m : mat4x2<f32> = load_a_0_a_2_m();
let l_a_1_a_3_m_0 : vec2<f32> = a[1u].a[3u].m_0;
- let l_a_2_a_0_m_1_2 : f32 = a[2u].a[0u].m_1[2u];
+ let l_a_2_a_0_m_1_0 : f32 = a[2u].a[0u].m_1[0u];
}
)";
@@ -2029,7 +2029,7 @@
@group(0) @binding(0) var<uniform> u : S;
fn f() {
- for (var i = u32(u.m[0][1]); i < u32(u.m[i][2]); i += u32(u.m[1][i])) {
+ for (var i = u32(u.m[0][0]); i < u32(u.m[i][1]); i += u32(u.m[1][i])) {
}
}
)";
@@ -2050,16 +2050,16 @@
@group(0) @binding(0) var<uniform> u : S_std140;
-fn load_u_m_p0_2(p0 : u32) -> f32 {
+fn load_u_m_p0_1(p0 : u32) -> f32 {
switch(p0) {
case 0u: {
- return u.m_0[2u];
+ return u.m_0[1u];
}
case 1u: {
- return u.m_1[2u];
+ return u.m_1[1u];
}
case 2u: {
- return u.m_2[2u];
+ return u.m_2[1u];
}
default: {
return f32();
@@ -2068,7 +2068,7 @@
}
fn f() {
- for(var i = u32(u.m_0[1u]); (i < u32(load_u_m_p0_2(u32(i)))); i += u32(u.m_1[i])) {
+ for(var i = u32(u.m_0[0u]); (i < u32(load_u_m_p0_1(u32(i)))); i += u32(u.m_1[i])) {
}
}
)";
diff --git a/src/tint/writer/spirv/builder_accessor_expression_test.cc b/src/tint/writer/spirv/builder_accessor_expression_test.cc
index ae755bb..5c06717 100644
--- a/src/tint/writer/spirv/builder_accessor_expression_test.cc
+++ b/src/tint/writer/spirv/builder_accessor_expression_test.cc
@@ -653,11 +653,11 @@
Validate(b);
}
-TEST_F(BuilderTest, Runtime_IndexAccessor_Nested_Array_f32) {
- // var pos : array<array<f32, 2>, 3u>;
+TEST_F(BuilderTest, Runtime_IndexAccessor_Array_Vec3_f32) {
+ // var pos : array<vec3<f32>, 3u>;
// var x = pos[1u][2u];
- auto* pos = Var("pos", ty.array(ty.vec2<f32>(), 3_u));
+ auto* pos = Var("pos", ty.array(ty.vec3<f32>(), 3_a));
auto* x = Var("x", IndexAccessor(IndexAccessor(pos, 1_u), 2_u));
WrapInFunction(pos, x);
@@ -668,7 +668,7 @@
EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeVoid
%1 = OpTypeFunction %2
%9 = OpTypeFloat 32
-%8 = OpTypeVector %9 2
+%8 = OpTypeVector %9 3
%10 = OpTypeInt 32 0
%11 = OpConstant %10 3
%7 = OpTypeArray %8 %11
@@ -693,11 +693,11 @@
}
TEST_F(BuilderTest, Dynamic_IndexAccessor_Nested_Array_f32) {
- // var pos : array<array<f32, 2>, 3u>;
+ // var pos : array<array<f32, 4>, 3u>;
// var one = 1u;
// var x = pos[one][2u];
- auto* pos = Var("pos", ty.array(ty.vec2<f32>(), 3_u));
+ auto* pos = Var("pos", ty.array(ty.array<f32, 4>(), 3_u));
auto* one = Var("one", Expr(2_u));
auto* x = Var("x", IndexAccessor(IndexAccessor(pos, "one"), 2_u));
WrapInFunction(pos, one, x);
@@ -709,27 +709,28 @@
EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeVoid
%1 = OpTypeFunction %2
%9 = OpTypeFloat 32
-%8 = OpTypeVector %9 2
%10 = OpTypeInt 32 0
-%11 = OpConstant %10 3
-%7 = OpTypeArray %8 %11
+%11 = OpConstant %10 4
+%8 = OpTypeArray %9 %11
+%12 = OpConstant %10 3
+%7 = OpTypeArray %8 %12
%6 = OpTypePointer Function %7
-%12 = OpConstantNull %7
-%13 = OpConstant %10 2
-%15 = OpTypePointer Function %10
-%16 = OpConstantNull %10
-%18 = OpTypePointer Function %9
-%22 = OpConstantNull %9
+%13 = OpConstantNull %7
+%14 = OpConstant %10 2
+%16 = OpTypePointer Function %10
+%17 = OpConstantNull %10
+%19 = OpTypePointer Function %9
+%23 = OpConstantNull %9
)");
- EXPECT_EQ(DumpInstructions(b.functions()[0].variables()), R"(%5 = OpVariable %6 Function %12
-%14 = OpVariable %15 Function %16
-%21 = OpVariable %18 Function %22
+ EXPECT_EQ(DumpInstructions(b.functions()[0].variables()), R"(%5 = OpVariable %6 Function %13
+%15 = OpVariable %16 Function %17
+%22 = OpVariable %19 Function %23
)");
- EXPECT_EQ(DumpInstructions(b.functions()[0].instructions()), R"(OpStore %14 %13
-%17 = OpLoad %10 %14
-%19 = OpAccessChain %18 %5 %17 %13
-%20 = OpLoad %9 %19
-OpStore %21 %20
+ EXPECT_EQ(DumpInstructions(b.functions()[0].instructions()), R"(OpStore %15 %14
+%18 = OpLoad %10 %15
+%20 = OpAccessChain %19 %5 %18 %14
+%21 = OpLoad %9 %20
+OpStore %22 %21
OpReturn
)");
