diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn index 6560b6f429..d2c69e1312 100644 --- a/src/tint/BUILD.gn +++ b/src/tint/BUILD.gn @@ -517,6 +517,8 @@ libtint_source_set("libtint_core_all_src") { "transform/multiplanar_external_texture.h", "transform/num_workgroups_from_uniform.cc", "transform/num_workgroups_from_uniform.h", + "transform/packed_vec3.cc", + "transform/packed_vec3.h", "transform/pad_structs.cc", "transform/pad_structs.h", "transform/promote_initializers_to_let.cc", @@ -1225,6 +1227,7 @@ if (tint_build_unittests) { "transform/module_scope_var_to_entry_point_param_test.cc", "transform/multiplanar_external_texture_test.cc", "transform/num_workgroups_from_uniform_test.cc", + "transform/packed_vec3_test.cc", "transform/pad_structs_test.cc", "transform/promote_initializers_to_let_test.cc", "transform/promote_side_effects_to_decl_test.cc", diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt index af4384ced9..47b44a59ad 100644 --- a/src/tint/CMakeLists.txt +++ b/src/tint/CMakeLists.txt @@ -430,6 +430,8 @@ set(TINT_LIB_SRCS transform/multiplanar_external_texture.h transform/num_workgroups_from_uniform.cc transform/num_workgroups_from_uniform.h + transform/packed_vec3.cc + transform/packed_vec3.h transform/pad_structs.cc transform/pad_structs.h transform/promote_initializers_to_let.cc @@ -1141,6 +1143,7 @@ if(TINT_BUILD_TESTS) transform/module_scope_var_to_entry_point_param_test.cc transform/multiplanar_external_texture_test.cc transform/num_workgroups_from_uniform_test.cc + transform/packed_vec3_test.cc transform/pad_structs_test.cc transform/promote_initializers_to_let_test.cc transform/promote_side_effects_to_decl_test.cc diff --git a/src/tint/transform/packed_vec3.cc b/src/tint/transform/packed_vec3.cc new file mode 100644 index 0000000000..dde5aca374 --- /dev/null +++ b/src/tint/transform/packed_vec3.cc @@ -0,0 +1,194 @@ +// 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/packed_vec3.h" + +#include +#include +#include + +#include "src/tint/program_builder.h" +#include "src/tint/sem/index_accessor_expression.h" +#include "src/tint/sem/member_accessor_expression.h" +#include "src/tint/sem/statement.h" +#include "src/tint/sem/variable.h" +#include "src/tint/utils/hashmap.h" +#include "src/tint/utils/hashset.h" + +TINT_INSTANTIATE_TYPEINFO(tint::transform::PackedVec3); +TINT_INSTANTIATE_TYPEINFO(tint::transform::PackedVec3::Attribute); + +using namespace tint::number_suffixes; // NOLINT + +namespace tint::transform { + +/// The PIMPL state for the PackedVec3 transform +struct PackedVec3::State { + /// Constructor + /// @param c the CloneContext + explicit State(CloneContext& c) : ctx(c) {} + + /// Runs the transform + void Run() { + // Packed vec3 struct members + utils::Hashset members; + + // Find all the packed vector struct members, and apply the @internal(packed_vector) + // attribute. + for (auto* decl : ctx.src->AST().GlobalDeclarations()) { + if (auto* str = sem.Get(decl)) { + if (str->IsHostShareable()) { + for (auto* member : str->Members()) { + if (auto* vec = member->Type()->As()) { + if (vec->Width() == 3) { + members.Add(member); + + // Apply the PackedVec3::Attribute to the member + auto* member_decl = member->Declaration(); + auto name = ctx.Clone(member_decl->symbol); + auto* type = ctx.Clone(member_decl->type); + utils::Vector attrs{ + b.ASTNodes().Create(b.ID(), b.AllocateNodeID()), + }; + for (auto* attr : member_decl->attributes) { + attrs.Push(ctx.Clone(attr)); + } + ctx.Replace(member_decl, b.Member(name, type, std::move(attrs))); + } + } + } + } + } + } + + // Walk the nodes, starting with the most deeply nested, finding all the AST expressions + // that load a whole packed vector (not a scalar / swizzle of the vector). + utils::Hashset refs; + for (auto* node : ctx.src->ASTNodes().Objects()) { + Switch( + sem.Get(node), // + [&](const sem::StructMemberAccess* access) { + if (members.Contains(access->Member())) { + // Access to a packed vector member. Seed the expression tracking. + refs.Add(access); + } + }, + [&](const sem::IndexAccessorExpression* access) { + // Not loading a whole packed vector. Ignore. + refs.Remove(access->Object()); + }, + [&](const sem::Swizzle* access) { + // Not loading a whole packed vector. Ignore. + refs.Remove(access->Object()); + }, + [&](const sem::VariableUser* user) { + auto* v = user->Variable(); + if (v->Declaration()->Is() && // if variable is let... + v->Type()->Is() && // and let is a pointer... + refs.Contains(v->Initializer())) { // and pointer is to a packed vector... + refs.Add(user); // then propagate tracking to pointer usage + } + }, + [&](const sem::Expression* expr) { + if (auto* unary = expr->Declaration()->As()) { + if (unary->op == ast::UnaryOp::kAddressOf || + unary->op == ast::UnaryOp::kIndirection) { + // Memory access on the packed vector. Track these. + auto* inner = sem.Get(unary->expr); + if (refs.Remove(inner)) { + refs.Add(expr); + } + } + // Note: non-memory ops (e.g. '-') are ignored, leaving any tracked + // reference at the inner expression, so we'd cast, then apply the unary op. + } + }, + [&](const sem::Statement* e) { + if (auto* assign = e->Declaration()->As()) { + // We don't want to cast packed_vectors if they're being assigned to. + refs.Remove(sem.Get(assign->lhs)); + } + }); + } + + // Wrap the load expressions with a cast to the unpacked type. + utils::Hashmap unpack_fns; + for (auto* ref : refs) { + // ref is either a packed vec3 that needs casting, or a pointer to a vec3 which we just + // leave alone. + if (auto* vec_ty = ref->Type()->UnwrapRef()->As()) { + auto* expr = ref->Declaration(); + ctx.Replace(expr, [this, vec_ty, expr] { // + auto* packed = ctx.CloneWithoutTransform(expr); + return b.Construct(CreateASTTypeFor(ctx, vec_ty), packed); + }); + } + } + + ctx.Clone(); + } + + /// @returns true if this transform should be run for the given program + /// @param program the program to inspect + static bool ShouldRun(const Program* program) { + for (auto* decl : program->AST().GlobalDeclarations()) { + if (auto* str = program->Sem().Get(decl)) { + if (str->IsHostShareable()) { + for (auto* member : str->Members()) { + if (auto* vec = member->Type()->As()) { + if (vec->Width() == 3) { + return true; + } + } + } + } + } + } + return false; + } + + private: + /// 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(); + /// Alias to the ctx.dst program builder + ProgramBuilder& b = *ctx.dst; +}; + +PackedVec3::Attribute::Attribute(ProgramID pid, ast::NodeID nid) : Base(pid, nid) {} +PackedVec3::Attribute::~Attribute() = default; + +const PackedVec3::Attribute* PackedVec3::Attribute::Clone(CloneContext* ctx) const { + return ctx->dst->ASTNodes().Create(ctx->dst->ID(), ctx->dst->AllocateNodeID()); +} + +std::string PackedVec3::Attribute::InternalName() const { + return "packed_vector"; +} + +PackedVec3::PackedVec3() = default; +PackedVec3::~PackedVec3() = default; + +bool PackedVec3::ShouldRun(const Program* program, const DataMap&) const { + return State::ShouldRun(program); +} + +void PackedVec3::Run(CloneContext& ctx, const DataMap&, DataMap&) const { + State(ctx).Run(); +} + +} // namespace tint::transform diff --git a/src/tint/transform/packed_vec3.h b/src/tint/transform/packed_vec3.h new file mode 100644 index 0000000000..9d899cbf5c --- /dev/null +++ b/src/tint/transform/packed_vec3.h @@ -0,0 +1,78 @@ +// 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_PACKED_VEC3_H_ +#define SRC_TINT_TRANSFORM_PACKED_VEC3_H_ + +#include + +#include "src/tint/ast/internal_attribute.h" +#include "src/tint/transform/transform.h" + +namespace tint::transform { + +/// A transform to be used by the MSL backend which will: +/// * Apply the `@internal('packed_vector')` attribute (PackedVec3::Attribute) to all host-sharable +/// structure members that have a vec3 type. +/// * Cast all direct (not sub-accessed) loads of these packed vectors to the 'unpacked' vec3 +/// type before usage. +/// +/// This transform papers over overload holes in the MSL standard library where an MSL +/// `packed_vector` type cannot be interchangable used as a regular `vec` type. +class PackedVec3 final : public Castable { + public: + /// Attribute is the attribute applied to padded vector structure members. + class Attribute final : public Castable { + public: + /// Constructor + /// @param pid the identifier of the program that owns this node + /// @param nid the unique node identifier + Attribute(ProgramID pid, ast::NodeID nid); + /// Destructor + ~Attribute() override; + + /// @returns "packed_vector". + std::string InternalName() const override; + + /// Performs a deep clone of this object using the CloneContext `ctx`. + /// @param ctx the clone context + /// @return the newly cloned object + const Attribute* Clone(CloneContext* ctx) const override; + }; + + /// Constructor + PackedVec3(); + /// Destructor + ~PackedVec3() override; + + /// @param program the program to inspect + /// @param data optional extra transform-specific input data + /// @returns true if this transform should be run for the given program + bool ShouldRun(const Program* program, const DataMap& data = {}) const override; + + private: + struct State; + + /// Runs the transform using the CloneContext built for transforming a + /// program. Run() is responsible for calling Clone() on the CloneContext. + /// @param ctx the CloneContext primed with the input program and + /// ProgramBuilder + /// @param inputs optional extra transform-specific input data + /// @param outputs optional extra transform-specific output data + void Run(CloneContext& ctx, const DataMap& inputs, DataMap& outputs) const override; +}; + +} // namespace tint::transform + +#endif // SRC_TINT_TRANSFORM_PACKED_VEC3_H_ diff --git a/src/tint/transform/packed_vec3_test.cc b/src/tint/transform/packed_vec3_test.cc new file mode 100644 index 0000000000..0f5c92e82a --- /dev/null +++ b/src/tint/transform/packed_vec3_test.cc @@ -0,0 +1,662 @@ +// 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/packed_vec3.h" + +#include +#include +#include + +#include "src/tint/transform/test_helper.h" +#include "src/tint/utils/string.h" + +namespace tint::transform { +namespace { + +using PackedVec3Test = TransformTest; + +TEST_F(PackedVec3Test, ShouldRun_EmptyModule) { + auto* src = R"()"; + + EXPECT_FALSE(ShouldRun(src)); +} + +TEST_F(PackedVec3Test, ShouldRun_NonHostSharableStruct) { + auto* src = R"( +struct S { + v : vec3, +} + +fn f() { + var v : S; // function address-space - not host sharable +} +)"; + + EXPECT_FALSE(ShouldRun(src)); +} + +TEST_F(PackedVec3Test, ShouldRun_Vec4Vec2) { + auto* src = R"( +struct S { + v4 : vec4, + v2 : vec2, +} + +@group(0) @binding(0) var P : S; // Host sharable +)"; + + EXPECT_FALSE(ShouldRun(src)); +} + +TEST_F(PackedVec3Test, ShouldRun_HostSharableStruct) { + auto* src = R"( +struct S { + v : vec3, +} + +@group(0) @binding(0) var P : S; // Host sharable +)"; + + EXPECT_TRUE(ShouldRun(src)); +} + +TEST_F(PackedVec3Test, UniformAddressSpace) { + auto* src = R"( +struct S { + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + let x = P.v; +} +)"; + + auto* expect = R"( +struct S { + @internal(packed_vector) + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + let x = vec3(P.v); +} +)"; + + DataMap data; + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(PackedVec3Test, StorageAddressSpace) { + auto* src = R"( +struct S { + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + let x = P.v; +} +)"; + + auto* expect = R"( +struct S { + @internal(packed_vector) + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + let x = vec3(P.v); +} +)"; + + DataMap data; + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(PackedVec3Test, ExistingMemberAttributes) { + auto* src = R"( +struct S { + @align(32) @size(64) v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + let x = P.v; +} +)"; + + auto* expect = R"( +struct S { + @internal(packed_vector) @align(32) @size(64) + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + let x = vec3(P.v); +} +)"; + + DataMap data; + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(PackedVec3Test, MultipleVectors) { + auto* src = R"( +struct S { + v2_a : vec2, + v3_a : vec3, + v4_a : vec4, + v2_b : vec2, + v3_b : vec3, + v4_b : vec4, +} + +@group(0) @binding(0) var P : S; + +fn f() { + let v2_a = P.v2_a; + let v3_a = P.v3_a; + let v4_a = P.v4_a; + let v2_b = P.v2_b; + let v3_b = P.v3_b; + let v4_b = P.v4_b; +} +)"; + + auto* expect = R"( +struct S { + v2_a : vec2, + @internal(packed_vector) + v3_a : vec3, + v4_a : vec4, + v2_b : vec2, + @internal(packed_vector) + v3_b : vec3, + v4_b : vec4, +} + +@group(0) @binding(0) var P : S; + +fn f() { + let v2_a = P.v2_a; + let v3_a = vec3(P.v3_a); + let v4_a = P.v4_a; + let v2_b = P.v2_b; + let v3_b = vec3(P.v3_b); + let v4_b = P.v4_b; +} +)"; + + DataMap data; + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(PackedVec3Test, MixedAddressSpace) { + auto* src = R"( +struct S { + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + var f : S; + let x = f.v; +} +)"; + + auto* expect = R"( +struct S { + @internal(packed_vector) + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + var f : S; + let x = vec3(f.v); +} +)"; + + DataMap data; + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(PackedVec3Test, ReadMemberAccessChain) { + auto* src = R"( +struct S { + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + let x = P.v.yz.x; +} +)"; + + auto* expect = R"( +struct S { + @internal(packed_vector) + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + let x = P.v.yz.x; +} +)"; + + DataMap data; + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(PackedVec3Test, ReadVector) { + auto* src = R"( +struct S { + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + let x = P.v; +} +)"; + + auto* expect = R"( +struct S { + @internal(packed_vector) + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + let x = vec3(P.v); +} +)"; + + DataMap data; + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(PackedVec3Test, ReadIndexAccessor) { + auto* src = R"( +struct S { + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + let x = P.v[1]; +} +)"; + + auto* expect = R"( +struct S { + @internal(packed_vector) + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + let x = P.v[1]; +} +)"; + + DataMap data; + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(PackedVec3Test, ReadViaStructPtrDirect) { + auto* src = R"( +struct S { + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + let x = (*(&(*(&P)))).v; +} +)"; + + auto* expect = R"( +struct S { + @internal(packed_vector) + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + let x = vec3((*(&(*(&(P))))).v); +} +)"; + + DataMap data; + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(PackedVec3Test, ReadViaVectorPtrDirect) { + auto* src = R"( +struct S { + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + let x = *(&(*(&(P.v)))); +} +)"; + + auto* expect = R"( +struct S { + @internal(packed_vector) + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + let x = vec3(*(&(*(&(P.v))))); +} +)"; + + DataMap data; + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(PackedVec3Test, ReadViaStructPtrViaLet) { + auto* src = R"( +struct S { + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + let p0 = &P; + let p1 = &(*(p0)); + let a = (*p1).v; + let p2 = &(*(p1)); + let b = (*p2).v; + let c = (*p2).v; +} +)"; + + auto* expect = R"( +struct S { + @internal(packed_vector) + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + let p0 = &(P); + let p1 = &(*(p0)); + let a = vec3((*(p1)).v); + let p2 = &(*(p1)); + let b = vec3((*(p2)).v); + let c = vec3((*(p2)).v); +} +)"; + + DataMap data; + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(PackedVec3Test, ReadViaVectorPtrViaLet) { + auto* src = R"( +struct S { + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + let p0 = &(P.v); + let p1 = &(*(p0)); + let a = *p1; + let p2 = &(*(p1)); + let b = *p2; + let c = *p2; +} +)"; + + auto* expect = R"( +struct S { + @internal(packed_vector) + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + let p0 = &(P.v); + let p1 = &(*(p0)); + let a = vec3(*(p1)); + let p2 = &(*(p1)); + let b = vec3(*(p2)); + let c = vec3(*(p2)); +} +)"; + + DataMap data; + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(PackedVec3Test, ReadUnaryOp) { + auto* src = R"( +struct S { + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + let x = -P.v; +} +)"; + + auto* expect = R"( +struct S { + @internal(packed_vector) + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + let x = -(vec3(P.v)); +} +)"; + + DataMap data; + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(PackedVec3Test, ReadBinaryOp) { + auto* src = R"( +struct S { + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + let x = P.v + P.v; +} +)"; + + auto* expect = R"( +struct S { + @internal(packed_vector) + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + let x = (vec3(P.v) + vec3(P.v)); +} +)"; + + DataMap data; + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(PackedVec3Test, WriteVector) { + auto* src = R"( +struct S { + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + P.v = vec3(1.23); +} +)"; + + auto* expect = R"( +struct S { + @internal(packed_vector) + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + P.v = vec3(1.23); +} +)"; + + DataMap data; + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(PackedVec3Test, WriteMemberAccess) { + auto* src = R"( +struct S { + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + P.v.y = 1.23; +} +)"; + + auto* expect = R"( +struct S { + @internal(packed_vector) + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + P.v.y = 1.23; +} +)"; + + DataMap data; + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(PackedVec3Test, WriteIndexAccessor) { + auto* src = R"( +struct S { + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + P.v[1] = 1.23; +} +)"; + + auto* expect = R"( +struct S { + @internal(packed_vector) + v : vec3, +} + +@group(0) @binding(0) var P : S; + +fn f() { + P.v[1] = 1.23; +} +)"; + + DataMap data; + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); +} + +} // namespace +} // namespace tint::transform diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc index 0bc8c39992..cf67e10443 100644 --- a/src/tint/writer/msl/generator_impl.cc +++ b/src/tint/writer/msl/generator_impl.cc @@ -66,6 +66,7 @@ #include "src/tint/transform/expand_compound_assignment.h" #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/promote_initializers_to_let.h" #include "src/tint/transform/promote_side_effects_to_decl.h" #include "src/tint/transform/remove_phonies.h" @@ -154,32 +155,6 @@ class ScopedBitCast { std::ostream& s; }; -class ScopedCast { - public: - ScopedCast(GeneratorImpl* generator, - std::ostream& stream, - const sem::Type* curr_type, - const sem::Type* target_type) - : s(stream) { - auto* target_vec_type = target_type->As(); - - // If we need to promote from scalar to vector, cast the scalar to the - // vector element type. - if (curr_type->is_scalar() && target_vec_type) { - target_type = target_vec_type->type(); - } - - // Cast - generator->EmitType(s, target_type, ""); - s << "("; - } - - ~ScopedCast() { s << ")"; } - - private: - std::ostream& s; -}; - } // namespace SanitizedResult::SanitizedResult() = default; @@ -259,6 +234,7 @@ SanitizedResult Sanitize(const Program* in, const Options& options) { // it assumes that the form of the array length argument is &var.array. manager.Add(); manager.Add(); + manager.Add(); data.Add(std::move(array_length_from_uniform_cfg)); data.Add(std::move(entry_point_io_cfg)); auto out = manager.Run(in, data); @@ -554,18 +530,8 @@ bool GeneratorImpl::EmitBinary(std::ostream& out, const ast::BinaryExpression* e ScopedParen sp(out); { ScopedBitCast lhs_uint_cast(this, out, lhs_type, unsigned_type_of(target_type)); - - // In case the type is packed, cast to our own type in order to remove the packing. - // Otherwise, this just casts to itself. - if (lhs_type->is_signed_integer_vector()) { - ScopedCast lhs_self_cast(this, out, lhs_type, lhs_type); - if (!EmitExpression(out, expr->lhs)) { - return false; - } - } else { - if (!EmitExpression(out, expr->lhs)) { - return false; - } + if (!EmitExpression(out, expr->lhs)) { + return false; } } if (!emit_op()) { @@ -573,18 +539,8 @@ bool GeneratorImpl::EmitBinary(std::ostream& out, const ast::BinaryExpression* e } { ScopedBitCast rhs_uint_cast(this, out, rhs_type, unsigned_type_of(target_type)); - - // In case the type is packed, cast to our own type in order to remove the packing. - // Otherwise, this just casts to itself. - if (rhs_type->is_signed_integer_vector()) { - ScopedCast rhs_self_cast(this, out, rhs_type, rhs_type); - if (!EmitExpression(out, expr->rhs)) { - return false; - } - } else { - if (!EmitExpression(out, expr->rhs)) { - return false; - } + if (!EmitExpression(out, expr->rhs)) { + return false; } } return true; @@ -601,18 +557,8 @@ bool GeneratorImpl::EmitBinary(std::ostream& out, const ast::BinaryExpression* e ScopedParen sp(out); { ScopedBitCast lhs_uint_cast(this, out, lhs_type, unsigned_type_of(lhs_type)); - - // In case the type is packed, cast to our own type in order to remove the packing. - // Otherwise, this just casts to itself. - if (lhs_type->is_signed_integer_vector()) { - ScopedCast lhs_self_cast(this, out, lhs_type, lhs_type); - if (!EmitExpression(out, expr->lhs)) { - return false; - } - } else { - if (!EmitExpression(out, expr->lhs)) { - return false; - } + if (!EmitExpression(out, expr->lhs)) { + return false; } } if (!emit_op()) { @@ -2780,41 +2726,6 @@ bool GeneratorImpl::EmitAddressSpace(std::ostream& out, ast::AddressSpace sc) { return false; } -bool GeneratorImpl::EmitPackedType(std::ostream& out, - const sem::Type* type, - const std::string& name) { - auto* vec = type->As(); - if (vec && vec->Width() == 3) { - out << "packed_"; - if (!EmitType(out, vec, "")) { - return false; - } - - if (vec->is_float_vector() && !matrix_packed_vector_overloads_) { - // Overload operators for matrix-vector arithmetic where the vector - // operand is packed, as these overloads to not exist in the metal - // namespace. - TextBuffer b; - TINT_DEFER(helpers_.Append(b)); - line(&b) << R"(template -inline vec operator*(matrix lhs, packed_vec rhs) { - return lhs * vec(rhs); -} - -template -inline vec operator*(packed_vec lhs, matrix rhs) { - return vec(lhs) * rhs; -} -)"; - matrix_packed_vector_overloads_ = true; - } - - return true; - } - - return EmitType(out, type, name); -} - bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) { line(b) << "struct " << StructName(str) << " {"; @@ -2861,15 +2772,16 @@ bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) { } add_byte_offset_comment(out, msl_offset); + } - if (!EmitPackedType(out, mem->Type(), mem_name)) { - return false; - } - } else { - if (!EmitType(out, mem->Type(), mem_name)) { - return false; + if (auto* decl = mem->Declaration()) { + if (ast::HasAttribute(decl->attributes)) { + out << "packed_"; } } + if (!EmitType(out, mem->Type(), mem_name)) { + return false; + } auto* ty = mem->Type(); @@ -2934,6 +2846,7 @@ bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) { [&](const ast::StructMemberOffsetAttribute*) { return true; }, [&](const ast::StructMemberAlignAttribute*) { return true; }, [&](const ast::StructMemberSizeAttribute*) { return true; }, + [&](const transform::PackedVec3::Attribute*) { return true; }, [&](Default) { TINT_ICE(Writer, diagnostics_) << "unhandled struct member attribute: " << attr->Name(); diff --git a/src/tint/writer/msl/generator_impl.h b/src/tint/writer/msl/generator_impl.h index 188bfea69e..c3e1ac0486 100644 --- a/src/tint/writer/msl/generator_impl.h +++ b/src/tint/writer/msl/generator_impl.h @@ -328,14 +328,6 @@ class GeneratorImpl : public TextGenerator { /// @param sc the address space to generate /// @returns true if the address space is emitted bool EmitAddressSpace(std::ostream& out, ast::AddressSpace sc); - /// Handles generating an MSL-packed storage type. - /// If the type does not have a packed form, the standard non-packed form is - /// emitted. - /// @param out the output of the type stream - /// @param type the type to generate - /// @param name the name of the variable, only used for array emission - /// @returns true if the type is emitted - bool EmitPackedType(std::ostream& out, const sem::Type* type, const std::string& name); /// Handles generating a struct declaration /// @param buffer the text buffer that the type declaration will be written to /// @param str the struct to generate @@ -431,8 +423,8 @@ class GeneratorImpl : public TextGenerator { /// Non-empty only if an invariant attribute has been generated. std::string invariant_define_name_; - /// True if matrix-packed_vector operator overloads have been generated. - bool matrix_packed_vector_overloads_ = false; + /// The generated name for the packed vec3 type. + std::string packed_vec3_ty_; /// Unique name of the tint_array template. /// Non-empty only if the template has been generated. diff --git a/src/tint/writer/msl/generator_impl_binary_test.cc b/src/tint/writer/msl/generator_impl_binary_test.cc index e08c420c6b..fc03507b93 100644 --- a/src/tint/writer/msl/generator_impl_binary_test.cc +++ b/src/tint/writer/msl/generator_impl_binary_test.cc @@ -128,18 +128,14 @@ TEST_P(MslBinaryTest_SignedOverflowDefinedBehaviour_Chained, Emit) { } using Op = ast::BinaryOp; constexpr BinaryData signed_overflow_defined_behaviour_chained_cases[] = { - {"as_type((as_type(as_type((as_type(a) << b))) << " - "b))", + {R"(as_type((as_type(as_type((as_type(a) << b))) << b)))", Op::kShiftLeft}, - {"((a >> b) >> b)", Op::kShiftRight}, - {"as_type((as_type(as_type((as_type(a) + " - "as_type(b)))) + as_type(b)))", + {R"(((a >> b) >> b))", Op::kShiftRight}, + {R"(as_type((as_type(as_type((as_type(a) + as_type(b)))) + as_type(b))))", Op::kAdd}, - {"as_type((as_type(as_type((as_type(a) - " - "as_type(b)))) - as_type(b)))", + {R"(as_type((as_type(as_type((as_type(a) - as_type(b)))) - as_type(b))))", Op::kSubtract}, - {"as_type((as_type(as_type((as_type(a) * " - "as_type(b)))) * as_type(b)))", + {R"(as_type((as_type(as_type((as_type(a) * as_type(b)))) * as_type(b))))", Op::kMultiply}}; INSTANTIATE_TEST_SUITE_P(MslGeneratorImplTest, MslBinaryTest_SignedOverflowDefinedBehaviour_Chained, diff --git a/src/tint/writer/msl/generator_impl_type_test.cc b/src/tint/writer/msl/generator_impl_type_test.cc index bc9767a4c9..27de589c24 100644 --- a/src/tint/writer/msl/generator_impl_type_test.cc +++ b/src/tint/writer/msl/generator_impl_type_test.cc @@ -60,7 +60,7 @@ void FormatMSLField(std::stringstream& out, // Size and alignments taken from the MSL spec: // https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf DECLARE_TYPE(float2, 8, 8); -DECLARE_TYPE(packed_float3, 12, 4); +DECLARE_TYPE(float3, 12, 4); DECLARE_TYPE(float4, 16, 16); DECLARE_TYPE(float2x2, 16, 8); DECLARE_TYPE(float2x3, 32, 16); @@ -301,7 +301,7 @@ TEST_F(MslGeneratorImplTest, EmitType_Struct_Layout_NonComposites) { FIELD(0x0100, float2, 0, c) \ FIELD(0x0108, uint, 0, d) \ FIELD(0x010c, int8_t, 4, tint_pad_2) \ - FIELD(0x0110, packed_float3, 0, e) \ + FIELD(0x0110, float3, 0, e) \ FIELD(0x011c, uint, 0, f) \ FIELD(0x0120, float4, 0, g) \ FIELD(0x0130, uint, 0, h) \ @@ -641,7 +641,7 @@ TEST_F(MslGeneratorImplTest, AttemptTintPadSymbolCollision) { /* 0x0100 */ float2 tint_pad_33; /* 0x0108 */ uint tint_pad_1; /* 0x010c */ tint_array tint_pad_12; - /* 0x0110 */ packed_float3 tint_pad_3; + /* 0x0110 */ float3 tint_pad_3; /* 0x011c */ uint tint_pad_7; /* 0x0120 */ float4 tint_pad_25; /* 0x0130 */ uint tint_pad_5; diff --git a/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.msl b/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.msl index 549e70fa1a..52e5e7d94a 100644 --- a/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.msl +++ b/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.msl @@ -2,16 +2,6 @@ using namespace metal; -template -inline vec operator*(matrix lhs, packed_vec rhs) { - return lhs * vec(rhs); -} - -template -inline vec operator*(packed_vec lhs, matrix rhs) { - return vec(lhs) * rhs; -} - template struct tint_array { const constant T& operator[](size_t i) const constant { return elements[i]; } @@ -42,11 +32,11 @@ struct S { }; void tint_symbol_inner(uint idx, const device S* const tint_symbol_1) { - int3 const a = (*(tint_symbol_1)).arr[idx].a; + int3 const a = int3((*(tint_symbol_1)).arr[idx].a); int const b = (*(tint_symbol_1)).arr[idx].b; - uint3 const c = (*(tint_symbol_1)).arr[idx].c; + uint3 const c = uint3((*(tint_symbol_1)).arr[idx].c); uint const d = (*(tint_symbol_1)).arr[idx].d; - float3 const e = (*(tint_symbol_1)).arr[idx].e; + float3 const e = float3((*(tint_symbol_1)).arr[idx].e); float const f = (*(tint_symbol_1)).arr[idx].f; float2x3 const g = (*(tint_symbol_1)).arr[idx].g; float3x2 const h = (*(tint_symbol_1)).arr[idx].h; diff --git a/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.msl b/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.msl index e2430dfba9..bb82069e0e 100644 --- a/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.msl +++ b/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.msl @@ -2,16 +2,6 @@ using namespace metal; -template -inline vec operator*(matrix lhs, packed_vec rhs) { - return lhs * vec(rhs); -} - -template -inline vec operator*(packed_vec lhs, matrix rhs) { - return vec(lhs) * rhs; -} - template struct tint_array { const constant T& operator[](size_t i) const constant { return elements[i]; } diff --git a/test/tint/buffer/storage/static_index/read.wgsl.expected.msl b/test/tint/buffer/storage/static_index/read.wgsl.expected.msl index 8f891a2e3c..db51912bc3 100644 --- a/test/tint/buffer/storage/static_index/read.wgsl.expected.msl +++ b/test/tint/buffer/storage/static_index/read.wgsl.expected.msl @@ -2,16 +2,6 @@ using namespace metal; -template -inline vec operator*(matrix lhs, packed_vec rhs) { - return lhs * vec(rhs); -} - -template -inline vec operator*(packed_vec lhs, matrix rhs) { - return vec(lhs) * rhs; -} - template struct tint_array { const constant T& operator[](size_t i) const constant { return elements[i]; } @@ -43,11 +33,11 @@ struct S { }; kernel void tint_symbol(const device S* tint_symbol_1 [[buffer(0)]]) { - int3 const a = (*(tint_symbol_1)).a; + int3 const a = int3((*(tint_symbol_1)).a); int const b = (*(tint_symbol_1)).b; - uint3 const c = (*(tint_symbol_1)).c; + uint3 const c = uint3((*(tint_symbol_1)).c); uint const d = (*(tint_symbol_1)).d; - float3 const e = (*(tint_symbol_1)).e; + float3 const e = float3((*(tint_symbol_1)).e); float const f = (*(tint_symbol_1)).f; float2x3 const g = (*(tint_symbol_1)).g; float3x2 const h = (*(tint_symbol_1)).h; diff --git a/test/tint/buffer/storage/static_index/write.wgsl.expected.msl b/test/tint/buffer/storage/static_index/write.wgsl.expected.msl index 251e06f9ea..1b21710a4a 100644 --- a/test/tint/buffer/storage/static_index/write.wgsl.expected.msl +++ b/test/tint/buffer/storage/static_index/write.wgsl.expected.msl @@ -2,16 +2,6 @@ using namespace metal; -template -inline vec operator*(matrix lhs, packed_vec rhs) { - return lhs * vec(rhs); -} - -template -inline vec operator*(packed_vec lhs, matrix rhs) { - return vec(lhs) * rhs; -} - template struct tint_array { const constant T& operator[](size_t i) const constant { return elements[i]; } diff --git a/test/tint/buffer/uniform/dynamic_index/read.wgsl.expected.msl b/test/tint/buffer/uniform/dynamic_index/read.wgsl.expected.msl index e361022682..37817c39f9 100644 --- a/test/tint/buffer/uniform/dynamic_index/read.wgsl.expected.msl +++ b/test/tint/buffer/uniform/dynamic_index/read.wgsl.expected.msl @@ -2,16 +2,6 @@ using namespace metal; -template -inline vec operator*(matrix lhs, packed_vec rhs) { - return lhs * vec(rhs); -} - -template -inline vec operator*(packed_vec lhs, matrix rhs) { - return vec(lhs) * rhs; -} - template struct tint_array { const constant T& operator[](size_t i) const constant { return elements[i]; } @@ -44,11 +34,11 @@ struct S { }; void tint_symbol_inner(uint idx, const constant S* const tint_symbol_1) { - int3 const a = (*(tint_symbol_1)).arr[idx].a; + int3 const a = int3((*(tint_symbol_1)).arr[idx].a); int const b = (*(tint_symbol_1)).arr[idx].b; - uint3 const c = (*(tint_symbol_1)).arr[idx].c; + uint3 const c = uint3((*(tint_symbol_1)).arr[idx].c); uint const d = (*(tint_symbol_1)).arr[idx].d; - float3 const e = (*(tint_symbol_1)).arr[idx].e; + float3 const e = float3((*(tint_symbol_1)).arr[idx].e); float const f = (*(tint_symbol_1)).arr[idx].f; int2 const g = (*(tint_symbol_1)).arr[idx].g; int2 const h = (*(tint_symbol_1)).arr[idx].h; diff --git a/test/tint/buffer/uniform/static_index/read.wgsl.expected.msl b/test/tint/buffer/uniform/static_index/read.wgsl.expected.msl index 26497bd6d4..02d091d2b4 100644 --- a/test/tint/buffer/uniform/static_index/read.wgsl.expected.msl +++ b/test/tint/buffer/uniform/static_index/read.wgsl.expected.msl @@ -14,16 +14,6 @@ struct tint_array { T elements[N]; }; -template -inline vec operator*(matrix lhs, packed_vec rhs) { - return lhs * vec(rhs); -} - -template -inline vec operator*(packed_vec lhs, matrix rhs) { - return vec(lhs) * rhs; -} - struct Inner { /* 0x0000 */ int x; /* 0x0004 */ tint_array tint_pad; @@ -46,11 +36,11 @@ struct S { }; kernel void tint_symbol(const constant S* tint_symbol_1 [[buffer(0)]]) { - int3 const a = (*(tint_symbol_1)).a; + int3 const a = int3((*(tint_symbol_1)).a); int const b = (*(tint_symbol_1)).b; - uint3 const c = (*(tint_symbol_1)).c; + uint3 const c = uint3((*(tint_symbol_1)).c); uint const d = (*(tint_symbol_1)).d; - float3 const e = (*(tint_symbol_1)).e; + float3 const e = float3((*(tint_symbol_1)).e); float const f = (*(tint_symbol_1)).f; int2 const g = (*(tint_symbol_1)).g; int2 const h = (*(tint_symbol_1)).h; diff --git a/test/tint/bug/chromium/1273230.wgsl.expected.msl b/test/tint/bug/chromium/1273230.wgsl.expected.msl index 08f7854102..1e627d7253 100644 --- a/test/tint/bug/chromium/1273230.wgsl.expected.msl +++ b/test/tint/bug/chromium/1273230.wgsl.expected.msl @@ -2,16 +2,6 @@ using namespace metal; -template -inline vec operator*(matrix lhs, packed_vec rhs) { - return lhs * vec(rhs); -} - -template -inline vec operator*(packed_vec lhs, matrix rhs) { - return vec(lhs) * rhs; -} - template struct tint_array { const constant T& operator[](size_t i) const constant { return elements[i]; } diff --git a/test/tint/bug/fxc/indexed_assign_to_array_in_struct/1206.wgsl.expected.msl b/test/tint/bug/fxc/indexed_assign_to_array_in_struct/1206.wgsl.expected.msl index 6507f5803f..c63e9f1ef9 100644 --- a/test/tint/bug/fxc/indexed_assign_to_array_in_struct/1206.wgsl.expected.msl +++ b/test/tint/bug/fxc/indexed_assign_to_array_in_struct/1206.wgsl.expected.msl @@ -14,16 +14,6 @@ struct tint_array { T elements[N]; }; -template -inline vec operator*(matrix lhs, packed_vec rhs) { - return lhs * vec(rhs); -} - -template -inline vec operator*(packed_vec lhs, matrix rhs) { - return vec(lhs) * rhs; -} - struct Simulation { /* 0x0000 */ uint i; }; diff --git a/test/tint/bug/tint/1113.wgsl.expected.msl b/test/tint/bug/tint/1113.wgsl.expected.msl index b62d4c5f7d..bf546c194a 100644 --- a/test/tint/bug/tint/1113.wgsl.expected.msl +++ b/test/tint/bug/tint/1113.wgsl.expected.msl @@ -2,16 +2,6 @@ using namespace metal; -template -inline vec operator*(matrix lhs, packed_vec rhs) { - return lhs * vec(rhs); -} - -template -inline vec operator*(packed_vec lhs, matrix rhs) { - return vec(lhs) * rhs; -} - template struct tint_array { const constant T& operator[](size_t i) const constant { return elements[i]; } diff --git a/test/tint/bug/tint/1118.wgsl.expected.msl b/test/tint/bug/tint/1118.wgsl.expected.msl index 40e92e8efa..5539c67f1a 100644 --- a/test/tint/bug/tint/1118.wgsl.expected.msl +++ b/test/tint/bug/tint/1118.wgsl.expected.msl @@ -1,17 +1,6 @@ #include using namespace metal; - -template -inline vec operator*(matrix lhs, packed_vec rhs) { - return lhs * vec(rhs); -} - -template -inline vec operator*(packed_vec lhs, matrix rhs) { - return vec(lhs) * rhs; -} - struct Scene { /* 0x0000 */ float4 vEyePosition; }; @@ -77,12 +66,12 @@ void main_1(thread float* const tint_symbol_7, thread bool* const tint_symbol_8, shadow = 1.0f; refractionColor = float4(0.0f, 0.0f, 0.0f, 1.0f); reflectionColor = float4(0.0f, 0.0f, 0.0f, 1.0f); - float3 const x_94 = (*(tint_symbol_11)).vEmissiveColor; + float3 const x_94 = float3((*(tint_symbol_11)).vEmissiveColor); emissiveColor = x_94; float3 const x_96 = diffuseBase; float3 const x_97 = diffuseColor; float3 const x_99 = emissiveColor; - float3 const x_103 = (*(tint_symbol_11)).vAmbientColor; + float3 const x_103 = float3((*(tint_symbol_11)).vAmbientColor); float4 const x_108 = baseColor; finalDiffuse = (clamp((((x_96 * x_97) + x_99) + x_103), float3(0.0f), float3(1.0f)) * float3(x_108[0], x_108[1], x_108[2])); finalSpecular = float3(0.0f); diff --git a/test/tint/bug/tint/1121.wgsl.expected.msl b/test/tint/bug/tint/1121.wgsl.expected.msl index 34c70b2d60..22cc6ddb86 100644 --- a/test/tint/bug/tint/1121.wgsl.expected.msl +++ b/test/tint/bug/tint/1121.wgsl.expected.msl @@ -2,16 +2,6 @@ using namespace metal; -template -inline vec operator*(matrix lhs, packed_vec rhs) { - return lhs * vec(rhs); -} - -template -inline vec operator*(packed_vec lhs, matrix rhs) { - return vec(lhs) * rhs; -} - template struct tint_array { const constant T& operator[](size_t i) const constant { return elements[i]; } @@ -88,7 +78,7 @@ void tint_symbol_inner(uint3 GlobalInvocationID, const constant Config* const ti for(int x_1 = 0; (x_1 < TILE_COUNT_X); x_1 = as_type((as_type(x_1) + as_type(1)))) { int2 tilePixel0Idx = int2(as_type((as_type(x_1) * as_type(TILE_SIZE))), as_type((as_type(y_1) * as_type(TILE_SIZE)))); float2 floorCoord = (((2.0f * float2(tilePixel0Idx)) / float4((*(tint_symbol_3)).fullScreenSize).xy) - float2(1.0f)); - float2 ceilCoord = (((2.0f * float2(as_type((as_type(int2(tilePixel0Idx)) + as_type(int2(int2(TILE_SIZE))))))) / float4((*(tint_symbol_3)).fullScreenSize).xy) - float2(1.0f)); + float2 ceilCoord = (((2.0f * float2(as_type((as_type(tilePixel0Idx) + as_type(int2(TILE_SIZE)))))) / float4((*(tint_symbol_3)).fullScreenSize).xy) - float2(1.0f)); float2 viewFloorCoord = float2((((-(viewNear) * floorCoord[0]) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * floorCoord[1]) - (M[2][1] * viewNear)) / M[1][1])); float2 viewCeilCoord = float2((((-(viewNear) * ceilCoord[0]) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * ceilCoord[1]) - (M[2][1] * viewNear)) / M[1][1])); frustumPlanes[0] = float4(1.0f, 0.0f, (-(viewFloorCoord[0]) / viewNear), 0.0f); diff --git a/test/tint/bug/tint/1520.spvasm.expected.msl b/test/tint/bug/tint/1520.spvasm.expected.msl index 3b8e4f7a59..d9ff17cfb6 100644 --- a/test/tint/bug/tint/1520.spvasm.expected.msl +++ b/test/tint/bug/tint/1520.spvasm.expected.msl @@ -45,13 +45,13 @@ bool test_int_S1_c0_b(const constant UniformBuffer* const tint_symbol_5) { ok = x_41; int4 const x_44 = int4(x_27, x_27, x_27, x_27); val = x_44; - int4 const x_47 = as_type((as_type(int4(x_44)) + as_type(int4(int4(1))))); + int4 const x_47 = as_type((as_type(x_44) + as_type(int4(1)))); val = x_47; - int4 const x_48 = as_type((as_type(int4(x_47)) - as_type(int4(int4(1))))); + int4 const x_48 = as_type((as_type(x_47) - as_type(int4(1)))); val = x_48; - int4 const x_49 = as_type((as_type(int4(x_48)) + as_type(int4(int4(1))))); + int4 const x_49 = as_type((as_type(x_48) + as_type(int4(1)))); val = x_49; - int4 const x_50 = as_type((as_type(int4(x_49)) - as_type(int4(int4(1))))); + int4 const x_50 = as_type((as_type(x_49) - as_type(int4(1)))); val = x_50; x_55 = false; if (x_41) { @@ -59,11 +59,11 @@ bool test_int_S1_c0_b(const constant UniformBuffer* const tint_symbol_5) { x_55 = x_54; } ok = x_55; - int4 const x_58 = as_type((as_type(int4(x_50)) * as_type(int4(int4(2))))); + int4 const x_58 = as_type((as_type(x_50) * as_type(int4(2)))); val = x_58; int4 const x_59 = (x_58 / int4(2)); val = x_59; - int4 const x_60 = as_type((as_type(int4(x_59)) * as_type(int4(int4(2))))); + int4 const x_60 = as_type((as_type(x_59) * as_type(int4(2)))); val = x_60; int4 const x_61 = (x_60 / int4(2)); val = x_61; diff --git a/test/tint/bug/tint/1534.wgsl b/test/tint/bug/tint/1534.wgsl new file mode 100644 index 0000000000..1b0c2a6148 --- /dev/null +++ b/test/tint/bug/tint/1534.wgsl @@ -0,0 +1,17 @@ + +struct g { + a : vec3, +} + +struct h { + a : u32, +} + +@group(0) @binding(0) var i : g; + +@group(0) @binding(1) var j : h; + +@compute @workgroup_size(1) fn main() { + let l = dot(i.a, i.a); + j.a = i.a.x; +} diff --git a/test/tint/bug/tint/1534.wgsl.expected.msl b/test/tint/bug/tint/1534.wgsl.expected.msl new file mode 100644 index 0000000000..6f8a3debda --- /dev/null +++ b/test/tint/bug/tint/1534.wgsl.expected.msl @@ -0,0 +1,35 @@ +#include + +using namespace metal; + +template +struct tint_array { + const constant T& operator[](size_t i) const constant { return elements[i]; } + device T& operator[](size_t i) device { return elements[i]; } + const device T& operator[](size_t i) const device { return elements[i]; } + thread T& operator[](size_t i) thread { return elements[i]; } + const thread T& operator[](size_t i) const thread { return elements[i]; } + threadgroup T& operator[](size_t i) threadgroup { return elements[i]; } + const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; } + T elements[N]; +}; + +template +T tint_dot3(vec a, vec b) { + return a[0]*b[0] + a[1]*b[1] + a[2]*b[2]; +} +struct g { + /* 0x0000 */ packed_uint3 a; + /* 0x000c */ tint_array tint_pad; +}; + +struct h { + /* 0x0000 */ uint a; +}; + +kernel void tint_symbol(const constant g* tint_symbol_1 [[buffer(0)]], device h* tint_symbol_2 [[buffer(1)]]) { + uint const l = tint_dot3(uint3((*(tint_symbol_1)).a), uint3((*(tint_symbol_1)).a)); + (*(tint_symbol_2)).a = (*(tint_symbol_1)).a[0]; + return; +} + diff --git a/test/tint/bug/tint/1677.wgsl.expected.msl b/test/tint/bug/tint/1677.wgsl.expected.msl index 6ee8bbda91..580eba4d82 100644 --- a/test/tint/bug/tint/1677.wgsl.expected.msl +++ b/test/tint/bug/tint/1677.wgsl.expected.msl @@ -20,7 +20,7 @@ struct Input { }; void tint_symbol_inner(uint3 id, const device Input* const tint_symbol_1) { - int3 const pos = as_type((as_type(int3((*(tint_symbol_1)).position)) - as_type(int3(int3(0))))); + int3 const pos = as_type((as_type(int3((*(tint_symbol_1)).position)) - as_type(int3(0)))); } kernel void tint_symbol(const device Input* tint_symbol_2 [[buffer(0)]], uint3 id [[thread_position_in_grid]]) { diff --git a/test/tint/bug/tint/948.wgsl.expected.msl b/test/tint/bug/tint/948.wgsl.expected.msl index 6113640630..e77abf2fae 100644 --- a/test/tint/bug/tint/948.wgsl.expected.msl +++ b/test/tint/bug/tint/948.wgsl.expected.msl @@ -14,16 +14,6 @@ struct tint_array { T elements[N]; }; -template -inline vec operator*(matrix lhs, packed_vec rhs) { - return lhs * vec(rhs); -} - -template -inline vec operator*(packed_vec lhs, matrix rhs) { - return vec(lhs) * rhs; -} - struct LeftOver { /* 0x0000 */ float time; /* 0x0004 */ uint padding; @@ -193,7 +183,7 @@ void main_1(thread float2* const tint_symbol_8, const constant LeftOver* const t i = as_type((as_type(x_304) + as_type(1))); } } - float3 const x_310 = (*(tint_symbol_9)).colorMul; + float3 const x_310 = float3((*(tint_symbol_9)).colorMul); float4 const x_311 = color; float3 const x_313 = (float3(x_311[0], x_311[1], x_311[2]) * x_310); float4 const x_314 = color; diff --git a/test/tint/bug/tint/949.wgsl.expected.msl b/test/tint/bug/tint/949.wgsl.expected.msl index c7bdad37fe..44f53207f4 100644 --- a/test/tint/bug/tint/949.wgsl.expected.msl +++ b/test/tint/bug/tint/949.wgsl.expected.msl @@ -14,16 +14,6 @@ struct tint_array { T elements[N]; }; -template -inline vec operator*(matrix lhs, packed_vec rhs) { - return lhs * vec(rhs); -} - -template -inline vec operator*(packed_vec lhs, matrix rhs) { - return vec(lhs) * rhs; -} - struct lightingInfo { float3 diffuse; float3 specular; @@ -248,7 +238,7 @@ void main_1(thread float2* const tint_symbol_7, texture2d float4 const x_264 = tempTextureRead; float const x_273 = (*(tint_symbol_10)).textureInfoName; rgb = (float3(x_264[0], x_264[1], x_264[2]) * x_273); - float3 const x_279 = (*(tint_symbol_10)).u_cameraPosition; + float3 const x_279 = float3((*(tint_symbol_10)).u_cameraPosition); float4 const x_282 = *(tint_symbol_11); output5 = normalize((x_279 - float3(x_282[0], x_282[1], x_282[2]))); output4 = float4(0.0f); @@ -380,7 +370,7 @@ void main_1(thread float2* const tint_symbol_7, texture2d tempTextureRead1 = x_475; float4 const x_477 = tempTextureRead1; rgb1 = float3(x_477[0], x_477[1], x_477[2]); - float3 const x_481 = (*(tint_symbol_10)).u_cameraPosition; + float3 const x_481 = float3((*(tint_symbol_10)).u_cameraPosition); float4 const x_482 = *(tint_symbol_11); viewDirectionW_1 = normalize((x_481 - float3(x_482[0], x_482[1], x_482[2]))); shadow = 1.0f; @@ -400,7 +390,7 @@ void main_1(thread float2* const tint_symbol_7, texture2d param_14 = float3(x_510[0], x_510[1], x_510[2]); float4 const x_514 = (*(tint_symbol_17)).vLightSpecular; param_15 = float3(x_514[0], x_514[1], x_514[2]); - float3 const x_518 = (*(tint_symbol_17)).vLightGround; + float3 const x_518 = float3((*(tint_symbol_17)).vLightGround); param_16 = x_518; float const x_520 = glossiness_1; param_17 = x_520; diff --git a/test/tint/bug/tint/980.wgsl.expected.msl b/test/tint/bug/tint/980.wgsl.expected.msl index 2e23ef125b..36cce9d3c5 100644 --- a/test/tint/bug/tint/980.wgsl.expected.msl +++ b/test/tint/bug/tint/980.wgsl.expected.msl @@ -1,17 +1,6 @@ #include using namespace metal; - -template -inline vec operator*(matrix lhs, packed_vec rhs) { - return lhs * vec(rhs); -} - -template -inline vec operator*(packed_vec lhs, matrix rhs) { - return vec(lhs) * rhs; -} - float3 Bad(uint index, float3 rd) { float3 normal = float3(0.0f); normal[index] = -(sign(rd[index])); @@ -24,7 +13,7 @@ struct S { }; void tint_symbol_inner(uint idx, device S* const tint_symbol_2) { - float3 const tint_symbol_1 = Bad((*(tint_symbol_2)).i, (*(tint_symbol_2)).v); + float3 const tint_symbol_1 = Bad((*(tint_symbol_2)).i, float3((*(tint_symbol_2)).v)); (*(tint_symbol_2)).v = tint_symbol_1; } diff --git a/test/tint/expressions/binary/add/scalar-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/add/scalar-vec3/i32.wgsl.expected.msl index 8658c03602..46f1909da0 100644 --- a/test/tint/expressions/binary/add/scalar-vec3/i32.wgsl.expected.msl +++ b/test/tint/expressions/binary/add/scalar-vec3/i32.wgsl.expected.msl @@ -4,7 +4,7 @@ using namespace metal; kernel void f() { int const a = 4; int3 const b = int3(1, 2, 3); - int3 const r = as_type((as_type(a) + as_type(int3(b)))); + int3 const r = as_type((as_type(a) + as_type(b))); return; } diff --git a/test/tint/expressions/binary/add/vec3-scalar/i32.wgsl.expected.msl b/test/tint/expressions/binary/add/vec3-scalar/i32.wgsl.expected.msl index de4195f275..89306cc895 100644 --- a/test/tint/expressions/binary/add/vec3-scalar/i32.wgsl.expected.msl +++ b/test/tint/expressions/binary/add/vec3-scalar/i32.wgsl.expected.msl @@ -4,7 +4,7 @@ using namespace metal; kernel void f() { int3 const a = int3(1, 2, 3); int const b = 4; - int3 const r = as_type((as_type(int3(a)) + as_type(b))); + int3 const r = as_type((as_type(a) + as_type(b))); return; } diff --git a/test/tint/expressions/binary/add/vec3-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/add/vec3-vec3/i32.wgsl.expected.msl index 009e87ebc9..54d88c782f 100644 --- a/test/tint/expressions/binary/add/vec3-vec3/i32.wgsl.expected.msl +++ b/test/tint/expressions/binary/add/vec3-vec3/i32.wgsl.expected.msl @@ -4,7 +4,7 @@ using namespace metal; kernel void f() { int3 const a = int3(1, 2, 3); int3 const b = int3(4, 5, 6); - int3 const r = as_type((as_type(int3(a)) + as_type(int3(b)))); + int3 const r = as_type((as_type(a) + as_type(b))); return; } diff --git a/test/tint/expressions/binary/div_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/div_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl index 6d2f8bc123..72cb90c2b1 100644 --- a/test/tint/expressions/binary/div_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl +++ b/test/tint/expressions/binary/div_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl @@ -4,7 +4,7 @@ using namespace metal; kernel void f() { int a = 4; int3 b = int3(0, 2, 0); - int3 const r = (a / as_type((as_type(int3(b)) + as_type(int3(b))))); + int3 const r = (a / as_type((as_type(b) + as_type(b)))); return; } diff --git a/test/tint/expressions/binary/div_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/div_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl index c90c00e080..e2a2a76026 100644 --- a/test/tint/expressions/binary/div_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl +++ b/test/tint/expressions/binary/div_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl @@ -4,7 +4,7 @@ using namespace metal; kernel void f() { int3 a = int3(1, 2, 3); int3 b = int3(0, 5, 0); - int3 const r = (a / as_type((as_type(int3(b)) + as_type(int3(b))))); + int3 const r = (a / as_type((as_type(b) + as_type(b)))); return; } diff --git a/test/tint/expressions/binary/left-shift/vector-vector/i32.wgsl.expected.msl b/test/tint/expressions/binary/left-shift/vector-vector/i32.wgsl.expected.msl index 1e4e146d2f..f97db25d1b 100644 --- a/test/tint/expressions/binary/left-shift/vector-vector/i32.wgsl.expected.msl +++ b/test/tint/expressions/binary/left-shift/vector-vector/i32.wgsl.expected.msl @@ -4,7 +4,7 @@ using namespace metal; kernel void f() { int3 const a = int3(1, 2, 3); uint3 const b = uint3(4u, 5u, 6u); - int3 const r = as_type((as_type(int3(a)) << b)); + int3 const r = as_type((as_type(a) << b)); return; } diff --git a/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl index 22f6994df9..12f1a72cc2 100644 --- a/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl +++ b/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl @@ -4,7 +4,7 @@ using namespace metal; kernel void f() { int a = 4; int3 b = int3(0, 2, 0); - int3 const r = (a % as_type((as_type(int3(b)) + as_type(int3(b))))); + int3 const r = (a % as_type((as_type(b) + as_type(b)))); return; } diff --git a/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl index 2d83cd38eb..61356dc074 100644 --- a/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl +++ b/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl @@ -4,7 +4,7 @@ using namespace metal; kernel void f() { int3 a = int3(1, 2, 3); int3 b = int3(0, 5, 0); - int3 const r = (a % as_type((as_type(int3(b)) + as_type(int3(b))))); + int3 const r = (a % as_type((as_type(b) + as_type(b)))); return; } diff --git a/test/tint/expressions/binary/mul/mat3x2-vec3/f32.wgsl.expected.msl b/test/tint/expressions/binary/mul/mat3x2-vec3/f32.wgsl.expected.msl index e210c3cd6e..a36de59fd7 100644 --- a/test/tint/expressions/binary/mul/mat3x2-vec3/f32.wgsl.expected.msl +++ b/test/tint/expressions/binary/mul/mat3x2-vec3/f32.wgsl.expected.msl @@ -14,16 +14,6 @@ struct tint_array { T elements[N]; }; -template -inline vec operator*(matrix lhs, packed_vec rhs) { - return lhs * vec(rhs); -} - -template -inline vec operator*(packed_vec lhs, matrix rhs) { - return vec(lhs) * rhs; -} - struct S { /* 0x0000 */ float3x2 tint_symbol; /* 0x0018 */ tint_array tint_pad; @@ -32,7 +22,7 @@ struct S { }; fragment void tint_symbol_1(const constant S* tint_symbol_2 [[buffer(0)]]) { - float2 const x = ((*(tint_symbol_2)).tint_symbol * (*(tint_symbol_2)).vector); + float2 const x = ((*(tint_symbol_2)).tint_symbol * float3((*(tint_symbol_2)).vector)); return; } diff --git a/test/tint/expressions/binary/mul/mat3x3-vec3/f32.wgsl.expected.msl b/test/tint/expressions/binary/mul/mat3x3-vec3/f32.wgsl.expected.msl index 8f4ff550d7..79c3e09dbf 100644 --- a/test/tint/expressions/binary/mul/mat3x3-vec3/f32.wgsl.expected.msl +++ b/test/tint/expressions/binary/mul/mat3x3-vec3/f32.wgsl.expected.msl @@ -2,16 +2,6 @@ using namespace metal; -template -inline vec operator*(matrix lhs, packed_vec rhs) { - return lhs * vec(rhs); -} - -template -inline vec operator*(packed_vec lhs, matrix rhs) { - return vec(lhs) * rhs; -} - template struct tint_array { const constant T& operator[](size_t i) const constant { return elements[i]; } @@ -31,7 +21,7 @@ struct S { }; fragment void tint_symbol_1(const constant S* tint_symbol_2 [[buffer(0)]]) { - float3 const x = ((*(tint_symbol_2)).tint_symbol * (*(tint_symbol_2)).vector); + float3 const x = ((*(tint_symbol_2)).tint_symbol * float3((*(tint_symbol_2)).vector)); return; } diff --git a/test/tint/expressions/binary/mul/scalar-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/mul/scalar-vec3/i32.wgsl.expected.msl index 0a492f4854..216de39caa 100644 --- a/test/tint/expressions/binary/mul/scalar-vec3/i32.wgsl.expected.msl +++ b/test/tint/expressions/binary/mul/scalar-vec3/i32.wgsl.expected.msl @@ -4,7 +4,7 @@ using namespace metal; kernel void f() { int const a = 4; int3 const b = int3(1, 2, 3); - int3 const r = as_type((as_type(a) * as_type(int3(b)))); + int3 const r = as_type((as_type(a) * as_type(b))); return; } diff --git a/test/tint/expressions/binary/mul/vec3-mat3x3/f32.wgsl.expected.msl b/test/tint/expressions/binary/mul/vec3-mat3x3/f32.wgsl.expected.msl index d9fc80bf60..321a379624 100644 --- a/test/tint/expressions/binary/mul/vec3-mat3x3/f32.wgsl.expected.msl +++ b/test/tint/expressions/binary/mul/vec3-mat3x3/f32.wgsl.expected.msl @@ -2,16 +2,6 @@ using namespace metal; -template -inline vec operator*(matrix lhs, packed_vec rhs) { - return lhs * vec(rhs); -} - -template -inline vec operator*(packed_vec lhs, matrix rhs) { - return vec(lhs) * rhs; -} - template struct tint_array { const constant T& operator[](size_t i) const constant { return elements[i]; } @@ -31,7 +21,7 @@ struct S { }; fragment void tint_symbol_1(const constant S* tint_symbol_2 [[buffer(0)]]) { - float3 const x = ((*(tint_symbol_2)).vector * (*(tint_symbol_2)).tint_symbol); + float3 const x = (float3((*(tint_symbol_2)).vector) * (*(tint_symbol_2)).tint_symbol); return; } diff --git a/test/tint/expressions/binary/mul/vec3-mat4x3/f32.wgsl.expected.msl b/test/tint/expressions/binary/mul/vec3-mat4x3/f32.wgsl.expected.msl index 3f66659e22..b1f6165dff 100644 --- a/test/tint/expressions/binary/mul/vec3-mat4x3/f32.wgsl.expected.msl +++ b/test/tint/expressions/binary/mul/vec3-mat4x3/f32.wgsl.expected.msl @@ -2,16 +2,6 @@ using namespace metal; -template -inline vec operator*(matrix lhs, packed_vec rhs) { - return lhs * vec(rhs); -} - -template -inline vec operator*(packed_vec lhs, matrix rhs) { - return vec(lhs) * rhs; -} - template struct tint_array { const constant T& operator[](size_t i) const constant { return elements[i]; } @@ -31,7 +21,7 @@ struct S { }; fragment void tint_symbol_1(const constant S* tint_symbol_2 [[buffer(0)]]) { - float4 const x = ((*(tint_symbol_2)).vector * (*(tint_symbol_2)).tint_symbol); + float4 const x = (float3((*(tint_symbol_2)).vector) * (*(tint_symbol_2)).tint_symbol); return; } diff --git a/test/tint/expressions/binary/mul/vec3-scalar/i32.wgsl.expected.msl b/test/tint/expressions/binary/mul/vec3-scalar/i32.wgsl.expected.msl index 134ba04d8f..f320bd7520 100644 --- a/test/tint/expressions/binary/mul/vec3-scalar/i32.wgsl.expected.msl +++ b/test/tint/expressions/binary/mul/vec3-scalar/i32.wgsl.expected.msl @@ -4,7 +4,7 @@ using namespace metal; kernel void f() { int3 const a = int3(1, 2, 3); int const b = 4; - int3 const r = as_type((as_type(int3(a)) * as_type(b))); + int3 const r = as_type((as_type(a) * as_type(b))); return; } diff --git a/test/tint/expressions/binary/mul/vec3-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/mul/vec3-vec3/i32.wgsl.expected.msl index dbbcac50ca..f4511b2010 100644 --- a/test/tint/expressions/binary/mul/vec3-vec3/i32.wgsl.expected.msl +++ b/test/tint/expressions/binary/mul/vec3-vec3/i32.wgsl.expected.msl @@ -4,7 +4,7 @@ using namespace metal; kernel void f() { int3 const a = int3(1, 2, 3); int3 const b = int3(4, 5, 6); - int3 const r = as_type((as_type(int3(a)) * as_type(int3(b)))); + int3 const r = as_type((as_type(a) * as_type(b))); return; } diff --git a/test/tint/expressions/binary/sub/scalar-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/sub/scalar-vec3/i32.wgsl.expected.msl index 5403240cbc..8837be0e92 100644 --- a/test/tint/expressions/binary/sub/scalar-vec3/i32.wgsl.expected.msl +++ b/test/tint/expressions/binary/sub/scalar-vec3/i32.wgsl.expected.msl @@ -4,7 +4,7 @@ using namespace metal; kernel void f() { int const a = 4; int3 const b = int3(1, 2, 3); - int3 const r = as_type((as_type(a) - as_type(int3(b)))); + int3 const r = as_type((as_type(a) - as_type(b))); return; } diff --git a/test/tint/expressions/binary/sub/vec3-scalar/i32.wgsl.expected.msl b/test/tint/expressions/binary/sub/vec3-scalar/i32.wgsl.expected.msl index 9e59fcd91b..6149a0c537 100644 --- a/test/tint/expressions/binary/sub/vec3-scalar/i32.wgsl.expected.msl +++ b/test/tint/expressions/binary/sub/vec3-scalar/i32.wgsl.expected.msl @@ -4,7 +4,7 @@ using namespace metal; kernel void f() { int3 const a = int3(1, 2, 3); int const b = 4; - int3 const r = as_type((as_type(int3(a)) - as_type(b))); + int3 const r = as_type((as_type(a) - as_type(b))); return; } diff --git a/test/tint/expressions/binary/sub/vec3-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/sub/vec3-vec3/i32.wgsl.expected.msl index a42ec6ce13..142bb72d1a 100644 --- a/test/tint/expressions/binary/sub/vec3-vec3/i32.wgsl.expected.msl +++ b/test/tint/expressions/binary/sub/vec3-vec3/i32.wgsl.expected.msl @@ -4,7 +4,7 @@ using namespace metal; kernel void f() { int3 const a = int3(1, 2, 3); int3 const b = int3(4, 5, 6); - int3 const r = as_type((as_type(int3(a)) - as_type(int3(b)))); + int3 const r = as_type((as_type(a) - as_type(b))); return; } diff --git a/test/tint/expressions/swizzle/read/packed_vec3/f32.wgsl.expected.msl b/test/tint/expressions/swizzle/read/packed_vec3/f32.wgsl.expected.msl index e2bbca322e..1974ac0f87 100644 --- a/test/tint/expressions/swizzle/read/packed_vec3/f32.wgsl.expected.msl +++ b/test/tint/expressions/swizzle/read/packed_vec3/f32.wgsl.expected.msl @@ -2,16 +2,6 @@ using namespace metal; -template -inline vec operator*(matrix lhs, packed_vec rhs) { - return lhs * vec(rhs); -} - -template -inline vec operator*(packed_vec lhs, matrix rhs) { - return vec(lhs) * rhs; -} - template struct tint_array { const constant T& operator[](size_t i) const constant { return elements[i]; } @@ -30,7 +20,7 @@ struct S { }; void f(const constant S* const tint_symbol) { - float3 v = (*(tint_symbol)).v; + float3 v = float3((*(tint_symbol)).v); float x = (*(tint_symbol)).v[0]; float y = (*(tint_symbol)).v[1]; float z = (*(tint_symbol)).v[2]; diff --git a/test/tint/expressions/swizzle/read/packed_vec3/i32.wgsl.expected.msl b/test/tint/expressions/swizzle/read/packed_vec3/i32.wgsl.expected.msl index 7a93a594a7..cfabec5277 100644 --- a/test/tint/expressions/swizzle/read/packed_vec3/i32.wgsl.expected.msl +++ b/test/tint/expressions/swizzle/read/packed_vec3/i32.wgsl.expected.msl @@ -20,7 +20,7 @@ struct S { }; void f(const constant S* const tint_symbol) { - int3 v = (*(tint_symbol)).v; + int3 v = int3((*(tint_symbol)).v); int x = (*(tint_symbol)).v[0]; int y = (*(tint_symbol)).v[1]; int z = (*(tint_symbol)).v[2]; diff --git a/test/tint/expressions/swizzle/read/packed_vec3/u32.wgsl.expected.msl b/test/tint/expressions/swizzle/read/packed_vec3/u32.wgsl.expected.msl index 7444f75464..3f1899cbf2 100644 --- a/test/tint/expressions/swizzle/read/packed_vec3/u32.wgsl.expected.msl +++ b/test/tint/expressions/swizzle/read/packed_vec3/u32.wgsl.expected.msl @@ -20,7 +20,7 @@ struct S { }; void f(const constant S* const tint_symbol) { - uint3 v = (*(tint_symbol)).v; + uint3 v = uint3((*(tint_symbol)).v); uint x = (*(tint_symbol)).v[0]; uint y = (*(tint_symbol)).v[1]; uint z = (*(tint_symbol)).v[2]; diff --git a/test/tint/expressions/swizzle/write/packed_vec3/f32.wgsl.expected.msl b/test/tint/expressions/swizzle/write/packed_vec3/f32.wgsl.expected.msl index 7599cbc49e..8640cf6e75 100644 --- a/test/tint/expressions/swizzle/write/packed_vec3/f32.wgsl.expected.msl +++ b/test/tint/expressions/swizzle/write/packed_vec3/f32.wgsl.expected.msl @@ -2,16 +2,6 @@ using namespace metal; -template -inline vec operator*(matrix lhs, packed_vec rhs) { - return lhs * vec(rhs); -} - -template -inline vec operator*(packed_vec lhs, matrix rhs) { - return vec(lhs) * rhs; -} - template struct tint_array { const constant T& operator[](size_t i) const constant { return elements[i]; } diff --git a/test/tint/statements/compound_assign/vector/minus.wgsl.expected.msl b/test/tint/statements/compound_assign/vector/minus.wgsl.expected.msl index d6ddc129df..1a225d385b 100644 --- a/test/tint/statements/compound_assign/vector/minus.wgsl.expected.msl +++ b/test/tint/statements/compound_assign/vector/minus.wgsl.expected.msl @@ -6,6 +6,6 @@ struct S { }; void foo(device S* const tint_symbol) { - (*(tint_symbol)).a = as_type((as_type(int4((*(tint_symbol)).a)) - as_type(int4(int4(2))))); + (*(tint_symbol)).a = as_type((as_type((*(tint_symbol)).a) - as_type(int4(2)))); } diff --git a/test/tint/statements/compound_assign/vector/plus.wgsl.expected.msl b/test/tint/statements/compound_assign/vector/plus.wgsl.expected.msl index 455e61d7ea..7ea366b4d1 100644 --- a/test/tint/statements/compound_assign/vector/plus.wgsl.expected.msl +++ b/test/tint/statements/compound_assign/vector/plus.wgsl.expected.msl @@ -6,6 +6,6 @@ struct S { }; void foo(device S* const tint_symbol) { - (*(tint_symbol)).a = as_type((as_type(int4((*(tint_symbol)).a)) + as_type(int4(int4(2))))); + (*(tint_symbol)).a = as_type((as_type((*(tint_symbol)).a) + as_type(int4(2)))); } diff --git a/test/tint/statements/compound_assign/vector/shift_left.wgsl.expected.msl b/test/tint/statements/compound_assign/vector/shift_left.wgsl.expected.msl index 84a60be660..2dd819cebe 100644 --- a/test/tint/statements/compound_assign/vector/shift_left.wgsl.expected.msl +++ b/test/tint/statements/compound_assign/vector/shift_left.wgsl.expected.msl @@ -6,6 +6,6 @@ struct S { }; void foo(device S* const tint_symbol) { - (*(tint_symbol)).a = as_type((as_type(int4((*(tint_symbol)).a)) << uint4(2u))); + (*(tint_symbol)).a = as_type((as_type((*(tint_symbol)).a) << uint4(2u))); } diff --git a/test/tint/statements/compound_assign/vector/times.wgsl.expected.msl b/test/tint/statements/compound_assign/vector/times.wgsl.expected.msl index 3dfa83f58c..ef7fd0626a 100644 --- a/test/tint/statements/compound_assign/vector/times.wgsl.expected.msl +++ b/test/tint/statements/compound_assign/vector/times.wgsl.expected.msl @@ -6,6 +6,6 @@ struct S { }; void foo(device S* const tint_symbol) { - (*(tint_symbol)).a = as_type((as_type(int4((*(tint_symbol)).a)) * as_type(int4(int4(2))))); + (*(tint_symbol)).a = as_type((as_type((*(tint_symbol)).a) * as_type(int4(2)))); }