encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

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>
This commit is contained in:
Ben Clayton 2021-05-13 12:38:12 +00:00 committed by Commit Bot service account
parent 9b1ee6bdea
commit 688fe4477d
8 changed files with 49 additions and 73 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

7
src/reader/spirv/function.cc
View File

@ -1042,9 +1042,10 @@ bool FunctionEmitter::EmitEntryPointAsWrapper() {
}
// 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>(

2
src/writer/wgsl/generator.cc
View File

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

74
src/writer/wgsl/generator_impl.cc
View File

@ -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(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<ast::Type>()) {
@ -83,25 +70,10 @@ bool GeneratorImpl::Generate(const ast::Function* entry) {
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::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<ast::StructBlockDecoration>()) {
out_ << "block";
} else if (auto* stage = deco->As<ast::StageDecoration>()) {
out_ << "stage(" << stage->value() << ")";
} else if (auto* binding = deco->As<ast::BindingDecoration>()) {

5
src/writer/wgsl/generator_impl.h
View File

@ -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

22
src/writer/wgsl/generator_impl_constructor_test.cc
View File

@ -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<f32>(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<f32>(vec3<f32>(1.0, 2.0, 3.0), "
"vec3<f32>(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<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))"));

2
src/writer/wgsl/generator_impl_function_test.cc
View File

@ -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;

8
src/writer/wgsl/generator_impl_global_decl_test.cc
View File

@ -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<private> 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<f32>;\n");

2
src/writer/wgsl/generator_impl_test.cc
View File

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