diff --git a/include/tint/.clang-format b/include/tint/.clang-format
new file mode 100644
index 0000000..2fb833a
--- /dev/null
+++ b/include/tint/.clang-format
@@ -0,0 +1,2 @@
+# http://clang.llvm.org/docs/ClangFormatStyleOptions.html
+BasedOnStyle: Chromium
diff --git a/include/tint/tint.h b/include/tint/tint.h
index 1a04196..2b8430e 100644
--- a/include/tint/tint.h
+++ b/include/tint/tint.h
@@ -61,7 +61,6 @@
 #endif  // TINT_BUILD_HLSL_WRITER
 
 #if TINT_BUILD_GLSL_WRITER
-#include "src/tint/transform/glsl.h"
 #include "src/tint/writer/glsl/generator.h"
 #endif  // TINT_BUILD_GLSL_WRITER
 
diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index 79beb3d..5594bf8 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -688,6 +688,8 @@
     "writer/spirv/function.h",
     "writer/spirv/generator.cc",
     "writer/spirv/generator.h",
+    "writer/spirv/generator_impl.cc",
+    "writer/spirv/generator_impl.h",
     "writer/spirv/instruction.cc",
     "writer/spirv/instruction.h",
     "writer/spirv/operand.cc",
@@ -749,8 +751,6 @@
 
 libtint_source_set("libtint_glsl_writer_src") {
   sources = [
-    "transform/glsl.cc",
-    "transform/glsl.h",
     "writer/glsl/generator.cc",
     "writer/glsl/generator.h",
     "writer/glsl/generator_impl.cc",
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index 046b814..527eeaf 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -339,8 +339,6 @@
   transform/for_loop_to_loop.h
   transform/expand_compound_assignment.cc
   transform/expand_compound_assignment.h
-  transform/glsl.cc
-  transform/glsl.h
   transform/loop_to_for_loop.cc
   transform/loop_to_for_loop.h
   transform/manager.cc
@@ -519,6 +517,8 @@
     writer/spirv/function.h
     writer/spirv/generator.cc
     writer/spirv/generator.h
+    writer/spirv/generator_impl.cc
+    writer/spirv/generator_impl.h
     writer/spirv/instruction.cc
     writer/spirv/instruction.h
     writer/spirv/operand.cc
diff --git a/src/tint/fuzzers/BUILD.gn b/src/tint/fuzzers/BUILD.gn
index 3253e94..afd1e3a 100644
--- a/src/tint/fuzzers/BUILD.gn
+++ b/src/tint/fuzzers/BUILD.gn
@@ -22,14 +22,16 @@
   import("//testing/libfuzzer/fuzzer_test.gni")
 
   fuzzer_corpus_wgsl_dir = "${target_gen_dir}/fuzzer_corpus_wgsl"
+  fuzzer_corpus_wgsl_stamp = "${fuzzer_corpus_wgsl_dir}.stamp"
   action("tint_generate_wgsl_corpus") {
     script = "generate_wgsl_corpus.py"
     sources = [ "generate_wgsl_corpus.py" ]
     args = [
+      "--stamp=${fuzzer_corpus_wgsl_stamp}",
       rebase_path("${tint_root_dir}/test", root_build_dir),
       rebase_path(fuzzer_corpus_wgsl_dir, root_build_dir),
     ]
-    outputs = [ fuzzer_corpus_wgsl_dir ]
+    outputs = [ fuzzer_corpus_wgsl_stamp ]
   }
 
   tint_fuzzer_common_libfuzzer_options = [
diff --git a/src/tint/fuzzers/generate_wgsl_corpus.py b/src/tint/fuzzers/generate_wgsl_corpus.py
index 65c564f..983acd9 100644
--- a/src/tint/fuzzers/generate_wgsl_corpus.py
+++ b/src/tint/fuzzers/generate_wgsl_corpus.py
@@ -26,6 +26,7 @@
 # Usage:
 #    generate_wgsl_corpus.py <input_dir> <corpus_dir>
 
+import optparse
 import os
 import pathlib
 import shutil
@@ -40,11 +41,14 @@
 
 
 def main():
-    if len(sys.argv) != 3:
-        print("Usage: " + sys.argv[0] + " <input dir> <output dir>")
-        return 1
-    input_dir: str = os.path.abspath(sys.argv[1].rstrip(os.sep))
-    corpus_dir: str = os.path.abspath(sys.argv[2])
+    parser = optparse.OptionParser(
+        usage="usage: %prog [option] input-dir output-dir")
+    parser.add_option('--stamp', dest='stamp', help='stamp file')
+    options, args = parser.parse_args(sys.argv[1:])
+    if len(args) != 2:
+        parser.error("incorrect number of arguments")
+    input_dir: str = os.path.abspath(args[0].rstrip(os.sep))
+    corpus_dir: str = os.path.abspath(args[1])
     if os.path.exists(corpus_dir):
         shutil.rmtree(corpus_dir)
     os.makedirs(corpus_dir)
@@ -53,6 +57,8 @@
             continue
         out_file = in_file[len(input_dir) + 1:].replace(os.sep, '_')
         shutil.copy(in_file, corpus_dir + os.sep + out_file)
+    if options.stamp:
+        pathlib.Path(options.stamp).touch(mode=0o644, exist_ok=True)
 
 
 if __name__ == "__main__":
diff --git a/src/tint/fuzzers/tint_regex_fuzzer/CPPLINT.cfg b/src/tint/fuzzers/tint_regex_fuzzer/CPPLINT.cfg
deleted file mode 100644
index 96988ad..0000000
--- a/src/tint/fuzzers/tint_regex_fuzzer/CPPLINT.cfg
+++ /dev/null
@@ -1 +0,0 @@
-filter=-build/c++11
\ No newline at end of file
diff --git a/src/tint/fuzzers/tint_regex_fuzzer/wgsl_mutator.cc b/src/tint/fuzzers/tint_regex_fuzzer/wgsl_mutator.cc
index 2397e7f..d507944 100644
--- a/src/tint/fuzzers/tint_regex_fuzzer/wgsl_mutator.cc
+++ b/src/tint/fuzzers/tint_regex_fuzzer/wgsl_mutator.cc
@@ -17,7 +17,7 @@
 #include <cassert>
 #include <cstring>
 #include <map>
-#include <regex>
+#include <regex>  // NOLINT(build/c++11)
 #include <string>
 #include <utility>
 #include <vector>
diff --git a/src/tint/transform/glsl.cc b/src/tint/transform/glsl.cc
deleted file mode 100644
index f0120e8..0000000
--- a/src/tint/transform/glsl.cc
+++ /dev/null
@@ -1,145 +0,0 @@
-// Copyright 2021 The Tint Authors.
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-//     http://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "src/tint/transform/glsl.h"
-
-#include <utility>
-
-#include "src/tint/program_builder.h"
-#include "src/tint/transform/add_empty_entry_point.h"
-#include "src/tint/transform/add_spirv_block_attribute.h"
-#include "src/tint/transform/binding_remapper.h"
-#include "src/tint/transform/builtin_polyfill.h"
-#include "src/tint/transform/canonicalize_entry_point_io.h"
-#include "src/tint/transform/combine_samplers.h"
-#include "src/tint/transform/decompose_memory_access.h"
-#include "src/tint/transform/expand_compound_assignment.h"
-#include "src/tint/transform/fold_trivial_single_use_lets.h"
-#include "src/tint/transform/loop_to_for_loop.h"
-#include "src/tint/transform/manager.h"
-#include "src/tint/transform/promote_initializers_to_const_var.h"
-#include "src/tint/transform/promote_side_effects_to_decl.h"
-#include "src/tint/transform/remove_phonies.h"
-#include "src/tint/transform/renamer.h"
-#include "src/tint/transform/simplify_pointers.h"
-#include "src/tint/transform/single_entry_point.h"
-#include "src/tint/transform/unshadow.h"
-#include "src/tint/transform/unwind_discard_functions.h"
-#include "src/tint/transform/zero_init_workgroup_memory.h"
-#include "src/tint/writer/generate_external_texture_bindings.h"
-
-TINT_INSTANTIATE_TYPEINFO(tint::transform::Glsl);
-TINT_INSTANTIATE_TYPEINFO(tint::transform::Glsl::Config);
-
-namespace tint::transform {
-
-Glsl::Glsl() = default;
-Glsl::~Glsl() = default;
-
-Output Glsl::Run(const Program* in, const DataMap& inputs) const {
-  Manager manager;
-  DataMap data;
-
-  auto* cfg = inputs.Get<Config>();
-
-  {  // Builtin polyfills
-    BuiltinPolyfill::Builtins polyfills;
-    polyfills.count_leading_zeros = true;
-    polyfills.count_trailing_zeros = true;
-    polyfills.extract_bits = BuiltinPolyfill::Level::kClampParameters;
-    polyfills.first_leading_bit = true;
-    polyfills.first_trailing_bit = true;
-    polyfills.insert_bits = BuiltinPolyfill::Level::kClampParameters;
-    data.Add<BuiltinPolyfill::Config>(polyfills);
-    manager.Add<BuiltinPolyfill>();
-  }
-
-  if (cfg && !cfg->entry_point.empty()) {
-    manager.Add<SingleEntryPoint>();
-    data.Add<SingleEntryPoint::Config>(cfg->entry_point);
-  }
-  manager.Add<Renamer>();
-  data.Add<Renamer::Config>(Renamer::Target::kGlslKeywords,
-                            /* preserve_unicode */ false);
-  manager.Add<Unshadow>();
-
-  // Attempt to convert `loop`s into for-loops. This is to try and massage the
-  // output into something that will not cause FXC to choke or misbehave.
-  manager.Add<FoldTrivialSingleUseLets>();
-  manager.Add<LoopToForLoop>();
-
-  if (!cfg || !cfg->disable_workgroup_init) {
-    // ZeroInitWorkgroupMemory must come before CanonicalizeEntryPointIO as
-    // ZeroInitWorkgroupMemory may inject new builtin parameters.
-    manager.Add<ZeroInitWorkgroupMemory>();
-  }
-  manager.Add<CanonicalizeEntryPointIO>();
-  manager.Add<ExpandCompoundAssignment>();
-  manager.Add<PromoteSideEffectsToDecl>();
-  manager.Add<UnwindDiscardFunctions>();
-  manager.Add<SimplifyPointers>();
-
-  manager.Add<RemovePhonies>();
-
-  if (cfg && cfg->generate_external_texture_bindings) {
-    auto new_bindings_map = writer::GenerateExternalTextureBindings(in);
-    data.Add<MultiplanarExternalTexture::NewBindingPoints>(new_bindings_map);
-  }
-  manager.Add<MultiplanarExternalTexture>();
-
-  manager.Add<CombineSamplers>();
-  if (auto* binding_info = inputs.Get<CombineSamplers::BindingInfo>()) {
-    data.Add<CombineSamplers::BindingInfo>(*binding_info);
-  } else {
-    data.Add<CombineSamplers::BindingInfo>(CombineSamplers::BindingMap(),
-                                           sem::BindingPoint());
-  }
-  manager.Add<BindingRemapper>();
-  if (auto* remappings = inputs.Get<BindingRemapper::Remappings>()) {
-    data.Add<BindingRemapper::Remappings>(*remappings);
-  } else {
-    BindingRemapper::BindingPoints bp;
-    BindingRemapper::AccessControls ac;
-    data.Add<BindingRemapper::Remappings>(bp, ac, /* mayCollide */ true);
-  }
-
-  manager.Add<PromoteInitializersToConstVar>();
-
-  manager.Add<AddEmptyEntryPoint>();
-  manager.Add<AddSpirvBlockAttribute>();
-
-  data.Add<CanonicalizeEntryPointIO::Config>(
-      CanonicalizeEntryPointIO::ShaderStyle::kGlsl);
-  auto out = manager.Run(in, data);
-  if (!out.program.IsValid()) {
-    return out;
-  }
-
-  ProgramBuilder builder;
-  CloneContext ctx(&builder, &out.program);
-  ctx.Clone();
-  return Output{Program(std::move(builder))};
-}
-
-Glsl::Config::Config(const std::string& entry_point_in,
-                     bool disable_workgroup_init_in,
-                     bool generate_external_texture_bindings_in)
-    : entry_point(entry_point_in),
-      disable_workgroup_init(disable_workgroup_init_in),
-      generate_external_texture_bindings(
-          generate_external_texture_bindings_in) {}
-Glsl::Config::Config(const Config&) = default;
-Glsl::Config::~Config() = default;
-
-}  // namespace tint::transform
diff --git a/src/tint/transform/glsl.h b/src/tint/transform/glsl.h
deleted file mode 100644
index c84ff29..0000000
--- a/src/tint/transform/glsl.h
+++ /dev/null
@@ -1,75 +0,0 @@
-// Copyright 2021 The Tint Authors.
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-//     http://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#ifndef SRC_TINT_TRANSFORM_GLSL_H_
-#define SRC_TINT_TRANSFORM_GLSL_H_
-
-#include <string>
-
-#include "src/tint/transform/transform.h"
-
-// Forward declarations
-namespace tint {
-class CloneContext;
-}  // namespace tint
-
-namespace tint::transform {
-
-/// Glsl is a transform used to sanitize a Program for use with the Glsl writer.
-/// Passing a non-sanitized Program to the Glsl writer will result in undefined
-/// behavior.
-class Glsl final : public Castable<Glsl, Transform> {
- public:
-  /// Configuration options for the Glsl sanitizer transform.
-  struct Config final : public Castable<Data, transform::Data> {
-    /// Constructor
-    /// @param entry_point the root entry point function to generate
-    /// @param disable_workgroup_init `true` to disable workgroup memory zero
-    ///        initialization
-    /// @param generate_external_texture_bindings 'true' to generates binding
-    /// mappings for external textures
-    explicit Config(const std::string& entry_point,
-                    bool disable_workgroup_init,
-                    bool generate_external_texture_bindings);
-
-    /// Copy constructor
-    Config(const Config&);
-
-    /// Destructor
-    ~Config() override;
-
-    /// GLSL generator wraps a single entry point in a main() function.
-    std::string entry_point;
-
-    /// Set to `true` to disable workgroup memory zero initialization
-    bool disable_workgroup_init = false;
-
-    /// Set to 'true' to generates binding mappings for external textures
-    bool generate_external_texture_bindings = false;
-  };
-
-  /// Constructor
-  Glsl();
-  ~Glsl() override;
-
-  /// Runs the transform on `program`, returning the transformation result.
-  /// @param program the source program to transform
-  /// @param data optional extra transform-specific data
-  /// @returns the transformation result
-  Output Run(const Program* program, const DataMap& data = {}) const override;
-};
-
-}  // namespace tint::transform
-
-#endif  // SRC_TINT_TRANSFORM_GLSL_H_
diff --git a/src/tint/transform/glsl_test.cc b/src/tint/transform/glsl_test.cc
deleted file mode 100644
index 555701b..0000000
--- a/src/tint/transform/glsl_test.cc
+++ /dev/null
@@ -1,39 +0,0 @@
-// Copyright 2021 The Tint Authors.
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-//     http://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "src/tint/transform/glsl.h"
-
-#include "src/tint/transform/test_helper.h"
-
-namespace tint::transform {
-namespace {
-
-using GlslTest = TransformTest;
-
-TEST_F(GlslTest, AddEmptyEntryPoint) {
-  auto* src = R"()";
-
-  auto* expect = R"(
-@stage(compute) @workgroup_size(1)
-fn unused_entry_point() {
-}
-)";
-
-  auto got = Run<Glsl>(src);
-
-  EXPECT_EQ(expect, str(got));
-}
-
-}  // namespace
-}  // namespace tint::transform
diff --git a/src/tint/writer/glsl/generator.cc b/src/tint/writer/glsl/generator.cc
index ad06ac4..83ae8c4 100644
--- a/src/tint/writer/glsl/generator.cc
+++ b/src/tint/writer/glsl/generator.cc
@@ -16,7 +16,6 @@
 
 #include "src/tint/transform/binding_remapper.h"
 #include "src/tint/transform/combine_samplers.h"
-#include "src/tint/transform/glsl.h"
 #include "src/tint/writer/glsl/generator_impl.h"
 
 namespace tint::writer::glsl {
@@ -34,34 +33,25 @@
                 const std::string& entry_point) {
   Result result;
 
-  // Run the GLSL sanitizer.
-  transform::DataMap data;
-  data.Add<transform::BindingRemapper::Remappings>(options.binding_points,
-                                                   options.access_controls,
-                                                   options.allow_collisions);
-  data.Add<transform::CombineSamplers::BindingInfo>(
-      options.binding_map, options.placeholder_binding_point);
-  data.Add<transform::Glsl::Config>(entry_point,
-                                    /* disable_workgroup_init */ false,
-                                    options.generate_external_texture_bindings);
-  transform::Glsl sanitizer;
-  auto output = sanitizer.Run(program, data);
-  if (!output.program.IsValid()) {
+  // Sanitize the program.
+  auto sanitized_result = Sanitize(program, options, entry_point);
+  if (!sanitized_result.program.IsValid()) {
     result.success = false;
-    result.error = output.program.Diagnostics().str();
+    result.error = sanitized_result.program.Diagnostics().str();
     return result;
   }
 
   // Generate the GLSL code.
-  auto impl = std::make_unique<GeneratorImpl>(&output.program, options.version);
+  auto impl = std::make_unique<GeneratorImpl>(&sanitized_result.program,
+                                              options.version);
   result.success = impl->Generate();
   result.error = impl->error();
   result.glsl = impl->result();
 
   // Collect the list of entry points in the sanitized program.
-  for (auto* func : output.program.AST().Functions()) {
+  for (auto* func : sanitized_result.program.AST().Functions()) {
     if (func->IsEntryPoint()) {
-      auto name = output.program.Symbols().NameFor(func->symbol);
+      auto name = sanitized_result.program.Symbols().NameFor(func->symbol);
       result.entry_points.push_back({name, func->PipelineStage()});
     }
   }
diff --git a/src/tint/writer/glsl/generator.h b/src/tint/writer/glsl/generator.h
index ab5c9bb..11206cf 100644
--- a/src/tint/writer/glsl/generator.h
+++ b/src/tint/writer/glsl/generator.h
@@ -70,6 +70,9 @@
   /// generated by the BindingRemapper transform
   bool allow_collisions = false;
 
+  /// Set to `true` to disable workgroup memory zero initialization
+  bool disable_workgroup_init = false;
+
   /// Set to 'true' to generates binding mappings for external textures
   bool generate_external_texture_bindings = false;
 
@@ -106,6 +109,7 @@
 /// information.
 /// @param program the program to translate to GLSL
 /// @param options the configuration options to use when generating GLSL
+/// @param entry_point the entry point to generate GLSL for
 /// @returns the resulting GLSL and supplementary information
 Result Generate(const Program* program,
                 const Options& options,
diff --git a/src/tint/writer/glsl/generator_impl.cc b/src/tint/writer/glsl/generator_impl.cc
index 729a8c0..ef7467d 100644
--- a/src/tint/writer/glsl/generator_impl.cc
+++ b/src/tint/writer/glsl/generator_impl.cc
@@ -45,12 +45,32 @@
 #include "src/tint/sem/type_constructor.h"
 #include "src/tint/sem/type_conversion.h"
 #include "src/tint/sem/variable.h"
-#include "src/tint/transform/glsl.h"
+#include "src/tint/transform/add_empty_entry_point.h"
+#include "src/tint/transform/add_spirv_block_attribute.h"
+#include "src/tint/transform/binding_remapper.h"
+#include "src/tint/transform/builtin_polyfill.h"
+#include "src/tint/transform/canonicalize_entry_point_io.h"
+#include "src/tint/transform/combine_samplers.h"
+#include "src/tint/transform/decompose_memory_access.h"
+#include "src/tint/transform/expand_compound_assignment.h"
+#include "src/tint/transform/fold_trivial_single_use_lets.h"
+#include "src/tint/transform/loop_to_for_loop.h"
+#include "src/tint/transform/manager.h"
+#include "src/tint/transform/promote_initializers_to_const_var.h"
+#include "src/tint/transform/promote_side_effects_to_decl.h"
+#include "src/tint/transform/remove_phonies.h"
+#include "src/tint/transform/renamer.h"
+#include "src/tint/transform/simplify_pointers.h"
+#include "src/tint/transform/single_entry_point.h"
+#include "src/tint/transform/unshadow.h"
+#include "src/tint/transform/unwind_discard_functions.h"
+#include "src/tint/transform/zero_init_workgroup_memory.h"
 #include "src/tint/utils/defer.h"
 #include "src/tint/utils/map.h"
 #include "src/tint/utils/scoped_assignment.h"
 #include "src/tint/writer/append_vector.h"
 #include "src/tint/writer/float_to_string.h"
+#include "src/tint/writer/generate_external_texture_bindings.h"
 
 namespace {
 
@@ -127,6 +147,86 @@
 
 }  // namespace
 
+SanitizedResult::SanitizedResult() = default;
+SanitizedResult::~SanitizedResult() = default;
+SanitizedResult::SanitizedResult(SanitizedResult&&) = default;
+
+SanitizedResult Sanitize(const Program* in,
+                         const Options& options,
+                         const std::string& entry_point) {
+  transform::Manager manager;
+  transform::DataMap data;
+
+  {  // Builtin polyfills
+    transform::BuiltinPolyfill::Builtins polyfills;
+    polyfills.count_leading_zeros = true;
+    polyfills.count_trailing_zeros = true;
+    polyfills.extract_bits =
+        transform::BuiltinPolyfill::Level::kClampParameters;
+    polyfills.first_leading_bit = true;
+    polyfills.first_trailing_bit = true;
+    polyfills.insert_bits = transform::BuiltinPolyfill::Level::kClampParameters;
+    data.Add<transform::BuiltinPolyfill::Config>(polyfills);
+    manager.Add<transform::BuiltinPolyfill>();
+  }
+
+  if (!entry_point.empty()) {
+    manager.Add<transform::SingleEntryPoint>();
+    data.Add<transform::SingleEntryPoint::Config>(entry_point);
+  }
+  manager.Add<transform::Renamer>();
+  data.Add<transform::Renamer::Config>(
+      transform::Renamer::Target::kGlslKeywords,
+      /* preserve_unicode */ false);
+  manager.Add<transform::Unshadow>();
+
+  // Attempt to convert `loop`s into for-loops. This is to try and massage the
+  // output into something that will not cause FXC to choke or misbehave.
+  manager.Add<transform::FoldTrivialSingleUseLets>();
+  manager.Add<transform::LoopToForLoop>();
+
+  if (!options.disable_workgroup_init) {
+    // ZeroInitWorkgroupMemory must come before CanonicalizeEntryPointIO as
+    // ZeroInitWorkgroupMemory may inject new builtin parameters.
+    manager.Add<transform::ZeroInitWorkgroupMemory>();
+  }
+  manager.Add<transform::CanonicalizeEntryPointIO>();
+  manager.Add<transform::ExpandCompoundAssignment>();
+  manager.Add<transform::PromoteSideEffectsToDecl>();
+  manager.Add<transform::UnwindDiscardFunctions>();
+  manager.Add<transform::SimplifyPointers>();
+
+  manager.Add<transform::RemovePhonies>();
+
+  if (options.generate_external_texture_bindings) {
+    auto new_bindings_map = writer::GenerateExternalTextureBindings(in);
+    data.Add<transform::MultiplanarExternalTexture::NewBindingPoints>(
+        new_bindings_map);
+  }
+  manager.Add<transform::MultiplanarExternalTexture>();
+
+  data.Add<transform::CombineSamplers::BindingInfo>(
+      options.binding_map, options.placeholder_binding_point);
+  manager.Add<transform::CombineSamplers>();
+
+  data.Add<transform::BindingRemapper::Remappings>(options.binding_points,
+                                                   options.access_controls,
+                                                   options.allow_collisions);
+  manager.Add<transform::BindingRemapper>();
+
+  manager.Add<transform::PromoteInitializersToConstVar>();
+  manager.Add<transform::AddEmptyEntryPoint>();
+  manager.Add<transform::AddSpirvBlockAttribute>();
+  data.Add<transform::CanonicalizeEntryPointIO::Config>(
+      transform::CanonicalizeEntryPointIO::ShaderStyle::kGlsl);
+
+  auto out = manager.Run(in, data);
+
+  SanitizedResult result;
+  result.program = std::move(out.program);
+  return result;
+}
+
 GeneratorImpl::GeneratorImpl(const Program* program, const Version& version)
     : TextGenerator(program), version_(version) {}
 
diff --git a/src/tint/writer/glsl/generator_impl.h b/src/tint/writer/glsl/generator_impl.h
index 37be658..9104182 100644
--- a/src/tint/writer/glsl/generator_impl.h
+++ b/src/tint/writer/glsl/generator_impl.h
@@ -35,6 +35,7 @@
 #include "src/tint/scope_stack.h"
 #include "src/tint/transform/decompose_memory_access.h"
 #include "src/tint/utils/hash.h"
+#include "src/tint/writer/glsl/generator.h"
 #include "src/tint/writer/glsl/version.h"
 #include "src/tint/writer/text_generator.h"
 
@@ -48,6 +49,28 @@
 
 namespace tint::writer::glsl {
 
+/// The result of sanitizing a program for generation.
+struct SanitizedResult {
+  /// Constructor
+  SanitizedResult();
+  /// Destructor
+  ~SanitizedResult();
+  /// Move constructor
+  SanitizedResult(SanitizedResult&&);
+
+  /// The sanitized program.
+  Program program;
+};
+
+/// Sanitize a program in preparation for generating GLSL.
+/// @program The program to sanitize
+/// @param options The HLSL generator options.
+/// @param entry_point the entry point to generate GLSL for
+/// @returns the sanitized program and any supplementary information
+SanitizedResult Sanitize(const Program* program,
+                         const Options& options,
+                         const std::string& entry_point);
+
 /// Implementation class for GLSL generator
 class GeneratorImpl : public TextGenerator {
  public:
diff --git a/src/tint/writer/glsl/test_helper.h b/src/tint/writer/glsl/test_helper.h
index 969490a..7266d0a 100644
--- a/src/tint/writer/glsl/test_helper.h
+++ b/src/tint/writer/glsl/test_helper.h
@@ -20,7 +20,6 @@
 #include <utility>
 
 #include "gtest/gtest.h"
-#include "src/tint/transform/glsl.h"
 #include "src/tint/transform/manager.h"
 #include "src/tint/writer/glsl/generator_impl.h"
 
@@ -60,8 +59,10 @@
   /// @note The generator is only built once. Multiple calls to Build() will
   /// return the same GeneratorImpl without rebuilding.
   /// @param version the GLSL version
+  /// @param options the GLSL backend options
   /// @return the built generator
-  GeneratorImpl& SanitizeAndBuild(Version version = Version()) {
+  GeneratorImpl& SanitizeAndBuild(Version version = Version(),
+                                  const Options& options = {}) {
     if (gen_) {
       return *gen_;
     }
@@ -76,15 +77,14 @@
           << formatter.format(program->Diagnostics());
     }();
 
-    transform::Manager transform_manager;
-    transform::DataMap transform_data;
-    transform_manager.Add<tint::transform::Glsl>();
-    auto result = transform_manager.Run(program.get(), transform_data);
+    auto sanitized_result =
+        Sanitize(program.get(), options, /* entry_point */ "");
     [&]() {
-      ASSERT_TRUE(result.program.IsValid())
-          << formatter.format(result.program.Diagnostics());
+      ASSERT_TRUE(sanitized_result.program.IsValid())
+          << formatter.format(sanitized_result.program.Diagnostics());
     }();
-    *program = std::move(result.program);
+
+    *program = std::move(sanitized_result.program);
     gen_ = std::make_unique<GeneratorImpl>(program.get(), version);
     return *gen_;
   }
diff --git a/src/tint/writer/spirv/binary_writer.h b/src/tint/writer/spirv/binary_writer.h
index 89789fd..a071d7c 100644
--- a/src/tint/writer/spirv/binary_writer.h
+++ b/src/tint/writer/spirv/binary_writer.h
@@ -45,6 +45,9 @@
   /// @returns the assembled SPIR-V
   const std::vector<uint32_t>& result() const { return out_; }
 
+  /// @returns the assembled SPIR-V
+  std::vector<uint32_t>& result() { return out_; }
+
  private:
   void process_instruction(const Instruction& inst);
   void process_op(const Operand& op);
diff --git a/src/tint/writer/spirv/builder.cc b/src/tint/writer/spirv/builder.cc
index 436f977..c033344 100644
--- a/src/tint/writer/spirv/builder.cc
+++ b/src/tint/writer/spirv/builder.cc
@@ -41,26 +41,10 @@
 #include "src/tint/sem/type_conversion.h"
 #include "src/tint/sem/variable.h"
 #include "src/tint/sem/vector_type.h"
-#include "src/tint/transform/add_empty_entry_point.h"
 #include "src/tint/transform/add_spirv_block_attribute.h"
-#include "src/tint/transform/builtin_polyfill.h"
-#include "src/tint/transform/canonicalize_entry_point_io.h"
-#include "src/tint/transform/expand_compound_assignment.h"
-#include "src/tint/transform/fold_constants.h"
-#include "src/tint/transform/for_loop_to_loop.h"
-#include "src/tint/transform/manager.h"
-#include "src/tint/transform/promote_side_effects_to_decl.h"
-#include "src/tint/transform/remove_unreachable_statements.h"
-#include "src/tint/transform/simplify_pointers.h"
-#include "src/tint/transform/unshadow.h"
-#include "src/tint/transform/unwind_discard_functions.h"
-#include "src/tint/transform/var_for_dynamic_index.h"
-#include "src/tint/transform/vectorize_scalar_matrix_constructors.h"
-#include "src/tint/transform/zero_init_workgroup_memory.h"
 #include "src/tint/utils/defer.h"
 #include "src/tint/utils/map.h"
 #include "src/tint/writer/append_vector.h"
-#include "src/tint/writer/generate_external_texture_bindings.h"
 
 namespace tint::writer::spirv {
 namespace {
@@ -255,61 +239,6 @@
 
 }  // namespace
 
-SanitizedResult Sanitize(const Program* in, const Options& options) {
-  transform::Manager manager;
-  transform::DataMap data;
-
-  {  // Builtin polyfills
-    transform::BuiltinPolyfill::Builtins polyfills;
-    polyfills.count_leading_zeros = true;
-    polyfills.count_trailing_zeros = true;
-    polyfills.extract_bits =
-        transform::BuiltinPolyfill::Level::kClampParameters;
-    polyfills.first_leading_bit = true;
-    polyfills.first_trailing_bit = true;
-    polyfills.insert_bits = transform::BuiltinPolyfill::Level::kClampParameters;
-    data.Add<transform::BuiltinPolyfill::Config>(polyfills);
-    manager.Add<transform::BuiltinPolyfill>();
-  }
-
-  if (options.generate_external_texture_bindings) {
-    auto new_bindings_map = GenerateExternalTextureBindings(in);
-    data.Add<transform::MultiplanarExternalTexture::NewBindingPoints>(
-        new_bindings_map);
-  }
-  manager.Add<transform::MultiplanarExternalTexture>();
-
-  manager.Add<transform::Unshadow>();
-  bool disable_workgroup_init_in_sanitizer =
-      options.disable_workgroup_init ||
-      options.use_zero_initialize_workgroup_memory_extension;
-  if (!disable_workgroup_init_in_sanitizer) {
-    manager.Add<transform::ZeroInitWorkgroupMemory>();
-  }
-  manager.Add<transform::RemoveUnreachableStatements>();
-  manager.Add<transform::ExpandCompoundAssignment>();
-  manager.Add<transform::PromoteSideEffectsToDecl>();
-  manager.Add<transform::UnwindDiscardFunctions>();
-  manager.Add<transform::SimplifyPointers>();  // Required for arrayLength()
-  manager.Add<transform::FoldConstants>();
-  manager.Add<transform::VectorizeScalarMatrixConstructors>();
-  manager.Add<transform::ForLoopToLoop>();  // Must come after
-                                            // ZeroInitWorkgroupMemory
-  manager.Add<transform::CanonicalizeEntryPointIO>();
-  manager.Add<transform::AddEmptyEntryPoint>();
-  manager.Add<transform::AddSpirvBlockAttribute>();
-  manager.Add<transform::VarForDynamicIndex>();
-
-  data.Add<transform::CanonicalizeEntryPointIO::Config>(
-      transform::CanonicalizeEntryPointIO::Config(
-          transform::CanonicalizeEntryPointIO::ShaderStyle::kSpirv, 0xFFFFFFFF,
-          options.emit_vertex_point_size));
-
-  SanitizedResult result;
-  result.program = std::move(manager.Run(in, data).program);
-  return result;
-}
-
 Builder::AccessorInfo::AccessorInfo() : source_id(0), source_type(nullptr) {}
 
 Builder::AccessorInfo::~AccessorInfo() {}
diff --git a/src/tint/writer/spirv/builder.h b/src/tint/writer/spirv/builder.h
index c36efcc..7ad9cf2 100644
--- a/src/tint/writer/spirv/builder.h
+++ b/src/tint/writer/spirv/builder.h
@@ -38,7 +38,6 @@
 #include "src/tint/sem/builtin.h"
 #include "src/tint/sem/storage_texture_type.h"
 #include "src/tint/writer/spirv/function.h"
-#include "src/tint/writer/spirv/generator.h"
 #include "src/tint/writer/spirv/scalar_constant.h"
 
 // Forward declarations
@@ -51,17 +50,6 @@
 
 namespace tint::writer::spirv {
 
-/// The result of sanitizing a program for generation.
-struct SanitizedResult {
-  /// The sanitized program.
-  Program program;
-};
-
-/// Sanitize a program in preparation for generating SPIR-V.
-/// @program The program to sanitize
-/// @param options The SPIR-V generator options.
-SanitizedResult Sanitize(const Program* program, const Options& options);
-
 /// Builder class to create SPIR-V instructions from a module.
 class Builder {
  public:
diff --git a/src/tint/writer/spirv/generator.cc b/src/tint/writer/spirv/generator.cc
index 938f8f7..04f6e15 100644
--- a/src/tint/writer/spirv/generator.cc
+++ b/src/tint/writer/spirv/generator.cc
@@ -14,7 +14,9 @@
 
 #include "src/tint/writer/spirv/generator.h"
 
-#include "src/tint/writer/spirv/binary_writer.h"
+#include <utility>
+
+#include "src/tint/writer/spirv/generator_impl.h"
 
 namespace tint::writer::spirv {
 
@@ -37,20 +39,12 @@
   bool zero_initialize_workgroup_memory =
       !options.disable_workgroup_init &&
       options.use_zero_initialize_workgroup_memory_extension;
-  auto builder = std::make_unique<Builder>(&sanitized_result.program,
-                                           zero_initialize_workgroup_memory);
-  auto writer = std::make_unique<BinaryWriter>();
-  if (!builder->Build()) {
-    result.success = false;
-    result.error = builder->error();
-    return result;
-  }
 
-  writer->WriteHeader(builder->id_bound());
-  writer->WriteBuilder(builder.get());
-
-  result.success = true;
-  result.spirv = writer->result();
+  auto impl = std::make_unique<GeneratorImpl>(&sanitized_result.program,
+                                              zero_initialize_workgroup_memory);
+  result.success = impl->Generate();
+  result.error = impl->error();
+  result.spirv = std::move(impl->result());
 
   return result;
 }
diff --git a/src/tint/writer/spirv/generator_impl.cc b/src/tint/writer/spirv/generator_impl.cc
new file mode 100644
index 0000000..3318808
--- /dev/null
+++ b/src/tint/writer/spirv/generator_impl.cc
@@ -0,0 +1,120 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "src/tint/writer/spirv/generator_impl.h"
+
+#include <utility>
+#include <vector>
+
+#include "src/tint/transform/add_empty_entry_point.h"
+#include "src/tint/transform/add_spirv_block_attribute.h"
+#include "src/tint/transform/builtin_polyfill.h"
+#include "src/tint/transform/canonicalize_entry_point_io.h"
+#include "src/tint/transform/expand_compound_assignment.h"
+#include "src/tint/transform/fold_constants.h"
+#include "src/tint/transform/for_loop_to_loop.h"
+#include "src/tint/transform/manager.h"
+#include "src/tint/transform/promote_side_effects_to_decl.h"
+#include "src/tint/transform/remove_unreachable_statements.h"
+#include "src/tint/transform/simplify_pointers.h"
+#include "src/tint/transform/unshadow.h"
+#include "src/tint/transform/unwind_discard_functions.h"
+#include "src/tint/transform/var_for_dynamic_index.h"
+#include "src/tint/transform/vectorize_scalar_matrix_constructors.h"
+#include "src/tint/transform/zero_init_workgroup_memory.h"
+#include "src/tint/writer/generate_external_texture_bindings.h"
+
+namespace tint::writer::spirv {
+
+SanitizedResult Sanitize(const Program* in, const Options& options) {
+  transform::Manager manager;
+  transform::DataMap data;
+
+  {  // Builtin polyfills
+    transform::BuiltinPolyfill::Builtins polyfills;
+    polyfills.count_leading_zeros = true;
+    polyfills.count_trailing_zeros = true;
+    polyfills.extract_bits =
+        transform::BuiltinPolyfill::Level::kClampParameters;
+    polyfills.first_leading_bit = true;
+    polyfills.first_trailing_bit = true;
+    polyfills.insert_bits = transform::BuiltinPolyfill::Level::kClampParameters;
+    data.Add<transform::BuiltinPolyfill::Config>(polyfills);
+    manager.Add<transform::BuiltinPolyfill>();
+  }
+
+  if (options.generate_external_texture_bindings) {
+    auto new_bindings_map = GenerateExternalTextureBindings(in);
+    data.Add<transform::MultiplanarExternalTexture::NewBindingPoints>(
+        new_bindings_map);
+  }
+  manager.Add<transform::MultiplanarExternalTexture>();
+
+  manager.Add<transform::Unshadow>();
+  bool disable_workgroup_init_in_sanitizer =
+      options.disable_workgroup_init ||
+      options.use_zero_initialize_workgroup_memory_extension;
+  if (!disable_workgroup_init_in_sanitizer) {
+    manager.Add<transform::ZeroInitWorkgroupMemory>();
+  }
+  manager.Add<transform::RemoveUnreachableStatements>();
+  manager.Add<transform::ExpandCompoundAssignment>();
+  manager.Add<transform::PromoteSideEffectsToDecl>();
+  manager.Add<transform::UnwindDiscardFunctions>();
+  manager.Add<transform::SimplifyPointers>();  // Required for arrayLength()
+  manager.Add<transform::FoldConstants>();
+  manager.Add<transform::VectorizeScalarMatrixConstructors>();
+  manager.Add<transform::ForLoopToLoop>();  // Must come after
+                                            // ZeroInitWorkgroupMemory
+  manager.Add<transform::CanonicalizeEntryPointIO>();
+  manager.Add<transform::AddEmptyEntryPoint>();
+  manager.Add<transform::AddSpirvBlockAttribute>();
+  manager.Add<transform::VarForDynamicIndex>();
+
+  data.Add<transform::CanonicalizeEntryPointIO::Config>(
+      transform::CanonicalizeEntryPointIO::Config(
+          transform::CanonicalizeEntryPointIO::ShaderStyle::kSpirv, 0xFFFFFFFF,
+          options.emit_vertex_point_size));
+
+  SanitizedResult result;
+  result.program = std::move(manager.Run(in, data).program);
+  return result;
+}
+
+GeneratorImpl::GeneratorImpl(const Program* program,
+                             bool zero_initialize_workgroup_memory)
+    : builder_(program, zero_initialize_workgroup_memory) {}
+
+bool GeneratorImpl::Generate() {
+  if (builder_.Build()) {
+    writer_.WriteHeader(builder_.id_bound());
+    writer_.WriteBuilder(&builder_);
+    return true;
+  }
+  return false;
+}
+
+const std::vector<uint32_t>& GeneratorImpl::result() const {
+  return writer_.result();
+}
+
+std::vector<uint32_t>& GeneratorImpl::result() {
+  return writer_.result();
+}
+
+std::string GeneratorImpl::error() const {
+  return builder_.error();
+}
+
+}  // namespace tint::writer::spirv
diff --git a/src/tint/writer/spirv/generator_impl.h b/src/tint/writer/spirv/generator_impl.h
new file mode 100644
index 0000000..2e02af1
--- /dev/null
+++ b/src/tint/writer/spirv/generator_impl.h
@@ -0,0 +1,67 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef SRC_TINT_WRITER_SPIRV_GENERATOR_IMPL_H_
+#define SRC_TINT_WRITER_SPIRV_GENERATOR_IMPL_H_
+
+#include <string>
+#include <vector>
+
+#include "src/tint/program.h"
+#include "src/tint/writer/spirv/binary_writer.h"
+#include "src/tint/writer/spirv/builder.h"
+#include "src/tint/writer/spirv/generator.h"
+
+namespace tint::writer::spirv {
+
+/// The result of sanitizing a program for generation.
+struct SanitizedResult {
+  /// The sanitized program.
+  Program program;
+};
+
+/// Sanitize a program in preparation for generating SPIR-V.
+/// @program The program to sanitize
+/// @param options The SPIR-V generator options.
+SanitizedResult Sanitize(const Program* program, const Options& options);
+
+/// Implementation class for SPIR-V generator
+class GeneratorImpl {
+ public:
+  /// Constructor
+  /// @param program the program to generate
+  /// @param zero_initialize_workgroup_memory `true` to initialize all the
+  /// variables in the Workgroup storage class with OpConstantNull
+  GeneratorImpl(const Program* program, bool zero_initialize_workgroup_memory);
+
+  /// @returns true on successful generation; false otherwise
+  bool Generate();
+
+  /// @returns the result data
+  const std::vector<uint32_t>& result() const;
+
+  /// @returns the result data
+  std::vector<uint32_t>& result();
+
+  /// @returns the error
+  std::string error() const;
+
+ private:
+  Builder builder_;
+  BinaryWriter writer_;
+};
+
+}  // namespace tint::writer::spirv
+
+#endif  // SRC_TINT_WRITER_SPIRV_GENERATOR_IMPL_H_
diff --git a/src/tint/writer/spirv/test_helper.h b/src/tint/writer/spirv/test_helper.h
index ecf2a21..c8052b2 100644
--- a/src/tint/writer/spirv/test_helper.h
+++ b/src/tint/writer/spirv/test_helper.h
@@ -22,6 +22,7 @@
 #include "gtest/gtest.h"
 #include "spirv-tools/libspirv.hpp"
 #include "src/tint/writer/spirv/binary_writer.h"
+#include "src/tint/writer/spirv/generator_impl.h"
 
 namespace tint::writer::spirv {
 
diff --git a/test/tint/BUILD.gn b/test/tint/BUILD.gn
index d2ca288..bb217df 100644
--- a/test/tint/BUILD.gn
+++ b/test/tint/BUILD.gn
@@ -1,4 +1,4 @@
-# Copyright 2021 The Tint Authors
+# Copyright 2022 The Dawn & Tint Authors
 #
 # Licensed under the Apache License, Version 2.0 (the "License");
 # you may not use this file except in compliance with the License.
@@ -175,6 +175,7 @@
     "../../src/tint/ast/id_attribute_test.cc",
     "../../src/tint/ast/identifier_expression_test.cc",
     "../../src/tint/ast/if_statement_test.cc",
+    "../../src/tint/ast/increment_decrement_statement_test.cc",
     "../../src/tint/ast/index_accessor_expression_test.cc",
     "../../src/tint/ast/int_literal_expression_test.cc",
     "../../src/tint/ast/interpolate_attribute_test.cc",
@@ -252,6 +253,7 @@
     "../../src/tint/resolver/entry_point_validation_test.cc",
     "../../src/tint/resolver/function_validation_test.cc",
     "../../src/tint/resolver/host_shareable_validation_test.cc",
+    "../../src/tint/resolver/increment_decrement_validation_test.cc",
     "../../src/tint/resolver/is_host_shareable_test.cc",
     "../../src/tint/resolver/is_storeable_test.cc",
     "../../src/tint/resolver/pipeline_overridable_constant_test.cc",
@@ -320,11 +322,11 @@
     "../../src/tint/transform/decompose_memory_access_test.cc",
     "../../src/tint/transform/decompose_strided_array_test.cc",
     "../../src/tint/transform/decompose_strided_matrix_test.cc",
+    "../../src/tint/transform/expand_compound_assignment_test.cc",
     "../../src/tint/transform/first_index_offset_test.cc",
     "../../src/tint/transform/fold_constants_test.cc",
     "../../src/tint/transform/fold_trivial_single_use_lets_test.cc",
     "../../src/tint/transform/for_loop_to_loop_test.cc",
-    "../../src/tint/transform/expand_compound_assignment_test.cc",
     "../../src/tint/transform/localize_struct_array_assignment_test.cc",
     "../../src/tint/transform/loop_to_for_loop_test.cc",
     "../../src/tint/transform/module_scope_var_to_entry_point_param_test.cc",
@@ -497,6 +499,7 @@
     "../../src/tint/reader/wgsl/parser_impl_global_variable_decl_test.cc",
     "../../src/tint/reader/wgsl/parser_impl_if_stmt_test.cc",
     "../../src/tint/reader/wgsl/parser_impl_inclusive_or_expression_test.cc",
+    "../../src/tint/reader/wgsl/parser_impl_increment_decrement_stmt_test.cc",
     "../../src/tint/reader/wgsl/parser_impl_logical_and_expression_test.cc",
     "../../src/tint/reader/wgsl/parser_impl_logical_or_expression_test.cc",
     "../../src/tint/reader/wgsl/parser_impl_loop_stmt_test.cc",
@@ -664,7 +667,6 @@
 
 tint_unittests_source_set("tint_unittests_glsl_writer_src") {
   sources = [
-    "../../src/tint/transform/glsl_test.cc",
     "../../src/tint/writer/glsl/generator_impl_array_accessor_test.cc",
     "../../src/tint/writer/glsl/generator_impl_assign_test.cc",
     "../../src/tint/writer/glsl/generator_impl_binary_test.cc",