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>

src/BUILD.gn

@@ -463,6 +463,8 @@ libtint_source_set("libtint_core_all_src") {
     "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",

src/CMakeLists.txt

@@ -329,6 +329,8 @@ set(TINT_LIB_SRCS
   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 @@ if(${TINT_BUILD_TESTS})
       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

src/program_builder.h

@@ -1762,6 +1762,18 @@ class ProgramBuilder {
                                                 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`

src/reader/wgsl/parser_impl.cc

@@ -2792,7 +2792,7 @@ Maybe<const ast::Expression*> ParserImpl::logical_or_expression() {
 }
 
 // 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 @@ Maybe<const ast::AssignmentStatement*> ParserImpl::assignment_stmt() {
   }
 
   auto lhs = unary_expression();
-  if (lhs.errored)
+  if (lhs.errored) {
     return Failure::kErrored;
-  if (!lhs.matched)
+  }
+  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);
 }

src/reader/wgsl/parser_impl_assignment_stmt_test.cc

@@ -98,6 +98,28 @@ TEST_F(ParserImplTest, AssignmentStmt_Parses_ToMember) {
   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();

src/resolver/assignment_validation_test.cc

@@ -293,6 +293,111 @@ TEST_F(ResolverAssignmentValidationTest, AssignNonConstructible_RuntimeArray) {
             "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

src/resolver/resolver.cc

@@ -2303,6 +2303,8 @@ bool Resolver::Expression(const ast::Expression* root) {
       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 @@ bool Resolver::Assignment(const ast::AssignmentStatement* a) {
   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;

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 @@ Output Glsl::Run(const Program* in, const DataMap& inputs) {
   }
   manager.Add<CanonicalizeEntryPointIO>();
   manager.Add<InlinePointerLets>();
+  manager.Add<RemovePhonies>();
   // Simplify cleans up messy `*(&(expr))` expressions from InlinePointerLets.
   manager.Add<Simplify>();
   manager.Add<CalculateArrayLength>();

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

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_

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

src/transform/transform.cc

@@ -75,7 +75,7 @@ bool Transform::Requires(CloneContext& ctx,
   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);

src/transform/transform.h

@@ -190,7 +190,7 @@ class Transform : public Castable<Transform> {
   /// 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`.

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 @@ SanitizedResult Sanitize(const Program* in,
   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
+  manager.Add<transform::RemovePhonies>();
-  // the address of calls to DecomposeMemoryAccess::Intrinsic. Must also come
+  // DecomposeMemoryAccess must come after:
-  // after Simplify, as we need to fold away the address-of and defers of
+  // * 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

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 @@ SanitizedResult Sanitize(const Program* in,
   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.

src/writer/spirv/builder.cc

@@ -388,6 +388,13 @@ bool Builder::GenerateLabel(uint32_t id) {
 }
 
 bool Builder::GenerateAssignStatement(const ast::AssignmentStatement* assign) {
+  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;
@@ -403,6 +410,7 @@ bool Builder::GenerateAssignStatement(const ast::AssignmentStatement* assign) {
 
     return GenerateStore(lhs_id, rhs_id);
   }
+}
 
 bool Builder::GenerateBreakStatement(const ast::BreakStatement*) {
   if (merge_stack_.empty()) {

src/writer/wgsl/generator_impl.cc

@@ -137,6 +137,10 @@ bool GeneratorImpl::EmitExpression(std::ostream& out,
   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);
   }

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

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

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

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

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);
+}

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

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

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

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

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);
+}

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);
+}

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

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

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

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);
+}

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);
+}

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

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

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

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)));
+}

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

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

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

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

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

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

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

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

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

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