From 8b30ca3efa5e35ed735c6cfa2b90ef5e0528c840 Mon Sep 17 00:00:00 2001 From: Ben Clayton Date: Fri, 2 Sep 2022 13:58:00 +0000 Subject: [PATCH] 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 Commit-Queue: Ben Clayton Kokoro: Kokoro --- .../transform/fold_trivial_single_use_lets.cc | 19 ++++- .../fold_trivial_single_use_lets_test.cc | 81 ++++++++++++++++--- test/tint/bug/tint/1664.wgsl | 11 +++ .../tint/bug/tint/1664.wgsl.expected.dxc.hlsl | 10 +++ .../tint/bug/tint/1664.wgsl.expected.fxc.hlsl | 10 +++ test/tint/bug/tint/1664.wgsl.expected.glsl | 12 +++ test/tint/bug/tint/1664.wgsl.expected.msl | 15 ++++ test/tint/bug/tint/1664.wgsl.expected.spvasm | 27 +++++++ test/tint/bug/tint/1664.wgsl.expected.wgsl | 11 +++ 9 files changed, 185 insertions(+), 11 deletions(-) create mode 100644 test/tint/bug/tint/1664.wgsl create mode 100644 test/tint/bug/tint/1664.wgsl.expected.dxc.hlsl create mode 100644 test/tint/bug/tint/1664.wgsl.expected.fxc.hlsl create mode 100644 test/tint/bug/tint/1664.wgsl.expected.glsl create mode 100644 test/tint/bug/tint/1664.wgsl.expected.msl create mode 100644 test/tint/bug/tint/1664.wgsl.expected.spvasm create mode 100644 test/tint/bug/tint/1664.wgsl.expected.wgsl diff --git a/src/tint/transform/fold_trivial_single_use_lets.cc b/src/tint/transform/fold_trivial_single_use_lets.cc index 8a949ecc6c..4d6a2a587a 100644 --- a/src/tint/transform/fold_trivial_single_use_lets.cc +++ b/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(); 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()) { 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()) { + // 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 diff --git a/src/tint/transform/fold_trivial_single_use_lets_test.cc b/src/tint/transform/fold_trivial_single_use_lets_test.cc index 00159e9e93..80e9841a1e 100644 --- a/src/tint/transform/fold_trivial_single_use_lets_test.cc +++ b/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(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(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(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; } )"; diff --git a/test/tint/bug/tint/1664.wgsl b/test/tint/bug/tint/1664.wgsl new file mode 100644 index 0000000000..eb5d3af673 --- /dev/null +++ b/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; +} diff --git a/test/tint/bug/tint/1664.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/1664.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..f80897a882 --- /dev/null +++ b/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; +} diff --git a/test/tint/bug/tint/1664.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/1664.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..f80897a882 --- /dev/null +++ b/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; +} diff --git a/test/tint/bug/tint/1664.wgsl.expected.glsl b/test/tint/bug/tint/1664.wgsl.expected.glsl new file mode 100644 index 0000000000..51da551324 --- /dev/null +++ b/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; +} diff --git a/test/tint/bug/tint/1664.wgsl.expected.msl b/test/tint/bug/tint/1664.wgsl.expected.msl new file mode 100644 index 0000000000..a915cc8d10 --- /dev/null +++ b/test/tint/bug/tint/1664.wgsl.expected.msl @@ -0,0 +1,15 @@ +#include + +using namespace metal; +kernel void f0() { + int const a = 2147483647; + int const b = 1; + int const c = as_type((as_type(a) + as_type(1))); + return; +} + +void f1() { + int const a = 1; + int const b = as_type((as_type((-2147483647 - 1)) - as_type(a))); +} + diff --git a/test/tint/bug/tint/1664.wgsl.expected.spvasm b/test/tint/bug/tint/1664.wgsl.expected.spvasm new file mode 100644 index 0000000000..f95ebadad2 --- /dev/null +++ b/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 diff --git a/test/tint/bug/tint/1664.wgsl.expected.wgsl b/test/tint/bug/tint/1664.wgsl.expected.wgsl new file mode 100644 index 0000000000..a6442d71af --- /dev/null +++ b/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); +}