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() {
}
)");
