diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index 9e1736f07a..b6158abcae 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -16,7 +16,6 @@
 
 #include <algorithm>
 #include <iomanip>
-#include <iosfwd>
 #include <set>
 #include <utility>
 #include <vector>
@@ -123,7 +122,7 @@ bool GeneratorImpl::Generate() {
   }
 
   const TypeInfo* last_kind = nullptr;
-  std::streampos last_padding_pos;
+  size_t last_padding_line = 0;
 
   if (!FindAndEmitVectorAssignmentInLoopFunctions()) {
     return false;
@@ -137,10 +136,10 @@ bool GeneratorImpl::Generate() {
     // 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 (out_.tellp() != last_padding_pos) {
+    if (current_buffer_->lines.size() != last_padding_line) {
       if (last_kind && (last_kind != kind || decl->Is<ast::Function>())) {
-        out_ << std::endl;
-        last_padding_pos = out_.tellp();
+        line();
+        last_padding_line = current_buffer_->lines.size();
       }
     }
     last_kind = kind;
diff --git a/src/writer/text_generator.cc b/src/writer/text_generator.cc
index 6a6c36423f..4dc86d4899 100644
--- a/src/writer/text_generator.cc
+++ b/src/writer/text_generator.cc
@@ -14,6 +14,7 @@
 
 #include "src/writer/text_generator.h"
 
+#include <algorithm>
 #include <limits>
 
 namespace tint {
@@ -24,20 +25,20 @@ TextGenerator::TextGenerator(const Program* program)
 
 TextGenerator::~TextGenerator() = default;
 
-void TextGenerator::make_indent() {
-  make_indent(out_);
-}
-
-void TextGenerator::make_indent(std::ostream& out) const {
-  for (size_t i = 0; i < indent_; i++) {
-    out << " ";
-  }
-}
-
 std::string TextGenerator::UniqueIdentifier(const std::string& prefix) {
   return builder_.Symbols().NameFor(builder_.Symbols().New(prefix));
 }
 
+std::string TextGenerator::TrimSuffix(std::string str,
+                                      const std::string& suffix) {
+  if (str.size() >= suffix.size()) {
+    if (str.substr(str.size() - suffix.size(), suffix.size()) == suffix) {
+      return str.substr(0, str.size() - suffix.size());
+    }
+  }
+  return str;
+}
+
 TextGenerator::LineWriter::LineWriter(TextGenerator* generator)
     : gen(generator) {}
 
@@ -48,15 +49,45 @@ TextGenerator::LineWriter::LineWriter(LineWriter&& other) {
 
 TextGenerator::LineWriter::~LineWriter() {
   if (gen) {
-    auto str = os.str();
-    if (!str.empty()) {
-      gen->make_indent();
-      gen->out_ << str;
-    }
-    gen->out_ << std::endl;
+    gen->current_buffer_->Append(os.str());
   }
 }
 
+TextGenerator::TextBuffer::TextBuffer() = default;
+TextGenerator::TextBuffer::~TextBuffer() = default;
+
+void TextGenerator::TextBuffer::IncrementIndent() {
+  current_indent += 2;
+}
+
+void TextGenerator::TextBuffer::DecrementIndent() {
+  current_indent = std::max(2u, current_indent) - 2u;
+}
+
+void TextGenerator::TextBuffer::Append(const std::string& line) {
+  lines.emplace_back(Line{current_indent, line});
+}
+
+void TextGenerator::TextBuffer::Append(const TextBuffer& tb) {
+  for (auto& line : tb.lines) {
+    lines.emplace_back(Line{current_indent + line.indent, line.content});
+  }
+}
+
+std::string TextGenerator::TextBuffer::String() const {
+  std::stringstream ss;
+  for (auto& line : lines) {
+    if (!line.content.empty()) {
+      for (uint32_t i = 0; i < line.indent; i++) {
+        ss << " ";
+      }
+      ss << line.content;
+    }
+    ss << std::endl;
+  }
+  return ss.str();
+}
+
 TextGenerator::ScopedParen::ScopedParen(std::ostream& stream) : s(stream) {
   s << "(";
 }
diff --git a/src/writer/text_generator.h b/src/writer/text_generator.h
index 6d4f51050a..885c319661 100644
--- a/src/writer/text_generator.h
+++ b/src/writer/text_generator.h
@@ -18,6 +18,7 @@
 #include <sstream>
 #include <string>
 #include <utility>
+#include <vector>
 
 #include "src/diagnostic/diagnostic.h"
 #include "src/program_builder.h"
@@ -34,25 +35,12 @@ class TextGenerator {
   ~TextGenerator();
 
   /// Increment the emitter indent level
-  void increment_indent() { indent_ += 2; }
-  /// Decrement the emiter indent level
-  void decrement_indent() {
-    if (indent_ < 2) {
-      indent_ = 0;
-      return;
-    }
-    indent_ -= 2;
-  }
-
-  /// Writes the current indent to the output stream
-  void make_indent();
-
-  /// Writes the current indent to `out`
-  /// @param out the stream to write the indent to
-  void make_indent(std::ostream& out) const;
+  void increment_indent() { current_buffer_->IncrementIndent(); }
+  /// Decrement the emitter indent level
+  void decrement_indent() { current_buffer_->DecrementIndent(); }
 
   /// @returns the result data
-  std::string result() const { return out_.str(); }
+  std::string result() const { return main_buffer_.String(); }
 
   /// @returns the list of diagnostics raised by the generator.
   const diag::List& Diagnostics() const { return diagnostics_; }
@@ -65,6 +53,12 @@ class TextGenerator {
   /// empty "tint_symbol" will be used.
   std::string UniqueIdentifier(const std::string& prefix = "");
 
+  /// @param str the string
+  /// @param suffix the suffix to remove
+  /// @return returns str without the provided trailing suffix string. If str
+  /// doesn't end with suffix, str is returned unchanged.
+  std::string TrimSuffix(std::string str, const std::string& suffix);
+
  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.
@@ -81,7 +75,7 @@ class TextGenerator {
     ~LineWriter();
 
     /// @returns the ostringstream
-    operator std::ostream&() { return os; }
+    operator std::ostream &() { return os; }
 
     /// @param rhs the value to write to the line
     /// @returns the ostream so calls can be chained
@@ -98,6 +92,49 @@ class TextGenerator {
     TextGenerator* gen;
   };
 
+  /// Line holds a single line of text
+  struct Line {
+    /// The indentation of the line in whitespaces
+    uint32_t indent = 0;
+    /// The content of the line, without a trailing newline character
+    std::string content;
+  };
+
+  /// TextBuffer holds a list of lines of text.
+  struct TextBuffer {
+    // Constructor
+    TextBuffer();
+
+    // Destructor
+    ~TextBuffer();
+
+    /// IncrementIndent increases the indentation of lines that will be written
+    /// to the TextBuffer
+    void IncrementIndent();
+
+    /// DecrementIndent decreases the indentation of lines that will be written
+    /// to the TextBuffer
+    void DecrementIndent();
+
+    /// Appends the line to the end of the TextBuffer
+    /// @param line the line to append to the TextBuffer
+    void Append(const std::string& line);
+
+    /// Appends the lines of `tb` to the end of this TextBuffer
+    /// @param tb the TextBuffer to append to the end of this TextBuffer
+    void Append(const TextBuffer& tb);
+
+    /// @returns the buffer's content as a single string
+    std::string String() const;
+
+    /// The current indentation of the TextBuffer. Lines appended to the
+    /// TextBuffer will use this indentation.
+    uint32_t current_indent = 0;
+
+    /// The lines
+    std::vector<Line> lines;
+  };
+
   /// Helper for writing a '(' on construction and a ')' destruction.
   struct ScopedParen {
     /// Constructor
@@ -154,13 +191,14 @@ class TextGenerator {
   Program const* const program_;
   /// A ProgramBuilder that thinly wraps program_
   ProgramBuilder builder_;
-  /// The text output stream
-  std::ostringstream out_;
   /// Diagnostics generated by the generator
   diag::List diagnostics_;
+  /// The buffer the TextGenerator is currently appending lines to
+  TextBuffer* current_buffer_ = &main_buffer_;
 
  private:
-  size_t indent_ = 0;
+  /// The primary text buffer that the generator will emit
+  TextBuffer main_buffer_;
 };
 
 }  // namespace writer
diff --git a/src/writer/wgsl/generator_impl.cc b/src/writer/wgsl/generator_impl.cc
index 762f832d46..bb96dd6ffe 100644
--- a/src/writer/wgsl/generator_impl.cc
+++ b/src/writer/wgsl/generator_impl.cc
@@ -53,6 +53,7 @@
 #include "src/ast/workgroup_decoration.h"
 #include "src/sem/struct.h"
 #include "src/utils/math.h"
+#include "src/utils/scoped_assignment.h"
 #include "src/writer/float_to_string.h"
 
 namespace tint {
@@ -75,18 +76,16 @@ bool GeneratorImpl::Generate() {
         return false;
       }
     } else if (auto* var = decl->As<ast::Variable>()) {
-      make_indent();
-      if (!EmitVariable(var)) {
+      if (!EmitVariable(line(), var)) {
         return false;
       }
-      out_ << ";" << std::endl;
     } else {
       TINT_UNREACHABLE(Writer, diagnostics_);
       return false;
     }
 
     if (decl != program_->AST().GlobalDeclarations().back()) {
-      out_ << std::endl;
+      line();
     }
   }
 
@@ -94,14 +93,13 @@ bool GeneratorImpl::Generate() {
 }
 
 bool GeneratorImpl::EmitTypeDecl(const ast::TypeDecl* ty) {
-  make_indent();
-
   if (auto* alias = ty->As<ast::Alias>()) {
-    out_ << "type " << program_->Symbols().NameFor(alias->symbol()) << " = ";
-    if (!EmitType(alias->type())) {
+    auto out = line();
+    out << "type " << program_->Symbols().NameFor(alias->symbol()) << " = ";
+    if (!EmitType(out, alias->type())) {
       return false;
     }
-    out_ << ";" << std::endl;
+    out << ";";
   } else if (auto* str = ty->As<ast::Struct>()) {
     if (!EmitStructType(str)) {
       return false;
@@ -115,165 +113,171 @@ bool GeneratorImpl::EmitTypeDecl(const ast::TypeDecl* ty) {
   return true;
 }
 
-bool GeneratorImpl::EmitExpression(ast::Expression* expr) {
+bool GeneratorImpl::EmitExpression(std::ostream& out, ast::Expression* expr) {
   if (auto* a = expr->As<ast::ArrayAccessorExpression>()) {
-    return EmitArrayAccessor(a);
+    return EmitArrayAccessor(out, a);
   }
   if (auto* b = expr->As<ast::BinaryExpression>()) {
-    return EmitBinary(b);
+    return EmitBinary(out, b);
   }
   if (auto* b = expr->As<ast::BitcastExpression>()) {
-    return EmitBitcast(b);
+    return EmitBitcast(out, b);
   }
   if (auto* c = expr->As<ast::CallExpression>()) {
-    return EmitCall(c);
+    return EmitCall(out, c);
   }
   if (auto* i = expr->As<ast::IdentifierExpression>()) {
-    return EmitIdentifier(i);
+    return EmitIdentifier(out, i);
   }
   if (auto* c = expr->As<ast::ConstructorExpression>()) {
-    return EmitConstructor(c);
+    return EmitConstructor(out, c);
   }
   if (auto* m = expr->As<ast::MemberAccessorExpression>()) {
-    return EmitMemberAccessor(m);
+    return EmitMemberAccessor(out, m);
   }
   if (auto* u = expr->As<ast::UnaryOpExpression>()) {
-    return EmitUnaryOp(u);
+    return EmitUnaryOp(out, u);
   }
 
   diagnostics_.add_error(diag::System::Writer, "unknown expression type");
   return false;
 }
 
-bool GeneratorImpl::EmitArrayAccessor(ast::ArrayAccessorExpression* expr) {
+bool GeneratorImpl::EmitArrayAccessor(std::ostream& out,
+                                      ast::ArrayAccessorExpression* expr) {
   bool paren_lhs =
       !expr->array()
            ->IsAnyOf<ast::ArrayAccessorExpression, ast::CallExpression,
                      ast::IdentifierExpression, ast::MemberAccessorExpression,
                      ast::TypeConstructorExpression>();
   if (paren_lhs) {
-    out_ << "(";
+    out << "(";
   }
-  if (!EmitExpression(expr->array())) {
+  if (!EmitExpression(out, expr->array())) {
     return false;
   }
   if (paren_lhs) {
-    out_ << ")";
+    out << ")";
   }
-  out_ << "[";
+  out << "[";
 
-  if (!EmitExpression(expr->idx_expr())) {
+  if (!EmitExpression(out, expr->idx_expr())) {
     return false;
   }
-  out_ << "]";
+  out << "]";
 
   return true;
 }
 
-bool GeneratorImpl::EmitMemberAccessor(ast::MemberAccessorExpression* expr) {
+bool GeneratorImpl::EmitMemberAccessor(std::ostream& out,
+                                       ast::MemberAccessorExpression* expr) {
   bool paren_lhs =
       !expr->structure()
            ->IsAnyOf<ast::ArrayAccessorExpression, ast::CallExpression,
                      ast::IdentifierExpression, ast::MemberAccessorExpression,
                      ast::TypeConstructorExpression>();
   if (paren_lhs) {
-    out_ << "(";
+    out << "(";
   }
-  if (!EmitExpression(expr->structure())) {
+  if (!EmitExpression(out, expr->structure())) {
     return false;
   }
   if (paren_lhs) {
-    out_ << ")";
+    out << ")";
   }
 
-  out_ << ".";
+  out << ".";
 
-  return EmitExpression(expr->member());
+  return EmitExpression(out, expr->member());
 }
 
-bool GeneratorImpl::EmitBitcast(ast::BitcastExpression* expr) {
-  out_ << "bitcast<";
-  if (!EmitType(expr->type())) {
+bool GeneratorImpl::EmitBitcast(std::ostream& out,
+                                ast::BitcastExpression* expr) {
+  out << "bitcast<";
+  if (!EmitType(out, expr->type())) {
     return false;
   }
 
-  out_ << ">(";
-  if (!EmitExpression(expr->expr())) {
+  out << ">(";
+  if (!EmitExpression(out, expr->expr())) {
     return false;
   }
 
-  out_ << ")";
+  out << ")";
   return true;
 }
 
-bool GeneratorImpl::EmitCall(ast::CallExpression* expr) {
-  if (!EmitExpression(expr->func())) {
+bool GeneratorImpl::EmitCall(std::ostream& out, ast::CallExpression* expr) {
+  if (!EmitExpression(out, expr->func())) {
     return false;
   }
-  out_ << "(";
+  out << "(";
 
   bool first = true;
   const auto& params = expr->params();
   for (auto* param : params) {
     if (!first) {
-      out_ << ", ";
+      out << ", ";
     }
     first = false;
 
-    if (!EmitExpression(param)) {
+    if (!EmitExpression(out, param)) {
       return false;
     }
   }
 
-  out_ << ")";
+  out << ")";
 
   return true;
 }
 
-bool GeneratorImpl::EmitConstructor(ast::ConstructorExpression* expr) {
+bool GeneratorImpl::EmitConstructor(std::ostream& out,
+                                    ast::ConstructorExpression* expr) {
   if (auto* scalar = expr->As<ast::ScalarConstructorExpression>()) {
-    return EmitScalarConstructor(scalar);
+    return EmitScalarConstructor(out, scalar);
   }
-  return EmitTypeConstructor(expr->As<ast::TypeConstructorExpression>());
+  return EmitTypeConstructor(out, expr->As<ast::TypeConstructorExpression>());
 }
 
-bool GeneratorImpl::EmitTypeConstructor(ast::TypeConstructorExpression* expr) {
-  if (!EmitType(expr->type())) {
+bool GeneratorImpl::EmitTypeConstructor(std::ostream& out,
+                                        ast::TypeConstructorExpression* expr) {
+  if (!EmitType(out, expr->type())) {
     return false;
   }
 
-  out_ << "(";
+  out << "(";
 
   bool first = true;
   for (auto* e : expr->values()) {
     if (!first) {
-      out_ << ", ";
+      out << ", ";
     }
     first = false;
 
-    if (!EmitExpression(e)) {
+    if (!EmitExpression(out, e)) {
       return false;
     }
   }
 
-  out_ << ")";
+  out << ")";
   return true;
 }
 
 bool GeneratorImpl::EmitScalarConstructor(
+    std::ostream& out,
     ast::ScalarConstructorExpression* expr) {
-  return EmitLiteral(expr->literal());
+  return EmitLiteral(out, expr->literal());
 }
 
-bool GeneratorImpl::EmitLiteral(ast::Literal* lit) {
+bool GeneratorImpl::EmitLiteral(std::ostream& out, ast::Literal* lit) {
   if (auto* bl = lit->As<ast::BoolLiteral>()) {
-    out_ << (bl->IsTrue() ? "true" : "false");
+    out << (bl->IsTrue() ? "true" : "false");
   } else if (auto* fl = lit->As<ast::FloatLiteral>()) {
-    out_ << FloatToBitPreservingString(fl->value());
+    out << FloatToBitPreservingString(fl->value());
   } else if (auto* sl = lit->As<ast::SintLiteral>()) {
-    out_ << sl->value();
+    out << sl->value();
   } else if (auto* ul = lit->As<ast::UintLiteral>()) {
-    out_ << ul->value() << "u";
+    out << ul->value() << "u";
   } else {
     diagnostics_.add_error(diag::System::Writer, "unknown literal type");
     return false;
@@ -281,94 +285,98 @@ bool GeneratorImpl::EmitLiteral(ast::Literal* lit) {
   return true;
 }
 
-bool GeneratorImpl::EmitIdentifier(ast::IdentifierExpression* expr) {
-  auto* ident = expr->As<ast::IdentifierExpression>();
-  out_ << program_->Symbols().NameFor(ident->symbol());
+bool GeneratorImpl::EmitIdentifier(std::ostream& out,
+                                   ast::IdentifierExpression* expr) {
+  out << program_->Symbols().NameFor(expr->symbol());
   return true;
 }
 
 bool GeneratorImpl::EmitFunction(ast::Function* func) {
   if (func->decorations().size()) {
-    make_indent();
-    if (!EmitDecorations(func->decorations())) {
-      return false;
-    }
-    out_ << std::endl;
-  }
-
-  make_indent();
-  out_ << "fn " << program_->Symbols().NameFor(func->symbol()) << "(";
-
-  bool first = true;
-  for (auto* v : func->params()) {
-    if (!first) {
-      out_ << ", ";
-    }
-    first = false;
-
-    if (!v->decorations().empty()) {
-      if (!EmitDecorations(v->decorations())) {
-        return false;
-      }
-      out_ << " ";
-    }
-
-    out_ << program_->Symbols().NameFor(v->symbol()) << " : ";
-
-    if (!EmitType(v->type())) {
+    if (!EmitDecorations(line(), func->decorations())) {
       return false;
     }
   }
+  {
+    auto out = line();
+    out << "fn " << program_->Symbols().NameFor(func->symbol()) << "(";
 
-  out_ << ")";
+    bool first = true;
+    for (auto* v : func->params()) {
+      if (!first) {
+        out << ", ";
+      }
+      first = false;
 
-  if (!func->return_type()->Is<ast::Void>() ||
-      !func->return_type_decorations().empty()) {
-    out_ << " -> ";
+      if (!v->decorations().empty()) {
+        if (!EmitDecorations(out, v->decorations())) {
+          return false;
+        }
+        out << " ";
+      }
 
-    if (!func->return_type_decorations().empty()) {
-      if (!EmitDecorations(func->return_type_decorations())) {
+      out << program_->Symbols().NameFor(v->symbol()) << " : ";
+
+      if (!EmitType(out, v->type())) {
         return false;
       }
-      out_ << " ";
     }
 
-    if (!EmitType(func->return_type())) {
-      return false;
+    out << ")";
+
+    if (!func->return_type()->Is<ast::Void>() ||
+        !func->return_type_decorations().empty()) {
+      out << " -> ";
+
+      if (!func->return_type_decorations().empty()) {
+        if (!EmitDecorations(out, func->return_type_decorations())) {
+          return false;
+        }
+        out << " ";
+      }
+
+      if (!EmitType(out, func->return_type())) {
+        return false;
+      }
+    }
+
+    if (func->body()) {
+      out << " {";
     }
   }
 
   if (func->body()) {
-    out_ << " ";
-    return EmitBlockAndNewline(func->body());
-  } else {
-    out_ << std::endl;
+    if (!EmitStatementsWithIndent(func->body()->statements())) {
+      return false;
+    }
+    line() << "}";
   }
 
   return true;
 }
 
-bool GeneratorImpl::EmitImageFormat(const ast::ImageFormat fmt) {
+bool GeneratorImpl::EmitImageFormat(std::ostream& out,
+                                    const ast::ImageFormat fmt) {
   switch (fmt) {
     case ast::ImageFormat::kNone:
       diagnostics_.add_error(diag::System::Writer, "unknown image format");
       return false;
     default:
-      out_ << fmt;
+      out << fmt;
   }
   return true;
 }
 
-bool GeneratorImpl::EmitAccess(const ast::Access access) {
+bool GeneratorImpl::EmitAccess(std::ostream& out, const ast::Access access) {
   switch (access) {
     case ast::Access::kRead:
-      out_ << "read";
+      out << "read";
       return true;
     case ast::Access::kWrite:
-      out_ << "write";
+      out << "write";
       return true;
     case ast::Access::kReadWrite:
-      out_ << "read_write";
+      out << "read_write";
       return true;
     default:
       break;
@@ -377,65 +385,65 @@ bool GeneratorImpl::EmitAccess(const ast::Access access) {
   return false;
 }
 
-bool GeneratorImpl::EmitType(const ast::Type* ty) {
+bool GeneratorImpl::EmitType(std::ostream& out, const ast::Type* ty) {
   if (auto* ary = ty->As<ast::Array>()) {
     for (auto* deco : ary->decorations()) {
       if (auto* stride = deco->As<ast::StrideDecoration>()) {
-        out_ << "[[stride(" << stride->stride() << ")]] ";
+        out << "[[stride(" << stride->stride() << ")]] ";
       }
     }
 
-    out_ << "array<";
-    if (!EmitType(ary->type())) {
+    out << "array<";
+    if (!EmitType(out, ary->type())) {
       return false;
     }
 
     if (!ary->IsRuntimeArray())
-      out_ << ", " << ary->size();
+      out << ", " << ary->size();
 
-    out_ << ">";
+    out << ">";
   } else if (ty->Is<ast::Bool>()) {
-    out_ << "bool";
+    out << "bool";
   } else if (ty->Is<ast::F32>()) {
-    out_ << "f32";
+    out << "f32";
   } else if (ty->Is<ast::I32>()) {
-    out_ << "i32";
+    out << "i32";
   } else if (auto* mat = ty->As<ast::Matrix>()) {
-    out_ << "mat" << mat->columns() << "x" << mat->rows() << "<";
-    if (!EmitType(mat->type())) {
+    out << "mat" << mat->columns() << "x" << mat->rows() << "<";
+    if (!EmitType(out, mat->type())) {
       return false;
     }
-    out_ << ">";
+    out << ">";
   } else if (auto* ptr = ty->As<ast::Pointer>()) {
-    out_ << "ptr<" << ptr->storage_class() << ", ";
-    if (!EmitType(ptr->type())) {
+    out << "ptr<" << ptr->storage_class() << ", ";
+    if (!EmitType(out, ptr->type())) {
       return false;
     }
-    out_ << ">";
+    out << ">";
   } else if (auto* atomic = ty->As<ast::Atomic>()) {
-    out_ << "atomic<";
-    if (!EmitType(atomic->type())) {
+    out << "atomic<";
+    if (!EmitType(out, atomic->type())) {
       return false;
     }
-    out_ << ">";
+    out << ">";
   } else if (auto* sampler = ty->As<ast::Sampler>()) {
-    out_ << "sampler";
+    out << "sampler";
 
     if (sampler->IsComparison()) {
-      out_ << "_comparison";
+      out << "_comparison";
     }
   } else if (ty->Is<ast::ExternalTexture>()) {
-    out_ << "external_texture";
+    out << "external_texture";
   } else if (auto* texture = ty->As<ast::Texture>()) {
-    out_ << "texture_";
+    out << "texture_";
     if (texture->Is<ast::DepthTexture>()) {
-      out_ << "depth_";
+      out << "depth_";
     } else if (texture->Is<ast::SampledTexture>()) {
       /* nothing to emit */
     } else if (texture->Is<ast::MultisampledTexture>()) {
-      out_ << "multisampled_";
+      out << "multisampled_";
     } else if (texture->Is<ast::StorageTexture>()) {
-      out_ << "storage_";
+      out << "storage_";
     } else {
       diagnostics_.add_error(diag::System::Writer, "unknown texture type");
       return false;
@@ -443,22 +451,22 @@ bool GeneratorImpl::EmitType(const ast::Type* ty) {
 
     switch (texture->dim()) {
       case ast::TextureDimension::k1d:
-        out_ << "1d";
+        out << "1d";
         break;
       case ast::TextureDimension::k2d:
-        out_ << "2d";
+        out << "2d";
         break;
       case ast::TextureDimension::k2dArray:
-        out_ << "2d_array";
+        out << "2d_array";
         break;
       case ast::TextureDimension::k3d:
-        out_ << "3d";
+        out << "3d";
         break;
       case ast::TextureDimension::kCube:
-        out_ << "cube";
+        out << "cube";
         break;
       case ast::TextureDimension::kCubeArray:
-        out_ << "cube_array";
+        out << "cube_array";
         break;
       default:
         diagnostics_.add_error(diag::System::Writer,
@@ -467,41 +475,41 @@ bool GeneratorImpl::EmitType(const ast::Type* ty) {
     }
 
     if (auto* sampled = texture->As<ast::SampledTexture>()) {
-      out_ << "<";
-      if (!EmitType(sampled->type())) {
+      out << "<";
+      if (!EmitType(out, sampled->type())) {
         return false;
       }
-      out_ << ">";
+      out << ">";
     } else if (auto* ms = texture->As<ast::MultisampledTexture>()) {
-      out_ << "<";
-      if (!EmitType(ms->type())) {
+      out << "<";
+      if (!EmitType(out, ms->type())) {
         return false;
       }
-      out_ << ">";
+      out << ">";
     } else if (auto* storage = texture->As<ast::StorageTexture>()) {
-      out_ << "<";
-      if (!EmitImageFormat(storage->image_format())) {
+      out << "<";
+      if (!EmitImageFormat(out, storage->image_format())) {
         return false;
       }
-      out_ << ", ";
-      if (!EmitAccess(storage->access())) {
+      out << ", ";
+      if (!EmitAccess(out, storage->access())) {
         return false;
       }
-      out_ << ">";
+      out << ">";
     }
 
   } else if (ty->Is<ast::U32>()) {
-    out_ << "u32";
+    out << "u32";
   } else if (auto* vec = ty->As<ast::Vector>()) {
-    out_ << "vec" << vec->size() << "<";
-    if (!EmitType(vec->type())) {
+    out << "vec" << vec->size() << "<";
+    if (!EmitType(out, vec->type())) {
       return false;
     }
-    out_ << ">";
+    out << ">";
   } else if (ty->Is<ast::Void>()) {
-    out_ << "void";
+    out << "void";
   } else if (auto* tn = ty->As<ast::TypeName>()) {
-    out_ << program_->Symbols().NameFor(tn->name());
+    out << program_->Symbols().NameFor(tn->name());
   } else {
     diagnostics_.add_error(diag::System::Writer,
                            "unknown type in EmitType: " + ty->type_name());
@@ -512,21 +520,18 @@ bool GeneratorImpl::EmitType(const ast::Type* ty) {
 
 bool GeneratorImpl::EmitStructType(const ast::Struct* str) {
   if (str->decorations().size()) {
-    if (!EmitDecorations(str->decorations())) {
+    if (!EmitDecorations(line(), str->decorations())) {
       return false;
     }
-    out_ << std::endl;
   }
-  out_ << "struct " << program_->Symbols().NameFor(str->name()) << " {"
-       << std::endl;
+  line() << "struct " << program_->Symbols().NameFor(str->name()) << " {";
 
   auto add_padding = [&](uint32_t size) {
-    make_indent();
-    out_ << "[[size(" << size << ")]]" << std::endl;
-    make_indent();
+    line() << "[[size(" << size << ")]]";
+
     // Note: u32 is the smallest primitive we currently support. When WGSL
     // supports smaller types, this will need to be updated.
-    out_ << UniqueIdentifier("padding") << " : u32;" << std::endl;
+    line() << UniqueIdentifier("padding") << " : u32;";
   };
 
   increment_indent();
@@ -555,96 +560,95 @@ bool GeneratorImpl::EmitStructType(const ast::Struct* str) {
     }
 
     if (!decorations_sanitized.empty()) {
-      make_indent();
-      if (!EmitDecorations(decorations_sanitized)) {
+      if (!EmitDecorations(line(), decorations_sanitized)) {
         return false;
       }
-      out_ << std::endl;
     }
 
-    make_indent();
-    out_ << program_->Symbols().NameFor(mem->symbol()) << " : ";
-    if (!EmitType(mem->type())) {
+    auto out = line();
+    out << program_->Symbols().NameFor(mem->symbol()) << " : ";
+    if (!EmitType(out, mem->type())) {
       return false;
     }
-    out_ << ";" << std::endl;
+    out << ";";
   }
   decrement_indent();
-  make_indent();
 
-  out_ << "};" << std::endl;
+  line() << "};";
   return true;
 }
 
-bool GeneratorImpl::EmitVariable(ast::Variable* var) {
+bool GeneratorImpl::EmitVariable(std::ostream& out, ast::Variable* var) {
   if (!var->decorations().empty()) {
-    if (!EmitDecorations(var->decorations())) {
+    if (!EmitDecorations(out, var->decorations())) {
       return false;
     }
-    out_ << " ";
+    out << " ";
   }
 
   if (var->is_const()) {
-    out_ << "let";
+    out << "let";
   } else {
-    out_ << "var";
+    out << "var";
     auto sc = var->declared_storage_class();
     auto ac = var->declared_access();
     if (sc != ast::StorageClass::kNone || ac != ast::Access::kUndefined) {
-      out_ << "<" << sc;
+      out << "<" << sc;
       if (ac != ast::Access::kUndefined) {
-        out_ << ", ";
-        if (!EmitAccess(ac)) {
+        out << ", ";
+        if (!EmitAccess(out, ac)) {
           return false;
         }
       }
-      out_ << ">";
+      out << ">";
     }
   }
 
-  out_ << " " << program_->Symbols().NameFor(var->symbol());
+  out << " " << program_->Symbols().NameFor(var->symbol());
 
   if (auto* ty = var->type()) {
-    out_ << " : ";
-    if (!EmitType(ty)) {
+    out << " : ";
+    if (!EmitType(out, ty)) {
       return false;
     }
   }
 
   if (var->constructor() != nullptr) {
-    out_ << " = ";
-    if (!EmitExpression(var->constructor())) {
+    out << " = ";
+    if (!EmitExpression(out, var->constructor())) {
       return false;
     }
   }
+  out << ";";
 
   return true;
 }
 
-bool GeneratorImpl::EmitDecorations(const ast::DecorationList& decos) {
-  out_ << "[[";
+bool GeneratorImpl::EmitDecorations(std::ostream& out,
+                                    const ast::DecorationList& decos) {
+  out << "[[";
   bool first = true;
   for (auto* deco : decos) {
     if (!first) {
-      out_ << ", ";
+      out << ", ";
     }
     first = false;
 
     if (auto* workgroup = deco->As<ast::WorkgroupDecoration>()) {
       auto values = workgroup->values();
-      out_ << "workgroup_size(";
+      out << "workgroup_size(";
       for (int i = 0; i < 3; i++) {
         if (values[i]) {
           if (i > 0) {
-            out_ << ", ";
+            out << ", ";
           }
           if (auto* ident = values[i]->As<ast::IdentifierExpression>()) {
-            if (!EmitIdentifier(ident)) {
+            if (!EmitIdentifier(out, ident)) {
               return false;
             }
           } else if (auto* scalar =
                          values[i]->As<ast::ScalarConstructorExpression>()) {
-            if (!EmitScalarConstructor(scalar)) {
+            if (!EmitScalarConstructor(out, scalar)) {
               return false;
             }
           } else {
@@ -653,206 +657,182 @@ bool GeneratorImpl::EmitDecorations(const ast::DecorationList& decos) {
           }
         }
       }
-      out_ << ")";
+      out << ")";
     } else if (deco->Is<ast::StructBlockDecoration>()) {
-      out_ << "block";
+      out << "block";
     } else if (auto* stage = deco->As<ast::StageDecoration>()) {
-      out_ << "stage(" << stage->value() << ")";
+      out << "stage(" << stage->value() << ")";
     } else if (auto* binding = deco->As<ast::BindingDecoration>()) {
-      out_ << "binding(" << binding->value() << ")";
+      out << "binding(" << binding->value() << ")";
     } else if (auto* group = deco->As<ast::GroupDecoration>()) {
-      out_ << "group(" << group->value() << ")";
+      out << "group(" << group->value() << ")";
     } else if (auto* location = deco->As<ast::LocationDecoration>()) {
-      out_ << "location(" << location->value() << ")";
+      out << "location(" << location->value() << ")";
     } else if (auto* builtin = deco->As<ast::BuiltinDecoration>()) {
-      out_ << "builtin(" << builtin->value() << ")";
+      out << "builtin(" << builtin->value() << ")";
     } else if (auto* interpolate = deco->As<ast::InterpolateDecoration>()) {
-      out_ << "interpolate(" << interpolate->type();
+      out << "interpolate(" << interpolate->type();
       if (interpolate->sampling() != ast::InterpolationSampling::kNone) {
-        out_ << ", " << interpolate->sampling();
+        out << ", " << interpolate->sampling();
       }
-      out_ << ")";
+      out << ")";
     } else if (auto* override_deco = deco->As<ast::OverrideDecoration>()) {
-      out_ << "override";
+      out << "override";
       if (override_deco->HasValue()) {
-        out_ << "(" << override_deco->value() << ")";
+        out << "(" << override_deco->value() << ")";
       }
     } else if (auto* size = deco->As<ast::StructMemberSizeDecoration>()) {
-      out_ << "size(" << size->size() << ")";
+      out << "size(" << size->size() << ")";
     } else if (auto* align = deco->As<ast::StructMemberAlignDecoration>()) {
-      out_ << "align(" << align->align() << ")";
+      out << "align(" << align->align() << ")";
     } else if (auto* internal = deco->As<ast::InternalDecoration>()) {
-      out_ << "internal(" << internal->InternalName() << ")";
+      out << "internal(" << internal->InternalName() << ")";
     } else {
       TINT_ICE(Writer, diagnostics_)
           << "Unsupported decoration '" << deco->TypeInfo().name << "'";
       return false;
     }
   }
-  out_ << "]]";
+  out << "]]";
 
   return true;
 }
 
-bool GeneratorImpl::EmitBinary(ast::BinaryExpression* expr) {
-  out_ << "(";
+bool GeneratorImpl::EmitBinary(std::ostream& out, ast::BinaryExpression* expr) {
+  out << "(";
 
-  if (!EmitExpression(expr->lhs())) {
+  if (!EmitExpression(out, expr->lhs())) {
     return false;
   }
-  out_ << " ";
+  out << " ";
 
   switch (expr->op()) {
     case ast::BinaryOp::kAnd:
-      out_ << "&";
+      out << "&";
       break;
     case ast::BinaryOp::kOr:
-      out_ << "|";
+      out << "|";
       break;
     case ast::BinaryOp::kXor:
-      out_ << "^";
+      out << "^";
       break;
     case ast::BinaryOp::kLogicalAnd:
-      out_ << "&&";
+      out << "&&";
       break;
     case ast::BinaryOp::kLogicalOr:
-      out_ << "||";
+      out << "||";
       break;
     case ast::BinaryOp::kEqual:
-      out_ << "==";
+      out << "==";
       break;
     case ast::BinaryOp::kNotEqual:
-      out_ << "!=";
+      out << "!=";
       break;
     case ast::BinaryOp::kLessThan:
-      out_ << "<";
+      out << "<";
       break;
     case ast::BinaryOp::kGreaterThan:
-      out_ << ">";
+      out << ">";
       break;
     case ast::BinaryOp::kLessThanEqual:
-      out_ << "<=";
+      out << "<=";
       break;
     case ast::BinaryOp::kGreaterThanEqual:
-      out_ << ">=";
+      out << ">=";
       break;
     case ast::BinaryOp::kShiftLeft:
-      out_ << "<<";
+      out << "<<";
       break;
     case ast::BinaryOp::kShiftRight:
-      out_ << ">>";
+      out << ">>";
       break;
     case ast::BinaryOp::kAdd:
-      out_ << "+";
+      out << "+";
       break;
     case ast::BinaryOp::kSubtract:
-      out_ << "-";
+      out << "-";
       break;
     case ast::BinaryOp::kMultiply:
-      out_ << "*";
+      out << "*";
       break;
     case ast::BinaryOp::kDivide:
-      out_ << "/";
+      out << "/";
       break;
     case ast::BinaryOp::kModulo:
-      out_ << "%";
+      out << "%";
       break;
     case ast::BinaryOp::kNone:
       diagnostics_.add_error(diag::System::Writer,
                              "missing binary operation type");
       return false;
   }
-  out_ << " ";
+  out << " ";
 
-  if (!EmitExpression(expr->rhs())) {
+  if (!EmitExpression(out, expr->rhs())) {
     return false;
   }
 
-  out_ << ")";
+  out << ")";
   return true;
 }
 
-bool GeneratorImpl::EmitUnaryOp(ast::UnaryOpExpression* expr) {
+bool GeneratorImpl::EmitUnaryOp(std::ostream& out,
+                                ast::UnaryOpExpression* expr) {
   switch (expr->op()) {
     case ast::UnaryOp::kAddressOf:
-      out_ << "&";
+      out << "&";
       break;
     case ast::UnaryOp::kComplement:
-      out_ << "~";
+      out << "~";
       break;
     case ast::UnaryOp::kIndirection:
-      out_ << "*";
+      out << "*";
       break;
     case ast::UnaryOp::kNot:
-      out_ << "!";
+      out << "!";
       break;
     case ast::UnaryOp::kNegation:
-      out_ << "-";
+      out << "-";
       break;
   }
-  out_ << "(";
+  out << "(";
 
-  if (!EmitExpression(expr->expr())) {
+  if (!EmitExpression(out, expr->expr())) {
     return false;
   }
 
-  out_ << ")";
+  out << ")";
 
   return true;
 }
 
 bool GeneratorImpl::EmitBlock(const ast::BlockStatement* stmt) {
-  out_ << "{" << std::endl;
-  increment_indent();
-
-  for (auto* s : *stmt) {
-    if (!EmitStatement(s)) {
-      return false;
-    }
+  line() << "{";
+  if (!EmitStatementsWithIndent(stmt->statements())) {
+    return false;
   }
-
-  decrement_indent();
-  make_indent();
-  out_ << "}";
+  line() << "}";
 
   return true;
 }
 
-bool GeneratorImpl::EmitBlockAndNewline(const ast::BlockStatement* stmt) {
-  const bool result = EmitBlock(stmt);
-  if (result) {
-    out_ << std::endl;
-  }
-  return result;
-}
-
 bool GeneratorImpl::EmitStatement(ast::Statement* stmt) {
-  make_indent();
-
-  if (!EmitRawStatement(stmt)) {
-    return false;
-  }
-
-  if (!stmt->IsAnyOf<ast::BlockStatement, ast::IfStatement,
-                     ast::SwitchStatement, ast::LoopStatement,
-                     ast::ForLoopStatement>()) {
-    out_ << ";" << std::endl;
-  }
-  return true;
-}
-
-bool GeneratorImpl::EmitRawStatement(ast::Statement* stmt) {
   if (auto* a = stmt->As<ast::AssignmentStatement>()) {
     return EmitAssign(a);
   }
   if (auto* b = stmt->As<ast::BlockStatement>()) {
-    return EmitBlockAndNewline(b);
+    return EmitBlock(b);
   }
   if (auto* b = stmt->As<ast::BreakStatement>()) {
     return EmitBreak(b);
   }
   if (auto* c = stmt->As<ast::CallStatement>()) {
-    return EmitCall(c->expr());
+    auto out = line();
+    if (!EmitCall(out, c->expr())) {
+      return false;
+    }
+    out << ";";
+    return true;
   }
   if (auto* c = stmt->As<ast::ContinueStatement>()) {
     return EmitContinue(c);
@@ -879,7 +859,7 @@ bool GeneratorImpl::EmitRawStatement(ast::Statement* stmt) {
     return EmitSwitch(s);
   }
   if (auto* v = stmt->As<ast::VariableDeclStatement>()) {
-    return EmitVariable(v->variable());
+    return EmitVariable(line(), v->variable());
   }
 
   diagnostics_.add_error(diag::System::Writer,
@@ -887,201 +867,247 @@ bool GeneratorImpl::EmitRawStatement(ast::Statement* stmt) {
   return false;
 }
 
+bool GeneratorImpl::EmitStatements(const ast::StatementList& stmts) {
+  for (auto* s : stmts) {
+    if (!EmitStatement(s)) {
+      return false;
+    }
+  }
+  return true;
+}
+
+bool GeneratorImpl::EmitStatementsWithIndent(const ast::StatementList& stmts) {
+  ScopedIndent si(this);
+  return EmitStatements(stmts);
+}
+
 bool GeneratorImpl::EmitAssign(ast::AssignmentStatement* stmt) {
-  if (!EmitExpression(stmt->lhs())) {
+  auto out = line();
+
+  if (!EmitExpression(out, stmt->lhs())) {
     return false;
   }
 
-  out_ << " = ";
+  out << " = ";
 
-  if (!EmitExpression(stmt->rhs())) {
+  if (!EmitExpression(out, stmt->rhs())) {
     return false;
   }
 
+  out << ";";
+
   return true;
 }
 
 bool GeneratorImpl::EmitBreak(ast::BreakStatement*) {
-  out_ << "break";
+  line() << "break;";
   return true;
 }
 
 bool GeneratorImpl::EmitCase(ast::CaseStatement* stmt) {
-  make_indent();
-
   if (stmt->IsDefault()) {
-    out_ << "default";
+    line() << "default: {";
   } else {
-    out_ << "case ";
+    auto out = line();
+    out << "case ";
 
     bool first = true;
     for (auto* selector : stmt->selectors()) {
       if (!first) {
-        out_ << ", ";
+        out << ", ";
       }
 
       first = false;
-      if (!EmitLiteral(selector)) {
+      if (!EmitLiteral(out, selector)) {
         return false;
       }
     }
+    out << ": {";
   }
-  out_ << ": ";
 
-  return EmitBlockAndNewline(stmt->body());
-}
+  if (!EmitStatementsWithIndent(stmt->body()->statements())) {
+    return false;
+  }
 
-bool GeneratorImpl::EmitContinue(ast::ContinueStatement*) {
-  out_ << "continue";
+  line() << "}";
   return true;
 }
 
-bool GeneratorImpl::EmitElse(ast::ElseStatement* stmt) {
-  if (stmt->HasCondition()) {
-    out_ << " elseif (";
-    if (!EmitExpression(stmt->condition())) {
-      return false;
-    }
-    out_ << ") ";
-  } else {
-    out_ << " else ";
-  }
-
-  return EmitBlock(stmt->body());
+bool GeneratorImpl::EmitContinue(ast::ContinueStatement*) {
+  line() << "continue;";
+  return true;
 }
 
 bool GeneratorImpl::EmitFallthrough(ast::FallthroughStatement*) {
-  out_ << "fallthrough";
+  line() << "fallthrough;";
   return true;
 }
 
 bool GeneratorImpl::EmitIf(ast::IfStatement* stmt) {
-  out_ << "if (";
-  if (!EmitExpression(stmt->condition())) {
-    return false;
+  {
+    auto out = line();
+    out << "if (";
+    if (!EmitExpression(out, stmt->condition())) {
+      return false;
+    }
+    out << ") {";
   }
-  out_ << ") ";
 
-  if (!EmitBlock(stmt->body())) {
+  if (!EmitStatementsWithIndent(stmt->body()->statements())) {
     return false;
   }
 
   for (auto* e : stmt->else_statements()) {
-    if (!EmitElse(e)) {
+    if (e->HasCondition()) {
+      auto out = line();
+      out << "} elseif (";
+      if (!EmitExpression(out, e->condition())) {
+        return false;
+      }
+      out << ") {";
+    } else {
+      line() << "} else {";
+    }
+
+    if (!EmitStatementsWithIndent(e->body()->statements())) {
       return false;
     }
   }
-  out_ << std::endl;
+
+  line() << "}";
 
   return true;
 }
 
 bool GeneratorImpl::EmitDiscard(ast::DiscardStatement*) {
-  out_ << "discard";
+  line() << "discard;";
   return true;
 }
 
 bool GeneratorImpl::EmitLoop(ast::LoopStatement* stmt) {
-  out_ << "loop {" << std::endl;
+  line() << "loop {";
   increment_indent();
 
-  for (auto* s : *(stmt->body())) {
-    if (!EmitStatement(s)) {
-      return false;
-    }
+  if (!EmitStatements(stmt->body()->statements())) {
+    return false;
   }
 
   if (stmt->has_continuing()) {
-    out_ << std::endl;
-
-    make_indent();
-    out_ << "continuing ";
-
-    if (!EmitBlockAndNewline(stmt->continuing())) {
+    line();
+    line() << "continuing {";
+    if (!EmitStatementsWithIndent(stmt->continuing()->statements())) {
       return false;
     }
+    line() << "}";
   }
 
   decrement_indent();
-  make_indent();
-  out_ << "}" << std::endl;
+  line() << "}";
 
   return true;
 }
 
 bool GeneratorImpl::EmitForLoop(ast::ForLoopStatement* stmt) {
-  out_ << "for";
-  {
-    ScopedParen sp(out_);
-    if (auto* init = stmt->initializer()) {
-      if (!EmitRawStatement(init)) {
-        return false;
-      }
-    }
-
-    out_ << "; ";
-
-    if (auto* cond = stmt->condition()) {
-      if (!EmitExpression(cond)) {
-        return false;
-      }
-    }
-
-    out_ << "; ";
-
-    if (auto* cont = stmt->continuing()) {
-      if (!EmitRawStatement(cont)) {
-        return false;
-      }
-    }
-  }
-  out_ << " {" << std::endl;
-
-  {
-    ScopedIndent si(this);
-    for (auto* s : stmt->body()->statements()) {
-      if (!EmitStatement(s)) {
-        return false;
-      }
+  TextBuffer init_buf;
+  if (auto* init = stmt->initializer()) {
+    TINT_SCOPED_ASSIGNMENT(current_buffer_, &init_buf);
+    if (!EmitStatement(init)) {
+      return false;
     }
   }
 
-  make_indent();
-  out_ << "}" << std::endl;
+  TextBuffer cont_buf;
+  if (auto* cont = stmt->continuing()) {
+    TINT_SCOPED_ASSIGNMENT(current_buffer_, &cont_buf);
+    if (!EmitStatement(cont)) {
+      return false;
+    }
+  }
+
+  {
+    auto out = line();
+    out << "for";
+    {
+      ScopedParen sp(out);
+      switch (init_buf.lines.size()) {
+        case 0:  // No initializer
+          out << ";";
+          break;
+        case 1:  // Single line initializer statement
+          out << init_buf.lines[0].content;
+          break;
+        default:  // Block initializer statement
+          current_buffer_->Append(init_buf);
+          break;
+      }
+
+      out << " ";
+
+      if (auto* cond = stmt->condition()) {
+        if (!EmitExpression(out, cond)) {
+          return false;
+        }
+      }
+
+      out << "; ";
+
+      switch (cont_buf.lines.size()) {
+        case 0:  // No continuing
+          out << ";";
+          break;
+        case 1:  // Single line continuing statement
+          out << TrimSuffix(cont_buf.lines[0].content, ";");
+          break;
+        default:  // Block continuing statement
+          current_buffer_->Append(cont_buf);
+          break;
+      }
+    }
+    out << " {";
+  }
+
+  if (!EmitStatementsWithIndent(stmt->body()->statements())) {
+    return false;
+  }
+
+  line() << "}";
 
   return true;
 }
 
 bool GeneratorImpl::EmitReturn(ast::ReturnStatement* stmt) {
-  out_ << "return";
+  auto out = line();
+  out << "return";
   if (stmt->has_value()) {
-    out_ << " ";
-    if (!EmitExpression(stmt->value())) {
+    out << " ";
+    if (!EmitExpression(out, stmt->value())) {
       return false;
     }
   }
+  out << ";";
   return true;
 }
 
 bool GeneratorImpl::EmitSwitch(ast::SwitchStatement* stmt) {
-  out_ << "switch(";
-  if (!EmitExpression(stmt->condition())) {
-    return false;
-  }
-  out_ << ") {" << std::endl;
-
-  increment_indent();
-
-  for (auto* s : stmt->body()) {
-    if (!EmitCase(s)) {
+  {
+    auto out = line();
+    out << "switch(";
+    if (!EmitExpression(out, stmt->condition())) {
       return false;
     }
+    out << ") {";
+  }
+
+  {
+    ScopedIndent si(this);
+    for (auto* s : stmt->body()) {
+      if (!EmitCase(s)) {
+        return false;
+      }
+    }
   }
 
-  decrement_indent();
-  make_indent();
-  out_ << "}" << std::endl;
-
+  line() << "}";
   return true;
 }
 
diff --git a/src/writer/wgsl/generator_impl.h b/src/writer/wgsl/generator_impl.h
index 34d5fa1a51..c413ee2081 100644
--- a/src/writer/wgsl/generator_impl.h
+++ b/src/writer/wgsl/generator_impl.h
@@ -60,57 +60,56 @@ class GeneratorImpl : public TextGenerator {
   /// @returns true if the declared type was emitted
   bool EmitTypeDecl(const ast::TypeDecl* ty);
   /// Handles an array accessor expression
+  /// @param out the output of the expression stream
   /// @param expr the expression to emit
   /// @returns true if the array accessor was emitted
-  bool EmitArrayAccessor(ast::ArrayAccessorExpression* expr);
+  bool EmitArrayAccessor(std::ostream& out, ast::ArrayAccessorExpression* expr);
   /// Handles an assignment statement
   /// @param stmt the statement to emit
   /// @returns true if the statement was emitted successfully
   bool EmitAssign(ast::AssignmentStatement* stmt);
   /// Handles generating a binary expression
+  /// @param out the output of the expression stream
   /// @param expr the binary expression
   /// @returns true if the expression was emitted, false otherwise
-  bool EmitBinary(ast::BinaryExpression* expr);
+  bool EmitBinary(std::ostream& out, ast::BinaryExpression* expr);
   /// Handles generating a bitcast expression
+  /// @param out the output of the expression stream
   /// @param expr the bitcast expression
   /// @returns true if the bitcast was emitted
-  bool EmitBitcast(ast::BitcastExpression* expr);
+  bool EmitBitcast(std::ostream& out, ast::BitcastExpression* expr);
   /// Handles a block statement
   /// @param stmt the statement to emit
   /// @returns true if the statement was emitted successfully
   bool EmitBlock(const ast::BlockStatement* stmt);
-  /// Handles a block statement with a newline at the end
-  /// @param stmt the statement to emit
-  /// @returns true if the statement was emitted successfully
-  bool EmitBlockAndNewline(const ast::BlockStatement* stmt);
   /// Handles a break statement
   /// @param stmt the statement to emit
   /// @returns true if the statement was emitted successfully
   bool EmitBreak(ast::BreakStatement* stmt);
   /// Handles generating a call expression
+  /// @param out the output of the expression stream
   /// @param expr the call expression
   /// @returns true if the call expression is emitted
-  bool EmitCall(ast::CallExpression* expr);
+  bool EmitCall(std::ostream& out, ast::CallExpression* expr);
   /// Handles a case statement
   /// @param stmt the statement
   /// @returns true if the statment was emitted successfully
   bool EmitCase(ast::CaseStatement* stmt);
   /// Handles generating a scalar constructor
+  /// @param out the output of the expression stream
   /// @param expr the scalar constructor expression
   /// @returns true if the scalar constructor is emitted
-  bool EmitScalarConstructor(ast::ScalarConstructorExpression* expr);
+  bool EmitScalarConstructor(std::ostream& out,
+                             ast::ScalarConstructorExpression* expr);
   /// Handles a continue statement
   /// @param stmt the statement to emit
   /// @returns true if the statement was emitted successfully
   bool EmitContinue(ast::ContinueStatement* stmt);
-  /// Handles generating an else statement
-  /// @param stmt the statement to emit
-  /// @returns true if the statement was emitted
-  bool EmitElse(ast::ElseStatement* stmt);
   /// Handles generate an Expression
+  /// @param out the output of the expression stream
   /// @param expr the expression
   /// @returns true if the expression was emitted
-  bool EmitExpression(ast::Expression* expr);
+  bool EmitExpression(std::ostream& out, ast::Expression* expr);
   /// Handles generating a fallthrough statement
   /// @param stmt the fallthrough statement
   /// @returns true if the statement was successfully emitted
@@ -120,25 +119,28 @@ class GeneratorImpl : public TextGenerator {
   /// @returns true if the function was emitted
   bool EmitFunction(ast::Function* func);
   /// Handles generating an identifier expression
+  /// @param out the output of the expression stream
   /// @param expr the identifier expression
   /// @returns true if the identifeir was emitted
-  bool EmitIdentifier(ast::IdentifierExpression* expr);
+  bool EmitIdentifier(std::ostream& out, ast::IdentifierExpression* expr);
   /// Handles an if statement
   /// @param stmt the statement to emit
   /// @returns true if the statement was successfully emitted
   bool EmitIf(ast::IfStatement* stmt);
   /// Handles generating constructor expressions
+  /// @param out the output of the expression stream
   /// @param expr the constructor expression
   /// @returns true if the expression was emitted
-  bool EmitConstructor(ast::ConstructorExpression* expr);
+  bool EmitConstructor(std::ostream& out, ast::ConstructorExpression* expr);
   /// Handles generating a discard statement
   /// @param stmt the discard statement
   /// @returns true if the statement was successfully emitted
   bool EmitDiscard(ast::DiscardStatement* stmt);
   /// Handles a literal
+  /// @param out the output of the expression stream
   /// @param lit the literal to emit
   /// @returns true if the literal was successfully emitted
-  bool EmitLiteral(ast::Literal* lit);
+  bool EmitLiteral(std::ostream& out, ast::Literal* lit);
   /// Handles a loop statement
   /// @param stmt the statement to emit
   /// @returns true if the statement was emtited
@@ -148,9 +150,11 @@ class GeneratorImpl : public TextGenerator {
   /// @returns true if the statement was emtited
   bool EmitForLoop(ast::ForLoopStatement* stmt);
   /// Handles a member accessor expression
+  /// @param out the output of the expression stream
   /// @param expr the member accessor expression
   /// @returns true if the member accessor was emitted
-  bool EmitMemberAccessor(ast::MemberAccessorExpression* expr);
+  bool EmitMemberAccessor(std::ostream& out,
+                          ast::MemberAccessorExpression* expr);
   /// Handles return statements
   /// @param stmt the statement to emit
   /// @returns true if the statement was successfully emitted
@@ -159,47 +163,58 @@ class GeneratorImpl : public TextGenerator {
   /// @param stmt the statement to emit
   /// @returns true if the statement was emitted
   bool EmitStatement(ast::Statement* stmt);
-  /// Emits a statement without an indentation or trailing semi-colon and
-  /// newline
-  /// @param stmt the statement to emit
-  /// @returns true if the statement was emitted
-  bool EmitRawStatement(ast::Statement* stmt);
+  /// Handles a statement list
+  /// @param stmts the statements to emit
+  /// @returns true if the statements were emitted
+  bool EmitStatements(const ast::StatementList& stmts);
+  /// Handles a statement list with an increased indentation
+  /// @param stmts the statements to emit
+  /// @returns true if the statements were emitted
+  bool EmitStatementsWithIndent(const ast::StatementList& stmts);
   /// Handles generating a switch statement
   /// @param stmt the statement to emit
   /// @returns true if the statement was emitted
   bool EmitSwitch(ast::SwitchStatement* stmt);
   /// Handles generating type
+  /// @param out the output of the expression stream
   /// @param type the type to generate
   /// @returns true if the type is emitted
-  bool EmitType(const ast::Type* type);
+  bool EmitType(std::ostream& out, const ast::Type* type);
   /// Handles generating a struct declaration
   /// @param str the struct
   /// @returns true if the struct is emitted
   bool EmitStructType(const ast::Struct* str);
   /// Handles emitting an image format
+  /// @param out the output of the expression stream
   /// @param fmt the format to generate
   /// @returns true if the format is emitted
-  bool EmitImageFormat(const ast::ImageFormat fmt);
+  bool EmitImageFormat(std::ostream& out, const ast::ImageFormat fmt);
   /// Handles emitting an access control
+  /// @param out the output of the expression stream
   /// @param access the access to generate
   /// @returns true if the access is emitted
-  bool EmitAccess(const ast::Access access);
+  bool EmitAccess(std::ostream& out, const ast::Access access);
   /// Handles emitting a type constructor
+  /// @param out the output of the expression stream
   /// @param expr the type constructor expression
   /// @returns true if the constructor is emitted
-  bool EmitTypeConstructor(ast::TypeConstructorExpression* expr);
+  bool EmitTypeConstructor(std::ostream& out,
+                           ast::TypeConstructorExpression* expr);
   /// Handles a unary op expression
+  /// @param out the output of the expression stream
   /// @param expr the expression to emit
   /// @returns true if the expression was emitted
-  bool EmitUnaryOp(ast::UnaryOpExpression* expr);
+  bool EmitUnaryOp(std::ostream& out, ast::UnaryOpExpression* expr);
   /// Handles generating a variable
+  /// @param out the output of the expression stream
   /// @param var the variable to generate
   /// @returns true if the variable was emitted
-  bool EmitVariable(ast::Variable* var);
+  bool EmitVariable(std::ostream& out, ast::Variable* var);
   /// Handles generating a decoration list
+  /// @param out the output of the expression stream
   /// @param decos the decoration list
   /// @returns true if the decorations were emitted
-  bool EmitDecorations(const ast::DecorationList& decos);
+  bool EmitDecorations(std::ostream& out, const ast::DecorationList& decos);
 };
 
 }  // namespace wgsl
diff --git a/src/writer/wgsl/generator_impl_array_accessor_test.cc b/src/writer/wgsl/generator_impl_array_accessor_test.cc
index db18d16b83..4eff51a529 100644
--- a/src/writer/wgsl/generator_impl_array_accessor_test.cc
+++ b/src/writer/wgsl/generator_impl_array_accessor_test.cc
@@ -28,8 +28,9 @@ TEST_F(WgslGeneratorImplTest, ArrayAccessor) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitExpression(expr)) << gen.error();
-  EXPECT_EQ(gen.result(), "ary[5]");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(), "ary[5]");
 }
 
 TEST_F(WgslGeneratorImplTest, ArrayAccessor_OfDref) {
@@ -41,8 +42,9 @@ TEST_F(WgslGeneratorImplTest, ArrayAccessor_OfDref) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitExpression(expr)) << gen.error();
-  EXPECT_EQ(gen.result(), "(*(p))[5]");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(), "(*(p))[5]");
 }
 
 }  // namespace
diff --git a/src/writer/wgsl/generator_impl_binary_test.cc b/src/writer/wgsl/generator_impl_binary_test.cc
index 7c37d9b99d..3432f83cff 100644
--- a/src/writer/wgsl/generator_impl_binary_test.cc
+++ b/src/writer/wgsl/generator_impl_binary_test.cc
@@ -50,8 +50,9 @@ TEST_P(WgslBinaryTest, Emit) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitExpression(expr)) << gen.error();
-  EXPECT_EQ(gen.result(), params.result);
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(), params.result);
 }
 INSTANTIATE_TEST_SUITE_P(
     WgslGeneratorImplTest,
diff --git a/src/writer/wgsl/generator_impl_bitcast_test.cc b/src/writer/wgsl/generator_impl_bitcast_test.cc
index 9608196e31..31d0f6a6d0 100644
--- a/src/writer/wgsl/generator_impl_bitcast_test.cc
+++ b/src/writer/wgsl/generator_impl_bitcast_test.cc
@@ -27,8 +27,9 @@ TEST_F(WgslGeneratorImplTest, EmitExpression_Bitcast) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitExpression(bitcast)) << gen.error();
-  EXPECT_EQ(gen.result(), "bitcast<f32>(1)");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, bitcast)) << gen.error();
+  EXPECT_EQ(out.str(), "bitcast<f32>(1)");
 }
 
 }  // namespace
diff --git a/src/writer/wgsl/generator_impl_block_test.cc b/src/writer/wgsl/generator_impl_block_test.cc
index 27a0608dc3..c1bc7b5961 100644
--- a/src/writer/wgsl/generator_impl_block_test.cc
+++ b/src/writer/wgsl/generator_impl_block_test.cc
@@ -36,20 +36,6 @@ TEST_F(WgslGeneratorImplTest, Emit_Block) {
 )");
 }
 
-TEST_F(WgslGeneratorImplTest, Emit_Block_WithoutNewline) {
-  auto* b = Block(create<ast::DiscardStatement>());
-  WrapInFunction(b);
-
-  GeneratorImpl& gen = Build();
-
-  gen.increment_indent();
-
-  ASSERT_TRUE(gen.EmitBlock(b)) << gen.error();
-  EXPECT_EQ(gen.result(), R"({
-    discard;
-  })");
-}
-
 }  // namespace
 }  // namespace wgsl
 }  // namespace writer
diff --git a/src/writer/wgsl/generator_impl_call_test.cc b/src/writer/wgsl/generator_impl_call_test.cc
index ecce2726a8..42c360d365 100644
--- a/src/writer/wgsl/generator_impl_call_test.cc
+++ b/src/writer/wgsl/generator_impl_call_test.cc
@@ -31,8 +31,9 @@ TEST_F(WgslGeneratorImplTest, EmitExpression_Call_WithoutParams) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitExpression(call)) << gen.error();
-  EXPECT_EQ(gen.result(), "my_func()");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+  EXPECT_EQ(out.str(), "my_func()");
 }
 
 TEST_F(WgslGeneratorImplTest, EmitExpression_Call_WithParams) {
@@ -50,8 +51,9 @@ TEST_F(WgslGeneratorImplTest, EmitExpression_Call_WithParams) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitExpression(call)) << gen.error();
-  EXPECT_EQ(gen.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(WgslGeneratorImplTest, EmitStatement_Call) {
diff --git a/src/writer/wgsl/generator_impl_cast_test.cc b/src/writer/wgsl/generator_impl_cast_test.cc
index 7ecc446857..1af5088407 100644
--- a/src/writer/wgsl/generator_impl_cast_test.cc
+++ b/src/writer/wgsl/generator_impl_cast_test.cc
@@ -27,8 +27,9 @@ TEST_F(WgslGeneratorImplTest, EmitExpression_Cast_Scalar) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitExpression(cast)) << gen.error();
-  EXPECT_EQ(gen.result(), "f32(1)");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, cast)) << gen.error();
+  EXPECT_EQ(out.str(), "f32(1)");
 }
 
 TEST_F(WgslGeneratorImplTest, EmitExpression_Cast_Vector) {
@@ -37,8 +38,9 @@ TEST_F(WgslGeneratorImplTest, EmitExpression_Cast_Vector) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitExpression(cast)) << gen.error();
-  EXPECT_EQ(gen.result(), "vec3<f32>(vec3<i32>(1, 2, 3))");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, cast)) << gen.error();
+  EXPECT_EQ(out.str(), "vec3<f32>(vec3<i32>(1, 2, 3))");
 }
 
 }  // namespace
diff --git a/src/writer/wgsl/generator_impl_identifier_test.cc b/src/writer/wgsl/generator_impl_identifier_test.cc
index 2a66db11b6..58dbb27037 100644
--- a/src/writer/wgsl/generator_impl_identifier_test.cc
+++ b/src/writer/wgsl/generator_impl_identifier_test.cc
@@ -28,8 +28,9 @@ TEST_F(WgslGeneratorImplTest, EmitIdentifierExpression_Single) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitExpression(i)) << gen.error();
-  EXPECT_EQ(gen.result(), "glsl");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, i)) << gen.error();
+  EXPECT_EQ(out.str(), "glsl");
 }
 
 }  // namespace
diff --git a/src/writer/wgsl/generator_impl_literal_test.cc b/src/writer/wgsl/generator_impl_literal_test.cc
index 6ded4c8d40..e2244b1092 100644
--- a/src/writer/wgsl/generator_impl_literal_test.cc
+++ b/src/writer/wgsl/generator_impl_literal_test.cc
@@ -58,8 +58,9 @@ TEST_P(WgslGenerator_FloatLiteralTest, Emit) {
   SetResolveOnBuild(false);
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitScalarConstructor(v)) << gen.error();
-  EXPECT_EQ(gen.result(), GetParam().expected);
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitScalarConstructor(out, v)) << gen.error();
+  EXPECT_EQ(out.str(), GetParam().expected);
 }
 
 INSTANTIATE_TEST_SUITE_P(Zero,
diff --git a/src/writer/wgsl/generator_impl_member_accessor_test.cc b/src/writer/wgsl/generator_impl_member_accessor_test.cc
index 92ad0b8fca..787cd94707 100644
--- a/src/writer/wgsl/generator_impl_member_accessor_test.cc
+++ b/src/writer/wgsl/generator_impl_member_accessor_test.cc
@@ -30,8 +30,9 @@ TEST_F(WgslGeneratorImplTest, EmitExpression_MemberAccessor) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitExpression(expr)) << gen.error();
-  EXPECT_EQ(gen.result(), "str.mem");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(), "str.mem");
 }
 
 TEST_F(WgslGeneratorImplTest, EmitExpression_MemberAccessor_OfDref) {
@@ -44,8 +45,9 @@ TEST_F(WgslGeneratorImplTest, EmitExpression_MemberAccessor_OfDref) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitExpression(expr)) << gen.error();
-  EXPECT_EQ(gen.result(), "(*(p)).mem");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(), "(*(p)).mem");
 }
 
 }  // namespace
diff --git a/src/writer/wgsl/generator_impl_type_test.cc b/src/writer/wgsl/generator_impl_type_test.cc
index 3273e9206a..3c4d4c3e43 100644
--- a/src/writer/wgsl/generator_impl_type_test.cc
+++ b/src/writer/wgsl/generator_impl_type_test.cc
@@ -32,8 +32,9 @@ TEST_F(WgslGeneratorImplTest, EmitType_Alias) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitType(alias_ty)) << gen.error();
-  EXPECT_EQ(gen.result(), "alias");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, alias_ty)) << gen.error();
+  EXPECT_EQ(out.str(), "alias");
 }
 
 TEST_F(WgslGeneratorImplTest, EmitType_Array) {
@@ -42,8 +43,9 @@ TEST_F(WgslGeneratorImplTest, EmitType_Array) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitType(arr)) << gen.error();
-  EXPECT_EQ(gen.result(), "array<bool, 4>");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, arr)) << gen.error();
+  EXPECT_EQ(out.str(), "array<bool, 4>");
 }
 
 TEST_F(WgslGeneratorImplTest, EmitType_Array_Decoration) {
@@ -52,8 +54,9 @@ TEST_F(WgslGeneratorImplTest, EmitType_Array_Decoration) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitType(a)) << gen.error();
-  EXPECT_EQ(gen.result(), "[[stride(16)]] array<bool, 4>");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, a)) << gen.error();
+  EXPECT_EQ(out.str(), "[[stride(16)]] array<bool, 4>");
 }
 
 TEST_F(WgslGeneratorImplTest, EmitType_RuntimeArray) {
@@ -62,8 +65,9 @@ TEST_F(WgslGeneratorImplTest, EmitType_RuntimeArray) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitType(a)) << gen.error();
-  EXPECT_EQ(gen.result(), "array<bool>");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, a)) << gen.error();
+  EXPECT_EQ(out.str(), "array<bool>");
 }
 
 TEST_F(WgslGeneratorImplTest, EmitType_Bool) {
@@ -72,8 +76,9 @@ TEST_F(WgslGeneratorImplTest, EmitType_Bool) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitType(bool_)) << gen.error();
-  EXPECT_EQ(gen.result(), "bool");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, bool_)) << gen.error();
+  EXPECT_EQ(out.str(), "bool");
 }
 
 TEST_F(WgslGeneratorImplTest, EmitType_F32) {
@@ -82,8 +87,9 @@ TEST_F(WgslGeneratorImplTest, EmitType_F32) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitType(f32)) << gen.error();
-  EXPECT_EQ(gen.result(), "f32");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, f32)) << gen.error();
+  EXPECT_EQ(out.str(), "f32");
 }
 
 TEST_F(WgslGeneratorImplTest, EmitType_I32) {
@@ -92,8 +98,9 @@ TEST_F(WgslGeneratorImplTest, EmitType_I32) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitType(i32)) << gen.error();
-  EXPECT_EQ(gen.result(), "i32");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, i32)) << gen.error();
+  EXPECT_EQ(out.str(), "i32");
 }
 
 TEST_F(WgslGeneratorImplTest, EmitType_Matrix) {
@@ -102,8 +109,9 @@ TEST_F(WgslGeneratorImplTest, EmitType_Matrix) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitType(mat2x3)) << gen.error();
-  EXPECT_EQ(gen.result(), "mat2x3<f32>");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, mat2x3)) << gen.error();
+  EXPECT_EQ(out.str(), "mat2x3<f32>");
 }
 
 TEST_F(WgslGeneratorImplTest, EmitType_Pointer) {
@@ -112,8 +120,9 @@ TEST_F(WgslGeneratorImplTest, EmitType_Pointer) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitType(p)) << gen.error();
-  EXPECT_EQ(gen.result(), "ptr<workgroup, f32>");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, p)) << gen.error();
+  EXPECT_EQ(out.str(), "ptr<workgroup, f32>");
 }
 
 TEST_F(WgslGeneratorImplTest, EmitType_Struct) {
@@ -126,8 +135,9 @@ TEST_F(WgslGeneratorImplTest, EmitType_Struct) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitType(s_ty)) << gen.error();
-  EXPECT_EQ(gen.result(), "S");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, s_ty)) << gen.error();
+  EXPECT_EQ(out.str(), "S");
 }
 
 TEST_F(WgslGeneratorImplTest, EmitType_StructOffsetDecl) {
@@ -257,8 +267,9 @@ TEST_F(WgslGeneratorImplTest, EmitType_U32) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitType(u32)) << gen.error();
-  EXPECT_EQ(gen.result(), "u32");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, u32)) << gen.error();
+  EXPECT_EQ(out.str(), "u32");
 }
 
 TEST_F(WgslGeneratorImplTest, EmitType_Vector) {
@@ -267,8 +278,9 @@ TEST_F(WgslGeneratorImplTest, EmitType_Vector) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitType(vec3)) << gen.error();
-  EXPECT_EQ(gen.result(), "vec3<f32>");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, vec3)) << gen.error();
+  EXPECT_EQ(out.str(), "vec3<f32>");
 }
 
 struct TextureData {
@@ -289,8 +301,9 @@ TEST_P(WgslGenerator_DepthTextureTest, EmitType_DepthTexture) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitType(d)) << gen.error();
-  EXPECT_EQ(gen.result(), param.name);
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, d)) << gen.error();
+  EXPECT_EQ(out.str(), param.name);
 }
 INSTANTIATE_TEST_SUITE_P(
     WgslGeneratorImplTest,
@@ -311,8 +324,9 @@ TEST_P(WgslGenerator_SampledTextureTest, EmitType_SampledTexture_F32) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitType(t)) << gen.error();
-  EXPECT_EQ(gen.result(), std::string(param.name) + "<f32>");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, t)) << gen.error();
+  EXPECT_EQ(out.str(), std::string(param.name) + "<f32>");
 }
 
 TEST_P(WgslGenerator_SampledTextureTest, EmitType_SampledTexture_I32) {
@@ -323,8 +337,9 @@ TEST_P(WgslGenerator_SampledTextureTest, EmitType_SampledTexture_I32) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitType(t)) << gen.error();
-  EXPECT_EQ(gen.result(), std::string(param.name) + "<i32>");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, t)) << gen.error();
+  EXPECT_EQ(out.str(), std::string(param.name) + "<i32>");
 }
 
 TEST_P(WgslGenerator_SampledTextureTest, EmitType_SampledTexture_U32) {
@@ -335,8 +350,9 @@ TEST_P(WgslGenerator_SampledTextureTest, EmitType_SampledTexture_U32) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitType(t)) << gen.error();
-  EXPECT_EQ(gen.result(), std::string(param.name) + "<u32>");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, t)) << gen.error();
+  EXPECT_EQ(out.str(), std::string(param.name) + "<u32>");
 }
 INSTANTIATE_TEST_SUITE_P(
     WgslGeneratorImplTest,
@@ -358,8 +374,9 @@ TEST_P(WgslGenerator_MultiampledTextureTest, EmitType_MultisampledTexture_F32) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitType(t)) << gen.error();
-  EXPECT_EQ(gen.result(), std::string(param.name) + "<f32>");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, t)) << gen.error();
+  EXPECT_EQ(out.str(), std::string(param.name) + "<f32>");
 }
 
 TEST_P(WgslGenerator_MultiampledTextureTest, EmitType_MultisampledTexture_I32) {
@@ -370,8 +387,9 @@ TEST_P(WgslGenerator_MultiampledTextureTest, EmitType_MultisampledTexture_I32) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitType(t)) << gen.error();
-  EXPECT_EQ(gen.result(), std::string(param.name) + "<i32>");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, t)) << gen.error();
+  EXPECT_EQ(out.str(), std::string(param.name) + "<i32>");
 }
 
 TEST_P(WgslGenerator_MultiampledTextureTest, EmitType_MultisampledTexture_U32) {
@@ -382,8 +400,9 @@ TEST_P(WgslGenerator_MultiampledTextureTest, EmitType_MultisampledTexture_U32) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitType(t)) << gen.error();
-  EXPECT_EQ(gen.result(), std::string(param.name) + "<u32>");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, t)) << gen.error();
+  EXPECT_EQ(out.str(), std::string(param.name) + "<u32>");
 }
 INSTANTIATE_TEST_SUITE_P(WgslGeneratorImplTest,
                          WgslGenerator_MultiampledTextureTest,
@@ -414,8 +433,9 @@ TEST_P(WgslGenerator_StorageTextureTest, EmitType_StorageTexture) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitType(t)) << gen.error();
-  EXPECT_EQ(gen.result(), param.name);
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, t)) << gen.error();
+  EXPECT_EQ(out.str(), param.name);
 }
 INSTANTIATE_TEST_SUITE_P(
     WgslGeneratorImplTest,
@@ -460,8 +480,9 @@ TEST_P(WgslGenerator_ImageFormatTest, EmitType_StorageTexture_ImageFormat) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitImageFormat(param.fmt)) << gen.error();
-  EXPECT_EQ(gen.result(), param.name);
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitImageFormat(out, param.fmt)) << gen.error();
+  EXPECT_EQ(out.str(), param.name);
 }
 
 INSTANTIATE_TEST_SUITE_P(
@@ -510,8 +531,9 @@ TEST_F(WgslGeneratorImplTest, EmitType_Sampler) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitType(sampler)) << gen.error();
-  EXPECT_EQ(gen.result(), "sampler");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, sampler)) << gen.error();
+  EXPECT_EQ(out.str(), "sampler");
 }
 
 TEST_F(WgslGeneratorImplTest, EmitType_SamplerComparison) {
@@ -520,8 +542,9 @@ TEST_F(WgslGeneratorImplTest, EmitType_SamplerComparison) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitType(sampler)) << gen.error();
-  EXPECT_EQ(gen.result(), "sampler_comparison");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, sampler)) << gen.error();
+  EXPECT_EQ(out.str(), "sampler_comparison");
 }
 
 }  // namespace
diff --git a/src/writer/wgsl/generator_impl_unary_op_test.cc b/src/writer/wgsl/generator_impl_unary_op_test.cc
index 16b1a62a17..f771359b13 100644
--- a/src/writer/wgsl/generator_impl_unary_op_test.cc
+++ b/src/writer/wgsl/generator_impl_unary_op_test.cc
@@ -29,8 +29,9 @@ TEST_F(WgslUnaryOpTest, AddressOf) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitExpression(op)) << gen.error();
-  EXPECT_EQ(gen.result(), "&(expr)");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, op)) << gen.error();
+  EXPECT_EQ(out.str(), "&(expr)");
 }
 
 TEST_F(WgslUnaryOpTest, Complement) {
@@ -41,8 +42,9 @@ TEST_F(WgslUnaryOpTest, Complement) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitExpression(op)) << gen.error();
-  EXPECT_EQ(gen.result(), "~(expr)");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, op)) << gen.error();
+  EXPECT_EQ(out.str(), "~(expr)");
 }
 
 TEST_F(WgslUnaryOpTest, Indirection) {
@@ -56,8 +58,9 @@ TEST_F(WgslUnaryOpTest, Indirection) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitExpression(op)) << gen.error();
-  EXPECT_EQ(gen.result(), "*(expr)");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, op)) << gen.error();
+  EXPECT_EQ(out.str(), "*(expr)");
 }
 
 TEST_F(WgslUnaryOpTest, Not) {
@@ -67,8 +70,9 @@ TEST_F(WgslUnaryOpTest, Not) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitExpression(op)) << gen.error();
-  EXPECT_EQ(gen.result(), "!(expr)");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, op)) << gen.error();
+  EXPECT_EQ(out.str(), "!(expr)");
 }
 
 TEST_F(WgslUnaryOpTest, Negation) {
@@ -79,8 +83,9 @@ TEST_F(WgslUnaryOpTest, Negation) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitExpression(op)) << gen.error();
-  EXPECT_EQ(gen.result(), "-(expr)");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, op)) << gen.error();
+  EXPECT_EQ(out.str(), "-(expr)");
 }
 
 }  // namespace
diff --git a/src/writer/wgsl/generator_impl_variable_test.cc b/src/writer/wgsl/generator_impl_variable_test.cc
index 5cc340e37b..7e54432b92 100644
--- a/src/writer/wgsl/generator_impl_variable_test.cc
+++ b/src/writer/wgsl/generator_impl_variable_test.cc
@@ -28,8 +28,9 @@ TEST_F(WgslGeneratorImplTest, EmitVariable) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitVariable(v)) << gen.error();
-  EXPECT_EQ(gen.result(), R"(var<private> a : f32)");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitVariable(out, v)) << gen.error();
+  EXPECT_EQ(out.str(), R"(var<private> a : f32;)");
 }
 
 TEST_F(WgslGeneratorImplTest, EmitVariable_StorageClass) {
@@ -37,8 +38,9 @@ TEST_F(WgslGeneratorImplTest, EmitVariable_StorageClass) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitVariable(v)) << gen.error();
-  EXPECT_EQ(gen.result(), R"(var<private> a : f32)");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitVariable(out, v)) << gen.error();
+  EXPECT_EQ(out.str(), R"(var<private> a : f32;)");
 }
 
 TEST_F(WgslGeneratorImplTest, EmitVariable_Access_Read) {
@@ -53,9 +55,9 @@ TEST_F(WgslGeneratorImplTest, EmitVariable_Access_Read) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitVariable(v)) << gen.error();
-  EXPECT_EQ(gen.result(),
-            R"([[binding(0), group(0)]] var<storage, read> a : S)");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitVariable(out, v)) << gen.error();
+  EXPECT_EQ(out.str(), R"([[binding(0), group(0)]] var<storage, read> a : S;)");
 }
 
 TEST_F(WgslGeneratorImplTest, EmitVariable_Access_Write) {
@@ -70,9 +72,10 @@ TEST_F(WgslGeneratorImplTest, EmitVariable_Access_Write) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitVariable(v)) << gen.error();
-  EXPECT_EQ(gen.result(),
-            R"([[binding(0), group(0)]] var<storage, write> a : S)");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitVariable(out, v)) << gen.error();
+  EXPECT_EQ(out.str(),
+            R"([[binding(0), group(0)]] var<storage, write> a : S;)");
 }
 
 TEST_F(WgslGeneratorImplTest, EmitVariable_Access_ReadWrite) {
@@ -87,9 +90,10 @@ TEST_F(WgslGeneratorImplTest, EmitVariable_Access_ReadWrite) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitVariable(v)) << gen.error();
-  EXPECT_EQ(gen.result(),
-            R"([[binding(0), group(0)]] var<storage, read_write> a : S)");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitVariable(out, v)) << gen.error();
+  EXPECT_EQ(out.str(),
+            R"([[binding(0), group(0)]] var<storage, read_write> a : S;)");
 }
 
 TEST_F(WgslGeneratorImplTest, EmitVariable_Decorated) {
@@ -102,8 +106,9 @@ TEST_F(WgslGeneratorImplTest, EmitVariable_Decorated) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitVariable(v)) << gen.error();
-  EXPECT_EQ(gen.result(), R"([[group(1), binding(2)]] var a : sampler)");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitVariable(out, v)) << gen.error();
+  EXPECT_EQ(out.str(), R"([[group(1), binding(2)]] var a : sampler;)");
 }
 
 TEST_F(WgslGeneratorImplTest, EmitVariable_Constructor) {
@@ -111,8 +116,9 @@ TEST_F(WgslGeneratorImplTest, EmitVariable_Constructor) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitVariable(v)) << gen.error();
-  EXPECT_EQ(gen.result(), R"(var<private> a : f32 = 1.0)");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitVariable(out, v)) << gen.error();
+  EXPECT_EQ(out.str(), R"(var<private> a : f32 = 1.0;)");
 }
 
 TEST_F(WgslGeneratorImplTest, EmitVariable_Const) {
@@ -121,8 +127,9 @@ TEST_F(WgslGeneratorImplTest, EmitVariable_Const) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_TRUE(gen.EmitVariable(v)) << gen.error();
-  EXPECT_EQ(gen.result(), R"(let a : f32 = 1.0)");
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitVariable(out, v)) << gen.error();
+  EXPECT_EQ(out.str(), R"(let a : f32 = 1.0;)");
 }
 
 }  // namespace