Implement phony assignment
Bug: tint:1213
Change-Id: Ib1ebc4947405c4ada7a9bdbc6bd5a36447bbd234
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/67064
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
diff --git a/src/BUILD.gn b/src/BUILD.gn
index b400b27..d8b7baa 100644
--- a/src/BUILD.gn
+++ b/src/BUILD.gn
@@ -463,6 +463,8 @@
"transform/pad_array_elements.h",
"transform/promote_initializers_to_const_var.cc",
"transform/promote_initializers_to_const_var.h",
+ "transform/remove_phonies.cc",
+ "transform/remove_phonies.h",
"transform/renamer.cc",
"transform/renamer.h",
"transform/robustness.cc",
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index a77bb9b..fd9bfbe 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -329,6 +329,8 @@
transform/pad_array_elements.h
transform/promote_initializers_to_const_var.cc
transform/promote_initializers_to_const_var.h
+ transform/remove_phonies.cc
+ transform/remove_phonies.h
transform/renamer.cc
transform/renamer.h
transform/robustness.cc
@@ -955,6 +957,7 @@
transform/num_workgroups_from_uniform_test.cc
transform/pad_array_elements_test.cc
transform/promote_initializers_to_const_var_test.cc
+ transform/remove_phonies_test.cc
transform/renamer_test.cc
transform/robustness_test.cc
transform/simplify_test.cc
diff --git a/src/program_builder.h b/src/program_builder.h
index 01ca389..fd3e685 100644
--- a/src/program_builder.h
+++ b/src/program_builder.h
@@ -1762,6 +1762,18 @@
Expr(std::forward<IDX>(idx)));
}
+ /// @param source the source information
+ /// @param obj the object for the member accessor expression
+ /// @param idx the index argument for the array accessor expression
+ /// @returns a `ast::MemberAccessorExpression` that indexes `obj` with `idx`
+ template <typename OBJ, typename IDX>
+ const ast::MemberAccessorExpression* MemberAccessor(const Source& source,
+ OBJ&& obj,
+ IDX&& idx) {
+ return create<ast::MemberAccessorExpression>(
+ source, Expr(std::forward<OBJ>(obj)), Expr(std::forward<IDX>(idx)));
+ }
+
/// @param obj the object for the member accessor expression
/// @param idx the index argument for the array accessor expression
/// @returns a `ast::MemberAccessorExpression` that indexes `obj` with `idx`
diff --git a/src/reader/wgsl/parser_impl.cc b/src/reader/wgsl/parser_impl.cc
index 6ab154a..fdcb027 100644
--- a/src/reader/wgsl/parser_impl.cc
+++ b/src/reader/wgsl/parser_impl.cc
@@ -2792,7 +2792,7 @@
}
// assignment_stmt
-// : unary_expression EQUAL logical_or_expression
+// : (unary_expression | underscore) EQUAL logical_or_expression
Maybe<const ast::AssignmentStatement*> ParserImpl::assignment_stmt() {
auto t = peek();
auto source = t.source();
@@ -2806,19 +2806,27 @@
}
auto lhs = unary_expression();
- if (lhs.errored)
+ if (lhs.errored) {
return Failure::kErrored;
- if (!lhs.matched)
- return Failure::kNoMatch;
+ }
+ if (!lhs.matched) {
+ if (!match(Token::Type::kUnderscore, &source)) {
+ return Failure::kNoMatch;
+ }
+ lhs = create<ast::PhonyExpression>(source);
+ }
- if (!expect("assignment", Token::Type::kEqual))
+ if (!expect("assignment", Token::Type::kEqual)) {
return Failure::kErrored;
+ }
auto rhs = logical_or_expression();
- if (rhs.errored)
+ if (rhs.errored) {
return Failure::kErrored;
- if (!rhs.matched)
+ }
+ if (!rhs.matched) {
return add_error(peek(), "unable to parse right side of assignment");
+ }
return create<ast::AssignmentStatement>(source, lhs.value, rhs.value);
}
diff --git a/src/reader/wgsl/parser_impl_assignment_stmt_test.cc b/src/reader/wgsl/parser_impl_assignment_stmt_test.cc
index ef4ec97..7fe7097 100644
--- a/src/reader/wgsl/parser_impl_assignment_stmt_test.cc
+++ b/src/reader/wgsl/parser_impl_assignment_stmt_test.cc
@@ -98,6 +98,28 @@
EXPECT_EQ(ident->symbol, p->builder().Symbols().Get("b"));
}
+TEST_F(ParserImplTest, AssignmentStmt_Parses_ToPhony) {
+ auto p = parser("_ = 123");
+ auto e = p->assignment_stmt();
+ EXPECT_TRUE(e.matched);
+ EXPECT_FALSE(e.errored);
+ EXPECT_FALSE(p->has_error()) << p->error();
+ ASSERT_NE(e.value, nullptr);
+
+ ASSERT_TRUE(e->Is<ast::AssignmentStatement>());
+ ASSERT_NE(e->lhs, nullptr);
+ ASSERT_NE(e->rhs, nullptr);
+
+ ASSERT_TRUE(e->rhs->Is<ast::ConstructorExpression>());
+ ASSERT_TRUE(e->rhs->Is<ast::ScalarConstructorExpression>());
+ auto* init = e->rhs->As<ast::ScalarConstructorExpression>();
+ ASSERT_NE(init->literal, nullptr);
+ ASSERT_TRUE(init->literal->Is<ast::SintLiteral>());
+ EXPECT_EQ(init->literal->As<ast::SintLiteral>()->value, 123);
+
+ ASSERT_TRUE(e->lhs->Is<ast::PhonyExpression>());
+}
+
TEST_F(ParserImplTest, AssignmentStmt_MissingEqual) {
auto p = parser("a.b.c[2].d 123");
auto e = p->assignment_stmt();
diff --git a/src/resolver/assignment_validation_test.cc b/src/resolver/assignment_validation_test.cc
index c0ed563..1691864 100644
--- a/src/resolver/assignment_validation_test.cc
+++ b/src/resolver/assignment_validation_test.cc
@@ -293,6 +293,111 @@
"56:78 error: storage type of assignment must be constructible");
}
+TEST_F(ResolverAssignmentValidationTest,
+ AssignToPhony_NonConstructableStruct_Fail) {
+ // [[block]]
+ // struct S {
+ // arr: array<i32>;
+ // };
+ // [[group(0), binding(0)]] var<storage, read_write> s : S;
+ // fn f() {
+ // _ = s;
+ // }
+ auto* s = Structure("S", {Member("arr", ty.array<i32>())}, {StructBlock()});
+ Global("s", ty.Of(s), ast::StorageClass::kStorage, GroupAndBinding(0, 0));
+
+ WrapInFunction(Assign(Phony(), Expr(Source{{12, 34}}, "s")));
+
+ EXPECT_FALSE(r()->Resolve());
+ EXPECT_EQ(r()->error(),
+ "12:34 error: cannot assign 'S' to '_'. "
+ "'_' can only be assigned a constructable, pointer, texture or "
+ "sampler type");
+}
+
+TEST_F(ResolverAssignmentValidationTest, AssignToPhony_DynamicArray_Fail) {
+ // [[block]]
+ // struct S {
+ // arr: array<i32>;
+ // };
+ // [[group(0), binding(0)]] var<storage, read_write> s : S;
+ // fn f() {
+ // _ = s.arr;
+ // }
+ auto* s = Structure("S", {Member("arr", ty.array<i32>())}, {StructBlock()});
+ Global("s", ty.Of(s), ast::StorageClass::kStorage, GroupAndBinding(0, 0));
+
+ WrapInFunction(Assign(Phony(), MemberAccessor(Source{{12, 34}}, "s", "arr")));
+
+ EXPECT_FALSE(r()->Resolve());
+ EXPECT_EQ(
+ r()->error(),
+ "12:34 error: cannot assign 'ref<storage, array<i32>, read>' to '_'. "
+ "'_' can only be assigned a constructable, pointer, texture or sampler "
+ "type");
+}
+
+TEST_F(ResolverAssignmentValidationTest, AssignToPhony_Pass) {
+ // [[block]]
+ // struct S {
+ // i: i32;
+ // arr: array<i32>;
+ // };
+ // [[block]]
+ // struct U {
+ // i: i32;
+ // };
+ // [[group(0), binding(0)]] var tex texture_2d;
+ // [[group(0), binding(1)]] var smp sampler;
+ // [[group(0), binding(2)]] var<uniform> u : U;
+ // [[group(0), binding(3)]] var<storage, read_write> s : S;
+ // var<workgroup> wg : array<f32, 10>
+ // fn f() {
+ // _ = 1;
+ // _ = 2u;
+ // _ = 3.0;
+ // _ = vec2<bool>();
+ // _ = tex;
+ // _ = smp;
+ // _ = &s;
+ // _ = s.i;
+ // _ = &s.arr;
+ // _ = u;
+ // _ = u.i;
+ // _ = wg;
+ // _ = wg[3];
+ // }
+ auto* S = Structure("S",
+ {
+ Member("i", ty.i32()),
+ Member("arr", ty.array<i32>()),
+ },
+ {StructBlock()});
+ auto* U = Structure("U", {Member("i", ty.i32())}, {StructBlock()});
+ Global("tex", ty.sampled_texture(ast::TextureDimension::k2d, ty.f32()),
+ GroupAndBinding(0, 0));
+ Global("smp", ty.sampler(ast::SamplerKind::kSampler), GroupAndBinding(0, 1));
+ Global("u", ty.Of(U), ast::StorageClass::kUniform, GroupAndBinding(0, 2));
+ Global("s", ty.Of(S), ast::StorageClass::kStorage, GroupAndBinding(0, 3));
+ Global("wg", ty.array<f32, 10>(), ast::StorageClass::kWorkgroup);
+
+ WrapInFunction(Assign(Phony(), 1), //
+ Assign(Phony(), 2), //
+ Assign(Phony(), 3), //
+ Assign(Phony(), vec2<bool>()), //
+ Assign(Phony(), "tex"), //
+ Assign(Phony(), "smp"), //
+ Assign(Phony(), AddressOf("s")), //
+ Assign(Phony(), MemberAccessor("s", "i")), //
+ Assign(Phony(), AddressOf(MemberAccessor("s", "arr"))), //
+ Assign(Phony(), "u"), //
+ Assign(Phony(), MemberAccessor("u", "i")), //
+ Assign(Phony(), "wg"), //
+ Assign(Phony(), IndexAccessor("wg", 3)));
+
+ EXPECT_TRUE(r()->Resolve()) << r()->error();
+}
+
} // namespace
} // namespace resolver
} // namespace tint
diff --git a/src/resolver/resolver.cc b/src/resolver/resolver.cc
index aee387b..49a5564 100644
--- a/src/resolver/resolver.cc
+++ b/src/resolver/resolver.cc
@@ -2303,6 +2303,8 @@
ok = MemberAccessor(member);
} else if (auto* unary = expr->As<ast::UnaryOpExpression>()) {
ok = UnaryOp(unary);
+ } else if (expr->Is<ast::PhonyExpression>()) {
+ ok = true; // No-op
} else {
TINT_ICE(Resolver, diagnostics_)
<< "unhandled expression type: " << expr->TypeInfo().name;
@@ -4393,13 +4395,30 @@
if (!Expression(a->lhs) || !Expression(a->rhs)) {
return false;
}
+
return ValidateAssignment(a);
}
bool Resolver::ValidateAssignment(const ast::AssignmentStatement* a) {
+ auto const* rhs_type = TypeOf(a->rhs);
+
+ if (a->lhs->Is<ast::PhonyExpression>()) {
+ // https://www.w3.org/TR/WGSL/#phony-assignment-section
+ auto* ty = rhs_type->UnwrapRef();
+ if (!ty->IsConstructible() &&
+ !ty->IsAnyOf<sem::Pointer, sem::Texture, sem::Sampler>()) {
+ AddError(
+ "cannot assign '" + TypeNameOf(a->rhs) +
+ "' to '_'. '_' can only be assigned a constructable, pointer, "
+ "texture or sampler type",
+ a->rhs->source);
+ return false;
+ }
+ return true; // RHS can be anything.
+ }
+
// https://gpuweb.github.io/gpuweb/wgsl/#assignment-statement
auto const* lhs_type = TypeOf(a->lhs);
- auto const* rhs_type = TypeOf(a->rhs);
if (auto* ident = a->lhs->As<ast::IdentifierExpression>()) {
VariableInfo* var;
diff --git a/src/transform/glsl.cc b/src/transform/glsl.cc
index 1023e4f..b3b5972 100644
--- a/src/transform/glsl.cc
+++ b/src/transform/glsl.cc
@@ -27,6 +27,7 @@
#include "src/transform/manager.h"
#include "src/transform/pad_array_elements.h"
#include "src/transform/promote_initializers_to_const_var.h"
+#include "src/transform/remove_phonies.h"
#include "src/transform/simplify.h"
#include "src/transform/single_entry_point.h"
#include "src/transform/zero_init_workgroup_memory.h"
@@ -58,6 +59,7 @@
}
manager.Add<CanonicalizeEntryPointIO>();
manager.Add<InlinePointerLets>();
+ manager.Add<RemovePhonies>();
// Simplify cleans up messy `*(&(expr))` expressions from InlinePointerLets.
manager.Add<Simplify>();
manager.Add<CalculateArrayLength>();
diff --git a/src/transform/remove_phonies.cc b/src/transform/remove_phonies.cc
new file mode 100644
index 0000000..fe8b75f
--- /dev/null
+++ b/src/transform/remove_phonies.cc
@@ -0,0 +1,138 @@
+// Copyright 2021 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/transform/remove_phonies.h"
+
+#include <memory>
+#include <unordered_map>
+#include <utility>
+#include <vector>
+
+#include "src/ast/traverse_expressions.h"
+#include "src/program_builder.h"
+#include "src/sem/block_statement.h"
+#include "src/sem/function.h"
+#include "src/sem/statement.h"
+#include "src/sem/variable.h"
+#include "src/utils/get_or_create.h"
+#include "src/utils/scoped_assignment.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::transform::RemovePhonies);
+
+namespace tint {
+namespace transform {
+namespace {
+
+struct SinkSignature {
+ std::vector<const sem::Type*> types;
+
+ bool operator==(const SinkSignature& other) const {
+ if (types.size() != other.types.size()) {
+ return false;
+ }
+ for (size_t i = 0; i < types.size(); i++) {
+ if (types[i] != other.types[i]) {
+ return false;
+ }
+ }
+ return true;
+ }
+
+ struct Hasher {
+ /// @param sig the CallTargetSignature to hash
+ /// @return the hash value
+ std::size_t operator()(const SinkSignature& sig) const {
+ size_t hash = tint::utils::Hash(sig.types.size());
+ for (auto* ty : sig.types) {
+ tint::utils::HashCombine(&hash, ty);
+ }
+ return hash;
+ }
+ };
+};
+
+} // namespace
+
+RemovePhonies::RemovePhonies() = default;
+
+RemovePhonies::~RemovePhonies() = default;
+
+void RemovePhonies::Run(CloneContext& ctx, const DataMap&, DataMap&) {
+ auto& sem = ctx.src->Sem();
+
+ std::unordered_map<SinkSignature, Symbol, SinkSignature::Hasher> sinks;
+
+ for (auto* node : ctx.src->ASTNodes().Objects()) {
+ if (auto* stmt = node->As<ast::AssignmentStatement>()) {
+ if (stmt->lhs->Is<ast::PhonyExpression>()) {
+ std::vector<const ast::Expression*> side_effects;
+ if (!ast::TraverseExpressions(stmt->rhs, ctx.dst->Diagnostics(),
+ [&](const ast::CallExpression* call) {
+ side_effects.push_back(call);
+ return ast::TraverseAction::Skip;
+ })) {
+ return;
+ }
+
+ if (side_effects.empty()) {
+ // Phony assignment with no side effects.
+ // Just remove it.
+ RemoveStatement(ctx, stmt);
+ continue;
+ }
+
+ if (side_effects.size() == 1) {
+ if (auto* call = side_effects[0]->As<ast::CallExpression>()) {
+ // Phony assignment with single call side effect.
+ // Replace phony assignment with call.
+ ctx.Replace(
+ stmt, [&, call] { return ctx.dst->CallStmt(ctx.Clone(call)); });
+ continue;
+ }
+ }
+
+ // Phony assignment with multiple side effects.
+ // Generate a call to a dummy function with the side effects as
+ // arguments.
+ ctx.Replace(stmt, [&, side_effects] {
+ SinkSignature sig;
+ for (auto* arg : side_effects) {
+ sig.types.push_back(sem.Get(arg)->Type()->UnwrapRef());
+ }
+ auto sink = utils::GetOrCreate(sinks, sig, [&] {
+ auto name = ctx.dst->Symbols().New("phony_sink");
+ ast::VariableList params;
+ for (auto* ty : sig.types) {
+ auto* ast_ty = CreateASTTypeFor(ctx, ty);
+ params.push_back(
+ ctx.dst->Param("p" + std::to_string(params.size()), ast_ty));
+ }
+ ctx.dst->Func(name, params, ctx.dst->ty.void_(), {});
+ return name;
+ });
+ ast::ExpressionList args;
+ for (auto* arg : side_effects) {
+ args.push_back(ctx.Clone(arg));
+ }
+ return ctx.dst->CallStmt(ctx.dst->Call(sink, args));
+ });
+ }
+ }
+ }
+
+ ctx.Clone();
+}
+
+} // namespace transform
+} // namespace tint
diff --git a/src/transform/remove_phonies.h b/src/transform/remove_phonies.h
new file mode 100644
index 0000000..c94959b
--- /dev/null
+++ b/src/transform/remove_phonies.h
@@ -0,0 +1,50 @@
+// Copyright 2021 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_TRANSFORM_REMOVE_PHONIES_H_
+#define SRC_TRANSFORM_REMOVE_PHONIES_H_
+
+#include <string>
+#include <unordered_map>
+
+#include "src/transform/transform.h"
+
+namespace tint {
+namespace transform {
+
+/// RemovePhonies is a Transform that removes all phony-assignment statements,
+/// while preserving function call expressions in the RHS of the assignment that
+/// may have side-effects.
+class RemovePhonies : public Castable<RemovePhonies, Transform> {
+ public:
+ /// Constructor
+ RemovePhonies();
+
+ /// Destructor
+ ~RemovePhonies() override;
+
+ protected:
+ /// 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) override;
+};
+
+} // namespace transform
+} // namespace tint
+
+#endif // SRC_TRANSFORM_REMOVE_PHONIES_H_
diff --git a/src/transform/remove_phonies_test.cc b/src/transform/remove_phonies_test.cc
new file mode 100644
index 0000000..9629ab8
--- /dev/null
+++ b/src/transform/remove_phonies_test.cc
@@ -0,0 +1,226 @@
+// Copyright 2021 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/transform/remove_phonies.h"
+
+#include <memory>
+#include <utility>
+#include <vector>
+
+#include "src/transform/test_helper.h"
+
+namespace tint {
+namespace transform {
+namespace {
+
+using RemovePhoniesTest = TransformTest;
+
+TEST_F(RemovePhoniesTest, EmptyModule) {
+ auto* src = "";
+ auto* expect = "";
+
+ auto got = Run<RemovePhonies>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(RemovePhoniesTest, NoSideEffects) {
+ auto* src = R"(
+[[group(0), binding(0)]] var t : texture_2d<f32>;
+
+fn f() {
+ var v : i32;
+ _ = &v;
+ _ = 1;
+ _ = 1 + 2;
+ _ = t;
+}
+)";
+
+ auto* expect = R"(
+[[group(0), binding(0)]] var t : texture_2d<f32>;
+
+fn f() {
+ var v : i32;
+}
+)";
+
+ auto got = Run<RemovePhonies>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(RemovePhoniesTest, SingleSideEffects) {
+ auto* src = R"(
+fn neg(a : i32) -> i32 {
+ return -(a);
+}
+
+fn add(a : i32, b : i32) -> i32 {
+ return (a + b);
+}
+
+fn f() {
+ _ = neg(1);
+ _ = add(2, 3);
+ _ = add(neg(4), neg(5));
+}
+)";
+
+ auto* expect = R"(
+fn neg(a : i32) -> i32 {
+ return -(a);
+}
+
+fn add(a : i32, b : i32) -> i32 {
+ return (a + b);
+}
+
+fn f() {
+ neg(1);
+ add(2, 3);
+ add(neg(4), neg(5));
+}
+)";
+
+ auto got = Run<RemovePhonies>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(RemovePhoniesTest, MultipleSideEffects) {
+ auto* src = R"(
+fn neg(a : i32) -> i32 {
+ return -(a);
+}
+
+fn add(a : i32, b : i32) -> i32 {
+ return (a + b);
+}
+
+fn xor(a : u32, b : u32) -> u32 {
+ return (a ^ b);
+}
+
+fn f() {
+ _ = (1 + add(2 + add(3, 4), 5)) * add(6, 7) * neg(8);
+ _ = add(9, neg(10)) + neg(11);
+ _ = xor(12u, 13u) + xor(14u, 15u);
+ _ = neg(16) / neg(17) + add(18, 19);
+}
+)";
+
+ auto* expect = R"(
+fn neg(a : i32) -> i32 {
+ return -(a);
+}
+
+fn add(a : i32, b : i32) -> i32 {
+ return (a + b);
+}
+
+fn xor(a : u32, b : u32) -> u32 {
+ return (a ^ b);
+}
+
+fn phony_sink(p0 : i32, p1 : i32, p2 : i32) {
+}
+
+fn phony_sink_1(p0 : i32, p1 : i32) {
+}
+
+fn phony_sink_2(p0 : u32, p1 : u32) {
+}
+
+fn f() {
+ phony_sink(add((2 + add(3, 4)), 5), add(6, 7), neg(8));
+ phony_sink_1(add(9, neg(10)), neg(11));
+ phony_sink_2(xor(12u, 13u), xor(14u, 15u));
+ phony_sink(neg(16), neg(17), add(18, 19));
+}
+)";
+
+ auto got = Run<RemovePhonies>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(RemovePhoniesTest, ForLoop) {
+ auto* src = R"(
+[[block]]
+struct S {
+ arr : array<i32>;
+};
+
+[[group(0), binding(0)]] var<storage, read_write> s : S;
+
+fn x() -> i32 {
+ return 0;
+}
+
+fn y() -> i32 {
+ return 0;
+}
+
+fn z() -> i32 {
+ return 0;
+}
+
+fn f() {
+ for (_ = &s.arr; ;_ = &s.arr) {
+ }
+ for (_ = x(); ;_ = y() + z()) {
+ }
+}
+)";
+
+ auto* expect = R"(
+[[block]]
+struct S {
+ arr : array<i32>;
+};
+
+[[group(0), binding(0)]] var<storage, read_write> s : S;
+
+fn x() -> i32 {
+ return 0;
+}
+
+fn y() -> i32 {
+ return 0;
+}
+
+fn z() -> i32 {
+ return 0;
+}
+
+fn phony_sink(p0 : i32, p1 : i32) {
+}
+
+fn f() {
+ for(; ; ) {
+ }
+ for(x(); ; phony_sink(y(), z())) {
+ }
+}
+)";
+
+ auto got = Run<RemovePhonies>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+} // namespace
+} // namespace transform
+} // namespace tint
diff --git a/src/transform/transform.cc b/src/transform/transform.cc
index 34f5ec6..bf94f08 100644
--- a/src/transform/transform.cc
+++ b/src/transform/transform.cc
@@ -75,7 +75,7 @@
return true;
}
-void Transform::RemoveStatement(CloneContext& ctx, ast::Statement* stmt) {
+void Transform::RemoveStatement(CloneContext& ctx, const ast::Statement* stmt) {
auto* sem = ctx.src->Sem().Get(stmt);
if (auto* block = tint::As<sem::BlockStatement>(sem->Parent())) {
ctx.Remove(block->Declaration()->statements, stmt);
diff --git a/src/transform/transform.h b/src/transform/transform.h
index 6d34997..59a04e7 100644
--- a/src/transform/transform.h
+++ b/src/transform/transform.h
@@ -190,7 +190,7 @@
/// continuing of for-loops.
/// @param ctx the clone context
/// @param stmt the statement to remove when the program is cloned
- static void RemoveStatement(CloneContext& ctx, ast::Statement* stmt);
+ static void RemoveStatement(CloneContext& ctx, const ast::Statement* stmt);
/// CreateASTTypeFor constructs new ast::Type nodes that reconstructs the
/// semantic type `ty`.
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index f0dee54..bedb2ae 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -54,6 +54,7 @@
#include "src/transform/num_workgroups_from_uniform.h"
#include "src/transform/pad_array_elements.h"
#include "src/transform/promote_initializers_to_const_var.h"
+#include "src/transform/remove_phonies.h"
#include "src/transform/simplify.h"
#include "src/transform/zero_init_workgroup_memory.h"
#include "src/utils/defer.h"
@@ -138,10 +139,14 @@
manager.Add<transform::InlinePointerLets>();
// Simplify cleans up messy `*(&(expr))` expressions from InlinePointerLets.
manager.Add<transform::Simplify>();
- // DecomposeMemoryAccess must come after InlinePointerLets as we cannot take
- // the address of calls to DecomposeMemoryAccess::Intrinsic. Must also come
- // after Simplify, as we need to fold away the address-of and defers of
+ manager.Add<transform::RemovePhonies>();
+ // DecomposeMemoryAccess must come after:
+ // * InlinePointerLets, as we cannot take the address of calls to
+ // DecomposeMemoryAccess::Intrinsic.
+ // * Simplify, as we need to fold away the address-of and dereferences of
// `*(&(intrinsic_load()))` expressions.
+ // * RemovePhonies, as phonies can be assigned a pointer to a
+ // non-constructable buffer, or dynamic array, which DMA cannot cope with.
manager.Add<transform::DecomposeMemoryAccess>();
// CalculateArrayLength must come after DecomposeMemoryAccess, as
// DecomposeMemoryAccess special-cases the arrayLength() intrinsic, which
diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc
index ecc63ca..715acd7 100644
--- a/src/writer/msl/generator_impl.cc
+++ b/src/writer/msl/generator_impl.cc
@@ -63,6 +63,7 @@
#include "src/transform/module_scope_var_to_entry_point_param.h"
#include "src/transform/pad_array_elements.h"
#include "src/transform/promote_initializers_to_const_var.h"
+#include "src/transform/remove_phonies.h"
#include "src/transform/simplify.h"
#include "src/transform/wrap_arrays_in_structs.h"
#include "src/transform/zero_init_workgroup_memory.h"
@@ -145,6 +146,7 @@
manager.Add<transform::PadArrayElements>();
manager.Add<transform::ModuleScopeVarToEntryPointParam>();
manager.Add<transform::InlinePointerLets>();
+ manager.Add<transform::RemovePhonies>();
manager.Add<transform::Simplify>();
// ArrayLengthFromUniform must come after InlinePointerLets and Simplify, as
// it assumes that the form of the array length argument is &var.array.
diff --git a/src/writer/spirv/builder.cc b/src/writer/spirv/builder.cc
index 363378c..bbd6db9 100644
--- a/src/writer/spirv/builder.cc
+++ b/src/writer/spirv/builder.cc
@@ -388,20 +388,28 @@
}
bool Builder::GenerateAssignStatement(const ast::AssignmentStatement* assign) {
- auto lhs_id = GenerateExpression(assign->lhs);
- if (lhs_id == 0) {
- return false;
- }
- auto rhs_id = GenerateExpression(assign->rhs);
- if (rhs_id == 0) {
- return false;
- }
+ if (assign->lhs->Is<ast::PhonyExpression>()) {
+ auto rhs_id = GenerateExpression(assign->rhs);
+ if (rhs_id == 0) {
+ return false;
+ }
+ return true;
+ } else {
+ auto lhs_id = GenerateExpression(assign->lhs);
+ if (lhs_id == 0) {
+ return false;
+ }
+ auto rhs_id = GenerateExpression(assign->rhs);
+ if (rhs_id == 0) {
+ return false;
+ }
- // If the thing we're assigning is a reference then we must load it first.
- auto* type = TypeOf(assign->rhs);
- rhs_id = GenerateLoadIfNeeded(type, rhs_id);
+ // If the thing we're assigning is a reference then we must load it first.
+ auto* type = TypeOf(assign->rhs);
+ rhs_id = GenerateLoadIfNeeded(type, rhs_id);
- return GenerateStore(lhs_id, rhs_id);
+ return GenerateStore(lhs_id, rhs_id);
+ }
}
bool Builder::GenerateBreakStatement(const ast::BreakStatement*) {
diff --git a/src/writer/wgsl/generator_impl.cc b/src/writer/wgsl/generator_impl.cc
index 3c3a06d..c97669f 100644
--- a/src/writer/wgsl/generator_impl.cc
+++ b/src/writer/wgsl/generator_impl.cc
@@ -137,6 +137,10 @@
if (auto* m = expr->As<ast::MemberAccessorExpression>()) {
return EmitMemberAccessor(out, m);
}
+ if (expr->Is<ast::PhonyExpression>()) {
+ out << "_";
+ return true;
+ }
if (auto* u = expr->As<ast::UnaryOpExpression>()) {
return EmitUnaryOp(out, u);
}
diff --git a/test/statements/assign/phony/addr_of_non_constructable.wgsl b/test/statements/assign/phony/addr_of_non_constructable.wgsl
new file mode 100644
index 0000000..dc2b494
--- /dev/null
+++ b/test/statements/assign/phony/addr_of_non_constructable.wgsl
@@ -0,0 +1,11 @@
+[[block]]
+struct S {
+ arr : array<i32>;
+};
+
+[[binding(0), group(0)]] var<storage, read_write> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = &s;
+}
diff --git a/test/statements/assign/phony/addr_of_non_constructable.wgsl.expected.hlsl b/test/statements/assign/phony/addr_of_non_constructable.wgsl.expected.hlsl
new file mode 100644
index 0000000..146b286
--- /dev/null
+++ b/test/statements/assign/phony/addr_of_non_constructable.wgsl.expected.hlsl
@@ -0,0 +1,6 @@
+RWByteAddressBuffer s : register(u0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+ return;
+}
diff --git a/test/statements/assign/phony/addr_of_non_constructable.wgsl.expected.msl b/test/statements/assign/phony/addr_of_non_constructable.wgsl.expected.msl
new file mode 100644
index 0000000..10d815f
--- /dev/null
+++ b/test/statements/assign/phony/addr_of_non_constructable.wgsl.expected.msl
@@ -0,0 +1,11 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+ /* 0x0000 */ int arr[1];
+};
+
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/statements/assign/phony/addr_of_non_constructable.wgsl.expected.spvasm b/test/statements/assign/phony/addr_of_non_constructable.wgsl.expected.spvasm
new file mode 100644
index 0000000..36a0962
--- /dev/null
+++ b/test/statements/assign/phony/addr_of_non_constructable.wgsl.expected.spvasm
@@ -0,0 +1,29 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 11
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %S "S"
+ OpMemberName %S 0 "arr"
+ OpName %s "s"
+ OpName %main "main"
+ OpDecorate %S Block
+ OpMemberDecorate %S 0 Offset 0
+ OpDecorate %_runtimearr_int ArrayStride 4
+ OpDecorate %s Binding 0
+ OpDecorate %s DescriptorSet 0
+ %int = OpTypeInt 32 1
+%_runtimearr_int = OpTypeRuntimeArray %int
+ %S = OpTypeStruct %_runtimearr_int
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+ %s = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+ %void = OpTypeVoid
+ %6 = OpTypeFunction %void
+ %main = OpFunction %void None %6
+ %9 = OpLabel
+ OpReturn
+ OpFunctionEnd
diff --git a/test/statements/assign/phony/addr_of_non_constructable.wgsl.expected.wgsl b/test/statements/assign/phony/addr_of_non_constructable.wgsl.expected.wgsl
new file mode 100644
index 0000000..31c4b7c
--- /dev/null
+++ b/test/statements/assign/phony/addr_of_non_constructable.wgsl.expected.wgsl
@@ -0,0 +1,11 @@
+[[block]]
+struct S {
+ arr : array<i32>;
+};
+
+[[binding(0), group(0)]] var<storage, read_write> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = &(s);
+}
diff --git a/test/statements/assign/phony/addr_of_runtime_array.wgsl b/test/statements/assign/phony/addr_of_runtime_array.wgsl
new file mode 100644
index 0000000..d146bc3
--- /dev/null
+++ b/test/statements/assign/phony/addr_of_runtime_array.wgsl
@@ -0,0 +1,11 @@
+[[block]]
+struct S {
+ arr : array<i32>;
+};
+
+[[binding(0), group(0)]] var<storage, read_write> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = &s.arr;
+}
diff --git a/test/statements/assign/phony/addr_of_runtime_array.wgsl.expected.hlsl b/test/statements/assign/phony/addr_of_runtime_array.wgsl.expected.hlsl
new file mode 100644
index 0000000..146b286
--- /dev/null
+++ b/test/statements/assign/phony/addr_of_runtime_array.wgsl.expected.hlsl
@@ -0,0 +1,6 @@
+RWByteAddressBuffer s : register(u0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+ return;
+}
diff --git a/test/statements/assign/phony/addr_of_runtime_array.wgsl.expected.msl b/test/statements/assign/phony/addr_of_runtime_array.wgsl.expected.msl
new file mode 100644
index 0000000..10d815f
--- /dev/null
+++ b/test/statements/assign/phony/addr_of_runtime_array.wgsl.expected.msl
@@ -0,0 +1,11 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+ /* 0x0000 */ int arr[1];
+};
+
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/statements/assign/phony/addr_of_runtime_array.wgsl.expected.spvasm b/test/statements/assign/phony/addr_of_runtime_array.wgsl.expected.spvasm
new file mode 100644
index 0000000..8a33594
--- /dev/null
+++ b/test/statements/assign/phony/addr_of_runtime_array.wgsl.expected.spvasm
@@ -0,0 +1,33 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 15
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %S "S"
+ OpMemberName %S 0 "arr"
+ OpName %s "s"
+ OpName %main "main"
+ OpDecorate %S Block
+ OpMemberDecorate %S 0 Offset 0
+ OpDecorate %_runtimearr_int ArrayStride 4
+ OpDecorate %s Binding 0
+ OpDecorate %s DescriptorSet 0
+ %int = OpTypeInt 32 1
+%_runtimearr_int = OpTypeRuntimeArray %int
+ %S = OpTypeStruct %_runtimearr_int
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+ %s = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+ %void = OpTypeVoid
+ %6 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer__runtimearr_int = OpTypePointer StorageBuffer %_runtimearr_int
+ %main = OpFunction %void None %6
+ %9 = OpLabel
+ %14 = OpAccessChain %_ptr_StorageBuffer__runtimearr_int %s %uint_0
+ OpReturn
+ OpFunctionEnd
diff --git a/test/statements/assign/phony/addr_of_runtime_array.wgsl.expected.wgsl b/test/statements/assign/phony/addr_of_runtime_array.wgsl.expected.wgsl
new file mode 100644
index 0000000..1bdbbb5
--- /dev/null
+++ b/test/statements/assign/phony/addr_of_runtime_array.wgsl.expected.wgsl
@@ -0,0 +1,11 @@
+[[block]]
+struct S {
+ arr : array<i32>;
+};
+
+[[binding(0), group(0)]] var<storage, read_write> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = &(s.arr);
+}
diff --git a/test/statements/assign/phony/call.wgsl b/test/statements/assign/phony/call.wgsl
new file mode 100644
index 0000000..8be76ec
--- /dev/null
+++ b/test/statements/assign/phony/call.wgsl
@@ -0,0 +1,8 @@
+fn f(a: i32, b: i32, c: i32) -> i32 {
+ return a * b + c;
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = f(1, 2, 3);
+}
diff --git a/test/statements/assign/phony/call.wgsl.expected.hlsl b/test/statements/assign/phony/call.wgsl.expected.hlsl
new file mode 100644
index 0000000..2df8bdf
--- /dev/null
+++ b/test/statements/assign/phony/call.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+int f(int a, int b, int c) {
+ return ((a * b) + c);
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ (void) f(1, 2, 3);
+ return;
+}
diff --git a/test/statements/assign/phony/call.wgsl.expected.msl b/test/statements/assign/phony/call.wgsl.expected.msl
new file mode 100644
index 0000000..dcc3993
--- /dev/null
+++ b/test/statements/assign/phony/call.wgsl.expected.msl
@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int f(int a, int b, int c) {
+ return as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(a) * as_type<uint>(b)))) + as_type<uint>(c)));
+}
+
+kernel void tint_symbol() {
+ f(1, 2, 3);
+ return;
+}
+
diff --git a/test/statements/assign/phony/call.wgsl.expected.spvasm b/test/statements/assign/phony/call.wgsl.expected.spvasm
new file mode 100644
index 0000000..6c0fb76
--- /dev/null
+++ b/test/statements/assign/phony/call.wgsl.expected.spvasm
@@ -0,0 +1,35 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 18
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %f "f"
+ OpName %a "a"
+ OpName %b "b"
+ OpName %c "c"
+ OpName %main "main"
+ %int = OpTypeInt 32 1
+ %1 = OpTypeFunction %int %int %int %int
+ %void = OpTypeVoid
+ %10 = OpTypeFunction %void
+ %int_1 = OpConstant %int 1
+ %int_2 = OpConstant %int 2
+ %int_3 = OpConstant %int 3
+ %f = OpFunction %int None %1
+ %a = OpFunctionParameter %int
+ %b = OpFunctionParameter %int
+ %c = OpFunctionParameter %int
+ %7 = OpLabel
+ %8 = OpIMul %int %a %b
+ %9 = OpIAdd %int %8 %c
+ OpReturnValue %9
+ OpFunctionEnd
+ %main = OpFunction %void None %10
+ %13 = OpLabel
+ %14 = OpFunctionCall %int %f %int_1 %int_2 %int_3
+ OpReturn
+ OpFunctionEnd
diff --git a/test/statements/assign/phony/call.wgsl.expected.wgsl b/test/statements/assign/phony/call.wgsl.expected.wgsl
new file mode 100644
index 0000000..c650359
--- /dev/null
+++ b/test/statements/assign/phony/call.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+fn f(a : i32, b : i32, c : i32) -> i32 {
+ return ((a * b) + c);
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = f(1, 2, 3);
+}
diff --git a/test/statements/assign/phony/multiple_side_effects.wgsl b/test/statements/assign/phony/multiple_side_effects.wgsl
new file mode 100644
index 0000000..fcfccd1
--- /dev/null
+++ b/test/statements/assign/phony/multiple_side_effects.wgsl
@@ -0,0 +1,8 @@
+fn f(a: i32, b: i32, c: i32) -> i32 {
+ return a * b + c;
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = f(1, 2, 3) + f(4, 5, 6) * f(7, f(8, 9, 10), 11);
+}
diff --git a/test/statements/assign/phony/multiple_side_effects.wgsl.expected.hlsl b/test/statements/assign/phony/multiple_side_effects.wgsl.expected.hlsl
new file mode 100644
index 0000000..a295582
--- /dev/null
+++ b/test/statements/assign/phony/multiple_side_effects.wgsl.expected.hlsl
@@ -0,0 +1,12 @@
+int f(int a, int b, int c) {
+ return ((a * b) + c);
+}
+
+void phony_sink(int p0, int p1, int p2) {
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ phony_sink(f(1, 2, 3), f(4, 5, 6), f(7, f(8, 9, 10), 11));
+ return;
+}
diff --git a/test/statements/assign/phony/multiple_side_effects.wgsl.expected.msl b/test/statements/assign/phony/multiple_side_effects.wgsl.expected.msl
new file mode 100644
index 0000000..303cc5a
--- /dev/null
+++ b/test/statements/assign/phony/multiple_side_effects.wgsl.expected.msl
@@ -0,0 +1,15 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int f(int a, int b, int c) {
+ return as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(a) * as_type<uint>(b)))) + as_type<uint>(c)));
+}
+
+void phony_sink(int p0, int p1, int p2) {
+}
+
+kernel void tint_symbol() {
+ phony_sink(f(1, 2, 3), f(4, 5, 6), f(7, f(8, 9, 10), 11));
+ return;
+}
+
diff --git a/test/statements/assign/phony/multiple_side_effects.wgsl.expected.spvasm b/test/statements/assign/phony/multiple_side_effects.wgsl.expected.spvasm
new file mode 100644
index 0000000..2776d06
--- /dev/null
+++ b/test/statements/assign/phony/multiple_side_effects.wgsl.expected.spvasm
@@ -0,0 +1,48 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 31
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %f "f"
+ OpName %a "a"
+ OpName %b "b"
+ OpName %c "c"
+ OpName %main "main"
+ %int = OpTypeInt 32 1
+ %1 = OpTypeFunction %int %int %int %int
+ %void = OpTypeVoid
+ %10 = OpTypeFunction %void
+ %int_1 = OpConstant %int 1
+ %int_2 = OpConstant %int 2
+ %int_3 = OpConstant %int 3
+ %int_4 = OpConstant %int 4
+ %int_5 = OpConstant %int 5
+ %int_6 = OpConstant %int 6
+ %int_7 = OpConstant %int 7
+ %int_8 = OpConstant %int 8
+ %int_9 = OpConstant %int 9
+ %int_10 = OpConstant %int 10
+ %int_11 = OpConstant %int 11
+ %f = OpFunction %int None %1
+ %a = OpFunctionParameter %int
+ %b = OpFunctionParameter %int
+ %c = OpFunctionParameter %int
+ %7 = OpLabel
+ %8 = OpIMul %int %a %b
+ %9 = OpIAdd %int %8 %c
+ OpReturnValue %9
+ OpFunctionEnd
+ %main = OpFunction %void None %10
+ %13 = OpLabel
+ %14 = OpFunctionCall %int %f %int_1 %int_2 %int_3
+ %18 = OpFunctionCall %int %f %int_4 %int_5 %int_6
+ %24 = OpFunctionCall %int %f %int_8 %int_9 %int_10
+ %22 = OpFunctionCall %int %f %int_7 %24 %int_11
+ %29 = OpIMul %int %18 %22
+ %30 = OpIAdd %int %14 %29
+ OpReturn
+ OpFunctionEnd
diff --git a/test/statements/assign/phony/multiple_side_effects.wgsl.expected.wgsl b/test/statements/assign/phony/multiple_side_effects.wgsl.expected.wgsl
new file mode 100644
index 0000000..0c7cc23
--- /dev/null
+++ b/test/statements/assign/phony/multiple_side_effects.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+fn f(a : i32, b : i32, c : i32) -> i32 {
+ return ((a * b) + c);
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = (f(1, 2, 3) + (f(4, 5, 6) * f(7, f(8, 9, 10), 11)));
+}
diff --git a/test/statements/assign/phony/storage_buffer.wgsl b/test/statements/assign/phony/storage_buffer.wgsl
new file mode 100644
index 0000000..b4edd05
--- /dev/null
+++ b/test/statements/assign/phony/storage_buffer.wgsl
@@ -0,0 +1,12 @@
+[[block]]
+struct S {
+ i : i32;
+};
+
+[[binding(0), group(0)]] var<storage, read_write> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = s;
+ _ = s.i;
+}
diff --git a/test/statements/assign/phony/storage_buffer.wgsl.expected.hlsl b/test/statements/assign/phony/storage_buffer.wgsl.expected.hlsl
new file mode 100644
index 0000000..146b286
--- /dev/null
+++ b/test/statements/assign/phony/storage_buffer.wgsl.expected.hlsl
@@ -0,0 +1,6 @@
+RWByteAddressBuffer s : register(u0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+ return;
+}
diff --git a/test/statements/assign/phony/storage_buffer.wgsl.expected.msl b/test/statements/assign/phony/storage_buffer.wgsl.expected.msl
new file mode 100644
index 0000000..6da18d2
--- /dev/null
+++ b/test/statements/assign/phony/storage_buffer.wgsl.expected.msl
@@ -0,0 +1,11 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+ /* 0x0000 */ int i;
+};
+
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/statements/assign/phony/storage_buffer.wgsl.expected.spvasm b/test/statements/assign/phony/storage_buffer.wgsl.expected.spvasm
new file mode 100644
index 0000000..ff08f0d
--- /dev/null
+++ b/test/statements/assign/phony/storage_buffer.wgsl.expected.spvasm
@@ -0,0 +1,31 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 13
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %S "S"
+ OpMemberName %S 0 "i"
+ OpName %s "s"
+ OpName %main "main"
+ OpDecorate %S Block
+ OpMemberDecorate %S 0 Offset 0
+ OpDecorate %s Binding 0
+ OpDecorate %s DescriptorSet 0
+ %int = OpTypeInt 32 1
+ %S = OpTypeStruct %int
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+ %s = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+ %void = OpTypeVoid
+ %5 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+ %main = OpFunction %void None %5
+ %8 = OpLabel
+ %12 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_0
+ OpReturn
+ OpFunctionEnd
diff --git a/test/statements/assign/phony/storage_buffer.wgsl.expected.wgsl b/test/statements/assign/phony/storage_buffer.wgsl.expected.wgsl
new file mode 100644
index 0000000..da77a76
--- /dev/null
+++ b/test/statements/assign/phony/storage_buffer.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+[[block]]
+struct S {
+ i : i32;
+};
+
+[[binding(0), group(0)]] var<storage, read_write> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = s;
+ _ = s.i;
+}
diff --git a/test/statements/assign/phony/uniform_buffer.wgsl b/test/statements/assign/phony/uniform_buffer.wgsl
new file mode 100644
index 0000000..bc59e1b
--- /dev/null
+++ b/test/statements/assign/phony/uniform_buffer.wgsl
@@ -0,0 +1,12 @@
+[[block]]
+struct S {
+ i : i32;
+};
+
+[[binding(0), group(0)]] var<uniform> u : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = u;
+ _ = u.i;
+}
diff --git a/test/statements/assign/phony/uniform_buffer.wgsl.expected.hlsl b/test/statements/assign/phony/uniform_buffer.wgsl.expected.hlsl
new file mode 100644
index 0000000..35f86d8
--- /dev/null
+++ b/test/statements/assign/phony/uniform_buffer.wgsl.expected.hlsl
@@ -0,0 +1,8 @@
+cbuffer cbuffer_u : register(b0, space0) {
+ uint4 u[1];
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+ return;
+}
diff --git a/test/statements/assign/phony/uniform_buffer.wgsl.expected.msl b/test/statements/assign/phony/uniform_buffer.wgsl.expected.msl
new file mode 100644
index 0000000..6da18d2
--- /dev/null
+++ b/test/statements/assign/phony/uniform_buffer.wgsl.expected.msl
@@ -0,0 +1,11 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+ /* 0x0000 */ int i;
+};
+
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/statements/assign/phony/uniform_buffer.wgsl.expected.spvasm b/test/statements/assign/phony/uniform_buffer.wgsl.expected.spvasm
new file mode 100644
index 0000000..aa242e7
--- /dev/null
+++ b/test/statements/assign/phony/uniform_buffer.wgsl.expected.spvasm
@@ -0,0 +1,32 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 13
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %S "S"
+ OpMemberName %S 0 "i"
+ OpName %u "u"
+ OpName %main "main"
+ OpDecorate %S Block
+ OpMemberDecorate %S 0 Offset 0
+ OpDecorate %u NonWritable
+ OpDecorate %u Binding 0
+ OpDecorate %u DescriptorSet 0
+ %int = OpTypeInt 32 1
+ %S = OpTypeStruct %int
+%_ptr_Uniform_S = OpTypePointer Uniform %S
+ %u = OpVariable %_ptr_Uniform_S Uniform
+ %void = OpTypeVoid
+ %5 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_int = OpTypePointer Uniform %int
+ %main = OpFunction %void None %5
+ %8 = OpLabel
+ %12 = OpAccessChain %_ptr_Uniform_int %u %uint_0
+ OpReturn
+ OpFunctionEnd
diff --git a/test/statements/assign/phony/uniform_buffer.wgsl.expected.wgsl b/test/statements/assign/phony/uniform_buffer.wgsl.expected.wgsl
new file mode 100644
index 0000000..15763ac
--- /dev/null
+++ b/test/statements/assign/phony/uniform_buffer.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+[[block]]
+struct S {
+ i : i32;
+};
+
+[[binding(0), group(0)]] var<uniform> u : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = u;
+ _ = u.i;
+}
