tint/writer/msl: Move packed_vector hacks to transform
Attempting to paper over all the MSL standard library holes for packed_vector in the MSL writer added complexity to the writer, produced messy output, and didn't actually catch all the cases where casts were needed.
Add a new PackedVec3 transform that applies the packed_vector -> vec casts in a smarter, more precise way.
Fixed: tint:1534
Change-Id: I73ce7e5a62fbc9cb04e1093133070f5fb8965dce
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/107340
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
Commit-Queue: Ben Clayton <bclayton@chromium.org>
diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index 6560b6f..d2c69e1 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -517,6 +517,8 @@
"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 @@
"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 af4384c..47b44a5 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -430,6 +430,8 @@
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 @@
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 0000000..dde5aca
--- /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 <algorithm>
+#include <string>
+#include <utility>
+
+#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<T> struct members
+ utils::Hashset<const sem::StructMember*, 8> 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<sem::Struct>(decl)) {
+ if (str->IsHostShareable()) {
+ for (auto* member : str->Members()) {
+ if (auto* vec = member->Type()->As<sem::Vector>()) {
+ 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<const ast::Attribute*, 4> attrs{
+ b.ASTNodes().Create<Attribute>(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<const sem::Expression*, 16> 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<ast::Let>() && // if variable is let...
+ v->Type()->Is<sem::Pointer>() && // 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<ast::UnaryOpExpression>()) {
+ 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<ast::AssignmentStatement>()) {
+ // 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<const sem::Vector*, Symbol, 3> 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<sem::Vector>()) {
+ 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<sem::Struct>(decl)) {
+ if (str->IsHostShareable()) {
+ for (auto* member : str->Members()) {
+ if (auto* vec = member->Type()->As<sem::Vector>()) {
+ 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<Attribute>(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 0000000..9d899cb
--- /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 <string>
+
+#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<T> type.
+/// * Cast all direct (not sub-accessed) loads of these packed vectors to the 'unpacked' vec3<T>
+/// 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<PackedVec3, Transform> {
+ public:
+ /// Attribute is the attribute applied to padded vector structure members.
+ class Attribute final : public Castable<Attribute, ast::InternalAttribute> {
+ 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 0000000..0f5c92e
--- /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 <string>
+#include <utility>
+#include <vector>
+
+#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<PackedVec3>(src));
+}
+
+TEST_F(PackedVec3Test, ShouldRun_NonHostSharableStruct) {
+ auto* src = R"(
+struct S {
+ v : vec3<f32>,
+}
+
+fn f() {
+ var v : S; // function address-space - not host sharable
+}
+)";
+
+ EXPECT_FALSE(ShouldRun<PackedVec3>(src));
+}
+
+TEST_F(PackedVec3Test, ShouldRun_Vec4Vec2) {
+ auto* src = R"(
+struct S {
+ v4 : vec4<f32>,
+ v2 : vec2<f32>,
+}
+
+@group(0) @binding(0) var<uniform> P : S; // Host sharable
+)";
+
+ EXPECT_FALSE(ShouldRun<PackedVec3>(src));
+}
+
+TEST_F(PackedVec3Test, ShouldRun_HostSharableStruct) {
+ auto* src = R"(
+struct S {
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<uniform> P : S; // Host sharable
+)";
+
+ EXPECT_TRUE(ShouldRun<PackedVec3>(src));
+}
+
+TEST_F(PackedVec3Test, UniformAddressSpace) {
+ auto* src = R"(
+struct S {
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<uniform> P : S;
+
+fn f() {
+ let x = P.v;
+}
+)";
+
+ auto* expect = R"(
+struct S {
+ @internal(packed_vector)
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<uniform> P : S;
+
+fn f() {
+ let x = vec3<f32>(P.v);
+}
+)";
+
+ DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, StorageAddressSpace) {
+ auto* src = R"(
+struct S {
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+ let x = P.v;
+}
+)";
+
+ auto* expect = R"(
+struct S {
+ @internal(packed_vector)
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+ let x = vec3<f32>(P.v);
+}
+)";
+
+ DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, ExistingMemberAttributes) {
+ auto* src = R"(
+struct S {
+ @align(32) @size(64) v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+ let x = P.v;
+}
+)";
+
+ auto* expect = R"(
+struct S {
+ @internal(packed_vector) @align(32) @size(64)
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+ let x = vec3<f32>(P.v);
+}
+)";
+
+ DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, MultipleVectors) {
+ auto* src = R"(
+struct S {
+ v2_a : vec2<f32>,
+ v3_a : vec3<f32>,
+ v4_a : vec4<f32>,
+ v2_b : vec2<f32>,
+ v3_b : vec3<f32>,
+ v4_b : vec4<f32>,
+}
+
+@group(0) @binding(0) var<storage> 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<f32>,
+ @internal(packed_vector)
+ v3_a : vec3<f32>,
+ v4_a : vec4<f32>,
+ v2_b : vec2<f32>,
+ @internal(packed_vector)
+ v3_b : vec3<f32>,
+ v4_b : vec4<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+ let v2_a = P.v2_a;
+ let v3_a = vec3<f32>(P.v3_a);
+ let v4_a = P.v4_a;
+ let v2_b = P.v2_b;
+ let v3_b = vec3<f32>(P.v3_b);
+ let v4_b = P.v4_b;
+}
+)";
+
+ DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, MixedAddressSpace) {
+ auto* src = R"(
+struct S {
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+ var f : S;
+ let x = f.v;
+}
+)";
+
+ auto* expect = R"(
+struct S {
+ @internal(packed_vector)
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+ var f : S;
+ let x = vec3<f32>(f.v);
+}
+)";
+
+ DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, ReadMemberAccessChain) {
+ auto* src = R"(
+struct S {
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+ let x = P.v.yz.x;
+}
+)";
+
+ auto* expect = R"(
+struct S {
+ @internal(packed_vector)
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+ let x = P.v.yz.x;
+}
+)";
+
+ DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, ReadVector) {
+ auto* src = R"(
+struct S {
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+ let x = P.v;
+}
+)";
+
+ auto* expect = R"(
+struct S {
+ @internal(packed_vector)
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+ let x = vec3<f32>(P.v);
+}
+)";
+
+ DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, ReadIndexAccessor) {
+ auto* src = R"(
+struct S {
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+ let x = P.v[1];
+}
+)";
+
+ auto* expect = R"(
+struct S {
+ @internal(packed_vector)
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+ let x = P.v[1];
+}
+)";
+
+ DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, ReadViaStructPtrDirect) {
+ auto* src = R"(
+struct S {
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+ let x = (*(&(*(&P)))).v;
+}
+)";
+
+ auto* expect = R"(
+struct S {
+ @internal(packed_vector)
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+ let x = vec3<f32>((*(&(*(&(P))))).v);
+}
+)";
+
+ DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, ReadViaVectorPtrDirect) {
+ auto* src = R"(
+struct S {
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+ let x = *(&(*(&(P.v))));
+}
+)";
+
+ auto* expect = R"(
+struct S {
+ @internal(packed_vector)
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+ let x = vec3<f32>(*(&(*(&(P.v)))));
+}
+)";
+
+ DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, ReadViaStructPtrViaLet) {
+ auto* src = R"(
+struct S {
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> 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<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+ let p0 = &(P);
+ let p1 = &(*(p0));
+ let a = vec3<f32>((*(p1)).v);
+ let p2 = &(*(p1));
+ let b = vec3<f32>((*(p2)).v);
+ let c = vec3<f32>((*(p2)).v);
+}
+)";
+
+ DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, ReadViaVectorPtrViaLet) {
+ auto* src = R"(
+struct S {
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> 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<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+ let p0 = &(P.v);
+ let p1 = &(*(p0));
+ let a = vec3<f32>(*(p1));
+ let p2 = &(*(p1));
+ let b = vec3<f32>(*(p2));
+ let c = vec3<f32>(*(p2));
+}
+)";
+
+ DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, ReadUnaryOp) {
+ auto* src = R"(
+struct S {
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+ let x = -P.v;
+}
+)";
+
+ auto* expect = R"(
+struct S {
+ @internal(packed_vector)
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+ let x = -(vec3<f32>(P.v));
+}
+)";
+
+ DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, ReadBinaryOp) {
+ auto* src = R"(
+struct S {
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+ let x = P.v + P.v;
+}
+)";
+
+ auto* expect = R"(
+struct S {
+ @internal(packed_vector)
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+ let x = (vec3<f32>(P.v) + vec3<f32>(P.v));
+}
+)";
+
+ DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, WriteVector) {
+ auto* src = R"(
+struct S {
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> P : S;
+
+fn f() {
+ P.v = vec3(1.23);
+}
+)";
+
+ auto* expect = R"(
+struct S {
+ @internal(packed_vector)
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> P : S;
+
+fn f() {
+ P.v = vec3(1.23);
+}
+)";
+
+ DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, WriteMemberAccess) {
+ auto* src = R"(
+struct S {
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> P : S;
+
+fn f() {
+ P.v.y = 1.23;
+}
+)";
+
+ auto* expect = R"(
+struct S {
+ @internal(packed_vector)
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> P : S;
+
+fn f() {
+ P.v.y = 1.23;
+}
+)";
+
+ DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, WriteIndexAccessor) {
+ auto* src = R"(
+struct S {
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> P : S;
+
+fn f() {
+ P.v[1] = 1.23;
+}
+)";
+
+ auto* expect = R"(
+struct S {
+ @internal(packed_vector)
+ v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> P : S;
+
+fn f() {
+ P.v[1] = 1.23;
+}
+)";
+
+ DataMap data;
+ auto got = Run<PackedVec3>(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 0bc8c39..cf67e10 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 @@
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<sem::Vector>();
-
- // 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 @@
// it assumes that the form of the array length argument is &var.array.
manager.Add<transform::ArrayLengthFromUniform>();
manager.Add<transform::ModuleScopeVarToEntryPointParam>();
+ manager.Add<transform::PackedVec3>();
data.Add<transform::ArrayLengthFromUniform::Config>(std::move(array_length_from_uniform_cfg));
data.Add<transform::CanonicalizeEntryPointIO::Config>(std::move(entry_point_io_cfg));
auto out = manager.Run(in, data);
@@ -554,18 +530,8 @@
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 @@
}
{
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 @@
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 @@
return false;
}
-bool GeneratorImpl::EmitPackedType(std::ostream& out,
- const sem::Type* type,
- const std::string& name) {
- auto* vec = type->As<sem::Vector>();
- 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<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
- return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
- return vec<T, M>(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,14 +2772,15 @@
}
add_byte_offset_comment(out, msl_offset);
+ }
- if (!EmitPackedType(out, mem->Type(), mem_name)) {
- return false;
+ if (auto* decl = mem->Declaration()) {
+ if (ast::HasAttribute<transform::PackedVec3::Attribute>(decl->attributes)) {
+ out << "packed_";
}
- } else {
- if (!EmitType(out, mem->Type(), mem_name)) {
- return false;
- }
+ }
+ if (!EmitType(out, mem->Type(), mem_name)) {
+ return false;
}
auto* ty = mem->Type();
@@ -2934,6 +2846,7 @@
[&](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 188bfea..c3e1ac0 100644
--- a/src/tint/writer/msl/generator_impl.h
+++ b/src/tint/writer/msl/generator_impl.h
@@ -328,14 +328,6 @@
/// @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 @@
/// 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<T, N> 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 e08c420..fc03507 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 @@
}
using Op = ast::BinaryOp;
constexpr BinaryData signed_overflow_defined_behaviour_chained_cases[] = {
- {"as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(a) << b))) << "
- "b))",
+ {R"(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(a) << b))) << b)))",
Op::kShiftLeft},
- {"((a >> b) >> b)", Op::kShiftRight},
- {"as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(a) + "
- "as_type<uint>(b)))) + as_type<uint>(b)))",
+ {R"(((a >> b) >> b))", Op::kShiftRight},
+ {R"(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(a) + as_type<uint>(b)))) + as_type<uint>(b))))",
Op::kAdd},
- {"as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(a) - "
- "as_type<uint>(b)))) - as_type<uint>(b)))",
+ {R"(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(a) - as_type<uint>(b)))) - as_type<uint>(b))))",
Op::kSubtract},
- {"as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(a) * "
- "as_type<uint>(b)))) * as_type<uint>(b)))",
+ {R"(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(a) * as_type<uint>(b)))) * as_type<uint>(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 bc9767a..27de589 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 @@
// 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 @@
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 @@
/* 0x0100 */ float2 tint_pad_33;
/* 0x0108 */ uint tint_pad_1;
/* 0x010c */ tint_array<int8_t, 4> 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 549e70f..52e5e7d 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<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
- return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
- return vec<T, M>(lhs) * rhs;
-}
-
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
@@ -42,11 +32,11 @@
};
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 e2430df..bb82069 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<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
- return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
- return vec<T, M>(lhs) * rhs;
-}
-
template<typename T, size_t N>
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 8f891a2..db51912 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<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
- return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
- return vec<T, M>(lhs) * rhs;
-}
-
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
@@ -43,11 +33,11 @@
};
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 251e06f..1b21710 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<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
- return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
- return vec<T, M>(lhs) * rhs;
-}
-
template<typename T, size_t N>
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 e361022..37817c3 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<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
- return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
- return vec<T, M>(lhs) * rhs;
-}
-
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
@@ -44,11 +34,11 @@
};
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 26497bd..02d091d 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 @@
T elements[N];
};
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
- return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
- return vec<T, M>(lhs) * rhs;
-}
-
struct Inner {
/* 0x0000 */ int x;
/* 0x0004 */ tint_array<int8_t, 12> tint_pad;
@@ -46,11 +36,11 @@
};
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 08f7854..1e627d7 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<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
- return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
- return vec<T, M>(lhs) * rhs;
-}
-
template<typename T, size_t N>
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 6507f58..c63e9f1 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 @@
T elements[N];
};
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
- return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
- return vec<T, M>(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 b62d4c5..bf546c1 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<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
- return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
- return vec<T, M>(lhs) * rhs;
-}
-
template<typename T, size_t N>
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 40e92e8..5539c67 100644
--- a/test/tint/bug/tint/1118.wgsl.expected.msl
+++ b/test/tint/bug/tint/1118.wgsl.expected.msl
@@ -1,17 +1,6 @@
#include <metal_stdlib>
using namespace metal;
-
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
- return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
- return vec<T, M>(lhs) * rhs;
-}
-
struct Scene {
/* 0x0000 */ float4 vEyePosition;
};
@@ -77,12 +66,12 @@
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 34c70b2..22cc6dd 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<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
- return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
- return vec<T, M>(lhs) * rhs;
-}
-
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
@@ -88,7 +78,7 @@
for(int x_1 = 0; (x_1 < TILE_COUNT_X); x_1 = as_type<int>((as_type<uint>(x_1) + as_type<uint>(1)))) {
int2 tilePixel0Idx = int2(as_type<int>((as_type<uint>(x_1) * as_type<uint>(TILE_SIZE))), as_type<int>((as_type<uint>(y_1) * as_type<uint>(TILE_SIZE))));
float2 floorCoord = (((2.0f * float2(tilePixel0Idx)) / float4((*(tint_symbol_3)).fullScreenSize).xy) - float2(1.0f));
- float2 ceilCoord = (((2.0f * float2(as_type<int2>((as_type<uint2>(int2(tilePixel0Idx)) + as_type<uint2>(int2(int2(TILE_SIZE))))))) / float4((*(tint_symbol_3)).fullScreenSize).xy) - float2(1.0f));
+ float2 ceilCoord = (((2.0f * float2(as_type<int2>((as_type<uint2>(tilePixel0Idx) + as_type<uint2>(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 3b8e4f7..d9ff17c 100644
--- a/test/tint/bug/tint/1520.spvasm.expected.msl
+++ b/test/tint/bug/tint/1520.spvasm.expected.msl
@@ -45,13 +45,13 @@
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<int4>((as_type<uint4>(int4(x_44)) + as_type<uint4>(int4(int4(1)))));
+ int4 const x_47 = as_type<int4>((as_type<uint4>(x_44) + as_type<uint4>(int4(1))));
val = x_47;
- int4 const x_48 = as_type<int4>((as_type<uint4>(int4(x_47)) - as_type<uint4>(int4(int4(1)))));
+ int4 const x_48 = as_type<int4>((as_type<uint4>(x_47) - as_type<uint4>(int4(1))));
val = x_48;
- int4 const x_49 = as_type<int4>((as_type<uint4>(int4(x_48)) + as_type<uint4>(int4(int4(1)))));
+ int4 const x_49 = as_type<int4>((as_type<uint4>(x_48) + as_type<uint4>(int4(1))));
val = x_49;
- int4 const x_50 = as_type<int4>((as_type<uint4>(int4(x_49)) - as_type<uint4>(int4(int4(1)))));
+ int4 const x_50 = as_type<int4>((as_type<uint4>(x_49) - as_type<uint4>(int4(1))));
val = x_50;
x_55 = false;
if (x_41) {
@@ -59,11 +59,11 @@
x_55 = x_54;
}
ok = x_55;
- int4 const x_58 = as_type<int4>((as_type<uint4>(int4(x_50)) * as_type<uint4>(int4(int4(2)))));
+ int4 const x_58 = as_type<int4>((as_type<uint4>(x_50) * as_type<uint4>(int4(2))));
val = x_58;
int4 const x_59 = (x_58 / int4(2));
val = x_59;
- int4 const x_60 = as_type<int4>((as_type<uint4>(int4(x_59)) * as_type<uint4>(int4(int4(2)))));
+ int4 const x_60 = as_type<int4>((as_type<uint4>(x_59) * as_type<uint4>(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 0000000..1b0c2a6
--- /dev/null
+++ b/test/tint/bug/tint/1534.wgsl
@@ -0,0 +1,17 @@
+
+struct g {
+ a : vec3<u32>,
+}
+
+struct h {
+ a : u32,
+}
+
+@group(0) @binding(0) var<uniform> i : g;
+
+@group(0) @binding(1) var<storage, read_write> 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 0000000..6f8a3de
--- /dev/null
+++ b/test/tint/bug/tint/1534.wgsl.expected.msl
@@ -0,0 +1,35 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+template<typename T>
+T tint_dot3(vec<T,3> a, vec<T,3> b) {
+ return a[0]*b[0] + a[1]*b[1] + a[2]*b[2];
+}
+struct g {
+ /* 0x0000 */ packed_uint3 a;
+ /* 0x000c */ tint_array<int8_t, 4> 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 6ee8bbd..580eba4 100644
--- a/test/tint/bug/tint/1677.wgsl.expected.msl
+++ b/test/tint/bug/tint/1677.wgsl.expected.msl
@@ -20,7 +20,7 @@
};
void tint_symbol_inner(uint3 id, const device Input* const tint_symbol_1) {
- int3 const pos = as_type<int3>((as_type<uint3>(int3((*(tint_symbol_1)).position)) - as_type<uint3>(int3(int3(0)))));
+ int3 const pos = as_type<int3>((as_type<uint3>(int3((*(tint_symbol_1)).position)) - as_type<uint3>(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 6113640..e77abf2 100644
--- a/test/tint/bug/tint/948.wgsl.expected.msl
+++ b/test/tint/bug/tint/948.wgsl.expected.msl
@@ -14,16 +14,6 @@
T elements[N];
};
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
- return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
- return vec<T, M>(lhs) * rhs;
-}
-
struct LeftOver {
/* 0x0000 */ float time;
/* 0x0004 */ uint padding;
@@ -193,7 +183,7 @@
i = as_type<int>((as_type<uint>(x_304) + as_type<uint>(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 c7bdad3..44f5320 100644
--- a/test/tint/bug/tint/949.wgsl.expected.msl
+++ b/test/tint/bug/tint/949.wgsl.expected.msl
@@ -14,16 +14,6 @@
T elements[N];
};
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
- return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
- return vec<T, M>(lhs) * rhs;
-}
-
struct lightingInfo {
float3 diffuse;
float3 specular;
@@ -248,7 +238,7 @@
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 @@
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 @@
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 2e23ef1..36cce9d 100644
--- a/test/tint/bug/tint/980.wgsl.expected.msl
+++ b/test/tint/bug/tint/980.wgsl.expected.msl
@@ -1,17 +1,6 @@
#include <metal_stdlib>
using namespace metal;
-
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
- return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
- return vec<T, M>(lhs) * rhs;
-}
-
float3 Bad(uint index, float3 rd) {
float3 normal = float3(0.0f);
normal[index] = -(sign(rd[index]));
@@ -24,7 +13,7 @@
};
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 8658c03..46f1909 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 @@
kernel void f() {
int const a = 4;
int3 const b = int3(1, 2, 3);
- int3 const r = as_type<int3>((as_type<uint>(a) + as_type<uint3>(int3(b))));
+ int3 const r = as_type<int3>((as_type<uint>(a) + as_type<uint3>(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 de4195f..89306cc 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 @@
kernel void f() {
int3 const a = int3(1, 2, 3);
int const b = 4;
- int3 const r = as_type<int3>((as_type<uint3>(int3(a)) + as_type<uint>(b)));
+ int3 const r = as_type<int3>((as_type<uint3>(a) + as_type<uint>(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 009e87e..54d88c7 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 @@
kernel void f() {
int3 const a = int3(1, 2, 3);
int3 const b = int3(4, 5, 6);
- int3 const r = as_type<int3>((as_type<uint3>(int3(a)) + as_type<uint3>(int3(b))));
+ int3 const r = as_type<int3>((as_type<uint3>(a) + as_type<uint3>(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 6d2f8bc..72cb90c 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 @@
kernel void f() {
int a = 4;
int3 b = int3(0, 2, 0);
- int3 const r = (a / as_type<int3>((as_type<uint3>(int3(b)) + as_type<uint3>(int3(b)))));
+ int3 const r = (a / as_type<int3>((as_type<uint3>(b) + as_type<uint3>(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 c90c00e..e2a2a76 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 @@
kernel void f() {
int3 a = int3(1, 2, 3);
int3 b = int3(0, 5, 0);
- int3 const r = (a / as_type<int3>((as_type<uint3>(int3(b)) + as_type<uint3>(int3(b)))));
+ int3 const r = (a / as_type<int3>((as_type<uint3>(b) + as_type<uint3>(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 1e4e146..f97db25 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 @@
kernel void f() {
int3 const a = int3(1, 2, 3);
uint3 const b = uint3(4u, 5u, 6u);
- int3 const r = as_type<int3>((as_type<uint3>(int3(a)) << b));
+ int3 const r = as_type<int3>((as_type<uint3>(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 22f6994..12f1a72 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 @@
kernel void f() {
int a = 4;
int3 b = int3(0, 2, 0);
- int3 const r = (a % as_type<int3>((as_type<uint3>(int3(b)) + as_type<uint3>(int3(b)))));
+ int3 const r = (a % as_type<int3>((as_type<uint3>(b) + as_type<uint3>(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 2d83cd3..61356dc 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 @@
kernel void f() {
int3 a = int3(1, 2, 3);
int3 b = int3(0, 5, 0);
- int3 const r = (a % as_type<int3>((as_type<uint3>(int3(b)) + as_type<uint3>(int3(b)))));
+ int3 const r = (a % as_type<int3>((as_type<uint3>(b) + as_type<uint3>(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 e210c3c..a36de59 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 @@
T elements[N];
};
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
- return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
- return vec<T, M>(lhs) * rhs;
-}
-
struct S {
/* 0x0000 */ float3x2 tint_symbol;
/* 0x0018 */ tint_array<int8_t, 8> tint_pad;
@@ -32,7 +22,7 @@
};
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 8f4ff55..79c3e09 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<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
- return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
- return vec<T, M>(lhs) * rhs;
-}
-
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
@@ -31,7 +21,7 @@
};
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 0a492f4..216de39 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 @@
kernel void f() {
int const a = 4;
int3 const b = int3(1, 2, 3);
- int3 const r = as_type<int3>((as_type<uint>(a) * as_type<uint3>(int3(b))));
+ int3 const r = as_type<int3>((as_type<uint>(a) * as_type<uint3>(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 d9fc80b..321a379 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<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
- return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
- return vec<T, M>(lhs) * rhs;
-}
-
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
@@ -31,7 +21,7 @@
};
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 3f66659..b1f6165 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<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
- return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
- return vec<T, M>(lhs) * rhs;
-}
-
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
@@ -31,7 +21,7 @@
};
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 134ba04..f320bd7 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 @@
kernel void f() {
int3 const a = int3(1, 2, 3);
int const b = 4;
- int3 const r = as_type<int3>((as_type<uint3>(int3(a)) * as_type<uint>(b)));
+ int3 const r = as_type<int3>((as_type<uint3>(a) * as_type<uint>(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 dbbcac5..f4511b2 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 @@
kernel void f() {
int3 const a = int3(1, 2, 3);
int3 const b = int3(4, 5, 6);
- int3 const r = as_type<int3>((as_type<uint3>(int3(a)) * as_type<uint3>(int3(b))));
+ int3 const r = as_type<int3>((as_type<uint3>(a) * as_type<uint3>(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 5403240..8837be0 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 @@
kernel void f() {
int const a = 4;
int3 const b = int3(1, 2, 3);
- int3 const r = as_type<int3>((as_type<uint>(a) - as_type<uint3>(int3(b))));
+ int3 const r = as_type<int3>((as_type<uint>(a) - as_type<uint3>(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 9e59fcd..6149a0c 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 @@
kernel void f() {
int3 const a = int3(1, 2, 3);
int const b = 4;
- int3 const r = as_type<int3>((as_type<uint3>(int3(a)) - as_type<uint>(b)));
+ int3 const r = as_type<int3>((as_type<uint3>(a) - as_type<uint>(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 a42ec6c..142bb72 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 @@
kernel void f() {
int3 const a = int3(1, 2, 3);
int3 const b = int3(4, 5, 6);
- int3 const r = as_type<int3>((as_type<uint3>(int3(a)) - as_type<uint3>(int3(b))));
+ int3 const r = as_type<int3>((as_type<uint3>(a) - as_type<uint3>(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 e2bbca3..1974ac0 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<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
- return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
- return vec<T, M>(lhs) * rhs;
-}
-
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
@@ -30,7 +20,7 @@
};
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 7a93a59..cfabec5 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 @@
};
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 7444f75..3f1899c 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 @@
};
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 7599cbc..8640cf6 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<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
- return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
- return vec<T, M>(lhs) * rhs;
-}
-
template<typename T, size_t N>
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 d6ddc12..1a225d3 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 @@
};
void foo(device S* const tint_symbol) {
- (*(tint_symbol)).a = as_type<int4>((as_type<uint4>(int4((*(tint_symbol)).a)) - as_type<uint4>(int4(int4(2)))));
+ (*(tint_symbol)).a = as_type<int4>((as_type<uint4>((*(tint_symbol)).a) - as_type<uint4>(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 455e61d..7ea366b 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 @@
};
void foo(device S* const tint_symbol) {
- (*(tint_symbol)).a = as_type<int4>((as_type<uint4>(int4((*(tint_symbol)).a)) + as_type<uint4>(int4(int4(2)))));
+ (*(tint_symbol)).a = as_type<int4>((as_type<uint4>((*(tint_symbol)).a) + as_type<uint4>(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 84a60be..2dd819c 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 @@
};
void foo(device S* const tint_symbol) {
- (*(tint_symbol)).a = as_type<int4>((as_type<uint4>(int4((*(tint_symbol)).a)) << uint4(2u)));
+ (*(tint_symbol)).a = as_type<int4>((as_type<uint4>((*(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 3dfa83f..ef7fd06 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 @@
};
void foo(device S* const tint_symbol) {
- (*(tint_symbol)).a = as_type<int4>((as_type<uint4>(int4((*(tint_symbol)).a)) * as_type<uint4>(int4(int4(2)))));
+ (*(tint_symbol)).a = as_type<int4>((as_type<uint4>((*(tint_symbol)).a) * as_type<uint4>(int4(2))));
}
