encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

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>
This commit is contained in:
Ben Clayton 2022-10-27 14:36:49 +00:00 committed by Dawn LUCI CQ
parent 59c0982426
commit a92f4259d5
53 changed files with 1070 additions and 379 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
src/tint/BUILD.gn
View File

@ -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",

3
src/tint/CMakeLists.txt
View File

@ -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

194
src/tint/transform/packed_vec3.cc Normal file
View File

@ -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

78
src/tint/transform/packed_vec3.h Normal file
View File

@ -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_

662
src/tint/transform/packed_vec3_test.cc Normal file
View File

@ -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

105
src/tint/writer/msl/generator_impl.cc
View File

@ -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<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 @@ SanitizedResult Sanitize(const Program* in, const Options& options) {
// 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,38 +530,18 @@ 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 (!emit_op()) {
return false;
}
{
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;
}
}
}
return true;
}
@ -601,19 +557,9 @@ 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 (!emit_op()) {
return false;
@ -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<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,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 (auto* decl = mem->Declaration()) {
if (ast::HasAttribute<transform::PackedVec3::Attribute>(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();

12
src/tint/writer/msl/generator_impl.h
View File

@ -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<T, N> template.
/// Non-empty only if the template has been generated.

14
src/tint/writer/msl/generator_impl_binary_test.cc
View File

@ -128,18 +128,14 @@ TEST_P(MslBinaryTest_SignedOverflowDefinedBehaviour_Chained, Emit) {
}
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,

6
src/tint/writer/msl/generator_impl_type_test.cc
View File

@ -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<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;

16
test/tint/buffer/storage/dynamic_index/read.wgsl.expected.msl
View File

@ -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 @@ 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;

10
test/tint/buffer/storage/dynamic_index/write.wgsl.expected.msl
View File

@ -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]; }

16
test/tint/buffer/storage/static_index/read.wgsl.expected.msl
View File

@ -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 @@ 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;

10
test/tint/buffer/storage/static_index/write.wgsl.expected.msl
View File

@ -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]; }

16
test/tint/buffer/uniform/dynamic_index/read.wgsl.expected.msl
View File

@ -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 @@ 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;

16
test/tint/buffer/uniform/static_index/read.wgsl.expected.msl
View File

@ -14,16 +14,6 @@ struct tint_array {
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 @@ 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;

10
test/tint/bug/chromium/1273230.wgsl.expected.msl
View File

@ -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]; }

10
test/tint/bug/fxc/indexed_assign_to_array_in_struct/1206.wgsl.expected.msl
View File

@ -14,16 +14,6 @@ struct tint_array {
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;
};

10
test/tint/bug/tint/1113.wgsl.expected.msl
View File

@ -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]; }

15
test/tint/bug/tint/1118.wgsl.expected.msl
View File

@ -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 @@ 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);

12
test/tint/bug/tint/1121.wgsl.expected.msl
View File

@ -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 @@ 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<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);

12
test/tint/bug/tint/1520.spvasm.expected.msl
View File

@ -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<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 @@ 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<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;

17
test/tint/bug/tint/1534.wgsl Normal file
View File

@ -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;
}

35
test/tint/bug/tint/1534.wgsl.expected.msl Normal file
View File

@ -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;
}

2
test/tint/bug/tint/1677.wgsl.expected.msl
View File

@ -20,7 +20,7 @@ struct Input {
};
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]]) {

12
test/tint/bug/tint/948.wgsl.expected.msl
View File

@ -14,16 +14,6 @@ struct tint_array {
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 @@ void main_1(thread float2* const tint_symbol_8, const constant LeftOver* const t
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;

16
test/tint/bug/tint/949.wgsl.expected.msl
View File

@ -14,16 +14,6 @@ struct tint_array {
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 @@ void main_1(thread float2* const tint_symbol_7, texture2d<float, access::sample>
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<float, access::sample>
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<float, access::sample>
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;

13
test/tint/bug/tint/980.wgsl.expected.msl
View File

@ -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 @@ 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;
}

2
test/tint/expressions/binary/add/scalar-vec3/i32.wgsl.expected.msl
View File

@ -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<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;
}

2
test/tint/expressions/binary/add/vec3-scalar/i32.wgsl.expected.msl
View File

@ -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<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;
}

2
test/tint/expressions/binary/add/vec3-vec3/i32.wgsl.expected.msl
View File

@ -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<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;
}

2
test/tint/expressions/binary/div_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl
View File

@ -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<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;
}

2
test/tint/expressions/binary/div_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl
View File

@ -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<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;
}

2
test/tint/expressions/binary/left-shift/vector-vector/i32.wgsl.expected.msl
View File

@ -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<int3>((as_type<uint3>(int3(a)) << b));
int3 const r = as_type<int3>((as_type<uint3>(a) << b));
return;
}

2
test/tint/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl
View File

@ -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<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;
}

2
test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl
View File

@ -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<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;
}

12
test/tint/expressions/binary/mul/mat3x2-vec3/f32.wgsl.expected.msl
View File

@ -14,16 +14,6 @@ struct tint_array {
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 @@ 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;
}

12
test/tint/expressions/binary/mul/mat3x3-vec3/f32.wgsl.expected.msl
View File

@ -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 @@ 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;
}

2
test/tint/expressions/binary/mul/scalar-vec3/i32.wgsl.expected.msl
View File

@ -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<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;
}

12
test/tint/expressions/binary/mul/vec3-mat3x3/f32.wgsl.expected.msl
View File

@ -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 @@ 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;
}

12
test/tint/expressions/binary/mul/vec3-mat4x3/f32.wgsl.expected.msl
View File

@ -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 @@ 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;
}

2
test/tint/expressions/binary/mul/vec3-scalar/i32.wgsl.expected.msl
View File

@ -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<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;
}

2
test/tint/expressions/binary/mul/vec3-vec3/i32.wgsl.expected.msl
View File

@ -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<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;
}

2
test/tint/expressions/binary/sub/scalar-vec3/i32.wgsl.expected.msl
View File

@ -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<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;
}

2
test/tint/expressions/binary/sub/vec3-scalar/i32.wgsl.expected.msl
View File

@ -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<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;
}

2
test/tint/expressions/binary/sub/vec3-vec3/i32.wgsl.expected.msl
View File

@ -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<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;
}

12
test/tint/expressions/swizzle/read/packed_vec3/f32.wgsl.expected.msl
View File

@ -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 @@ 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];

2
test/tint/expressions/swizzle/read/packed_vec3/i32.wgsl.expected.msl
View File

@ -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];

2
test/tint/expressions/swizzle/read/packed_vec3/u32.wgsl.expected.msl
View File

@ -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];

10
test/tint/expressions/swizzle/write/packed_vec3/f32.wgsl.expected.msl
View File

@ -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]; }

2
test/tint/statements/compound_assign/vector/minus.wgsl.expected.msl
View File

@ -6,6 +6,6 @@ struct S {
};
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))));
}

2
test/tint/statements/compound_assign/vector/plus.wgsl.expected.msl
View File

@ -6,6 +6,6 @@ struct S {
};
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))));
}

2
test/tint/statements/compound_assign/vector/shift_left.wgsl.expected.msl
View File

@ -6,6 +6,6 @@ struct S {
};
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)));
}

2
test/tint/statements/compound_assign/vector/times.wgsl.expected.msl
View File

@ -6,6 +6,6 @@ struct S {
};
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))));
}