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