From b05e185a3633846e599c1c588fb7bbe367827585 Mon Sep 17 00:00:00 2001 From: Ben Clayton Date: Tue, 23 Nov 2021 20:45:51 +0000 Subject: [PATCH] resolver: Support shadowing Add transform::Unshadow to renamed shadowed symbols. Required by a number of other transforms. Replace Resolver symbol resolution with dep-graph. The dependency graph now performs full symbol resolution before the regular resolver pass. Make use of this instead of duplicating the effort. Simplfies code, and actually performs variable shadowing consistently. Fixed: tint:819 Bug: tint:1266 Change-Id: I595d1812aebe1d79d2d32e724ff90de36e74cf4b Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/70523 Reviewed-by: David Neto Kokoro: Kokoro --- docs/origin-trial-changes.md | 1 + src/BUILD.gn | 2 + src/CMakeLists.txt | 5 +- src/reader/spirv/parser.cc | 2 + src/resolver/dependency_graph.cc | 6 + src/resolver/dependency_graph.h | 6 + src/resolver/dependency_graph_test.cc | 58 +++ src/resolver/function_validation_test.cc | 65 ++- src/resolver/resolver.cc | 128 +++--- src/resolver/resolver.h | 32 +- src/resolver/resolver_validation.cc | 87 +--- src/resolver/validation_test.cc | 78 ++-- src/resolver/var_let_test.cc | 388 ++++++++++++++++- src/resolver/var_let_validation_test.cc | 15 +- src/sem/type_mappings.h | 4 +- src/sem/variable.cc | 8 +- src/sem/variable.h | 24 ++ .../array_length_from_uniform_test.cc | 15 +- src/transform/calculate_array_length.cc | 6 +- src/transform/calculate_array_length_test.cc | 80 +++- src/transform/canonicalize_entry_point_io.cc | 5 + .../canonicalize_entry_point_io_test.cc | 97 +++-- .../decompose_strided_matrix_test.cc | 43 +- .../fold_trivial_single_use_lets_test.cc | 37 +- src/transform/glsl.cc | 3 + .../num_workgroups_from_uniform_test.cc | 19 +- src/transform/simplify_pointers.cc | 4 + src/transform/simplify_pointers.h | 6 +- src/transform/simplify_pointers_test.cc | 71 ++-- src/transform/unshadow.cc | 99 +++++ src/transform/unshadow.h | 48 +++ src/transform/unshadow_test.cc | 397 ++++++++++++++++++ src/writer/hlsl/generator_impl.cc | 3 + src/writer/msl/generator_impl.cc | 3 + src/writer/spirv/builder.cc | 2 + test.wgsl | 15 + test/BUILD.gn | 1 + test/bug/tint/825.wgsl.expected.hlsl | 2 +- test/bug/tint/825.wgsl.expected.msl | 2 +- test/shadowing/alias/let.wgsl | 10 + test/shadowing/alias/let.wgsl.expected.hlsl | 13 + test/shadowing/alias/let.wgsl.expected.msl | 13 + test/shadowing/alias/let.wgsl.expected.spvasm | 23 + test/shadowing/alias/let.wgsl.expected.wgsl | 10 + test/shadowing/alias/param.wgsl | 5 + test/shadowing/alias/param.wgsl.expected.hlsl | 8 + test/shadowing/alias/param.wgsl.expected.msl | 8 + .../alias/param.wgsl.expected.spvasm | 25 ++ test/shadowing/alias/param.wgsl.expected.wgsl | 5 + test/shadowing/alias/var.wgsl | 10 + test/shadowing/alias/var.wgsl.expected.hlsl | 13 + test/shadowing/alias/var.wgsl.expected.msl | 13 + test/shadowing/alias/var.wgsl.expected.spvasm | 38 ++ test/shadowing/alias/var.wgsl.expected.wgsl | 10 + test/shadowing/function/let.wgsl | 8 + .../shadowing/function/let.wgsl.expected.hlsl | 12 + test/shadowing/function/let.wgsl.expected.msl | 12 + .../function/let.wgsl.expected.spvasm | 32 ++ .../shadowing/function/let.wgsl.expected.wgsl | 8 + test/shadowing/function/param.wgsl | 3 + .../function/param.wgsl.expected.hlsl | 8 + .../function/param.wgsl.expected.msl | 7 + .../function/param.wgsl.expected.spvasm | 25 ++ .../function/param.wgsl.expected.wgsl | 3 + test/shadowing/function/var.wgsl | 10 + .../shadowing/function/var.wgsl.expected.hlsl | 17 + test/shadowing/function/var.wgsl.expected.msl | 16 + .../function/var.wgsl.expected.spvasm | 42 ++ .../shadowing/function/var.wgsl.expected.wgsl | 12 + test/shadowing/param/function.wgsl | 6 + .../param/function.wgsl.expected.hlsl | 11 + .../param/function.wgsl.expected.msl | 10 + .../param/function.wgsl.expected.spvasm | 34 ++ .../param/function.wgsl.expected.wgsl | 6 + test/shadowing/param/let.wgsl | 6 + test/shadowing/param/let.wgsl.expected.hlsl | 10 + test/shadowing/param/let.wgsl.expected.msl | 10 + test/shadowing/param/let.wgsl.expected.spvasm | 25 ++ test/shadowing/param/let.wgsl.expected.wgsl | 6 + test/shadowing/param/var.wgsl | 6 + test/shadowing/param/var.wgsl.expected.hlsl | 11 + test/shadowing/param/var.wgsl.expected.msl | 10 + test/shadowing/param/var.wgsl.expected.spvasm | 34 ++ test/shadowing/param/var.wgsl.expected.wgsl | 6 + test/shadowing/struct/let.wgsl | 10 + test/shadowing/struct/let.wgsl.expected.hlsl | 17 + test/shadowing/struct/let.wgsl.expected.msl | 16 + .../shadowing/struct/let.wgsl.expected.spvasm | 27 ++ test/shadowing/struct/let.wgsl.expected.wgsl | 12 + test/shadowing/struct/param.wgsl | 5 + .../shadowing/struct/param.wgsl.expected.hlsl | 12 + test/shadowing/struct/param.wgsl.expected.msl | 11 + .../struct/param.wgsl.expected.spvasm | 29 ++ .../shadowing/struct/param.wgsl.expected.wgsl | 7 + test/shadowing/struct/var.wgsl | 10 + test/shadowing/struct/var.wgsl.expected.hlsl | 17 + test/shadowing/struct/var.wgsl.expected.msl | 16 + .../shadowing/struct/var.wgsl.expected.spvasm | 42 ++ test/shadowing/struct/var.wgsl.expected.wgsl | 12 + 99 files changed, 2357 insertions(+), 363 deletions(-) create mode 100644 src/transform/unshadow.cc create mode 100644 src/transform/unshadow.h create mode 100644 src/transform/unshadow_test.cc create mode 100644 test.wgsl create mode 100644 test/shadowing/alias/let.wgsl create mode 100644 test/shadowing/alias/let.wgsl.expected.hlsl create mode 100644 test/shadowing/alias/let.wgsl.expected.msl create mode 100644 test/shadowing/alias/let.wgsl.expected.spvasm create mode 100644 test/shadowing/alias/let.wgsl.expected.wgsl create mode 100644 test/shadowing/alias/param.wgsl create mode 100644 test/shadowing/alias/param.wgsl.expected.hlsl create mode 100644 test/shadowing/alias/param.wgsl.expected.msl create mode 100644 test/shadowing/alias/param.wgsl.expected.spvasm create mode 100644 test/shadowing/alias/param.wgsl.expected.wgsl create mode 100644 test/shadowing/alias/var.wgsl create mode 100644 test/shadowing/alias/var.wgsl.expected.hlsl create mode 100644 test/shadowing/alias/var.wgsl.expected.msl create mode 100644 test/shadowing/alias/var.wgsl.expected.spvasm create mode 100644 test/shadowing/alias/var.wgsl.expected.wgsl create mode 100644 test/shadowing/function/let.wgsl create mode 100644 test/shadowing/function/let.wgsl.expected.hlsl create mode 100644 test/shadowing/function/let.wgsl.expected.msl create mode 100644 test/shadowing/function/let.wgsl.expected.spvasm create mode 100644 test/shadowing/function/let.wgsl.expected.wgsl create mode 100644 test/shadowing/function/param.wgsl create mode 100644 test/shadowing/function/param.wgsl.expected.hlsl create mode 100644 test/shadowing/function/param.wgsl.expected.msl create mode 100644 test/shadowing/function/param.wgsl.expected.spvasm create mode 100644 test/shadowing/function/param.wgsl.expected.wgsl create mode 100644 test/shadowing/function/var.wgsl create mode 100644 test/shadowing/function/var.wgsl.expected.hlsl create mode 100644 test/shadowing/function/var.wgsl.expected.msl create mode 100644 test/shadowing/function/var.wgsl.expected.spvasm create mode 100644 test/shadowing/function/var.wgsl.expected.wgsl create mode 100644 test/shadowing/param/function.wgsl create mode 100644 test/shadowing/param/function.wgsl.expected.hlsl create mode 100644 test/shadowing/param/function.wgsl.expected.msl create mode 100644 test/shadowing/param/function.wgsl.expected.spvasm create mode 100644 test/shadowing/param/function.wgsl.expected.wgsl create mode 100644 test/shadowing/param/let.wgsl create mode 100644 test/shadowing/param/let.wgsl.expected.hlsl create mode 100644 test/shadowing/param/let.wgsl.expected.msl create mode 100644 test/shadowing/param/let.wgsl.expected.spvasm create mode 100644 test/shadowing/param/let.wgsl.expected.wgsl create mode 100644 test/shadowing/param/var.wgsl create mode 100644 test/shadowing/param/var.wgsl.expected.hlsl create mode 100644 test/shadowing/param/var.wgsl.expected.msl create mode 100644 test/shadowing/param/var.wgsl.expected.spvasm create mode 100644 test/shadowing/param/var.wgsl.expected.wgsl create mode 100644 test/shadowing/struct/let.wgsl create mode 100644 test/shadowing/struct/let.wgsl.expected.hlsl create mode 100644 test/shadowing/struct/let.wgsl.expected.msl create mode 100644 test/shadowing/struct/let.wgsl.expected.spvasm create mode 100644 test/shadowing/struct/let.wgsl.expected.wgsl create mode 100644 test/shadowing/struct/param.wgsl create mode 100644 test/shadowing/struct/param.wgsl.expected.hlsl create mode 100644 test/shadowing/struct/param.wgsl.expected.msl create mode 100644 test/shadowing/struct/param.wgsl.expected.spvasm create mode 100644 test/shadowing/struct/param.wgsl.expected.wgsl create mode 100644 test/shadowing/struct/var.wgsl create mode 100644 test/shadowing/struct/var.wgsl.expected.hlsl create mode 100644 test/shadowing/struct/var.wgsl.expected.msl create mode 100644 test/shadowing/struct/var.wgsl.expected.spvasm create mode 100644 test/shadowing/struct/var.wgsl.expected.wgsl diff --git a/docs/origin-trial-changes.md b/docs/origin-trial-changes.md index d5d69f115f..3437a84c7b 100644 --- a/docs/origin-trial-changes.md +++ b/docs/origin-trial-changes.md @@ -12,6 +12,7 @@ ### New Features +* Shadowing is now fully supported. [tint:819](https://crbug.com/tint/819) * The `dot()` builtin now supports integer vector types. * Identifiers can now start with a single leading underscore. [tint:1292](https://crbug.com/tint/1292) diff --git a/src/BUILD.gn b/src/BUILD.gn index b225b19dfe..2fd70f53ef 100644 --- a/src/BUILD.gn +++ b/src/BUILD.gn @@ -474,6 +474,8 @@ libtint_source_set("libtint_core_all_src") { "transform/single_entry_point.h", "transform/transform.cc", "transform/transform.h", + "transform/unshadow.cc", + "transform/unshadow.h", "transform/vectorize_scalar_matrix_constructors.cc", "transform/vectorize_scalar_matrix_constructors.h", "transform/vertex_pulling.cc", diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 8d098349ff..0e8c98eeda 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -338,6 +338,8 @@ set(TINT_LIB_SRCS transform/single_entry_point.h transform/transform.cc transform/transform.h + transform/unshadow.cc + transform/unshadow.h transform/vectorize_scalar_matrix_constructors.cc transform/vectorize_scalar_matrix_constructors.h transform/vertex_pulling.cc @@ -969,7 +971,8 @@ if(${TINT_BUILD_TESTS}) transform/simplify_pointers_test.cc transform/single_entry_point_test.cc transform/test_helper.h - transform/vectorize_scalar_matrix_constructors.cc + transform/unshadow_test.cc + transform/vectorize_scalar_matrix_constructors_test.cc transform/vertex_pulling_test.cc transform/wrap_arrays_in_structs_test.cc transform/zero_init_workgroup_memory_test.cc diff --git a/src/reader/spirv/parser.cc b/src/reader/spirv/parser.cc index c5ffebbbe3..6500966d9d 100644 --- a/src/reader/spirv/parser.cc +++ b/src/reader/spirv/parser.cc @@ -20,6 +20,7 @@ #include "src/transform/decompose_strided_matrix.h" #include "src/transform/manager.h" #include "src/transform/simplify_pointers.h" +#include "src/transform/unshadow.h" namespace tint { namespace reader { @@ -52,6 +53,7 @@ Program Parse(const std::vector& input) { // attribute then we need to decompose these into an array of vectors if (transform::DecomposeStridedMatrix::ShouldRun(&program)) { transform::Manager manager; + manager.Add(); manager.Add(); manager.Add(); return manager.Run(&program).program; diff --git a/src/resolver/dependency_graph.cc b/src/resolver/dependency_graph.cc index 2de794fce9..4cfb45dda3 100644 --- a/src/resolver/dependency_graph.cc +++ b/src/resolver/dependency_graph.cc @@ -177,6 +177,9 @@ class DependencyScanner { TINT_DEFER(scope_stack_.Pop()); for (auto* param : func->params) { + if (auto* shadows = scope_stack_.Get(param->symbol)) { + graph_.shadows.emplace(param, shadows); + } Declare(param->symbol, param); TraverseType(param->type); } @@ -255,6 +258,9 @@ class DependencyScanner { return; } if (auto* v = stmt->As()) { + if (auto* shadows = scope_stack_.Get(v->variable->symbol)) { + graph_.shadows.emplace(v->variable, shadows); + } TraverseType(v->variable->type); TraverseExpression(v->variable->constructor); Declare(v->variable->symbol, v->variable); diff --git a/src/resolver/dependency_graph.h b/src/resolver/dependency_graph.h index b97e3fdfaf..6235ad74cb 100644 --- a/src/resolver/dependency_graph.h +++ b/src/resolver/dependency_graph.h @@ -55,6 +55,12 @@ struct DependencyGraph { /// Map of ast::IdentifierExpression or ast::TypeName to a type, function, or /// variable that declares the symbol. std::unordered_map resolved_symbols; + + /// Map of ast::Variable to a type, function, or variable that is shadowed by + /// the variable key. A declaration (X) shadows another (Y) if X and Y use + /// the same symbol, and X is declared in a sub-scope of the scope that + /// declares Y. + std::unordered_map shadows; }; } // namespace resolver diff --git a/src/resolver/dependency_graph_test.cc b/src/resolver/dependency_graph_test.cc index d256197289..e222b4382d 100644 --- a/src/resolver/dependency_graph_test.cc +++ b/src/resolver/dependency_graph_test.cc @@ -88,6 +88,17 @@ static constexpr SymbolDeclKind kValueDeclKinds[] = { SymbolDeclKind::NestedLocalLet, }; +static constexpr SymbolDeclKind kGlobalDeclKinds[] = { + SymbolDeclKind::GlobalVar, SymbolDeclKind::GlobalLet, SymbolDeclKind::Alias, + SymbolDeclKind::Struct, SymbolDeclKind::Function, +}; + +static constexpr SymbolDeclKind kLocalDeclKinds[] = { + SymbolDeclKind::Parameter, SymbolDeclKind::LocalVar, + SymbolDeclKind::LocalLet, SymbolDeclKind::NestedLocalVar, + SymbolDeclKind::NestedLocalLet, +}; + static constexpr SymbolDeclKind kGlobalValueDeclKinds[] = { SymbolDeclKind::GlobalVar, SymbolDeclKind::GlobalLet, @@ -1171,6 +1182,53 @@ INSTANTIATE_TEST_SUITE_P(Functions, } // namespace resolved_symbols +//////////////////////////////////////////////////////////////////////////////// +// Shadowing tests +//////////////////////////////////////////////////////////////////////////////// +namespace shadowing { + +using ResolverDependencyShadowTest = ResolverDependencyGraphTestWithParam< + std::tuple>; + +TEST_P(ResolverDependencyShadowTest, Test) { + const Symbol symbol = Sym("SYMBOL"); + const auto outer_kind = std::get<0>(GetParam()); + const auto inner_kind = std::get<1>(GetParam()); + + // Build a symbol declaration and a use of that symbol + SymbolTestHelper helper(this); + auto* outer = helper.Add(outer_kind, symbol, Source{{12, 34}}); + helper.Add(inner_kind, symbol, Source{{56, 78}}); + auto* inner_var = helper.nested_statements.size() + ? helper.nested_statements[0] + ->As() + ->variable + : helper.statements.size() + ? helper.statements[0] + ->As() + ->variable + : helper.parameters[0]; + helper.Build(); + + EXPECT_EQ(Build().shadows[inner_var], outer); +} + +INSTANTIATE_TEST_SUITE_P(LocalShadowGlobal, + ResolverDependencyShadowTest, + testing::Combine(testing::ValuesIn(kGlobalDeclKinds), + testing::ValuesIn(kLocalDeclKinds))); + +INSTANTIATE_TEST_SUITE_P( + NestedLocalShadowLocal, + ResolverDependencyShadowTest, + testing::Combine(testing::Values(SymbolDeclKind::Parameter, + SymbolDeclKind::LocalVar, + SymbolDeclKind::LocalLet), + testing::Values(SymbolDeclKind::NestedLocalVar, + SymbolDeclKind::NestedLocalLet))); + +} // namespace shadowing + //////////////////////////////////////////////////////////////////////////////// // AST traversal tests //////////////////////////////////////////////////////////////////////////////// diff --git a/src/resolver/function_validation_test.cc b/src/resolver/function_validation_test.cc index 09dab00592..af3b0975b1 100644 --- a/src/resolver/function_validation_test.cc +++ b/src/resolver/function_validation_test.cc @@ -26,25 +26,46 @@ namespace { class ResolverFunctionValidationTest : public resolver::TestHelper, public testing::Test {}; -TEST_F(ResolverFunctionValidationTest, ParameterNamesMustBeUnique_pass) { +TEST_F(ResolverFunctionValidationTest, DuplicateParameterName) { // fn func_a(common_name : f32) { } // fn func_b(common_name : f32) { } Func("func_a", {Param("common_name", ty.f32())}, ty.void_(), {}); Func("func_b", {Param("common_name", ty.f32())}, ty.void_(), {}); - EXPECT_TRUE(r()->Resolve()); - EXPECT_EQ(r()->error(), ""); + ASSERT_TRUE(r()->Resolve()) << r()->error(); } -TEST_F(ResolverFunctionValidationTest, - ParameterNamesMustBeUniqueShadowsGlobal_pass) { +TEST_F(ResolverFunctionValidationTest, ParameterMayShadowGlobal) { // var common_name : f32; // fn func(common_name : f32) { } Global("common_name", ty.f32(), ast::StorageClass::kPrivate); Func("func", {Param("common_name", ty.f32())}, ty.void_(), {}); - EXPECT_TRUE(r()->Resolve()); - EXPECT_EQ(r()->error(), ""); + ASSERT_TRUE(r()->Resolve()) << r()->error(); +} + +TEST_F(ResolverFunctionValidationTest, LocalConflictsWithParameter) { + // fn func(common_name : f32) { + // let common_name = 1; + // } + Func("func", {Param(Source{{12, 34}}, "common_name", ty.f32())}, ty.void_(), + {Decl(Const(Source{{56, 78}}, "common_name", nullptr, Expr(1)))}); + + EXPECT_FALSE(r()->Resolve()); + EXPECT_EQ(r()->error(), R"(56:78 error: redeclaration of 'common_name' +12:34 note: 'common_name' previously declared here)"); +} + +TEST_F(ResolverFunctionValidationTest, NestedLocalMayShadowParameter) { + // fn func(common_name : f32) { + // { + // let common_name = 1; + // } + // } + Func("func", {Param(Source{{12, 34}}, "common_name", ty.f32())}, ty.void_(), + {Block(Decl(Const(Source{{56, 78}}, "common_name", nullptr, Expr(1))))}); + + ASSERT_TRUE(r()->Resolve()) << r()->error(); } TEST_F(ResolverFunctionValidationTest, @@ -57,7 +78,7 @@ TEST_F(ResolverFunctionValidationTest, Decl(var), }); - EXPECT_TRUE(r()->Resolve()) << r()->error(); + ASSERT_TRUE(r()->Resolve()) << r()->error(); } TEST_F(ResolverFunctionValidationTest, FunctionUsingSameVariableName_Pass) { @@ -74,7 +95,7 @@ TEST_F(ResolverFunctionValidationTest, FunctionUsingSameVariableName_Pass) { }, ast::DecorationList{}); - EXPECT_TRUE(r()->Resolve()) << r()->error(); + ASSERT_TRUE(r()->Resolve()) << r()->error(); } TEST_F(ResolverFunctionValidationTest, @@ -95,7 +116,7 @@ TEST_F(ResolverFunctionValidationTest, }, ast::DecorationList{}); - EXPECT_TRUE(r()->Resolve()) << r()->error(); + ASSERT_TRUE(r()->Resolve()) << r()->error(); } TEST_F(ResolverFunctionValidationTest, UnreachableCode_return) { @@ -189,7 +210,7 @@ TEST_F(ResolverFunctionValidationTest, Func(Source{{12, 34}}, "func", ast::VariableList{}, ty.void_(), ast::StatementList{}); - EXPECT_TRUE(r()->Resolve()) << r()->error(); + ASSERT_TRUE(r()->Resolve()) << r()->error(); } TEST_F(ResolverFunctionValidationTest, @@ -213,7 +234,7 @@ TEST_F(ResolverFunctionValidationTest, Return(), }); - EXPECT_TRUE(r()->Resolve()) << r()->error(); + ASSERT_TRUE(r()->Resolve()) << r()->error(); } TEST_F(ResolverFunctionValidationTest, @@ -269,7 +290,7 @@ TEST_F(ResolverFunctionValidationTest, }, ast::DecorationList{}); - EXPECT_TRUE(r()->Resolve()) << r()->error(); + ASSERT_TRUE(r()->Resolve()) << r()->error(); } TEST_F(ResolverFunctionValidationTest, @@ -298,7 +319,7 @@ TEST_F(ResolverFunctionValidationTest, }, ast::DecorationList{}); - EXPECT_TRUE(r()->Resolve()) << r()->error(); + ASSERT_TRUE(r()->Resolve()) << r()->error(); } TEST_F(ResolverFunctionValidationTest, @@ -362,7 +383,7 @@ TEST_F(ResolverFunctionValidationTest, NoPipelineEntryPoints) { }, ast::DecorationList{}); - EXPECT_TRUE(r()->Resolve()) << r()->error(); + ASSERT_TRUE(r()->Resolve()) << r()->error(); } TEST_F(ResolverFunctionValidationTest, FunctionVarInitWithParam) { @@ -376,7 +397,7 @@ TEST_F(ResolverFunctionValidationTest, FunctionVarInitWithParam) { Func("foo", ast::VariableList{bar}, ty.void_(), ast::StatementList{Decl(baz)}, ast::DecorationList{}); - EXPECT_TRUE(r()->Resolve()) << r()->error(); + ASSERT_TRUE(r()->Resolve()) << r()->error(); } TEST_F(ResolverFunctionValidationTest, FunctionConstInitWithParam) { @@ -390,7 +411,7 @@ TEST_F(ResolverFunctionValidationTest, FunctionConstInitWithParam) { Func("foo", ast::VariableList{bar}, ty.void_(), ast::StatementList{Decl(baz)}, ast::DecorationList{}); - EXPECT_TRUE(r()->Resolve()) << r()->error(); + ASSERT_TRUE(r()->Resolve()) << r()->error(); } TEST_F(ResolverFunctionValidationTest, FunctionParamsConst) { @@ -414,7 +435,7 @@ TEST_F(ResolverFunctionValidationTest, WorkgroupSize_GoodType_ConstU32) { {Stage(ast::PipelineStage::kCompute), WorkgroupSize(Expr("x"), Expr("y"), Expr(16u))}); - EXPECT_TRUE(r()->Resolve()) << r()->error(); + ASSERT_TRUE(r()->Resolve()) << r()->error(); } TEST_F(ResolverFunctionValidationTest, WorkgroupSize_GoodType_U32) { @@ -425,7 +446,7 @@ TEST_F(ResolverFunctionValidationTest, WorkgroupSize_GoodType_U32) { {Stage(ast::PipelineStage::kCompute), WorkgroupSize(Source{{12, 34}}, Expr(1u), Expr(2u), Expr(3u))}); - EXPECT_TRUE(r()->Resolve()) << r()->error(); + ASSERT_TRUE(r()->Resolve()) << r()->error(); } TEST_F(ResolverFunctionValidationTest, WorkgroupSize_MismatchTypeU32) { @@ -696,7 +717,7 @@ TEST_F(ResolverFunctionValidationTest, ParameterSotreType_AtomicFree) { auto* bar = Param(Source{{12, 34}}, "bar", ret_type); Func("f", ast::VariableList{bar}, ty.void_(), {}); - EXPECT_TRUE(r()->Resolve()) << r()->error(); + ASSERT_TRUE(r()->Resolve()) << r()->error(); } TEST_F(ResolverFunctionValidationTest, ParametersAtLimit) { @@ -706,7 +727,7 @@ TEST_F(ResolverFunctionValidationTest, ParametersAtLimit) { } Func(Source{{12, 34}}, "f", params, ty.void_(), {}); - EXPECT_TRUE(r()->Resolve()) << r()->error(); + ASSERT_TRUE(r()->Resolve()) << r()->error(); } TEST_F(ResolverFunctionValidationTest, ParametersOverLimit) { @@ -736,7 +757,7 @@ TEST_P(ResolverFunctionParameterValidationTest, StorageClass) { Func("f", ast::VariableList{arg}, ty.void_(), {}); if (param.should_pass) { - EXPECT_TRUE(r()->Resolve()) << r()->error(); + ASSERT_TRUE(r()->Resolve()) << r()->error(); } else { std::stringstream ss; ss << param.storage_class; diff --git a/src/resolver/resolver.cc b/src/resolver/resolver.cc index b075644caa..0e0abc9e7f 100644 --- a/src/resolver/resolver.cc +++ b/src/resolver/resolver.cc @@ -74,7 +74,6 @@ #include "src/sem/type_conversion.h" #include "src/sem/variable.h" #include "src/utils/defer.h" -#include "src/utils/map.h" #include "src/utils/math.h" #include "src/utils/reverse.h" #include "src/utils/scoped_assignment.h" @@ -142,6 +141,8 @@ bool Resolver::ResolveInternal() { AllocateOverridableConstantIds(); + SetShadows(); + if (!ValidatePipelineStages()) { return false; } @@ -258,14 +259,8 @@ sem::Type* Resolver::Type(const ast::Type* ty) { if (ty->As()) { return builder_->create(); } - if (auto* t = ty->As()) { - auto it = named_type_info_.find(t->name); - if (it == named_type_info_.end()) { - AddError("unknown type '" + builder_->Symbols().NameFor(t->name) + "'", - t->source); - return nullptr; - } - return it->second.sem; + if (auto* type = ResolvedSymbol(ty)) { + return type; } TINT_UNREACHABLE(Resolver, diagnostics_) << "Unhandled ast::Type: " << ty->TypeInfo().name; @@ -423,7 +418,7 @@ sem::Variable* Resolver::Variable(const ast::Variable* var, } case VariableKind::kLocal: { auto* local = builder_->create( - var, var_ty, storage_class, access, + var, var_ty, storage_class, access, current_statement_, (rhs && var->is_const) ? rhs->ConstantValue() : sem::Constant{}); builder_->Sem().Add(var, local); return local; @@ -496,17 +491,23 @@ void Resolver::AllocateOverridableConstantIds() { } } -bool Resolver::GlobalVariable(const ast::Variable* var) { - if (!ValidateNoDuplicateDefinition(var->symbol, var->source, - /* check_global_scope_only */ true)) { - return false; +void Resolver::SetShadows() { + for (auto it : dependencies_.shadows) { + auto* var = Sem(it.first); + if (auto* local = var->As()) { + local->SetShadows(Sem(it.second)); + } + if (auto* param = var->As()) { + param->SetShadows(Sem(it.second)); + } } +} // namespace resolver +bool Resolver::GlobalVariable(const ast::Variable* var) { auto* sem = Variable(var, VariableKind::kGlobal); if (!sem) { return false; } - variable_stack_.Set(var->symbol, sem); auto storage_class = sem->StorageClass(); if (!var->is_const && storage_class == ast::StorageClass::kNone) { @@ -547,9 +548,6 @@ bool Resolver::GlobalVariable(const ast::Variable* var) { } sem::Function* Resolver::Function(const ast::Function* decl) { - variable_stack_.Push(); - TINT_DEFER(variable_stack_.Pop()); - uint32_t parameter_index = 0; std::unordered_map parameter_names; std::vector parameters; @@ -581,7 +579,6 @@ sem::Function* Resolver::Function(const ast::Function* decl) { return nullptr; } - variable_stack_.Set(param->symbol, var); parameters.emplace_back(var); auto* var_ty = const_cast(var->Type()); @@ -684,11 +681,6 @@ sem::Function* Resolver::Function(const ast::Function* decl) { return nullptr; } - // Register the function information _after_ processing the statements. This - // allows us to catch a function calling itself when determining the call - // information as this function doesn't exist until it's finished. - symbol_to_function_[decl->symbol] = func; - // If this is an entry point, mark all transitively called functions as being // used by this entry point. if (decl->IsEntryPoint()) { @@ -1241,22 +1233,28 @@ sem::Call* Resolver::Call(const ast::CallExpression* expr) { auto* ident = expr->target.name; Mark(ident); - auto it = named_type_info_.find(ident->symbol); - if (it != named_type_info_.end()) { - // We have a type. - return type_ctor_or_conv(it->second.sem); + auto* resolved = ResolvedSymbol(ident); + if (auto* ty = As(resolved)) { + return type_ctor_or_conv(ty); + } + + if (auto* fn = As(resolved)) { + return FunctionCall(expr, fn, std::move(args)); } - // Not a type, treat as a intrinsic / function call. auto name = builder_->Symbols().NameFor(ident->symbol); auto intrinsic_type = sem::ParseIntrinsicType(name); - auto* call = (intrinsic_type != sem::IntrinsicType::kNone) - ? IntrinsicCall(expr, intrinsic_type, std::move(args), - std::move(arg_tys)) - : FunctionCall(expr, std::move(args)); + if (intrinsic_type != sem::IntrinsicType::kNone) { + return IntrinsicCall(expr, intrinsic_type, std::move(args), + std::move(arg_tys)); + } - current_function_->AddDirectCall(call); - return call; + TINT_ICE(Resolver, diagnostics_) + << expr->source << " unresolved CallExpression target:\n" + << "resolved: " << (resolved ? resolved->TypeInfo().name : "") + << "\n" + << "name: " << builder_->Symbols().NameFor(ident->symbol); + return nullptr; } sem::Call* Resolver::IntrinsicCall( @@ -1288,37 +1286,27 @@ sem::Call* Resolver::IntrinsicCall( return nullptr; } + current_function_->AddDirectCall(call); + return call; } sem::Call* Resolver::FunctionCall( const ast::CallExpression* expr, + sem::Function* target, const std::vector args) { auto sym = expr->target.name->symbol; auto name = builder_->Symbols().NameFor(sym); - auto target_it = symbol_to_function_.find(sym); - if (target_it == symbol_to_function_.end()) { - if (current_function_ && current_function_->Declaration()->symbol == sym) { - AddError("recursion is not permitted. '" + name + - "' attempted to call itself.", - expr->source); - } else { - AddError("unable to find called function: " + name, expr->source); - } - return nullptr; - } - auto* target = target_it->second; auto* call = builder_->create(expr, target, std::move(args), current_statement_, sem::Constant{}); if (current_function_) { - target->AddCallSite(call); - // Note: Requires called functions to be resolved first. // This is currently guaranteed as functions must be declared before // use. current_function_->AddTransitivelyCalledFunction(target); + current_function_->AddDirectCall(call); for (auto* transitive_call : target->TransitivelyCalledFunctions()) { current_function_->AddTransitivelyCalledFunction(transitive_call); } @@ -1329,6 +1317,8 @@ sem::Call* Resolver::FunctionCall( } } + target->AddCallSite(call); + if (!ValidateFunctionCall(call)) { return nullptr; } @@ -1476,7 +1466,8 @@ sem::Expression* Resolver::Literal(const ast::LiteralExpression* literal) { sem::Expression* Resolver::Identifier(const ast::IdentifierExpression* expr) { auto symbol = expr->symbol; - if (auto* var = variable_stack_.Get(symbol)) { + auto* resolved = ResolvedSymbol(expr); + if (auto* var = As(resolved)) { auto* user = builder_->create(expr, current_statement_, var); @@ -1526,18 +1517,26 @@ sem::Expression* Resolver::Identifier(const ast::IdentifierExpression* expr) { return user; } - if (symbol_to_function_.count(symbol)) { + if (Is(resolved)) { AddError("missing '(' for function call", expr->source.End()); return nullptr; } - std::string name = builder_->Symbols().NameFor(symbol); - if (sem::ParseIntrinsicType(name) != sem::IntrinsicType::kNone) { + if (IsIntrinsic(symbol)) { AddError("missing '(' for intrinsic call", expr->source.End()); return nullptr; } - AddError("unknown identifier: '" + name + "'", expr->source); + if (resolved->Is()) { + AddError("missing '(' for type constructor or cast", expr->source.End()); + return nullptr; + } + + TINT_ICE(Resolver, diagnostics_) + << expr->source << " unresolved identifier:\n" + << "resolved: " << (resolved ? resolved->TypeInfo().name : "") + << "\n" + << "name: " << builder_->Symbols().NameFor(symbol); return nullptr; } @@ -1934,11 +1933,6 @@ sem::Expression* Resolver::UnaryOp(const ast::UnaryOpExpression* unary) { bool Resolver::VariableDeclStatement(const ast::VariableDeclStatement* stmt) { Mark(stmt->variable); - if (!ValidateNoDuplicateDefinition(stmt->variable->symbol, - stmt->variable->source)) { - return false; - } - auto* var = Variable(stmt->variable, VariableKind::kLocal); if (!var) { return false; @@ -1952,7 +1946,6 @@ bool Resolver::VariableDeclStatement(const ast::VariableDeclStatement* stmt) { } } - variable_stack_.Set(stmt->variable->symbol, var); if (current_block_) { // Not all statements are inside a block current_block_->AddDecl(stmt->variable); } @@ -1978,12 +1971,6 @@ sem::Type* Resolver::TypeDecl(const ast::TypeDecl* named_type) { return nullptr; } - named_type_info_.emplace(named_type->name, TypeDeclInfo{named_type, result}); - - if (!ValidateTypeDecl(named_type)) { - return nullptr; - } - builder_->Sem().Add(named_type, result); return result; } @@ -2081,7 +2068,7 @@ sem::Array* Resolver::Array(const ast::Array* arr) { if (auto* ident = count_expr->As()) { // Make sure the identifier is a non-overridable module-scope constant. - auto* var = variable_stack_.Get(ident->symbol); + auto* var = ResolvedSymbol(ident); if (!var || !var->Is() || !var->Declaration()->is_const) { AddError("array size identifier must be a module-scope constant", @@ -2244,13 +2231,11 @@ bool Resolver::Scope(sem::CompoundStatement* stmt, F&& callback) { current_statement_ = stmt; current_compound_statement_ = stmt; current_block_ = stmt->As(); - variable_stack_.Push(); TINT_DEFER({ current_block_ = prev_current_block; current_compound_statement_ = prev_current_compound_statement; current_statement_ = prev_current_statement; - variable_stack_.Pop(); }); return callback(); @@ -2330,6 +2315,11 @@ bool Resolver::IsHostShareable(const sem::Type* type) const { return false; } +bool Resolver::IsIntrinsic(Symbol symbol) const { + std::string name = builder_->Symbols().NameFor(symbol); + return sem::ParseIntrinsicType(name) != sem::IntrinsicType::kNone; +} + //////////////////////////////////////////////////////////////////////////////// // Resolver::TypeConversionSig //////////////////////////////////////////////////////////////////////////////// diff --git a/src/resolver/resolver.h b/src/resolver/resolver.h index a55c38fb11..cca4440250 100644 --- a/src/resolver/resolver.h +++ b/src/resolver/resolver.h @@ -32,6 +32,7 @@ #include "src/sem/constant.h" #include "src/sem/function.h" #include "src/sem/struct.h" +#include "src/utils/map.h" #include "src/utils/unique_vector.h" namespace tint { @@ -176,6 +177,7 @@ class Resolver { sem::Expression* Expression(const ast::Expression*); sem::Function* Function(const ast::Function*); sem::Call* FunctionCall(const ast::CallExpression*, + sem::Function* target, const std::vector args); sem::Expression* Identifier(const ast::IdentifierExpression*); sem::Call* IntrinsicCall(const ast::CallExpression*, @@ -239,9 +241,6 @@ class Resolver { bool ValidateMatrix(const sem::Matrix* ty, const Source& source); bool ValidateFunctionParameter(const ast::Function* func, const sem::Variable* var); - bool ValidateNoDuplicateDefinition(Symbol sym, - const Source& source, - bool check_global_scope_only = false); bool ValidateParameter(const ast::Function* func, const sem::Variable* var); bool ValidateReturn(const ast::ReturnStatement* ret); bool ValidateStatements(const ast::StatementList& stmts); @@ -264,7 +263,6 @@ class Resolver { const sem::Type* type); bool ValidateArrayConstructorOrCast(const ast::CallExpression* ctor, const sem::Array* arr_type); - bool ValidateTypeDecl(const ast::TypeDecl* named_type) const; bool ValidateTextureIntrinsicFunction(const sem::Call* call); bool ValidateNoDuplicateDecorations(const ast::DecorationList& decorations); // sem::Struct is assumed to have at least one member @@ -343,6 +341,10 @@ class Resolver { /// Allocate constant IDs for pipeline-overridable constants. void AllocateOverridableConstantIds(); + /// Set the shadowing information on variable declarations. + /// @note this method must only be called after all semantic nodes are built. + void SetShadows(); + /// @returns the resolved type of the ast::Expression `expr` /// @param expr the expression sem::Type* TypeOf(const ast::Expression* expr); @@ -410,14 +412,28 @@ class Resolver { template auto* Sem(const AST_OR_TYPE* ast) { - auto* sem = builder_->Sem().Get(ast); + using T = sem::Info::GetResultType; + auto* sem = builder_->Sem().Get(ast); if (!sem) { TINT_ICE(Resolver, diagnostics_) << "AST node '" << ast->TypeInfo().name << "' had no semantic info\n" << "At: " << ast->source << "\n" << "Pointer: " << ast; } - return sem; + return const_cast(As(sem)); + } + + /// @returns true if the symbol is the name of an intrinsic (builtin) + /// function. + bool IsIntrinsic(Symbol) const; + + /// @returns the resolved symbol (function, type or variable) for the given + /// ast::Identifier or ast::TypeName cast to the given semantic type. + template + SEM* ResolvedSymbol(const ast::Node* node) { + auto* resolved = utils::Lookup(dependencies_.resolved_symbols, node); + return resolved ? const_cast(builder_->Sem().Get(resolved)) + : nullptr; } struct TypeConversionSig { @@ -456,12 +472,8 @@ class Resolver { diag::List& diagnostics_; std::unique_ptr const intrinsic_table_; DependencyGraph dependencies_; - ScopeStack variable_stack_; - std::unordered_map symbol_to_function_; std::vector entry_points_; std::unordered_map atomic_composite_info_; - std::unordered_map named_type_info_; - std::unordered_set marked_; std::unordered_map constant_ids_; std::unordered_mapDeclaration(); - if (!ValidateNoDuplicateDefinition(decl->symbol, decl->source, - /* check_global_scope_only */ true)) { - return false; - } - auto workgroup_deco_count = 0; for (auto* deco : decl->decorations) { if (deco->Is()) { @@ -1444,7 +1439,7 @@ bool Resolver::ValidateFunctionCall(const sem::Call* call) { if (param_type->Is()) { auto is_valid = false; if (auto* ident_expr = arg_expr->As()) { - auto* var = variable_stack_.Get(ident_expr->symbol); + auto* var = ResolvedSymbol(ident_expr); if (!var) { TINT_ICE(Resolver, diagnostics_) << "failed to resolve identifier"; return false; @@ -1456,7 +1451,7 @@ bool Resolver::ValidateFunctionCall(const sem::Call* call) { if (unary->op == ast::UnaryOp::kAddressOf) { if (auto* ident_unary = unary->expr->As()) { - auto* var = variable_stack_.Get(ident_unary->symbol); + auto* var = ResolvedSymbol(ident_unary); if (!var) { TINT_ICE(Resolver, diagnostics_) << "failed to resolve identifier"; @@ -1755,23 +1750,6 @@ bool Resolver::ValidateScalarConstructorOrCast(const ast::CallExpression* ctor, return true; } -bool Resolver::ValidateTypeDecl(const ast::TypeDecl* named_type) const { - auto iter = named_type_info_.find(named_type->name); - if (iter == named_type_info_.end()) { - TINT_ICE(Resolver, diagnostics_) - << "ValidateTypeDecl called() before TypeDecl()"; - } - if (iter->second.ast != named_type) { - AddError("type with the name '" + - builder_->Symbols().NameFor(named_type->name) + - "' was already declared", - named_type->source); - AddNote("first declared here", iter->second.ast->source); - return false; - } - return true; -} - bool Resolver::ValidatePipelineStages() { auto check_workgroup_storage = [&](const sem::Function* func, const sem::Function* entry_point) { @@ -2338,22 +2316,21 @@ bool Resolver::ValidateAssignment(const ast::AssignmentStatement* a) { // https://gpuweb.github.io/gpuweb/wgsl/#assignment-statement auto const* lhs_ty = TypeOf(a->lhs); - if (auto* ident = a->lhs->As()) { - if (auto* var = variable_stack_.Get(ident->symbol)) { - if (var->Is()) { - AddError("cannot assign to function parameter", a->lhs->source); - AddNote("'" + builder_->Symbols().NameFor(ident->symbol) + - "' is declared here:", - var->Declaration()->source); - return false; - } - if (var->Declaration()->is_const) { - AddError("cannot assign to const", a->lhs->source); - AddNote("'" + builder_->Symbols().NameFor(ident->symbol) + - "' is declared here:", - var->Declaration()->source); - return false; - } + if (auto* var = ResolvedSymbol(a->lhs)) { + auto* decl = var->Declaration(); + if (var->Is()) { + AddError("cannot assign to function parameter", a->lhs->source); + AddNote("'" + builder_->Symbols().NameFor(decl->symbol) + + "' is declared here:", + decl->source); + return false; + } + if (decl->is_const) { + AddError("cannot assign to const", a->lhs->source); + AddNote("'" + builder_->Symbols().NameFor(decl->symbol) + + "' is declared here:", + decl->source); + return false; } } @@ -2388,36 +2365,6 @@ bool Resolver::ValidateAssignment(const ast::AssignmentStatement* a) { return true; } -bool Resolver::ValidateNoDuplicateDefinition(Symbol sym, - const Source& source, - bool check_global_scope_only) { - if (check_global_scope_only) { - if (auto* var = variable_stack_.Get(sym)) { - if (var->Is()) { - AddError("redefinition of '" + builder_->Symbols().NameFor(sym) + "'", - source); - AddNote("previous definition is here", var->Declaration()->source); - return false; - } - } - auto it = symbol_to_function_.find(sym); - if (it != symbol_to_function_.end()) { - AddError("redefinition of '" + builder_->Symbols().NameFor(sym) + "'", - source); - AddNote("previous definition is here", it->second->Declaration()->source); - return false; - } - } else { - if (auto* var = variable_stack_.Get(sym)) { - AddError("redefinition of '" + builder_->Symbols().NameFor(sym) + "'", - source); - AddNote("previous definition is here", var->Declaration()->source); - return false; - } - } - return true; -} - bool Resolver::ValidateNoDuplicateDecorations( const ast::DecorationList& decorations) { std::unordered_map seen; diff --git a/src/resolver/validation_test.cc b/src/resolver/validation_test.cc index 790fde5631..96557d9cdd 100644 --- a/src/resolver/validation_test.cc +++ b/src/resolver/validation_test.cc @@ -150,25 +150,28 @@ TEST_F(ResolverValidationTest, Expr_ErrUnknownExprType) { TEST_F(ResolverValidationTest, Expr_DontCall_Function) { Func("func", {}, ty.void_(), {}, {}); - auto* ident = create( - Source{{Source::Location{3, 3}, Source::Location{3, 8}}}, - Symbols().Register("func")); - WrapInFunction(ident); + WrapInFunction(Expr(Source{{{3, 3}, {3, 8}}}, "func")); EXPECT_FALSE(r()->Resolve()); EXPECT_EQ(r()->error(), "3:8 error: missing '(' for function call"); } TEST_F(ResolverValidationTest, Expr_DontCall_Intrinsic) { - auto* ident = create( - Source{{Source::Location{3, 3}, Source::Location{3, 8}}}, - Symbols().Register("round")); - WrapInFunction(ident); + WrapInFunction(Expr(Source{{{3, 3}, {3, 8}}}, "round")); EXPECT_FALSE(r()->Resolve()); EXPECT_EQ(r()->error(), "3:8 error: missing '(' for intrinsic call"); } +TEST_F(ResolverValidationTest, Expr_DontCall_Type) { + Alias("T", ty.u32()); + WrapInFunction(Expr(Source{{{3, 3}, {3, 8}}}, "T")); + + EXPECT_FALSE(r()->Resolve()); + EXPECT_EQ(r()->error(), + "3:8 error: missing '(' for type constructor or cast"); +} + TEST_F(ResolverValidationTest, AssignmentStmt_InvalidLHS_IntrinsicFunctionName) { // normalize = 2; @@ -220,7 +223,7 @@ TEST_F(ResolverValidationTest, UsingUndefinedVariableGlobalVariable_Pass) { Func("my_func", ast::VariableList{}, ty.void_(), { - Assign(Expr(Source{Source::Location{12, 34}}, "global_var"), 3.14f), + Assign(Expr(Source{{12, 34}}, "global_var"), 3.14f), Return(), }); @@ -237,7 +240,7 @@ TEST_F(ResolverValidationTest, UsingUndefinedVariableInnerScope_Fail) { auto* cond = Expr(true); auto* body = Block(Decl(var)); - SetSource(Source{Source::Location{12, 34}}); + SetSource(Source{{12, 34}}); auto* lhs = Expr(Source{{12, 34}}, "a"); auto* rhs = Expr(3.14f); @@ -350,9 +353,7 @@ TEST_F(ResolverValidationTest, StorageClass_TextureExplicitStorageClass) { TEST_F(ResolverValidationTest, Expr_MemberAccessor_VectorSwizzle_BadChar) { Global("my_vec", ty.vec3(), ast::StorageClass::kPrivate); - auto* ident = create( - Source{{Source::Location{3, 3}, Source::Location{3, 7}}}, - Symbols().Register("xyqz")); + auto* ident = Expr(Source{{{3, 3}, {3, 7}}}, "xyqz"); auto* mem = MemberAccessor("my_vec", ident); WrapInFunction(mem); @@ -364,9 +365,7 @@ TEST_F(ResolverValidationTest, Expr_MemberAccessor_VectorSwizzle_BadChar) { TEST_F(ResolverValidationTest, Expr_MemberAccessor_VectorSwizzle_MixedChars) { Global("my_vec", ty.vec4(), ast::StorageClass::kPrivate); - auto* ident = create( - Source{{Source::Location{3, 3}, Source::Location{3, 7}}}, - Symbols().Register("rgyw")); + auto* ident = Expr(Source{{{3, 3}, {3, 7}}}, "rgyw"); auto* mem = MemberAccessor("my_vec", ident); WrapInFunction(mem); @@ -380,9 +379,7 @@ TEST_F(ResolverValidationTest, Expr_MemberAccessor_VectorSwizzle_MixedChars) { TEST_F(ResolverValidationTest, Expr_MemberAccessor_VectorSwizzle_BadLength) { Global("my_vec", ty.vec3(), ast::StorageClass::kPrivate); - auto* ident = create( - Source{{Source::Location{3, 3}, Source::Location{3, 8}}}, - Symbols().Register("zzzzz")); + auto* ident = Expr(Source{{{3, 3}, {3, 8}}}, "zzzzz"); auto* mem = MemberAccessor("my_vec", ident); WrapInFunction(mem); @@ -393,8 +390,7 @@ TEST_F(ResolverValidationTest, Expr_MemberAccessor_VectorSwizzle_BadLength) { TEST_F(ResolverValidationTest, Expr_MemberAccessor_VectorSwizzle_BadIndex) { Global("my_vec", ty.vec2(), ast::StorageClass::kPrivate); - auto* ident = create(Source{{3, 3}}, - Symbols().Register("z")); + auto* ident = Expr(Source{{3, 3}}, "z"); auto* mem = MemberAccessor("my_vec", ident); WrapInFunction(mem); @@ -406,9 +402,7 @@ TEST_F(ResolverValidationTest, Expr_MemberAccessor_BadParent) { // var param: vec4 // let ret: f32 = *(¶m).x; auto* param = Var("param", ty.vec4()); - auto* x = create( - Source{{Source::Location{3, 3}, Source::Location{3, 8}}}, - Symbols().Register("x")); + auto* x = Expr(Source{{{3, 3}, {3, 8}}}, "x"); auto* addressOf_expr = AddressOf(Source{{12, 34}}, param); auto* accessor_expr = MemberAccessor(addressOf_expr, x); @@ -430,9 +424,7 @@ TEST_F(ResolverValidationTest, EXpr_MemberAccessor_FuncGoodParent) { auto* p = Param("p", ty.pointer(ty.vec4(), ast::StorageClass::kFunction)); auto* star_p = Deref(p); - auto* z = create( - Source{{Source::Location{3, 3}, Source::Location{3, 8}}}, - Symbols().Register("z")); + auto* z = Expr(Source{{{3, 3}, {3, 8}}}, "z"); auto* accessor_expr = MemberAccessor(star_p, z); auto* x = Var("x", ty.f32(), accessor_expr); Func("func", {p}, ty.f32(), {Decl(x), Return(x)}); @@ -446,9 +438,7 @@ TEST_F(ResolverValidationTest, EXpr_MemberAccessor_FuncBadParent) { // } auto* p = Param("p", ty.pointer(ty.vec4(), ast::StorageClass::kFunction)); - auto* z = create( - Source{{Source::Location{3, 3}, Source::Location{3, 8}}}, - Symbols().Register("z")); + auto* z = Expr(Source{{{3, 3}, {3, 8}}}, "z"); auto* accessor_expr = MemberAccessor(p, z); auto* star_p = Deref(accessor_expr); auto* x = Var("x", ty.f32(), star_p); @@ -473,7 +463,7 @@ TEST_F(ResolverValidationTest, // } // } - auto error_loc = Source{Source::Location{12, 34}}; + auto error_loc = Source{{12, 34}}; auto* body = Block(create(), Decl(error_loc, Var("z", ty.i32(), ast::StorageClass::kNone)), @@ -524,9 +514,9 @@ TEST_F(ResolverValidationTest, // } // } - auto cont_loc = Source{Source::Location{12, 34}}; - auto decl_loc = Source{Source::Location{56, 78}}; - auto ref_loc = Source{Source::Location{90, 12}}; + auto cont_loc = Source{{12, 34}}; + auto decl_loc = Source{{56, 78}}; + auto ref_loc = Source{{90, 12}}; auto* body = Block(If(Expr(true), Block(create(cont_loc))), Decl(Var(decl_loc, "z", ty.i32(), ast::StorageClass::kNone))); @@ -556,9 +546,9 @@ TEST_F( // } // } - auto cont_loc = Source{Source::Location{12, 34}}; - auto decl_loc = Source{Source::Location{56, 78}}; - auto ref_loc = Source{Source::Location{90, 12}}; + auto cont_loc = Source{{12, 34}}; + auto decl_loc = Source{{56, 78}}; + auto ref_loc = Source{{90, 12}}; auto* body = Block(If(Expr(true), Block(create(cont_loc))), Decl(Var(decl_loc, "z", ty.i32(), ast::StorageClass::kNone))); @@ -590,9 +580,9 @@ TEST_F(ResolverValidationTest, // } // } - auto cont_loc = Source{Source::Location{12, 34}}; - auto decl_loc = Source{Source::Location{56, 78}}; - auto ref_loc = Source{Source::Location{90, 12}}; + auto cont_loc = Source{{12, 34}}; + auto decl_loc = Source{{56, 78}}; + auto ref_loc = Source{{90, 12}}; auto* body = Block(If(Expr(true), Block(create(cont_loc))), Decl(Var(decl_loc, "z", ty.i32(), ast::StorageClass::kNone))); @@ -623,9 +613,9 @@ TEST_F(ResolverValidationTest, // } // } - auto cont_loc = Source{Source::Location{12, 34}}; - auto decl_loc = Source{Source::Location{56, 78}}; - auto ref_loc = Source{Source::Location{90, 12}}; + auto cont_loc = Source{{12, 34}}; + auto decl_loc = Source{{56, 78}}; + auto ref_loc = Source{{90, 12}}; auto* body = Block(If(Expr(true), Block(create(cont_loc))), Decl(Var(decl_loc, "z", ty.i32(), ast::StorageClass::kNone))); @@ -724,7 +714,7 @@ TEST_F(ResolverTest, Stmt_Loop_ContinueInLoopBodyAfterDecl_UsageInContinuing) { // } // } - auto error_loc = Source{Source::Location{12, 34}}; + auto error_loc = Source{{12, 34}}; auto* body = Block(Decl(Var("z", ty.i32(), ast::StorageClass::kNone)), create()); auto* continuing = Block(Assign(Expr(error_loc, "z"), 2)); diff --git a/src/resolver/var_let_test.cc b/src/resolver/var_let_test.cc index 1b45fac41b..57d7a148c5 100644 --- a/src/resolver/var_let_test.cc +++ b/src/resolver/var_let_test.cc @@ -58,7 +58,7 @@ TEST_F(ResolverVarLetTest, TypeOfVar) { Decl(a), }); - EXPECT_TRUE(r()->Resolve()) << r()->error(); + ASSERT_TRUE(r()->Resolve()) << r()->error(); // `var` declarations are always of reference type ASSERT_TRUE(TypeOf(i)->Is()); @@ -114,7 +114,7 @@ TEST_F(ResolverVarLetTest, TypeOfLet) { Decl(p), }); - EXPECT_TRUE(r()->Resolve()) << r()->error(); + ASSERT_TRUE(r()->Resolve()) << r()->error(); // `let` declarations are always of the storage type EXPECT_TRUE(TypeOf(i)->Is()); @@ -153,7 +153,7 @@ TEST_F(ResolverVarLetTest, DefaultVarStorageClass) { WrapInFunction(function); - EXPECT_TRUE(r()->Resolve()) << r()->error(); + ASSERT_TRUE(r()->Resolve()) << r()->error(); ASSERT_TRUE(TypeOf(function)->Is()); ASSERT_TRUE(TypeOf(private_)->Is()); @@ -187,7 +187,7 @@ TEST_F(ResolverVarLetTest, ExplicitVarStorageClass) { create(0), }); - EXPECT_TRUE(r()->Resolve()) << r()->error(); + ASSERT_TRUE(r()->Resolve()) << r()->error(); ASSERT_TRUE(TypeOf(storage)->Is()); @@ -222,7 +222,7 @@ TEST_F(ResolverVarLetTest, LetInheritsAccessFromOriginatingVariable) { WrapInFunction(ptr); - EXPECT_TRUE(r()->Resolve()) << r()->error(); + ASSERT_TRUE(r()->Resolve()) << r()->error(); ASSERT_TRUE(TypeOf(expr)->Is()); ASSERT_TRUE(TypeOf(ptr)->Is()); @@ -232,6 +232,384 @@ TEST_F(ResolverVarLetTest, LetInheritsAccessFromOriginatingVariable) { EXPECT_EQ(TypeOf(ptr)->As()->Access(), ast::Access::kReadWrite); } +TEST_F(ResolverVarLetTest, LocalShadowsAlias) { + // type a = i32; + // + // fn X() { + // var a = false; + // } + // + // fn Y() { + // let a = true; + // } + + auto* t = Alias("a", ty.i32()); + auto* v = Var("a", nullptr, Expr(false)); + auto* l = Const("a", nullptr, Expr(false)); + Func("X", {}, ty.void_(), {Decl(v)}); + Func("Y", {}, ty.void_(), {Decl(l)}); + + ASSERT_TRUE(r()->Resolve()) << r()->error(); + + auto* type_t = Sem().Get(t); + auto* local_v = Sem().Get(v); + auto* local_l = Sem().Get(l); + + ASSERT_NE(local_v, nullptr); + ASSERT_NE(local_l, nullptr); + + EXPECT_EQ(local_v->Shadows(), type_t); + EXPECT_EQ(local_l->Shadows(), type_t); +} + +TEST_F(ResolverVarLetTest, LocalShadowsStruct) { + // struct a { + // m : i32; + // }; + // + // fn X() { + // var a = true; + // } + // + // fn Y() { + // let a = false; + // } + + auto* t = Structure("a", {Member("m", ty.i32())}); + auto* v = Var("a", nullptr, Expr(false)); + auto* l = Const("a", nullptr, Expr(false)); + Func("X", {}, ty.void_(), {Decl(v)}); + Func("Y", {}, ty.void_(), {Decl(l)}); + + ASSERT_TRUE(r()->Resolve()) << r()->error(); + + auto* type_t = Sem().Get(t); + auto* local_v = Sem().Get(v); + auto* local_l = Sem().Get(l); + + ASSERT_NE(local_v, nullptr); + ASSERT_NE(local_l, nullptr); + + EXPECT_EQ(local_v->Shadows(), type_t); + EXPECT_EQ(local_l->Shadows(), type_t); +} + +TEST_F(ResolverVarLetTest, LocalShadowsFunction) { + // fn a() { + // var a = true; + // } + // + // fn b() { + // let b = false; + // } + + auto* v = Var("a", nullptr, Expr(false)); + auto* l = Const("b", nullptr, Expr(false)); + auto* fa = Func("a", {}, ty.void_(), {Decl(v)}); + auto* fb = Func("b", {}, ty.void_(), {Decl(l)}); + + ASSERT_TRUE(r()->Resolve()) << r()->error(); + + auto* local_v = Sem().Get(v); + auto* local_l = Sem().Get(l); + auto* func_a = Sem().Get(fa); + auto* func_b = Sem().Get(fb); + + ASSERT_NE(local_v, nullptr); + ASSERT_NE(local_l, nullptr); + ASSERT_NE(func_a, nullptr); + ASSERT_NE(func_b, nullptr); + + EXPECT_EQ(local_v->Shadows(), func_a); + EXPECT_EQ(local_l->Shadows(), func_b); +} + +TEST_F(ResolverVarLetTest, LocalShadowsGlobalVar) { + // var a : i32; + // + // fn X() { + // var a = a; + // } + // + // fn Y() { + // let a = a; + // } + + auto* g = Global("a", ty.i32(), ast::StorageClass::kPrivate); + auto* v = Var("a", nullptr, Expr("a")); + auto* l = Const("a", nullptr, Expr("a")); + Func("X", {}, ty.void_(), {Decl(v)}); + Func("Y", {}, ty.void_(), {Decl(l)}); + + ASSERT_TRUE(r()->Resolve()) << r()->error(); + + auto* global = Sem().Get(g); + auto* local_v = Sem().Get(v); + auto* local_l = Sem().Get(l); + + ASSERT_NE(local_v, nullptr); + ASSERT_NE(local_l, nullptr); + + EXPECT_EQ(local_v->Shadows(), global); + EXPECT_EQ(local_l->Shadows(), global); + + auto* user_v = + Sem().Get(local_v->Declaration()->constructor); + auto* user_l = + Sem().Get(local_l->Declaration()->constructor); + + ASSERT_NE(user_v, nullptr); + ASSERT_NE(user_l, nullptr); + + EXPECT_EQ(user_v->Variable(), global); + EXPECT_EQ(user_l->Variable(), global); +} + +TEST_F(ResolverVarLetTest, LocalShadowsGlobalLet) { + // let a : i32 = 1; + // + // fn X() { + // var a = (a == 123); + // } + // + // fn Y() { + // let a = (a == 321); + // } + + auto* g = GlobalConst("a", ty.i32(), Expr(1)); + auto* v = Var("a", nullptr, Expr("a")); + auto* l = Const("a", nullptr, Expr("a")); + Func("X", {}, ty.void_(), {Decl(v)}); + Func("Y", {}, ty.void_(), {Decl(l)}); + + ASSERT_TRUE(r()->Resolve()) << r()->error(); + + auto* global = Sem().Get(g); + auto* local_v = Sem().Get(v); + auto* local_l = Sem().Get(l); + + ASSERT_NE(local_v, nullptr); + ASSERT_NE(local_l, nullptr); + + EXPECT_EQ(local_v->Shadows(), global); + EXPECT_EQ(local_l->Shadows(), global); + + auto* user_v = + Sem().Get(local_v->Declaration()->constructor); + auto* user_l = + Sem().Get(local_l->Declaration()->constructor); + + ASSERT_NE(user_v, nullptr); + ASSERT_NE(user_l, nullptr); + + EXPECT_EQ(user_v->Variable(), global); + EXPECT_EQ(user_l->Variable(), global); +} + +TEST_F(ResolverVarLetTest, LocalShadowsLocalVar) { + // fn X() { + // var a : i32; + // { + // var a = a; + // } + // { + // let a = a; + // } + // } + + auto* s = Var("a", ty.i32(), Expr(1)); + auto* v = Var("a", nullptr, Expr("a")); + auto* l = Const("a", nullptr, Expr("a")); + Func("X", {}, ty.void_(), {Decl(s), Block(Decl(v)), Block(Decl(l))}); + + ASSERT_TRUE(r()->Resolve()) << r()->error(); + + auto* local_s = Sem().Get(s); + auto* local_v = Sem().Get(v); + auto* local_l = Sem().Get(l); + + ASSERT_NE(local_s, nullptr); + ASSERT_NE(local_v, nullptr); + ASSERT_NE(local_l, nullptr); + + EXPECT_EQ(local_v->Shadows(), local_s); + EXPECT_EQ(local_l->Shadows(), local_s); + + auto* user_v = + Sem().Get(local_v->Declaration()->constructor); + auto* user_l = + Sem().Get(local_l->Declaration()->constructor); + + ASSERT_NE(user_v, nullptr); + ASSERT_NE(user_l, nullptr); + + EXPECT_EQ(user_v->Variable(), local_s); + EXPECT_EQ(user_l->Variable(), local_s); +} + +TEST_F(ResolverVarLetTest, LocalShadowsLocalLet) { + // fn X() { + // let a = 1; + // { + // var a = (a == 123); + // } + // { + // let a = (a == 321); + // } + // } + + auto* s = Const("a", ty.i32(), Expr(1)); + auto* v = Var("a", nullptr, Expr("a")); + auto* l = Const("a", nullptr, Expr("a")); + Func("X", {}, ty.void_(), {Decl(s), Block(Decl(v)), Block(Decl(l))}); + + ASSERT_TRUE(r()->Resolve()) << r()->error(); + + auto* local_s = Sem().Get(s); + auto* local_v = Sem().Get(v); + auto* local_l = Sem().Get(l); + + ASSERT_NE(local_s, nullptr); + ASSERT_NE(local_v, nullptr); + ASSERT_NE(local_l, nullptr); + + EXPECT_EQ(local_v->Shadows(), local_s); + EXPECT_EQ(local_l->Shadows(), local_s); + + auto* user_v = + Sem().Get(local_v->Declaration()->constructor); + auto* user_l = + Sem().Get(local_l->Declaration()->constructor); + + ASSERT_NE(user_v, nullptr); + ASSERT_NE(user_l, nullptr); + + EXPECT_EQ(user_v->Variable(), local_s); + EXPECT_EQ(user_l->Variable(), local_s); +} + +TEST_F(ResolverVarLetTest, LocalShadowsParam) { + // fn F(a : i32) { + // { + // var a = a; + // } + // { + // let a = a; + // } + // } + + auto* p = Param("a", ty.i32()); + auto* v = Var("a", nullptr, Expr("a")); + auto* l = Const("a", nullptr, Expr("a")); + Func("X", {p}, ty.void_(), {Block(Decl(v)), Block(Decl(l))}); + + ASSERT_TRUE(r()->Resolve()) << r()->error(); + + auto* param = Sem().Get(p); + auto* local_v = Sem().Get(v); + auto* local_l = Sem().Get(l); + + ASSERT_NE(param, nullptr); + ASSERT_NE(local_v, nullptr); + ASSERT_NE(local_l, nullptr); + + EXPECT_EQ(local_v->Shadows(), param); + EXPECT_EQ(local_l->Shadows(), param); + + auto* user_v = + Sem().Get(local_v->Declaration()->constructor); + auto* user_l = + Sem().Get(local_l->Declaration()->constructor); + + ASSERT_NE(user_v, nullptr); + ASSERT_NE(user_l, nullptr); + + EXPECT_EQ(user_v->Variable(), param); + EXPECT_EQ(user_l->Variable(), param); +} + +TEST_F(ResolverVarLetTest, ParamShadowsFunction) { + // fn a(a : bool) { + // } + + auto* p = Param("a", ty.bool_()); + auto* f = Func("a", {p}, ty.void_(), {}); + + ASSERT_TRUE(r()->Resolve()) << r()->error(); + + auto* func = Sem().Get(f); + auto* param = Sem().Get(p); + + ASSERT_NE(func, nullptr); + ASSERT_NE(param, nullptr); + + EXPECT_EQ(param->Shadows(), func); +} + +TEST_F(ResolverVarLetTest, ParamShadowsGlobalVar) { + // var a : i32; + // + // fn F(a : bool) { + // } + + auto* g = Global("a", ty.i32(), ast::StorageClass::kPrivate); + auto* p = Param("a", ty.bool_()); + Func("F", {p}, ty.void_(), {}); + + ASSERT_TRUE(r()->Resolve()) << r()->error(); + + auto* global = Sem().Get(g); + auto* param = Sem().Get(p); + + ASSERT_NE(global, nullptr); + ASSERT_NE(param, nullptr); + + EXPECT_EQ(param->Shadows(), global); +} + +TEST_F(ResolverVarLetTest, ParamShadowsGlobalLet) { + // let a : i32 = 1; + // + // fn F(a : bool) { + // } + + auto* g = GlobalConst("a", ty.i32(), Expr(1)); + auto* p = Param("a", ty.bool_()); + Func("F", {p}, ty.void_(), {}); + + ASSERT_TRUE(r()->Resolve()) << r()->error(); + + auto* global = Sem().Get(g); + auto* param = Sem().Get(p); + + ASSERT_NE(global, nullptr); + ASSERT_NE(param, nullptr); + + EXPECT_EQ(param->Shadows(), global); +} + +TEST_F(ResolverVarLetTest, ParamShadowsAlias) { + // type a = i32; + // + // fn F(a : a) { + // } + + auto* a = Alias("a", ty.i32()); + auto* p = Param("a", ty.type_name("a")); + Func("F", {p}, ty.void_(), {}); + + ASSERT_TRUE(r()->Resolve()) << r()->error(); + + auto* alias = Sem().Get(a); + auto* param = Sem().Get(p); + + ASSERT_NE(alias, nullptr); + ASSERT_NE(param, nullptr); + + EXPECT_EQ(param->Shadows(), alias); + EXPECT_EQ(param->Type(), alias); +} + } // namespace } // namespace resolver } // namespace tint diff --git a/src/resolver/var_let_validation_test.cc b/src/resolver/var_let_validation_test.cc index b3a903e745..6326f3ccc4 100644 --- a/src/resolver/var_let_validation_test.cc +++ b/src/resolver/var_let_validation_test.cc @@ -175,10 +175,7 @@ TEST_F(ResolverVarLetValidationTest, GlobalVarRedeclaredAsLocal) { WrapInFunction(Var(Source{{12, 34}}, "v", ty.f32(), ast::StorageClass::kNone, Expr(2.0f))); - EXPECT_FALSE(r()->Resolve()) << r()->error(); - EXPECT_EQ( - r()->error(), - "12:34 error: redefinition of 'v'\nnote: previous definition is here"); + EXPECT_TRUE(r()->Resolve()) << r()->error(); } TEST_F(ResolverVarLetValidationTest, VarRedeclaredInInnerBlock) { @@ -194,10 +191,7 @@ TEST_F(ResolverVarLetValidationTest, VarRedeclaredInInnerBlock) { WrapInFunction(outer_body); - EXPECT_FALSE(r()->Resolve()); - EXPECT_EQ( - r()->error(), - "12:34 error: redefinition of 'v'\nnote: previous definition is here"); + EXPECT_TRUE(r()->Resolve()) << r()->error(); } TEST_F(ResolverVarLetValidationTest, VarRedeclaredInIfBlock) { @@ -219,10 +213,7 @@ TEST_F(ResolverVarLetValidationTest, VarRedeclaredInIfBlock) { WrapInFunction(outer_body); - EXPECT_FALSE(r()->Resolve()); - EXPECT_EQ( - r()->error(), - "12:34 error: redefinition of 'v'\nnote: previous definition is here"); + EXPECT_TRUE(r()->Resolve()) << r()->error(); } TEST_F(ResolverVarLetValidationTest, InferredPtrStorageAccessMismatch) { diff --git a/src/sem/type_mappings.h b/src/sem/type_mappings.h index 97a052b729..073ff13200 100644 --- a/src/sem/type_mappings.h +++ b/src/sem/type_mappings.h @@ -25,6 +25,7 @@ class CallExpression; class Expression; class Function; class MemberAccessorExpression; +class Node; class Statement; class Struct; class StructMember; @@ -56,11 +57,12 @@ struct TypeMappings { Expression* operator()(ast::Expression*); Function* operator()(ast::Function*); MemberAccessorExpression* operator()(ast::MemberAccessorExpression*); + Node* operator()(ast::Node*); Statement* operator()(ast::Statement*); + Struct* operator()(ast::Struct*); StructMember* operator()(ast::StructMember*); Type* operator()(ast::Type*); Type* operator()(ast::TypeDecl*); - Struct* operator()(ast::Struct*); Variable* operator()(ast::Variable*); //! @endcond }; diff --git a/src/sem/variable.cc b/src/sem/variable.cc index 8b5fdd98db..364a3cb032 100644 --- a/src/sem/variable.cc +++ b/src/sem/variable.cc @@ -45,12 +45,10 @@ LocalVariable::LocalVariable(const ast::Variable* declaration, const sem::Type* type, ast::StorageClass storage_class, ast::Access access, + const sem::Statement* statement, Constant constant_value) - : Base(declaration, - type, - storage_class, - access, - std::move(constant_value)) {} + : Base(declaration, type, storage_class, access, std::move(constant_value)), + statement_(statement) {} LocalVariable::~LocalVariable() = default; diff --git a/src/sem/variable.h b/src/sem/variable.h index 78e208ad75..a389eed70d 100644 --- a/src/sem/variable.h +++ b/src/sem/variable.h @@ -95,15 +95,31 @@ class LocalVariable : public Castable { /// @param type the variable type /// @param storage_class the variable storage class /// @param access the variable access control type + /// @param statement the statement that declared this local variable /// @param constant_value the constant value for the variable. May be invalid LocalVariable(const ast::Variable* declaration, const sem::Type* type, ast::StorageClass storage_class, ast::Access access, + const sem::Statement* statement, Constant constant_value); /// Destructor ~LocalVariable() override; + + /// @returns the statement that declares this local variable + const sem::Statement* Statement() const { return statement_; } + + /// @returns the Type, Function or Variable that this local variable shadows + const sem::Node* Shadows() const { return shadows_; } + + /// Sets the Type, Function or Variable that this local variable shadows + /// @param shadows the Type, Function or Variable that this variable shadows + void SetShadows(const sem::Node* shadows) { shadows_ = shadows; } + + private: + const sem::Statement* const statement_; + const sem::Node* shadows_ = nullptr; }; /// GlobalVariable is a module-scope variable @@ -185,10 +201,18 @@ class Parameter : public Castable { /// @param owner the CallTarget owner of this parameter void SetOwner(CallTarget const* owner) { owner_ = owner; } + /// @returns the Type, Function or Variable that this local variable shadows + const sem::Node* Shadows() const { return shadows_; } + + /// Sets the Type, Function or Variable that this local variable shadows + /// @param shadows the Type, Function or Variable that this variable shadows + void SetShadows(const sem::Node* shadows) { shadows_ = shadows; } + private: const uint32_t index_; const ParameterUsage usage_; CallTarget const* owner_ = nullptr; + const sem::Node* shadows_ = nullptr; }; /// ParameterList is a list of Parameter diff --git a/src/transform/array_length_from_uniform_test.cc b/src/transform/array_length_from_uniform_test.cc index 13aca0ae90..c31641210d 100644 --- a/src/transform/array_length_from_uniform_test.cc +++ b/src/transform/array_length_from_uniform_test.cc @@ -18,6 +18,7 @@ #include "src/transform/simplify_pointers.h" #include "src/transform/test_helper.h" +#include "src/transform/unshadow.h" namespace tint { namespace transform { @@ -32,7 +33,7 @@ TEST_F(ArrayLengthFromUniformTest, Error_MissingTransformData) { "error: missing transform data for " "tint::transform::ArrayLengthFromUniform"; - auto got = Run(src); + auto got = Run(src); EXPECT_EQ(expect, str(got)); } @@ -93,7 +94,7 @@ fn main() { DataMap data; data.Add(std::move(cfg)); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); EXPECT_EQ(std::unordered_set({0}), @@ -146,7 +147,7 @@ fn main() { DataMap data; data.Add(std::move(cfg)); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); EXPECT_EQ(std::unordered_set({0}), @@ -267,7 +268,7 @@ fn main() { DataMap data; data.Add(std::move(cfg)); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); EXPECT_EQ(std::unordered_set({0, 1, 2, 3, 4}), @@ -382,7 +383,7 @@ fn main() { DataMap data; data.Add(std::move(cfg)); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); EXPECT_EQ(std::unordered_set({0, 2}), @@ -411,7 +412,7 @@ fn main() { DataMap data; data.Add(std::move(cfg)); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(src, str(got)); EXPECT_EQ(std::unordered_set(), @@ -482,7 +483,7 @@ fn main() { DataMap data; data.Add(std::move(cfg)); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); EXPECT_EQ(std::unordered_set({0}), diff --git a/src/transform/calculate_array_length.cc b/src/transform/calculate_array_length.cc index 05076850ad..e06b95914e 100644 --- a/src/transform/calculate_array_length.cc +++ b/src/transform/calculate_array_length.cc @@ -25,8 +25,9 @@ #include "src/sem/statement.h" #include "src/sem/struct.h" #include "src/sem/variable.h" -#include "src/utils/map.h" +#include "src/transform/simplify_pointers.h" #include "src/utils/hash.h" +#include "src/utils/map.h" TINT_INSTANTIATE_TYPEINFO(tint::transform::CalculateArrayLength); TINT_INSTANTIATE_TYPEINFO( @@ -72,6 +73,9 @@ CalculateArrayLength::~CalculateArrayLength() = default; void CalculateArrayLength::Run(CloneContext& ctx, const DataMap&, DataMap&) { auto& sem = ctx.src->Sem(); + if (!Requires(ctx)) { + return; + } // get_buffer_size_intrinsic() emits the function decorated with // BufferSizeIntrinsic that is transformed by the HLSL writer into a call to diff --git a/src/transform/calculate_array_length_test.cc b/src/transform/calculate_array_length_test.cc index 1e72a40ffa..897213c0d3 100644 --- a/src/transform/calculate_array_length_test.cc +++ b/src/transform/calculate_array_length_test.cc @@ -14,7 +14,9 @@ #include "src/transform/calculate_array_length.h" +#include "src/transform/simplify_pointers.h" #include "src/transform/test_helper.h" +#include "src/transform/unshadow.h" namespace tint { namespace transform { @@ -22,6 +24,18 @@ namespace { using CalculateArrayLengthTest = TransformTest; +TEST_F(CalculateArrayLengthTest, Error_MissingCalculateArrayLength) { + auto* src = ""; + + auto* expect = + "error: tint::transform::CalculateArrayLength depends on " + "tint::transform::SimplifyPointers but the dependency was not run"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + TEST_F(CalculateArrayLengthTest, Basic) { auto* src = R"( [[block]] @@ -59,7 +73,7 @@ fn main() { } )"; - auto got = Run(src); + auto got = Run(src); EXPECT_EQ(expect, str(got)); } @@ -105,7 +119,7 @@ fn main() { } )"; - auto got = Run(src); + auto got = Run(src); EXPECT_EQ(expect, str(got)); } @@ -149,7 +163,7 @@ fn main() { } )"; - auto got = Run(src); + auto got = Run(src); EXPECT_EQ(expect, str(got)); } @@ -206,7 +220,7 @@ fn main() { } )"; - auto got = Run(src); + auto got = Run(src); EXPECT_EQ(expect, str(got)); } @@ -274,7 +288,63 @@ fn main() { } )"; - auto got = Run(src); + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(CalculateArrayLengthTest, Shadowing) { + auto* src = R"( +[[block]] +struct SB { + x : i32; + arr : array; +}; + +[[group(0), binding(0)]] var a : SB; +[[group(0), binding(1)]] var b : SB; + +[[stage(compute), workgroup_size(1)]] +fn main() { + let x = &a; + var a : u32 = arrayLength(&a.arr); + { + var b : u32 = arrayLength(&((*x).arr)); + } +} +)"; + + auto* expect = + R"( +[[block]] +struct SB { + x : i32; + arr : array; +}; + +[[internal(intrinsic_buffer_size)]] +fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, result : ptr) + +[[group(0), binding(0)]] var a : SB; + +[[group(0), binding(1)]] var b : SB; + +[[stage(compute), workgroup_size(1)]] +fn main() { + var tint_symbol_1 : u32 = 0u; + tint_symbol(a, &(tint_symbol_1)); + let tint_symbol_2 : u32 = ((tint_symbol_1 - 4u) / 4u); + var a_1 : u32 = tint_symbol_2; + { + var tint_symbol_3 : u32 = 0u; + tint_symbol(a, &(tint_symbol_3)); + let tint_symbol_4 : u32 = ((tint_symbol_3 - 4u) / 4u); + var b_1 : u32 = tint_symbol_4; + } +} +)"; + + auto got = Run(src); EXPECT_EQ(expect, str(got)); } diff --git a/src/transform/canonicalize_entry_point_io.cc b/src/transform/canonicalize_entry_point_io.cc index af54484501..c78ac19099 100644 --- a/src/transform/canonicalize_entry_point_io.cc +++ b/src/transform/canonicalize_entry_point_io.cc @@ -23,6 +23,7 @@ #include "src/ast/disable_validation_decoration.h" #include "src/program_builder.h" #include "src/sem/function.h" +#include "src/transform/unshadow.h" TINT_INSTANTIATE_TYPEINFO(tint::transform::CanonicalizeEntryPointIO); TINT_INSTANTIATE_TYPEINFO(tint::transform::CanonicalizeEntryPointIO::Config); @@ -550,6 +551,10 @@ struct CanonicalizeEntryPointIO::State { void CanonicalizeEntryPointIO::Run(CloneContext& ctx, const DataMap& inputs, DataMap&) { + if (!Requires(ctx)) { + return; + } + auto* cfg = inputs.Get(); if (cfg == nullptr) { ctx.dst->Diagnostics().add_error( diff --git a/src/transform/canonicalize_entry_point_io_test.cc b/src/transform/canonicalize_entry_point_io_test.cc index a011e85836..3ea25c1f4d 100644 --- a/src/transform/canonicalize_entry_point_io_test.cc +++ b/src/transform/canonicalize_entry_point_io_test.cc @@ -15,6 +15,7 @@ #include "src/transform/canonicalize_entry_point_io.h" #include "src/transform/test_helper.h" +#include "src/transform/unshadow.h" namespace tint { namespace transform { @@ -22,6 +23,18 @@ namespace { using CanonicalizeEntryPointIOTest = TransformTest; +TEST_F(CanonicalizeEntryPointIOTest, Error_MissingUnshadow) { + auto* src = ""; + + auto* expect = + "error: tint::transform::CanonicalizeEntryPointIO depends on " + "tint::transform::Unshadow but the dependency was not run"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + TEST_F(CanonicalizeEntryPointIOTest, Error_MissingTransformData) { auto* src = ""; @@ -29,7 +42,7 @@ TEST_F(CanonicalizeEntryPointIOTest, Error_MissingTransformData) { "error: missing transform data for " "tint::transform::CanonicalizeEntryPointIO"; - auto got = Run(src); + auto got = Run(src); EXPECT_EQ(expect, str(got)); } @@ -52,7 +65,7 @@ fn comp_main() { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kMsl); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -87,7 +100,7 @@ fn frag_main() { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kSpirv); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -123,7 +136,7 @@ fn frag_main([[builtin(position)]] coord : vec4, tint_symbol : tint_symbol_ DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kMsl); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -161,7 +174,7 @@ fn frag_main(tint_symbol : tint_symbol_1) { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kHlsl); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -197,7 +210,7 @@ fn frag_main(tint_symbol : tint_symbol_1) { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kMsl); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -251,7 +264,7 @@ fn frag_main() { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kSpirv); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -306,7 +319,7 @@ fn frag_main([[builtin(position)]] coord : vec4, tint_symbol : tint_symbol_ DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kMsl); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -363,7 +376,7 @@ fn frag_main(tint_symbol : tint_symbol_1) { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kHlsl); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -393,7 +406,7 @@ fn frag_main() { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kSpirv); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -428,7 +441,7 @@ fn frag_main() -> tint_symbol { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kMsl); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -463,7 +476,7 @@ fn frag_main() -> tint_symbol { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kHlsl); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -519,7 +532,7 @@ fn frag_main() { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kSpirv); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -580,7 +593,7 @@ fn frag_main() -> tint_symbol { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kMsl); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -641,7 +654,7 @@ fn frag_main() -> tint_symbol { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kHlsl); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -709,7 +722,7 @@ fn frag_main2() { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kSpirv); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -783,7 +796,7 @@ fn frag_main2(tint_symbol_2 : tint_symbol_3) { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kMsl); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -857,7 +870,7 @@ fn frag_main2(tint_symbol_2 : tint_symbol_3) { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kHlsl); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -925,7 +938,7 @@ fn frag_main1(tint_symbol : tint_symbol_1) { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kMsl); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -1012,7 +1025,7 @@ fn frag_main(tint_symbol : tint_symbol_1) -> tint_symbol_2 { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kMsl); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -1104,7 +1117,7 @@ fn frag_main(tint_symbol_1 : tint_symbol_2) { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kHlsl); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -1235,7 +1248,7 @@ fn frag_main() { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kSpirv); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -1300,7 +1313,7 @@ fn main2() -> tint_symbol_1 { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kHlsl); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -1370,7 +1383,7 @@ fn frag_main(tint_symbol : tint_symbol_1) -> tint_symbol_2 { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kHlsl); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -1475,7 +1488,7 @@ fn frag_main(tint_symbol_1 : tint_symbol_2) { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kHlsl); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -1505,7 +1518,7 @@ fn tint_symbol_1(tint_symbol : tint_symbol_2) { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kMsl); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -1538,7 +1551,7 @@ fn frag_main() -> tint_symbol { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kMsl, 0x03u); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -1573,7 +1586,7 @@ fn frag_main() -> tint_symbol { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kMsl, 0x03u); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -1608,7 +1621,7 @@ fn frag_main() -> tint_symbol { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kMsl, 0x03u); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -1646,7 +1659,7 @@ fn frag_main() -> tint_symbol { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kMsl, 0x03u); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -1699,7 +1712,7 @@ fn frag_main() -> tint_symbol { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kMsl, 0x03u); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -1751,7 +1764,7 @@ fn frag_main() -> tint_symbol { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kMsl, 0x03u); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -1841,7 +1854,7 @@ fn comp_main1() { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kMsl, 0x03u); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -1892,7 +1905,7 @@ fn frag_main() -> tint_symbol { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kMsl, 0x03); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -1926,7 +1939,7 @@ fn vert_main() { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kSpirv, 0xFFFFFFFF, true); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -1964,7 +1977,7 @@ fn vert_main() -> tint_symbol { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kMsl, 0xFFFFFFFF, true); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -2005,7 +2018,7 @@ fn vert_main() { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kSpirv, 0xFFFFFFFF, true); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -2051,7 +2064,7 @@ fn vert_main() -> tint_symbol { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kMsl, 0xFFFFFFFF, true); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -2129,7 +2142,7 @@ fn vert_main() { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kSpirv, 0xFFFFFFFF, true); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -2205,7 +2218,7 @@ fn vert_main(tint_symbol : tint_symbol_1) -> tint_symbol_2 { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kMsl, 0xFFFFFFFF, true); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -2281,7 +2294,7 @@ fn vert_main(tint_symbol : tint_symbol_1) -> tint_symbol_2 { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kHlsl, 0xFFFFFFFF, true); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } @@ -2317,7 +2330,7 @@ fn main() { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kSpirv); - auto got = Run(src, data); + auto got = Run(src, data); EXPECT_EQ(expect, str(got)); } diff --git a/src/transform/decompose_strided_matrix_test.cc b/src/transform/decompose_strided_matrix_test.cc index 714e8a1818..b7a22e4c5d 100644 --- a/src/transform/decompose_strided_matrix_test.cc +++ b/src/transform/decompose_strided_matrix_test.cc @@ -22,6 +22,7 @@ #include "src/program_builder.h" #include "src/transform/simplify_pointers.h" #include "src/transform/test_helper.h" +#include "src/transform/unshadow.h" namespace tint { namespace transform { @@ -34,7 +35,7 @@ TEST_F(DecomposeStridedMatrixTest, Empty) { auto* src = R"()"; auto* expect = src; - auto got = Run(src); + auto got = Run(src); EXPECT_EQ(expect, str(got)); } @@ -109,8 +110,8 @@ fn f() { } )"; - auto got = - Run(Program(std::move(b))); + auto got = Run( + Program(std::move(b))); EXPECT_EQ(expect, str(got)); } @@ -171,8 +172,8 @@ fn f() { } )"; - auto got = - Run(Program(std::move(b))); + auto got = Run( + Program(std::move(b))); EXPECT_EQ(expect, str(got)); } @@ -234,8 +235,8 @@ fn f() { } )"; - auto got = - Run(Program(std::move(b))); + auto got = Run( + Program(std::move(b))); EXPECT_EQ(expect, str(got)); } @@ -300,8 +301,8 @@ fn f() { } )"; - auto got = - Run(Program(std::move(b))); + auto got = Run( + Program(std::move(b))); EXPECT_EQ(expect, str(got)); } @@ -362,8 +363,8 @@ fn f() { } )"; - auto got = - Run(Program(std::move(b))); + auto got = Run( + Program(std::move(b))); EXPECT_EQ(expect, str(got)); } @@ -429,8 +430,8 @@ fn f() { } )"; - auto got = - Run(Program(std::move(b))); + auto got = Run( + Program(std::move(b))); EXPECT_EQ(expect, str(got)); } @@ -491,8 +492,8 @@ fn f() { } )"; - auto got = - Run(Program(std::move(b))); + auto got = Run( + Program(std::move(b))); EXPECT_EQ(expect, str(got)); } @@ -580,8 +581,8 @@ fn f() { } )"; - auto got = - Run(Program(std::move(b))); + auto got = Run( + Program(std::move(b))); EXPECT_EQ(expect, str(got)); } @@ -637,8 +638,8 @@ fn f() { } )"; - auto got = - Run(Program(std::move(b))); + auto got = Run( + Program(std::move(b))); EXPECT_EQ(expect, str(got)); } @@ -695,8 +696,8 @@ fn f() { } )"; - auto got = - Run(Program(std::move(b))); + auto got = Run( + Program(std::move(b))); EXPECT_EQ(expect, str(got)); } diff --git a/src/transform/fold_trivial_single_use_lets_test.cc b/src/transform/fold_trivial_single_use_lets_test.cc index b571ce8fdd..42aae29163 100644 --- a/src/transform/fold_trivial_single_use_lets_test.cc +++ b/src/transform/fold_trivial_single_use_lets_test.cc @@ -35,13 +35,13 @@ TEST_F(FoldTrivialSingleUseLetsTest, Single) { auto* src = R"( fn f() { let x = 1; - ignore(x); + _ = x; } )"; auto* expect = R"( fn f() { - ignore(1); + _ = 1; } )"; @@ -56,13 +56,13 @@ fn f() { let x = 1; let y = 2; let z = 3; - ignore(x + y + z); + _ = x + y + z; } )"; auto* expect = R"( fn f() { - ignore(((1 + 2) + 3)); + _ = ((1 + 2) + 3); } )"; @@ -77,13 +77,13 @@ fn f() { let x = 1; let y = x; let z = y; - ignore(z); + _ = z; } )"; auto* expect = R"( fn f() { - ignore(1); + _ = 1; } )"; @@ -101,7 +101,7 @@ fn function_with_posssible_side_effect() -> i32 { fn f() { let x = 1; let y = function_with_posssible_side_effect(); - ignore((x + y)); + _ = (x + y); } )"; @@ -117,7 +117,7 @@ TEST_F(FoldTrivialSingleUseLetsTest, NoFold_UseInSubBlock) { fn f() { let x = 1; { - ignore(x); + _ = x; } } )"; @@ -133,7 +133,26 @@ TEST_F(FoldTrivialSingleUseLetsTest, NoFold_MultipleUses) { auto* src = R"( fn f() { let x = 1; - ignore((x + x)); + _ = (x + x); +} +)"; + + auto* expect = src; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(FoldTrivialSingleUseLetsTest, NoFold_Shadowing) { + auto* src = R"( +fn f() { + var y = 1; + let x = y; + { + let y = false; + _ = (x + x); + } } )"; diff --git a/src/transform/glsl.cc b/src/transform/glsl.cc index 4b4512d196..100023465d 100644 --- a/src/transform/glsl.cc +++ b/src/transform/glsl.cc @@ -29,6 +29,7 @@ #include "src/transform/remove_phonies.h" #include "src/transform/simplify_pointers.h" #include "src/transform/single_entry_point.h" +#include "src/transform/unshadow.h" #include "src/transform/zero_init_workgroup_memory.h" TINT_INSTANTIATE_TYPEINFO(tint::transform::Glsl); @@ -46,6 +47,8 @@ Output Glsl::Run(const Program* in, const DataMap& inputs) { auto* cfg = inputs.Get(); + manager.Add(); + // Attempt to convert `loop`s into for-loops. This is to try and massage the // output into something that will not cause FXC to choke or misbehave. manager.Add(); diff --git a/src/transform/num_workgroups_from_uniform_test.cc b/src/transform/num_workgroups_from_uniform_test.cc index 16973a400b..1d719417ea 100644 --- a/src/transform/num_workgroups_from_uniform_test.cc +++ b/src/transform/num_workgroups_from_uniform_test.cc @@ -18,6 +18,7 @@ #include "src/transform/canonicalize_entry_point_io.h" #include "src/transform/test_helper.h" +#include "src/transform/unshadow.h" namespace tint { namespace transform { @@ -35,7 +36,8 @@ TEST_F(NumWorkgroupsFromUniformTest, Error_MissingTransformData) { DataMap data; data.Add( CanonicalizeEntryPointIO::ShaderStyle::kHlsl); - auto got = Run(src, data); + auto got = Run( + src, data); EXPECT_EQ(expect, str(got)); } @@ -87,7 +89,8 @@ fn main() { data.Add( CanonicalizeEntryPointIO::ShaderStyle::kHlsl); data.Add(sem::BindingPoint{0, 30u}); - auto got = Run(src, data); + auto got = Run( + src, data); EXPECT_EQ(expect, str(got)); } @@ -133,7 +136,8 @@ fn main() { data.Add( CanonicalizeEntryPointIO::ShaderStyle::kHlsl); data.Add(sem::BindingPoint{0, 30u}); - auto got = Run(src, data); + auto got = Run( + src, data); EXPECT_EQ(expect, str(got)); } @@ -190,7 +194,8 @@ fn main(tint_symbol : tint_symbol_1) { data.Add( CanonicalizeEntryPointIO::ShaderStyle::kHlsl); data.Add(sem::BindingPoint{0, 30u}); - auto got = Run(src, data); + auto got = Run( + src, data); EXPECT_EQ(expect, str(got)); } @@ -291,7 +296,8 @@ fn main3() { data.Add( CanonicalizeEntryPointIO::ShaderStyle::kHlsl); data.Add(sem::BindingPoint{0, 30u}); - auto got = Run(src, data); + auto got = Run( + src, data); EXPECT_EQ(expect, str(got)); } @@ -333,7 +339,8 @@ fn main(tint_symbol : tint_symbol_1) { data.Add( CanonicalizeEntryPointIO::ShaderStyle::kHlsl); data.Add(sem::BindingPoint{0, 30u}); - auto got = Run(src, data); + auto got = Run( + src, data); EXPECT_EQ(expect, str(got)); } diff --git a/src/transform/simplify_pointers.cc b/src/transform/simplify_pointers.cc index 6ab7b6a418..b789e87cc9 100644 --- a/src/transform/simplify_pointers.cc +++ b/src/transform/simplify_pointers.cc @@ -24,6 +24,7 @@ #include "src/sem/function.h" #include "src/sem/statement.h" #include "src/sem/variable.h" +#include "src/transform/unshadow.h" TINT_INSTANTIATE_TYPEINFO(tint::transform::SimplifyPointers); @@ -231,6 +232,9 @@ SimplifyPointers::SimplifyPointers() = default; SimplifyPointers::~SimplifyPointers() = default; void SimplifyPointers::Run(CloneContext& ctx, const DataMap&, DataMap&) { + if (!Requires(ctx)) { + return; + } State(ctx).Run(); } diff --git a/src/transform/simplify_pointers.h b/src/transform/simplify_pointers.h index 9f00f9c0e3..1f8cf71be1 100644 --- a/src/transform/simplify_pointers.h +++ b/src/transform/simplify_pointers.h @@ -15,16 +15,14 @@ #ifndef SRC_TRANSFORM_SIMPLIFY_POINTERS_H_ #define SRC_TRANSFORM_SIMPLIFY_POINTERS_H_ -#include -#include - #include "src/transform/transform.h" namespace tint { namespace transform { /// SimplifyPointers is a Transform that moves all usage of function-scope -/// `let` statements of a pointer type into their places of usage. +/// `let` statements of a pointer type into their places of usage, while also +/// simplifying any chains of address-of or indirections operators. /// /// Parameters of a pointer type are not adjusted. /// diff --git a/src/transform/simplify_pointers_test.cc b/src/transform/simplify_pointers_test.cc index b2816abfcd..d6fde98770 100644 --- a/src/transform/simplify_pointers_test.cc +++ b/src/transform/simplify_pointers_test.cc @@ -14,11 +14,8 @@ #include "src/transform/simplify_pointers.h" -#include -#include -#include - #include "src/transform/test_helper.h" +#include "src/transform/unshadow.h" namespace tint { namespace transform { @@ -26,11 +23,23 @@ namespace { using SimplifyPointersTest = TransformTest; +TEST_F(SimplifyPointersTest, Error_MissingSimplifyPointers) { + auto* src = ""; + + auto* expect = + "error: tint::transform::SimplifyPointers depends on " + "tint::transform::Unshadow but the dependency was not run"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + TEST_F(SimplifyPointersTest, EmptyModule) { auto* src = ""; auto* expect = ""; - auto got = Run(src); + auto got = Run(src); EXPECT_EQ(expect, str(got)); } @@ -51,7 +60,7 @@ fn f() { } )"; - auto got = Run(src); + auto got = Run(src); EXPECT_EQ(expect, str(got)); } @@ -79,7 +88,7 @@ fn f() { } )"; - auto got = Run(src); + auto got = Run(src); EXPECT_EQ(expect, str(got)); } @@ -103,7 +112,7 @@ fn f() { } )"; - auto got = Run(src); + auto got = Run(src); EXPECT_EQ(expect, str(got)); } @@ -126,7 +135,7 @@ fn f() { } )"; - auto got = Run(src); + auto got = Run(src); EXPECT_EQ(expect, str(got)); } @@ -180,7 +189,7 @@ fn matrix() { } )"; - auto got = Run(src); + auto got = Run(src); EXPECT_EQ(expect, str(got)); } @@ -201,7 +210,7 @@ fn f() { } )"; - auto got = Run(src); + auto got = Run(src); EXPECT_EQ(expect, str(got)); } @@ -229,7 +238,7 @@ fn f() { } )"; - auto got = Run(src); + auto got = Run(src); EXPECT_EQ(expect, str(got)); } @@ -264,7 +273,7 @@ fn main() { } )"; - auto got = Run(src); + auto got = Run(src); EXPECT_EQ(expect, str(got)); } @@ -330,34 +339,42 @@ fn f() { } )"; - auto got = Run(src); + auto got = Run(src); EXPECT_EQ(expect, str(got)); } -// TODO(crbug.com/tint/819): Enable when we support inter-scope shadowing. -TEST_F(SimplifyPointersTest, DISABLED_ModificationAfterInline) { +TEST_F(SimplifyPointersTest, ShadowPointer) { auto* src = R"( -fn x(p : ptr) -> i32 { - return *p; -} +var a : array; -fn f() { - var i : i32 = 1; - let p : ptr = &i; - if (true) { - var i : i32 = 2; - x(p); +[[stage(compute), workgroup_size(1)]] +fn main() { + let x = &a; + var a : i32 = (*x)[0]; + { + var a : i32 = (*x)[1]; } } )"; - auto* expect = R"()"; + auto* expect = R"( +var a : array; - auto got = Run(src); +[[stage(compute), workgroup_size(1)]] +fn main() { + var a_1 : i32 = a[0]; + { + var a_2 : i32 = a[1]; + } +} +)"; + + auto got = Run(src); EXPECT_EQ(expect, str(got)); } + } // namespace } // namespace transform } // namespace tint diff --git a/src/transform/unshadow.cc b/src/transform/unshadow.cc new file mode 100644 index 0000000000..a7bc66d3a1 --- /dev/null +++ b/src/transform/unshadow.cc @@ -0,0 +1,99 @@ +// 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/unshadow.h" + +#include +#include +#include + +#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" + +TINT_INSTANTIATE_TYPEINFO(tint::transform::Unshadow); + +namespace tint { +namespace transform { + +/// The PIMPL state for the Unshadow transform +struct Unshadow::State { + /// The clone context + CloneContext& ctx; + + /// Constructor + /// @param context the clone context + explicit State(CloneContext& context) : ctx(context) {} + + /// Performs the transformation + void Run() { + auto& sem = ctx.src->Sem(); + + // Maps a variable to its new name. + std::unordered_map renamed_to; + + auto rename = [&](const sem::Variable* var) -> const ast::Variable* { + auto* decl = var->Declaration(); + auto name = ctx.src->Symbols().NameFor(decl->symbol); + auto symbol = ctx.dst->Symbols().New(name); + renamed_to.emplace(var, symbol); + + auto source = ctx.Clone(decl->source); + auto* type = ctx.Clone(decl->type); + auto* constructor = ctx.Clone(decl->constructor); + auto decorations = ctx.Clone(decl->decorations); + return ctx.dst->create( + source, symbol, decl->declared_storage_class, decl->declared_access, + type, decl->is_const, constructor, decorations); + }; + + ctx.ReplaceAll([&](const ast::Variable* var) -> const ast::Variable* { + if (auto* local = sem.Get(var)) { + if (local->Shadows()) { + return rename(local); + } + } + if (auto* param = sem.Get(var)) { + if (param->Shadows()) { + return rename(param); + } + } + return nullptr; + }); + ctx.ReplaceAll([&](const ast::IdentifierExpression* ident) + -> const tint::ast::IdentifierExpression* { + if (auto* user = sem.Get(ident)) { + auto it = renamed_to.find(user->Variable()); + if (it != renamed_to.end()) { + return ctx.dst->Expr(it->second); + } + } + return nullptr; + }); + ctx.Clone(); + } +}; + +Unshadow::Unshadow() = default; + +Unshadow::~Unshadow() = default; + +void Unshadow::Run(CloneContext& ctx, const DataMap&, DataMap&) { + State(ctx).Run(); +} + +} // namespace transform +} // namespace tint diff --git a/src/transform/unshadow.h b/src/transform/unshadow.h new file mode 100644 index 0000000000..7c16bcf2c7 --- /dev/null +++ b/src/transform/unshadow.h @@ -0,0 +1,48 @@ +// 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_UNSHADOW_H_ +#define SRC_TRANSFORM_UNSHADOW_H_ + +#include "src/transform/transform.h" + +namespace tint { +namespace transform { + +/// Unshadow is a Transform that renames any variables that shadow another +/// variable. +class Unshadow : public Castable { + public: + /// Constructor + Unshadow(); + + /// Destructor + ~Unshadow() override; + + protected: + struct State; + + /// Runs the transform using the CloneContext built for transforming a + /// program. Run() is responsible for calling Clone() on the CloneContext. + /// @param ctx the CloneContext primed with the input program and + /// ProgramBuilder + /// @param inputs optional extra transform-specific input data + /// @param outputs optional extra transform-specific output data + void Run(CloneContext& ctx, const DataMap& inputs, DataMap& outputs) override; +}; + +} // namespace transform +} // namespace tint + +#endif // SRC_TRANSFORM_UNSHADOW_H_ diff --git a/src/transform/unshadow_test.cc b/src/transform/unshadow_test.cc new file mode 100644 index 0000000000..2a4cbed454 --- /dev/null +++ b/src/transform/unshadow_test.cc @@ -0,0 +1,397 @@ +// 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/unshadow.h" + +#include "src/transform/test_helper.h" + +namespace tint { +namespace transform { +namespace { + +using UnshadowTest = TransformTest; + +TEST_F(UnshadowTest, EmptyModule) { + auto* src = ""; + auto* expect = ""; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(UnshadowTest, Noop) { + auto* src = R"( +var a : i32; + +let b : i32 = 1; + +fn F(c : i32) { + var d : i32; + let e : i32 = 1; + { + var f : i32; + let g : i32 = 1; + } +} +)"; + + auto* expect = src; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(UnshadowTest, LocalShadowsAlias) { + auto* src = R"( +type a = i32; + +fn X() { + var a = false; +} + +fn Y() { + let a = true; +} +)"; + + auto* expect = R"( +type a = i32; + +fn X() { + var a_1 = false; +} + +fn Y() { + let a_2 = true; +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(UnshadowTest, LocalShadowsStruct) { + auto* src = R"( +struct a { + m : i32; +}; + +fn X() { + var a = true; +} + +fn Y() { + let a = false; +} +)"; + + auto* expect = R"( +struct a { + m : i32; +}; + +fn X() { + var a_1 = true; +} + +fn Y() { + let a_2 = false; +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(UnshadowTest, LocalShadowsFunction) { + auto* src = R"( +fn a() { + var a = true; +} + +fn b() { + let b = false; +} +)"; + + auto* expect = R"( +fn a() { + var a_1 = true; +} + +fn b() { + let b_1 = false; +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(UnshadowTest, LocalShadowsGlobalVar) { + auto* src = R"( +var a : i32; + +fn X() { + var a = (a == 123); +} + +fn Y() { + let a = (a == 321); +} +)"; + + auto* expect = R"( +var a : i32; + +fn X() { + var a_1 = (a == 123); +} + +fn Y() { + let a_2 = (a == 321); +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(UnshadowTest, LocalShadowsGlobalLet) { + auto* src = R"( +let a : i32 = 1; + +fn X() { + var a = (a == 123); +} + +fn Y() { + let a = (a == 321); +} +)"; + + auto* expect = R"( +let a : i32 = 1; + +fn X() { + var a_1 = (a == 123); +} + +fn Y() { + let a_2 = (a == 321); +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(UnshadowTest, LocalShadowsLocalVar) { + auto* src = R"( +fn X() { + var a : i32; + { + var a = (a == 123); + } + { + let a = (a == 321); + } +} +)"; + + auto* expect = R"( +fn X() { + var a : i32; + { + var a_1 = (a == 123); + } + { + let a_2 = (a == 321); + } +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(UnshadowTest, LocalShadowsLocalLet) { + auto* src = R"( +fn X() { + let a = 1; + { + var a = (a == 123); + } + { + let a = (a == 321); + } +} +)"; + + auto* expect = R"( +fn X() { + let a = 1; + { + var a_1 = (a == 123); + } + { + let a_2 = (a == 321); + } +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(UnshadowTest, LocalShadowsParam) { + auto* src = R"( +fn F(a : i32) { + { + var a = (a == 123); + } + { + let a = (a == 321); + } +} +)"; + + auto* expect = R"( +fn F(a : i32) { + { + var a_1 = (a == 123); + } + { + let a_2 = (a == 321); + } +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(UnshadowTest, ParamShadowsFunction) { + auto* src = R"( +fn a(a : i32) { + { + var a = (a == 123); + } + { + let a = (a == 321); + } +} +)"; + + auto* expect = R"( +fn a(a_1 : i32) { + { + var a_2 = (a_1 == 123); + } + { + let a_3 = (a_1 == 321); + } +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(UnshadowTest, ParamShadowsGlobalVar) { + auto* src = R"( +var a : i32; + +fn F(a : bool) { +} +)"; + + auto* expect = R"( +var a : i32; + +fn F(a_1 : bool) { +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(UnshadowTest, ParamShadowsGlobalLet) { + auto* src = R"( +let a : i32 = 1; + +fn F(a : bool) { +} +)"; + + auto* expect = R"( +let a : i32 = 1; + +fn F(a_1 : bool) { +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(UnshadowTest, ParamShadowsAlias) { + auto* src = R"( +type a = i32; + +fn F(a : a) { + { + var a = (a == 123); + } + { + let a = (a == 321); + } +} +)"; + + auto* expect = R"( +type a = i32; + +fn F(a_1 : a) { + { + var a_2 = (a_1 == 123); + } + { + let a_3 = (a_1 == 321); + } +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +} // namespace +} // namespace transform +} // namespace tint diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc index 3bf5394ba4..1a24ed2138 100644 --- a/src/writer/hlsl/generator_impl.cc +++ b/src/writer/hlsl/generator_impl.cc @@ -58,6 +58,7 @@ #include "src/transform/promote_initializers_to_const_var.h" #include "src/transform/remove_phonies.h" #include "src/transform/simplify_pointers.h" +#include "src/transform/unshadow.h" #include "src/transform/zero_init_workgroup_memory.h" #include "src/utils/defer.h" #include "src/utils/map.h" @@ -135,6 +136,8 @@ SanitizedResult Sanitize( array_length_from_uniform_cfg.bindpoint_to_size_index = array_length_from_uniform.bindpoint_to_size_index; + manager.Add(); + // Attempt to convert `loop`s into for-loops. This is to try and massage the // output into something that will not cause FXC to choke or misbehave. manager.Add(); diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc index 6054609a99..077beaee5b 100644 --- a/src/writer/msl/generator_impl.cc +++ b/src/writer/msl/generator_impl.cc @@ -66,6 +66,7 @@ #include "src/transform/promote_initializers_to_const_var.h" #include "src/transform/remove_phonies.h" #include "src/transform/simplify_pointers.h" +#include "src/transform/unshadow.h" #include "src/transform/vectorize_scalar_matrix_constructors.h" #include "src/transform/wrap_arrays_in_structs.h" #include "src/transform/zero_init_workgroup_memory.h" @@ -152,6 +153,8 @@ SanitizedResult Sanitize( transform::CanonicalizeEntryPointIO::ShaderStyle::kMsl, fixed_sample_mask, emit_vertex_point_size); + manager.Add(); + if (!disable_workgroup_init) { // ZeroInitWorkgroupMemory must come before CanonicalizeEntryPointIO as // ZeroInitWorkgroupMemory may inject new builtin parameters. diff --git a/src/writer/spirv/builder.cc b/src/writer/spirv/builder.cc index b911359576..1bbb2b6ba7 100644 --- a/src/writer/spirv/builder.cc +++ b/src/writer/spirv/builder.cc @@ -47,6 +47,7 @@ #include "src/transform/for_loop_to_loop.h" #include "src/transform/manager.h" #include "src/transform/simplify_pointers.h" +#include "src/transform/unshadow.h" #include "src/transform/vectorize_scalar_matrix_constructors.h" #include "src/transform/zero_init_workgroup_memory.h" #include "src/utils/defer.h" @@ -266,6 +267,7 @@ SanitizedResult Sanitize(const Program* in, transform::Manager manager; transform::DataMap data; + manager.Add(); if (!disable_workgroup_init) { manager.Add(); } diff --git a/test.wgsl b/test.wgsl new file mode 100644 index 0000000000..827d20065a --- /dev/null +++ b/test.wgsl @@ -0,0 +1,15 @@ +[[block]] +struct SB { + x : i32; + arr : array; +}; + +[[group(0), binding(0)]] var sb : SB; + +[[stage(compute), workgroup_size(1)]] +fn main() { + var a : u32 = arrayLength(&sb.arr); + let p = &sb; + let sb2 = p; + var b : u32 = arrayLength(&((*sb2).arr)); +} diff --git a/test/BUILD.gn b/test/BUILD.gn index 41e106913b..a634ea416e 100644 --- a/test/BUILD.gn +++ b/test/BUILD.gn @@ -327,6 +327,7 @@ tint_unittests_source_set("tint_unittests_transform_src") { "../src/transform/single_entry_point_test.cc", "../src/transform/test_helper.h", "../src/transform/transform_test.cc", + "../src/transform/unshadow_test.cc", "../src/transform/vectorize_scalar_matrix_constructors_test.cc", "../src/transform/vertex_pulling_test.cc", "../src/transform/wrap_arrays_in_structs_test.cc", diff --git a/test/bug/tint/825.wgsl.expected.hlsl b/test/bug/tint/825.wgsl.expected.hlsl index 1e1f290b43..febe29bd51 100644 --- a/test/bug/tint/825.wgsl.expected.hlsl +++ b/test/bug/tint/825.wgsl.expected.hlsl @@ -7,5 +7,5 @@ void f() { int i = 0; int j = 0; float2x2 m = float2x2(float2(1.0f, 2.0f), float2(3.0f, 4.0f)); - const float f = m[i][j]; + const float f_1 = m[i][j]; } diff --git a/test/bug/tint/825.wgsl.expected.msl b/test/bug/tint/825.wgsl.expected.msl index 58b800ab3d..463024342a 100644 --- a/test/bug/tint/825.wgsl.expected.msl +++ b/test/bug/tint/825.wgsl.expected.msl @@ -5,6 +5,6 @@ void f() { int i = 0; int j = 0; float2x2 m = float2x2(float2(1.0f, 2.0f), float2(3.0f, 4.0f)); - float const f = m[i][j]; + float const f_1 = m[i][j]; } diff --git a/test/shadowing/alias/let.wgsl b/test/shadowing/alias/let.wgsl new file mode 100644 index 0000000000..558efe64c5 --- /dev/null +++ b/test/shadowing/alias/let.wgsl @@ -0,0 +1,10 @@ +type a = i32; + +fn f() { + { + let a : a = a(); + let b = a; + } + let a : a = a(); + let b = a; +} diff --git a/test/shadowing/alias/let.wgsl.expected.hlsl b/test/shadowing/alias/let.wgsl.expected.hlsl new file mode 100644 index 0000000000..ad5e10450e --- /dev/null +++ b/test/shadowing/alias/let.wgsl.expected.hlsl @@ -0,0 +1,13 @@ +[numthreads(1, 1, 1)] +void unused_entry_point() { + return; +} + +void f() { + { + const int a_1 = 0; + const int b = a_1; + } + const int a_2 = 0; + const int b = a_2; +} diff --git a/test/shadowing/alias/let.wgsl.expected.msl b/test/shadowing/alias/let.wgsl.expected.msl new file mode 100644 index 0000000000..292a0ba14e --- /dev/null +++ b/test/shadowing/alias/let.wgsl.expected.msl @@ -0,0 +1,13 @@ +#include + +using namespace metal; + +void f() { + { + int const a_1 = int(); + int const b = a_1; + } + int const a_2 = int(); + int const b = a_2; +} + diff --git a/test/shadowing/alias/let.wgsl.expected.spvasm b/test/shadowing/alias/let.wgsl.expected.spvasm new file mode 100644 index 0000000000..56bd05ba41 --- /dev/null +++ b/test/shadowing/alias/let.wgsl.expected.spvasm @@ -0,0 +1,23 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 9 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %unused_entry_point "unused_entry_point" + OpExecutionMode %unused_entry_point LocalSize 1 1 1 + OpName %unused_entry_point "unused_entry_point" + OpName %f "f" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %8 = OpConstantNull %int +%unused_entry_point = OpFunction %void None %1 + %4 = OpLabel + OpReturn + OpFunctionEnd + %f = OpFunction %void None %1 + %6 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/shadowing/alias/let.wgsl.expected.wgsl b/test/shadowing/alias/let.wgsl.expected.wgsl new file mode 100644 index 0000000000..558efe64c5 --- /dev/null +++ b/test/shadowing/alias/let.wgsl.expected.wgsl @@ -0,0 +1,10 @@ +type a = i32; + +fn f() { + { + let a : a = a(); + let b = a; + } + let a : a = a(); + let b = a; +} diff --git a/test/shadowing/alias/param.wgsl b/test/shadowing/alias/param.wgsl new file mode 100644 index 0000000000..42d95bf49c --- /dev/null +++ b/test/shadowing/alias/param.wgsl @@ -0,0 +1,5 @@ +type a = i32; + +fn f(a : a) { + let b = a; +} diff --git a/test/shadowing/alias/param.wgsl.expected.hlsl b/test/shadowing/alias/param.wgsl.expected.hlsl new file mode 100644 index 0000000000..b48a6b1132 --- /dev/null +++ b/test/shadowing/alias/param.wgsl.expected.hlsl @@ -0,0 +1,8 @@ +[numthreads(1, 1, 1)] +void unused_entry_point() { + return; +} + +void f(int a_1) { + const int b = a_1; +} diff --git a/test/shadowing/alias/param.wgsl.expected.msl b/test/shadowing/alias/param.wgsl.expected.msl new file mode 100644 index 0000000000..770f18edc9 --- /dev/null +++ b/test/shadowing/alias/param.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; + +void f(int a_1) { + int const b = a_1; +} + diff --git a/test/shadowing/alias/param.wgsl.expected.spvasm b/test/shadowing/alias/param.wgsl.expected.spvasm new file mode 100644 index 0000000000..18181e653d --- /dev/null +++ b/test/shadowing/alias/param.wgsl.expected.spvasm @@ -0,0 +1,25 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 10 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %unused_entry_point "unused_entry_point" + OpExecutionMode %unused_entry_point LocalSize 1 1 1 + OpName %unused_entry_point "unused_entry_point" + OpName %f "f" + OpName %a_1 "a_1" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %5 = OpTypeFunction %void %int +%unused_entry_point = OpFunction %void None %1 + %4 = OpLabel + OpReturn + OpFunctionEnd + %f = OpFunction %void None %5 + %a_1 = OpFunctionParameter %int + %9 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/shadowing/alias/param.wgsl.expected.wgsl b/test/shadowing/alias/param.wgsl.expected.wgsl new file mode 100644 index 0000000000..42d95bf49c --- /dev/null +++ b/test/shadowing/alias/param.wgsl.expected.wgsl @@ -0,0 +1,5 @@ +type a = i32; + +fn f(a : a) { + let b = a; +} diff --git a/test/shadowing/alias/var.wgsl b/test/shadowing/alias/var.wgsl new file mode 100644 index 0000000000..10e0c25973 --- /dev/null +++ b/test/shadowing/alias/var.wgsl @@ -0,0 +1,10 @@ +type a = i32; + +fn f() { + { + var a : a = a(); + var b = a; + } + var a : a = a(); + var b = a; +} diff --git a/test/shadowing/alias/var.wgsl.expected.hlsl b/test/shadowing/alias/var.wgsl.expected.hlsl new file mode 100644 index 0000000000..7b1a2a095e --- /dev/null +++ b/test/shadowing/alias/var.wgsl.expected.hlsl @@ -0,0 +1,13 @@ +[numthreads(1, 1, 1)] +void unused_entry_point() { + return; +} + +void f() { + { + int a_1 = 0; + int b = a_1; + } + int a_2 = 0; + int b = a_2; +} diff --git a/test/shadowing/alias/var.wgsl.expected.msl b/test/shadowing/alias/var.wgsl.expected.msl new file mode 100644 index 0000000000..9958607f26 --- /dev/null +++ b/test/shadowing/alias/var.wgsl.expected.msl @@ -0,0 +1,13 @@ +#include + +using namespace metal; + +void f() { + { + int a_1 = int(); + int b = a_1; + } + int a_2 = int(); + int b = a_2; +} + diff --git a/test/shadowing/alias/var.wgsl.expected.spvasm b/test/shadowing/alias/var.wgsl.expected.spvasm new file mode 100644 index 0000000000..0fec746dbd --- /dev/null +++ b/test/shadowing/alias/var.wgsl.expected.spvasm @@ -0,0 +1,38 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 16 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %unused_entry_point "unused_entry_point" + OpExecutionMode %unused_entry_point LocalSize 1 1 1 + OpName %unused_entry_point "unused_entry_point" + OpName %f "f" + OpName %a_1 "a_1" + OpName %b "b" + OpName %a_2 "a_2" + OpName %b_0 "b" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %8 = OpConstantNull %int +%_ptr_Function_int = OpTypePointer Function %int +%unused_entry_point = OpFunction %void None %1 + %4 = OpLabel + OpReturn + OpFunctionEnd + %f = OpFunction %void None %1 + %6 = OpLabel + %a_1 = OpVariable %_ptr_Function_int Function %8 + %b = OpVariable %_ptr_Function_int Function %8 + %a_2 = OpVariable %_ptr_Function_int Function %8 + %b_0 = OpVariable %_ptr_Function_int Function %8 + OpStore %a_1 %8 + %11 = OpLoad %int %a_1 + OpStore %b %11 + OpStore %a_2 %8 + %14 = OpLoad %int %a_2 + OpStore %b_0 %14 + OpReturn + OpFunctionEnd diff --git a/test/shadowing/alias/var.wgsl.expected.wgsl b/test/shadowing/alias/var.wgsl.expected.wgsl new file mode 100644 index 0000000000..10e0c25973 --- /dev/null +++ b/test/shadowing/alias/var.wgsl.expected.wgsl @@ -0,0 +1,10 @@ +type a = i32; + +fn f() { + { + var a : a = a(); + var b = a; + } + var a : a = a(); + var b = a; +} diff --git a/test/shadowing/function/let.wgsl b/test/shadowing/function/let.wgsl new file mode 100644 index 0000000000..fb5f04afd3 --- /dev/null +++ b/test/shadowing/function/let.wgsl @@ -0,0 +1,8 @@ +fn a() { + { + var a = 1; + var b = a; + } + let a = 1; + let b = a; +} diff --git a/test/shadowing/function/let.wgsl.expected.hlsl b/test/shadowing/function/let.wgsl.expected.hlsl new file mode 100644 index 0000000000..33c805bb67 --- /dev/null +++ b/test/shadowing/function/let.wgsl.expected.hlsl @@ -0,0 +1,12 @@ +[numthreads(1, 1, 1)] +void unused_entry_point() { + return; +} + +void a() { + { + int a_1 = 1; + int b = a_1; + } + const int b = 1; +} diff --git a/test/shadowing/function/let.wgsl.expected.msl b/test/shadowing/function/let.wgsl.expected.msl new file mode 100644 index 0000000000..16f3280523 --- /dev/null +++ b/test/shadowing/function/let.wgsl.expected.msl @@ -0,0 +1,12 @@ +#include + +using namespace metal; +void a() { + { + int a_1 = 1; + int b = a_1; + } + int const a_2 = 1; + int const b = a_2; +} + diff --git a/test/shadowing/function/let.wgsl.expected.spvasm b/test/shadowing/function/let.wgsl.expected.spvasm new file mode 100644 index 0000000000..592904d0e2 --- /dev/null +++ b/test/shadowing/function/let.wgsl.expected.spvasm @@ -0,0 +1,32 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 14 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %unused_entry_point "unused_entry_point" + OpExecutionMode %unused_entry_point LocalSize 1 1 1 + OpName %unused_entry_point "unused_entry_point" + OpName %a "a" + OpName %a_1 "a_1" + OpName %b "b" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %int_1 = OpConstant %int 1 +%_ptr_Function_int = OpTypePointer Function %int + %11 = OpConstantNull %int +%unused_entry_point = OpFunction %void None %1 + %4 = OpLabel + OpReturn + OpFunctionEnd + %a = OpFunction %void None %1 + %6 = OpLabel + %a_1 = OpVariable %_ptr_Function_int Function %11 + %b = OpVariable %_ptr_Function_int Function %11 + OpStore %a_1 %int_1 + %12 = OpLoad %int %a_1 + OpStore %b %12 + OpReturn + OpFunctionEnd diff --git a/test/shadowing/function/let.wgsl.expected.wgsl b/test/shadowing/function/let.wgsl.expected.wgsl new file mode 100644 index 0000000000..fb5f04afd3 --- /dev/null +++ b/test/shadowing/function/let.wgsl.expected.wgsl @@ -0,0 +1,8 @@ +fn a() { + { + var a = 1; + var b = a; + } + let a = 1; + let b = a; +} diff --git a/test/shadowing/function/param.wgsl b/test/shadowing/function/param.wgsl new file mode 100644 index 0000000000..9415ae1703 --- /dev/null +++ b/test/shadowing/function/param.wgsl @@ -0,0 +1,3 @@ +fn a(a : i32) { + let b = a; +} diff --git a/test/shadowing/function/param.wgsl.expected.hlsl b/test/shadowing/function/param.wgsl.expected.hlsl new file mode 100644 index 0000000000..f06b78f7e6 --- /dev/null +++ b/test/shadowing/function/param.wgsl.expected.hlsl @@ -0,0 +1,8 @@ +[numthreads(1, 1, 1)] +void unused_entry_point() { + return; +} + +void a(int a_1) { + const int b = a_1; +} diff --git a/test/shadowing/function/param.wgsl.expected.msl b/test/shadowing/function/param.wgsl.expected.msl new file mode 100644 index 0000000000..1c6280dc86 --- /dev/null +++ b/test/shadowing/function/param.wgsl.expected.msl @@ -0,0 +1,7 @@ +#include + +using namespace metal; +void a(int a_1) { + int const b = a_1; +} + diff --git a/test/shadowing/function/param.wgsl.expected.spvasm b/test/shadowing/function/param.wgsl.expected.spvasm new file mode 100644 index 0000000000..de4e783dc9 --- /dev/null +++ b/test/shadowing/function/param.wgsl.expected.spvasm @@ -0,0 +1,25 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 10 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %unused_entry_point "unused_entry_point" + OpExecutionMode %unused_entry_point LocalSize 1 1 1 + OpName %unused_entry_point "unused_entry_point" + OpName %a "a" + OpName %a_1 "a_1" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %5 = OpTypeFunction %void %int +%unused_entry_point = OpFunction %void None %1 + %4 = OpLabel + OpReturn + OpFunctionEnd + %a = OpFunction %void None %5 + %a_1 = OpFunctionParameter %int + %9 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/shadowing/function/param.wgsl.expected.wgsl b/test/shadowing/function/param.wgsl.expected.wgsl new file mode 100644 index 0000000000..9415ae1703 --- /dev/null +++ b/test/shadowing/function/param.wgsl.expected.wgsl @@ -0,0 +1,3 @@ +fn a(a : i32) { + let b = a; +} diff --git a/test/shadowing/function/var.wgsl b/test/shadowing/function/var.wgsl new file mode 100644 index 0000000000..7b3d8eaa97 --- /dev/null +++ b/test/shadowing/function/var.wgsl @@ -0,0 +1,10 @@ +struct a { a : i32; }; + +fn f() { + { + var a : a = a(); + var b = a; + } + var a : a = a(); + var b = a; +} diff --git a/test/shadowing/function/var.wgsl.expected.hlsl b/test/shadowing/function/var.wgsl.expected.hlsl new file mode 100644 index 0000000000..96ed5410c6 --- /dev/null +++ b/test/shadowing/function/var.wgsl.expected.hlsl @@ -0,0 +1,17 @@ +[numthreads(1, 1, 1)] +void unused_entry_point() { + return; +} + +struct a { + int a; +}; + +void f() { + { + a a_1 = (a)0; + a b = a_1; + } + a a_2 = (a)0; + a b = a_2; +} diff --git a/test/shadowing/function/var.wgsl.expected.msl b/test/shadowing/function/var.wgsl.expected.msl new file mode 100644 index 0000000000..ad6b654015 --- /dev/null +++ b/test/shadowing/function/var.wgsl.expected.msl @@ -0,0 +1,16 @@ +#include + +using namespace metal; +struct a { + int a; +}; + +void f() { + { + a a_1 = {}; + a b = a_1; + } + a a_2 = {}; + a b = a_2; +} + diff --git a/test/shadowing/function/var.wgsl.expected.spvasm b/test/shadowing/function/var.wgsl.expected.spvasm new file mode 100644 index 0000000000..2cf9d6b099 --- /dev/null +++ b/test/shadowing/function/var.wgsl.expected.spvasm @@ -0,0 +1,42 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 17 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %unused_entry_point "unused_entry_point" + OpExecutionMode %unused_entry_point LocalSize 1 1 1 + OpName %unused_entry_point "unused_entry_point" + OpName %f "f" + OpName %a "a" + OpMemberName %a 0 "a" + OpName %a_1 "a_1" + OpName %b "b" + OpName %a_2 "a_2" + OpName %b_0 "b" + OpMemberDecorate %a 0 Offset 0 + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %a = OpTypeStruct %int + %9 = OpConstantNull %a +%_ptr_Function_a = OpTypePointer Function %a +%unused_entry_point = OpFunction %void None %1 + %4 = OpLabel + OpReturn + OpFunctionEnd + %f = OpFunction %void None %1 + %6 = OpLabel + %a_1 = OpVariable %_ptr_Function_a Function %9 + %b = OpVariable %_ptr_Function_a Function %9 + %a_2 = OpVariable %_ptr_Function_a Function %9 + %b_0 = OpVariable %_ptr_Function_a Function %9 + OpStore %a_1 %9 + %12 = OpLoad %a %a_1 + OpStore %b %12 + OpStore %a_2 %9 + %15 = OpLoad %a %a_2 + OpStore %b_0 %15 + OpReturn + OpFunctionEnd diff --git a/test/shadowing/function/var.wgsl.expected.wgsl b/test/shadowing/function/var.wgsl.expected.wgsl new file mode 100644 index 0000000000..feea3e6ce5 --- /dev/null +++ b/test/shadowing/function/var.wgsl.expected.wgsl @@ -0,0 +1,12 @@ +struct a { + a : i32; +}; + +fn f() { + { + var a : a = a(); + var b = a; + } + var a : a = a(); + var b = a; +} diff --git a/test/shadowing/param/function.wgsl b/test/shadowing/param/function.wgsl new file mode 100644 index 0000000000..581a6d9552 --- /dev/null +++ b/test/shadowing/param/function.wgsl @@ -0,0 +1,6 @@ +fn a(a : i32) { + { + var a = a; + var b = a; + } +} diff --git a/test/shadowing/param/function.wgsl.expected.hlsl b/test/shadowing/param/function.wgsl.expected.hlsl new file mode 100644 index 0000000000..831cce98b1 --- /dev/null +++ b/test/shadowing/param/function.wgsl.expected.hlsl @@ -0,0 +1,11 @@ +[numthreads(1, 1, 1)] +void unused_entry_point() { + return; +} + +void a(int a_1) { + { + int a_2 = a_1; + int b = a_2; + } +} diff --git a/test/shadowing/param/function.wgsl.expected.msl b/test/shadowing/param/function.wgsl.expected.msl new file mode 100644 index 0000000000..8be9be4237 --- /dev/null +++ b/test/shadowing/param/function.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +void a(int a_1) { + { + int a_2 = a_1; + int b = a_2; + } +} + diff --git a/test/shadowing/param/function.wgsl.expected.spvasm b/test/shadowing/param/function.wgsl.expected.spvasm new file mode 100644 index 0000000000..e6291f3e86 --- /dev/null +++ b/test/shadowing/param/function.wgsl.expected.spvasm @@ -0,0 +1,34 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 15 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %unused_entry_point "unused_entry_point" + OpExecutionMode %unused_entry_point LocalSize 1 1 1 + OpName %unused_entry_point "unused_entry_point" + OpName %a "a" + OpName %a_1 "a_1" + OpName %a_2 "a_2" + OpName %b "b" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %5 = OpTypeFunction %void %int +%_ptr_Function_int = OpTypePointer Function %int + %12 = OpConstantNull %int +%unused_entry_point = OpFunction %void None %1 + %4 = OpLabel + OpReturn + OpFunctionEnd + %a = OpFunction %void None %5 + %a_1 = OpFunctionParameter %int + %9 = OpLabel + %a_2 = OpVariable %_ptr_Function_int Function %12 + %b = OpVariable %_ptr_Function_int Function %12 + OpStore %a_2 %a_1 + %13 = OpLoad %int %a_2 + OpStore %b %13 + OpReturn + OpFunctionEnd diff --git a/test/shadowing/param/function.wgsl.expected.wgsl b/test/shadowing/param/function.wgsl.expected.wgsl new file mode 100644 index 0000000000..581a6d9552 --- /dev/null +++ b/test/shadowing/param/function.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +fn a(a : i32) { + { + var a = a; + var b = a; + } +} diff --git a/test/shadowing/param/let.wgsl b/test/shadowing/param/let.wgsl new file mode 100644 index 0000000000..7a7829593d --- /dev/null +++ b/test/shadowing/param/let.wgsl @@ -0,0 +1,6 @@ +fn f(a : i32) { + { + let a = a; + let b = a; + } +} diff --git a/test/shadowing/param/let.wgsl.expected.hlsl b/test/shadowing/param/let.wgsl.expected.hlsl new file mode 100644 index 0000000000..ad522f35a0 --- /dev/null +++ b/test/shadowing/param/let.wgsl.expected.hlsl @@ -0,0 +1,10 @@ +[numthreads(1, 1, 1)] +void unused_entry_point() { + return; +} + +void f(int a) { + { + const int b = a; + } +} diff --git a/test/shadowing/param/let.wgsl.expected.msl b/test/shadowing/param/let.wgsl.expected.msl new file mode 100644 index 0000000000..9f402b1ee4 --- /dev/null +++ b/test/shadowing/param/let.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +void f(int a) { + { + int const a_1 = a; + int const b = a_1; + } +} + diff --git a/test/shadowing/param/let.wgsl.expected.spvasm b/test/shadowing/param/let.wgsl.expected.spvasm new file mode 100644 index 0000000000..1cf92967eb --- /dev/null +++ b/test/shadowing/param/let.wgsl.expected.spvasm @@ -0,0 +1,25 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 10 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %unused_entry_point "unused_entry_point" + OpExecutionMode %unused_entry_point LocalSize 1 1 1 + OpName %unused_entry_point "unused_entry_point" + OpName %f "f" + OpName %a "a" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %5 = OpTypeFunction %void %int +%unused_entry_point = OpFunction %void None %1 + %4 = OpLabel + OpReturn + OpFunctionEnd + %f = OpFunction %void None %5 + %a = OpFunctionParameter %int + %9 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/shadowing/param/let.wgsl.expected.wgsl b/test/shadowing/param/let.wgsl.expected.wgsl new file mode 100644 index 0000000000..7a7829593d --- /dev/null +++ b/test/shadowing/param/let.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +fn f(a : i32) { + { + let a = a; + let b = a; + } +} diff --git a/test/shadowing/param/var.wgsl b/test/shadowing/param/var.wgsl new file mode 100644 index 0000000000..e4392147fa --- /dev/null +++ b/test/shadowing/param/var.wgsl @@ -0,0 +1,6 @@ +fn f(a : i32) { + { + var a = a; + var b = a; + } +} diff --git a/test/shadowing/param/var.wgsl.expected.hlsl b/test/shadowing/param/var.wgsl.expected.hlsl new file mode 100644 index 0000000000..33ab9f179f --- /dev/null +++ b/test/shadowing/param/var.wgsl.expected.hlsl @@ -0,0 +1,11 @@ +[numthreads(1, 1, 1)] +void unused_entry_point() { + return; +} + +void f(int a) { + { + int a_1 = a; + int b = a_1; + } +} diff --git a/test/shadowing/param/var.wgsl.expected.msl b/test/shadowing/param/var.wgsl.expected.msl new file mode 100644 index 0000000000..a26ff19ba0 --- /dev/null +++ b/test/shadowing/param/var.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +void f(int a) { + { + int a_1 = a; + int b = a_1; + } +} + diff --git a/test/shadowing/param/var.wgsl.expected.spvasm b/test/shadowing/param/var.wgsl.expected.spvasm new file mode 100644 index 0000000000..36f0e1594d --- /dev/null +++ b/test/shadowing/param/var.wgsl.expected.spvasm @@ -0,0 +1,34 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 15 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %unused_entry_point "unused_entry_point" + OpExecutionMode %unused_entry_point LocalSize 1 1 1 + OpName %unused_entry_point "unused_entry_point" + OpName %f "f" + OpName %a "a" + OpName %a_1 "a_1" + OpName %b "b" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %5 = OpTypeFunction %void %int +%_ptr_Function_int = OpTypePointer Function %int + %12 = OpConstantNull %int +%unused_entry_point = OpFunction %void None %1 + %4 = OpLabel + OpReturn + OpFunctionEnd + %f = OpFunction %void None %5 + %a = OpFunctionParameter %int + %9 = OpLabel + %a_1 = OpVariable %_ptr_Function_int Function %12 + %b = OpVariable %_ptr_Function_int Function %12 + OpStore %a_1 %a + %13 = OpLoad %int %a_1 + OpStore %b %13 + OpReturn + OpFunctionEnd diff --git a/test/shadowing/param/var.wgsl.expected.wgsl b/test/shadowing/param/var.wgsl.expected.wgsl new file mode 100644 index 0000000000..e4392147fa --- /dev/null +++ b/test/shadowing/param/var.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +fn f(a : i32) { + { + var a = a; + var b = a; + } +} diff --git a/test/shadowing/struct/let.wgsl b/test/shadowing/struct/let.wgsl new file mode 100644 index 0000000000..5379971180 --- /dev/null +++ b/test/shadowing/struct/let.wgsl @@ -0,0 +1,10 @@ +struct a { a : i32; }; + +fn f() { + { + let a : a = a(); + let b = a; + } + let a : a = a(); + let b = a; +} diff --git a/test/shadowing/struct/let.wgsl.expected.hlsl b/test/shadowing/struct/let.wgsl.expected.hlsl new file mode 100644 index 0000000000..c4cb81d50c --- /dev/null +++ b/test/shadowing/struct/let.wgsl.expected.hlsl @@ -0,0 +1,17 @@ +[numthreads(1, 1, 1)] +void unused_entry_point() { + return; +} + +struct a { + int a; +}; + +void f() { + { + const a a_1 = (a)0; + const a b = a_1; + } + const a a_2 = (a)0; + const a b = a_2; +} diff --git a/test/shadowing/struct/let.wgsl.expected.msl b/test/shadowing/struct/let.wgsl.expected.msl new file mode 100644 index 0000000000..4780a5ed31 --- /dev/null +++ b/test/shadowing/struct/let.wgsl.expected.msl @@ -0,0 +1,16 @@ +#include + +using namespace metal; +struct a { + int a; +}; + +void f() { + { + a const a_1 = {}; + a const b = a_1; + } + a const a_2 = {}; + a const b = a_2; +} + diff --git a/test/shadowing/struct/let.wgsl.expected.spvasm b/test/shadowing/struct/let.wgsl.expected.spvasm new file mode 100644 index 0000000000..5119099b60 --- /dev/null +++ b/test/shadowing/struct/let.wgsl.expected.spvasm @@ -0,0 +1,27 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 10 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %unused_entry_point "unused_entry_point" + OpExecutionMode %unused_entry_point LocalSize 1 1 1 + OpName %unused_entry_point "unused_entry_point" + OpName %f "f" + OpName %a "a" + OpMemberName %a 0 "a" + OpMemberDecorate %a 0 Offset 0 + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %a = OpTypeStruct %int + %9 = OpConstantNull %a +%unused_entry_point = OpFunction %void None %1 + %4 = OpLabel + OpReturn + OpFunctionEnd + %f = OpFunction %void None %1 + %6 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/shadowing/struct/let.wgsl.expected.wgsl b/test/shadowing/struct/let.wgsl.expected.wgsl new file mode 100644 index 0000000000..4425a838df --- /dev/null +++ b/test/shadowing/struct/let.wgsl.expected.wgsl @@ -0,0 +1,12 @@ +struct a { + a : i32; +}; + +fn f() { + { + let a : a = a(); + let b = a; + } + let a : a = a(); + let b = a; +} diff --git a/test/shadowing/struct/param.wgsl b/test/shadowing/struct/param.wgsl new file mode 100644 index 0000000000..6c792810f6 --- /dev/null +++ b/test/shadowing/struct/param.wgsl @@ -0,0 +1,5 @@ +struct a { a : i32; }; + +fn f(a : a) { + let b = a; +} diff --git a/test/shadowing/struct/param.wgsl.expected.hlsl b/test/shadowing/struct/param.wgsl.expected.hlsl new file mode 100644 index 0000000000..03c3814c29 --- /dev/null +++ b/test/shadowing/struct/param.wgsl.expected.hlsl @@ -0,0 +1,12 @@ +[numthreads(1, 1, 1)] +void unused_entry_point() { + return; +} + +struct a { + int a; +}; + +void f(a a_1) { + const a b = a_1; +} diff --git a/test/shadowing/struct/param.wgsl.expected.msl b/test/shadowing/struct/param.wgsl.expected.msl new file mode 100644 index 0000000000..e524263311 --- /dev/null +++ b/test/shadowing/struct/param.wgsl.expected.msl @@ -0,0 +1,11 @@ +#include + +using namespace metal; +struct a { + int a; +}; + +void f(a a_1) { + a const b = a_1; +} + diff --git a/test/shadowing/struct/param.wgsl.expected.spvasm b/test/shadowing/struct/param.wgsl.expected.spvasm new file mode 100644 index 0000000000..ceab72e282 --- /dev/null +++ b/test/shadowing/struct/param.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 %unused_entry_point "unused_entry_point" + OpExecutionMode %unused_entry_point LocalSize 1 1 1 + OpName %unused_entry_point "unused_entry_point" + OpName %a "a" + OpMemberName %a 0 "a" + OpName %f "f" + OpName %a_1 "a_1" + OpMemberDecorate %a 0 Offset 0 + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %a = OpTypeStruct %int + %5 = OpTypeFunction %void %a +%unused_entry_point = OpFunction %void None %1 + %4 = OpLabel + OpReturn + OpFunctionEnd + %f = OpFunction %void None %5 + %a_1 = OpFunctionParameter %a + %10 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/shadowing/struct/param.wgsl.expected.wgsl b/test/shadowing/struct/param.wgsl.expected.wgsl new file mode 100644 index 0000000000..2cb3206793 --- /dev/null +++ b/test/shadowing/struct/param.wgsl.expected.wgsl @@ -0,0 +1,7 @@ +struct a { + a : i32; +}; + +fn f(a : a) { + let b = a; +} diff --git a/test/shadowing/struct/var.wgsl b/test/shadowing/struct/var.wgsl new file mode 100644 index 0000000000..7b3d8eaa97 --- /dev/null +++ b/test/shadowing/struct/var.wgsl @@ -0,0 +1,10 @@ +struct a { a : i32; }; + +fn f() { + { + var a : a = a(); + var b = a; + } + var a : a = a(); + var b = a; +} diff --git a/test/shadowing/struct/var.wgsl.expected.hlsl b/test/shadowing/struct/var.wgsl.expected.hlsl new file mode 100644 index 0000000000..96ed5410c6 --- /dev/null +++ b/test/shadowing/struct/var.wgsl.expected.hlsl @@ -0,0 +1,17 @@ +[numthreads(1, 1, 1)] +void unused_entry_point() { + return; +} + +struct a { + int a; +}; + +void f() { + { + a a_1 = (a)0; + a b = a_1; + } + a a_2 = (a)0; + a b = a_2; +} diff --git a/test/shadowing/struct/var.wgsl.expected.msl b/test/shadowing/struct/var.wgsl.expected.msl new file mode 100644 index 0000000000..ad6b654015 --- /dev/null +++ b/test/shadowing/struct/var.wgsl.expected.msl @@ -0,0 +1,16 @@ +#include + +using namespace metal; +struct a { + int a; +}; + +void f() { + { + a a_1 = {}; + a b = a_1; + } + a a_2 = {}; + a b = a_2; +} + diff --git a/test/shadowing/struct/var.wgsl.expected.spvasm b/test/shadowing/struct/var.wgsl.expected.spvasm new file mode 100644 index 0000000000..2cf9d6b099 --- /dev/null +++ b/test/shadowing/struct/var.wgsl.expected.spvasm @@ -0,0 +1,42 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 17 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %unused_entry_point "unused_entry_point" + OpExecutionMode %unused_entry_point LocalSize 1 1 1 + OpName %unused_entry_point "unused_entry_point" + OpName %f "f" + OpName %a "a" + OpMemberName %a 0 "a" + OpName %a_1 "a_1" + OpName %b "b" + OpName %a_2 "a_2" + OpName %b_0 "b" + OpMemberDecorate %a 0 Offset 0 + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %a = OpTypeStruct %int + %9 = OpConstantNull %a +%_ptr_Function_a = OpTypePointer Function %a +%unused_entry_point = OpFunction %void None %1 + %4 = OpLabel + OpReturn + OpFunctionEnd + %f = OpFunction %void None %1 + %6 = OpLabel + %a_1 = OpVariable %_ptr_Function_a Function %9 + %b = OpVariable %_ptr_Function_a Function %9 + %a_2 = OpVariable %_ptr_Function_a Function %9 + %b_0 = OpVariable %_ptr_Function_a Function %9 + OpStore %a_1 %9 + %12 = OpLoad %a %a_1 + OpStore %b %12 + OpStore %a_2 %9 + %15 = OpLoad %a %a_2 + OpStore %b_0 %15 + OpReturn + OpFunctionEnd diff --git a/test/shadowing/struct/var.wgsl.expected.wgsl b/test/shadowing/struct/var.wgsl.expected.wgsl new file mode 100644 index 0000000000..feea3e6ce5 --- /dev/null +++ b/test/shadowing/struct/var.wgsl.expected.wgsl @@ -0,0 +1,12 @@ +struct a { + a : i32; +}; + +fn f() { + { + var a : a = a(); + var b = a; + } + var a : a = a(); + var b = a; +}