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 <bclayton@google.com> Reviewed-by: David Neto <dneto@google.com> Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/reader/spirv/function.cc b/src/reader/spirv/function.cc index 341e687..6aad051 100644 --- a/src/reader/spirv/function.cc +++ b/src/reader/spirv/function.cc
@@ -1042,9 +1042,10 @@ } // Create and register the result type. - return_type = create<ast::Struct>( - Source{}, return_struct_sym, return_members, ast::DecorationList{}); - parser_impl_.AddConstructedType(return_struct_sym, return_type->As<ast::NamedType>()); + return_type = create<ast::Struct>(Source{}, return_struct_sym, + return_members, ast::DecorationList{}); + parser_impl_.AddConstructedType(return_struct_sym, + return_type->As<ast::NamedType>()); // Add the return-value statement. stmts.push_back(create<ast::ReturnStatement>(
diff --git a/src/writer/wgsl/generator.cc b/src/writer/wgsl/generator.cc index 6ce05f6..f190463 100644 --- a/src/writer/wgsl/generator.cc +++ b/src/writer/wgsl/generator.cc
@@ -24,7 +24,7 @@ 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 ac0146e..f488a63 100644 --- a/src/writer/wgsl/generator_impl.cc +++ b/src/writer/wgsl/generator_impl.cc
@@ -17,6 +17,7 @@ #include <algorithm> #include <limits> +#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() = 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<ast::Type>()) { @@ -83,25 +70,10 @@ return false; } } else if (auto* func = decl->As<ast::Function>()) { - 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<ast::Variable>()) { - 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::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::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 @@ 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::EmitVariable(ast::Variable* var) { - auto* sem = program_->Sem().Get(var); - make_indent(); if (!var->decorations().empty()) { @@ -579,10 +554,9 @@ 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 @@ 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<ast::StructBlockDecoration>()) { + out_ << "block"; } else if (auto* stage = deco->As<ast::StageDecoration>()) { out_ << "stage(" << stage->value() << ")"; } else if (auto* binding = deco->As<ast::BindingDecoration>()) {
diff --git a/src/writer/wgsl/generator_impl.h b/src/writer/wgsl/generator_impl.h index e7ed3de..ecb17a5 100644 --- a/src/writer/wgsl/generator_impl.h +++ b/src/writer/wgsl/generator_impl.h
@@ -50,10 +50,9 @@ 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 65bdafb..aaf382c 100644 --- a/src/writer/wgsl/generator_impl_constructor_test.cc +++ b/src/writer/wgsl/generator_impl_constructor_test.cc
@@ -29,7 +29,7 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ GeneratorImpl& gen = Build(); - ASSERT_TRUE(gen.Generate(nullptr)) << gen.error(); + ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_THAT(gen.result(), HasSubstr("vec3<f32>(1.0, 2.0, 3.0)")); } @@ -112,7 +112,7 @@ GeneratorImpl& gen = Build(); - ASSERT_TRUE(gen.Generate(nullptr)) << gen.error(); + ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_THAT(gen.result(), HasSubstr("mat2x3<f32>(vec3<f32>(1.0, 2.0, 3.0), " "vec3<f32>(3.0, 4.0, 5.0))")); } @@ -124,7 +124,7 @@ GeneratorImpl& gen = Build(); - ASSERT_TRUE(gen.Generate(nullptr)) << gen.error(); + ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_THAT(gen.result(), HasSubstr("array<vec3<f32>, 3>(vec3<f32>(1.0, 2.0, 3.0), " "vec3<f32>(4.0, 5.0, 6.0), vec3<f32>(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 6a11f84..f9e79d2 100644 --- a/src/writer/wgsl/generator_impl_function_test.cc +++ b/src/writer/wgsl/generator_impl_function_test.cc
@@ -242,7 +242,7 @@ 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 342c358..4a91097 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 @@ 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 @@ gen.increment_indent(); - ASSERT_TRUE(gen.Generate(nullptr)) << gen.error(); + ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"( var<private> a0 : f32; struct S0 { @@ -109,7 +109,7 @@ 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 @@ 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<f32>;\n");
diff --git a/src/writer/wgsl/generator_impl_test.cc b/src/writer/wgsl/generator_impl_test.cc index 57298ac..f61f0c7 100644 --- a/src/writer/wgsl/generator_impl_test.cc +++ b/src/writer/wgsl/generator_impl_test.cc
@@ -28,7 +28,7 @@ GeneratorImpl& gen = Build(); - ASSERT_TRUE(gen.Generate(nullptr)) << gen.error(); + ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"(fn my_func() { } )");