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 {