writer/hlsl: Rework string printing
Remove `pre` and `out` parameters from most generator methods.
Use the `out_` string stream in TextGenerator, add helpers to TextGenerator to simplify line printing.
Remove the `pre` and `out` fields from TestHelper.
Cleans up the `pre` aspects of the HLSL writer, so the same concept can be used by the MSL writer.
Fixes indentation bugs in formatting.
Change-Id: Ia35daf632c7c7b84a6fbf3b9ae42baaeb3c97649
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/55960
Auto-Submit: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: David Neto <dneto@google.com>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/writer/hlsl/generator.cc b/src/writer/hlsl/generator.cc
index d55e70c..297788e 100644
--- a/src/writer/hlsl/generator.cc
+++ b/src/writer/hlsl/generator.cc
@@ -26,7 +26,7 @@
Generator::~Generator() = default;
bool Generator::Generate() {
- auto ret = impl_->Generate(out_);
+ auto ret = impl_->Generate();
if (!ret) {
error_ = impl_->error();
}
@@ -34,7 +34,7 @@
}
std::string Generator::result() const {
- return out_.str();
+ return impl_->result();
}
std::string Generator::error() const {
diff --git a/src/writer/hlsl/generator.h b/src/writer/hlsl/generator.h
index 3623dd5..a3bef00 100644
--- a/src/writer/hlsl/generator.h
+++ b/src/writer/hlsl/generator.h
@@ -48,7 +48,6 @@
std::string error() const;
private:
- std::ostringstream out_;
std::unique_ptr<GeneratorImpl> impl_;
};
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index f7677e4..496be2d 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -16,6 +16,7 @@
#include <algorithm>
#include <iomanip>
+#include <iosfwd>
#include <utility>
#include <vector>
@@ -100,13 +101,6 @@
return s;
}
-// Helper for writting a '(' on construction and a ')' destruction.
-struct ScopedParen {
- std::ostream& s_;
- explicit ScopedParen(std::ostream& s) : s_(s) { s << "("; }
- ~ScopedParen() { s_ << ")"; }
-};
-
} // namespace
GeneratorImpl::GeneratorImpl(const Program* program)
@@ -114,7 +108,7 @@
GeneratorImpl::~GeneratorImpl() = default;
-bool GeneratorImpl::Generate(std::ostream& out) {
+bool GeneratorImpl::Generate() {
if (!builder_.HasTransformApplied<transform::Hlsl>()) {
diagnostics_.add_error(
diag::System::Writer,
@@ -123,8 +117,8 @@
return false;
}
- std::stringstream pending;
const TypeInfo* last_kind = nullptr;
+ std::streampos last_padding_pos;
for (auto* decl : builder_.AST().GlobalDeclarations()) {
if (decl->Is<ast::Alias>()) {
@@ -134,29 +128,29 @@
// Emit a new line between declarations if the type of declaration has
// changed, or we're about to emit a function
auto* kind = &decl->TypeInfo();
- if (pending.str().length() &&
- (last_kind != kind || decl->Is<ast::Function>())) {
- out << pending.str() << std::endl;
- pending.str(std::string());
- make_indent(out);
+ if (out_.tellp() != last_padding_pos) {
+ if (last_kind && (last_kind != kind || decl->Is<ast::Function>())) {
+ out_ << std::endl;
+ last_padding_pos = out_.tellp();
+ }
}
last_kind = kind;
if (auto* global = decl->As<ast::Variable>()) {
- if (!EmitGlobalVariable(pending, global)) {
+ if (!EmitGlobalVariable(global)) {
return false;
}
} else if (auto* str = decl->As<ast::Struct>()) {
- if (!EmitStructType(pending, builder_.Sem().Get(str))) {
+ if (!EmitStructType(builder_.Sem().Get(str))) {
return false;
}
} else if (auto* func = decl->As<ast::Function>()) {
if (func->IsEntryPoint()) {
- if (!EmitEntryPointFunction(pending, func)) {
+ if (!EmitEntryPointFunction(func)) {
return false;
}
} else {
- if (!EmitFunction(pending, func)) {
+ if (!EmitFunction(func)) {
return false;
}
}
@@ -167,8 +161,6 @@
}
}
- out << pending.str();
-
return true;
}
@@ -176,15 +168,14 @@
return builder_.Symbols().NameFor(builder_.Symbols().New(prefix));
}
-bool GeneratorImpl::EmitArrayAccessor(std::ostream& pre,
- std::ostream& out,
+bool GeneratorImpl::EmitArrayAccessor(std::ostream& out,
ast::ArrayAccessorExpression* expr) {
- if (!EmitExpression(pre, out, expr->array())) {
+ if (!EmitExpression(out, expr->array())) {
return false;
}
out << "[";
- if (!EmitExpression(pre, out, expr->idx_expr())) {
+ if (!EmitExpression(out, expr->idx_expr())) {
return false;
}
out << "]";
@@ -192,8 +183,7 @@
return true;
}
-bool GeneratorImpl::EmitBitcast(std::ostream& pre,
- std::ostream& out,
+bool GeneratorImpl::EmitBitcast(std::ostream& out,
ast::BitcastExpression* expr) {
auto* type = TypeOf(expr);
if (!type->is_integer_scalar() && !type->is_float_scalar()) {
@@ -208,67 +198,57 @@
return false;
}
out << "(";
- if (!EmitExpression(pre, out, expr->expr())) {
+ if (!EmitExpression(out, expr->expr())) {
return false;
}
out << ")";
return true;
}
-bool GeneratorImpl::EmitAssign(std::ostream& out,
- ast::AssignmentStatement* stmt) {
- make_indent(out);
-
- std::ostringstream pre;
-
- std::ostringstream lhs_out;
- if (!EmitExpression(pre, lhs_out, stmt->lhs())) {
+bool GeneratorImpl::EmitAssign(ast::AssignmentStatement* stmt) {
+ auto out = line();
+ if (!EmitExpression(out, stmt->lhs())) {
return false;
}
- std::ostringstream rhs_out;
- if (!EmitExpression(pre, rhs_out, stmt->rhs())) {
+ out << " = ";
+ if (!EmitExpression(out, stmt->rhs())) {
return false;
}
-
- out << pre.str();
- out << lhs_out.str() << " = " << rhs_out.str() << ";" << std::endl;
-
+ out << ";";
return true;
}
-bool GeneratorImpl::EmitBinary(std::ostream& pre,
- std::ostream& out,
- ast::BinaryExpression* expr) {
+bool GeneratorImpl::EmitBinary(std::ostream& out, ast::BinaryExpression* expr) {
if (expr->op() == ast::BinaryOp::kLogicalAnd ||
expr->op() == ast::BinaryOp::kLogicalOr) {
- std::ostringstream lhs_out;
- if (!EmitExpression(pre, lhs_out, expr->lhs())) {
- return false;
- }
-
auto name = generate_name(kTempNamePrefix);
- make_indent(pre);
- pre << "bool " << name << " = " << lhs_out.str() << ";" << std::endl;
- make_indent(pre);
- pre << "if (";
+ {
+ auto pre = line();
+ pre << "bool " << name << " = ";
+ if (!EmitExpression(pre, expr->lhs())) {
+ return false;
+ }
+ pre << ";";
+ }
+
if (expr->op() == ast::BinaryOp::kLogicalOr) {
- pre << "!";
- }
- pre << name << ") {" << std::endl;
- increment_indent();
-
- std::ostringstream rhs_out;
- if (!EmitExpression(pre, rhs_out, expr->rhs())) {
- return false;
+ line() << "if (!" << name << ") {";
+ } else {
+ line() << "if (" << name << ") {";
}
- make_indent(pre);
- pre << name << " = " << rhs_out.str() << ";" << std::endl;
+ {
+ ScopedIndent si(this);
+ auto pre = line();
+ pre << name << " = ";
+ if (!EmitExpression(pre, expr->rhs())) {
+ return false;
+ }
+ pre << ";";
+ }
- decrement_indent();
- make_indent(pre);
- pre << "}" << std::endl;
+ line() << "}";
out << "(" << name << ")";
return true;
@@ -284,11 +264,11 @@
(lhs_type->Is<sem::Matrix>() && rhs_type->Is<sem::Matrix>()))) {
// Matrices are transposed, so swap LHS and RHS.
out << "mul(";
- if (!EmitExpression(pre, out, expr->rhs())) {
+ if (!EmitExpression(out, expr->rhs())) {
return false;
}
out << ", ";
- if (!EmitExpression(pre, out, expr->lhs())) {
+ if (!EmitExpression(out, expr->lhs())) {
return false;
}
out << ")";
@@ -297,7 +277,7 @@
}
out << "(";
- if (!EmitExpression(pre, out, expr->lhs())) {
+ if (!EmitExpression(out, expr->lhs())) {
return false;
}
out << " ";
@@ -369,7 +349,7 @@
}
out << " ";
- if (!EmitExpression(pre, out, expr->rhs())) {
+ if (!EmitExpression(out, expr->rhs())) {
return false;
}
@@ -377,46 +357,35 @@
return true;
}
-bool GeneratorImpl::EmitBlock(std::ostream& out,
- const ast::BlockStatement* stmt) {
- return EmitBlockBraces(out, [&] {
- for (auto* s : *stmt) {
- if (!EmitStatement(out, s)) {
- return false;
- }
+bool GeneratorImpl::EmitStatements(const ast::StatementList& stmts) {
+ for (auto* s : stmts) {
+ if (!EmitStatement(s)) {
+ return false;
}
- return true;
- });
-}
-
-bool GeneratorImpl::EmitBlockAndNewline(std::ostream& out,
- const ast::BlockStatement* stmt) {
- const bool result = EmitBlock(out, stmt);
- if (result) {
- out << std::endl;
}
- return result;
-}
-
-bool GeneratorImpl::EmitIndentedBlockAndNewline(std::ostream& out,
- ast::BlockStatement* stmt) {
- make_indent(out);
- const bool result = EmitBlock(out, stmt);
- if (result) {
- out << std::endl;
- }
- return result;
-}
-
-bool GeneratorImpl::EmitBreak(std::ostream& out, ast::BreakStatement*) {
- make_indent(out);
- out << "break;" << std::endl;
return true;
}
-bool GeneratorImpl::EmitCall(std::ostream& pre,
- std::ostream& out,
- ast::CallExpression* expr) {
+bool GeneratorImpl::EmitStatementsWithIndent(const ast::StatementList& stmts) {
+ ScopedIndent si(this);
+ return EmitStatements(stmts);
+}
+
+bool GeneratorImpl::EmitBlock(const ast::BlockStatement* stmt) {
+ line() << "{";
+ if (!EmitStatementsWithIndent(stmt->statements())) {
+ return false;
+ }
+ line() << "}";
+ return true;
+}
+
+bool GeneratorImpl::EmitBreak(ast::BreakStatement*) {
+ line() << "break;";
+ return true;
+}
+
+bool GeneratorImpl::EmitCall(std::ostream& out, ast::CallExpression* expr) {
const auto& params = expr->params();
auto* ident = expr->func();
auto* call = builder_.Sem().Get(expr);
@@ -428,11 +397,11 @@
func->Declaration()->decorations())) {
// Special function generated by the CalculateArrayLength transform for
// calling X.GetDimensions(Y)
- if (!EmitExpression(pre, out, params[0])) {
+ if (!EmitExpression(out, params[0])) {
return false;
}
out << ".GetDimensions(";
- if (!EmitExpression(pre, out, params[1])) {
+ if (!EmitExpression(out, params[1])) {
return false;
}
out << ")";
@@ -444,9 +413,9 @@
func->Declaration()->decorations())) {
switch (intrinsic->storage_class) {
case ast::StorageClass::kUniform:
- return EmitUniformBufferAccess(pre, out, expr, intrinsic);
+ return EmitUniformBufferAccess(out, expr, intrinsic);
case ast::StorageClass::kStorage:
- return EmitStorageBufferAccess(pre, out, expr, intrinsic);
+ return EmitStorageBufferAccess(out, expr, intrinsic);
default:
TINT_UNREACHABLE(Writer, diagnostics_)
<< "unsupported DecomposeMemoryAccess::Intrinsic storage class:"
@@ -458,23 +427,23 @@
if (auto* intrinsic = call->Target()->As<sem::Intrinsic>()) {
if (intrinsic->IsTexture()) {
- return EmitTextureCall(pre, out, expr, intrinsic);
+ return EmitTextureCall(out, expr, intrinsic);
} else if (intrinsic->Type() == sem::IntrinsicType::kSelect) {
- return EmitSelectCall(pre, out, expr);
+ return EmitSelectCall(out, expr);
} else if (intrinsic->Type() == sem::IntrinsicType::kFrexp) {
- return EmitFrexpCall(pre, out, expr, intrinsic);
+ return EmitFrexpCall(out, expr, intrinsic);
} else if (intrinsic->Type() == sem::IntrinsicType::kIsNormal) {
- return EmitIsNormalCall(pre, out, expr, intrinsic);
+ return EmitIsNormalCall(out, expr, intrinsic);
} else if (intrinsic->Type() == sem::IntrinsicType::kIgnore) {
- return EmitExpression(pre, out, expr->params()[0]);
+ return EmitExpression(out, expr->params()[0]);
} else if (intrinsic->IsDataPacking()) {
- return EmitDataPackingCall(pre, out, expr, intrinsic);
+ return EmitDataPackingCall(out, expr, intrinsic);
} else if (intrinsic->IsDataUnpacking()) {
- return EmitDataUnpackingCall(pre, out, expr, intrinsic);
+ return EmitDataUnpackingCall(out, expr, intrinsic);
} else if (intrinsic->IsBarrier()) {
- return EmitBarrierCall(pre, out, intrinsic);
+ return EmitBarrierCall(out, intrinsic);
} else if (intrinsic->IsAtomic()) {
- return EmitWorkgroupAtomicCall(pre, out, expr, intrinsic);
+ return EmitWorkgroupAtomicCall(out, expr, intrinsic);
}
auto name = generate_builtin_name(intrinsic);
if (name.empty()) {
@@ -490,7 +459,7 @@
}
first = false;
- if (!EmitExpression(pre, out, param)) {
+ if (!EmitExpression(out, param)) {
return false;
}
}
@@ -519,7 +488,7 @@
}
first = false;
- if (!EmitExpression(pre, out, param)) {
+ if (!EmitExpression(out, param)) {
return false;
}
}
@@ -530,7 +499,6 @@
}
bool GeneratorImpl::EmitUniformBufferAccess(
- std::ostream& pre,
std::ostream& out,
ast::CallExpression* expr,
const transform::DecomposeMemoryAccess::Intrinsic* intrinsic) {
@@ -538,13 +506,12 @@
std::string scalar_offset = generate_name("scalar_offset");
{
- std::stringstream ss;
- ss << "const int " << scalar_offset << " = (";
- if (!EmitExpression(pre, ss, params[1])) { // offset
+ auto pre = line();
+ pre << "const int " << scalar_offset << " = (";
+ if (!EmitExpression(pre, params[1])) { // offset
return false;
}
- make_indent(ss << ") / 4;" << std::endl);
- pre << ss.str();
+ pre << ") / 4;";
}
using Op = transform::DecomposeMemoryAccess::Intrinsic::Op;
@@ -558,7 +525,7 @@
return result;
};
auto load_scalar = [&]() {
- if (!EmitExpression(pre, out, params[0])) { // buffer
+ if (!EmitExpression(out, params[0])) { // buffer
return false;
}
out << "[" << scalar_offset << " / 4][" << scalar_offset << " % 4]";
@@ -567,14 +534,15 @@
// Has a minimum alignment of 8 bytes, so is either .xy or .zw
auto load_vec2 = [&] {
std::string ubo_load = generate_name("ubo_load");
- std::stringstream ss;
- ss << "uint4 " << ubo_load << " = ";
- if (!EmitExpression(pre, ss, params[0])) { // buffer
- return false;
+
+ {
+ auto pre = line();
+ pre << "uint4 " << ubo_load << " = ";
+ if (!EmitExpression(pre, params[0])) { // buffer
+ return false;
+ }
+ pre << "[" << scalar_offset << " / 4];";
}
- ss << "[" << scalar_offset << " / 4]";
- make_indent(ss << ";" << std::endl);
- pre << ss.str();
out << "((" << scalar_offset << " & 2) ? " << ubo_load
<< ".zw : " << ubo_load << ".xy)";
@@ -582,7 +550,7 @@
};
// vec3 has a minimum alignment of 16 bytes, so is just a .xyz swizzle
auto load_vec3 = [&] {
- if (!EmitExpression(pre, out, params[0])) { // buffer
+ if (!EmitExpression(out, params[0])) { // buffer
return false;
}
out << "[" << scalar_offset << " / 4].xyz";
@@ -590,7 +558,7 @@
};
// vec4 has a minimum alignment of 16 bytes, easiest case
auto load_vec4 = [&] {
- if (!EmitExpression(pre, out, params[0])) { // buffer
+ if (!EmitExpression(out, params[0])) { // buffer
return false;
}
out << "[" << scalar_offset << " / 4]";
@@ -637,7 +605,6 @@
}
bool GeneratorImpl::EmitStorageBufferAccess(
- std::ostream& pre,
std::ostream& out,
ast::CallExpression* expr,
const transform::DecomposeMemoryAccess::Intrinsic* intrinsic) {
@@ -651,7 +618,7 @@
if (cast) {
out << cast << "(";
}
- if (!EmitExpression(pre, out, params[0])) { // buffer
+ if (!EmitExpression(out, params[0])) { // buffer
return false;
}
out << ".Load";
@@ -659,7 +626,7 @@
out << n;
}
ScopedParen sp(out);
- if (!EmitExpression(pre, out, params[1])) { // offset
+ if (!EmitExpression(out, params[1])) { // offset
return false;
}
if (cast) {
@@ -701,7 +668,7 @@
case Op::kStore: {
auto store = [&](int n) {
- if (!EmitExpression(pre, out, params[0])) { // buffer
+ if (!EmitExpression(out, params[0])) { // buffer
return false;
}
out << ".Store";
@@ -709,12 +676,12 @@
out << n;
}
ScopedParen sp1(out);
- if (!EmitExpression(pre, out, params[1])) { // offset
+ if (!EmitExpression(out, params[1])) { // offset
return false;
}
out << ", asuint";
ScopedParen sp2(out);
- if (!EmitExpression(pre, out, params[2])) { // value
+ if (!EmitExpression(out, params[2])) { // value
return false;
}
return true;
@@ -761,7 +728,7 @@
case Op::kAtomicXor:
case Op::kAtomicExchange:
case Op::kAtomicCompareExchangeWeak:
- return EmitStorageAtomicCall(pre, out, expr, intrinsic->op);
+ return EmitStorageAtomicCall(out, expr, intrinsic->op);
}
TINT_UNREACHABLE(Writer, diagnostics_)
@@ -771,55 +738,83 @@
}
bool GeneratorImpl::EmitStorageAtomicCall(
- std::ostream& pre,
std::ostream& out,
ast::CallExpression* expr,
transform::DecomposeMemoryAccess::Intrinsic::Op op) {
using Op = transform::DecomposeMemoryAccess::Intrinsic::Op;
- std::stringstream ss;
std::string result = generate_name("atomic_result");
auto* result_ty = TypeOf(expr);
if (!result_ty->Is<sem::Void>()) {
- if (!EmitTypeAndName(ss, TypeOf(expr), ast::StorageClass::kNone,
+ auto pre = line();
+ if (!EmitTypeAndName(pre, TypeOf(expr), ast::StorageClass::kNone,
ast::Access::kUndefined, result)) {
return false;
}
- ss << " = ";
- if (!EmitZeroValue(ss, result_ty)) {
+ pre << " = ";
+ if (!EmitZeroValue(pre, result_ty)) {
return false;
}
- make_indent(ss << ";" << std::endl);
+ pre << ";";
}
auto* buffer = expr->params()[0];
auto* offset = expr->params()[1];
+ auto call_buffer_method = [&](const char* name) {
+ auto pre = line();
+ if (!EmitExpression(pre, buffer)) {
+ return false;
+ }
+ pre << "." << name;
+ {
+ ScopedParen sp(pre);
+ if (!EmitExpression(pre, offset)) {
+ return false;
+ }
+
+ for (size_t i = 1; i < expr->params().size() - 1; i++) {
+ auto* arg = expr->params()[i];
+ pre << ", ";
+ if (!EmitExpression(pre, arg)) {
+ return false;
+ }
+ }
+
+ pre << ", " << result;
+ }
+ pre << ";";
+
+ out << result;
+ return true;
+ };
+
switch (op) {
case Op::kAtomicLoad: {
// HLSL does not have an InterlockedLoad, so we emulate it with
// InterlockedOr using 0 as the OR value
- if (!EmitExpression(pre, ss, buffer)) {
+ auto pre = line();
+ if (!EmitExpression(pre, buffer)) {
return false;
}
- ss << ".InterlockedOr";
+ pre << ".InterlockedOr";
{
- ScopedParen sp(ss);
- if (!EmitExpression(pre, ss, offset)) {
+ ScopedParen sp(pre);
+ if (!EmitExpression(pre, offset)) {
return false;
}
- ss << ", 0, " << result;
+ pre << ", 0, " << result;
}
- make_indent(ss << ";" << std::endl);
- pre << ss.str();
+ pre << ";";
out << result;
return true;
}
case Op::kAtomicStore: {
// HLSL does not have an InterlockedStore, so we emulate it with
// InterlockedExchange and discard the returned value
+ auto pre = line();
auto* value = expr->params()[2];
auto* value_ty = TypeOf(value);
if (!EmitTypeAndName(pre, value_ty, ast::StorageClass::kNone,
@@ -830,19 +825,19 @@
if (!EmitZeroValue(pre, value_ty)) {
return false;
}
- make_indent(pre << ";" << std::endl);
+ pre << ";";
- if (!EmitExpression(pre, out, buffer)) {
+ if (!EmitExpression(out, buffer)) {
return false;
}
out << ".InterlockedExchange";
{
ScopedParen sp(out);
- if (!EmitExpression(pre, out, offset)) {
+ if (!EmitExpression(out, offset)) {
return false;
}
out << ", ";
- if (!EmitExpression(pre, out, value)) {
+ if (!EmitExpression(out, value)) {
return false;
}
out << ", " << result;
@@ -854,174 +849,163 @@
auto* value = expr->params()[3];
std::string compare = generate_name("atomic_compare_value");
- if (!EmitTypeAndName(ss, TypeOf(compare_value), ast::StorageClass::kNone,
- ast::Access::kUndefined, compare)) {
- return false;
- }
- ss << " = ";
- if (!EmitExpression(pre, ss, compare_value)) {
- return false;
- }
- make_indent(ss << ";" << std::endl);
-
- if (!EmitExpression(pre, ss, buffer)) {
- return false;
- }
- ss << ".InterlockedCompareExchange";
- {
- ScopedParen sp(ss);
- if (!EmitExpression(pre, ss, offset)) {
+ { // T atomic_compare_value = compare_value;
+ auto pre = line();
+ if (!EmitTypeAndName(pre, TypeOf(compare_value),
+ ast::StorageClass::kNone, ast::Access::kUndefined,
+ compare)) {
return false;
}
- ss << ", " << compare << ", ";
- if (!EmitExpression(pre, ss, value)) {
+ pre << " = ";
+ if (!EmitExpression(pre, compare_value)) {
return false;
}
- ss << ", " << result << ".x";
+ pre << ";";
}
- make_indent(ss << ";" << std::endl);
+ { // buffer.InterlockedCompareExchange(offset, compare, value, result.x);
+ auto pre = line();
+ if (!EmitExpression(pre, buffer)) {
+ return false;
+ }
+ pre << ".InterlockedCompareExchange";
+ {
+ ScopedParen sp(pre);
+ if (!EmitExpression(pre, offset)) {
+ return false;
+ }
+ pre << ", " << compare << ", ";
+ if (!EmitExpression(pre, value)) {
+ return false;
+ }
+ pre << ", " << result << ".x";
+ }
+ pre << ";";
+ }
+ { // result.y = result.x == compare;
+ line() << result << ".y = " << result << ".x == " << compare << ";";
+ }
- ss << result << ".y = " << result << ".x == " << compare;
- make_indent(ss << ";" << std::endl);
-
- pre << ss.str();
out << result;
return true;
}
case Op::kAtomicAdd:
- if (!EmitExpression(pre, ss, buffer)) {
- return false;
- }
- ss << ".InterlockedAdd";
- break;
+ return call_buffer_method("InterlockedAdd");
+
case Op::kAtomicMax:
- if (!EmitExpression(pre, ss, buffer)) {
- return false;
- }
- ss << ".InterlockedMax";
- break;
+ return call_buffer_method("InterlockedMax");
+
case Op::kAtomicMin:
- if (!EmitExpression(pre, ss, buffer)) {
- return false;
- }
- ss << ".InterlockedMin";
- break;
+ return call_buffer_method("InterlockedMin");
+
case Op::kAtomicAnd:
- if (!EmitExpression(pre, ss, buffer)) {
- return false;
- }
- ss << ".InterlockedAnd";
- break;
+ return call_buffer_method("InterlockedAnd");
+
case Op::kAtomicOr:
- if (!EmitExpression(pre, ss, buffer)) {
- return false;
- }
- ss << ".InterlockedOr";
- break;
+ return call_buffer_method("InterlockedOr");
+
case Op::kAtomicXor:
- if (!EmitExpression(pre, ss, buffer)) {
- return false;
- }
- ss << ".InterlockedXor";
- break;
+ return call_buffer_method("InterlockedXor");
+
case Op::kAtomicExchange:
- if (!EmitExpression(pre, ss, buffer)) {
- return false;
- }
- ss << ".InterlockedExchange";
- break;
+ return call_buffer_method("InterlockedExchange");
default:
- TINT_UNREACHABLE(Writer, diagnostics_)
- << "unsupported atomic DecomposeMemoryAccess::Intrinsic::Op: "
- << static_cast<int>(op);
- return false;
+ break;
}
- {
- ScopedParen sp(ss);
- if (!EmitExpression(pre, ss, offset)) {
- return false;
- }
-
- for (size_t i = 1; i < expr->params().size() - 1; i++) {
- auto* arg = expr->params()[i];
- ss << ", ";
- if (!EmitExpression(pre, ss, arg)) {
- return false;
- }
- }
-
- ss << ", " << result;
- }
-
- make_indent(ss << ";" << std::endl);
- pre << ss.str();
- out << result;
-
- return true;
+ TINT_UNREACHABLE(Writer, diagnostics_)
+ << "unsupported atomic DecomposeMemoryAccess::Intrinsic::Op: "
+ << static_cast<int>(op);
+ return false;
}
-bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& pre,
- std::ostream& out,
+bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out,
ast::CallExpression* expr,
const sem::Intrinsic* intrinsic) {
- std::stringstream ss;
std::string result = generate_name("atomic_result");
if (!intrinsic->ReturnType()->Is<sem::Void>()) {
- if (!EmitTypeAndName(ss, intrinsic->ReturnType(), ast::StorageClass::kNone,
+ auto pre = line();
+ if (!EmitTypeAndName(pre, intrinsic->ReturnType(), ast::StorageClass::kNone,
ast::Access::kUndefined, result)) {
return false;
}
- ss << " = ";
- if (!EmitZeroValue(ss, intrinsic->ReturnType())) {
+ pre << " = ";
+ if (!EmitZeroValue(pre, intrinsic->ReturnType())) {
return false;
}
- make_indent(ss << ";" << std::endl);
+ pre << ";";
}
+ auto call = [&](const char* name) {
+ auto pre = line();
+ pre << name;
+
+ {
+ ScopedParen sp(pre);
+ for (size_t i = 0; i < expr->params().size(); i++) {
+ auto* arg = expr->params()[i];
+ if (i > 0) {
+ pre << ", ";
+ }
+ if (!EmitExpression(pre, arg)) {
+ return false;
+ }
+ }
+
+ pre << ", " << result;
+ }
+
+ pre << ";";
+
+ out << result;
+ return true;
+ };
+
switch (intrinsic->Type()) {
case sem::IntrinsicType::kAtomicLoad: {
// HLSL does not have an InterlockedLoad, so we emulate it with
// InterlockedOr using 0 as the OR value
- ss << "InterlockedOr";
+ auto pre = line();
+ pre << "InterlockedOr";
{
- ScopedParen sp(ss);
- if (!EmitExpression(pre, ss, expr->params()[0])) {
+ ScopedParen sp(pre);
+ if (!EmitExpression(pre, expr->params()[0])) {
return false;
}
- ss << ", 0, " << result;
+ pre << ", 0, " << result;
}
- make_indent(ss << ";" << std::endl);
+ pre << ";";
- pre << ss.str();
out << result;
return true;
}
case sem::IntrinsicType::kAtomicStore: {
// HLSL does not have an InterlockedStore, so we emulate it with
// InterlockedExchange and discard the returned value
- auto* value_ty = intrinsic->Parameters()[1].type;
- if (!EmitTypeAndName(pre, value_ty, ast::StorageClass::kNone,
- ast::Access::kUndefined, result)) {
- return false;
+ { // T result = 0;
+ auto pre = line();
+ auto* value_ty = intrinsic->Parameters()[1].type;
+ if (!EmitTypeAndName(pre, value_ty, ast::StorageClass::kNone,
+ ast::Access::kUndefined, result)) {
+ return false;
+ }
+ pre << " = ";
+ if (!EmitZeroValue(pre, value_ty)) {
+ return false;
+ }
+ pre << ";";
}
- pre << " = ";
- if (!EmitZeroValue(pre, value_ty)) {
- return false;
- }
- make_indent(pre << ";" << std::endl);
out << "InterlockedExchange";
{
ScopedParen sp(out);
- if (!EmitExpression(pre, out, expr->params()[0])) {
+ if (!EmitExpression(out, expr->params()[0])) {
return false;
}
out << ", ";
- if (!EmitExpression(pre, out, expr->params()[1])) {
+ if (!EmitExpression(out, expr->params()[1])) {
return false;
}
out << ", " << result;
@@ -1034,116 +1018,102 @@
auto* value = expr->params()[2];
std::string compare = generate_name("atomic_compare_value");
- if (!EmitTypeAndName(ss, TypeOf(compare_value), ast::StorageClass::kNone,
- ast::Access::kUndefined, compare)) {
- return false;
- }
- ss << " = ";
- if (!EmitExpression(pre, ss, compare_value)) {
- return false;
- }
- make_indent(ss << ";" << std::endl);
- ss << "InterlockedCompareExchange";
- {
- ScopedParen sp(ss);
- if (!EmitExpression(pre, ss, dest)) {
+ { // T compare_value = <compare_value>;
+ auto pre = line();
+ if (!EmitTypeAndName(pre, TypeOf(compare_value),
+ ast::StorageClass::kNone, ast::Access::kUndefined,
+ compare)) {
return false;
}
- ss << ", " << compare << ", ";
- if (!EmitExpression(pre, ss, value)) {
+ pre << " = ";
+ if (!EmitExpression(pre, compare_value)) {
return false;
}
- ss << ", " << result << ".x";
+ pre << ";";
}
- make_indent(ss << ";" << std::endl);
- ss << result << ".y = " << result << ".x == " << compare;
- make_indent(ss << ";" << std::endl);
+ { // InterlockedCompareExchange(dst, compare, value, result.x);
+ auto pre = line();
+ pre << "InterlockedCompareExchange";
+ {
+ ScopedParen sp(pre);
+ if (!EmitExpression(pre, dest)) {
+ return false;
+ }
+ pre << ", " << compare << ", ";
+ if (!EmitExpression(pre, value)) {
+ return false;
+ }
+ pre << ", " << result << ".x";
+ }
+ pre << ";";
+ }
- pre << ss.str();
+ { // result.y = result.x == compare;
+ line() << result << ".y = " << result << ".x == " << compare << ";";
+ }
+
out << result;
return true;
}
case sem::IntrinsicType::kAtomicAdd:
- ss << "InterlockedAdd";
- break;
+ return call("InterlockedAdd");
+
case sem::IntrinsicType::kAtomicMax:
- ss << "InterlockedMax";
- break;
+ return call("InterlockedMax");
+
case sem::IntrinsicType::kAtomicMin:
- ss << "InterlockedMin";
- break;
+ return call("InterlockedMin");
+
case sem::IntrinsicType::kAtomicAnd:
- ss << "InterlockedAnd";
- break;
+ return call("InterlockedAnd");
+
case sem::IntrinsicType::kAtomicOr:
- ss << "InterlockedOr";
- break;
+ return call("InterlockedOr");
+
case sem::IntrinsicType::kAtomicXor:
- ss << "InterlockedXor";
- break;
+ return call("InterlockedXor");
+
case sem::IntrinsicType::kAtomicExchange:
- ss << "InterlockedExchange";
- break;
+ return call("InterlockedExchange");
default:
- TINT_UNREACHABLE(Writer, diagnostics_)
- << "unsupported atomic intrinsic: " << intrinsic->Type();
- return false;
+ break;
}
- {
- ScopedParen sp(ss);
- for (size_t i = 0; i < expr->params().size(); i++) {
- auto* arg = expr->params()[i];
- if (i > 0) {
- ss << ", ";
- }
- if (!EmitExpression(pre, ss, arg)) {
- return false;
- }
- }
-
- ss << ", " << result;
- }
- make_indent(ss << ";" << std::endl);
-
- pre << ss.str();
- out << result;
-
- return true;
+ TINT_UNREACHABLE(Writer, diagnostics_)
+ << "unsupported atomic intrinsic: " << intrinsic->Type();
+ return false;
}
-bool GeneratorImpl::EmitSelectCall(std::ostream& pre,
- std::ostream& out,
+bool GeneratorImpl::EmitSelectCall(std::ostream& out,
ast::CallExpression* expr) {
auto* expr_true = expr->params()[0];
auto* expr_false = expr->params()[1];
auto* expr_cond = expr->params()[2];
ScopedParen paren(out);
- if (!EmitExpression(pre, out, expr_cond)) {
+ if (!EmitExpression(out, expr_cond)) {
return false;
}
out << " ? ";
- if (!EmitExpression(pre, out, expr_true)) {
+ if (!EmitExpression(out, expr_true)) {
return false;
}
out << " : ";
- if (!EmitExpression(pre, out, expr_false)) {
+ if (!EmitExpression(out, expr_false)) {
return false;
}
return true;
}
-bool GeneratorImpl::EmitFrexpCall(std::ostream& pre,
- std::ostream& out,
+bool GeneratorImpl::EmitFrexpCall(std::ostream& out,
ast::CallExpression* expr,
const sem::Intrinsic* intrinsic) {
// Exponent is an integer in WGSL, but HLSL wants a float.
@@ -1159,37 +1129,35 @@
// Exponent is an integer, which HLSL does not have an overload for.
// We need to cast from a float.
- std::stringstream ss;
auto float_exp = generate_name(kTempNamePrefix);
- ss << "float" << width << " " << float_exp << ";";
-
- make_indent(ss << std::endl);
auto significand = generate_name(kTempNamePrefix);
- ss << "float" << width << " " << significand << " = frexp(";
- if (!EmitExpression(pre, ss, expr->params()[0])) {
- return false;
+ line() << "float" << width << " " << float_exp << ";";
+ {
+ auto pre = line();
+ pre << "float" << width << " " << significand << " = frexp(";
+ if (!EmitExpression(pre, expr->params()[0])) {
+ return false;
+ }
+ pre << ", " << float_exp << ");";
}
- ss << ", " << float_exp << ");";
+ {
+ auto pre = line();
+ if (!EmitExpression(pre, expr->params()[1])) {
+ return false;
+ }
+ pre << " = ";
+ if (!EmitType(pre, exponent.type->UnwrapPtr(), ast::StorageClass::kNone,
+ ast::Access::kUndefined, "")) {
+ return false;
+ }
+ pre << "(" << float_exp << ");";
+ }
- make_indent(ss << std::endl);
- if (!EmitExpression(pre, ss, expr->params()[1])) {
- return false;
- }
- ss << " = ";
- if (!EmitType(ss, exponent.type->UnwrapPtr(), ast::StorageClass::kNone,
- ast::Access::kUndefined, "")) {
- return false;
- }
- ss << "(" << float_exp << ");";
-
- make_indent(ss << std::endl);
- pre << ss.str();
out << significand;
return true;
}
-bool GeneratorImpl::EmitIsNormalCall(std::ostream& pre,
- std::ostream& out,
+bool GeneratorImpl::EmitIsNormalCall(std::ostream& out,
ast::CallExpression* expr,
const sem::Intrinsic* intrinsic) {
// HLSL doesn't have a isNormal intrinsic, we need to emulate
@@ -1207,33 +1175,29 @@
auto exponent = generate_name("tint_isnormal_exponent");
auto clamped = generate_name("tint_isnormal_clamped");
- std::stringstream ss;
- ss << "uint" << width << " " << exponent << " = asuint(";
- if (!EmitExpression(pre, ss, expr->params()[0])) {
- return false;
+ {
+ auto pre = line();
+ pre << "uint" << width << " " << exponent << " = asuint(";
+ if (!EmitExpression(pre, expr->params()[0])) {
+ return false;
+ }
+ pre << ") & " << kExponentMask << ";";
}
- ss << ") & " << kExponentMask << ";";
-
- make_indent(ss << std::endl);
- ss << "uint" << width << " " << clamped << " = "
- << "clamp(" << exponent << ", " << kMinNormalExponent << ", "
- << kMaxNormalExponent << ");";
-
- make_indent(ss << std::endl);
- pre << ss.str();
+ line() << "uint" << width << " " << clamped << " = "
+ << "clamp(" << exponent << ", " << kMinNormalExponent << ", "
+ << kMaxNormalExponent << ");";
out << "(" << clamped << " == " << exponent << ")";
return true;
}
-bool GeneratorImpl::EmitDataPackingCall(std::ostream& pre,
- std::ostream& out,
+bool GeneratorImpl::EmitDataPackingCall(std::ostream& out,
ast::CallExpression* expr,
const sem::Intrinsic* intrinsic) {
auto* param = expr->params()[0];
auto tmp_name = generate_name(kTempNamePrefix);
std::ostringstream expr_out;
- if (!EmitExpression(pre, expr_out, param)) {
+ if (!EmitExpression(expr_out, param)) {
return false;
}
uint32_t dims = 2;
@@ -1253,15 +1217,18 @@
case sem::IntrinsicType::kPack4x8snorm:
case sem::IntrinsicType::kPack4x8unorm:
case sem::IntrinsicType::kPack2x16snorm:
- case sem::IntrinsicType::kPack2x16unorm:
- pre << (is_signed ? "" : "u") << "int" << dims << " " << tmp_name << " = "
- << (is_signed ? "" : "u") << "int" << dims << "(round(clamp("
- << expr_out.str() << ", " << (is_signed ? "-1.0" : "0.0")
- << ", 1.0) * " << scale << ".0))";
- if (is_signed) {
- pre << " & " << (dims == 4 ? "0xff" : "0xffff");
+ case sem::IntrinsicType::kPack2x16unorm: {
+ {
+ auto pre = line();
+ pre << (is_signed ? "" : "u") << "int" << dims << " " << tmp_name
+ << " = " << (is_signed ? "" : "u") << "int" << dims
+ << "(round(clamp(" << expr_out.str() << ", "
+ << (is_signed ? "-1.0" : "0.0") << ", 1.0) * " << scale << ".0))";
+ if (is_signed) {
+ pre << " & " << (dims == 4 ? "0xff" : "0xffff");
+ }
+ pre << ";";
}
- pre << ";\n";
if (is_signed) {
out << "asuint";
}
@@ -1272,10 +1239,13 @@
}
out << ")";
break;
- case sem::IntrinsicType::kPack2x16float:
- pre << "uint2 " << tmp_name << " = f32tof16(" << expr_out.str() << ");\n";
+ }
+ case sem::IntrinsicType::kPack2x16float: {
+ line() << "uint2 " << tmp_name << " = f32tof16(" << expr_out.str()
+ << ");";
out << "(" << tmp_name << ".x | " << tmp_name << ".y << 16)";
break;
+ }
default:
diagnostics_.add_error(
diag::System::Writer,
@@ -1286,14 +1256,13 @@
return true;
}
-bool GeneratorImpl::EmitDataUnpackingCall(std::ostream& pre,
- std::ostream& out,
+bool GeneratorImpl::EmitDataUnpackingCall(std::ostream& out,
ast::CallExpression* expr,
const sem::Intrinsic* intrinsic) {
auto* param = expr->params()[0];
auto tmp_name = generate_name(kTempNamePrefix);
std::ostringstream expr_out;
- if (!EmitExpression(pre, expr_out, param)) {
+ if (!EmitExpression(expr_out, param)) {
return false;
}
uint32_t dims = 2;
@@ -1313,16 +1282,19 @@
case sem::IntrinsicType::kUnpack4x8snorm:
case sem::IntrinsicType::kUnpack2x16snorm: {
auto tmp_name2 = generate_name(kTempNamePrefix);
- pre << "int " << tmp_name2 << " = int(" << expr_out.str() << ");\n";
- // Perform sign extension on the converted values.
- pre << "int" << dims << " " << tmp_name << " = int" << dims << "(";
- if (dims == 2) {
- pre << tmp_name2 << " << 16, " << tmp_name2 << ") >> 16";
- } else {
- pre << tmp_name2 << " << 24, " << tmp_name2 << " << 16, " << tmp_name2
- << " << 8, " << tmp_name2 << ") >> 24";
+ line() << "int " << tmp_name2 << " = int(" << expr_out.str() << ");";
+ { // Perform sign extension on the converted values.
+ auto pre = line();
+ pre << "int" << dims << " " << tmp_name << " = int" << dims << "(";
+ if (dims == 2) {
+ pre << tmp_name2 << " << 16, " << tmp_name2 << ") >> 16";
+ } else {
+ pre << tmp_name2 << " << 24, " << tmp_name2 << " << 16, " << tmp_name2
+ << " << 8, " << tmp_name2 << ") >> 24";
+ }
+ pre << ";";
}
- pre << ";\n";
+
out << "clamp(float" << dims << "(" << tmp_name << ") / " << scale
<< ".0, " << (is_signed ? "-1.0" : "0.0") << ", 1.0)";
break;
@@ -1330,21 +1302,24 @@
case sem::IntrinsicType::kUnpack4x8unorm:
case sem::IntrinsicType::kUnpack2x16unorm: {
auto tmp_name2 = generate_name(kTempNamePrefix);
- pre << "uint " << tmp_name2 << " = " << expr_out.str() << ";\n";
- pre << "uint" << dims << " " << tmp_name << " = uint" << dims << "(";
- pre << tmp_name2 << " & " << (dims == 2 ? "0xffff" : "0xff") << ", ";
- if (dims == 4) {
- pre << "(" << tmp_name2 << " >> " << (32 / dims) << ") & 0xff, ("
- << tmp_name2 << " >> 16) & 0xff, " << tmp_name2 << " >> 24";
- } else {
- pre << tmp_name2 << " >> " << (32 / dims);
+ line() << "uint " << tmp_name2 << " = " << expr_out.str() << ";";
+ {
+ auto pre = line();
+ pre << "uint" << dims << " " << tmp_name << " = uint" << dims << "(";
+ pre << tmp_name2 << " & " << (dims == 2 ? "0xffff" : "0xff") << ", ";
+ if (dims == 4) {
+ pre << "(" << tmp_name2 << " >> " << (32 / dims) << ") & 0xff, ("
+ << tmp_name2 << " >> 16) & 0xff, " << tmp_name2 << " >> 24";
+ } else {
+ pre << tmp_name2 << " >> " << (32 / dims);
+ }
+ pre << ");";
}
- pre << ");\n";
out << "float" << dims << "(" << tmp_name << ") / " << scale << ".0";
break;
}
case sem::IntrinsicType::kUnpack2x16float:
- pre << "uint " << tmp_name << " = " << expr_out.str() << ";\n";
+ line() << "uint " << tmp_name << " = " << expr_out.str() << ";";
out << "f16tof32(uint2(" << tmp_name << " & 0xffff, " << tmp_name
<< " >> 16))";
break;
@@ -1358,8 +1333,7 @@
return true;
}
-bool GeneratorImpl::EmitBarrierCall(std::ostream&,
- std::ostream& out,
+bool GeneratorImpl::EmitBarrierCall(std::ostream& out,
const sem::Intrinsic* intrinsic) {
// TODO(crbug.com/tint/661): Combine sequential barriers to a single
// instruction.
@@ -1375,8 +1349,7 @@
return true;
}
-bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
- std::ostream& out,
+bool GeneratorImpl::EmitTextureCall(std::ostream& out,
ast::CallExpression* expr,
const sem::Intrinsic* intrinsic) {
using Usage = sem::ParameterUsage;
@@ -1519,49 +1492,47 @@
// Declare a variable to hold the queried texture info
auto dims = generate_name(kTempNamePrefix);
-
if (num_dimensions == 1) {
- pre << "int " << dims << ";";
+ line() << "int " << dims << ";";
} else {
- pre << "int" << num_dimensions << " " << dims << ";";
+ line() << "int" << num_dimensions << " " << dims << ";";
}
- pre << std::endl;
- make_indent(pre);
-
- if (!EmitExpression(pre, pre, texture)) {
- return false;
- }
- pre << ".GetDimensions(";
-
- if (level_arg) {
- if (!EmitExpression(pre, pre, level_arg)) {
+ { // texture.GetDimensions(...)
+ auto pre = line();
+ if (!EmitExpression(pre, texture)) {
return false;
}
- pre << ", ";
- } else if (intrinsic->Type() == sem::IntrinsicType::kTextureNumLevels) {
- pre << "0, ";
- }
+ pre << ".GetDimensions(";
- if (num_dimensions == 1) {
- pre << dims;
- } else {
- static constexpr char xyzw[] = {'x', 'y', 'z', 'w'};
- if (num_dimensions < 0 || num_dimensions > 4) {
- TINT_ICE(Writer, diagnostics_)
- << "vector dimensions are " << num_dimensions;
- return false;
- }
- for (int i = 0; i < num_dimensions; i++) {
- if (i > 0) {
- pre << ", ";
+ if (level_arg) {
+ if (!EmitExpression(pre, level_arg)) {
+ return false;
}
- pre << dims << "." << xyzw[i];
+ pre << ", ";
+ } else if (intrinsic->Type() == sem::IntrinsicType::kTextureNumLevels) {
+ pre << "0, ";
}
- }
- pre << ");" << std::endl;
- make_indent(pre);
+ if (num_dimensions == 1) {
+ pre << dims;
+ } else {
+ static constexpr char xyzw[] = {'x', 'y', 'z', 'w'};
+ if (num_dimensions < 0 || num_dimensions > 4) {
+ TINT_ICE(Writer, diagnostics_)
+ << "vector dimensions are " << num_dimensions;
+ return false;
+ }
+ for (int i = 0; i < num_dimensions; i++) {
+ if (i > 0) {
+ pre << ", ";
+ }
+ pre << dims << "." << xyzw[i];
+ }
+ }
+
+ pre << ");";
+ }
// The out parameters of the GetDimensions() call is now in temporary
// `dims` variable. This may be packed with other data, so the final
@@ -1573,7 +1544,7 @@
break;
}
- if (!EmitExpression(pre, out, texture))
+ if (!EmitExpression(out, texture))
return false;
bool pack_mip_in_coords = false;
@@ -1616,7 +1587,7 @@
}
if (auto* sampler = arg(Usage::kSampler)) {
- if (!EmitExpression(pre, out, sampler))
+ if (!EmitExpression(out, sampler))
return false;
out << ", ";
}
@@ -1633,7 +1604,7 @@
auto* stmt = builder_.Sem().Get(vector)->Stmt();
builder_.Sem().Add(zero, builder_.create<sem::Expression>(zero, i32, stmt));
auto* packed = AppendVector(&builder_, vector, zero);
- return EmitExpression(pre, out, packed);
+ return EmitExpression(out, packed);
};
if (auto* array_index = arg(Usage::kArrayIndex)) {
@@ -1644,7 +1615,7 @@
return false;
}
} else {
- if (!EmitExpression(pre, out, packed)) {
+ if (!EmitExpression(out, packed)) {
return false;
}
}
@@ -1654,7 +1625,7 @@
return false;
}
} else {
- if (!EmitExpression(pre, out, param_coords)) {
+ if (!EmitExpression(out, param_coords)) {
return false;
}
}
@@ -1663,7 +1634,7 @@
Usage::kDdy, Usage::kSampleIndex, Usage::kOffset}) {
if (auto* e = arg(usage)) {
out << ", ";
- if (!EmitExpression(pre, out, e)) {
+ if (!EmitExpression(out, e)) {
return false;
}
}
@@ -1671,7 +1642,7 @@
if (intrinsic->Type() == sem::IntrinsicType::kTextureStore) {
out << "] = ";
- if (!EmitExpression(pre, out, arg(Usage::kValue))) {
+ if (!EmitExpression(out, arg(Usage::kValue))) {
return false;
}
} else {
@@ -1792,69 +1763,53 @@
return "";
}
-bool GeneratorImpl::EmitCase(std::ostream& out, ast::CaseStatement* stmt) {
- make_indent(out);
-
+bool GeneratorImpl::EmitCase(ast::CaseStatement* stmt) {
if (stmt->IsDefault()) {
- out << "default:";
+ line() << "default: {";
} else {
- bool first = true;
for (auto* selector : stmt->selectors()) {
- if (!first) {
- out << std::endl;
- make_indent(out);
- }
- first = false;
-
+ auto out = line();
out << "case ";
if (!EmitLiteral(out, selector)) {
return false;
}
out << ":";
+ if (selector == stmt->selectors().back()) {
+ out << " {";
+ }
}
}
- out << " {" << std::endl;
-
- increment_indent();
-
- for (auto* s : *stmt->body()) {
- if (!EmitStatement(out, s)) {
+ {
+ ScopedIndent si(this);
+ if (!EmitStatements(stmt->body()->statements())) {
return false;
}
+ if (!last_is_break_or_fallthrough(stmt->body())) {
+ line() << "break;";
+ }
}
- if (!last_is_break_or_fallthrough(stmt->body())) {
- make_indent(out);
- out << "break;" << std::endl;
- }
-
- decrement_indent();
- make_indent(out);
- out << "}" << std::endl;
+ line() << "}";
return true;
}
-bool GeneratorImpl::EmitConstructor(std::ostream& pre,
- std::ostream& out,
+bool GeneratorImpl::EmitConstructor(std::ostream& out,
ast::ConstructorExpression* expr) {
if (auto* scalar = expr->As<ast::ScalarConstructorExpression>()) {
- return EmitScalarConstructor(pre, out, scalar);
+ return EmitScalarConstructor(out, scalar);
}
- return EmitTypeConstructor(pre, out,
- expr->As<ast::TypeConstructorExpression>());
+ return EmitTypeConstructor(out, expr->As<ast::TypeConstructorExpression>());
}
bool GeneratorImpl::EmitScalarConstructor(
- std::ostream&,
std::ostream& out,
ast::ScalarConstructorExpression* expr) {
return EmitLiteral(out, expr->literal());
}
-bool GeneratorImpl::EmitTypeConstructor(std::ostream& pre,
- std::ostream& out,
+bool GeneratorImpl::EmitTypeConstructor(std::ostream& out,
ast::TypeConstructorExpression* expr) {
auto* type = TypeOf(expr)->UnwrapRef();
@@ -1897,7 +1852,7 @@
}
first = false;
- if (!EmitExpression(pre, out, e)) {
+ if (!EmitExpression(out, e)) {
return false;
}
}
@@ -1910,49 +1865,45 @@
return true;
}
-bool GeneratorImpl::EmitContinue(std::ostream& out, ast::ContinueStatement*) {
- if (!emit_continuing_(out)) {
+bool GeneratorImpl::EmitContinue(ast::ContinueStatement*) {
+ if (!emit_continuing_()) {
return false;
}
- make_indent(out);
- out << "continue;" << std::endl;
+ line() << "continue;";
return true;
}
-bool GeneratorImpl::EmitDiscard(std::ostream& out, ast::DiscardStatement*) {
- make_indent(out);
+bool GeneratorImpl::EmitDiscard(ast::DiscardStatement*) {
// TODO(dsinclair): Verify this is correct when the discard semantics are
// defined for WGSL (https://github.com/gpuweb/gpuweb/issues/361)
- out << "discard;" << std::endl;
+ line() << "discard;";
return true;
}
-bool GeneratorImpl::EmitExpression(std::ostream& pre,
- std::ostream& out,
- ast::Expression* expr) {
+bool GeneratorImpl::EmitExpression(std::ostream& out, ast::Expression* expr) {
if (auto* a = expr->As<ast::ArrayAccessorExpression>()) {
- return EmitArrayAccessor(pre, out, a);
+ return EmitArrayAccessor(out, a);
}
if (auto* b = expr->As<ast::BinaryExpression>()) {
- return EmitBinary(pre, out, b);
+ return EmitBinary(out, b);
}
if (auto* b = expr->As<ast::BitcastExpression>()) {
- return EmitBitcast(pre, out, b);
+ return EmitBitcast(out, b);
}
if (auto* c = expr->As<ast::CallExpression>()) {
- return EmitCall(pre, out, c);
+ return EmitCall(out, c);
}
if (auto* c = expr->As<ast::ConstructorExpression>()) {
- return EmitConstructor(pre, out, c);
+ return EmitConstructor(out, c);
}
if (auto* i = expr->As<ast::IdentifierExpression>()) {
- return EmitIdentifier(pre, out, i);
+ return EmitIdentifier(out, i);
}
if (auto* m = expr->As<ast::MemberAccessorExpression>()) {
- return EmitMemberAccessor(pre, out, m);
+ return EmitMemberAccessor(out, m);
}
if (auto* u = expr->As<ast::UnaryOpExpression>()) {
- return EmitUnaryOp(pre, out, u);
+ return EmitUnaryOp(out, u);
}
diagnostics_.add_error(diag::System::Writer,
@@ -1960,71 +1911,60 @@
return false;
}
-bool GeneratorImpl::EmitIdentifier(std::ostream&,
- std::ostream& out,
+bool GeneratorImpl::EmitIdentifier(std::ostream& out,
ast::IdentifierExpression* expr) {
out << builder_.Symbols().NameFor(expr->symbol());
return true;
}
-bool GeneratorImpl::EmitIf(std::ostream& out, ast::IfStatement* stmt) {
- make_indent(out);
-
- std::ostringstream pre;
- std::ostringstream cond;
- if (!EmitExpression(pre, cond, stmt->condition())) {
- return false;
+bool GeneratorImpl::EmitIf(ast::IfStatement* stmt) {
+ {
+ auto out = line();
+ out << "if (";
+ if (!EmitExpression(out, stmt->condition())) {
+ return false;
+ }
+ out << ") {";
}
- std::ostringstream if_out;
- if_out << "if (" << cond.str() << ") ";
- if (!EmitBlock(if_out, stmt->body())) {
+ if (!EmitStatementsWithIndent(stmt->body()->statements())) {
return false;
}
for (auto* e : stmt->else_statements()) {
if (e->HasCondition()) {
- if_out << " else {" << std::endl;
-
+ line() << "} else {";
increment_indent();
- std::ostringstream else_pre;
- std::ostringstream else_cond_out;
- if (!EmitExpression(else_pre, else_cond_out, e->condition())) {
- return false;
+ {
+ auto out = line();
+ out << "if (";
+ if (!EmitExpression(out, e->condition())) {
+ return false;
+ }
+ out << ") {";
}
- if_out << else_pre.str();
-
- make_indent(if_out);
- if_out << "if (" << else_cond_out.str() << ") ";
} else {
- if_out << " else ";
+ line() << "} else {";
}
- if (!EmitBlock(if_out, e->body())) {
+ if (!EmitStatementsWithIndent(e->body()->statements())) {
return false;
}
}
- if_out << std::endl;
+
+ line() << "}";
for (auto* e : stmt->else_statements()) {
- if (!e->HasCondition()) {
- continue;
+ if (e->HasCondition()) {
+ decrement_indent();
+ line() << "}";
}
-
- decrement_indent();
- make_indent(if_out);
- if_out << "}" << std::endl;
}
-
- out << pre.str();
- out << if_out.str();
return true;
}
-bool GeneratorImpl::EmitFunction(std::ostream& out, ast::Function* func) {
- make_indent(out);
-
+bool GeneratorImpl::EmitFunction(ast::Function* func) {
auto* sem = builder_.Sem().Get(func);
if (ast::HasDecoration<ast::InternalDecoration>(func->decorations())) {
@@ -2032,74 +1972,75 @@
return true;
}
- if (!EmitType(out, sem->ReturnType(), ast::StorageClass::kNone,
- ast::Access::kReadWrite, "")) {
- return false;
- }
-
- out << " ";
-
- out << builder_.Symbols().NameFor(func->symbol()) << "(";
-
- bool first = true;
-
- for (auto* v : sem->Parameters()) {
- if (!first) {
- out << ", ";
- }
- first = false;
-
- auto const* type = v->Type();
-
- if (auto* ptr = type->As<sem::Pointer>()) {
- // Transform pointer parameters in to `inout` parameters.
- // The WGSL spec is highly restrictive in what can be passed in pointer
- // parameters, which allows for this transformation. See:
- // https://gpuweb.github.io/gpuweb/wgsl/#function-restriction
- out << "inout ";
- type = ptr->StoreType();
- }
-
- // Note: WGSL only allows for StorageClass::kNone on parameters, however the
- // sanitizer transforms generates load / store functions for storage or
- // uniform buffers. These functions have a buffer parameter with
- // StorageClass::kStorage or StorageClass::kUniform. This is required to
- // correctly translate the parameter to a [RW]ByteAddressBuffer for storage
- // buffers and a uint4[N] for uniform buffers.
- if (!EmitTypeAndName(
- out, type, v->StorageClass(), v->Access(),
- builder_.Symbols().NameFor(v->Declaration()->symbol()))) {
+ {
+ auto out = line();
+ if (!EmitType(out, sem->ReturnType(), ast::StorageClass::kNone,
+ ast::Access::kReadWrite, "")) {
return false;
}
+
+ out << " " << builder_.Symbols().NameFor(func->symbol()) << "(";
+
+ bool first = true;
+
+ for (auto* v : sem->Parameters()) {
+ if (!first) {
+ out << ", ";
+ }
+ first = false;
+
+ auto const* type = v->Type();
+
+ if (auto* ptr = type->As<sem::Pointer>()) {
+ // Transform pointer parameters in to `inout` parameters.
+ // The WGSL spec is highly restrictive in what can be passed in pointer
+ // parameters, which allows for this transformation. See:
+ // https://gpuweb.github.io/gpuweb/wgsl/#function-restriction
+ out << "inout ";
+ type = ptr->StoreType();
+ }
+
+ // Note: WGSL only allows for StorageClass::kNone on parameters, however
+ // the sanitizer transforms generates load / store functions for storage
+ // or uniform buffers. These functions have a buffer parameter with
+ // StorageClass::kStorage or StorageClass::kUniform. This is required to
+ // correctly translate the parameter to a [RW]ByteAddressBuffer for
+ // storage buffers and a uint4[N] for uniform buffers.
+ if (!EmitTypeAndName(
+ out, type, v->StorageClass(), v->Access(),
+ builder_.Symbols().NameFor(v->Declaration()->symbol()))) {
+ return false;
+ }
+ }
+ out << ") {";
}
- out << ") ";
-
- if (!EmitBlockAndNewline(out, func->body())) {
+ if (!EmitStatementsWithIndent(func->body()->statements())) {
return false;
}
+ line() << "}";
+
return true;
}
-bool GeneratorImpl::EmitGlobalVariable(std::ostream& out,
- ast::Variable* global) {
+bool GeneratorImpl::EmitGlobalVariable(ast::Variable* global) {
if (global->is_const()) {
- return EmitProgramConstVariable(out, global);
+ return EmitProgramConstVariable(global);
}
auto* sem = builder_.Sem().Get(global);
switch (sem->StorageClass()) {
case ast::StorageClass::kUniform:
- return EmitUniformVariable(out, sem);
+ return EmitUniformVariable(sem);
case ast::StorageClass::kStorage:
- return EmitStorageVariable(out, sem);
+ return EmitStorageVariable(sem);
case ast::StorageClass::kUniformConstant:
- return EmitHandleVariable(out, sem);
+ return EmitHandleVariable(sem);
case ast::StorageClass::kPrivate:
- return EmitPrivateVariable(out, sem);
+ return EmitPrivateVariable(sem);
case ast::StorageClass::kWorkgroup:
- return EmitWorkgroupVariable(out, sem);
+ return EmitWorkgroupVariable(sem);
default:
break;
}
@@ -2109,10 +2050,7 @@
return false;
}
-bool GeneratorImpl::EmitUniformVariable(std::ostream& out,
- const sem::Variable* var) {
- make_indent(out);
-
+bool GeneratorImpl::EmitUniformVariable(const sem::Variable* var) {
auto* decl = var->Declaration();
auto binding_point = decl->binding_point();
auto* type = var->Type()->UnwrapRef();
@@ -2125,29 +2063,28 @@
}
auto name = builder_.Symbols().NameFor(decl->symbol());
- out << "cbuffer cbuffer_" << name << RegisterAndSpace('b', binding_point)
- << " {" << std::endl;
+ line() << "cbuffer cbuffer_" << name << RegisterAndSpace('b', binding_point)
+ << " {";
- increment_indent();
- make_indent(out);
-
- if (!EmitTypeAndName(out, type, ast::StorageClass::kUniform, var->Access(),
- name)) {
- return false;
+ {
+ ScopedIndent si(this);
+ auto out = line();
+ if (!EmitTypeAndName(out, type, ast::StorageClass::kUniform, var->Access(),
+ name)) {
+ return false;
+ }
+ out << ";";
}
- out << ";" << std::endl;
- decrement_indent();
- out << "};" << std::endl;
+
+ line() << "};";
return true;
}
-bool GeneratorImpl::EmitStorageVariable(std::ostream& out,
- const sem::Variable* var) {
- make_indent(out);
-
+bool GeneratorImpl::EmitStorageVariable(const sem::Variable* var) {
auto* decl = var->Declaration();
auto* type = var->Type()->UnwrapRef();
+ auto out = line();
if (!EmitTypeAndName(out, type, ast::StorageClass::kStorage, var->Access(),
builder_.Symbols().NameFor(decl->symbol()))) {
return false;
@@ -2155,17 +2092,15 @@
out << RegisterAndSpace(var->Access() == ast::Access::kRead ? 't' : 'u',
decl->binding_point())
- << ";" << std::endl;
+ << ";";
return true;
}
-bool GeneratorImpl::EmitHandleVariable(std::ostream& out,
- const sem::Variable* var) {
- make_indent(out);
-
+bool GeneratorImpl::EmitHandleVariable(const sem::Variable* var) {
auto* decl = var->Declaration();
auto* unwrapped_type = var->Type()->UnwrapRef();
+ auto out = line();
auto name = builder_.Symbols().NameFor(decl->symbol());
auto* type = var->Type()->UnwrapRef();
@@ -2192,26 +2127,13 @@
<< bp.group->value() << ")";
}
- out << ";" << std::endl;
+ out << ";";
return true;
}
-bool GeneratorImpl::EmitPrivateVariable(std::ostream& out,
- const sem::Variable* var) {
- make_indent(out);
-
+bool GeneratorImpl::EmitPrivateVariable(const sem::Variable* var) {
auto* decl = var->Declaration();
-
- std::ostringstream constructor_out;
- if (auto* constructor = decl->constructor()) {
- if (!EmitExpression(out, constructor_out, constructor)) {
- return false;
- }
- } else {
- if (!EmitZeroValue(constructor_out, var->Type()->UnwrapRef())) {
- return false;
- }
- }
+ auto out = line();
out << "static ";
@@ -2221,27 +2143,25 @@
return false;
}
- if (constructor_out.str().length()) {
- out << " = " << constructor_out.str();
- }
-
- out << ";" << std::endl;
- return true;
-}
-
-bool GeneratorImpl::EmitWorkgroupVariable(std::ostream& out,
- const sem::Variable* var) {
- make_indent(out);
-
- auto* decl = var->Declaration();
-
- std::ostringstream constructor_out;
+ out << " = ";
if (auto* constructor = decl->constructor()) {
- if (!EmitExpression(out, constructor_out, constructor)) {
+ if (!EmitExpression(out, constructor)) {
+ return false;
+ }
+ } else {
+ if (!EmitZeroValue(out, var->Type()->UnwrapRef())) {
return false;
}
}
+ out << ";";
+ return true;
+}
+
+bool GeneratorImpl::EmitWorkgroupVariable(const sem::Variable* var) {
+ auto* decl = var->Declaration();
+ auto out = line();
+
out << "groupshared ";
auto name = builder_.Symbols().NameFor(decl->symbol());
@@ -2250,11 +2170,14 @@
return false;
}
- if (constructor_out.str().length()) {
- out << " = " << constructor_out.str();
+ if (auto* constructor = decl->constructor()) {
+ out << " = ";
+ if (!EmitExpression(out, constructor)) {
+ return false;
+ }
}
- out << ";" << std::endl;
+ out << ";";
return true;
}
@@ -2288,85 +2211,81 @@
return "";
}
-bool GeneratorImpl::EmitEntryPointFunction(std::ostream& out,
- ast::Function* func) {
- make_indent(out);
-
+bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) {
auto* func_sem = builder_.Sem().Get(func);
- if (func->pipeline_stage() == ast::PipelineStage::kCompute) {
- // Emit the workgroup_size attribute.
- auto wgsize = func_sem->workgroup_size();
- out << "[numthreads(";
- for (int i = 0; i < 3; i++) {
- if (i > 0) {
+ {
+ auto out = line();
+ if (func->pipeline_stage() == ast::PipelineStage::kCompute) {
+ // Emit the workgroup_size attribute.
+ auto wgsize = func_sem->workgroup_size();
+ out << "[numthreads(";
+ for (int i = 0; i < 3; i++) {
+ if (i > 0) {
+ out << ", ";
+ }
+
+ if (wgsize[i].overridable_const) {
+ auto* sem_const = builder_.Sem().Get(wgsize[i].overridable_const);
+ if (!sem_const->IsPipelineConstant()) {
+ TINT_ICE(Writer, builder_.Diagnostics())
+ << "expected a pipeline-overridable constant";
+ }
+ out << kSpecConstantPrefix << sem_const->ConstantId();
+ } else {
+ out << std::to_string(wgsize[i].value);
+ }
+ }
+ out << ")]" << std::endl;
+ }
+
+ out << func->return_type()->FriendlyName(builder_.Symbols());
+
+ out << " " << builder_.Symbols().NameFor(func->symbol()) << "(";
+
+ bool first = true;
+
+ // Emit entry point parameters.
+ for (auto* var : func->params()) {
+ auto* sem = builder_.Sem().Get(var);
+ auto* type = sem->Type();
+ if (!type->Is<sem::Struct>()) {
+ // ICE likely indicates that the CanonicalizeEntryPointIO transform was
+ // not run, or a builtin parameter was added after it was run.
+ TINT_ICE(Writer, diagnostics_)
+ << "Unsupported non-struct entry point parameter";
+ }
+
+ if (!first) {
out << ", ";
}
+ first = false;
- if (wgsize[i].overridable_const) {
- auto* sem_const = builder_.Sem().Get(wgsize[i].overridable_const);
- if (!sem_const->IsPipelineConstant()) {
- TINT_ICE(Writer, builder_.Diagnostics())
- << "expected a pipeline-overridable constant";
- }
- out << kSpecConstantPrefix << sem_const->ConstantId();
- } else {
- out << std::to_string(wgsize[i].value);
+ if (!EmitTypeAndName(out, type, sem->StorageClass(), sem->Access(),
+ builder_.Symbols().NameFor(var->symbol()))) {
+ return false;
}
}
- out << ")]" << std::endl;
- make_indent(out);
+
+ out << ") {";
}
- out << func->return_type()->FriendlyName(builder_.Symbols());
+ {
+ ScopedIndent si(this);
- out << " " << builder_.Symbols().NameFor(func->symbol()) << "(";
-
- bool first = true;
-
- // Emit entry point parameters.
- for (auto* var : func->params()) {
- auto* sem = builder_.Sem().Get(var);
- auto* type = sem->Type();
- if (!type->Is<sem::Struct>()) {
- // ICE likely indicates that the CanonicalizeEntryPointIO transform was
- // not run, or a builtin parameter was added after it was run.
- TINT_ICE(Writer, diagnostics_)
- << "Unsupported non-struct entry point parameter";
- }
-
- if (!first) {
- out << ", ";
- }
- first = false;
-
- if (!EmitTypeAndName(out, type, sem->StorageClass(), sem->Access(),
- builder_.Symbols().NameFor(var->symbol()))) {
+ if (!EmitStatements(func->body()->statements())) {
return false;
}
- }
- out << ") {" << std::endl;
-
- increment_indent();
-
- for (auto* s : *func->body()) {
- if (!EmitStatement(out, s)) {
- return false;
- }
- }
- auto* last_statement = func->get_last_statement();
- if (last_statement == nullptr ||
- !last_statement->Is<ast::ReturnStatement>()) {
- ast::ReturnStatement ret(ProgramID(), Source{});
- if (!EmitStatement(out, &ret)) {
- return false;
+ if (!Is<ast::ReturnStatement>(func->get_last_statement())) {
+ ast::ReturnStatement ret(ProgramID(), Source{});
+ if (!EmitStatement(&ret)) {
+ return false;
+ }
}
}
- decrement_indent();
- make_indent(out);
- out << "}" << std::endl;
+ line() << "}";
return true;
}
@@ -2465,37 +2384,35 @@
return true;
}
-bool GeneratorImpl::EmitLoop(std::ostream& out, ast::LoopStatement* stmt) {
- make_indent(out);
-
- auto emit_continuing = [this, stmt](std::ostream& o) {
+bool GeneratorImpl::EmitLoop(ast::LoopStatement* stmt) {
+ auto emit_continuing = [this, stmt]() {
if (stmt->has_continuing()) {
- make_indent(o);
- if (!EmitBlock(o, stmt->continuing())) {
+ if (!EmitBlock(stmt->continuing())) {
return false;
}
- o << std::endl;
}
return true;
};
TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
- bool ok = EmitBlockBraces(out, "while (true)", [&] {
- for (auto* s : stmt->body()->statements()) {
- if (!EmitStatement(out, s)) {
- return false;
- }
+ line() << "while (true) {";
+ {
+ ScopedIndent si(this);
+ if (!EmitStatements(stmt->body()->statements())) {
+ return false;
}
- return emit_continuing(out);
- });
- out << std::endl;
- return ok;
+ if (!emit_continuing()) {
+ return false;
+ }
+ }
+ line() << "}";
+
+ return true;
}
-bool GeneratorImpl::EmitMemberAccessor(std::ostream& pre,
- std::ostream& out,
+bool GeneratorImpl::EmitMemberAccessor(std::ostream& out,
ast::MemberAccessorExpression* expr) {
- if (!EmitExpression(pre, out, expr->structure())) {
+ if (!EmitExpression(out, expr->structure())) {
return false;
}
out << ".";
@@ -2503,80 +2420,72 @@
// Swizzles output the name directly
if (builder_.Sem().Get(expr)->Is<sem::Swizzle>()) {
out << builder_.Symbols().NameFor(expr->member()->symbol());
- } else if (!EmitExpression(pre, out, expr->member())) {
+ } else if (!EmitExpression(out, expr->member())) {
return false;
}
return true;
}
-bool GeneratorImpl::EmitReturn(std::ostream& out, ast::ReturnStatement* stmt) {
- make_indent(out);
-
+bool GeneratorImpl::EmitReturn(ast::ReturnStatement* stmt) {
if (stmt->has_value()) {
- std::ostringstream pre;
- std::ostringstream ret_out;
- if (!EmitExpression(pre, ret_out, stmt->value())) {
+ auto out = line();
+ out << "return ";
+ if (!EmitExpression(out, stmt->value())) {
return false;
}
- out << pre.str();
- out << "return " << ret_out.str();
+ out << ";";
} else {
- out << "return";
+ line() << "return;";
}
- out << ";" << std::endl;
return true;
}
-bool GeneratorImpl::EmitStatement(std::ostream& out, ast::Statement* stmt) {
+bool GeneratorImpl::EmitStatement(ast::Statement* stmt) {
if (auto* a = stmt->As<ast::AssignmentStatement>()) {
- return EmitAssign(out, a);
+ return EmitAssign(a);
}
if (auto* b = stmt->As<ast::BlockStatement>()) {
- return EmitIndentedBlockAndNewline(out, b);
+ return EmitBlock(b);
}
if (auto* b = stmt->As<ast::BreakStatement>()) {
- return EmitBreak(out, b);
+ return EmitBreak(b);
}
if (auto* c = stmt->As<ast::CallStatement>()) {
- make_indent(out);
- std::ostringstream pre;
- std::ostringstream call_out;
- if (!EmitCall(pre, call_out, c->expr())) {
- return false;
- }
- out << pre.str();
+ auto out = line();
if (!TypeOf(c->expr())->Is<sem::Void>()) {
out << "(void) ";
}
- out << call_out.str() << ";" << std::endl;
+ if (!EmitCall(out, c->expr())) {
+ return false;
+ }
+ out << ";";
return true;
}
if (auto* c = stmt->As<ast::ContinueStatement>()) {
- return EmitContinue(out, c);
+ return EmitContinue(c);
}
if (auto* d = stmt->As<ast::DiscardStatement>()) {
- return EmitDiscard(out, d);
+ return EmitDiscard(d);
}
if (stmt->As<ast::FallthroughStatement>()) {
- make_indent(out);
- out << "/* fallthrough */" << std::endl;
+ line() << "/* fallthrough */";
return true;
}
if (auto* i = stmt->As<ast::IfStatement>()) {
- return EmitIf(out, i);
+ return EmitIf(i);
}
if (auto* l = stmt->As<ast::LoopStatement>()) {
- return EmitLoop(out, l);
+ return EmitLoop(l);
}
if (auto* r = stmt->As<ast::ReturnStatement>()) {
- return EmitReturn(out, r);
+ return EmitReturn(r);
}
if (auto* s = stmt->As<ast::SwitchStatement>()) {
- return EmitSwitch(out, s);
+ return EmitSwitch(s);
}
if (auto* v = stmt->As<ast::VariableDeclStatement>()) {
- return EmitVariable(out, v->variable());
+ return EmitVariable(v->variable());
}
diagnostics_.add_error(diag::System::Writer,
@@ -2584,29 +2493,26 @@
return false;
}
-bool GeneratorImpl::EmitSwitch(std::ostream& out, ast::SwitchStatement* stmt) {
- make_indent(out);
-
- std::ostringstream pre;
- std::ostringstream cond;
- if (!EmitExpression(pre, cond, stmt->condition())) {
- return false;
+bool GeneratorImpl::EmitSwitch(ast::SwitchStatement* stmt) {
+ { // switch(expr) {
+ auto out = line();
+ out << "switch(";
+ if (!EmitExpression(out, stmt->condition())) {
+ return false;
+ }
+ out << ") {";
}
- out << pre.str();
- out << "switch(" << cond.str() << ") {" << std::endl;
-
- increment_indent();
-
- for (auto* s : stmt->body()) {
- if (!EmitCase(out, s)) {
- return false;
+ {
+ ScopedIndent si(this);
+ for (auto* s : stmt->body()) {
+ if (!EmitCase(s)) {
+ return false;
+ }
}
}
- decrement_indent();
- make_indent(out);
- out << "}" << std::endl;
+ line() << "}";
return true;
}
@@ -2806,7 +2712,7 @@
return true;
}
-bool GeneratorImpl::EmitStructType(std::ostream& out, const sem::Struct* str) {
+bool GeneratorImpl::EmitStructType(const sem::Struct* str) {
auto storage_class_uses = str->StorageClassUsage();
if (storage_class_uses.size() ==
(storage_class_uses.count(ast::StorageClass::kStorage) +
@@ -2821,72 +2727,69 @@
}
auto struct_name = builder_.Symbols().NameFor(str->Declaration()->name());
- out << "struct " << struct_name << " {" << std::endl;
+ line() << "struct " << struct_name << " {";
+ {
+ ScopedIndent si(this);
+ for (auto* mem : str->Members()) {
+ auto name = builder_.Symbols().NameFor(mem->Declaration()->symbol());
- increment_indent();
- for (auto* mem : str->Members()) {
- make_indent(out);
+ auto* ty = mem->Type();
- auto name = builder_.Symbols().NameFor(mem->Declaration()->symbol());
-
- auto* ty = mem->Type();
-
- if (!EmitTypeAndName(out, ty, ast::StorageClass::kNone,
- ast::Access::kReadWrite, name)) {
- return false;
- }
-
- for (auto* deco : mem->Declaration()->decorations()) {
- if (auto* location = deco->As<ast::LocationDecoration>()) {
- auto& pipeline_stage_uses = str->PipelineStageUses();
- if (pipeline_stage_uses.size() != 1) {
- TINT_ICE(Writer, diagnostics_)
- << "invalid entry point IO struct uses";
- }
-
- if (pipeline_stage_uses.count(sem::PipelineStageUsage::kVertexInput)) {
- out << " : TEXCOORD" + std::to_string(location->value());
- } else if (pipeline_stage_uses.count(
- sem::PipelineStageUsage::kVertexOutput)) {
- out << " : TEXCOORD" + std::to_string(location->value());
- } else if (pipeline_stage_uses.count(
- sem::PipelineStageUsage::kFragmentInput)) {
- out << " : TEXCOORD" + std::to_string(location->value());
- } else if (pipeline_stage_uses.count(
- sem::PipelineStageUsage::kFragmentOutput)) {
- out << " : SV_Target" + std::to_string(location->value());
- } else {
- TINT_ICE(Writer, diagnostics_)
- << "invalid use of location decoration";
- }
- } else if (auto* builtin = deco->As<ast::BuiltinDecoration>()) {
- auto attr = builtin_to_attribute(builtin->value());
- if (attr.empty()) {
- diagnostics_.add_error(diag::System::Writer, "unsupported builtin");
- return false;
- }
- out << " : " << attr;
+ auto out = line();
+ if (!EmitTypeAndName(out, ty, ast::StorageClass::kNone,
+ ast::Access::kReadWrite, name)) {
+ return false;
}
- }
- out << ";" << std::endl;
+ for (auto* deco : mem->Declaration()->decorations()) {
+ if (auto* location = deco->As<ast::LocationDecoration>()) {
+ auto& pipeline_stage_uses = str->PipelineStageUses();
+ if (pipeline_stage_uses.size() != 1) {
+ TINT_ICE(Writer, diagnostics_)
+ << "invalid entry point IO struct uses";
+ }
+
+ if (pipeline_stage_uses.count(
+ sem::PipelineStageUsage::kVertexInput)) {
+ out << " : TEXCOORD" + std::to_string(location->value());
+ } else if (pipeline_stage_uses.count(
+ sem::PipelineStageUsage::kVertexOutput)) {
+ out << " : TEXCOORD" + std::to_string(location->value());
+ } else if (pipeline_stage_uses.count(
+ sem::PipelineStageUsage::kFragmentInput)) {
+ out << " : TEXCOORD" + std::to_string(location->value());
+ } else if (pipeline_stage_uses.count(
+ sem::PipelineStageUsage::kFragmentOutput)) {
+ out << " : SV_Target" + std::to_string(location->value());
+ } else {
+ TINT_ICE(Writer, diagnostics_)
+ << "invalid use of location decoration";
+ }
+ } else if (auto* builtin = deco->As<ast::BuiltinDecoration>()) {
+ auto attr = builtin_to_attribute(builtin->value());
+ if (attr.empty()) {
+ diagnostics_.add_error(diag::System::Writer, "unsupported builtin");
+ return false;
+ }
+ out << " : " << attr;
+ }
+ }
+
+ out << ";";
+ }
}
- decrement_indent();
- make_indent(out);
-
- out << "};" << std::endl;
+ line() << "};";
return true;
}
-bool GeneratorImpl::EmitUnaryOp(std::ostream& pre,
- std::ostream& out,
+bool GeneratorImpl::EmitUnaryOp(std::ostream& out,
ast::UnaryOpExpression* expr) {
switch (expr->op()) {
case ast::UnaryOp::kIndirection:
case ast::UnaryOp::kAddressOf:
- return EmitExpression(pre, out, expr->expr());
+ return EmitExpression(out, expr->expr());
case ast::UnaryOp::kComplement:
out << "~";
break;
@@ -2899,7 +2802,7 @@
}
out << "(";
- if (!EmitExpression(pre, out, expr->expr())) {
+ if (!EmitExpression(out, expr->expr())) {
return false;
}
@@ -2908,9 +2811,7 @@
return true;
}
-bool GeneratorImpl::EmitVariable(std::ostream& out, ast::Variable* var) {
- make_indent(out);
-
+bool GeneratorImpl::EmitVariable(ast::Variable* var) {
auto* sem = builder_.Sem().Get(var);
auto* type = sem->Type()->UnwrapRef();
@@ -2921,21 +2822,7 @@
return false;
}
- std::ostringstream constructor_out;
- constructor_out << " = ";
-
- if (var->constructor()) {
- std::ostringstream pre;
- if (!EmitExpression(pre, constructor_out, var->constructor())) {
- return false;
- }
- out << pre.str();
- } else {
- if (!EmitZeroValue(constructor_out, type)) {
- return false;
- }
- }
-
+ auto out = line();
if (var->is_const()) {
out << "const ";
}
@@ -2943,15 +2830,24 @@
builder_.Symbols().NameFor(var->symbol()))) {
return false;
}
- out << constructor_out.str() << ";" << std::endl;
+
+ out << " = ";
+
+ if (var->constructor()) {
+ if (!EmitExpression(out, var->constructor())) {
+ return false;
+ }
+ } else {
+ if (!EmitZeroValue(out, type)) {
+ return false;
+ }
+ }
+ out << ";";
return true;
}
-bool GeneratorImpl::EmitProgramConstVariable(std::ostream& out,
- const ast::Variable* var) {
- make_indent(out);
-
+bool GeneratorImpl::EmitProgramConstVariable(const ast::Variable* var) {
for (auto* d : var->decorations()) {
if (!d->Is<ast::OverrideDecoration>()) {
diagnostics_.add_error(diag::System::Writer,
@@ -2964,48 +2860,45 @@
return false;
}
- std::ostringstream constructor_out;
- if (var->constructor() != nullptr) {
- std::ostringstream pre;
- if (!EmitExpression(pre, constructor_out, var->constructor())) {
- return false;
- }
- out << pre.str();
- }
-
auto* sem = builder_.Sem().Get(var);
auto* type = sem->Type();
if (sem->IsPipelineConstant()) {
auto const_id = sem->ConstantId();
- out << "#ifndef " << kSpecConstantPrefix << const_id << std::endl;
+ line() << "#ifndef " << kSpecConstantPrefix << const_id;
if (var->constructor() != nullptr) {
- out << "#define " << kSpecConstantPrefix << const_id << " "
- << constructor_out.str() << std::endl;
+ auto out = line();
+ out << "#define " << kSpecConstantPrefix << const_id << " ";
+ if (!EmitExpression(out, var->constructor())) {
+ return false;
+ }
} else {
- out << "#error spec constant required for constant id " << const_id
- << std::endl;
+ line() << "#error spec constant required for constant id " << const_id;
}
- out << "#endif" << std::endl;
- out << "static const ";
- if (!EmitTypeAndName(out, type, sem->StorageClass(), sem->Access(),
- builder_.Symbols().NameFor(var->symbol()))) {
- return false;
+ line() << "#endif";
+ {
+ auto out = line();
+ out << "static const ";
+ if (!EmitTypeAndName(out, type, sem->StorageClass(), sem->Access(),
+ builder_.Symbols().NameFor(var->symbol()))) {
+ return false;
+ }
+ out << " = " << kSpecConstantPrefix << const_id << ";";
}
- out << " = " << kSpecConstantPrefix << const_id << ";" << std::endl;
} else {
+ auto out = line();
out << "static const ";
if (!EmitTypeAndName(out, type, sem->StorageClass(), sem->Access(),
builder_.Symbols().NameFor(var->symbol()))) {
return false;
}
-
- if (var->constructor() != nullptr) {
- out << " = " << constructor_out.str();
+ out << " = ";
+ if (!EmitExpression(out, var->constructor())) {
+ return false;
}
- out << ";" << std::endl;
+ out << ";";
}
return true;
@@ -3026,23 +2919,6 @@
return "";
}
-template <typename F>
-bool GeneratorImpl::EmitBlockBraces(std::ostream& out,
- const std::string& prefix,
- F&& cb) {
- out << prefix << (prefix.empty() ? "{" : " {") << std::endl;
- increment_indent();
-
- if (!cb()) {
- return false;
- }
-
- decrement_indent();
- make_indent(out);
- out << "}";
- return true;
-}
-
} // namespace hlsl
} // namespace writer
} // namespace tint
diff --git a/src/writer/hlsl/generator_impl.h b/src/writer/hlsl/generator_impl.h
index e46a6bd..fd3c6bd 100644
--- a/src/writer/hlsl/generator_impl.h
+++ b/src/writer/hlsl/generator_impl.h
@@ -54,320 +54,246 @@
explicit GeneratorImpl(const Program* program);
~GeneratorImpl();
- /// @param out the output stream
/// @returns true on successful generation; false otherwise
- bool Generate(std::ostream& out);
+ bool Generate();
/// Handles an array accessor expression
- /// @param pre the preamble for the expression stream
/// @param out the output of the expression stream
/// @param expr the expression to emit
/// @returns true if the array accessor was emitted
- bool EmitArrayAccessor(std::ostream& pre,
- std::ostream& out,
- ast::ArrayAccessorExpression* expr);
+ bool EmitArrayAccessor(std::ostream& out, ast::ArrayAccessorExpression* expr);
/// Handles an assignment statement
- /// @param out the output stream
/// @param stmt the statement to emit
/// @returns true if the statement was emitted successfully
- bool EmitAssign(std::ostream& out, ast::AssignmentStatement* stmt);
+ bool EmitAssign(ast::AssignmentStatement* stmt);
/// Handles generating a binary expression
- /// @param pre the preamble for the expression stream
/// @param out the output of the expression stream
/// @param expr the binary expression
/// @returns true if the expression was emitted, false otherwise
- bool EmitBinary(std::ostream& pre,
- std::ostream& out,
- ast::BinaryExpression* expr);
+ bool EmitBinary(std::ostream& out, ast::BinaryExpression* expr);
/// Handles generating a bitcast expression
- /// @param pre the preamble for the expression stream
/// @param out the output of the expression stream
/// @param expr the as expression
/// @returns true if the bitcast was emitted
- bool EmitBitcast(std::ostream& pre,
- std::ostream& out,
- ast::BitcastExpression* expr);
+ bool EmitBitcast(std::ostream& out, ast::BitcastExpression* expr);
+ /// Emits a list of statements
+ /// @param stmts the statement list
+ /// @returns true if the statements were emitted successfully
+ bool EmitStatements(const ast::StatementList& stmts);
+ /// Emits a list of statements with an indentation
+ /// @param stmts the statement list
+ /// @returns true if the statements were emitted successfully
+ bool EmitStatementsWithIndent(const ast::StatementList& stmts);
/// Handles a block statement
- /// @param out the output stream
/// @param stmt the statement to emit
/// @returns true if the statement was emitted successfully
- bool EmitBlock(std::ostream& out, const ast::BlockStatement* stmt);
- /// Handles a block statement with a newline at the end
- /// @param out the output stream
- /// @param stmt the statement to emit
- /// @returns true if the statement was emitted successfully
- bool EmitIndentedBlockAndNewline(std::ostream& out,
- ast::BlockStatement* stmt);
- /// Handles a block statement with a newline at the end
- /// @param out the output stream
- /// @param stmt the statement to emit
- /// @returns true if the statement was emitted successfully
- bool EmitBlockAndNewline(std::ostream& out, const ast::BlockStatement* stmt);
+ bool EmitBlock(const ast::BlockStatement* stmt);
/// Handles a break statement
- /// @param out the output stream
/// @param stmt the statement to emit
/// @returns true if the statement was emitted successfully
- bool EmitBreak(std::ostream& out, ast::BreakStatement* stmt);
+ bool EmitBreak(ast::BreakStatement* stmt);
/// Handles generating a call expression
- /// @param pre the preamble for the expression stream
/// @param out the output of the expression stream
/// @param expr the call expression
/// @returns true if the call expression is emitted
- bool EmitCall(std::ostream& pre,
- std::ostream& out,
- ast::CallExpression* expr);
+ bool EmitCall(std::ostream& out, ast::CallExpression* expr);
/// Handles generating a call expression to a
/// transform::DecomposeMemoryAccess::Intrinsic for a uniform buffer
- /// @param pre the preamble for the expression stream
/// @param out the output of the expression stream
/// @param expr the call expression
/// @param intrinsic the transform::DecomposeMemoryAccess::Intrinsic
/// @returns true if the call expression is emitted
bool EmitUniformBufferAccess(
- std::ostream& pre,
std::ostream& out,
ast::CallExpression* expr,
const transform::DecomposeMemoryAccess::Intrinsic* intrinsic);
/// Handles generating a call expression to a
/// transform::DecomposeMemoryAccess::Intrinsic for a storage buffer
- /// @param pre the preamble for the expression stream
/// @param out the output of the expression stream
/// @param expr the call expression
/// @param intrinsic the transform::DecomposeMemoryAccess::Intrinsic
/// @returns true if the call expression is emitted
bool EmitStorageBufferAccess(
- std::ostream& pre,
std::ostream& out,
ast::CallExpression* expr,
const transform::DecomposeMemoryAccess::Intrinsic* intrinsic);
/// Handles generating a barrier intrinsic call
- /// @param pre the preamble for the expression stream
/// @param out the output of the expression stream
/// @param intrinsic the semantic information for the barrier intrinsic
/// @returns true if the call expression is emitted
- bool EmitBarrierCall(std::ostream& pre,
- std::ostream& out,
- const sem::Intrinsic* intrinsic);
+ bool EmitBarrierCall(std::ostream& out, const sem::Intrinsic* intrinsic);
/// Handles generating an atomic intrinsic call for a storage buffer variable
- /// @param pre the preamble for the expression stream
/// @param out the output of the expression stream
/// @param expr the call expression
/// @param op the atomic op
/// @returns true if the call expression is emitted
bool EmitStorageAtomicCall(
- std::ostream& pre,
std::ostream& out,
ast::CallExpression* expr,
transform::DecomposeMemoryAccess::Intrinsic::Op op);
/// Handles generating an atomic intrinsic call for a workgroup variable
- /// @param pre the preamble for the expression stream
/// @param out the output of the expression stream
/// @param expr the call expression
/// @param intrinsic the semantic information for the atomic intrinsic
/// @returns true if the call expression is emitted
- bool EmitWorkgroupAtomicCall(std::ostream& pre,
- std::ostream& out,
+ bool EmitWorkgroupAtomicCall(std::ostream& out,
ast::CallExpression* expr,
const sem::Intrinsic* intrinsic);
/// Handles generating a call to a texture function (`textureSample`,
/// `textureSampleGrad`, etc)
- /// @param pre the preamble for the expression stream
/// @param out the output of the expression stream
/// @param expr the call expression
/// @param intrinsic the semantic information for the texture intrinsic
/// @returns true if the call expression is emitted
- bool EmitTextureCall(std::ostream& pre,
- std::ostream& out,
+ bool EmitTextureCall(std::ostream& out,
ast::CallExpression* expr,
const sem::Intrinsic* intrinsic);
/// Handles generating a call to the `select()` intrinsic
- /// @param pre the preamble of the expression stream
/// @param out the output of the expression stream
/// @param expr the call expression
/// @returns true if the call expression is emitted
- bool EmitSelectCall(std::ostream& pre,
- std::ostream& out,
- ast::CallExpression* expr);
+ bool EmitSelectCall(std::ostream& out, ast::CallExpression* expr);
/// Handles generating a call to the `frexp()` intrinsic
- /// @param pre the preamble of the expression stream
/// @param out the output of the expression stream
/// @param expr the call expression
/// @param intrinsic the semantic information for the intrinsic
/// @returns true if the call expression is emitted
- bool EmitFrexpCall(std::ostream& pre,
- std::ostream& out,
+ bool EmitFrexpCall(std::ostream& out,
ast::CallExpression* expr,
const sem::Intrinsic* intrinsic);
/// Handles generating a call to the `isNormal()` intrinsic
- /// @param pre the preamble of the expression stream
/// @param out the output of the expression stream
/// @param expr the call expression
/// @param intrinsic the semantic information for the intrinsic
/// @returns true if the call expression is emitted
- bool EmitIsNormalCall(std::ostream& pre,
- std::ostream& out,
+ bool EmitIsNormalCall(std::ostream& out,
ast::CallExpression* expr,
const sem::Intrinsic* intrinsic);
/// Handles generating a call to data packing intrinsic
- /// @param pre the preamble of the expression stream
/// @param out the output of the expression stream
/// @param expr the call expression
/// @param intrinsic the semantic information for the texture intrinsic
/// @returns true if the call expression is emitted
- bool EmitDataPackingCall(std::ostream& pre,
- std::ostream& out,
+ bool EmitDataPackingCall(std::ostream& out,
ast::CallExpression* expr,
const sem::Intrinsic* intrinsic);
/// Handles generating a call to data unpacking intrinsic
- /// @param pre the preamble of the expression stream
/// @param out the output of the expression stream
/// @param expr the call expression
/// @param intrinsic the semantic information for the texture intrinsic
/// @returns true if the call expression is emitted
- bool EmitDataUnpackingCall(std::ostream& pre,
- std::ostream& out,
+ bool EmitDataUnpackingCall(std::ostream& out,
ast::CallExpression* expr,
const sem::Intrinsic* intrinsic);
/// Handles a case statement
- /// @param out the output stream
/// @param stmt the statement
- /// @returns true if the statment was emitted successfully
- bool EmitCase(std::ostream& out, ast::CaseStatement* stmt);
+ /// @returns true if the statement was emitted successfully
+ bool EmitCase(ast::CaseStatement* stmt);
/// Handles generating constructor expressions
- /// @param pre the preamble for the expression stream
/// @param out the output of the expression stream
/// @param expr the constructor expression
/// @returns true if the expression was emitted
- bool EmitConstructor(std::ostream& pre,
- std::ostream& out,
- ast::ConstructorExpression* expr);
+ bool EmitConstructor(std::ostream& out, ast::ConstructorExpression* expr);
/// Handles generating a discard statement
- /// @param out the output stream
/// @param stmt the discard statement
/// @returns true if the statement was successfully emitted
- bool EmitDiscard(std::ostream& out, ast::DiscardStatement* stmt);
+ bool EmitDiscard(ast::DiscardStatement* stmt);
/// Handles generating a scalar constructor
- /// @param pre the preamble for the expression stream
/// @param out the output of the expression stream
/// @param expr the scalar constructor expression
/// @returns true if the scalar constructor is emitted
- bool EmitScalarConstructor(std::ostream& pre,
- std::ostream& out,
+ bool EmitScalarConstructor(std::ostream& out,
ast::ScalarConstructorExpression* expr);
/// Handles emitting a type constructor
- /// @param pre the preamble for the expression stream
/// @param out the output of the expression stream
/// @param expr the type constructor expression
/// @returns true if the constructor is emitted
- bool EmitTypeConstructor(std::ostream& pre,
- std::ostream& out,
+ bool EmitTypeConstructor(std::ostream& out,
ast::TypeConstructorExpression* expr);
/// Handles a continue statement
- /// @param out the output stream
/// @param stmt the statement to emit
/// @returns true if the statement was emitted successfully
- bool EmitContinue(std::ostream& out, ast::ContinueStatement* stmt);
+ bool EmitContinue(ast::ContinueStatement* stmt);
/// Handles generate an Expression
- /// @param pre the preamble for the expression stream
/// @param out the output of the expression stream
/// @param expr the expression
/// @returns true if the expression was emitted
- bool EmitExpression(std::ostream& pre,
- std::ostream& out,
- ast::Expression* expr);
+ bool EmitExpression(std::ostream& out, ast::Expression* expr);
/// Handles generating a function
- /// @param out the output stream
/// @param func the function to generate
/// @returns true if the function was emitted
- bool EmitFunction(std::ostream& out, ast::Function* func);
+ bool EmitFunction(ast::Function* func);
/// Handles emitting a global variable
- /// @param out the output stream
/// @param global the global variable
/// @returns true on success
- bool EmitGlobalVariable(std::ostream& out, ast::Variable* global);
+ bool EmitGlobalVariable(ast::Variable* global);
/// Handles emitting a global variable with the uniform storage class
- /// @param out the output stream
/// @param var the global variable
/// @returns true on success
- bool EmitUniformVariable(std::ostream& out, const sem::Variable* var);
+ bool EmitUniformVariable(const sem::Variable* var);
/// Handles emitting a global variable with the storage storage class
- /// @param out the output stream
/// @param var the global variable
/// @returns true on success
- bool EmitStorageVariable(std::ostream& out, const sem::Variable* var);
+ bool EmitStorageVariable(const sem::Variable* var);
/// Handles emitting a global variable with the handle storage class
- /// @param out the output stream
/// @param var the global variable
/// @returns true on success
- bool EmitHandleVariable(std::ostream& out, const sem::Variable* var);
+ bool EmitHandleVariable(const sem::Variable* var);
/// Handles emitting a global variable with the private storage class
- /// @param out the output stream
/// @param var the global variable
/// @returns true on success
- bool EmitPrivateVariable(std::ostream& out, const sem::Variable* var);
+ bool EmitPrivateVariable(const sem::Variable* var);
/// Handles emitting a global variable with the workgroup storage class
- /// @param out the output stream
/// @param var the global variable
/// @returns true on success
- bool EmitWorkgroupVariable(std::ostream& out, const sem::Variable* var);
+ bool EmitWorkgroupVariable(const sem::Variable* var);
/// Handles emitting the entry point function
- /// @param out the output stream
/// @param func the entry point
/// @returns true if the entry point function was emitted
- bool EmitEntryPointFunction(std::ostream& out, ast::Function* func);
+ bool EmitEntryPointFunction(ast::Function* func);
/// Handles an if statement
- /// @param out the output stream
/// @param stmt the statement to emit
/// @returns true if the statement was successfully emitted
- bool EmitIf(std::ostream& out, ast::IfStatement* stmt);
+ bool EmitIf(ast::IfStatement* stmt);
/// Handles a literal
/// @param out the output stream
/// @param lit the literal to emit
/// @returns true if the literal was successfully emitted
bool EmitLiteral(std::ostream& out, ast::Literal* lit);
/// Handles a loop statement
- /// @param out the output stream
/// @param stmt the statement to emit
/// @returns true if the statement was emitted
- bool EmitLoop(std::ostream& out, ast::LoopStatement* stmt);
+ bool EmitLoop(ast::LoopStatement* stmt);
/// Handles generating an identifier expression
- /// @param pre the preamble for the expression stream
/// @param out the output of the expression stream
/// @param expr the identifier expression
/// @returns true if the identifeir was emitted
- bool EmitIdentifier(std::ostream& pre,
- std::ostream& out,
- ast::IdentifierExpression* expr);
+ bool EmitIdentifier(std::ostream& out, ast::IdentifierExpression* expr);
/// Handles a member accessor expression
- /// @param pre the preamble for the expression stream
/// @param out the output of the expression stream
/// @param expr the member accessor expression
/// @returns true if the member accessor was emitted
- bool EmitMemberAccessor(std::ostream& pre,
- std::ostream& out,
+ bool EmitMemberAccessor(std::ostream& out,
ast::MemberAccessorExpression* expr);
/// Handles return statements
- /// @param out the output stream
/// @param stmt the statement to emit
/// @returns true if the statement was successfully emitted
- bool EmitReturn(std::ostream& out, ast::ReturnStatement* stmt);
+ bool EmitReturn(ast::ReturnStatement* stmt);
/// Handles statement
- /// @param out the output stream
/// @param stmt the statement to emit
/// @returns true if the statement was emitted
- bool EmitStatement(std::ostream& out, ast::Statement* stmt);
+ bool EmitStatement(ast::Statement* stmt);
/// Handles generating a switch statement
- /// @param out the output stream
/// @param stmt the statement to emit
/// @returns true if the statement was emitted
- bool EmitSwitch(std::ostream& out, ast::SwitchStatement* stmt);
+ bool EmitSwitch(ast::SwitchStatement* stmt);
/// Handles generating type
/// @param out the output stream
/// @param type the type to generate
@@ -396,33 +322,27 @@
ast::Access access,
const std::string& name);
/// Handles generating a structure declaration
- /// @param out the output stream
/// @param ty the struct to generate
/// @returns true if the struct is emitted
- bool EmitStructType(std::ostream& out, const sem::Struct* ty);
+ bool EmitStructType(const sem::Struct* ty);
/// Handles a unary op expression
- /// @param pre the preamble for the expression stream
/// @param out the output of the expression stream
/// @param expr the expression to emit
/// @returns true if the expression was emitted
- bool EmitUnaryOp(std::ostream& pre,
- std::ostream& out,
- ast::UnaryOpExpression* expr);
+ bool EmitUnaryOp(std::ostream& out, ast::UnaryOpExpression* expr);
/// Emits the zero value for the given type
/// @param out the output stream
/// @param type the type to emit the value for
/// @returns true if the zero value was successfully emitted.
bool EmitZeroValue(std::ostream& out, const sem::Type* type);
/// Handles generating a variable
- /// @param out the output stream
/// @param var the variable to generate
/// @returns true if the variable was emitted
- bool EmitVariable(std::ostream& out, ast::Variable* var);
+ bool EmitVariable(ast::Variable* var);
/// Handles generating a program scope constant variable
- /// @param out the output stream
/// @param var the variable to emit
/// @returns true if the variable was emitted
- bool EmitProgramConstVariable(std::ostream& out, const ast::Variable* var);
+ bool EmitProgramConstVariable(const ast::Variable* var);
/// Handles generating a builtin method name
/// @param intrinsic the semantic info for the intrinsic
@@ -466,29 +386,8 @@
return builder_.TypeOf(type_decl);
}
- /// Emits `prefix`, followed by an opening brace `{`, then calls `cb` to emit
- /// the block body, then finally emits the closing brace `}`.
- /// @param out the output stream
- /// @param prefix the string to emit before the opening brace
- /// @param cb a function or function-like object with the signature `bool()`
- /// that emits the block body.
- /// @returns the return value of `cb`.
- template <typename F>
- bool EmitBlockBraces(std::ostream& out, const std::string& prefix, F&& cb);
-
- /// Emits an opening brace `{`, then calls `cb` to emit the block body, then
- /// finally emits the closing brace `}`.
- /// @param out the output stream
- /// @param cb a function or function-like object with the signature `bool()`
- /// that emits the block body.
- /// @returns the return value of `cb`.
- template <typename F>
- bool EmitBlockBraces(std::ostream& out, F&& cb) {
- return EmitBlockBraces(out, "", std::forward<F>(cb));
- }
-
ProgramBuilder builder_;
- std::function<bool(std::ostream& out)> emit_continuing_;
+ std::function<bool()> emit_continuing_;
std::unordered_map<const sem::Struct*, std::string> structure_builders_;
};
diff --git a/src/writer/hlsl/generator_impl_array_accessor_test.cc b/src/writer/hlsl/generator_impl_array_accessor_test.cc
index 53c42ad..b2fb021 100644
--- a/src/writer/hlsl/generator_impl_array_accessor_test.cc
+++ b/src/writer/hlsl/generator_impl_array_accessor_test.cc
@@ -28,8 +28,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitExpression(pre, out, expr)) << gen.error();
- EXPECT_EQ(result(), "ary[5]");
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+ EXPECT_EQ(out.str(), "ary[5]");
}
} // namespace
diff --git a/src/writer/hlsl/generator_impl_assign_test.cc b/src/writer/hlsl/generator_impl_assign_test.cc
index 54e957a..f5a0711 100644
--- a/src/writer/hlsl/generator_impl_assign_test.cc
+++ b/src/writer/hlsl/generator_impl_assign_test.cc
@@ -31,8 +31,8 @@
gen.increment_indent();
- ASSERT_TRUE(gen.EmitStatement(out, assign)) << gen.error();
- EXPECT_EQ(result(), " lhs = rhs;\n");
+ ASSERT_TRUE(gen.EmitStatement(assign)) << gen.error();
+ EXPECT_EQ(gen.result(), " lhs = rhs;\n");
}
} // namespace
diff --git a/src/writer/hlsl/generator_impl_binary_test.cc b/src/writer/hlsl/generator_impl_binary_test.cc
index 743bb2b..435ce1d 100644
--- a/src/writer/hlsl/generator_impl_binary_test.cc
+++ b/src/writer/hlsl/generator_impl_binary_test.cc
@@ -56,8 +56,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitExpression(pre, out, expr)) << gen.error();
- EXPECT_EQ(result(), params.result);
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+ EXPECT_EQ(out.str(), params.result);
}
TEST_P(HlslBinaryTest, Emit_u32) {
auto params = GetParam();
@@ -74,8 +75,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitExpression(pre, out, expr)) << gen.error();
- EXPECT_EQ(result(), params.result);
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+ EXPECT_EQ(out.str(), params.result);
}
TEST_P(HlslBinaryTest, Emit_i32) {
auto params = GetParam();
@@ -98,8 +100,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitExpression(pre, out, expr)) << gen.error();
- EXPECT_EQ(result(), params.result);
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+ EXPECT_EQ(out.str(), params.result);
}
INSTANTIATE_TEST_SUITE_P(
HlslGeneratorImplTest,
@@ -133,8 +136,9 @@
GeneratorImpl& gen = Build();
- EXPECT_TRUE(gen.EmitExpression(pre, out, expr)) << gen.error();
- EXPECT_EQ(result(),
+ std::stringstream out;
+ EXPECT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+ EXPECT_EQ(out.str(),
"(float3(1.0f, 1.0f, 1.0f) * "
"1.0f)");
}
@@ -150,8 +154,9 @@
GeneratorImpl& gen = Build();
- EXPECT_TRUE(gen.EmitExpression(pre, out, expr)) << gen.error();
- EXPECT_EQ(result(),
+ std::stringstream out;
+ EXPECT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+ EXPECT_EQ(out.str(),
"(1.0f * float3(1.0f, 1.0f, "
"1.0f))");
}
@@ -167,8 +172,9 @@
GeneratorImpl& gen = Build();
- EXPECT_TRUE(gen.EmitExpression(pre, out, expr)) << gen.error();
- EXPECT_EQ(result(), "(mat * 1.0f)");
+ std::stringstream out;
+ EXPECT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+ EXPECT_EQ(out.str(), "(mat * 1.0f)");
}
TEST_F(HlslGeneratorImplTest_Binary, Multiply_ScalarMatrix) {
@@ -182,8 +188,9 @@
GeneratorImpl& gen = Build();
- EXPECT_TRUE(gen.EmitExpression(pre, out, expr)) << gen.error();
- EXPECT_EQ(result(), "(1.0f * mat)");
+ std::stringstream out;
+ EXPECT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+ EXPECT_EQ(out.str(), "(1.0f * mat)");
}
TEST_F(HlslGeneratorImplTest_Binary, Multiply_MatrixVector) {
@@ -197,8 +204,9 @@
GeneratorImpl& gen = Build();
- EXPECT_TRUE(gen.EmitExpression(pre, out, expr)) << gen.error();
- EXPECT_EQ(result(), "mul(float3(1.0f, 1.0f, 1.0f), mat)");
+ std::stringstream out;
+ EXPECT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+ EXPECT_EQ(out.str(), "mul(float3(1.0f, 1.0f, 1.0f), mat)");
}
TEST_F(HlslGeneratorImplTest_Binary, Multiply_VectorMatrix) {
@@ -212,8 +220,9 @@
GeneratorImpl& gen = Build();
- EXPECT_TRUE(gen.EmitExpression(pre, out, expr)) << gen.error();
- EXPECT_EQ(result(), "mul(mat, float3(1.0f, 1.0f, 1.0f))");
+ std::stringstream out;
+ EXPECT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+ EXPECT_EQ(out.str(), "mul(mat, float3(1.0f, 1.0f, 1.0f))");
}
TEST_F(HlslGeneratorImplTest_Binary, Multiply_MatrixMatrix) {
@@ -226,8 +235,9 @@
GeneratorImpl& gen = Build();
- EXPECT_TRUE(gen.EmitExpression(pre, out, expr)) << gen.error();
- EXPECT_EQ(result(), "mul(rhs, lhs)");
+ std::stringstream out;
+ EXPECT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+ EXPECT_EQ(out.str(), "mul(rhs, lhs)");
}
TEST_F(HlslGeneratorImplTest_Binary, Logical_And) {
@@ -240,9 +250,10 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitExpression(pre, out, expr)) << gen.error();
- EXPECT_EQ(result(), "(tint_tmp)");
- EXPECT_EQ(pre_result(), R"(bool tint_tmp = a;
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+ EXPECT_EQ(out.str(), "(tint_tmp)");
+ EXPECT_EQ(gen.result(), R"(bool tint_tmp = a;
if (tint_tmp) {
tint_tmp = b;
}
@@ -266,19 +277,20 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitExpression(pre, out, expr)) << gen.error();
- EXPECT_EQ(result(), "(tint_tmp_1)");
- EXPECT_EQ(pre_result(), R"(bool tint_tmp = a;
-if (tint_tmp) {
- tint_tmp = b;
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+ EXPECT_EQ(out.str(), "(tint_tmp)");
+ EXPECT_EQ(gen.result(), R"(bool tint_tmp_1 = a;
+if (tint_tmp_1) {
+ tint_tmp_1 = b;
}
-bool tint_tmp_1 = (tint_tmp);
-if (!tint_tmp_1) {
+bool tint_tmp = (tint_tmp_1);
+if (!tint_tmp) {
bool tint_tmp_2 = c;
if (!tint_tmp_2) {
tint_tmp_2 = d;
}
- tint_tmp_1 = (tint_tmp_2);
+ tint_tmp = (tint_tmp_2);
}
)");
}
@@ -293,9 +305,10 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitExpression(pre, out, expr)) << gen.error();
- EXPECT_EQ(result(), "(tint_tmp)");
- EXPECT_EQ(pre_result(), R"(bool tint_tmp = a;
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+ EXPECT_EQ(out.str(), "(tint_tmp)");
+ EXPECT_EQ(gen.result(), R"(bool tint_tmp = a;
if (!tint_tmp) {
tint_tmp = b;
}
@@ -338,8 +351,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitStatement(out, expr)) << gen.error();
- EXPECT_EQ(result(), R"(bool tint_tmp = a;
+ ASSERT_TRUE(gen.EmitStatement(expr)) << gen.error();
+ EXPECT_EQ(gen.result(), R"(bool tint_tmp = a;
if (tint_tmp) {
tint_tmp = b;
}
@@ -375,16 +388,16 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitStatement(out, expr)) << gen.error();
- EXPECT_EQ(result(), R"(bool tint_tmp = a;
-if (tint_tmp) {
- tint_tmp = b;
+ ASSERT_TRUE(gen.EmitStatement(expr)) << gen.error();
+ EXPECT_EQ(gen.result(), R"(bool tint_tmp_1 = a;
+if (tint_tmp_1) {
+ tint_tmp_1 = b;
}
-bool tint_tmp_1 = (tint_tmp);
-if (!tint_tmp_1) {
- tint_tmp_1 = c;
+bool tint_tmp = (tint_tmp_1);
+if (!tint_tmp) {
+ tint_tmp = c;
}
-return (tint_tmp_1);
+return (tint_tmp);
)");
}
@@ -406,16 +419,16 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitStatement(out, expr)) << gen.error();
- EXPECT_EQ(result(), R"(bool tint_tmp = b;
-if (!tint_tmp) {
- tint_tmp = c;
+ ASSERT_TRUE(gen.EmitStatement(expr)) << gen.error();
+ EXPECT_EQ(gen.result(), R"(bool tint_tmp_1 = b;
+if (!tint_tmp_1) {
+ tint_tmp_1 = c;
}
-bool tint_tmp_1 = (tint_tmp);
-if (tint_tmp_1) {
- tint_tmp_1 = d;
+bool tint_tmp = (tint_tmp_1);
+if (tint_tmp) {
+ tint_tmp = d;
}
-a = (tint_tmp_1);
+a = (tint_tmp);
)");
}
@@ -438,16 +451,16 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitStatement(out, decl)) << gen.error();
- EXPECT_EQ(result(), R"(bool tint_tmp = b;
-if (tint_tmp) {
- tint_tmp = c;
+ ASSERT_TRUE(gen.EmitStatement(decl)) << gen.error();
+ EXPECT_EQ(gen.result(), R"(bool tint_tmp_1 = b;
+if (tint_tmp_1) {
+ tint_tmp_1 = c;
}
-bool tint_tmp_1 = (tint_tmp);
-if (!tint_tmp_1) {
- tint_tmp_1 = d;
+bool tint_tmp = (tint_tmp_1);
+if (!tint_tmp) {
+ tint_tmp = d;
}
-bool a = (tint_tmp_1);
+bool a = (tint_tmp);
)");
}
@@ -467,8 +480,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitExpression(pre, out, expr)) << gen.error();
- EXPECT_EQ(pre_result(), R"(bool tint_tmp = a;
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+ EXPECT_EQ(gen.result(), R"(bool tint_tmp = a;
if (tint_tmp) {
bool tint_tmp_1 = b;
if (!tint_tmp_1) {
@@ -477,7 +491,7 @@
tint_tmp = (tint_tmp_1);
}
)");
- EXPECT_EQ(result(), R"(asint((tint_tmp)))");
+ EXPECT_EQ(out.str(), R"(asint((tint_tmp)))");
}
TEST_F(HlslGeneratorImplTest_Binary, Call_WithLogical) {
@@ -512,8 +526,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitStatement(out, expr)) << gen.error();
- EXPECT_EQ(result(), R"(bool tint_tmp = a;
+ ASSERT_TRUE(gen.EmitStatement(expr)) << gen.error();
+ EXPECT_EQ(gen.result(), R"(bool tint_tmp = a;
if (tint_tmp) {
tint_tmp = b;
}
@@ -521,19 +535,19 @@
if (!tint_tmp_1) {
tint_tmp_1 = d;
}
-bool tint_tmp_2 = a;
-if (!tint_tmp_2) {
- tint_tmp_2 = c;
+bool tint_tmp_3 = a;
+if (!tint_tmp_3) {
+ tint_tmp_3 = c;
}
-bool tint_tmp_3 = (tint_tmp_2);
-if (tint_tmp_3) {
+bool tint_tmp_2 = (tint_tmp_3);
+if (tint_tmp_2) {
bool tint_tmp_4 = b;
if (!tint_tmp_4) {
tint_tmp_4 = d;
}
- tint_tmp_3 = (tint_tmp_4);
+ tint_tmp_2 = (tint_tmp_4);
}
-foo((tint_tmp), (tint_tmp_1), (tint_tmp_3));
+foo((tint_tmp), (tint_tmp_1), (tint_tmp_2));
)");
}
diff --git a/src/writer/hlsl/generator_impl_bitcast_test.cc b/src/writer/hlsl/generator_impl_bitcast_test.cc
index 45de152..9cb97d8 100644
--- a/src/writer/hlsl/generator_impl_bitcast_test.cc
+++ b/src/writer/hlsl/generator_impl_bitcast_test.cc
@@ -27,8 +27,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitExpression(pre, out, bitcast)) << gen.error();
- EXPECT_EQ(result(), "asfloat(1)");
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, bitcast)) << gen.error();
+ EXPECT_EQ(out.str(), "asfloat(1)");
}
TEST_F(HlslGeneratorImplTest_Bitcast, EmitExpression_Bitcast_Int) {
@@ -37,8 +38,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitExpression(pre, out, bitcast)) << gen.error();
- EXPECT_EQ(result(), "asint(1u)");
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, bitcast)) << gen.error();
+ EXPECT_EQ(out.str(), "asint(1u)");
}
TEST_F(HlslGeneratorImplTest_Bitcast, EmitExpression_Bitcast_Uint) {
@@ -47,8 +49,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitExpression(pre, out, bitcast)) << gen.error();
- EXPECT_EQ(result(), "asuint(1)");
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, bitcast)) << gen.error();
+ EXPECT_EQ(out.str(), "asuint(1)");
}
} // namespace
diff --git a/src/writer/hlsl/generator_impl_block_test.cc b/src/writer/hlsl/generator_impl_block_test.cc
index 535304b..4aec6a3 100644
--- a/src/writer/hlsl/generator_impl_block_test.cc
+++ b/src/writer/hlsl/generator_impl_block_test.cc
@@ -29,27 +29,13 @@
gen.increment_indent();
- ASSERT_TRUE(gen.EmitStatement(out, b)) << gen.error();
- EXPECT_EQ(result(), R"( {
+ ASSERT_TRUE(gen.EmitStatement(b)) << gen.error();
+ EXPECT_EQ(gen.result(), R"( {
discard;
}
)");
}
-TEST_F(HlslGeneratorImplTest_Block, Emit_Block_WithoutNewline) {
- auto* b = Block(create<ast::DiscardStatement>());
- WrapInFunction(b);
-
- GeneratorImpl& gen = Build();
-
- gen.increment_indent();
-
- ASSERT_TRUE(gen.EmitBlock(out, b)) << gen.error();
- EXPECT_EQ(result(), R"({
- discard;
- })");
-}
-
} // namespace
} // namespace hlsl
} // namespace writer
diff --git a/src/writer/hlsl/generator_impl_break_test.cc b/src/writer/hlsl/generator_impl_break_test.cc
index dbf3241..b2f532f 100644
--- a/src/writer/hlsl/generator_impl_break_test.cc
+++ b/src/writer/hlsl/generator_impl_break_test.cc
@@ -29,8 +29,8 @@
gen.increment_indent();
- ASSERT_TRUE(gen.EmitStatement(out, b)) << gen.error();
- EXPECT_EQ(result(), " break;\n");
+ ASSERT_TRUE(gen.EmitStatement(b)) << gen.error();
+ EXPECT_EQ(gen.result(), " break;\n");
}
} // namespace
diff --git a/src/writer/hlsl/generator_impl_call_test.cc b/src/writer/hlsl/generator_impl_call_test.cc
index eadfcd4..a0391e5 100644
--- a/src/writer/hlsl/generator_impl_call_test.cc
+++ b/src/writer/hlsl/generator_impl_call_test.cc
@@ -31,8 +31,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error();
- EXPECT_EQ(result(), "my_func()");
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+ EXPECT_EQ(out.str(), "my_func()");
}
TEST_F(HlslGeneratorImplTest_Call, EmitExpression_Call_WithParams) {
@@ -50,8 +51,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error();
- EXPECT_EQ(result(), "my_func(param1, param2)");
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+ EXPECT_EQ(out.str(), "my_func(param1, param2)");
}
TEST_F(HlslGeneratorImplTest_Call, EmitStatement_Call) {
@@ -70,8 +72,8 @@
GeneratorImpl& gen = Build();
gen.increment_indent();
- ASSERT_TRUE(gen.EmitStatement(out, call)) << gen.error();
- EXPECT_EQ(result(), " my_func(param1, param2);\n");
+ ASSERT_TRUE(gen.EmitStatement(call)) << gen.error();
+ EXPECT_EQ(gen.result(), " my_func(param1, param2);\n");
}
} // namespace
diff --git a/src/writer/hlsl/generator_impl_case_test.cc b/src/writer/hlsl/generator_impl_case_test.cc
index b892acb..93e80a6 100644
--- a/src/writer/hlsl/generator_impl_case_test.cc
+++ b/src/writer/hlsl/generator_impl_case_test.cc
@@ -33,8 +33,8 @@
gen.increment_indent();
- ASSERT_TRUE(gen.EmitCase(out, c)) << gen.error();
- EXPECT_EQ(result(), R"( case 5: {
+ ASSERT_TRUE(gen.EmitCase(c)) << gen.error();
+ EXPECT_EQ(gen.result(), R"( case 5: {
break;
}
)");
@@ -50,8 +50,8 @@
gen.increment_indent();
- ASSERT_TRUE(gen.EmitCase(out, c)) << gen.error();
- EXPECT_EQ(result(), R"( case 5: {
+ ASSERT_TRUE(gen.EmitCase(c)) << gen.error();
+ EXPECT_EQ(gen.result(), R"( case 5: {
break;
}
)");
@@ -68,8 +68,8 @@
gen.increment_indent();
- ASSERT_TRUE(gen.EmitCase(out, c)) << gen.error();
- EXPECT_EQ(result(), R"( case 5: {
+ ASSERT_TRUE(gen.EmitCase(c)) << gen.error();
+ EXPECT_EQ(gen.result(), R"( case 5: {
/* fallthrough */
}
)");
@@ -87,8 +87,8 @@
gen.increment_indent();
- ASSERT_TRUE(gen.EmitCase(out, c)) << gen.error();
- EXPECT_EQ(result(), R"( case 5:
+ ASSERT_TRUE(gen.EmitCase(c)) << gen.error();
+ EXPECT_EQ(gen.result(), R"( case 5:
case 6: {
break;
}
@@ -104,8 +104,8 @@
gen.increment_indent();
- ASSERT_TRUE(gen.EmitCase(out, c)) << gen.error();
- EXPECT_EQ(result(), R"( default: {
+ ASSERT_TRUE(gen.EmitCase(c)) << gen.error();
+ EXPECT_EQ(gen.result(), R"( default: {
break;
}
)");
diff --git a/src/writer/hlsl/generator_impl_cast_test.cc b/src/writer/hlsl/generator_impl_cast_test.cc
index f123536..b281f4b 100644
--- a/src/writer/hlsl/generator_impl_cast_test.cc
+++ b/src/writer/hlsl/generator_impl_cast_test.cc
@@ -27,8 +27,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitExpression(pre, out, cast)) << gen.error();
- EXPECT_EQ(result(), "float(1)");
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, cast)) << gen.error();
+ EXPECT_EQ(out.str(), "float(1)");
}
TEST_F(HlslGeneratorImplTest_Cast, EmitExpression_Cast_Vector) {
@@ -37,8 +38,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitExpression(pre, out, cast)) << gen.error();
- EXPECT_EQ(result(), "float3(int3(1, 2, 3))");
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, cast)) << gen.error();
+ EXPECT_EQ(out.str(), "float3(int3(1, 2, 3))");
}
} // namespace
diff --git a/src/writer/hlsl/generator_impl_constructor_test.cc b/src/writer/hlsl/generator_impl_constructor_test.cc
index 9d18226..429b86d 100644
--- a/src/writer/hlsl/generator_impl_constructor_test.cc
+++ b/src/writer/hlsl/generator_impl_constructor_test.cc
@@ -29,8 +29,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr("false"));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("false"));
}
TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Int) {
@@ -38,8 +38,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr("-12345"));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("-12345"));
}
TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_UInt) {
@@ -47,8 +47,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr("56779u"));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("56779u"));
}
TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Float) {
@@ -57,8 +57,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr("1073741824.0f"));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("1073741824.0f"));
}
TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Float) {
@@ -66,8 +66,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr("float(-0.000012f)"));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("float(-0.000012f)"));
}
TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Bool) {
@@ -75,8 +75,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr("bool(true)"));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("bool(true)"));
}
TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Int) {
@@ -84,8 +84,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr("int(-12345)"));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("int(-12345)"));
}
TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Uint) {
@@ -93,8 +93,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr("uint(12345u)"));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("uint(12345u)"));
}
TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Vec) {
@@ -102,8 +102,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr("float3(1.0f, 2.0f, 3.0f)"));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("float3(1.0f, 2.0f, 3.0f)"));
}
TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Vec_Empty) {
@@ -111,8 +111,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr("float3(0.0f, 0.0f, 0.0f)"));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("float3(0.0f, 0.0f, 0.0f)"));
}
TEST_F(HlslGeneratorImplTest_Constructor,
@@ -121,8 +121,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr("float3((2.0f).xxx)"));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("float3((2.0f).xxx)"));
}
TEST_F(HlslGeneratorImplTest_Constructor,
@@ -131,8 +131,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr("bool3((true).xxx)"));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("bool3((true).xxx)"));
}
TEST_F(HlslGeneratorImplTest_Constructor,
@@ -141,8 +141,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr("int3((2).xxx)"));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("int3((2).xxx)"));
}
TEST_F(HlslGeneratorImplTest_Constructor,
@@ -151,8 +151,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr("uint3((2u).xxx)"));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("uint3((2u).xxx)"));
}
TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Mat) {
@@ -161,10 +161,10 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
+ ASSERT_TRUE(gen.Generate()) << gen.error();
EXPECT_THAT(
- result(),
+ gen.result(),
HasSubstr(
"float2x3(float3(1.0f, 2.0f, 3.0f), float3(3.0f, 4.0f, 5.0f))"));
}
@@ -174,9 +174,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
+ ASSERT_TRUE(gen.Generate()) << gen.error();
- EXPECT_THAT(result(),
+ EXPECT_THAT(gen.result(),
HasSubstr("float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)"));
}
@@ -187,8 +187,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(),
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(),
HasSubstr("{float3(1.0f, 2.0f, 3.0f), float3(4.0f, 5.0f, 6.0f),"
" float3(7.0f, 8.0f, 9.0f)}"));
}
@@ -200,8 +200,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(),
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(),
HasSubstr("{float3(0.0f, 0.0f, 0.0f), float3(0.0f, 0.0f, 0.0f),"
" float3(0.0f, 0.0f, 0.0f)}"));
}
@@ -217,8 +217,8 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr("{1, 2.0f, int3(3, 4, 5)}"));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("{1, 2.0f, int3(3, 4, 5)}"));
}
TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Struct_Empty) {
@@ -232,8 +232,8 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr("{0, 0.0f, int3(0, 0, 0)}"));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("{0, 0.0f, int3(0, 0, 0)}"));
}
} // namespace
diff --git a/src/writer/hlsl/generator_impl_continue_test.cc b/src/writer/hlsl/generator_impl_continue_test.cc
index 6a2482b..e312b82 100644
--- a/src/writer/hlsl/generator_impl_continue_test.cc
+++ b/src/writer/hlsl/generator_impl_continue_test.cc
@@ -29,8 +29,8 @@
gen.increment_indent();
- ASSERT_TRUE(gen.EmitStatement(out, loop)) << gen.error();
- EXPECT_EQ(result(), R"( while (true) {
+ ASSERT_TRUE(gen.EmitStatement(loop)) << gen.error();
+ EXPECT_EQ(gen.result(), R"( while (true) {
continue;
}
)");
diff --git a/src/writer/hlsl/generator_impl_discard_test.cc b/src/writer/hlsl/generator_impl_discard_test.cc
index e738ef7..e7a5cca 100644
--- a/src/writer/hlsl/generator_impl_discard_test.cc
+++ b/src/writer/hlsl/generator_impl_discard_test.cc
@@ -29,8 +29,8 @@
gen.increment_indent();
- ASSERT_TRUE(gen.EmitStatement(out, stmt)) << gen.error();
- EXPECT_EQ(result(), " discard;\n");
+ ASSERT_TRUE(gen.EmitStatement(stmt)) << gen.error();
+ EXPECT_EQ(gen.result(), " discard;\n");
}
} // namespace
diff --git a/src/writer/hlsl/generator_impl_function_test.cc b/src/writer/hlsl/generator_impl_function_test.cc
index 7af36c4..6150010 100644
--- a/src/writer/hlsl/generator_impl_function_test.cc
+++ b/src/writer/hlsl/generator_impl_function_test.cc
@@ -38,8 +38,8 @@
gen.increment_indent();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_EQ(result(), R"( void my_func() {
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"( void my_func() {
return;
}
)");
@@ -55,8 +55,8 @@
gen.increment_indent();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr(R"( void tint_symbol() {
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr(R"( void tint_symbol() {
return;
})"));
}
@@ -72,8 +72,8 @@
gen.increment_indent();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_EQ(result(), R"( void my_func(float a, int b) {
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"( void my_func(float a, int b) {
return;
}
)");
@@ -88,8 +88,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_EQ(result(), R"(void main() {
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(void main() {
return;
}
)");
@@ -104,8 +104,8 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr(R"(float f(inout float foo) {
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr(R"(float f(inout float foo) {
return foo;
}
)"));
@@ -122,8 +122,8 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_EQ(result(), R"(struct tint_symbol_1 {
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(struct tint_symbol_1 {
float foo : TEXCOORD0;
};
struct tint_symbol_2 {
@@ -152,8 +152,8 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_EQ(result(), R"(struct tint_symbol_1 {
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(struct tint_symbol_1 {
float4 coord : SV_Position;
};
struct tint_symbol_2 {
@@ -206,8 +206,8 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_EQ(result(), R"(struct Interface {
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(struct Interface {
float4 pos;
float col1;
float col2;
@@ -275,8 +275,8 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_EQ(result(), R"(struct VertexOutput {
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(struct VertexOutput {
float4 pos;
};
@@ -340,8 +340,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_EQ(result(), R"(cbuffer cbuffer_ubo : register(b0, space1) {
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(cbuffer cbuffer_ubo : register(b0, space1) {
uint4 ubo[1];
};
@@ -381,8 +381,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_EQ(result(), R"(cbuffer cbuffer_uniforms : register(b0, space1) {
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(cbuffer cbuffer_uniforms : register(b0, space1) {
uint4 uniforms[1];
};
@@ -423,8 +423,8 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_EQ(result(),
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(),
R"(RWByteAddressBuffer coord : register(u0, space1);
void frag_main() {
@@ -463,8 +463,8 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_EQ(result(),
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(),
R"(ByteAddressBuffer coord : register(t0, space1);
void frag_main() {
@@ -500,8 +500,8 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_EQ(result(),
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(),
R"(RWByteAddressBuffer coord : register(u0, space1);
void frag_main() {
@@ -538,8 +538,8 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_EQ(result(),
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(),
R"(RWByteAddressBuffer coord : register(u0, space1);
void frag_main() {
@@ -578,8 +578,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_EQ(result(), R"(cbuffer cbuffer_coord : register(b0, space1) {
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(cbuffer cbuffer_coord : register(b0, space1) {
uint4 coord[1];
};
@@ -624,8 +624,8 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_EQ(result(),
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(),
R"(RWByteAddressBuffer coord : register(u0, space1);
float sub_func(float param) {
@@ -648,8 +648,8 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_EQ(result(), R"(void tint_symbol() {
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(void tint_symbol() {
return;
}
)");
@@ -666,8 +666,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_EQ(result(), R"([numthreads(1, 1, 1)]
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"([numthreads(1, 1, 1)]
void main() {
return;
}
@@ -684,8 +684,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_EQ(result(), R"([numthreads(2, 4, 6)]
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"([numthreads(2, 4, 6)]
void main() {
return;
}
@@ -705,8 +705,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_EQ(result(), R"(static const int width = int(2);
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(static const int width = int(2);
static const int height = int(3);
static const int depth = int(4);
@@ -730,8 +730,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_EQ(result(), R"(#ifndef WGSL_SPEC_CONSTANT_7
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(#ifndef WGSL_SPEC_CONSTANT_7
#define WGSL_SPEC_CONSTANT_7 int(2)
#endif
static const int width = WGSL_SPEC_CONSTANT_7;
@@ -759,8 +759,8 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr(R"(
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr(R"(
struct tint_array_wrapper {
float arr[5];
};
@@ -779,8 +779,8 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr(R"(
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr(R"(
struct tint_array_wrapper {
float arr[5];
};
@@ -851,8 +851,8 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_EQ(result(), R"(RWByteAddressBuffer data : register(u0, space0);
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(RWByteAddressBuffer data : register(u0, space0);
[numthreads(1, 1, 1)]
void a() {
diff --git a/src/writer/hlsl/generator_impl_identifier_test.cc b/src/writer/hlsl/generator_impl_identifier_test.cc
index 161cab0..4929bc9 100644
--- a/src/writer/hlsl/generator_impl_identifier_test.cc
+++ b/src/writer/hlsl/generator_impl_identifier_test.cc
@@ -29,8 +29,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitExpression(pre, out, i)) << gen.error();
- EXPECT_EQ(result(), "foo");
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, i)) << gen.error();
+ EXPECT_EQ(out.str(), "foo");
}
} // namespace
diff --git a/src/writer/hlsl/generator_impl_if_test.cc b/src/writer/hlsl/generator_impl_if_test.cc
index f2aece3..e40f842 100644
--- a/src/writer/hlsl/generator_impl_if_test.cc
+++ b/src/writer/hlsl/generator_impl_if_test.cc
@@ -32,8 +32,8 @@
GeneratorImpl& gen = Build();
gen.increment_indent();
- ASSERT_TRUE(gen.EmitStatement(out, i)) << gen.error();
- EXPECT_EQ(result(), R"( if (cond) {
+ ASSERT_TRUE(gen.EmitStatement(i)) << gen.error();
+ EXPECT_EQ(gen.result(), R"( if (cond) {
return;
}
)");
@@ -57,8 +57,8 @@
gen.increment_indent();
- ASSERT_TRUE(gen.EmitStatement(out, i)) << gen.error();
- EXPECT_EQ(result(), R"( if (cond) {
+ ASSERT_TRUE(gen.EmitStatement(i)) << gen.error();
+ EXPECT_EQ(gen.result(), R"( if (cond) {
return;
} else {
if (else_cond) {
@@ -84,8 +84,8 @@
gen.increment_indent();
- ASSERT_TRUE(gen.EmitStatement(out, i)) << gen.error();
- EXPECT_EQ(result(), R"( if (cond) {
+ ASSERT_TRUE(gen.EmitStatement(i)) << gen.error();
+ EXPECT_EQ(gen.result(), R"( if (cond) {
return;
} else {
return;
@@ -116,8 +116,8 @@
gen.increment_indent();
- ASSERT_TRUE(gen.EmitStatement(out, i)) << gen.error();
- EXPECT_EQ(result(), R"( if (cond) {
+ ASSERT_TRUE(gen.EmitStatement(i)) << gen.error();
+ EXPECT_EQ(gen.result(), R"( if (cond) {
return;
} else {
if (else_cond) {
diff --git a/src/writer/hlsl/generator_impl_import_test.cc b/src/writer/hlsl/generator_impl_import_test.cc
index 201f345..ec81c41 100644
--- a/src/writer/hlsl/generator_impl_import_test.cc
+++ b/src/writer/hlsl/generator_impl_import_test.cc
@@ -40,8 +40,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitCall(pre, out, expr)) << gen.error();
- EXPECT_EQ(result(), std::string(param.hlsl_name) + "(1.0f)");
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitCall(out, expr)) << gen.error();
+ EXPECT_EQ(out.str(), std::string(param.hlsl_name) + "(1.0f)");
}
INSTANTIATE_TEST_SUITE_P(HlslGeneratorImplTest_Import,
HlslImportData_SingleParamTest,
@@ -78,8 +79,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitCall(pre, out, expr)) << gen.error();
- EXPECT_EQ(result(), std::string(param.hlsl_name) + "(1)");
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitCall(out, expr)) << gen.error();
+ EXPECT_EQ(out.str(), std::string(param.hlsl_name) + "(1)");
}
INSTANTIATE_TEST_SUITE_P(HlslGeneratorImplTest_Import,
HlslImportData_SingleIntParamTest,
@@ -95,8 +97,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitCall(pre, out, expr)) << gen.error();
- EXPECT_EQ(result(),
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitCall(out, expr)) << gen.error();
+ EXPECT_EQ(out.str(),
std::string(param.hlsl_name) + "(float3(1.0f, 2.0f, 3.0f))");
}
INSTANTIATE_TEST_SUITE_P(HlslGeneratorImplTest_Import,
@@ -136,8 +139,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitCall(pre, out, expr)) << gen.error();
- EXPECT_EQ(result(), std::string(param.hlsl_name) + "(1.0f, 2.0f)");
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitCall(out, expr)) << gen.error();
+ EXPECT_EQ(out.str(), std::string(param.hlsl_name) + "(1.0f, 2.0f)");
}
INSTANTIATE_TEST_SUITE_P(HlslGeneratorImplTest_Import,
HlslImportData_DualParamTest,
@@ -159,8 +163,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitCall(pre, out, expr)) << gen.error();
- EXPECT_EQ(result(),
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitCall(out, expr)) << gen.error();
+ EXPECT_EQ(out.str(),
std::string(param.hlsl_name) +
"(float3(1.0f, 2.0f, 3.0f), float3(4.0f, 5.0f, 6.0f))");
}
@@ -177,8 +182,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitCall(pre, out, expr)) << gen.error();
- EXPECT_EQ(result(), std::string(param.hlsl_name) + "(1, 2)");
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitCall(out, expr)) << gen.error();
+ EXPECT_EQ(out.str(), std::string(param.hlsl_name) + "(1, 2)");
}
INSTANTIATE_TEST_SUITE_P(HlslGeneratorImplTest_Import,
HlslImportData_DualParam_Int_Test,
@@ -194,8 +200,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitCall(pre, out, expr)) << gen.error();
- EXPECT_EQ(result(), std::string(param.hlsl_name) + "(1.0f, 2.0f, 3.0f)");
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitCall(out, expr)) << gen.error();
+ EXPECT_EQ(out.str(), std::string(param.hlsl_name) + "(1.0f, 2.0f, 3.0f)");
}
INSTANTIATE_TEST_SUITE_P(
HlslGeneratorImplTest_Import,
@@ -218,8 +225,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitCall(pre, out, expr)) << gen.error();
- EXPECT_EQ(result(), std::string(param.hlsl_name) + "(1, 2, 3)");
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitCall(out, expr)) << gen.error();
+ EXPECT_EQ(out.str(), std::string(param.hlsl_name) + "(1, 2, 3)");
}
INSTANTIATE_TEST_SUITE_P(HlslGeneratorImplTest_Import,
HlslImportData_TripleParam_Int_Test,
@@ -233,8 +241,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitCall(pre, out, expr)) << gen.error();
- EXPECT_EQ(result(), std::string("determinant(var)"));
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitCall(out, expr)) << gen.error();
+ EXPECT_EQ(out.str(), std::string("determinant(var)"));
}
} // namespace
diff --git a/src/writer/hlsl/generator_impl_intrinsic_test.cc b/src/writer/hlsl/generator_impl_intrinsic_test.cc
index 0e5a619..4ec4192 100644
--- a/src/writer/hlsl/generator_impl_intrinsic_test.cc
+++ b/src/writer/hlsl/generator_impl_intrinsic_test.cc
@@ -268,8 +268,9 @@
GeneratorImpl& gen = Build();
gen.increment_indent();
- ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error();
- EXPECT_EQ(result(), "dot(param1, param2)");
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+ EXPECT_EQ(out.str(), "dot(param1, param2)");
}
TEST_F(HlslGeneratorImplTest_Intrinsic, Select_Scalar) {
@@ -278,8 +279,9 @@
GeneratorImpl& gen = Build();
gen.increment_indent();
- ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error();
- EXPECT_EQ(result(), "(true ? 1.0f : 2.0f)");
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+ EXPECT_EQ(out.str(), "(true ? 1.0f : 2.0f)");
}
TEST_F(HlslGeneratorImplTest_Intrinsic, Select_Vector) {
@@ -289,8 +291,9 @@
GeneratorImpl& gen = Build();
gen.increment_indent();
- ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error();
- EXPECT_EQ(result(), "(bool2(true, false) ? int2(1, 2) : int2(3, 4))");
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+ EXPECT_EQ(out.str(), "(bool2(true, false) ? int2(1, 2) : int2(3, 4))");
}
TEST_F(HlslGeneratorImplTest_Intrinsic, Modf_Scalar) {
@@ -300,8 +303,8 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr("modf(1.0f, res)"));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("modf(1.0f, res)"));
}
TEST_F(HlslGeneratorImplTest_Intrinsic, Modf_Vector) {
@@ -311,8 +314,8 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr("modf(float3(0.0f, 0.0f, 0.0f), res)"));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("modf(float3(0.0f, 0.0f, 0.0f), res)"));
}
TEST_F(HlslGeneratorImplTest_Intrinsic, Frexp_Scalar_i32) {
@@ -322,8 +325,8 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr(R"(
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr(R"(
float tint_tmp;
float tint_tmp_1 = frexp(1.0f, tint_tmp);
exp = int(tint_tmp);
@@ -338,8 +341,8 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr(R"(
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr(R"(
float3 tint_tmp;
float3 tint_tmp_1 = frexp(float3(0.0f, 0.0f, 0.0f), tint_tmp);
res = int3(tint_tmp);
@@ -354,8 +357,8 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr(R"(
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr(R"(
uint tint_isnormal_exponent = asuint(val) & 0x7f80000;
uint tint_isnormal_clamped = clamp(tint_isnormal_exponent, 0x0080000, 0x7f00000);
(tint_isnormal_clamped == tint_isnormal_exponent);
@@ -369,8 +372,8 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr(R"(
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr(R"(
uint3 tint_isnormal_exponent = asuint(val) & 0x7f80000;
uint3 tint_isnormal_clamped = clamp(tint_isnormal_exponent, 0x0080000, 0x7f00000);
(tint_isnormal_clamped == tint_isnormal_exponent);
@@ -384,11 +387,12 @@
GeneratorImpl& gen = Build();
gen.increment_indent();
- ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error();
- EXPECT_THAT(pre_result(), HasSubstr("int4 tint_tmp = int4(round(clamp(p1, "
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("int4 tint_tmp = int4(round(clamp(p1, "
"-1.0, 1.0) * 127.0)) & 0xff;"));
- EXPECT_THAT(result(), HasSubstr("asuint(tint_tmp.x | tint_tmp.y << 8 | "
- "tint_tmp.z << 16 | tint_tmp.w << 24)"));
+ EXPECT_THAT(out.str(), HasSubstr("asuint(tint_tmp.x | tint_tmp.y << 8 | "
+ "tint_tmp.z << 16 | tint_tmp.w << 24)"));
}
TEST_F(HlslGeneratorImplTest_Intrinsic, Pack4x8Unorm) {
@@ -398,11 +402,12 @@
GeneratorImpl& gen = Build();
gen.increment_indent();
- ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error();
- EXPECT_THAT(pre_result(), HasSubstr("uint4 tint_tmp = uint4(round(clamp(p1, "
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("uint4 tint_tmp = uint4(round(clamp(p1, "
"0.0, 1.0) * 255.0));"));
- EXPECT_THAT(result(), HasSubstr("(tint_tmp.x | tint_tmp.y << 8 | "
- "tint_tmp.z << 16 | tint_tmp.w << 24)"));
+ EXPECT_THAT(out.str(), HasSubstr("(tint_tmp.x | tint_tmp.y << 8 | "
+ "tint_tmp.z << 16 | tint_tmp.w << 24)"));
}
TEST_F(HlslGeneratorImplTest_Intrinsic, Pack2x16Snorm) {
@@ -412,10 +417,11 @@
GeneratorImpl& gen = Build();
gen.increment_indent();
- ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error();
- EXPECT_THAT(pre_result(), HasSubstr("int2 tint_tmp = int2(round(clamp(p1, "
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("int2 tint_tmp = int2(round(clamp(p1, "
"-1.0, 1.0) * 32767.0)) & 0xffff;"));
- EXPECT_THAT(result(), HasSubstr("asuint(tint_tmp.x | tint_tmp.y << 16)"));
+ EXPECT_THAT(out.str(), HasSubstr("asuint(tint_tmp.x | tint_tmp.y << 16)"));
}
TEST_F(HlslGeneratorImplTest_Intrinsic, Pack2x16Unorm) {
@@ -425,10 +431,11 @@
GeneratorImpl& gen = Build();
gen.increment_indent();
- ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error();
- EXPECT_THAT(pre_result(), HasSubstr("uint2 tint_tmp = uint2(round(clamp(p1, "
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("uint2 tint_tmp = uint2(round(clamp(p1, "
"0.0, 1.0) * 65535.0));"));
- EXPECT_THAT(result(), HasSubstr("(tint_tmp.x | tint_tmp.y << 16)"));
+ EXPECT_THAT(out.str(), HasSubstr("(tint_tmp.x | tint_tmp.y << 16)"));
}
TEST_F(HlslGeneratorImplTest_Intrinsic, Pack2x16Float) {
@@ -438,9 +445,10 @@
GeneratorImpl& gen = Build();
gen.increment_indent();
- ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error();
- EXPECT_THAT(pre_result(), HasSubstr("uint2 tint_tmp = f32tof16(p1);"));
- EXPECT_THAT(result(), HasSubstr("(tint_tmp.x | tint_tmp.y << 16)"));
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("uint2 tint_tmp = f32tof16(p1);"));
+ EXPECT_THAT(out.str(), HasSubstr("(tint_tmp.x | tint_tmp.y << 16)"));
}
TEST_F(HlslGeneratorImplTest_Intrinsic, Unpack4x8Snorm) {
@@ -450,12 +458,13 @@
GeneratorImpl& gen = Build();
gen.increment_indent();
- ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error();
- EXPECT_THAT(pre_result(), HasSubstr("int tint_tmp_1 = int(p1);"));
- EXPECT_THAT(pre_result(),
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("int tint_tmp_1 = int(p1);"));
+ EXPECT_THAT(gen.result(),
HasSubstr("int4 tint_tmp = int4(tint_tmp_1 << 24, tint_tmp_1 "
"<< 16, tint_tmp_1 << 8, tint_tmp_1) >> 24;"));
- EXPECT_THAT(result(),
+ EXPECT_THAT(out.str(),
HasSubstr("clamp(float4(tint_tmp) / 127.0, -1.0, 1.0)"));
}
@@ -466,13 +475,14 @@
GeneratorImpl& gen = Build();
gen.increment_indent();
- ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error();
- EXPECT_THAT(pre_result(), HasSubstr("uint tint_tmp_1 = p1;"));
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("uint tint_tmp_1 = p1;"));
EXPECT_THAT(
- pre_result(),
+ gen.result(),
HasSubstr("uint4 tint_tmp = uint4(tint_tmp_1 & 0xff, (tint_tmp_1 >> "
"8) & 0xff, (tint_tmp_1 >> 16) & 0xff, tint_tmp_1 >> 24);"));
- EXPECT_THAT(result(), HasSubstr("float4(tint_tmp) / 255.0"));
+ EXPECT_THAT(out.str(), HasSubstr("float4(tint_tmp) / 255.0"));
}
TEST_F(HlslGeneratorImplTest_Intrinsic, Unpack2x16Snorm) {
@@ -482,12 +492,13 @@
GeneratorImpl& gen = Build();
gen.increment_indent();
- ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error();
- EXPECT_THAT(pre_result(), HasSubstr("int tint_tmp_1 = int(p1);"));
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("int tint_tmp_1 = int(p1);"));
EXPECT_THAT(
- pre_result(),
+ gen.result(),
HasSubstr("int2 tint_tmp = int2(tint_tmp_1 << 16, tint_tmp_1) >> 16;"));
- EXPECT_THAT(result(),
+ EXPECT_THAT(out.str(),
HasSubstr("clamp(float2(tint_tmp) / 32767.0, -1.0, 1.0)"));
}
@@ -498,12 +509,13 @@
GeneratorImpl& gen = Build();
gen.increment_indent();
- ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error();
- EXPECT_THAT(pre_result(), HasSubstr("uint tint_tmp_1 = p1;"));
- EXPECT_THAT(pre_result(),
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("uint tint_tmp_1 = p1;"));
+ EXPECT_THAT(gen.result(),
HasSubstr("uint2 tint_tmp = uint2(tint_tmp_1 & 0xffff, "
"tint_tmp_1 >> 16);"));
- EXPECT_THAT(result(), HasSubstr("float2(tint_tmp) / 65535.0"));
+ EXPECT_THAT(out.str(), HasSubstr("float2(tint_tmp) / 65535.0"));
}
TEST_F(HlslGeneratorImplTest_Intrinsic, Unpack2x16Float) {
@@ -513,9 +525,10 @@
GeneratorImpl& gen = Build();
gen.increment_indent();
- ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error();
- EXPECT_THAT(pre_result(), HasSubstr("uint tint_tmp = p1;"));
- EXPECT_THAT(result(),
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("uint tint_tmp = p1;"));
+ EXPECT_THAT(out.str(),
HasSubstr("f16tof32(uint2(tint_tmp & 0xffff, tint_tmp >> 16))"));
}
@@ -526,8 +539,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_EQ(result(), R"([numthreads(1, 1, 1)]
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"([numthreads(1, 1, 1)]
void main() {
DeviceMemoryBarrierWithGroupSync();
return;
@@ -542,8 +555,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_EQ(result(), R"([numthreads(1, 1, 1)]
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"([numthreads(1, 1, 1)]
void main() {
GroupMemoryBarrierWithGroupSync();
return;
@@ -561,8 +574,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_EQ(result(), R"(int f(int a, int b, int c) {
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(int f(int a, int b, int c) {
return ((a + b) * c);
}
diff --git a/src/writer/hlsl/generator_impl_intrinsic_texture_test.cc b/src/writer/hlsl/generator_impl_intrinsic_texture_test.cc
index 82475fc..b868e74 100644
--- a/src/writer/hlsl/generator_impl_intrinsic_texture_test.cc
+++ b/src/writer/hlsl/generator_impl_intrinsic_texture_test.cc
@@ -381,12 +381,12 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
+ ASSERT_TRUE(gen.Generate()) << gen.error();
auto expected = expected_texture_overload(param.overload);
- EXPECT_THAT(result(), HasSubstr(expected.pre));
- EXPECT_THAT(result(), HasSubstr(expected.out));
+ EXPECT_THAT(gen.result(), HasSubstr(expected.pre));
+ EXPECT_THAT(gen.result(), HasSubstr(expected.out));
}
INSTANTIATE_TEST_SUITE_P(
diff --git a/src/writer/hlsl/generator_impl_loop_test.cc b/src/writer/hlsl/generator_impl_loop_test.cc
index 906d56e..a775429 100644
--- a/src/writer/hlsl/generator_impl_loop_test.cc
+++ b/src/writer/hlsl/generator_impl_loop_test.cc
@@ -33,8 +33,8 @@
gen.increment_indent();
- ASSERT_TRUE(gen.EmitStatement(out, l)) << gen.error();
- EXPECT_EQ(result(), R"( while (true) {
+ ASSERT_TRUE(gen.EmitStatement(l)) << gen.error();
+ EXPECT_EQ(gen.result(), R"( while (true) {
discard;
}
)");
@@ -51,8 +51,8 @@
gen.increment_indent();
- ASSERT_TRUE(gen.EmitStatement(out, l)) << gen.error();
- EXPECT_EQ(result(), R"( while (true) {
+ ASSERT_TRUE(gen.EmitStatement(l)) << gen.error();
+ EXPECT_EQ(gen.result(), R"( while (true) {
discard;
{
return;
@@ -83,8 +83,8 @@
gen.increment_indent();
- ASSERT_TRUE(gen.EmitStatement(out, outer)) << gen.error();
- EXPECT_EQ(result(), R"( while (true) {
+ ASSERT_TRUE(gen.EmitStatement(outer)) << gen.error();
+ EXPECT_EQ(gen.result(), R"( while (true) {
while (true) {
discard;
{
@@ -137,8 +137,8 @@
gen.increment_indent();
- ASSERT_TRUE(gen.EmitStatement(out, outer)) << gen.error();
- EXPECT_EQ(result(), R"( while (true) {
+ ASSERT_TRUE(gen.EmitStatement(outer)) << gen.error();
+ EXPECT_EQ(gen.result(), R"( while (true) {
float lhs = 2.400000095f;
float other = 0.0f;
{
diff --git a/src/writer/hlsl/generator_impl_member_accessor_test.cc b/src/writer/hlsl/generator_impl_member_accessor_test.cc
index 203551d..2afc81a 100644
--- a/src/writer/hlsl/generator_impl_member_accessor_test.cc
+++ b/src/writer/hlsl/generator_impl_member_accessor_test.cc
@@ -131,8 +131,8 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_EQ(result(), R"(struct Data {
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(struct Data {
float mem;
};
@@ -181,8 +181,8 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr(p.expected));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr(p.expected));
}
INSTANTIATE_TEST_SUITE_P(
@@ -254,8 +254,8 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr(p.expected));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr(p.expected));
}
INSTANTIATE_TEST_SUITE_P(
@@ -339,7 +339,7 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
+ ASSERT_TRUE(gen.Generate()) << gen.error();
auto* expected =
R"(void tint_symbol_1(RWByteAddressBuffer buffer, uint offset, float2x3 value) {
buffer.Store3((offset + 0u), asuint(value[0u]));
@@ -353,7 +353,7 @@
return;
}
)";
- EXPECT_EQ(result(), expected);
+ EXPECT_EQ(gen.result(), expected);
}
TEST_F(HlslGeneratorImplTest_MemberAccessor,
@@ -378,7 +378,7 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
+ ASSERT_TRUE(gen.Generate()) << gen.error();
auto* expected =
R"(RWByteAddressBuffer data : register(u0, space1);
@@ -387,7 +387,7 @@
return;
}
)";
- EXPECT_EQ(result(), expected);
+ EXPECT_EQ(gen.result(), expected);
}
TEST_F(HlslGeneratorImplTest_MemberAccessor,
@@ -410,7 +410,7 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
+ ASSERT_TRUE(gen.Generate()) << gen.error();
auto* expected =
R"(RWByteAddressBuffer data : register(u0, space1);
@@ -419,7 +419,7 @@
return;
}
)";
- EXPECT_EQ(result(), expected);
+ EXPECT_EQ(gen.result(), expected);
}
TEST_F(HlslGeneratorImplTest_MemberAccessor,
@@ -443,7 +443,7 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
+ ASSERT_TRUE(gen.Generate()) << gen.error();
auto* expected =
R"(RWByteAddressBuffer data : register(u0, space1);
@@ -452,7 +452,7 @@
return;
}
)";
- EXPECT_EQ(result(), expected);
+ EXPECT_EQ(gen.result(), expected);
}
TEST_F(HlslGeneratorImplTest_MemberAccessor, StorageBuffer_Store_ToArray) {
@@ -473,7 +473,7 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
+ ASSERT_TRUE(gen.Generate()) << gen.error();
auto* expected =
R"(RWByteAddressBuffer data : register(u0, space1);
@@ -482,7 +482,7 @@
return;
}
)";
- EXPECT_EQ(result(), expected);
+ EXPECT_EQ(gen.result(), expected);
}
TEST_F(HlslGeneratorImplTest_MemberAccessor, StorageBuffer_Load_MultiLevel) {
@@ -514,7 +514,7 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
+ ASSERT_TRUE(gen.Generate()) << gen.error();
auto* expected =
R"(RWByteAddressBuffer data : register(u0, space1);
@@ -523,7 +523,7 @@
return;
}
)";
- EXPECT_EQ(result(), expected);
+ EXPECT_EQ(gen.result(), expected);
}
TEST_F(HlslGeneratorImplTest_MemberAccessor,
@@ -558,7 +558,7 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
+ ASSERT_TRUE(gen.Generate()) << gen.error();
auto* expected =
R"(RWByteAddressBuffer data : register(u0, space1);
@@ -567,7 +567,7 @@
return;
}
)";
- EXPECT_EQ(result(), expected);
+ EXPECT_EQ(gen.result(), expected);
}
TEST_F(HlslGeneratorImplTest_MemberAccessor,
@@ -602,7 +602,7 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
+ ASSERT_TRUE(gen.Generate()) << gen.error();
auto* expected =
R"(RWByteAddressBuffer data : register(u0, space1);
@@ -611,7 +611,7 @@
return;
}
)";
- EXPECT_EQ(result(), expected);
+ EXPECT_EQ(gen.result(), expected);
}
TEST_F(HlslGeneratorImplTest_MemberAccessor,
@@ -646,7 +646,7 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
+ ASSERT_TRUE(gen.Generate()) << gen.error();
auto* expected =
R"(RWByteAddressBuffer data : register(u0, space1);
@@ -655,7 +655,7 @@
return;
}
)";
- EXPECT_EQ(result(), expected);
+ EXPECT_EQ(gen.result(), expected);
}
TEST_F(HlslGeneratorImplTest_MemberAccessor, StorageBuffer_Store_MultiLevel) {
@@ -686,7 +686,7 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
+ ASSERT_TRUE(gen.Generate()) << gen.error();
auto* expected =
R"(RWByteAddressBuffer data : register(u0, space1);
@@ -695,7 +695,7 @@
return;
}
)";
- EXPECT_EQ(result(), expected);
+ EXPECT_EQ(gen.result(), expected);
}
TEST_F(HlslGeneratorImplTest_MemberAccessor,
@@ -730,7 +730,7 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
+ ASSERT_TRUE(gen.Generate()) << gen.error();
auto* expected =
R"(RWByteAddressBuffer data : register(u0, space1);
@@ -739,7 +739,7 @@
return;
}
)";
- EXPECT_EQ(result(), expected);
+ EXPECT_EQ(gen.result(), expected);
}
TEST_F(HlslGeneratorImplTest_MemberAccessor, Swizzle_xyz) {
@@ -749,8 +749,8 @@
WrapInFunction(var, expr);
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr("my_vec.xyz"));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("my_vec.xyz"));
}
TEST_F(HlslGeneratorImplTest_MemberAccessor, Swizzle_gbr) {
@@ -760,8 +760,8 @@
WrapInFunction(var, expr);
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr("my_vec.gbr"));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("my_vec.gbr"));
}
} // namespace
diff --git a/src/writer/hlsl/generator_impl_module_constant_test.cc b/src/writer/hlsl/generator_impl_module_constant_test.cc
index 8ff9c88..f84f31d 100644
--- a/src/writer/hlsl/generator_impl_module_constant_test.cc
+++ b/src/writer/hlsl/generator_impl_module_constant_test.cc
@@ -28,8 +28,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitProgramConstVariable(out, var)) << gen.error();
- EXPECT_EQ(result(), "static const float pos[3] = {1.0f, 2.0f, 3.0f};\n");
+ ASSERT_TRUE(gen.EmitProgramConstVariable(var)) << gen.error();
+ EXPECT_EQ(gen.result(), "static const float pos[3] = {1.0f, 2.0f, 3.0f};\n");
}
TEST_F(HlslGeneratorImplTest_ModuleConstant, Emit_SpecConstant) {
@@ -40,8 +40,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitProgramConstVariable(out, var)) << gen.error();
- EXPECT_EQ(result(), R"(#ifndef WGSL_SPEC_CONSTANT_23
+ ASSERT_TRUE(gen.EmitProgramConstVariable(var)) << gen.error();
+ EXPECT_EQ(gen.result(), R"(#ifndef WGSL_SPEC_CONSTANT_23
#define WGSL_SPEC_CONSTANT_23 3.0f
#endif
static const float pos = WGSL_SPEC_CONSTANT_23;
@@ -56,8 +56,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitProgramConstVariable(out, var)) << gen.error();
- EXPECT_EQ(result(), R"(#ifndef WGSL_SPEC_CONSTANT_23
+ ASSERT_TRUE(gen.EmitProgramConstVariable(var)) << gen.error();
+ EXPECT_EQ(gen.result(), R"(#ifndef WGSL_SPEC_CONSTANT_23
#error spec constant required for constant id 23
#endif
static const float pos = WGSL_SPEC_CONSTANT_23;
@@ -76,9 +76,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitProgramConstVariable(out, a)) << gen.error();
- ASSERT_TRUE(gen.EmitProgramConstVariable(out, b)) << gen.error();
- EXPECT_EQ(result(), R"(#ifndef WGSL_SPEC_CONSTANT_0
+ ASSERT_TRUE(gen.EmitProgramConstVariable(a)) << gen.error();
+ ASSERT_TRUE(gen.EmitProgramConstVariable(b)) << gen.error();
+ EXPECT_EQ(gen.result(), R"(#ifndef WGSL_SPEC_CONSTANT_0
#define WGSL_SPEC_CONSTANT_0 3.0f
#endif
static const float a = WGSL_SPEC_CONSTANT_0;
diff --git a/src/writer/hlsl/generator_impl_return_test.cc b/src/writer/hlsl/generator_impl_return_test.cc
index 5a4bf74..177ac25 100644
--- a/src/writer/hlsl/generator_impl_return_test.cc
+++ b/src/writer/hlsl/generator_impl_return_test.cc
@@ -29,8 +29,8 @@
gen.increment_indent();
- ASSERT_TRUE(gen.EmitStatement(out, r)) << gen.error();
- EXPECT_EQ(result(), " return;\n");
+ ASSERT_TRUE(gen.EmitStatement(r)) << gen.error();
+ EXPECT_EQ(gen.result(), " return;\n");
}
TEST_F(HlslGeneratorImplTest_Return, Emit_ReturnWithValue) {
@@ -41,8 +41,8 @@
gen.increment_indent();
- ASSERT_TRUE(gen.EmitStatement(out, r)) << gen.error();
- EXPECT_EQ(result(), " return 123;\n");
+ ASSERT_TRUE(gen.EmitStatement(r)) << gen.error();
+ EXPECT_EQ(gen.result(), " return 123;\n");
}
} // namespace
diff --git a/src/writer/hlsl/generator_impl_sanitizer_test.cc b/src/writer/hlsl/generator_impl_sanitizer_test.cc
index e4e85c0..99cd52a 100644
--- a/src/writer/hlsl/generator_impl_sanitizer_test.cc
+++ b/src/writer/hlsl/generator_impl_sanitizer_test.cc
@@ -45,9 +45,9 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
+ ASSERT_TRUE(gen.Generate()) << gen.error();
- auto got = result();
+ auto got = gen.result();
auto* expect = R"(ByteAddressBuffer b : register(t1, space2);
void a_func() {
@@ -85,9 +85,9 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
+ ASSERT_TRUE(gen.Generate()) << gen.error();
- auto got = result();
+ auto got = gen.result();
auto* expect = R"(ByteAddressBuffer b : register(t1, space2);
void a_func() {
@@ -127,9 +127,9 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
+ ASSERT_TRUE(gen.Generate()) << gen.error();
- auto got = result();
+ auto got = gen.result();
auto* expect = R"(ByteAddressBuffer b : register(t1, space2);
void a_func() {
@@ -159,9 +159,9 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
+ ASSERT_TRUE(gen.Generate()) << gen.error();
- auto got = result();
+ auto got = gen.result();
auto* expect = R"(struct tint_array_wrapper {
int arr[4];
};
@@ -196,9 +196,9 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
+ ASSERT_TRUE(gen.Generate()) << gen.error();
- auto got = result();
+ auto got = gen.result();
auto* expect = R"(struct S {
int a;
float3 b;
@@ -235,9 +235,9 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
+ ASSERT_TRUE(gen.Generate()) << gen.error();
- auto got = result();
+ auto got = gen.result();
auto* expect = R"(void main() {
int v = 0;
int x = v;
@@ -278,9 +278,9 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
+ ASSERT_TRUE(gen.Generate()) << gen.error();
- auto got = result();
+ auto got = gen.result();
auto* expect = R"(void main() {
float4x4 m = float4x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
float f = m[2][1];
@@ -322,9 +322,9 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
+ ASSERT_TRUE(gen.Generate()) << gen.error();
- auto got = result();
+ auto got = gen.result();
auto* expect = R"(int x(inout int p) {
return p;
}
diff --git a/src/writer/hlsl/generator_impl_switch_test.cc b/src/writer/hlsl/generator_impl_switch_test.cc
index 6412524..c0a92cd 100644
--- a/src/writer/hlsl/generator_impl_switch_test.cc
+++ b/src/writer/hlsl/generator_impl_switch_test.cc
@@ -46,8 +46,8 @@
gen.increment_indent();
- ASSERT_TRUE(gen.EmitStatement(out, s)) << gen.error();
- EXPECT_EQ(result(), R"( switch(cond) {
+ ASSERT_TRUE(gen.EmitStatement(s)) << gen.error();
+ EXPECT_EQ(gen.result(), R"( switch(cond) {
case 5: {
break;
}
diff --git a/src/writer/hlsl/generator_impl_test.cc b/src/writer/hlsl/generator_impl_test.cc
index adf73e4..30cf659 100644
--- a/src/writer/hlsl/generator_impl_test.cc
+++ b/src/writer/hlsl/generator_impl_test.cc
@@ -24,7 +24,7 @@
TEST_F(HlslGeneratorImplTest, ErrorIfSanitizerNotRun) {
auto program = std::make_unique<Program>(std::move(*this));
GeneratorImpl gen(program.get());
- EXPECT_FALSE(gen.Generate(out));
+ EXPECT_FALSE(gen.Generate());
EXPECT_EQ(
gen.error(),
"error: HLSL writer requires the transform::Hlsl sanitizer to have been "
@@ -37,8 +37,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_EQ(result(), R"(void my_func() {
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(void my_func() {
}
)");
}
diff --git a/src/writer/hlsl/generator_impl_type_test.cc b/src/writer/hlsl/generator_impl_type_test.cc
index 98709a0..feedf4e 100644
--- a/src/writer/hlsl/generator_impl_type_test.cc
+++ b/src/writer/hlsl/generator_impl_type_test.cc
@@ -37,10 +37,11 @@
GeneratorImpl& gen = Build();
+ std::stringstream out;
ASSERT_TRUE(gen.EmitType(out, program->TypeOf(arr), ast::StorageClass::kNone,
ast::Access::kReadWrite, "ary"))
<< gen.error();
- EXPECT_EQ(result(), "bool ary[4]");
+ EXPECT_EQ(out.str(), "bool ary[4]");
}
TEST_F(HlslGeneratorImplTest_Type, EmitType_ArrayOfArray) {
@@ -49,10 +50,11 @@
GeneratorImpl& gen = Build();
+ std::stringstream out;
ASSERT_TRUE(gen.EmitType(out, program->TypeOf(arr), ast::StorageClass::kNone,
ast::Access::kReadWrite, "ary"))
<< gen.error();
- EXPECT_EQ(result(), "bool ary[5][4]");
+ EXPECT_EQ(out.str(), "bool ary[5][4]");
}
// TODO(dsinclair): Is this possible? What order should it output in?
@@ -63,10 +65,11 @@
GeneratorImpl& gen = Build();
+ std::stringstream out;
ASSERT_TRUE(gen.EmitType(out, program->TypeOf(arr), ast::StorageClass::kNone,
ast::Access::kReadWrite, "ary"))
<< gen.error();
- EXPECT_EQ(result(), "bool ary[5][4][1]");
+ EXPECT_EQ(out.str(), "bool ary[5][4][1]");
}
TEST_F(HlslGeneratorImplTest_Type, EmitType_ArrayOfArrayOfArray) {
@@ -75,10 +78,11 @@
GeneratorImpl& gen = Build();
+ std::stringstream out;
ASSERT_TRUE(gen.EmitType(out, program->TypeOf(arr), ast::StorageClass::kNone,
ast::Access::kReadWrite, "ary"))
<< gen.error();
- EXPECT_EQ(result(), "bool ary[6][5][4]");
+ EXPECT_EQ(out.str(), "bool ary[6][5][4]");
}
TEST_F(HlslGeneratorImplTest_Type, EmitType_Array_WithoutName) {
@@ -87,10 +91,11 @@
GeneratorImpl& gen = Build();
+ std::stringstream out;
ASSERT_TRUE(gen.EmitType(out, program->TypeOf(arr), ast::StorageClass::kNone,
ast::Access::kReadWrite, ""))
<< gen.error();
- EXPECT_EQ(result(), "bool[4]");
+ EXPECT_EQ(out.str(), "bool[4]");
}
TEST_F(HlslGeneratorImplTest_Type, EmitType_Bool) {
@@ -98,10 +103,11 @@
GeneratorImpl& gen = Build();
+ std::stringstream out;
ASSERT_TRUE(gen.EmitType(out, bool_, ast::StorageClass::kNone,
ast::Access::kReadWrite, ""))
<< gen.error();
- EXPECT_EQ(result(), "bool");
+ EXPECT_EQ(out.str(), "bool");
}
TEST_F(HlslGeneratorImplTest_Type, EmitType_F32) {
@@ -109,10 +115,11 @@
GeneratorImpl& gen = Build();
+ std::stringstream out;
ASSERT_TRUE(gen.EmitType(out, f32, ast::StorageClass::kNone,
ast::Access::kReadWrite, ""))
<< gen.error();
- EXPECT_EQ(result(), "float");
+ EXPECT_EQ(out.str(), "float");
}
TEST_F(HlslGeneratorImplTest_Type, EmitType_I32) {
@@ -120,10 +127,11 @@
GeneratorImpl& gen = Build();
+ std::stringstream out;
ASSERT_TRUE(gen.EmitType(out, i32, ast::StorageClass::kNone,
ast::Access::kReadWrite, ""))
<< gen.error();
- EXPECT_EQ(result(), "int");
+ EXPECT_EQ(out.str(), "int");
}
TEST_F(HlslGeneratorImplTest_Type, EmitType_Matrix) {
@@ -133,10 +141,11 @@
GeneratorImpl& gen = Build();
+ std::stringstream out;
ASSERT_TRUE(gen.EmitType(out, mat2x3, ast::StorageClass::kNone,
ast::Access::kReadWrite, ""))
<< gen.error();
- EXPECT_EQ(result(), "float2x3");
+ EXPECT_EQ(out.str(), "float2x3");
}
// TODO(dsinclair): How to annotate as workgroup?
@@ -147,10 +156,11 @@
GeneratorImpl& gen = Build();
+ std::stringstream out;
ASSERT_TRUE(gen.EmitType(out, p, ast::StorageClass::kNone,
ast::Access::kReadWrite, ""))
<< gen.error();
- EXPECT_EQ(result(), "float*");
+ EXPECT_EQ(out.str(), "float*");
}
TEST_F(HlslGeneratorImplTest_Type, EmitType_StructDecl) {
@@ -163,8 +173,8 @@
GeneratorImpl& gen = Build();
auto* sem_s = program->TypeOf(s)->As<sem::Struct>();
- ASSERT_TRUE(gen.EmitStructType(out, sem_s)) << gen.error();
- EXPECT_EQ(result(), R"(struct S {
+ ASSERT_TRUE(gen.EmitStructType(sem_s)) << gen.error();
+ EXPECT_EQ(gen.result(), R"(struct S {
int a;
float b;
};
@@ -187,8 +197,8 @@
GeneratorImpl& gen = Build();
auto* sem_s = program->TypeOf(s)->As<sem::Struct>();
- ASSERT_TRUE(gen.EmitStructType(out, sem_s)) << gen.error();
- EXPECT_EQ(result(), "");
+ ASSERT_TRUE(gen.EmitStructType(sem_s)) << gen.error();
+ EXPECT_EQ(gen.result(), "");
}
TEST_F(HlslGeneratorImplTest_Type, EmitType_Struct) {
@@ -201,10 +211,11 @@
GeneratorImpl& gen = Build();
auto* sem_s = program->TypeOf(s)->As<sem::Struct>();
+ std::stringstream out;
ASSERT_TRUE(gen.EmitType(out, sem_s, ast::StorageClass::kNone,
ast::Access::kReadWrite, ""))
<< gen.error();
- EXPECT_EQ(result(), "S");
+ EXPECT_EQ(out.str(), "S");
}
TEST_F(HlslGeneratorImplTest_Type, EmitType_Struct_NameCollision) {
@@ -216,8 +227,8 @@
GeneratorImpl& gen = SanitizeAndBuild();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr(R"(struct S {
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr(R"(struct S {
int tint_symbol;
float tint_symbol_1;
};
@@ -229,10 +240,11 @@
GeneratorImpl& gen = Build();
+ std::stringstream out;
ASSERT_TRUE(gen.EmitType(out, u32, ast::StorageClass::kNone,
ast::Access::kReadWrite, ""))
<< gen.error();
- EXPECT_EQ(result(), "uint");
+ EXPECT_EQ(out.str(), "uint");
}
TEST_F(HlslGeneratorImplTest_Type, EmitType_Vector) {
@@ -241,10 +253,11 @@
GeneratorImpl& gen = Build();
+ std::stringstream out;
ASSERT_TRUE(gen.EmitType(out, vec3, ast::StorageClass::kNone,
ast::Access::kReadWrite, ""))
<< gen.error();
- EXPECT_EQ(result(), "float3");
+ EXPECT_EQ(out.str(), "float3");
}
TEST_F(HlslGeneratorImplTest_Type, EmitType_Void) {
@@ -252,10 +265,11 @@
GeneratorImpl& gen = Build();
+ std::stringstream out;
ASSERT_TRUE(gen.EmitType(out, void_, ast::StorageClass::kNone,
ast::Access::kReadWrite, ""))
<< gen.error();
- EXPECT_EQ(result(), "void");
+ EXPECT_EQ(out.str(), "void");
}
TEST_F(HlslGeneratorImplTest_Type, EmitSampler) {
@@ -263,10 +277,11 @@
GeneratorImpl& gen = Build();
+ std::stringstream out;
ASSERT_TRUE(gen.EmitType(out, sampler, ast::StorageClass::kNone,
ast::Access::kReadWrite, ""))
<< gen.error();
- EXPECT_EQ(result(), "SamplerState");
+ EXPECT_EQ(out.str(), "SamplerState");
}
TEST_F(HlslGeneratorImplTest_Type, EmitSamplerComparison) {
@@ -274,10 +289,11 @@
GeneratorImpl& gen = Build();
+ std::stringstream out;
ASSERT_TRUE(gen.EmitType(out, sampler, ast::StorageClass::kNone,
ast::Access::kReadWrite, ""))
<< gen.error();
- EXPECT_EQ(result(), "SamplerComparisonState");
+ EXPECT_EQ(out.str(), "SamplerComparisonState");
}
struct HlslDepthTextureData {
@@ -305,8 +321,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr(params.result));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr(params.result));
}
INSTANTIATE_TEST_SUITE_P(
HlslGeneratorImplTest_Type,
@@ -361,8 +377,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr(params.result));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr(params.result));
}
INSTANTIATE_TEST_SUITE_P(
HlslGeneratorImplTest_Type,
@@ -465,10 +481,11 @@
GeneratorImpl& gen = Build();
+ std::stringstream out;
ASSERT_TRUE(gen.EmitType(out, s, ast::StorageClass::kNone,
ast::Access::kReadWrite, ""))
<< gen.error();
- EXPECT_EQ(result(), "Texture2DMS<float4>");
+ EXPECT_EQ(out.str(), "Texture2DMS<float4>");
}
struct HlslStorageTextureData {
@@ -501,8 +518,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr(params.result));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr(params.result));
}
INSTANTIATE_TEST_SUITE_P(
HlslGeneratorImplTest_Type,
diff --git a/src/writer/hlsl/generator_impl_unary_op_test.cc b/src/writer/hlsl/generator_impl_unary_op_test.cc
index 83ead64..6674692 100644
--- a/src/writer/hlsl/generator_impl_unary_op_test.cc
+++ b/src/writer/hlsl/generator_impl_unary_op_test.cc
@@ -29,8 +29,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitExpression(pre, out, op)) << gen.error();
- EXPECT_EQ(result(), "expr");
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, op)) << gen.error();
+ EXPECT_EQ(out.str(), "expr");
}
TEST_F(HlslUnaryOpTest, Complement) {
@@ -41,8 +42,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitExpression(pre, out, op)) << gen.error();
- EXPECT_EQ(result(), "~(expr)");
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, op)) << gen.error();
+ EXPECT_EQ(out.str(), "~(expr)");
}
TEST_F(HlslUnaryOpTest, Indirection) {
@@ -56,8 +58,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitExpression(pre, out, op)) << gen.error();
- EXPECT_EQ(result(), "expr");
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, op)) << gen.error();
+ EXPECT_EQ(out.str(), "expr");
}
TEST_F(HlslUnaryOpTest, Not) {
@@ -67,8 +70,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitExpression(pre, out, op)) << gen.error();
- EXPECT_EQ(result(), "!(expr)");
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, op)) << gen.error();
+ EXPECT_EQ(out.str(), "!(expr)");
}
TEST_F(HlslUnaryOpTest, Negation) {
@@ -79,8 +83,9 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitExpression(pre, out, op)) << gen.error();
- EXPECT_EQ(result(), "-(expr)");
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitExpression(out, op)) << gen.error();
+ EXPECT_EQ(out.str(), "-(expr)");
}
} // namespace
} // namespace hlsl
diff --git a/src/writer/hlsl/generator_impl_variable_decl_statement_test.cc b/src/writer/hlsl/generator_impl_variable_decl_statement_test.cc
index 2ce7469..5fa6059 100644
--- a/src/writer/hlsl/generator_impl_variable_decl_statement_test.cc
+++ b/src/writer/hlsl/generator_impl_variable_decl_statement_test.cc
@@ -34,8 +34,8 @@
gen.increment_indent();
- ASSERT_TRUE(gen.EmitStatement(out, stmt)) << gen.error();
- EXPECT_EQ(result(), " float a = 0.0f;\n");
+ ASSERT_TRUE(gen.EmitStatement(stmt)) << gen.error();
+ EXPECT_EQ(gen.result(), " float a = 0.0f;\n");
}
TEST_F(HlslGeneratorImplTest_VariableDecl, Emit_VariableDeclStatement_Const) {
@@ -47,8 +47,8 @@
gen.increment_indent();
- ASSERT_TRUE(gen.EmitStatement(out, stmt)) << gen.error();
- EXPECT_EQ(result(), " const float a = 0.0f;\n");
+ ASSERT_TRUE(gen.EmitStatement(stmt)) << gen.error();
+ EXPECT_EQ(gen.result(), " const float a = 0.0f;\n");
}
TEST_F(HlslGeneratorImplTest_VariableDecl, Emit_VariableDeclStatement_Array) {
@@ -60,8 +60,8 @@
gen.increment_indent();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(),
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(),
HasSubstr(" float a[5] = {0.0f, 0.0f, 0.0f, 0.0f, 0.0f};\n"));
}
@@ -74,8 +74,8 @@
gen.increment_indent();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr(" static float a = 0.0f;\n"));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr(" static float a = 0.0f;\n"));
}
TEST_F(HlslGeneratorImplTest_VariableDecl,
@@ -87,8 +87,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr(R"(float a = initializer;
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr(R"(float a = initializer;
)"));
}
@@ -101,8 +101,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitStatement(out, stmt)) << gen.error();
- EXPECT_EQ(result(), R"(float3 a = float3(0.0f, 0.0f, 0.0f);
+ ASSERT_TRUE(gen.EmitStatement(stmt)) << gen.error();
+ EXPECT_EQ(gen.result(), R"(float3 a = float3(0.0f, 0.0f, 0.0f);
)");
}
@@ -116,8 +116,8 @@
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.EmitStatement(out, stmt)) << gen.error();
- EXPECT_EQ(result(),
+ ASSERT_TRUE(gen.EmitStatement(stmt)) << gen.error();
+ EXPECT_EQ(gen.result(),
R"(float2x3 a = float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
)");
}
diff --git a/src/writer/hlsl/generator_impl_workgroup_var_test.cc b/src/writer/hlsl/generator_impl_workgroup_var_test.cc
index a965be7..626fc90 100644
--- a/src/writer/hlsl/generator_impl_workgroup_var_test.cc
+++ b/src/writer/hlsl/generator_impl_workgroup_var_test.cc
@@ -32,8 +32,8 @@
{Stage(ast::PipelineStage::kCompute)});
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr("groupshared float wg;\n"));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("groupshared float wg;\n"));
}
TEST_F(HlslGeneratorImplTest_WorkgroupVar, Aliased) {
@@ -45,8 +45,8 @@
{Stage(ast::PipelineStage::kCompute)});
GeneratorImpl& gen = Build();
- ASSERT_TRUE(gen.Generate(out)) << gen.error();
- EXPECT_THAT(result(), HasSubstr("groupshared float wg;\n"));
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr("groupshared float wg;\n"));
}
} // namespace
diff --git a/src/writer/hlsl/test_helper.h b/src/writer/hlsl/test_helper.h
index 26c5563..dc9420f 100644
--- a/src/writer/hlsl/test_helper.h
+++ b/src/writer/hlsl/test_helper.h
@@ -96,16 +96,6 @@
return *gen_;
}
- /// @returns the result string
- std::string result() const { return out.str(); }
-
- /// @returns the pre result string
- std::string pre_result() const { return pre.str(); }
-
- /// The output stream
- std::ostringstream out;
- /// The pre-output stream
- std::ostringstream pre;
/// The program built with a call to Build()
std::unique_ptr<Program> program;
diff --git a/src/writer/text_generator.cc b/src/writer/text_generator.cc
index 2a26d13..624d72c 100644
--- a/src/writer/text_generator.cc
+++ b/src/writer/text_generator.cc
@@ -31,5 +31,38 @@
}
}
+TextGenerator::LineWriter::LineWriter(TextGenerator* generator)
+ : gen(generator) {}
+
+TextGenerator::LineWriter::LineWriter(LineWriter&& other) {
+ gen = other.gen;
+ other.gen = nullptr;
+}
+
+TextGenerator::LineWriter::~LineWriter() {
+ if (gen) {
+ auto str = os.str();
+ if (!str.empty()) {
+ gen->make_indent();
+ gen->out_ << str << std::endl;
+ }
+ }
+}
+
+TextGenerator::ScopedParen::ScopedParen(std::ostream& stream) : s(stream) {
+ s << "(";
+}
+TextGenerator::ScopedParen::~ScopedParen() {
+ s << ")";
+}
+
+TextGenerator::ScopedIndent::ScopedIndent(TextGenerator* generator)
+ : gen(generator) {
+ gen->increment_indent();
+}
+TextGenerator::ScopedIndent::~ScopedIndent() {
+ gen->decrement_indent();
+}
+
} // namespace writer
} // namespace tint
diff --git a/src/writer/text_generator.h b/src/writer/text_generator.h
index ab6d47a..8f68385 100644
--- a/src/writer/text_generator.h
+++ b/src/writer/text_generator.h
@@ -17,6 +17,7 @@
#include <sstream>
#include <string>
+#include <utility>
#include "src/diagnostic/diagnostic.h"
@@ -58,6 +59,72 @@
std::string error() const { return diagnostics_.str(); }
protected:
+ /// LineWriter is a helper that acts as a string buffer, who's content is
+ /// emitted to the TextGenerator as a single line on destruction.
+ struct LineWriter {
+ public:
+ /// Constructor
+ /// @param generator the TextGenerator that the LineWriter will append its
+ /// content to on destruction
+ explicit LineWriter(TextGenerator* generator);
+ /// Move constructor
+ /// @param rhs the LineWriter to move
+ LineWriter(LineWriter&& rhs);
+ /// Destructor
+ ~LineWriter();
+
+ /// @returns the ostringstream
+ operator std::ostream &() { return os; }
+
+ /// @param rhs the value to write to the line
+ /// @returns the ostream so calls can be chained
+ template <typename T>
+ std::ostream& operator<<(T&& rhs) {
+ return os << std::forward<T>(rhs);
+ }
+
+ private:
+ LineWriter(const LineWriter&) = delete;
+ LineWriter& operator=(const LineWriter&) = delete;
+
+ std::ostringstream os;
+ TextGenerator* gen;
+ };
+
+ /// Helper for writing a '(' on construction and a ')' destruction.
+ struct ScopedParen {
+ /// Constructor
+ /// @param stream the std::ostream that will be written to
+ explicit ScopedParen(std::ostream& stream);
+ /// Destructor
+ ~ScopedParen();
+
+ private:
+ ScopedParen(ScopedParen&& rhs) = delete;
+ ScopedParen(const ScopedParen&) = delete;
+ ScopedParen& operator=(const ScopedParen&) = delete;
+ std::ostream& s;
+ };
+
+ /// Helper for incrementing indentation on construction and decrementing
+ /// indentation on destruction.
+ struct ScopedIndent {
+ /// Constructor
+ /// @param generator the TextGenerator that the ScopedIndent will indent
+ explicit ScopedIndent(TextGenerator* generator);
+ /// Destructor
+ ~ScopedIndent();
+
+ private:
+ ScopedIndent(ScopedIndent&& rhs) = delete;
+ ScopedIndent(const ScopedIndent&) = delete;
+ ScopedIndent& operator=(const ScopedIndent&) = delete;
+ TextGenerator* gen;
+ };
+
+ /// @returns a new LineWriter, used for buffering and writing a line to out_
+ LineWriter line() { return LineWriter(this); }
+
/// The text output stream
std::ostringstream out_;
/// Diagnostics generated by the generator
diff --git a/test/bug/tint/913.wgsl.expected.hlsl b/test/bug/tint/913.wgsl.expected.hlsl
index ada977d..1ad9ee7 100644
--- a/test/bug/tint/913.wgsl.expected.hlsl
+++ b/test/bug/tint/913.wgsl.expected.hlsl
@@ -26,29 +26,29 @@
const float4 nonCoveredColor = float4(0.0f, 1.0f, 0.0f, 1.0f);
bool success = true;
const int scalar_offset = (16u) / 4;
- bool tint_tmp_2 = (dstTexCoord.x < uniforms[scalar_offset / 4][scalar_offset % 4]);
- if (!tint_tmp_2) {
-const int scalar_offset_1 = (20u) / 4;
- tint_tmp_2 = (dstTexCoord.y < uniforms[scalar_offset_1 / 4][scalar_offset_1 % 4]);
- }
- bool tint_tmp_3 = (tint_tmp_2);
- if (!tint_tmp_3) {
-const int scalar_offset_2 = (16u) / 4;
- const int scalar_offset_3 = (24u) / 4;
- tint_tmp_3 = (dstTexCoord.x >= (uniforms[scalar_offset_2 / 4][scalar_offset_2 % 4] + uniforms[scalar_offset_3 / 4][scalar_offset_3 % 4]));
- }
- bool tint_tmp_4 = (tint_tmp_3);
+ bool tint_tmp_4 = (dstTexCoord.x < uniforms[scalar_offset / 4][scalar_offset % 4]);
if (!tint_tmp_4) {
-const int scalar_offset_4 = (20u) / 4;
- const int scalar_offset_5 = (28u) / 4;
- tint_tmp_4 = (dstTexCoord.y >= (uniforms[scalar_offset_4 / 4][scalar_offset_4 % 4] + uniforms[scalar_offset_5 / 4][scalar_offset_5 % 4]));
+ const int scalar_offset_1 = (20u) / 4;
+ tint_tmp_4 = (dstTexCoord.y < uniforms[scalar_offset_1 / 4][scalar_offset_1 % 4]);
}
-if ((tint_tmp_4)) {
- bool tint_tmp_5 = success;
+ bool tint_tmp_3 = (tint_tmp_4);
+ if (!tint_tmp_3) {
+ const int scalar_offset_2 = (16u) / 4;
+ const int scalar_offset_3 = (24u) / 4;
+ tint_tmp_3 = (dstTexCoord.x >= (uniforms[scalar_offset_2 / 4][scalar_offset_2 % 4] + uniforms[scalar_offset_3 / 4][scalar_offset_3 % 4]));
+ }
+ bool tint_tmp_2 = (tint_tmp_3);
+ if (!tint_tmp_2) {
+ const int scalar_offset_4 = (20u) / 4;
+ const int scalar_offset_5 = (28u) / 4;
+ tint_tmp_2 = (dstTexCoord.y >= (uniforms[scalar_offset_4 / 4][scalar_offset_4 % 4] + uniforms[scalar_offset_5 / 4][scalar_offset_5 % 4]));
+ }
+ if ((tint_tmp_2)) {
+ bool tint_tmp_5 = success;
if (tint_tmp_5) {
tint_tmp_5 = all((tint_symbol.Load(int3(dstTexCoord, 0), 0) == nonCoveredColor));
}
-success = (tint_tmp_5);
+ success = (tint_tmp_5);
} else {
const int scalar_offset_6 = (16u) / 4;
uint4 ubo_load = uniforms[scalar_offset_6 / 4];
@@ -63,33 +63,33 @@
const float4 dstColor = tint_symbol.Load(int3(dstTexCoord, 0), 0);
const int scalar_offset_9 = (4u) / 4;
if ((uniforms[scalar_offset_9 / 4][scalar_offset_9 % 4] == 2u)) {
- bool tint_tmp_6 = success;
- if (tint_tmp_6) {
- tint_tmp_6 = aboutEqual(dstColor.r, srcColor.r);
- }
- bool tint_tmp_7 = (tint_tmp_6);
+ bool tint_tmp_7 = success;
if (tint_tmp_7) {
- tint_tmp_7 = aboutEqual(dstColor.g, srcColor.g);
+ tint_tmp_7 = aboutEqual(dstColor.r, srcColor.r);
}
-success = (tint_tmp_7);
+ bool tint_tmp_6 = (tint_tmp_7);
+ if (tint_tmp_6) {
+ tint_tmp_6 = aboutEqual(dstColor.g, srcColor.g);
+ }
+ success = (tint_tmp_6);
} else {
- bool tint_tmp_8 = success;
- if (tint_tmp_8) {
- tint_tmp_8 = aboutEqual(dstColor.r, srcColor.r);
- }
- bool tint_tmp_9 = (tint_tmp_8);
- if (tint_tmp_9) {
- tint_tmp_9 = aboutEqual(dstColor.g, srcColor.g);
- }
- bool tint_tmp_10 = (tint_tmp_9);
- if (tint_tmp_10) {
- tint_tmp_10 = aboutEqual(dstColor.b, srcColor.b);
- }
- bool tint_tmp_11 = (tint_tmp_10);
+ bool tint_tmp_11 = success;
if (tint_tmp_11) {
- tint_tmp_11 = aboutEqual(dstColor.a, srcColor.a);
+ tint_tmp_11 = aboutEqual(dstColor.r, srcColor.r);
}
-success = (tint_tmp_11);
+ bool tint_tmp_10 = (tint_tmp_11);
+ if (tint_tmp_10) {
+ tint_tmp_10 = aboutEqual(dstColor.g, srcColor.g);
+ }
+ bool tint_tmp_9 = (tint_tmp_10);
+ if (tint_tmp_9) {
+ tint_tmp_9 = aboutEqual(dstColor.b, srcColor.b);
+ }
+ bool tint_tmp_8 = (tint_tmp_9);
+ if (tint_tmp_8) {
+ tint_tmp_8 = aboutEqual(dstColor.a, srcColor.a);
+ }
+ success = (tint_tmp_8);
}
}
const uint outputIndex = ((GlobalInvocationID.y * uint(dstSize.x)) + GlobalInvocationID.x);
diff --git a/test/bug/tint/914.wgsl.expected.hlsl b/test/bug/tint/914.wgsl.expected.hlsl
index 3a6738e..33cd2bb 100644
--- a/test/bug/tint/914.wgsl.expected.hlsl
+++ b/test/bug/tint/914.wgsl.expected.hlsl
@@ -7,12 +7,12 @@
float mm_readA(uint row, uint col) {
const int scalar_offset = (0u) / 4;
- bool tint_tmp = (row < uniforms[scalar_offset / 4][scalar_offset % 4]);
+ bool tint_tmp = (row < uniforms[scalar_offset / 4][scalar_offset % 4]);
if (tint_tmp) {
-const int scalar_offset_1 = (4u) / 4;
- tint_tmp = (col < uniforms[scalar_offset_1 / 4][scalar_offset_1 % 4]);
+ const int scalar_offset_1 = (4u) / 4;
+ tint_tmp = (col < uniforms[scalar_offset_1 / 4][scalar_offset_1 % 4]);
}
-if ((tint_tmp)) {
+ if ((tint_tmp)) {
const int scalar_offset_2 = (4u) / 4;
const float result = asfloat(firstMatrix.Load((4u * ((row * uniforms[scalar_offset_2 / 4][scalar_offset_2 % 4]) + col))));
return result;
@@ -22,12 +22,12 @@
float mm_readB(uint row, uint col) {
const int scalar_offset_3 = (4u) / 4;
- bool tint_tmp_1 = (row < uniforms[scalar_offset_3 / 4][scalar_offset_3 % 4]);
+ bool tint_tmp_1 = (row < uniforms[scalar_offset_3 / 4][scalar_offset_3 % 4]);
if (tint_tmp_1) {
-const int scalar_offset_4 = (8u) / 4;
- tint_tmp_1 = (col < uniforms[scalar_offset_4 / 4][scalar_offset_4 % 4]);
+ const int scalar_offset_4 = (8u) / 4;
+ tint_tmp_1 = (col < uniforms[scalar_offset_4 / 4][scalar_offset_4 % 4]);
}
-if ((tint_tmp_1)) {
+ if ((tint_tmp_1)) {
const int scalar_offset_5 = (8u) / 4;
const float result = asfloat(secondMatrix.Load((4u * ((row * uniforms[scalar_offset_5 / 4][scalar_offset_5 % 4]) + col))));
return result;
@@ -37,12 +37,12 @@
void mm_write(uint row, uint col, float value) {
const int scalar_offset_6 = (0u) / 4;
- bool tint_tmp_2 = (row < uniforms[scalar_offset_6 / 4][scalar_offset_6 % 4]);
+ bool tint_tmp_2 = (row < uniforms[scalar_offset_6 / 4][scalar_offset_6 % 4]);
if (tint_tmp_2) {
-const int scalar_offset_7 = (8u) / 4;
- tint_tmp_2 = (col < uniforms[scalar_offset_7 / 4][scalar_offset_7 % 4]);
+ const int scalar_offset_7 = (8u) / 4;
+ tint_tmp_2 = (col < uniforms[scalar_offset_7 / 4][scalar_offset_7 % 4]);
}
-if ((tint_tmp_2)) {
+ if ((tint_tmp_2)) {
const int scalar_offset_8 = (8u) / 4;
const uint index = (col + (row * uniforms[scalar_offset_8 / 4][scalar_offset_8 % 4]));
resultMatrix.Store((4u * index), asuint(value));
diff --git a/test/intrinsics/gen/pack2x16float/0e97b3.wgsl.expected.hlsl b/test/intrinsics/gen/pack2x16float/0e97b3.wgsl.expected.hlsl
index 9132660..c06e032 100644
--- a/test/intrinsics/gen/pack2x16float/0e97b3.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/pack2x16float/0e97b3.wgsl.expected.hlsl
@@ -1,6 +1,6 @@
void pack2x16float_0e97b3() {
uint2 tint_tmp = f32tof16(float2(0.0f, 0.0f));
-uint res = (tint_tmp.x | tint_tmp.y << 16);
+ uint res = (tint_tmp.x | tint_tmp.y << 16);
}
struct tint_symbol {
diff --git a/test/intrinsics/gen/pack2x16snorm/6c169b.wgsl.expected.hlsl b/test/intrinsics/gen/pack2x16snorm/6c169b.wgsl.expected.hlsl
index c2d596b..cccbc2d 100644
--- a/test/intrinsics/gen/pack2x16snorm/6c169b.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/pack2x16snorm/6c169b.wgsl.expected.hlsl
@@ -1,6 +1,6 @@
void pack2x16snorm_6c169b() {
int2 tint_tmp = int2(round(clamp(float2(0.0f, 0.0f), -1.0, 1.0) * 32767.0)) & 0xffff;
-uint res = asuint(tint_tmp.x | tint_tmp.y << 16);
+ uint res = asuint(tint_tmp.x | tint_tmp.y << 16);
}
struct tint_symbol {
diff --git a/test/intrinsics/gen/pack2x16unorm/0f08e4.wgsl.expected.hlsl b/test/intrinsics/gen/pack2x16unorm/0f08e4.wgsl.expected.hlsl
index cf0dbe7..3495457 100644
--- a/test/intrinsics/gen/pack2x16unorm/0f08e4.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/pack2x16unorm/0f08e4.wgsl.expected.hlsl
@@ -1,6 +1,6 @@
void pack2x16unorm_0f08e4() {
uint2 tint_tmp = uint2(round(clamp(float2(0.0f, 0.0f), 0.0, 1.0) * 65535.0));
-uint res = (tint_tmp.x | tint_tmp.y << 16);
+ uint res = (tint_tmp.x | tint_tmp.y << 16);
}
struct tint_symbol {
diff --git a/test/intrinsics/gen/pack4x8snorm/4d22e7.wgsl.expected.hlsl b/test/intrinsics/gen/pack4x8snorm/4d22e7.wgsl.expected.hlsl
index d495e8d..93c5fcf 100644
--- a/test/intrinsics/gen/pack4x8snorm/4d22e7.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/pack4x8snorm/4d22e7.wgsl.expected.hlsl
@@ -1,6 +1,6 @@
void pack4x8snorm_4d22e7() {
int4 tint_tmp = int4(round(clamp(float4(0.0f, 0.0f, 0.0f, 0.0f), -1.0, 1.0) * 127.0)) & 0xff;
-uint res = asuint(tint_tmp.x | tint_tmp.y << 8 | tint_tmp.z << 16 | tint_tmp.w << 24);
+ uint res = asuint(tint_tmp.x | tint_tmp.y << 8 | tint_tmp.z << 16 | tint_tmp.w << 24);
}
struct tint_symbol {
diff --git a/test/intrinsics/gen/pack4x8unorm/95c456.wgsl.expected.hlsl b/test/intrinsics/gen/pack4x8unorm/95c456.wgsl.expected.hlsl
index 6fca709..f22557c 100644
--- a/test/intrinsics/gen/pack4x8unorm/95c456.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/pack4x8unorm/95c456.wgsl.expected.hlsl
@@ -1,6 +1,6 @@
void pack4x8unorm_95c456() {
uint4 tint_tmp = uint4(round(clamp(float4(0.0f, 0.0f, 0.0f, 0.0f), 0.0, 1.0) * 255.0));
-uint res = (tint_tmp.x | tint_tmp.y << 8 | tint_tmp.z << 16 | tint_tmp.w << 24);
+ uint res = (tint_tmp.x | tint_tmp.y << 8 | tint_tmp.z << 16 | tint_tmp.w << 24);
}
struct tint_symbol {
diff --git a/test/intrinsics/gen/unpack2x16float/32a5cf.wgsl.expected.hlsl b/test/intrinsics/gen/unpack2x16float/32a5cf.wgsl.expected.hlsl
index 6dc4bda..7107517 100644
--- a/test/intrinsics/gen/unpack2x16float/32a5cf.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/unpack2x16float/32a5cf.wgsl.expected.hlsl
@@ -1,6 +1,6 @@
void unpack2x16float_32a5cf() {
uint tint_tmp = 1u;
-float2 res = f16tof32(uint2(tint_tmp & 0xffff, tint_tmp >> 16));
+ float2 res = f16tof32(uint2(tint_tmp & 0xffff, tint_tmp >> 16));
}
struct tint_symbol {
diff --git a/test/intrinsics/gen/unpack2x16snorm/b4aea6.wgsl.expected.hlsl b/test/intrinsics/gen/unpack2x16snorm/b4aea6.wgsl.expected.hlsl
index d5314f5..ad0d7b3 100644
--- a/test/intrinsics/gen/unpack2x16snorm/b4aea6.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/unpack2x16snorm/b4aea6.wgsl.expected.hlsl
@@ -1,7 +1,7 @@
void unpack2x16snorm_b4aea6() {
int tint_tmp_1 = int(1u);
-int2 tint_tmp = int2(tint_tmp_1 << 16, tint_tmp_1) >> 16;
-float2 res = clamp(float2(tint_tmp) / 32767.0, -1.0, 1.0);
+ int2 tint_tmp = int2(tint_tmp_1 << 16, tint_tmp_1) >> 16;
+ float2 res = clamp(float2(tint_tmp) / 32767.0, -1.0, 1.0);
}
struct tint_symbol {
diff --git a/test/intrinsics/gen/unpack2x16unorm/7699c0.wgsl.expected.hlsl b/test/intrinsics/gen/unpack2x16unorm/7699c0.wgsl.expected.hlsl
index 3fdceee..1afbba3 100644
--- a/test/intrinsics/gen/unpack2x16unorm/7699c0.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/unpack2x16unorm/7699c0.wgsl.expected.hlsl
@@ -1,7 +1,7 @@
void unpack2x16unorm_7699c0() {
uint tint_tmp_1 = 1u;
-uint2 tint_tmp = uint2(tint_tmp_1 & 0xffff, tint_tmp_1 >> 16);
-float2 res = float2(tint_tmp) / 65535.0;
+ uint2 tint_tmp = uint2(tint_tmp_1 & 0xffff, tint_tmp_1 >> 16);
+ float2 res = float2(tint_tmp) / 65535.0;
}
struct tint_symbol {
diff --git a/test/intrinsics/gen/unpack4x8snorm/523fb3.wgsl.expected.hlsl b/test/intrinsics/gen/unpack4x8snorm/523fb3.wgsl.expected.hlsl
index e02cf52..721a798 100644
--- a/test/intrinsics/gen/unpack4x8snorm/523fb3.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/unpack4x8snorm/523fb3.wgsl.expected.hlsl
@@ -1,7 +1,7 @@
void unpack4x8snorm_523fb3() {
int tint_tmp_1 = int(1u);
-int4 tint_tmp = int4(tint_tmp_1 << 24, tint_tmp_1 << 16, tint_tmp_1 << 8, tint_tmp_1) >> 24;
-float4 res = clamp(float4(tint_tmp) / 127.0, -1.0, 1.0);
+ int4 tint_tmp = int4(tint_tmp_1 << 24, tint_tmp_1 << 16, tint_tmp_1 << 8, tint_tmp_1) >> 24;
+ float4 res = clamp(float4(tint_tmp) / 127.0, -1.0, 1.0);
}
struct tint_symbol {
diff --git a/test/intrinsics/gen/unpack4x8unorm/750c74.wgsl.expected.hlsl b/test/intrinsics/gen/unpack4x8unorm/750c74.wgsl.expected.hlsl
index af35ee2..e3a10fe 100644
--- a/test/intrinsics/gen/unpack4x8unorm/750c74.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/unpack4x8unorm/750c74.wgsl.expected.hlsl
@@ -1,7 +1,7 @@
void unpack4x8unorm_750c74() {
uint tint_tmp_1 = 1u;
-uint4 tint_tmp = uint4(tint_tmp_1 & 0xff, (tint_tmp_1 >> 8) & 0xff, (tint_tmp_1 >> 16) & 0xff, tint_tmp_1 >> 24);
-float4 res = float4(tint_tmp) / 255.0;
+ uint4 tint_tmp = uint4(tint_tmp_1 & 0xff, (tint_tmp_1 >> 8) & 0xff, (tint_tmp_1 >> 16) & 0xff, tint_tmp_1 >> 24);
+ float4 res = float4(tint_tmp) / 255.0;
}
struct tint_symbol {
