From 688fe4477dffbe0ff6d4d5edab8fe202ec0ca66d Mon Sep 17 00:00:00 2001 From: Ben Clayton Date: Thu, 13 May 2021 12:38:12 +0000 Subject: [PATCH] writer/wgsl: Cut dependencies on sem info There's now one remainin use of the semantic info in the WGSL writer - the transformation of [[offset]] into padded fields. This does not belong in the WGSL writer, but instead part of the transform::Wgsl sanitizer. Bug: tint:798 Change-Id: I95ba11f022c41150cc12de84a4085cd7d42019fe Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/50822 Commit-Queue: Ben Clayton Reviewed-by: David Neto Reviewed-by: James Price --- src/reader/spirv/function.cc | 7 +- src/writer/wgsl/generator.cc | 2 +- src/writer/wgsl/generator_impl.cc | 74 +++++++------------ src/writer/wgsl/generator_impl.h | 5 +- .../wgsl/generator_impl_constructor_test.cc | 22 +++--- .../wgsl/generator_impl_function_test.cc | 2 +- .../wgsl/generator_impl_global_decl_test.cc | 8 +- src/writer/wgsl/generator_impl_test.cc | 2 +- 8 files changed, 49 insertions(+), 73 deletions(-) diff --git a/src/reader/spirv/function.cc b/src/reader/spirv/function.cc index 341e68775f..6aad05194b 100644 --- a/src/reader/spirv/function.cc +++ b/src/reader/spirv/function.cc @@ -1042,9 +1042,10 @@ bool FunctionEmitter::EmitEntryPointAsWrapper() { } // Create and register the result type. - return_type = create( - Source{}, return_struct_sym, return_members, ast::DecorationList{}); - parser_impl_.AddConstructedType(return_struct_sym, return_type->As()); + return_type = create(Source{}, return_struct_sym, + return_members, ast::DecorationList{}); + parser_impl_.AddConstructedType(return_struct_sym, + return_type->As()); // Add the return-value statement. stmts.push_back(create( diff --git a/src/writer/wgsl/generator.cc b/src/writer/wgsl/generator.cc index 6ce05f67b5..f19046375c 100644 --- a/src/writer/wgsl/generator.cc +++ b/src/writer/wgsl/generator.cc @@ -24,7 +24,7 @@ Generator::Generator(const Program* program) Generator::~Generator() = default; bool Generator::Generate() { - auto ret = impl_->Generate(nullptr); + auto ret = impl_->Generate(); if (!ret) { error_ = impl_->error(); } diff --git a/src/writer/wgsl/generator_impl.cc b/src/writer/wgsl/generator_impl.cc index ac0146ed5b..f488a63c56 100644 --- a/src/writer/wgsl/generator_impl.cc +++ b/src/writer/wgsl/generator_impl.cc @@ -17,6 +17,7 @@ #include #include +#include "src/ast/access_control.h" #include "src/ast/alias.h" #include "src/ast/array.h" #include "src/ast/bool.h" @@ -37,6 +38,7 @@ #include "src/ast/stage_decoration.h" #include "src/ast/storage_texture.h" #include "src/ast/stride_decoration.h" +#include "src/ast/struct_block_decoration.h" #include "src/ast/struct_member_align_decoration.h" #include "src/ast/struct_member_offset_decoration.h" #include "src/ast/struct_member_size_decoration.h" @@ -47,22 +49,7 @@ #include "src/ast/vector.h" #include "src/ast/void.h" #include "src/ast/workgroup_decoration.h" -#include "src/sem/access_control_type.h" -#include "src/sem/array.h" -#include "src/sem/bool_type.h" -#include "src/sem/depth_texture_type.h" -#include "src/sem/f32_type.h" -#include "src/sem/function.h" -#include "src/sem/i32_type.h" -#include "src/sem/matrix_type.h" -#include "src/sem/multisampled_texture_type.h" -#include "src/sem/pointer_type.h" -#include "src/sem/sampled_texture_type.h" #include "src/sem/struct.h" -#include "src/sem/u32_type.h" -#include "src/sem/variable.h" -#include "src/sem/vector_type.h" -#include "src/sem/void_type.h" #include "src/utils/math.h" #include "src/writer/float_to_string.h" @@ -75,7 +62,7 @@ GeneratorImpl::GeneratorImpl(const Program* program) GeneratorImpl::~GeneratorImpl() = default; -bool GeneratorImpl::Generate(const ast::Function* entry) { +bool GeneratorImpl::Generate() { // Generate global declarations in the order they appear in the module. for (auto* decl : program_->AST().GlobalDeclarations()) { if (auto* ty = decl->As()) { @@ -83,25 +70,10 @@ bool GeneratorImpl::Generate(const ast::Function* entry) { return false; } } else if (auto* func = decl->As()) { - if (entry && func != entry) { - // Skip functions that are not reachable by the target entry point. - auto* sem = program_->Sem().Get(func); - if (!sem->HasAncestorEntryPoint(entry->symbol())) { - continue; - } - } if (!EmitFunction(func)) { return false; } } else if (auto* var = decl->As()) { - if (entry && !var->is_const()) { - // Skip variables that are not referenced by the target entry point. - auto& refs = program_->Sem().Get(entry)->ReferencedModuleVariables(); - if (std::find(refs.begin(), refs.end(), program_->Sem().Get(var)) == - refs.end()) { - continue; - } - } if (!EmitVariable(var)) { return false; } @@ -291,7 +263,9 @@ bool GeneratorImpl::EmitIdentifier(ast::IdentifierExpression* expr) { bool GeneratorImpl::EmitFunction(ast::Function* func) { if (func->decorations().size()) { make_indent(); - EmitDecorations(func->decorations()); + if (!EmitDecorations(func->decorations())) { + return false; + } out_ << std::endl; } @@ -501,10 +475,11 @@ bool GeneratorImpl::EmitType(const ast::Type* ty) { } bool GeneratorImpl::EmitStructType(const ast::Struct* str) { - for (auto* deco : str->decorations()) { - out_ << "[["; - program_->to_str(deco, out_, 0); - out_ << "]]" << std::endl; + if (str->decorations().size()) { + if (!EmitDecorations(str->decorations())) { + return false; + } + out_ << std::endl; } out_ << "struct " << program_->Symbols().NameFor(str->name()) << " {" << std::endl; @@ -521,14 +496,16 @@ bool GeneratorImpl::EmitStructType(const ast::Struct* str) { increment_indent(); uint32_t offset = 0; for (auto* mem : str->members()) { - auto* mem_sem = program_->Sem().Get(mem); - - offset = utils::RoundUp(mem_sem->Align(), offset); - if (uint32_t padding = mem_sem->Offset() - offset) { - add_padding(padding); - offset += padding; + // TODO(crbug.com/tint/798) move the [[offset]] decoration handling to the + // transform::Wgsl sanitizer. + if (auto* mem_sem = program_->Sem().Get(mem)) { + offset = utils::RoundUp(mem_sem->Align(), offset); + if (uint32_t padding = mem_sem->Offset() - offset) { + add_padding(padding); + offset += padding; + } + offset += mem_sem->Size(); } - offset += mem_sem->Size(); // Offset decorations no longer exist in the WGSL spec, but are emitted // by the SPIR-V reader and are consumed by the Resolver(). These should not @@ -564,8 +541,6 @@ bool GeneratorImpl::EmitStructType(const ast::Struct* str) { } bool GeneratorImpl::EmitVariable(ast::Variable* var) { - auto* sem = program_->Sem().Get(var); - make_indent(); if (!var->decorations().empty()) { @@ -579,10 +554,9 @@ bool GeneratorImpl::EmitVariable(ast::Variable* var) { out_ << "let"; } else { out_ << "var"; - if (sem->StorageClass() != ast::StorageClass::kNone && - sem->StorageClass() != ast::StorageClass::kFunction && - !sem->Type()->UnwrapAll()->is_handle()) { - out_ << "<" << sem->StorageClass() << ">"; + auto sc = var->declared_storage_class(); + if (sc != ast::StorageClass::kNone) { + out_ << "<" << sc << ">"; } } @@ -622,6 +596,8 @@ bool GeneratorImpl::EmitDecorations(const ast::DecorationList& decos) { std::tie(x, y, z) = workgroup->values(); out_ << "workgroup_size(" << std::to_string(x) << ", " << std::to_string(y) << ", " << std::to_string(z) << ")"; + } else if (deco->Is()) { + out_ << "block"; } else if (auto* stage = deco->As()) { out_ << "stage(" << stage->value() << ")"; } else if (auto* binding = deco->As()) { diff --git a/src/writer/wgsl/generator_impl.h b/src/writer/wgsl/generator_impl.h index e7ed3debea..ecb17a5be3 100644 --- a/src/writer/wgsl/generator_impl.h +++ b/src/writer/wgsl/generator_impl.h @@ -50,10 +50,9 @@ class GeneratorImpl : public TextGenerator { explicit GeneratorImpl(const Program* program); ~GeneratorImpl(); - /// Generates the result data, optionally restricted to a single entry point - /// @param entry entry point to target, or nullptr + /// Generates the result data /// @returns true on successful generation; false otherwise - bool Generate(const ast::Function* entry); + bool Generate(); /// Handles generating a constructed type /// @param ty the constructed to generate diff --git a/src/writer/wgsl/generator_impl_constructor_test.cc b/src/writer/wgsl/generator_impl_constructor_test.cc index 65bdafb825..aaf382c1cc 100644 --- a/src/writer/wgsl/generator_impl_constructor_test.cc +++ b/src/writer/wgsl/generator_impl_constructor_test.cc @@ -29,7 +29,7 @@ TEST_F(WgslGeneratorImplTest, EmitConstructor_Bool) { GeneratorImpl& gen = Build(); - ASSERT_TRUE(gen.Generate(nullptr)) << gen.error(); + ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_THAT(gen.result(), HasSubstr("false")); } @@ -38,7 +38,7 @@ TEST_F(WgslGeneratorImplTest, EmitConstructor_Int) { GeneratorImpl& gen = Build(); - ASSERT_TRUE(gen.Generate(nullptr)) << gen.error(); + ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_THAT(gen.result(), HasSubstr("-12345")); } @@ -47,7 +47,7 @@ TEST_F(WgslGeneratorImplTest, EmitConstructor_UInt) { GeneratorImpl& gen = Build(); - ASSERT_TRUE(gen.Generate(nullptr)) << gen.error(); + ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_THAT(gen.result(), HasSubstr("56779u")); } @@ -57,7 +57,7 @@ TEST_F(WgslGeneratorImplTest, EmitConstructor_Float) { GeneratorImpl& gen = Build(); - ASSERT_TRUE(gen.Generate(nullptr)) << gen.error(); + ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_THAT(gen.result(), HasSubstr("1073741824.0")); } @@ -66,7 +66,7 @@ TEST_F(WgslGeneratorImplTest, EmitConstructor_Type_Float) { GeneratorImpl& gen = Build(); - ASSERT_TRUE(gen.Generate(nullptr)) << gen.error(); + ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_THAT(gen.result(), HasSubstr("f32(-0.000012)")); } @@ -75,7 +75,7 @@ TEST_F(WgslGeneratorImplTest, EmitConstructor_Type_Bool) { GeneratorImpl& gen = Build(); - ASSERT_TRUE(gen.Generate(nullptr)) << gen.error(); + ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_THAT(gen.result(), HasSubstr("bool(true)")); } @@ -84,7 +84,7 @@ TEST_F(WgslGeneratorImplTest, EmitConstructor_Type_Int) { GeneratorImpl& gen = Build(); - ASSERT_TRUE(gen.Generate(nullptr)) << gen.error(); + ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_THAT(gen.result(), HasSubstr("i32(-12345)")); } @@ -93,7 +93,7 @@ TEST_F(WgslGeneratorImplTest, EmitConstructor_Type_Uint) { GeneratorImpl& gen = Build(); - ASSERT_TRUE(gen.Generate(nullptr)) << gen.error(); + ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_THAT(gen.result(), HasSubstr("u32(12345u)")); } @@ -102,7 +102,7 @@ TEST_F(WgslGeneratorImplTest, EmitConstructor_Type_Vec) { GeneratorImpl& gen = Build(); - ASSERT_TRUE(gen.Generate(nullptr)) << gen.error(); + ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_THAT(gen.result(), HasSubstr("vec3(1.0, 2.0, 3.0)")); } @@ -112,7 +112,7 @@ TEST_F(WgslGeneratorImplTest, EmitConstructor_Type_Mat) { GeneratorImpl& gen = Build(); - ASSERT_TRUE(gen.Generate(nullptr)) << gen.error(); + ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_THAT(gen.result(), HasSubstr("mat2x3(vec3(1.0, 2.0, 3.0), " "vec3(3.0, 4.0, 5.0))")); } @@ -124,7 +124,7 @@ TEST_F(WgslGeneratorImplTest, EmitConstructor_Type_Array) { GeneratorImpl& gen = Build(); - ASSERT_TRUE(gen.Generate(nullptr)) << gen.error(); + ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_THAT(gen.result(), HasSubstr("array, 3>(vec3(1.0, 2.0, 3.0), " "vec3(4.0, 5.0, 6.0), vec3(7.0, 8.0, 9.0))")); diff --git a/src/writer/wgsl/generator_impl_function_test.cc b/src/writer/wgsl/generator_impl_function_test.cc index 6a11f84570..f9e79d2865 100644 --- a/src/writer/wgsl/generator_impl_function_test.cc +++ b/src/writer/wgsl/generator_impl_function_test.cc @@ -242,7 +242,7 @@ TEST_F(WgslGeneratorImplTest, GeneratorImpl& gen = Build(); - ASSERT_TRUE(gen.Generate(nullptr)) << gen.error(); + ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"([[block]] struct Data { d : f32; diff --git a/src/writer/wgsl/generator_impl_global_decl_test.cc b/src/writer/wgsl/generator_impl_global_decl_test.cc index 342c358698..4a9109769c 100644 --- a/src/writer/wgsl/generator_impl_global_decl_test.cc +++ b/src/writer/wgsl/generator_impl_global_decl_test.cc @@ -35,7 +35,7 @@ TEST_F(WgslGeneratorImplTest, Emit_GlobalDeclAfterFunction) { gen.increment_indent(); - ASSERT_TRUE(gen.Generate(nullptr)) << gen.error(); + ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"( [[stage(compute)]] fn test_function() { var a : f32; @@ -74,7 +74,7 @@ TEST_F(WgslGeneratorImplTest, Emit_GlobalsInterleaved) { gen.increment_indent(); - ASSERT_TRUE(gen.Generate(nullptr)) << gen.error(); + ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"( var a0 : f32; struct S0 { @@ -109,7 +109,7 @@ TEST_F(WgslGeneratorImplTest, Emit_Global_Sampler) { gen.increment_indent(); - ASSERT_TRUE(gen.Generate(nullptr)) << gen.error(); + ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), " [[group(0), binding(0)]] var s : sampler;\n"); } @@ -123,7 +123,7 @@ TEST_F(WgslGeneratorImplTest, Emit_Global_Texture) { gen.increment_indent(); - ASSERT_TRUE(gen.Generate(nullptr)) << gen.error(); + ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ( gen.result(), " [[group(0), binding(0)]] var t : [[access(read)]] texture_1d;\n"); diff --git a/src/writer/wgsl/generator_impl_test.cc b/src/writer/wgsl/generator_impl_test.cc index 57298acc5a..f61f0c7f31 100644 --- a/src/writer/wgsl/generator_impl_test.cc +++ b/src/writer/wgsl/generator_impl_test.cc @@ -28,7 +28,7 @@ TEST_F(WgslGeneratorImplTest, Generate) { GeneratorImpl& gen = Build(); - ASSERT_TRUE(gen.Generate(nullptr)) << gen.error(); + ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"(fn my_func() { } )");