tint/transform: Fix FoldTrivialSingleUseLet for abstracts

If the let initializer was an abstract numeric that was implicitly materialized to a concrete type, then we could inlining the initializer into the use without the implicit materialization cast.
This could lead to validation errors, and subtly different results in the program.

In this situation, add an explicit cast to keep the types the same before and after inlining.

Fixed: tint:1664
Change-Id: Icca980cf8af74673906ad6c681a6b07d0c1932fd
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/101160
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>

src/tint/transform/fold_trivial_single_use_lets.cc

@@ -17,6 +17,7 @@
 #include "src/tint/program_builder.h"
 #include "src/tint/sem/block_statement.h"
 #include "src/tint/sem/function.h"
+#include "src/tint/sem/materialize.h"
 #include "src/tint/sem/statement.h"
 #include "src/tint/sem/variable.h"
 #include "src/tint/utils/scoped_assignment.h"
@@ -26,6 +27,8 @@ TINT_INSTANTIATE_TYPEINFO(tint::transform::FoldTrivialSingleUseLets);
 namespace tint::transform {
 namespace {
 
+/// @returns a @p stmt cast to a ast::VariableDeclStatement iff the statement is a let declaration,
+/// with an initializer that's an identifier or literal expression.
 const ast::VariableDeclStatement* AsTrivialLetDecl(const ast::Statement* stmt) {
     auto* var_decl = stmt->As<ast::VariableDeclStatement>();
     if (!var_decl) {
@@ -50,10 +53,12 @@ FoldTrivialSingleUseLets::~FoldTrivialSingleUseLets() = default;
 
 void FoldTrivialSingleUseLets::Run(CloneContext& ctx, const DataMap&, DataMap&) const {
     for (auto* node : ctx.src->ASTNodes().Objects()) {
+        // For each statement in each block of the entire program...
         if (auto* block = node->As<ast::BlockStatement>()) {
             auto& stmts = block->statements;
             for (size_t stmt_idx = 0; stmt_idx < stmts.Length(); stmt_idx++) {
                 auto* stmt = stmts[stmt_idx];
+                // Is the statement a trivial let declaration with a single user?
                 if (auto* let_decl = AsTrivialLetDecl(stmt)) {
                     auto* let = let_decl->variable;
                     auto* sem_let = ctx.src->Sem().Get(let);
@@ -65,11 +70,23 @@ void FoldTrivialSingleUseLets::Run(CloneContext& ctx, const DataMap&, DataMap&)
                     auto* user = users[0];
                     auto* user_stmt = user->Stmt()->Declaration();
 
+                    // Scan forward to find the statement of use. If there's a statement between the
+                    // declaration and the use, then we cannot safely fold.
                     for (size_t i = stmt_idx; i < stmts.Length(); i++) {
                         if (user_stmt == stmts[i]) {
                             auto* user_expr = user->Declaration();
                             ctx.Remove(stmts, let_decl);
-                            ctx.Replace(user_expr, ctx.Clone(let->constructor));
+                            auto* initializer = ctx.Clone(let->constructor);
+                            if (auto* materialize =
+                                    sem_let->Constructor()->As<sem::Materialize>()) {
+                                // The let initializer was an abstract numeric that was implicitly
+                                // materialized to a concrete type. As we're inlining the
+                                // initializer into the use, we need to make this cast explicit,
+                                // otherwise we'll have a different type for the substitution.
+                                auto* concrete_ty = CreateASTTypeFor(ctx, materialize->Type());
+                                initializer = ctx.dst->Construct(concrete_ty, initializer);
+                            }
+                            ctx.Replace(user_expr, initializer);
                         }
                         if (!AsTrivialLetDecl(stmts[i])) {
                             // Stop if we hit a statement that isn't the single use of the

src/tint/transform/fold_trivial_single_use_lets_test.cc

@@ -30,7 +30,26 @@ TEST_F(FoldTrivialSingleUseLetsTest, EmptyModule) {
     EXPECT_EQ(expect, str(got));
 }
 
-TEST_F(FoldTrivialSingleUseLetsTest, Single) {
+TEST_F(FoldTrivialSingleUseLetsTest, Single_Concrete) {
+    auto* src = R"(
+fn f() {
+  let x = 1i;
+  _ = x;
+}
+)";
+
+    auto* expect = R"(
+fn f() {
+  _ = 1i;
+}
+)";
+
+    auto got = Run<FoldTrivialSingleUseLets>(src);
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(FoldTrivialSingleUseLetsTest, Single_Abstract) {
     auto* src = R"(
 fn f() {
   let x = 1;
@@ -40,7 +59,7 @@ fn f() {
 
     auto* expect = R"(
 fn f() {
-  _ = 1;
+  _ = i32(1);
 }
 )";
 
@@ -49,7 +68,28 @@ fn f() {
     EXPECT_EQ(expect, str(got));
 }
 
-TEST_F(FoldTrivialSingleUseLetsTest, Multiple) {
+TEST_F(FoldTrivialSingleUseLetsTest, Multiple_Concrete) {
+    auto* src = R"(
+fn f() {
+  let x = 1u;
+  let y = 2u;
+  let z = 3u;
+  _ = x + y + z;
+}
+)";
+
+    auto* expect = R"(
+fn f() {
+  _ = ((1u + 2u) + 3u);
+}
+)";
+
+    auto got = Run<FoldTrivialSingleUseLets>(src);
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(FoldTrivialSingleUseLetsTest, Multiple_Abstract) {
     auto* src = R"(
 fn f() {
   let x = 1;
@@ -61,7 +101,7 @@ fn f() {
 
     auto* expect = R"(
 fn f() {
-  _ = ((1 + 2) + 3);
+  _ = ((i32(1) + i32(2)) + i32(3));
 }
 )";
 
@@ -70,7 +110,28 @@ fn f() {
     EXPECT_EQ(expect, str(got));
 }
 
-TEST_F(FoldTrivialSingleUseLetsTest, Chained) {
+TEST_F(FoldTrivialSingleUseLetsTest, Chained_Concrete) {
+    auto* src = R"(
+fn f() {
+  let x = 1i;
+  let y = x;
+  let z = y;
+  _ = z;
+}
+)";
+
+    auto* expect = R"(
+fn f() {
+  _ = 1i;
+}
+)";
+
+    auto got = Run<FoldTrivialSingleUseLets>(src);
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(FoldTrivialSingleUseLetsTest, Chained_Abstract) {
     auto* src = R"(
 fn f() {
   let x = 1;
@@ -82,7 +143,7 @@ fn f() {
 
     auto* expect = R"(
 fn f() {
-  _ = 1;
+  _ = i32(1);
 }
 )";
 
@@ -93,13 +154,13 @@ fn f() {
 
 TEST_F(FoldTrivialSingleUseLetsTest, NoFold_NonTrivialLet) {
     auto* src = R"(
-fn function_with_posssible_side_effect() -> i32 {
+fn function_with_possible_side_effect() -> i32 {
   return 1;
 }
 
 fn f() {
   let x = 1;
-  let y = function_with_posssible_side_effect();
+  let y = function_with_possible_side_effect();
   _ = (x + y);
 }
 )";
@@ -115,11 +176,11 @@ TEST_F(FoldTrivialSingleUseLetsTest, NoFold_NonTrivialLet_OutOfOrder) {
     auto* src = R"(
 fn f() {
   let x = 1;
-  let y = function_with_posssible_side_effect();
+  let y = function_with_possible_side_effect();
   _ = (x + y);
 }
 
-fn function_with_posssible_side_effect() -> i32 {
+fn function_with_possible_side_effect() -> i32 {
   return 1;
 }
 )";

test/tint/bug/tint/1664.wgsl

@@ -0,0 +1,11 @@
+@compute @workgroup_size(1)
+fn f0() {
+    let a = 2147483647;
+    let b = 1;
+    let c = a + 1;
+}
+
+fn f1() {
+    let a = 1;
+    let b = -2147483648 - a;
+}

test/tint/bug/tint/1664.wgsl.expected.dxc.hlsl

@@ -0,0 +1,10 @@
+[numthreads(1, 1, 1)]
+void f0() {
+  const int b = 1;
+  const int c = -2147483648;
+  return;
+}
+
+void f1() {
+  const int b = 2147483647;
+}

test/tint/bug/tint/1664.wgsl.expected.fxc.hlsl

@@ -0,0 +1,10 @@
+[numthreads(1, 1, 1)]
+void f0() {
+  const int b = 1;
+  const int c = -2147483648;
+  return;
+}
+
+void f1() {
+  const int b = 2147483647;
+}

test/tint/bug/tint/1664.wgsl.expected.glsl

@@ -0,0 +1,12 @@
+#version 310 es
+
+void f0() {
+  int b = 1;
+  int c = -2147483648;
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  f0();
+  return;
+}

test/tint/bug/tint/1664.wgsl.expected.msl

@@ -0,0 +1,15 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f0() {
+  int const a = 2147483647;
+  int const b = 1;
+  int const c = as_type<int>((as_type<uint>(a) + as_type<uint>(1)));
+  return;
+}
+
+void f1() {
+  int const a = 1;
+  int const b = as_type<int>((as_type<uint>((-2147483647 - 1)) - as_type<uint>(a)));
+}
+

test/tint/bug/tint/1664.wgsl.expected.spvasm

@@ -0,0 +1,27 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 13
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f0 "f0"
+               OpExecutionMode %f0 LocalSize 1 1 1
+               OpName %f0 "f0"
+               OpName %f1 "f1"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+%int_2147483647 = OpConstant %int 2147483647
+      %int_1 = OpConstant %int 1
+%int_n2147483648 = OpConstant %int -2147483648
+         %f0 = OpFunction %void None %1
+          %4 = OpLabel
+          %8 = OpIAdd %int %int_2147483647 %int_1
+               OpReturn
+               OpFunctionEnd
+         %f1 = OpFunction %void None %1
+         %10 = OpLabel
+         %12 = OpISub %int %int_n2147483648 %int_1
+               OpReturn
+               OpFunctionEnd

test/tint/bug/tint/1664.wgsl.expected.wgsl

@@ -0,0 +1,11 @@
+@compute @workgroup_size(1)
+fn f0() {
+  let a = 2147483647;
+  let b = 1;
+  let c = (a + 1);
+}
+
+fn f1() {
+  let a = 1;
+  let b = (-2147483648 - a);
+}