// 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/writer/wgsl/generator_impl.h" #include #include "src/ast/access.h" #include "src/ast/alias.h" #include "src/ast/array.h" #include "src/ast/atomic.h" #include "src/ast/bool.h" #include "src/ast/bool_literal_expression.h" #include "src/ast/call_statement.h" #include "src/ast/depth_texture.h" #include "src/ast/external_texture.h" #include "src/ast/f32.h" #include "src/ast/float_literal_expression.h" #include "src/ast/i32.h" #include "src/ast/internal_decoration.h" #include "src/ast/interpolate_decoration.h" #include "src/ast/invariant_decoration.h" #include "src/ast/matrix.h" #include "src/ast/module.h" #include "src/ast/multisampled_texture.h" #include "src/ast/override_decoration.h" #include "src/ast/pointer.h" #include "src/ast/sampled_texture.h" #include "src/ast/sint_literal_expression.h" #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" #include "src/ast/type_name.h" #include "src/ast/u32.h" #include "src/ast/uint_literal_expression.h" #include "src/ast/variable_decl_statement.h" #include "src/ast/vector.h" #include "src/ast/void.h" #include "src/ast/workgroup_decoration.h" #include "src/sem/struct.h" #include "src/utils/math.h" #include "src/utils/scoped_assignment.h" #include "src/writer/float_to_string.h" namespace tint { namespace writer { namespace wgsl { GeneratorImpl::GeneratorImpl(const Program* program) : TextGenerator(program) {} GeneratorImpl::~GeneratorImpl() = default; bool GeneratorImpl::Generate() { // Generate global declarations in the order they appear in the module. for (auto* decl : program_->AST().GlobalDeclarations()) { if (auto* td = decl->As()) { if (!EmitTypeDecl(td)) { return false; } } else if (auto* func = decl->As()) { if (!EmitFunction(func)) { return false; } } else if (auto* var = decl->As()) { if (!EmitVariable(line(), var)) { return false; } } else { TINT_UNREACHABLE(Writer, diagnostics_); return false; } if (decl != program_->AST().GlobalDeclarations().back()) { line(); } } return true; } bool GeneratorImpl::EmitTypeDecl(const ast::TypeDecl* ty) { if (auto* alias = ty->As()) { auto out = line(); out << "type " << program_->Symbols().NameFor(alias->name) << " = "; if (!EmitType(out, alias->type)) { return false; } out << ";"; } else if (auto* str = ty->As()) { if (!EmitStructType(str)) { return false; } } else { diagnostics_.add_error( diag::System::Writer, "unknown declared type: " + std::string(ty->TypeInfo().name)); return false; } return true; } bool GeneratorImpl::EmitExpression(std::ostream& out, const ast::Expression* expr) { if (auto* a = expr->As()) { return EmitIndexAccessor(out, a); } if (auto* b = expr->As()) { return EmitBinary(out, b); } if (auto* b = expr->As()) { return EmitBitcast(out, b); } if (auto* c = expr->As()) { return EmitCall(out, c); } if (auto* i = expr->As()) { return EmitIdentifier(out, i); } if (auto* l = expr->As()) { return EmitLiteral(out, l); } if (auto* c = expr->As()) { return EmitTypeConstructor(out, c); } if (auto* m = expr->As()) { return EmitMemberAccessor(out, m); } if (expr->Is()) { out << "_"; return true; } if (auto* u = expr->As()) { return EmitUnaryOp(out, u); } diagnostics_.add_error(diag::System::Writer, "unknown expression type"); return false; } bool GeneratorImpl::EmitIndexAccessor( std::ostream& out, const ast::IndexAccessorExpression* expr) { bool paren_lhs = !expr->array ->IsAnyOf(); if (paren_lhs) { out << "("; } if (!EmitExpression(out, expr->array)) { return false; } if (paren_lhs) { out << ")"; } out << "["; if (!EmitExpression(out, expr->index)) { return false; } out << "]"; return true; } bool GeneratorImpl::EmitMemberAccessor( std::ostream& out, const ast::MemberAccessorExpression* expr) { bool paren_lhs = !expr->structure ->IsAnyOf(); if (paren_lhs) { out << "("; } if (!EmitExpression(out, expr->structure)) { return false; } if (paren_lhs) { out << ")"; } out << "."; return EmitExpression(out, expr->member); } bool GeneratorImpl::EmitBitcast(std::ostream& out, const ast::BitcastExpression* expr) { out << "bitcast<"; if (!EmitType(out, expr->type)) { return false; } out << ">("; if (!EmitExpression(out, expr->expr)) { return false; } out << ")"; return true; } bool GeneratorImpl::EmitCall(std::ostream& out, const ast::CallExpression* expr) { if (!EmitExpression(out, expr->func)) { return false; } out << "("; bool first = true; const auto& args = expr->args; for (auto* arg : args) { if (!first) { out << ", "; } first = false; if (!EmitExpression(out, arg)) { return false; } } out << ")"; return true; } bool GeneratorImpl::EmitTypeConstructor( std::ostream& out, const ast::TypeConstructorExpression* expr) { if (!EmitType(out, expr->type)) { return false; } out << "("; bool first = true; for (auto* e : expr->values) { if (!first) { out << ", "; } first = false; if (!EmitExpression(out, e)) { return false; } } out << ")"; return true; } bool GeneratorImpl::EmitLiteral(std::ostream& out, const ast::Literal* lit) { if (auto* bl = lit->As()) { out << (bl->value ? "true" : "false"); } else if (auto* fl = lit->As()) { out << FloatToBitPreservingString(fl->value); } else if (auto* sl = lit->As()) { out << sl->value; } else if (auto* ul = lit->As()) { out << ul->value << "u"; } else { diagnostics_.add_error(diag::System::Writer, "unknown literal type"); return false; } return true; } bool GeneratorImpl::EmitIdentifier(std::ostream& out, const ast::IdentifierExpression* expr) { out << program_->Symbols().NameFor(expr->symbol); return true; } bool GeneratorImpl::EmitFunction(const ast::Function* func) { if (func->decorations.size()) { if (!EmitDecorations(line(), func->decorations)) { return false; } } { auto out = line(); out << "fn " << program_->Symbols().NameFor(func->symbol) << "("; bool first = true; for (auto* v : func->params) { if (!first) { out << ", "; } first = false; if (!v->decorations.empty()) { if (!EmitDecorations(out, v->decorations)) { return false; } out << " "; } out << program_->Symbols().NameFor(v->symbol) << " : "; if (!EmitType(out, v->type)) { return false; } } out << ")"; if (!func->return_type->Is() || !func->return_type_decorations.empty()) { out << " -> "; if (!func->return_type_decorations.empty()) { if (!EmitDecorations(out, func->return_type_decorations)) { return false; } out << " "; } if (!EmitType(out, func->return_type)) { return false; } } if (func->body) { out << " {"; } } if (func->body) { if (!EmitStatementsWithIndent(func->body->statements)) { return false; } line() << "}"; } return true; } bool GeneratorImpl::EmitImageFormat(std::ostream& out, const ast::ImageFormat fmt) { switch (fmt) { case ast::ImageFormat::kNone: diagnostics_.add_error(diag::System::Writer, "unknown image format"); return false; default: out << fmt; } return true; } bool GeneratorImpl::EmitAccess(std::ostream& out, const ast::Access access) { switch (access) { case ast::Access::kRead: out << "read"; return true; case ast::Access::kWrite: out << "write"; return true; case ast::Access::kReadWrite: out << "read_write"; return true; default: break; } diagnostics_.add_error(diag::System::Writer, "unknown access"); return false; } bool GeneratorImpl::EmitType(std::ostream& out, const ast::Type* ty) { if (auto* ary = ty->As()) { for (auto* deco : ary->decorations) { if (auto* stride = deco->As()) { out << "[[stride(" << stride->stride << ")]] "; } } out << "array<"; if (!EmitType(out, ary->type)) { return false; } if (!ary->IsRuntimeArray()) { out << ", "; if (!EmitExpression(out, ary->count)) { return false; } } out << ">"; } else if (ty->Is()) { out << "bool"; } else if (ty->Is()) { out << "f32"; } else if (ty->Is()) { out << "i32"; } else if (auto* mat = ty->As()) { out << "mat" << mat->columns << "x" << mat->rows << "<"; if (!EmitType(out, mat->type)) { return false; } out << ">"; } else if (auto* ptr = ty->As()) { out << "ptr<" << ptr->storage_class << ", "; if (!EmitType(out, ptr->type)) { return false; } if (ptr->access != ast::Access::kUndefined) { out << ", "; if (!EmitAccess(out, ptr->access)) { return false; } } out << ">"; } else if (auto* atomic = ty->As()) { out << "atomic<"; if (!EmitType(out, atomic->type)) { return false; } out << ">"; } else if (auto* sampler = ty->As()) { out << "sampler"; if (sampler->IsComparison()) { out << "_comparison"; } } else if (ty->Is()) { out << "texture_external"; } else if (auto* texture = ty->As()) { out << "texture_"; if (texture->Is()) { out << "depth_"; } else if (texture->Is()) { out << "depth_multisampled_"; } else if (texture->Is()) { /* nothing to emit */ } else if (texture->Is()) { out << "multisampled_"; } else if (texture->Is()) { out << "storage_"; } else { diagnostics_.add_error(diag::System::Writer, "unknown texture type"); return false; } switch (texture->dim) { case ast::TextureDimension::k1d: out << "1d"; break; case ast::TextureDimension::k2d: out << "2d"; break; case ast::TextureDimension::k2dArray: out << "2d_array"; break; case ast::TextureDimension::k3d: out << "3d"; break; case ast::TextureDimension::kCube: out << "cube"; break; case ast::TextureDimension::kCubeArray: out << "cube_array"; break; default: diagnostics_.add_error(diag::System::Writer, "unknown texture dimension"); return false; } if (auto* sampled = texture->As()) { out << "<"; if (!EmitType(out, sampled->type)) { return false; } out << ">"; } else if (auto* ms = texture->As()) { out << "<"; if (!EmitType(out, ms->type)) { return false; } out << ">"; } else if (auto* storage = texture->As()) { out << "<"; if (!EmitImageFormat(out, storage->format)) { return false; } out << ", "; if (!EmitAccess(out, storage->access)) { return false; } out << ">"; } } else if (ty->Is()) { out << "u32"; } else if (auto* vec = ty->As()) { out << "vec" << vec->width << "<"; if (!EmitType(out, vec->type)) { return false; } out << ">"; } else if (ty->Is()) { out << "void"; } else if (auto* tn = ty->As()) { out << program_->Symbols().NameFor(tn->name); } else { diagnostics_.add_error( diag::System::Writer, "unknown type in EmitType: " + std::string(ty->TypeInfo().name)); return false; } return true; } bool GeneratorImpl::EmitStructType(const ast::Struct* str) { if (str->decorations.size()) { if (!EmitDecorations(line(), str->decorations)) { return false; } } line() << "struct " << program_->Symbols().NameFor(str->name) << " {"; auto add_padding = [&](uint32_t size) { line() << "[[size(" << size << ")]]"; // Note: u32 is the smallest primitive we currently support. When WGSL // supports smaller types, this will need to be updated. line() << UniqueIdentifier("padding") << " : u32;"; }; increment_indent(); uint32_t offset = 0; for (auto* mem : str->members) { // 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 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 // be emitted, but instead struct padding fields should be emitted. ast::DecorationList decorations_sanitized; decorations_sanitized.reserve(mem->decorations.size()); for (auto* deco : mem->decorations) { if (!deco->Is()) { decorations_sanitized.emplace_back(deco); } } if (!decorations_sanitized.empty()) { if (!EmitDecorations(line(), decorations_sanitized)) { return false; } } auto out = line(); out << program_->Symbols().NameFor(mem->symbol) << " : "; if (!EmitType(out, mem->type)) { return false; } out << ";"; } decrement_indent(); line() << "};"; return true; } bool GeneratorImpl::EmitVariable(std::ostream& out, const ast::Variable* var) { if (!var->decorations.empty()) { if (!EmitDecorations(out, var->decorations)) { return false; } out << " "; } if (var->is_const) { out << "let"; } else { out << "var"; auto sc = var->declared_storage_class; auto ac = var->declared_access; if (sc != ast::StorageClass::kNone || ac != ast::Access::kUndefined) { out << "<" << sc; if (ac != ast::Access::kUndefined) { out << ", "; if (!EmitAccess(out, ac)) { return false; } } out << ">"; } } out << " " << program_->Symbols().NameFor(var->symbol); if (auto* ty = var->type) { out << " : "; if (!EmitType(out, ty)) { return false; } } if (var->constructor != nullptr) { out << " = "; if (!EmitExpression(out, var->constructor)) { return false; } } out << ";"; return true; } bool GeneratorImpl::EmitDecorations(std::ostream& out, const ast::DecorationList& decos) { out << "[["; bool first = true; for (auto* deco : decos) { if (!first) { out << ", "; } first = false; if (auto* workgroup = deco->As()) { auto values = workgroup->Values(); out << "workgroup_size("; for (int i = 0; i < 3; i++) { if (values[i]) { if (i > 0) { out << ", "; } if (!EmitExpression(out, values[i])) { return false; } } } out << ")"; } else if (deco->Is()) { out << "block"; } else if (auto* stage = deco->As()) { out << "stage(" << stage->stage << ")"; } else if (auto* binding = deco->As()) { out << "binding(" << binding->value << ")"; } else if (auto* group = deco->As()) { out << "group(" << group->value << ")"; } else if (auto* location = deco->As()) { out << "location(" << location->value << ")"; } else if (auto* builtin = deco->As()) { out << "builtin(" << builtin->builtin << ")"; } else if (auto* interpolate = deco->As()) { out << "interpolate(" << interpolate->type; if (interpolate->sampling != ast::InterpolationSampling::kNone) { out << ", " << interpolate->sampling; } out << ")"; } else if (deco->Is()) { out << "invariant"; } else if (auto* override_deco = deco->As()) { out << "override"; if (override_deco->has_value) { out << "(" << override_deco->value << ")"; } } else if (auto* size = deco->As()) { out << "size(" << size->size << ")"; } else if (auto* align = deco->As()) { out << "align(" << align->align << ")"; } else if (auto* stride = deco->As()) { out << "stride(" << stride->stride << ")"; } else if (auto* internal = deco->As()) { out << "internal(" << internal->InternalName() << ")"; } else { TINT_ICE(Writer, diagnostics_) << "Unsupported decoration '" << deco->TypeInfo().name << "'"; return false; } } out << "]]"; return true; } bool GeneratorImpl::EmitBinary(std::ostream& out, const ast::BinaryExpression* expr) { out << "("; if (!EmitExpression(out, expr->lhs)) { return false; } out << " "; switch (expr->op) { case ast::BinaryOp::kAnd: out << "&"; break; case ast::BinaryOp::kOr: out << "|"; break; case ast::BinaryOp::kXor: out << "^"; break; case ast::BinaryOp::kLogicalAnd: out << "&&"; break; case ast::BinaryOp::kLogicalOr: out << "||"; break; case ast::BinaryOp::kEqual: out << "=="; break; case ast::BinaryOp::kNotEqual: out << "!="; break; case ast::BinaryOp::kLessThan: out << "<"; break; case ast::BinaryOp::kGreaterThan: out << ">"; break; case ast::BinaryOp::kLessThanEqual: out << "<="; break; case ast::BinaryOp::kGreaterThanEqual: out << ">="; break; case ast::BinaryOp::kShiftLeft: out << "<<"; break; case ast::BinaryOp::kShiftRight: out << ">>"; break; case ast::BinaryOp::kAdd: out << "+"; break; case ast::BinaryOp::kSubtract: out << "-"; break; case ast::BinaryOp::kMultiply: out << "*"; break; case ast::BinaryOp::kDivide: out << "/"; break; case ast::BinaryOp::kModulo: out << "%"; break; case ast::BinaryOp::kNone: diagnostics_.add_error(diag::System::Writer, "missing binary operation type"); return false; } out << " "; if (!EmitExpression(out, expr->rhs)) { return false; } out << ")"; return true; } bool GeneratorImpl::EmitUnaryOp(std::ostream& out, const ast::UnaryOpExpression* expr) { switch (expr->op) { case ast::UnaryOp::kAddressOf: out << "&"; break; case ast::UnaryOp::kComplement: out << "~"; break; case ast::UnaryOp::kIndirection: out << "*"; break; case ast::UnaryOp::kNot: out << "!"; break; case ast::UnaryOp::kNegation: out << "-"; break; } out << "("; if (!EmitExpression(out, expr->expr)) { return false; } out << ")"; return true; } bool GeneratorImpl::EmitBlock(const ast::BlockStatement* stmt) { line() << "{"; if (!EmitStatementsWithIndent(stmt->statements)) { return false; } line() << "}"; return true; } bool GeneratorImpl::EmitStatement(const ast::Statement* stmt) { if (auto* a = stmt->As()) { return EmitAssign(a); } if (auto* b = stmt->As()) { return EmitBlock(b); } if (auto* b = stmt->As()) { return EmitBreak(b); } if (auto* c = stmt->As()) { auto out = line(); if (!EmitCall(out, c->expr)) { return false; } out << ";"; return true; } if (auto* c = stmt->As()) { return EmitContinue(c); } if (auto* d = stmt->As()) { return EmitDiscard(d); } if (auto* f = stmt->As()) { return EmitFallthrough(f); } if (auto* i = stmt->As()) { return EmitIf(i); } if (auto* l = stmt->As()) { return EmitLoop(l); } if (auto* l = stmt->As()) { return EmitForLoop(l); } if (auto* r = stmt->As()) { return EmitReturn(r); } if (auto* s = stmt->As()) { return EmitSwitch(s); } if (auto* v = stmt->As()) { return EmitVariable(line(), v->variable); } diagnostics_.add_error( diag::System::Writer, "unknown statement type: " + std::string(stmt->TypeInfo().name)); return false; } bool GeneratorImpl::EmitStatements(const ast::StatementList& stmts) { for (auto* s : stmts) { if (!EmitStatement(s)) { return false; } } return true; } bool GeneratorImpl::EmitStatementsWithIndent(const ast::StatementList& stmts) { ScopedIndent si(this); return EmitStatements(stmts); } bool GeneratorImpl::EmitAssign(const ast::AssignmentStatement* stmt) { auto out = line(); if (!EmitExpression(out, stmt->lhs)) { return false; } out << " = "; if (!EmitExpression(out, stmt->rhs)) { return false; } out << ";"; return true; } bool GeneratorImpl::EmitBreak(const ast::BreakStatement*) { line() << "break;"; return true; } bool GeneratorImpl::EmitCase(const ast::CaseStatement* stmt) { if (stmt->IsDefault()) { line() << "default: {"; } else { auto out = line(); out << "case "; bool first = true; for (auto* selector : stmt->selectors) { if (!first) { out << ", "; } first = false; if (!EmitLiteral(out, selector)) { return false; } } out << ": {"; } if (!EmitStatementsWithIndent(stmt->body->statements)) { return false; } line() << "}"; return true; } bool GeneratorImpl::EmitContinue(const ast::ContinueStatement*) { line() << "continue;"; return true; } bool GeneratorImpl::EmitFallthrough(const ast::FallthroughStatement*) { line() << "fallthrough;"; return true; } bool GeneratorImpl::EmitIf(const ast::IfStatement* stmt) { { auto out = line(); out << "if ("; if (!EmitExpression(out, stmt->condition)) { return false; } out << ") {"; } if (!EmitStatementsWithIndent(stmt->body->statements)) { return false; } for (auto* e : stmt->else_statements) { if (e->condition) { auto out = line(); out << "} elseif ("; if (!EmitExpression(out, e->condition)) { return false; } out << ") {"; } else { line() << "} else {"; } if (!EmitStatementsWithIndent(e->body->statements)) { return false; } } line() << "}"; return true; } bool GeneratorImpl::EmitDiscard(const ast::DiscardStatement*) { line() << "discard;"; return true; } bool GeneratorImpl::EmitLoop(const ast::LoopStatement* stmt) { line() << "loop {"; increment_indent(); if (!EmitStatements(stmt->body->statements)) { return false; } if (stmt->continuing && !stmt->continuing->Empty()) { line(); line() << "continuing {"; if (!EmitStatementsWithIndent(stmt->continuing->statements)) { return false; } line() << "}"; } decrement_indent(); line() << "}"; return true; } bool GeneratorImpl::EmitForLoop(const ast::ForLoopStatement* stmt) { TextBuffer init_buf; if (auto* init = stmt->initializer) { TINT_SCOPED_ASSIGNMENT(current_buffer_, &init_buf); if (!EmitStatement(init)) { return false; } } TextBuffer cont_buf; if (auto* cont = stmt->continuing) { TINT_SCOPED_ASSIGNMENT(current_buffer_, &cont_buf); if (!EmitStatement(cont)) { return false; } } { auto out = line(); out << "for"; { ScopedParen sp(out); switch (init_buf.lines.size()) { case 0: // No initializer break; case 1: // Single line initializer statement out << TrimSuffix(init_buf.lines[0].content, ";"); break; default: // Block initializer statement for (size_t i = 1; i < init_buf.lines.size(); i++) { // Indent all by the first line init_buf.lines[i].indent += current_buffer_->current_indent; } out << TrimSuffix(init_buf.String(), "\n"); break; } out << "; "; if (auto* cond = stmt->condition) { if (!EmitExpression(out, cond)) { return false; } } out << "; "; switch (cont_buf.lines.size()) { case 0: // No continuing break; case 1: // Single line continuing statement out << TrimSuffix(cont_buf.lines[0].content, ";"); break; default: // Block continuing statement for (size_t i = 1; i < cont_buf.lines.size(); i++) { // Indent all by the first line cont_buf.lines[i].indent += current_buffer_->current_indent; } out << TrimSuffix(cont_buf.String(), "\n"); break; } } out << " {"; } if (!EmitStatementsWithIndent(stmt->body->statements)) { return false; } line() << "}"; return true; } bool GeneratorImpl::EmitReturn(const ast::ReturnStatement* stmt) { auto out = line(); out << "return"; if (stmt->value) { out << " "; if (!EmitExpression(out, stmt->value)) { return false; } } out << ";"; return true; } bool GeneratorImpl::EmitSwitch(const ast::SwitchStatement* stmt) { { auto out = line(); out << "switch("; if (!EmitExpression(out, stmt->condition)) { return false; } out << ") {"; } { ScopedIndent si(this); for (auto* s : stmt->body) { if (!EmitCase(s)) { return false; } } } line() << "}"; return true; } } // namespace wgsl } // namespace writer } // namespace tint