blob: 0d626f2d1b328b24e7a128da0fbdf65816241707 [file] [log] [blame]
// 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/tint/writer/wgsl/generator_impl.h"
#include <algorithm>
#include "src/tint/ast/alias.h"
#include "src/tint/ast/bool_literal_expression.h"
#include "src/tint/ast/call_statement.h"
#include "src/tint/ast/float_literal_expression.h"
#include "src/tint/ast/id_attribute.h"
#include "src/tint/ast/internal_attribute.h"
#include "src/tint/ast/interpolate_attribute.h"
#include "src/tint/ast/invariant_attribute.h"
#include "src/tint/ast/module.h"
#include "src/tint/ast/stage_attribute.h"
#include "src/tint/ast/stride_attribute.h"
#include "src/tint/ast/struct_member_align_attribute.h"
#include "src/tint/ast/struct_member_offset_attribute.h"
#include "src/tint/ast/struct_member_size_attribute.h"
#include "src/tint/ast/variable_decl_statement.h"
#include "src/tint/ast/workgroup_attribute.h"
#include "src/tint/sem/struct.h"
#include "src/tint/sem/switch_statement.h"
#include "src/tint/switch.h"
#include "src/tint/utils/defer.h"
#include "src/tint/utils/math.h"
#include "src/tint/utils/scoped_assignment.h"
#include "src/tint/writer/float_to_string.h"
namespace tint::writer::wgsl {
GeneratorImpl::GeneratorImpl(const Program* program) : TextGenerator(program) {}
GeneratorImpl::~GeneratorImpl() = default;
void GeneratorImpl::Generate() {
// Generate directives before any other global declarations.
bool has_directives = false;
for (auto enable : program_->AST().Enables()) {
EmitEnable(enable);
has_directives = true;
}
for (auto diagnostic : program_->AST().DiagnosticDirectives()) {
auto out = line();
EmitDiagnosticControl(out, diagnostic->control);
out << ";";
has_directives = true;
}
if (has_directives) {
line();
}
// Generate global declarations in the order they appear in the module.
for (auto* decl : program_->AST().GlobalDeclarations()) {
if (decl->IsAnyOf<ast::DiagnosticDirective, ast::Enable>()) {
continue;
}
Switch(
decl, //
[&](const ast::TypeDecl* td) { return EmitTypeDecl(td); },
[&](const ast::Function* func) { return EmitFunction(func); },
[&](const ast::Variable* var) { return EmitVariable(line(), var); },
[&](const ast::ConstAssert* ca) { return EmitConstAssert(ca); },
[&](Default) { TINT_UNREACHABLE(Writer, diagnostics_); });
if (decl != program_->AST().GlobalDeclarations().Back()) {
line();
}
}
}
void GeneratorImpl::EmitDiagnosticControl(utils::StringStream& out,
const ast::DiagnosticControl& diagnostic) {
out << "diagnostic(" << diagnostic.severity << ", " << diagnostic.rule_name->String() << ")";
}
void GeneratorImpl::EmitEnable(const ast::Enable* enable) {
auto out = line();
out << "enable ";
for (auto* ext : enable->extensions) {
if (ext != enable->extensions.Front()) {
out << ", ";
}
out << ext->name;
}
out << ";";
}
void GeneratorImpl::EmitTypeDecl(const ast::TypeDecl* ty) {
Switch(
ty, //
[&](const ast::Alias* alias) {
auto out = line();
out << "alias " << alias->name->symbol.Name() << " = ";
EmitExpression(out, alias->type);
out << ";";
},
[&](const ast::Struct* str) { EmitStructType(str); },
[&](Default) {
diagnostics_.add_error(diag::System::Writer,
"unknown declared type: " + std::string(ty->TypeInfo().name));
});
}
void GeneratorImpl::EmitExpression(utils::StringStream& out, const ast::Expression* expr) {
Switch(
expr, //
[&](const ast::IndexAccessorExpression* a) { EmitIndexAccessor(out, a); },
[&](const ast::BinaryExpression* b) { EmitBinary(out, b); },
[&](const ast::BitcastExpression* b) { EmitBitcast(out, b); },
[&](const ast::CallExpression* c) { EmitCall(out, c); },
[&](const ast::IdentifierExpression* i) { EmitIdentifier(out, i); },
[&](const ast::LiteralExpression* l) { EmitLiteral(out, l); },
[&](const ast::MemberAccessorExpression* m) { EmitMemberAccessor(out, m); },
[&](const ast::PhonyExpression*) { out << "_"; },
[&](const ast::UnaryOpExpression* u) { EmitUnaryOp(out, u); },
[&](Default) { diagnostics_.add_error(diag::System::Writer, "unknown expression type"); });
}
void GeneratorImpl::EmitIndexAccessor(utils::StringStream& out,
const ast::IndexAccessorExpression* expr) {
bool paren_lhs =
!expr->object
->IsAnyOf<ast::AccessorExpression, ast::CallExpression, ast::IdentifierExpression>();
if (paren_lhs) {
out << "(";
}
EmitExpression(out, expr->object);
if (paren_lhs) {
out << ")";
}
out << "[";
EmitExpression(out, expr->index);
out << "]";
}
void GeneratorImpl::EmitMemberAccessor(utils::StringStream& out,
const ast::MemberAccessorExpression* expr) {
bool paren_lhs =
!expr->object
->IsAnyOf<ast::AccessorExpression, ast::CallExpression, ast::IdentifierExpression>();
if (paren_lhs) {
out << "(";
}
EmitExpression(out, expr->object);
if (paren_lhs) {
out << ")";
}
out << "." << expr->member->symbol.Name();
}
void GeneratorImpl::EmitBitcast(utils::StringStream& out, const ast::BitcastExpression* expr) {
out << "bitcast<";
EmitExpression(out, expr->type);
out << ">(";
EmitExpression(out, expr->expr);
out << ")";
}
void GeneratorImpl::EmitCall(utils::StringStream& out, const ast::CallExpression* expr) {
EmitExpression(out, expr->target);
out << "(";
bool first = true;
const auto& args = expr->args;
for (auto* arg : args) {
if (!first) {
out << ", ";
}
first = false;
EmitExpression(out, arg);
}
out << ")";
}
void GeneratorImpl::EmitLiteral(utils::StringStream& out, const ast::LiteralExpression* lit) {
Switch(
lit, //
[&](const ast::BoolLiteralExpression* l) { out << (l->value ? "true" : "false"); },
[&](const ast::FloatLiteralExpression* l) {
// f16 literals are also emitted as float value with suffix "h".
// Note that all normal and subnormal f16 values are normal f32 values, and since NaN
// and Inf are not allowed to be spelled in literal, it should be fine to emit f16
// literals in this way.
if (l->suffix == ast::FloatLiteralExpression::Suffix::kNone) {
out << DoubleToBitPreservingString(l->value);
} else {
out << FloatToBitPreservingString(static_cast<float>(l->value)) << l->suffix;
}
},
[&](const ast::IntLiteralExpression* l) { out << l->value << l->suffix; },
[&](Default) { diagnostics_.add_error(diag::System::Writer, "unknown literal type"); });
}
void GeneratorImpl::EmitIdentifier(utils::StringStream& out,
const ast::IdentifierExpression* expr) {
EmitIdentifier(out, expr->identifier);
}
void GeneratorImpl::EmitIdentifier(utils::StringStream& out, const ast::Identifier* ident) {
if (auto* tmpl_ident = ident->As<ast::TemplatedIdentifier>()) {
if (!tmpl_ident->attributes.IsEmpty()) {
EmitAttributes(out, tmpl_ident->attributes);
out << " ";
}
out << ident->symbol.Name() << "<";
TINT_DEFER(out << ">");
for (auto* expr : tmpl_ident->arguments) {
if (expr != tmpl_ident->arguments.Front()) {
out << ", ";
}
EmitExpression(out, expr);
}
} else {
out << ident->symbol.Name();
}
}
void GeneratorImpl::EmitFunction(const ast::Function* func) {
if (func->attributes.Length()) {
EmitAttributes(line(), func->attributes);
}
{
auto out = line();
out << "fn " << func->name->symbol.Name() << "(";
bool first = true;
for (auto* v : func->params) {
if (!first) {
out << ", ";
}
first = false;
if (!v->attributes.IsEmpty()) {
EmitAttributes(out, v->attributes);
out << " ";
}
out << v->name->symbol.Name() << " : ";
EmitExpression(out, v->type);
}
out << ")";
if (func->return_type || !func->return_type_attributes.IsEmpty()) {
out << " -> ";
if (!func->return_type_attributes.IsEmpty()) {
EmitAttributes(out, func->return_type_attributes);
out << " ";
}
EmitExpression(out, func->return_type);
}
if (func->body) {
out << " ";
EmitBlockHeader(out, func->body);
}
}
if (func->body) {
EmitStatementsWithIndent(func->body->statements);
line() << "}";
}
}
void GeneratorImpl::EmitImageFormat(utils::StringStream& out, const builtin::TexelFormat fmt) {
switch (fmt) {
case builtin::TexelFormat::kUndefined:
diagnostics_.add_error(diag::System::Writer, "unknown image format");
break;
default:
out << fmt;
break;
}
}
void GeneratorImpl::EmitStructType(const ast::Struct* str) {
if (str->attributes.Length()) {
EmitAttributes(line(), str->attributes);
}
line() << "struct " << str->name->symbol.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 attribute 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 attributes 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.
utils::Vector<const ast::Attribute*, 4> attributes_sanitized;
attributes_sanitized.Reserve(mem->attributes.Length());
for (auto* attr : mem->attributes) {
if (attr->Is<ast::StructMemberOffsetAttribute>()) {
auto l = line();
l << "/* ";
EmitAttributes(l, utils::Vector{attr});
l << " */";
} else {
attributes_sanitized.Push(attr);
}
}
if (!attributes_sanitized.IsEmpty()) {
EmitAttributes(line(), attributes_sanitized);
}
auto out = line();
out << mem->name->symbol.Name() << " : ";
EmitExpression(out, mem->type);
out << ",";
}
decrement_indent();
line() << "}";
}
void GeneratorImpl::EmitVariable(utils::StringStream& out, const ast::Variable* v) {
if (!v->attributes.IsEmpty()) {
EmitAttributes(out, v->attributes);
out << " ";
}
Switch(
v, //
[&](const ast::Var* var) {
out << "var";
if (var->declared_address_space || var->declared_access) {
out << "<";
TINT_DEFER(out << ">");
EmitExpression(out, var->declared_address_space);
if (var->declared_access) {
out << ", ";
EmitExpression(out, var->declared_access);
}
}
},
[&](const ast::Let*) { out << "let"; }, [&](const ast::Override*) { out << "override"; },
[&](const ast::Const*) { out << "const"; },
[&](Default) {
TINT_ICE(Writer, diagnostics_) << "unhandled variable type " << v->TypeInfo().name;
});
out << " " << v->name->symbol.Name();
if (auto ty = v->type) {
out << " : ";
EmitExpression(out, ty);
}
if (v->initializer != nullptr) {
out << " = ";
EmitExpression(out, v->initializer);
}
out << ";";
}
void GeneratorImpl::EmitAttributes(utils::StringStream& out,
utils::VectorRef<const ast::Attribute*> attrs) {
bool first = true;
for (auto* attr : attrs) {
if (!first) {
out << " ";
}
first = false;
out << "@";
Switch(
attr, //
[&](const ast::WorkgroupAttribute* workgroup) {
auto values = workgroup->Values();
out << "workgroup_size(";
for (size_t i = 0; i < 3; i++) {
if (values[i]) {
if (i > 0) {
out << ", ";
}
EmitExpression(out, values[i]);
}
}
out << ")";
},
[&](const ast::StageAttribute* stage) { out << stage->stage; },
[&](const ast::BindingAttribute* binding) {
out << "binding(";
EmitExpression(out, binding->expr);
out << ")";
},
[&](const ast::GroupAttribute* group) {
out << "group(";
EmitExpression(out, group->expr);
out << ")";
},
[&](const ast::LocationAttribute* location) {
out << "location(";
EmitExpression(out, location->expr);
out << ")";
},
[&](const ast::BuiltinAttribute* builtin) {
out << "builtin(";
EmitExpression(out, builtin->builtin);
out << ")";
},
[&](const ast::DiagnosticAttribute* diagnostic) {
EmitDiagnosticControl(out, diagnostic->control);
},
[&](const ast::InterpolateAttribute* interpolate) {
out << "interpolate(";
EmitExpression(out, interpolate->type);
if (interpolate->sampling) {
out << ", ";
EmitExpression(out, interpolate->sampling);
}
out << ")";
},
[&](const ast::InvariantAttribute*) { out << "invariant"; },
[&](const ast::IdAttribute* override_deco) {
out << "id(";
EmitExpression(out, override_deco->expr);
out << ")";
},
[&](const ast::MustUseAttribute*) { out << "must_use"; },
[&](const ast::StructMemberOffsetAttribute* offset) {
out << "offset(";
EmitExpression(out, offset->expr);
out << ")";
},
[&](const ast::StructMemberSizeAttribute* size) {
out << "size(";
EmitExpression(out, size->expr);
out << ")";
},
[&](const ast::StructMemberAlignAttribute* align) {
out << "align(";
EmitExpression(out, align->expr);
out << ")";
},
[&](const ast::StrideAttribute* stride) { out << "stride(" << stride->stride << ")"; },
[&](const ast::InternalAttribute* internal) {
out << "internal(" << internal->InternalName() << ")";
},
[&](Default) {
TINT_ICE(Writer, diagnostics_)
<< "Unsupported attribute '" << attr->TypeInfo().name << "'";
});
}
}
void GeneratorImpl::EmitBinary(utils::StringStream& out, const ast::BinaryExpression* expr) {
out << "(";
EmitExpression(out, expr->lhs);
out << " ";
EmitBinaryOp(out, expr->op);
out << " ";
EmitExpression(out, expr->rhs);
out << ")";
}
void GeneratorImpl::EmitBinaryOp(utils::StringStream& out, const ast::BinaryOp op) {
switch (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");
break;
}
}
void GeneratorImpl::EmitUnaryOp(utils::StringStream& 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 << "(";
EmitExpression(out, expr->expr);
out << ")";
}
void GeneratorImpl::EmitBlock(const ast::BlockStatement* stmt) {
{
auto out = line();
EmitBlockHeader(out, stmt);
}
EmitStatementsWithIndent(stmt->statements);
line() << "}";
}
void GeneratorImpl::EmitBlockHeader(utils::StringStream& out, const ast::BlockStatement* stmt) {
if (!stmt->attributes.IsEmpty()) {
EmitAttributes(out, stmt->attributes);
out << " ";
}
out << "{";
}
void GeneratorImpl::EmitStatement(const ast::Statement* stmt) {
Switch(
stmt, //
[&](const ast::AssignmentStatement* a) { EmitAssign(a); },
[&](const ast::BlockStatement* b) { EmitBlock(b); },
[&](const ast::BreakStatement* b) { EmitBreak(b); },
[&](const ast::BreakIfStatement* b) { EmitBreakIf(b); },
[&](const ast::CallStatement* c) {
auto out = line();
EmitCall(out, c->expr);
out << ";";
},
[&](const ast::CompoundAssignmentStatement* c) { EmitCompoundAssign(c); },
[&](const ast::ContinueStatement* c) { EmitContinue(c); },
[&](const ast::DiscardStatement* d) { EmitDiscard(d); },
[&](const ast::IfStatement* i) { EmitIf(i); },
[&](const ast::IncrementDecrementStatement* l) { EmitIncrementDecrement(l); },
[&](const ast::LoopStatement* l) { EmitLoop(l); },
[&](const ast::ForLoopStatement* l) { EmitForLoop(l); },
[&](const ast::WhileStatement* l) { EmitWhile(l); },
[&](const ast::ReturnStatement* r) { EmitReturn(r); },
[&](const ast::ConstAssert* c) { EmitConstAssert(c); },
[&](const ast::SwitchStatement* s) { EmitSwitch(s); },
[&](const ast::VariableDeclStatement* v) { EmitVariable(line(), v->variable); },
[&](Default) {
diagnostics_.add_error(diag::System::Writer,
"unknown statement type: " + std::string(stmt->TypeInfo().name));
});
}
void GeneratorImpl::EmitStatements(utils::VectorRef<const ast::Statement*> stmts) {
for (auto* s : stmts) {
EmitStatement(s);
}
}
void GeneratorImpl::EmitStatementsWithIndent(utils::VectorRef<const ast::Statement*> stmts) {
ScopedIndent si(this);
EmitStatements(stmts);
}
void GeneratorImpl::EmitAssign(const ast::AssignmentStatement* stmt) {
auto out = line();
EmitExpression(out, stmt->lhs);
out << " = ";
EmitExpression(out, stmt->rhs);
out << ";";
}
void GeneratorImpl::EmitBreak(const ast::BreakStatement*) {
line() << "break;";
}
void GeneratorImpl::EmitBreakIf(const ast::BreakIfStatement* b) {
auto out = line();
out << "break if ";
EmitExpression(out, b->condition);
out << ";";
}
void GeneratorImpl::EmitCase(const ast::CaseStatement* stmt) {
if (stmt->selectors.Length() == 1 && stmt->ContainsDefault()) {
auto out = line();
out << "default: ";
EmitBlockHeader(out, stmt->body);
} else {
auto out = line();
out << "case ";
bool first = true;
for (auto* sel : stmt->selectors) {
if (!first) {
out << ", ";
}
first = false;
if (sel->IsDefault()) {
out << "default";
} else {
EmitExpression(out, sel->expr);
}
}
out << ": ";
EmitBlockHeader(out, stmt->body);
}
EmitStatementsWithIndent(stmt->body->statements);
line() << "}";
}
void GeneratorImpl::EmitCompoundAssign(const ast::CompoundAssignmentStatement* stmt) {
auto out = line();
EmitExpression(out, stmt->lhs);
out << " ";
EmitBinaryOp(out, stmt->op);
out << "= ";
EmitExpression(out, stmt->rhs);
out << ";";
}
void GeneratorImpl::EmitContinue(const ast::ContinueStatement*) {
line() << "continue;";
}
void GeneratorImpl::EmitIf(const ast::IfStatement* stmt) {
{
auto out = line();
if (!stmt->attributes.IsEmpty()) {
EmitAttributes(out, stmt->attributes);
out << " ";
}
out << "if (";
EmitExpression(out, stmt->condition);
out << ") ";
EmitBlockHeader(out, stmt->body);
}
EmitStatementsWithIndent(stmt->body->statements);
const ast::Statement* e = stmt->else_statement;
while (e) {
if (auto* elseif = e->As<ast::IfStatement>()) {
{
auto out = line();
out << "} else if (";
EmitExpression(out, elseif->condition);
out << ") ";
EmitBlockHeader(out, elseif->body);
}
EmitStatementsWithIndent(elseif->body->statements);
e = elseif->else_statement;
} else {
auto* body = e->As<ast::BlockStatement>();
{
auto out = line();
out << "} else ";
EmitBlockHeader(out, body);
}
EmitStatementsWithIndent(body->statements);
break;
}
}
line() << "}";
}
void GeneratorImpl::EmitIncrementDecrement(const ast::IncrementDecrementStatement* stmt) {
auto out = line();
EmitExpression(out, stmt->lhs);
out << (stmt->increment ? "++" : "--") << ";";
}
void GeneratorImpl::EmitDiscard(const ast::DiscardStatement*) {
line() << "discard;";
}
void GeneratorImpl::EmitLoop(const ast::LoopStatement* stmt) {
{
auto out = line();
if (!stmt->attributes.IsEmpty()) {
EmitAttributes(out, stmt->attributes);
out << " ";
}
out << "loop ";
EmitBlockHeader(out, stmt->body);
}
increment_indent();
EmitStatements(stmt->body->statements);
if (stmt->continuing && !stmt->continuing->Empty()) {
line();
{
auto out = line();
out << "continuing ";
if (!stmt->continuing->attributes.IsEmpty()) {
EmitAttributes(out, stmt->continuing->attributes);
out << " ";
}
out << "{";
}
EmitStatementsWithIndent(stmt->continuing->statements);
line() << "}";
}
decrement_indent();
line() << "}";
}
void GeneratorImpl::EmitForLoop(const ast::ForLoopStatement* stmt) {
TextBuffer init_buf;
if (auto* init = stmt->initializer) {
TINT_SCOPED_ASSIGNMENT(current_buffer_, &init_buf);
EmitStatement(init);
}
TextBuffer cont_buf;
if (auto* cont = stmt->continuing) {
TINT_SCOPED_ASSIGNMENT(current_buffer_, &cont_buf);
EmitStatement(cont);
}
{
auto out = line();
if (!stmt->attributes.IsEmpty()) {
EmitAttributes(out, stmt->attributes);
out << " ";
}
out << "for";
{
ScopedParen sp(out);
switch (init_buf.lines.size()) {
case 0: // No initializer
break;
case 1: // Single line initializer statement
out << utils::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 << utils::TrimSuffix(init_buf.String(), "\n");
break;
}
out << "; ";
if (auto* cond = stmt->condition) {
EmitExpression(out, cond);
}
out << "; ";
switch (cont_buf.lines.size()) {
case 0: // No continuing
break;
case 1: // Single line continuing statement
out << utils::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 << utils::TrimSuffix(cont_buf.String(), "\n");
break;
}
}
out << " ";
EmitBlockHeader(out, stmt->body);
}
EmitStatementsWithIndent(stmt->body->statements);
line() << "}";
}
void GeneratorImpl::EmitWhile(const ast::WhileStatement* stmt) {
{
auto out = line();
if (!stmt->attributes.IsEmpty()) {
EmitAttributes(out, stmt->attributes);
out << " ";
}
out << "while";
{
ScopedParen sp(out);
auto* cond = stmt->condition;
EmitExpression(out, cond);
}
out << " ";
EmitBlockHeader(out, stmt->body);
}
EmitStatementsWithIndent(stmt->body->statements);
line() << "}";
}
void GeneratorImpl::EmitReturn(const ast::ReturnStatement* stmt) {
auto out = line();
out << "return";
if (stmt->value) {
out << " ";
EmitExpression(out, stmt->value);
}
out << ";";
}
void GeneratorImpl::EmitConstAssert(const ast::ConstAssert* stmt) {
auto out = line();
out << "const_assert ";
EmitExpression(out, stmt->condition);
out << ";";
}
void GeneratorImpl::EmitSwitch(const ast::SwitchStatement* stmt) {
{
auto out = line();
if (!stmt->attributes.IsEmpty()) {
EmitAttributes(out, stmt->attributes);
out << " ";
}
out << "switch(";
EmitExpression(out, stmt->condition);
out << ") ";
if (!stmt->body_attributes.IsEmpty()) {
EmitAttributes(out, stmt->body_attributes);
out << " ";
}
out << "{";
}
{
ScopedIndent si(this);
for (auto* s : stmt->body) {
EmitCase(s);
}
}
line() << "}";
}
} // namespace tint::writer::wgsl