Implement GLSL writer backend.

This is a modified version of the HLSL writer.
Basic types, arrays, entry points, reserved keywords, uniforms,
builtin uniforms, structs, some builtin functions, zero initialization
are implemented. Textures, SSBOs and storage textures in particular are
unimplemented. All the unit tests "pass", but the output is not correct
in many cases.

triangle.wgsl outputs correct vertex and fragment shaders that pass
GLSL validation via glslang. compute_boids.wgsl outputs a valid but not
correct compute shader.

Change-Id: I96c7aaf60cf2d4237e45d732e5f51b345aea0552
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57780
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Stephen White <senorblanco@chromium.org>
Reviewed-by: Ben Clayton <bclayton@google.com>

CMakeLists.txt

@@ -54,6 +54,7 @@ option(TINT_BUILD_DOCS "Build documentation" ${TINT_BUILD_DOCS_DEFAULT})
 option(TINT_DOCS_WARN_AS_ERROR "When building documentation, treat warnings as errors" OFF)
 option(TINT_BUILD_SPV_READER "Build the SPIR-V input reader" ON)
 option(TINT_BUILD_WGSL_READER "Build the WGSL input reader" ON)
+option(TINT_BUILD_GLSL_WRITER "Build the GLSL output writer" ON)
 option(TINT_BUILD_HLSL_WRITER "Build the HLSL output writer" ON)
 option(TINT_BUILD_MSL_WRITER "Build the MSL output writer" ON)
 option(TINT_BUILD_SPV_WRITER "Build the SPIR-V output writer" ON)
@@ -81,6 +82,7 @@ message(STATUS "Tint build docs: ${TINT_BUILD_DOCS}")
 message(STATUS "Tint build docs with warn as error: ${TINT_DOCS_WARN_AS_ERROR}")
 message(STATUS "Tint build SPIR-V reader: ${TINT_BUILD_SPV_READER}")
 message(STATUS "Tint build WGSL reader: ${TINT_BUILD_WGSL_READER}")
+message(STATUS "Tint build GLSL writer: ${TINT_BUILD_GLSL_WRITER}")
 message(STATUS "Tint build HLSL writer: ${TINT_BUILD_HLSL_WRITER}")
 message(STATUS "Tint build MSL writer: ${TINT_BUILD_MSL_WRITER}")
 message(STATUS "Tint build SPIR-V writer: ${TINT_BUILD_SPV_WRITER}")
@@ -110,6 +112,7 @@ if (${TINT_BUILD_SPIRV_TOOLS_FUZZER})
       TINT_BUILD_SPV_WRITER
       TINT_BUILD_WGSL_READER
       TINT_BUILD_WGSL_WRITER
+      TINT_BUILD_GLSL_WRITER
       TINT_BUILD_HLSL_WRITER
       TINT_BUILD_MSL_WRITER to ON")
   set(TINT_BUILD_FUZZERS ON CACHE BOOL "Build tint fuzzers" FORCE)
@@ -117,6 +120,7 @@ if (${TINT_BUILD_SPIRV_TOOLS_FUZZER})
   set(TINT_BUILD_SPV_WRITER ON CACHE BOOL "Build SPIR-V writer" FORCE)
   set(TINT_BUILD_WGSL_READER ON CACHE BOOL "Build WGSL reader" FORCE)
   set(TINT_BUILD_WGSL_WRITER ON CACHE BOOL "Build WGSL writer" FORCE)
+  set(TINT_BUILD_GLSL_WRITER ON CACHE BOOL "Build HLSL writer" FORCE)
   set(TINT_BUILD_HLSL_WRITER ON CACHE BOOL "Build HLSL writer" FORCE)
   set(TINT_BUILD_MSL_WRITER ON CACHE BOOL "Build MSL writer" FORCE)
 endif()
@@ -128,12 +132,14 @@ if (${TINT_BUILD_AST_FUZZER})
       TINT_BUILD_WGSL_WRITER
       TINT_BUILD_SPV_WRITER
       TINT_BUILD_MSL_WRITER
+      TINT_BUILD_GLSL_WRITER
       TINT_BUILD_HLSL_WRITER to ON")
   set(TINT_BUILD_FUZZERS ON CACHE BOOL "Build tint fuzzers" FORCE)
   set(TINT_BUILD_WGSL_READER ON CACHE BOOL "Build WGSL reader" FORCE)
   set(TINT_BUILD_WGSL_WRITER ON CACHE BOOL "Build WGSL writer" FORCE)
   set(TINT_BUILD_SPV_WRITER ON CACHE BOOL "Build SPIR-V writer" FORCE)
   set(TINT_BUILD_MSL_WRITER ON CACHE BOOL "Build MSL writer" FORCE)
+  set(TINT_BUILD_GLSL_WRITER ON CACHE BOOL "Build GLSL writer" FORCE)
   set(TINT_BUILD_HLSL_WRITER ON CACHE BOOL "Build HLSL writer" FORCE)
 endif()
 
@@ -144,12 +150,14 @@ if (${TINT_BUILD_REGEX_FUZZER})
       TINT_BUILD_WGSL_WRITER
       TINT_BUILD_SPV_WRITER
       TINT_BUILD_MSL_WRITER
+      TINT_BUILD_GLSL_WRITER
       TINT_BUILD_HLSL_WRITER to ON")
       set(TINT_BUILD_FUZZERS ON CACHE BOOL "Build tint fuzzers" FORCE)
       set(TINT_BUILD_WGSL_READER ON CACHE BOOL "Build WGSL reader" FORCE)
       set(TINT_BUILD_WGSL_WRITER ON CACHE BOOL "Build WGSL writer" FORCE)
       set(TINT_BUILD_SPV_WRITER ON CACHE BOOL "Build SPIR-V writer" FORCE)
       set(TINT_BUILD_MSL_WRITER ON CACHE BOOL "Build MSL writer" FORCE)
+      set(TINT_BUILD_GLSL_WRITER ON CACHE BOOL "Build GLSL writer" FORCE)
       set(TINT_BUILD_HLSL_WRITER ON CACHE BOOL "Build HLSL writer" FORCE)
 endif()
 
@@ -239,6 +247,8 @@ function(tint_default_compile_options TARGET)
       -DTINT_BUILD_SPV_READER=$<BOOL:${TINT_BUILD_SPV_READER}>)
   target_compile_definitions(${TARGET} PUBLIC
       -DTINT_BUILD_WGSL_READER=$<BOOL:${TINT_BUILD_WGSL_READER}>)
+  target_compile_definitions(${TARGET} PUBLIC
+      -DTINT_BUILD_GLSL_WRITER=$<BOOL:${TINT_BUILD_GLSL_WRITER}>)
   target_compile_definitions(${TARGET} PUBLIC
       -DTINT_BUILD_HLSL_WRITER=$<BOOL:${TINT_BUILD_HLSL_WRITER}>)
   target_compile_definitions(${TARGET} PUBLIC

include/tint/tint.h

@@ -59,4 +59,9 @@
 #include "src/writer/hlsl/generator.h"
 #endif  // TINT_BUILD_HLSL_WRITER
 
+#if TINT_BUILD_GLSL_WRITER
+#include "src/transform/glsl.h"
+#include "src/writer/glsl/generator.h"
+#endif  // TINT_BUILD_GLSL_WRITER
+
 #endif  // INCLUDE_TINT_TINT_H_

samples/main.cc

@@ -54,6 +54,7 @@ enum class Format {
   kWgsl,
   kMsl,
   kHlsl,
+  kGlsl,
 };
 
 struct Options {
@@ -142,6 +143,11 @@ Format parse_format(const std::string& fmt) {
     return Format::kHlsl;
 #endif  // TINT_BUILD_HLSL_WRITER
 
+#if TINT_BUILD_GLSL_WRITER
+  if (fmt == "glsl")
+    return Format::kGlsl;
+#endif  // TINT_BUILD_GLSL_WRITER
+
   return Format::kNone;
 }
 
@@ -844,6 +850,33 @@ bool GenerateHlsl(const tint::Program* program, const Options& options) {
 #endif  // TINT_BUILD_HLSL_WRITER
 }
 
+/// Generate GLSL code for a program.
+/// @param program the program to generate
+/// @param options the options that Tint was invoked with
+/// @returns true on success
+bool GenerateGlsl(const tint::Program* program, const Options& options) {
+#if TINT_BUILD_GLSL_WRITER
+  tint::writer::glsl::Options gen_options;
+  auto result = tint::writer::glsl::Generate(program, gen_options);
+  if (!result.success) {
+    PrintWGSL(std::cerr, *program);
+    std::cerr << "Failed to generate: " << result.error << std::endl;
+    return false;
+  }
+
+  if (!WriteFile(options.output_file, "w", result.glsl)) {
+    return false;
+  }
+
+  // TODO(senorblanco): implement GLSL validation
+
+  return true;
+#else
+  std::cerr << "GLSL writer not enabled in tint build" << std::endl;
+  return false;
+#endif  // TINT_BUILD_GLSL_WRITER
+}
+
 }  // namespace
 
 int main(int argc, const char** argv) {
@@ -996,6 +1029,14 @@ int main(int argc, const char** argv) {
 #endif  // TINT_BUILD_MSL_WRITER
       break;
     }
+#if TINT_BUILD_GLSL_WRITER
+    case Format::kGlsl: {
+      transform_inputs.Add<tint::transform::Renamer::Config>(
+          tint::transform::Renamer::Target::kGlslKeywords);
+      transform_manager.Add<tint::transform::Renamer>();
+      break;
+    }
+#endif  // TINT_BUILD_GLSL_WRITER
     case Format::kHlsl: {
 #if TINT_BUILD_HLSL_WRITER
       transform_inputs.Add<tint::transform::Renamer::Config>(
@@ -1066,6 +1107,9 @@ int main(int argc, const char** argv) {
     case Format::kHlsl:
       success = GenerateHlsl(program.get(), options);
       break;
+    case Format::kGlsl:
+      success = GenerateGlsl(program.get(), options);
+      break;
     default:
       std::cerr << "Unknown output format specified" << std::endl;
       return 1;

src/BUILD.gn

@@ -67,6 +67,12 @@ config("tint_public_config") {
     defines += [ "TINT_BUILD_HLSL_WRITER=0" ]
   }
 
+  if (tint_build_glsl_writer) {
+    defines += [ "TINT_BUILD_GLSL_WRITER=1" ]
+  } else {
+    defines += [ "TINT_BUILD_GLSL_WRITER=0" ]
+  }
+
   include_dirs = [
     "${tint_root_dir}/",
     "${tint_root_dir}/include/",
@@ -688,6 +694,19 @@ libtint_source_set("libtint_hlsl_writer_src") {
   public_deps = [ ":libtint_core_src" ]
 }
 
+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",
+    "writer/glsl/generator_impl.h",
+  ]
+
+  public_deps = [ ":libtint_core_src" ]
+}
+
 source_set("libtint") {
   public_deps = [ ":libtint_core_src" ]
 
@@ -715,6 +734,10 @@ source_set("libtint") {
     public_deps += [ ":libtint_hlsl_writer_src" ]
   }
 
+  if (tint_build_glsl_writer) {
+    public_deps += [ ":libtint_glsl_writer_src" ]
+  }
+
   configs += [ ":tint_common_config" ]
   public_configs = [ ":tint_public_config" ]
 

src/CMakeLists.txt

@@ -310,6 +310,8 @@ set(TINT_LIB_SRCS
   transform/fold_trivial_single_use_lets.h
   transform/for_loop_to_loop.cc
   transform/for_loop_to_loop.h
+  transform/glsl.cc
+  transform/glsl.h
   transform/inline_pointer_lets.cc
   transform/inline_pointer_lets.h
   transform/loop_to_for_loop.cc
@@ -484,6 +486,15 @@ if(${TINT_BUILD_MSL_WRITER})
   )
 endif()
 
+if(${TINT_BUILD_GLSL_WRITER})
+  list(APPEND TINT_LIB_SRCS
+    writer/glsl/generator.cc
+    writer/glsl/generator.h
+    writer/glsl/generator_impl.cc
+    writer/glsl/generator_impl.h
+  )
+endif()
+
 if(${TINT_BUILD_HLSL_WRITER})
   list(APPEND TINT_LIB_SRCS
     writer/hlsl/generator.cc
@@ -981,6 +992,41 @@ if(${TINT_BUILD_TESTS})
     )
   endif()
 
+  if (${TINT_BUILD_GLSL_WRITER})
+    list(APPEND TINT_TEST_SRCS
+      writer/glsl/generator_impl_array_accessor_test.cc
+      writer/glsl/generator_impl_assign_test.cc
+      writer/glsl/generator_impl_binary_test.cc
+      writer/glsl/generator_impl_bitcast_test.cc
+      writer/glsl/generator_impl_block_test.cc
+      writer/glsl/generator_impl_break_test.cc
+      writer/glsl/generator_impl_call_test.cc
+      writer/glsl/generator_impl_case_test.cc
+      writer/glsl/generator_impl_cast_test.cc
+      writer/glsl/generator_impl_constructor_test.cc
+      writer/glsl/generator_impl_continue_test.cc
+      writer/glsl/generator_impl_discard_test.cc
+      writer/glsl/generator_impl_function_test.cc
+      writer/glsl/generator_impl_identifier_test.cc
+      writer/glsl/generator_impl_if_test.cc
+      writer/glsl/generator_impl_intrinsic_test.cc
+      writer/glsl/generator_impl_intrinsic_texture_test.cc
+      writer/glsl/generator_impl_import_test.cc
+      writer/glsl/generator_impl_loop_test.cc
+      writer/glsl/generator_impl_member_accessor_test.cc
+      writer/glsl/generator_impl_module_constant_test.cc
+      writer/glsl/generator_impl_return_test.cc
+      writer/glsl/generator_impl_sanitizer_test.cc
+      writer/glsl/generator_impl_switch_test.cc
+      writer/glsl/generator_impl_test.cc
+      writer/glsl/generator_impl_type_test.cc
+      writer/glsl/generator_impl_unary_op_test.cc
+      writer/glsl/generator_impl_variable_decl_statement_test.cc
+      writer/glsl/generator_impl_workgroup_var_test.cc
+      writer/glsl/test_helper.h
+    )
+  endif()
+
   if (${TINT_BUILD_HLSL_WRITER})
     list(APPEND TINT_TEST_SRCS
       writer/hlsl/generator_impl_array_accessor_test.cc

src/transform/glsl.cc

@@ -0,0 +1,102 @@
+// 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/transform/glsl.h"
+
+#include <utility>
+
+#include "src/program_builder.h"
+#include "src/transform/calculate_array_length.h"
+#include "src/transform/canonicalize_entry_point_io.h"
+#include "src/transform/decompose_memory_access.h"
+#include "src/transform/external_texture_transform.h"
+#include "src/transform/fold_trivial_single_use_lets.h"
+#include "src/transform/inline_pointer_lets.h"
+#include "src/transform/loop_to_for_loop.h"
+#include "src/transform/manager.h"
+#include "src/transform/pad_array_elements.h"
+#include "src/transform/promote_initializers_to_const_var.h"
+#include "src/transform/simplify.h"
+#include "src/transform/zero_init_workgroup_memory.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::transform::Glsl);
+TINT_INSTANTIATE_TYPEINFO(tint::transform::Glsl::Config);
+
+namespace tint {
+namespace transform {
+
+Glsl::Glsl() = default;
+Glsl::~Glsl() = default;
+
+Output Glsl::Run(const Program* in, const DataMap& inputs) {
+  Manager manager;
+  DataMap data;
+
+  auto* cfg = inputs.Get<Config>();
+
+  // 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<InlinePointerLets>();
+  // Simplify cleans up messy `*(&(expr))` expressions from InlinePointerLets.
+  manager.Add<Simplify>();
+  manager.Add<CalculateArrayLength>();
+  manager.Add<ExternalTextureTransform>();
+  manager.Add<PromoteInitializersToConstVar>();
+  manager.Add<PadArrayElements>();
+
+  // For now, canonicalize to structs for all IO, as in HLSL.
+  // TODO(senorblanco): we could skip this by accessing global entry point
+  // variables directly.
+  data.Add<CanonicalizeEntryPointIO::Config>(
+      CanonicalizeEntryPointIO::ShaderStyle::kHlsl);
+  auto out = manager.Run(in, data);
+  if (!out.program.IsValid()) {
+    return out;
+  }
+
+  ProgramBuilder builder;
+  CloneContext ctx(&builder, &out.program);
+  AddEmptyEntryPoint(ctx);
+  ctx.Clone();
+  builder.SetTransformApplied(this);
+  return Output{Program(std::move(builder))};
+}
+
+void Glsl::AddEmptyEntryPoint(CloneContext& ctx) const {
+  for (auto* func : ctx.src->AST().Functions()) {
+    if (func->IsEntryPoint()) {
+      return;
+    }
+  }
+  ctx.dst->Func(ctx.dst->Symbols().New("unused_entry_point"), {},
+                ctx.dst->ty.void_(), {},
+                {ctx.dst->Stage(ast::PipelineStage::kCompute),
+                 ctx.dst->WorkgroupSize(1)});
+}
+
+Glsl::Config::Config(bool disable_wi) : disable_workgroup_init(disable_wi) {}
+Glsl::Config::Config(const Config&) = default;
+Glsl::Config::~Config() = default;
+
+}  // namespace transform
+}  // namespace tint

src/transform/glsl.h

@@ -0,0 +1,67 @@
+// 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_TRANSFORM_GLSL_H_
+#define SRC_TRANSFORM_GLSL_H_
+
+#include "src/transform/transform.h"
+
+namespace tint {
+
+// Forward declarations
+class CloneContext;
+
+namespace 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 : public Castable<Glsl, Transform> {
+ public:
+  /// Configuration options for the Glsl sanitizer transform.
+  struct Config : public Castable<Data, transform::Data> {
+    /// Constructor
+    /// @param disable_workgroup_init `true` to disable workgroup memory zero
+    ///        initialization
+    explicit Config(bool disable_workgroup_init = false);
+
+    /// Copy constructor
+    Config(const Config&);
+
+    /// Destructor
+    ~Config() override;
+
+    /// Set to `true` to disable workgroup memory zero initialization
+    bool disable_workgroup_init = 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 = {}) override;
+
+ private:
+  /// Add an empty shader entry point if none exist in the module.
+  void AddEmptyEntryPoint(CloneContext& ctx) const;
+};
+
+}  // namespace transform
+}  // namespace tint
+
+#endif  // SRC_TRANSFORM_GLSL_H_

src/transform/glsl_test.cc

@@ -0,0 +1,41 @@
+// 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/transform/glsl.h"
+
+#include "src/transform/test_helper.h"
+
+namespace tint {
+namespace 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 transform
+}  // namespace tint

src/transform/renamer.cc

@@ -31,6 +31,211 @@ namespace transform {
 
 namespace {
 
+// This list is used for a binary search and must be kept in sorted order.
+const char* kReservedKeywordsGLSL[] = {
+    "active",
+    "asm",
+    "atomic_uint",
+    "attribute",
+    "bool",
+    "break",
+    "buffer",
+    "bvec2",
+    "bvec3",
+    "bvec4",
+    "case",
+    "cast",
+    "centroid",
+    "class",
+    "coherent",
+    "common",
+    "const",
+    "continue",
+    "default",
+    "discard",
+    "dmat2",
+    "dmat2x2",
+    "dmat2x3",
+    "dmat2x4",
+    "dmat3",
+    "dmat3x2",
+    "dmat3x3",
+    "dmat3x4",
+    "dmat4",
+    "dmat4x2",
+    "dmat4x3",
+    "dmat4x4",
+    "do",
+    "double",
+    "dvec2",
+    "dvec3",
+    "dvec4",
+    "else",
+    "enum",
+    "extern",
+    "external",
+    "false",
+    "filter",
+    "fixed",
+    "flat",
+    "float",
+    "for",
+    "fvec2",
+    "fvec3",
+    "fvec4",
+    "goto",
+    "half",
+    "highp",
+    "hvec2",
+    "hvec3",
+    "hvec4",
+    "if",
+    "iimage1D",
+    "iimage1DArray",
+    "iimage2D",
+    "iimage2DArray",
+    "iimage2DMS",
+    "iimage2DMSArray",
+    "iimage2DRect",
+    "iimage3D",
+    "iimageBuffer",
+    "iimageCube",
+    "iimageCubeArray",
+    "image1D",
+    "image1DArray",
+    "image2D",
+    "image2DArray",
+    "image2DMS",
+    "image2DMSArray",
+    "image2DRect",
+    "image3D",
+    "imageBuffer",
+    "imageCube",
+    "imageCubeArray",
+    "in",
+    "inline",
+    "inout",
+    "input",
+    "int",
+    "interface",
+    "invariant",
+    "isampler1D",
+    "isampler1DArray",
+    "isampler2D",
+    "isampler2DArray",
+    "isampler2DMS",
+    "isampler2DMSArray",
+    "isampler2DRect",
+    "isampler3D",
+    "isamplerBuffer",
+    "isamplerCube",
+    "isamplerCubeArray",
+    "ivec2",
+    "ivec3",
+    "ivec4",
+    "layout",
+    "long",
+    "lowp",
+    "mat2",
+    "mat2x2",
+    "mat2x3",
+    "mat2x4",
+    "mat3",
+    "mat3x2",
+    "mat3x3",
+    "mat3x4",
+    "mat4",
+    "mat4x2",
+    "mat4x3",
+    "mat4x4",
+    "mediump",
+    "namespace",
+    "noinline",
+    "noperspective",
+    "out",
+    "output",
+    "partition",
+    "patch",
+    "precise",
+    "precision",
+    "public",
+    "readonly",
+    "resource",
+    "restrict",
+    "return",
+    "sample",
+    "sampler1D",
+    "sampler1DArray",
+    "sampler1DArrayShadow",
+    "sampler1DShadow",
+    "sampler2D",
+    "sampler2DArray",
+    "sampler2DArrayShadow",
+    "sampler2DMS",
+    "sampler2DMSArray",
+    "sampler2DRect",
+    "sampler2DRectShadow",
+    "sampler2DShadow",
+    "sampler3D",
+    "sampler3DRect",
+    "samplerBuffer",
+    "samplerCube",
+    "samplerCubeArray",
+    "samplerCubeArrayShadow",
+    "samplerCubeShadow",
+    "shared",
+    "short",
+    "sizeof",
+    "smooth",
+    "static",
+    "struct",
+    "subroutine",
+    "superp",
+    "switch",
+    "template",
+    "this",
+    "true",
+    "typedef",
+    "uimage1D",
+    "uimage1DArray",
+    "uimage2D",
+    "uimage2DArray",
+    "uimage2DMS",
+    "uimage2DMSArray",
+    "uimage2DRect",
+    "uimage3D",
+    "uimageBuffer",
+    "uimageCube",
+    "uimageCubeArray",
+    "uint",
+    "uniform",
+    "union",
+    "unsigned",
+    "usampler1D",
+    "usampler1DArray",
+    "usampler2D",
+    "usampler2DArray",
+    "usampler2DMS",
+    "usampler2DMSArray",
+    "usampler2DRect",
+    "usampler3D",
+    "usamplerBuffer",
+    "usamplerCube",
+    "usamplerCubeArray",
+    "using",
+    "uvec2",
+    "uvec3",
+    "uvec4",
+    "varying",
+    "vec2",
+    "vec3",
+    "vec4",
+    "void",
+    "volatile",
+    "while",
+    "writeonly",
+};
+
 // This list is used for a binary search and must be kept in sorted order.
 const char* kReservedKeywordsHLSL[] = {
     "AddressU",
@@ -944,6 +1149,16 @@ Output Renamer::Run(const Program* in, const DataMap& inputs) {
       case Target::kAll:
         // Always rename.
         break;
+      case Target::kGlslKeywords:
+        if (!std::binary_search(
+                kReservedKeywordsGLSL,
+                kReservedKeywordsGLSL +
+                    sizeof(kReservedKeywordsGLSL) / sizeof(const char*),
+                name_in)) {
+          // No match, just reuse the original name.
+          return ctx.dst->Symbols().New(name_in);
+        }
+        break;
       case Target::kHlslKeywords:
         if (!std::binary_search(
                 kReservedKeywordsHLSL,

src/transform/renamer.h

@@ -50,6 +50,8 @@ class Renamer : public Castable<Renamer, Transform> {
   enum class Target {
     /// Rename every symbol.
     kAll,
+    /// Only rename symbols that are reserved keywords in GLSL.
+    kGlslKeywords,
     /// Only rename symbols that are reserved keywords in HLSL.
     kHlslKeywords,
     /// Only rename symbols that are reserved keywords in MSL.

src/writer/glsl/generator.cc

@@ -0,0 +1,59 @@
+// 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/writer/glsl/generator.h"
+
+#include "src/transform/glsl.h"
+#include "src/writer/glsl/generator_impl.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+
+Result::Result() = default;
+Result::~Result() = default;
+Result::Result(const Result&) = default;
+
+Result Generate(const Program* program, const Options&) {
+  Result result;
+
+  // Run the GLSL sanitizer.
+  transform::Glsl sanitizer;
+  auto output = sanitizer.Run(program);
+  if (!output.program.IsValid()) {
+    result.success = false;
+    result.error = output.program.Diagnostics().str();
+    return result;
+  }
+
+  // Generate the GLSL code.
+  auto impl = std::make_unique<GeneratorImpl>(&output.program);
+  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()) {
+    if (func->IsEntryPoint()) {
+      auto name = output.program.Symbols().NameFor(func->symbol());
+      result.entry_points.push_back({name, func->pipeline_stage()});
+    }
+  }
+
+  return result;
+}
+
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator.h

@@ -0,0 +1,76 @@
+// 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_WRITER_GLSL_GENERATOR_H_
+#define SRC_WRITER_GLSL_GENERATOR_H_
+
+#include <memory>
+#include <string>
+#include <utility>
+#include <vector>
+
+#include "src/ast/pipeline_stage.h"
+#include "src/writer/text.h"
+
+namespace tint {
+
+// Forward declarations
+class Program;
+
+namespace writer {
+namespace glsl {
+
+// Forward declarations
+class GeneratorImpl;
+
+/// Configuration options used for generating GLSL.
+struct Options {};
+
+/// The result produced when generating GLSL.
+struct Result {
+  /// Constructor
+  Result();
+
+  /// Destructor
+  ~Result();
+
+  /// Copy constructor
+  Result(const Result&);
+
+  /// True if generation was successful.
+  bool success = false;
+
+  /// The errors generated during code generation, if any.
+  std::string error;
+
+  /// The generated GLSL.
+  std::string glsl = "";
+
+  /// The list of entry points in the generated GLSL.
+  std::vector<std::pair<std::string, ast::PipelineStage>> entry_points;
+};
+
+/// Generate GLSL for a program, according to a set of configuration options.
+/// The result will contain the GLSL, as well as success status and diagnostic
+/// information.
+/// @param program the program to translate to GLSL
+/// @param options the configuration options to use when generating GLSL
+/// @returns the resulting GLSL and supplementary information
+Result Generate(const Program* program, const Options& options);
+
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint
+
+#endif  // SRC_WRITER_GLSL_GENERATOR_H_

src/writer/glsl/generator_impl.h

@@ -0,0 +1,418 @@
+// 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_WRITER_GLSL_GENERATOR_IMPL_H_
+#define SRC_WRITER_GLSL_GENERATOR_IMPL_H_
+
+#include <string>
+#include <unordered_map>
+#include <unordered_set>
+#include <utility>
+
+#include "src/ast/assignment_statement.h"
+#include "src/ast/bitcast_expression.h"
+#include "src/ast/break_statement.h"
+#include "src/ast/continue_statement.h"
+#include "src/ast/discard_statement.h"
+#include "src/ast/for_loop_statement.h"
+#include "src/ast/if_statement.h"
+#include "src/ast/loop_statement.h"
+#include "src/ast/return_statement.h"
+#include "src/ast/switch_statement.h"
+#include "src/ast/unary_op_expression.h"
+#include "src/program_builder.h"
+#include "src/scope_stack.h"
+#include "src/transform/decompose_memory_access.h"
+#include "src/utils/hash.h"
+#include "src/writer/text_generator.h"
+
+namespace tint {
+
+// Forward declarations
+namespace sem {
+class Call;
+class Intrinsic;
+}  // namespace sem
+
+namespace writer {
+namespace glsl {
+
+/// Implementation class for GLSL generator
+class GeneratorImpl : public TextGenerator {
+ public:
+  /// Constructor
+  /// @param program the program to generate
+  explicit GeneratorImpl(const Program* program);
+  ~GeneratorImpl();
+
+  /// @returns true on successful generation; false otherwise
+  bool Generate();
+
+  /// 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(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(std::ostream& out, ast::BinaryExpression* expr);
+  /// Handles generating a bitcast expression
+  /// @param out the output of the expression stream
+  /// @param expr the as expression
+  /// @returns true if the bitcast was emitted
+  bool EmitBitcast(std::ostream& out, ast::BitcastExpression* expr);
+  /// Emits a list of statements
+  /// @param stmts the statement list
+  /// @returns true if the statements were emitted successfully
+  bool EmitStatements(const ast::StatementList& stmts);
+  /// Emits a list of statements with an indentation
+  /// @param stmts the statement list
+  /// @returns true if the statements were emitted successfully
+  bool EmitStatementsWithIndent(const ast::StatementList& stmts);
+  /// Handles a block statement
+  /// @param stmt the statement to emit
+  /// @returns true if the statement was emitted successfully
+  bool EmitBlock(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(std::ostream& out, ast::CallExpression* expr);
+  /// Handles generating a barrier intrinsic call
+  /// @param out the output of the expression stream
+  /// @param intrinsic the semantic information for the barrier intrinsic
+  /// @returns true if the call expression is emitted
+  bool EmitBarrierCall(std::ostream& out, const sem::Intrinsic* intrinsic);
+  /// Handles generating an atomic intrinsic call for a storage buffer variable
+  /// @param out the output of the expression stream
+  /// @param expr the call expression
+  /// @param intrinsic the atomic intrinsic
+  /// @returns true if the call expression is emitted
+  bool EmitStorageAtomicCall(
+      std::ostream& out,
+      ast::CallExpression* expr,
+      const transform::DecomposeMemoryAccess::Intrinsic* intrinsic);
+  /// Handles generating an atomic intrinsic call for a workgroup variable
+  /// @param out the output of the expression stream
+  /// @param expr the call expression
+  /// @param intrinsic the semantic information for the atomic intrinsic
+  /// @returns true if the call expression is emitted
+  bool EmitWorkgroupAtomicCall(std::ostream& out,
+                               ast::CallExpression* expr,
+                               const sem::Intrinsic* intrinsic);
+  /// Handles generating a call to a texture function (`textureSample`,
+  /// `textureSampleGrad`, etc)
+  /// @param out the output of the expression stream
+  /// @param expr the call expression
+  /// @param intrinsic the semantic information for the texture intrinsic
+  /// @returns true if the call expression is emitted
+  bool EmitTextureCall(std::ostream& out,
+                       ast::CallExpression* expr,
+                       const sem::Intrinsic* intrinsic);
+  /// Handles generating a call to the `select()` intrinsic
+  /// @param out the output of the expression stream
+  /// @param expr the call expression
+  /// @returns true if the call expression is emitted
+  bool EmitSelectCall(std::ostream& out, ast::CallExpression* expr);
+  /// Handles generating a call to the `modf()` intrinsic
+  /// @param out the output of the expression stream
+  /// @param expr the call expression
+  /// @param intrinsic the semantic information for the intrinsic
+  /// @returns true if the call expression is emitted
+  bool EmitModfCall(std::ostream& out,
+                    ast::CallExpression* expr,
+                    const sem::Intrinsic* intrinsic);
+  /// Handles generating a call to the `frexp()` intrinsic
+  /// @param out the output of the expression stream
+  /// @param expr the call expression
+  /// @param intrinsic the semantic information for the intrinsic
+  /// @returns true if the call expression is emitted
+  bool EmitFrexpCall(std::ostream& out,
+                     ast::CallExpression* expr,
+                     const sem::Intrinsic* intrinsic);
+  /// Handles generating a call to the `isNormal()` intrinsic
+  /// @param out the output of the expression stream
+  /// @param expr the call expression
+  /// @param intrinsic the semantic information for the intrinsic
+  /// @returns true if the call expression is emitted
+  bool EmitIsNormalCall(std::ostream& out,
+                        ast::CallExpression* expr,
+                        const sem::Intrinsic* intrinsic);
+  /// Handles generating a call to data packing intrinsic
+  /// @param out the output of the expression stream
+  /// @param expr the call expression
+  /// @param intrinsic the semantic information for the texture intrinsic
+  /// @returns true if the call expression is emitted
+  bool EmitDataPackingCall(std::ostream& out,
+                           ast::CallExpression* expr,
+                           const sem::Intrinsic* intrinsic);
+  /// Handles generating a call to data unpacking intrinsic
+  /// @param out the output of the expression stream
+  /// @param expr the call expression
+  /// @param intrinsic the semantic information for the texture intrinsic
+  /// @returns true if the call expression is emitted
+  bool EmitDataUnpackingCall(std::ostream& out,
+                             ast::CallExpression* expr,
+                             const sem::Intrinsic* intrinsic);
+  /// Handles a case statement
+  /// @param stmt the statement
+  /// @returns true if the statement was emitted successfully
+  bool EmitCase(ast::CaseStatement* 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(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 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(std::ostream& out,
+                             ast::ScalarConstructorExpression* expr);
+  /// 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(std::ostream& out,
+                           ast::TypeConstructorExpression* 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 generate an Expression
+  /// @param out the output of the expression stream
+  /// @param expr the expression
+  /// @returns true if the expression was emitted
+  bool EmitExpression(std::ostream& out, ast::Expression* expr);
+  /// Handles generating a function
+  /// @param func the function to generate
+  /// @returns true if the function was emitted
+  bool EmitFunction(ast::Function* func);
+
+  /// Handles emitting a global variable
+  /// @param global the global variable
+  /// @returns true on success
+  bool EmitGlobalVariable(ast::Variable* global);
+
+  /// Handles emitting a global variable with the uniform storage class
+  /// @param var the global variable
+  /// @returns true on success
+  bool EmitUniformVariable(const sem::Variable* var);
+
+  /// Handles emitting a global variable with the storage storage class
+  /// @param var the global variable
+  /// @returns true on success
+  bool EmitStorageVariable(const sem::Variable* var);
+
+  /// Handles emitting a global variable with the handle storage class
+  /// @param var the global variable
+  /// @returns true on success
+  bool EmitHandleVariable(const sem::Variable* var);
+
+  /// Handles emitting a global variable with the private storage class
+  /// @param var the global variable
+  /// @returns true on success
+  bool EmitPrivateVariable(const sem::Variable* var);
+
+  /// Handles emitting a global variable with the workgroup storage class
+  /// @param var the global variable
+  /// @returns true on success
+  bool EmitWorkgroupVariable(const sem::Variable* var);
+
+  /// Handles emitting the entry point function
+  /// @param func the entry point
+  /// @returns true if the entry point function was emitted
+  bool EmitEntryPointFunction(ast::Function* func);
+  /// Handles an if statement
+  /// @param stmt the statement to emit
+  /// @returns true if the statement was successfully emitted
+  bool EmitIf(ast::IfStatement* stmt);
+  /// Handles a literal
+  /// @param out the output stream
+  /// @param lit the literal to emit
+  /// @returns true if the literal was successfully emitted
+  bool EmitLiteral(std::ostream& out, ast::Literal* lit);
+  /// Handles a loop statement
+  /// @param stmt the statement to emit
+  /// @returns true if the statement was emitted
+  bool EmitLoop(ast::LoopStatement* stmt);
+  /// Handles a for loop statement
+  /// @param stmt the statement to emit
+  /// @returns true if the statement was emitted
+  bool EmitForLoop(ast::ForLoopStatement* stmt);
+  /// 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(std::ostream& out, ast::IdentifierExpression* expr);
+  /// 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(std::ostream& out,
+                          ast::MemberAccessorExpression* expr);
+  /// Handles return statements
+  /// @param stmt the statement to emit
+  /// @returns true if the statement was successfully emitted
+  bool EmitReturn(ast::ReturnStatement* stmt);
+  /// Handles statement
+  /// @param stmt the statement to emit
+  /// @returns true if the statement was emitted
+  bool EmitStatement(ast::Statement* stmt);
+  /// 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 stream
+  /// @param type the type to generate
+  /// @param storage_class the storage class of the variable
+  /// @param access the access control type of the variable
+  /// @param name the name of the variable, used for array emission.
+  /// @param name_printed (optional) if not nullptr and an array was printed
+  /// then the boolean is set to true.
+  /// @returns true if the type is emitted
+  bool EmitType(std::ostream& out,
+                const sem::Type* type,
+                ast::StorageClass storage_class,
+                ast::Access access,
+                const std::string& name,
+                bool* name_printed = nullptr);
+  /// Handles generating type and name
+  /// @param out the output stream
+  /// @param type the type to generate
+  /// @param storage_class the storage class of the variable
+  /// @param access the access control type of the variable
+  /// @param name the name to emit
+  /// @returns true if the type is emitted
+  bool EmitTypeAndName(std::ostream& out,
+                       const sem::Type* type,
+                       ast::StorageClass storage_class,
+                       ast::Access access,
+                       const std::string& name);
+  /// Handles generating a structure declaration
+  /// @param buffer the text buffer that the type declaration will be written to
+  /// @param ty the struct to generate
+  /// @returns true if the struct is emitted
+  bool EmitStructType(TextBuffer* buffer, const sem::Struct* ty);
+  /// 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(std::ostream& out, ast::UnaryOpExpression* expr);
+  /// Emits the zero value for the given type
+  /// @param out the output stream
+  /// @param type the type to emit the value for
+  /// @returns true if the zero value was successfully emitted.
+  bool EmitZeroValue(std::ostream& out, const sem::Type* type);
+  /// Handles generating a variable
+  /// @param var the variable to generate
+  /// @returns true if the variable was emitted
+  bool EmitVariable(ast::Variable* var);
+  /// Handles generating a program scope constant variable
+  /// @param var the variable to emit
+  /// @returns true if the variable was emitted
+  bool EmitProgramConstVariable(const ast::Variable* var);
+  /// Handles generating a builtin method name
+  /// @param intrinsic the semantic info for the intrinsic
+  /// @returns the name or "" if not valid
+  std::string generate_builtin_name(const sem::Intrinsic* intrinsic);
+  /// Converts a builtin to a gl_ string
+  /// @param builtin the builtin to convert
+  /// @returns the string name of the builtin or blank on error
+  const char* builtin_to_string(ast::Builtin builtin) const;
+  /// Converts a builtin to a sem::Type appropriate for GLSL.
+  /// @param builtin the builtin to convert
+  /// @returns the appropriate semantic type or null on error.
+  sem::Type* builtin_type(ast::Builtin builtin);
+
+  /// Converts interpolation attributes to a GLSL modifiers
+  /// @param type the interpolation type
+  /// @param sampling the interpolation sampling
+  /// @returns the string name of the attribute or blank on error
+  std::string interpolation_to_modifiers(
+      ast::InterpolationType type,
+      ast::InterpolationSampling sampling) const;
+
+ private:
+  enum class VarType { kIn, kOut };
+
+  struct EntryPointData {
+    std::string struct_name;
+    std::string var_name;
+  };
+
+  struct DMAIntrinsic {
+    transform::DecomposeMemoryAccess::Intrinsic::Op op;
+    transform::DecomposeMemoryAccess::Intrinsic::DataType type;
+    bool operator==(const DMAIntrinsic& rhs) const {
+      return op == rhs.op && type == rhs.type;
+    }
+    /// Hasher is a std::hash function for DMAIntrinsic
+    struct Hasher {
+      /// @param i the DMAIntrinsic to hash
+      /// @returns the hash of `i`
+      inline std::size_t operator()(const DMAIntrinsic& i) const {
+        return utils::Hash(i.op, i.type);
+      }
+    };
+  };
+
+  /// CallIntrinsicHelper will call the intrinsic helper function, creating it
+  /// if it hasn't been built already. If the intrinsic needs to be built then
+  /// CallIntrinsicHelper will generate the function signature and will call
+  /// `build` to emit the body of the function.
+  /// @param out the output of the expression stream
+  /// @param call the call expression
+  /// @param intrinsic the semantic information for the intrinsic
+  /// @param build a function with the signature:
+  ///        `bool(TextBuffer* buffer, const std::vector<std::string>& params)`
+  ///        Where:
+  ///          `buffer` is the body of the generated function
+  ///          `params` is the name of all the generated function parameters
+  /// @returns true if the call expression is emitted
+  template <typename F>
+  bool CallIntrinsicHelper(std::ostream& out,
+                           ast::CallExpression* call,
+                           const sem::Intrinsic* intrinsic,
+                           F&& build);
+
+  TextBuffer helpers_;  // Helper functions emitted at the top of the output
+  std::function<bool()> emit_continuing_;
+  std::unordered_map<DMAIntrinsic, std::string, DMAIntrinsic::Hasher>
+      dma_intrinsics_;
+  std::unordered_map<const sem::Intrinsic*, std::string> intrinsics_;
+  std::unordered_map<const sem::Struct*, std::string> structure_builders_;
+  std::unordered_map<const sem::Vector*, std::string> dynamic_vector_write_;
+};
+
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint
+
+#endif  // SRC_WRITER_GLSL_GENERATOR_IMPL_H_

src/writer/glsl/generator_impl_array_accessor_test.cc

@@ -0,0 +1,39 @@
+// 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/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using GlslGeneratorImplTest_Expression = TestHelper;
+
+TEST_F(GlslGeneratorImplTest_Expression, ArrayAccessor) {
+  Global("ary", ty.array<i32, 10>(), ast::StorageClass::kPrivate);
+  auto* expr = IndexAccessor("ary", 5);
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(), "ary[5]");
+}
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_assign_test.cc

@@ -0,0 +1,41 @@
+// 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/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using GlslGeneratorImplTest_Assign = TestHelper;
+
+TEST_F(GlslGeneratorImplTest_Assign, Emit_Assign) {
+  Global("lhs", ty.i32(), ast::StorageClass::kPrivate);
+  Global("rhs", ty.i32(), ast::StorageClass::kPrivate);
+  auto* assign = Assign("lhs", "rhs");
+  WrapInFunction(assign);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(assign)) << gen.error();
+  EXPECT_EQ(gen.result(), "  lhs = rhs;\n");
+}
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_binary_test.cc

@@ -0,0 +1,557 @@
+// 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/ast/call_statement.h"
+#include "src/ast/variable_decl_statement.h"
+#include "src/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using GlslGeneratorImplTest_Binary = TestHelper;
+
+struct BinaryData {
+  const char* result;
+  ast::BinaryOp op;
+};
+inline std::ostream& operator<<(std::ostream& out, BinaryData data) {
+  out << data.op;
+  return out;
+}
+
+using GlslBinaryTest = TestParamHelper<BinaryData>;
+TEST_P(GlslBinaryTest, Emit_f32) {
+  auto params = GetParam();
+
+  // Skip ops that are illegal for this type
+  if (params.op == ast::BinaryOp::kAnd || params.op == ast::BinaryOp::kOr ||
+      params.op == ast::BinaryOp::kXor ||
+      params.op == ast::BinaryOp::kShiftLeft ||
+      params.op == ast::BinaryOp::kShiftRight) {
+    return;
+  }
+
+  Global("left", ty.f32(), ast::StorageClass::kPrivate);
+  Global("right", ty.f32(), ast::StorageClass::kPrivate);
+
+  auto* left = Expr("left");
+  auto* right = Expr("right");
+
+  auto* expr = create<ast::BinaryExpression>(params.op, left, right);
+
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(), params.result);
+}
+TEST_P(GlslBinaryTest, Emit_u32) {
+  auto params = GetParam();
+
+  Global("left", ty.u32(), ast::StorageClass::kPrivate);
+  Global("right", ty.u32(), ast::StorageClass::kPrivate);
+
+  auto* left = Expr("left");
+  auto* right = Expr("right");
+
+  auto* expr = create<ast::BinaryExpression>(params.op, left, right);
+
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(), params.result);
+}
+TEST_P(GlslBinaryTest, Emit_i32) {
+  auto params = GetParam();
+
+  // Skip ops that are illegal for this type
+  if (params.op == ast::BinaryOp::kShiftLeft ||
+      params.op == ast::BinaryOp::kShiftRight) {
+    return;
+  }
+
+  Global("left", ty.i32(), ast::StorageClass::kPrivate);
+  Global("right", ty.i32(), ast::StorageClass::kPrivate);
+
+  auto* left = Expr("left");
+  auto* right = Expr("right");
+
+  auto* expr = create<ast::BinaryExpression>(params.op, left, right);
+
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(), params.result);
+}
+INSTANTIATE_TEST_SUITE_P(
+    GlslGeneratorImplTest,
+    GlslBinaryTest,
+    testing::Values(
+        BinaryData{"(left & right)", ast::BinaryOp::kAnd},
+        BinaryData{"(left | right)", ast::BinaryOp::kOr},
+        BinaryData{"(left ^ right)", ast::BinaryOp::kXor},
+        BinaryData{"(left == right)", ast::BinaryOp::kEqual},
+        BinaryData{"(left != right)", ast::BinaryOp::kNotEqual},
+        BinaryData{"(left < right)", ast::BinaryOp::kLessThan},
+        BinaryData{"(left > right)", ast::BinaryOp::kGreaterThan},
+        BinaryData{"(left <= right)", ast::BinaryOp::kLessThanEqual},
+        BinaryData{"(left >= right)", ast::BinaryOp::kGreaterThanEqual},
+        BinaryData{"(left << right)", ast::BinaryOp::kShiftLeft},
+        BinaryData{"(left >> right)", ast::BinaryOp::kShiftRight},
+        BinaryData{"(left + right)", ast::BinaryOp::kAdd},
+        BinaryData{"(left - right)", ast::BinaryOp::kSubtract},
+        BinaryData{"(left * right)", ast::BinaryOp::kMultiply},
+        BinaryData{"(left / right)", ast::BinaryOp::kDivide},
+        BinaryData{"(left % right)", ast::BinaryOp::kModulo}));
+
+TEST_F(GlslGeneratorImplTest_Binary, Multiply_VectorScalar) {
+  auto* lhs = vec3<f32>(1.f, 1.f, 1.f);
+  auto* rhs = Expr(1.f);
+
+  auto* expr =
+      create<ast::BinaryExpression>(ast::BinaryOp::kMultiply, lhs, rhs);
+
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  EXPECT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(),
+            "(vec3(1.0f, 1.0f, 1.0f) * "
+            "1.0f)");
+}
+
+TEST_F(GlslGeneratorImplTest_Binary, Multiply_ScalarVector) {
+  auto* lhs = Expr(1.f);
+  auto* rhs = vec3<f32>(1.f, 1.f, 1.f);
+
+  auto* expr =
+      create<ast::BinaryExpression>(ast::BinaryOp::kMultiply, lhs, rhs);
+
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  EXPECT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(),
+            "(1.0f * vec3(1.0f, 1.0f, "
+            "1.0f))");
+}
+
+TEST_F(GlslGeneratorImplTest_Binary, Multiply_MatrixScalar) {
+  Global("mat", ty.mat3x3<f32>(), ast::StorageClass::kPrivate);
+  auto* lhs = Expr("mat");
+  auto* rhs = Expr(1.f);
+
+  auto* expr =
+      create<ast::BinaryExpression>(ast::BinaryOp::kMultiply, lhs, rhs);
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  EXPECT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(), "(mat * 1.0f)");
+}
+
+TEST_F(GlslGeneratorImplTest_Binary, Multiply_ScalarMatrix) {
+  Global("mat", ty.mat3x3<f32>(), ast::StorageClass::kPrivate);
+  auto* lhs = Expr(1.f);
+  auto* rhs = Expr("mat");
+
+  auto* expr =
+      create<ast::BinaryExpression>(ast::BinaryOp::kMultiply, lhs, rhs);
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  EXPECT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(), "(1.0f * mat)");
+}
+
+TEST_F(GlslGeneratorImplTest_Binary, Multiply_MatrixVector) {
+  Global("mat", ty.mat3x3<f32>(), ast::StorageClass::kPrivate);
+  auto* lhs = Expr("mat");
+  auto* rhs = vec3<f32>(1.f, 1.f, 1.f);
+
+  auto* expr =
+      create<ast::BinaryExpression>(ast::BinaryOp::kMultiply, lhs, rhs);
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  EXPECT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(), "(mat * vec3(1.0f, 1.0f, 1.0f))");
+}
+
+TEST_F(GlslGeneratorImplTest_Binary, Multiply_VectorMatrix) {
+  Global("mat", ty.mat3x3<f32>(), ast::StorageClass::kPrivate);
+  auto* lhs = vec3<f32>(1.f, 1.f, 1.f);
+  auto* rhs = Expr("mat");
+
+  auto* expr =
+      create<ast::BinaryExpression>(ast::BinaryOp::kMultiply, lhs, rhs);
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  EXPECT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(), "(vec3(1.0f, 1.0f, 1.0f) * mat)");
+}
+
+TEST_F(GlslGeneratorImplTest_Binary, Multiply_MatrixMatrix) {
+  Global("lhs", ty.mat3x3<f32>(), ast::StorageClass::kPrivate);
+  Global("rhs", ty.mat3x3<f32>(), ast::StorageClass::kPrivate);
+
+  auto* expr = create<ast::BinaryExpression>(ast::BinaryOp::kMultiply,
+                                             Expr("lhs"), Expr("rhs"));
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  EXPECT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(), "(lhs * rhs)");
+}
+
+TEST_F(GlslGeneratorImplTest_Binary, Logical_And) {
+  Global("a", ty.bool_(), ast::StorageClass::kPrivate);
+  Global("b", ty.bool_(), ast::StorageClass::kPrivate);
+
+  auto* expr = create<ast::BinaryExpression>(ast::BinaryOp::kLogicalAnd,
+                                             Expr("a"), Expr("b"));
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(), "(tint_tmp)");
+  EXPECT_EQ(gen.result(), R"(bool tint_tmp = a;
+if (tint_tmp) {
+  tint_tmp = b;
+}
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Binary, Logical_Multi) {
+  // (a && b) || (c || d)
+  Global("a", ty.bool_(), ast::StorageClass::kPrivate);
+  Global("b", ty.bool_(), ast::StorageClass::kPrivate);
+  Global("c", ty.bool_(), ast::StorageClass::kPrivate);
+  Global("d", ty.bool_(), ast::StorageClass::kPrivate);
+
+  auto* expr = create<ast::BinaryExpression>(
+      ast::BinaryOp::kLogicalOr,
+      create<ast::BinaryExpression>(ast::BinaryOp::kLogicalAnd, Expr("a"),
+                                    Expr("b")),
+      create<ast::BinaryExpression>(ast::BinaryOp::kLogicalOr, Expr("c"),
+                                    Expr("d")));
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(), "(tint_tmp)");
+  EXPECT_EQ(gen.result(), R"(bool tint_tmp_1 = a;
+if (tint_tmp_1) {
+  tint_tmp_1 = b;
+}
+bool tint_tmp = (tint_tmp_1);
+if (!tint_tmp) {
+  bool tint_tmp_2 = c;
+  if (!tint_tmp_2) {
+    tint_tmp_2 = d;
+  }
+  tint_tmp = (tint_tmp_2);
+}
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Binary, Logical_Or) {
+  Global("a", ty.bool_(), ast::StorageClass::kPrivate);
+  Global("b", ty.bool_(), ast::StorageClass::kPrivate);
+
+  auto* expr = create<ast::BinaryExpression>(ast::BinaryOp::kLogicalOr,
+                                             Expr("a"), Expr("b"));
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(), "(tint_tmp)");
+  EXPECT_EQ(gen.result(), R"(bool tint_tmp = a;
+if (!tint_tmp) {
+  tint_tmp = b;
+}
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Binary, If_WithLogical) {
+  // if (a && b) {
+  //   return 1;
+  // } else if (b || c) {
+  //   return 2;
+  // } else {
+  //   return 3;
+  // }
+
+  Global("a", ty.bool_(), ast::StorageClass::kPrivate);
+  Global("b", ty.bool_(), ast::StorageClass::kPrivate);
+  Global("c", ty.bool_(), ast::StorageClass::kPrivate);
+
+  auto* body = Block(Return(3));
+  auto* else_stmt = create<ast::ElseStatement>(nullptr, body);
+
+  body = Block(Return(2));
+  auto* else_if_stmt = create<ast::ElseStatement>(
+      create<ast::BinaryExpression>(ast::BinaryOp::kLogicalOr, Expr("b"),
+                                    Expr("c")),
+      body);
+
+  body = Block(Return(1));
+
+  auto* expr = create<ast::IfStatement>(
+      create<ast::BinaryExpression>(ast::BinaryOp::kLogicalAnd, Expr("a"),
+                                    Expr("b")),
+      body,
+      ast::ElseStatementList{
+          else_if_stmt,
+          else_stmt,
+      });
+  Func("func", {}, ty.i32(), {WrapInStatement(expr), Return(0)});
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.EmitStatement(expr)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(bool tint_tmp = a;
+if (tint_tmp) {
+  tint_tmp = b;
+}
+if ((tint_tmp)) {
+  return 1;
+} else {
+  bool tint_tmp_1 = b;
+  if (!tint_tmp_1) {
+    tint_tmp_1 = c;
+  }
+  if ((tint_tmp_1)) {
+    return 2;
+  } else {
+    return 3;
+  }
+}
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Binary, Return_WithLogical) {
+  // return (a && b) || c;
+
+  Global("a", ty.bool_(), ast::StorageClass::kPrivate);
+  Global("b", ty.bool_(), ast::StorageClass::kPrivate);
+  Global("c", ty.bool_(), ast::StorageClass::kPrivate);
+
+  auto* expr = Return(create<ast::BinaryExpression>(
+      ast::BinaryOp::kLogicalOr,
+      create<ast::BinaryExpression>(ast::BinaryOp::kLogicalAnd, Expr("a"),
+                                    Expr("b")),
+      Expr("c")));
+  Func("func", {}, ty.bool_(), {WrapInStatement(expr)});
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.EmitStatement(expr)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(bool tint_tmp_1 = a;
+if (tint_tmp_1) {
+  tint_tmp_1 = b;
+}
+bool tint_tmp = (tint_tmp_1);
+if (!tint_tmp) {
+  tint_tmp = c;
+}
+return (tint_tmp);
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Binary, Assign_WithLogical) {
+  // a = (b || c) && d;
+
+  Global("a", ty.bool_(), ast::StorageClass::kPrivate);
+  Global("b", ty.bool_(), ast::StorageClass::kPrivate);
+  Global("c", ty.bool_(), ast::StorageClass::kPrivate);
+  Global("d", ty.bool_(), ast::StorageClass::kPrivate);
+
+  auto* expr = Assign(
+      Expr("a"), create<ast::BinaryExpression>(
+                     ast::BinaryOp::kLogicalAnd,
+                     create<ast::BinaryExpression>(ast::BinaryOp::kLogicalOr,
+                                                   Expr("b"), Expr("c")),
+                     Expr("d")));
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.EmitStatement(expr)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(bool tint_tmp_1 = b;
+if (!tint_tmp_1) {
+  tint_tmp_1 = c;
+}
+bool tint_tmp = (tint_tmp_1);
+if (tint_tmp) {
+  tint_tmp = d;
+}
+a = (tint_tmp);
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Binary, Decl_WithLogical) {
+  // var a : bool = (b && c) || d;
+
+  Global("b", ty.bool_(), ast::StorageClass::kPrivate);
+  Global("c", ty.bool_(), ast::StorageClass::kPrivate);
+  Global("d", ty.bool_(), ast::StorageClass::kPrivate);
+
+  auto* var = Var("a", ty.bool_(), ast::StorageClass::kNone,
+                  create<ast::BinaryExpression>(
+                      ast::BinaryOp::kLogicalOr,
+                      create<ast::BinaryExpression>(ast::BinaryOp::kLogicalAnd,
+                                                    Expr("b"), Expr("c")),
+                      Expr("d")));
+
+  auto* decl = Decl(var);
+  WrapInFunction(decl);
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.EmitStatement(decl)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(bool tint_tmp_1 = b;
+if (tint_tmp_1) {
+  tint_tmp_1 = c;
+}
+bool tint_tmp = (tint_tmp_1);
+if (!tint_tmp) {
+  tint_tmp = d;
+}
+bool a = (tint_tmp);
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Binary, Bitcast_WithLogical) {
+  // as<i32>(a && (b || c))
+
+  Global("a", ty.bool_(), ast::StorageClass::kPrivate);
+  Global("b", ty.bool_(), ast::StorageClass::kPrivate);
+  Global("c", ty.bool_(), ast::StorageClass::kPrivate);
+
+  auto* expr = create<ast::BitcastExpression>(
+      ty.i32(), create<ast::BinaryExpression>(
+                    ast::BinaryOp::kLogicalAnd, Expr("a"),
+                    create<ast::BinaryExpression>(ast::BinaryOp::kLogicalOr,
+                                                  Expr("b"), Expr("c"))));
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(bool tint_tmp = a;
+if (tint_tmp) {
+  bool tint_tmp_1 = b;
+  if (!tint_tmp_1) {
+    tint_tmp_1 = c;
+  }
+  tint_tmp = (tint_tmp_1);
+}
+)");
+  EXPECT_EQ(out.str(), R"(asint((tint_tmp)))");
+}
+
+TEST_F(GlslGeneratorImplTest_Binary, Call_WithLogical) {
+  // foo(a && b, c || d, (a || c) && (b || d))
+
+  Func("foo",
+       {
+           Param(Sym(), ty.bool_()),
+           Param(Sym(), ty.bool_()),
+           Param(Sym(), ty.bool_()),
+       },
+       ty.void_(), ast::StatementList{}, ast::DecorationList{});
+  Global("a", ty.bool_(), ast::StorageClass::kPrivate);
+  Global("b", ty.bool_(), ast::StorageClass::kPrivate);
+  Global("c", ty.bool_(), ast::StorageClass::kPrivate);
+  Global("d", ty.bool_(), ast::StorageClass::kPrivate);
+
+  ast::ExpressionList params;
+  params.push_back(create<ast::BinaryExpression>(ast::BinaryOp::kLogicalAnd,
+                                                 Expr("a"), Expr("b")));
+  params.push_back(create<ast::BinaryExpression>(ast::BinaryOp::kLogicalOr,
+                                                 Expr("c"), Expr("d")));
+  params.push_back(create<ast::BinaryExpression>(
+      ast::BinaryOp::kLogicalAnd,
+      create<ast::BinaryExpression>(ast::BinaryOp::kLogicalOr, Expr("a"),
+                                    Expr("c")),
+      create<ast::BinaryExpression>(ast::BinaryOp::kLogicalOr, Expr("b"),
+                                    Expr("d"))));
+
+  auto* expr = create<ast::CallStatement>(Call("foo", params));
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.EmitStatement(expr)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(bool tint_tmp = a;
+if (tint_tmp) {
+  tint_tmp = b;
+}
+bool tint_tmp_1 = c;
+if (!tint_tmp_1) {
+  tint_tmp_1 = d;
+}
+bool tint_tmp_3 = a;
+if (!tint_tmp_3) {
+  tint_tmp_3 = c;
+}
+bool tint_tmp_2 = (tint_tmp_3);
+if (tint_tmp_2) {
+  bool tint_tmp_4 = b;
+  if (!tint_tmp_4) {
+    tint_tmp_4 = d;
+  }
+  tint_tmp_2 = (tint_tmp_4);
+}
+foo((tint_tmp), (tint_tmp_1), (tint_tmp_2));
+)");
+}
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_bitcast_test.cc

@@ -0,0 +1,60 @@
+// 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/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using GlslGeneratorImplTest_Bitcast = TestHelper;
+
+TEST_F(GlslGeneratorImplTest_Bitcast, EmitExpression_Bitcast_Float) {
+  auto* bitcast = create<ast::BitcastExpression>(ty.f32(), Expr(1));
+  WrapInFunction(bitcast);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, bitcast)) << gen.error();
+  EXPECT_EQ(out.str(), "asfloat(1)");
+}
+
+TEST_F(GlslGeneratorImplTest_Bitcast, EmitExpression_Bitcast_Int) {
+  auto* bitcast = create<ast::BitcastExpression>(ty.i32(), Expr(1u));
+  WrapInFunction(bitcast);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, bitcast)) << gen.error();
+  EXPECT_EQ(out.str(), "asint(1u)");
+}
+
+TEST_F(GlslGeneratorImplTest_Bitcast, EmitExpression_Bitcast_Uint) {
+  auto* bitcast = create<ast::BitcastExpression>(ty.u32(), Expr(1));
+  WrapInFunction(bitcast);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, bitcast)) << gen.error();
+  EXPECT_EQ(out.str(), "asuint(1)");
+}
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_block_test.cc

@@ -0,0 +1,42 @@
+// 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/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using GlslGeneratorImplTest_Block = TestHelper;
+
+TEST_F(GlslGeneratorImplTest_Block, Emit_Block) {
+  auto* b = Block(create<ast::DiscardStatement>());
+  WrapInFunction(b);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(b)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  {
+    discard;
+  }
+)");
+}
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_break_test.cc

@@ -0,0 +1,39 @@
+// 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/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using GlslGeneratorImplTest_Break = TestHelper;
+
+TEST_F(GlslGeneratorImplTest_Break, Emit_Break) {
+  auto* b = create<ast::BreakStatement>();
+  WrapInFunction(Loop(Block(b)));
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(b)) << gen.error();
+  EXPECT_EQ(gen.result(), "  break;\n");
+}
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_call_test.cc

@@ -0,0 +1,81 @@
+// 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/ast/call_statement.h"
+#include "src/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using GlslGeneratorImplTest_Call = TestHelper;
+
+TEST_F(GlslGeneratorImplTest_Call, EmitExpression_Call_WithoutParams) {
+  Func("my_func", {}, ty.f32(), {Return(1.23f)});
+
+  auto* call = Call("my_func");
+  WrapInFunction(call);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+  EXPECT_EQ(out.str(), "my_func()");
+}
+
+TEST_F(GlslGeneratorImplTest_Call, EmitExpression_Call_WithParams) {
+  Func("my_func",
+       {
+           Param(Sym(), ty.f32()),
+           Param(Sym(), ty.f32()),
+       },
+       ty.f32(), {Return(1.23f)});
+  Global("param1", ty.f32(), ast::StorageClass::kPrivate);
+  Global("param2", ty.f32(), ast::StorageClass::kPrivate);
+
+  auto* call = Call("my_func", "param1", "param2");
+  WrapInFunction(call);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+  EXPECT_EQ(out.str(), "my_func(param1, param2)");
+}
+
+TEST_F(GlslGeneratorImplTest_Call, EmitStatement_Call) {
+  Func("my_func",
+       {
+           Param(Sym(), ty.f32()),
+           Param(Sym(), ty.f32()),
+       },
+       ty.void_(), ast::StatementList{}, ast::DecorationList{});
+  Global("param1", ty.f32(), ast::StorageClass::kPrivate);
+  Global("param2", ty.f32(), ast::StorageClass::kPrivate);
+
+  auto* call = create<ast::CallStatement>(Call("my_func", "param1", "param2"));
+  WrapInFunction(call);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+  ASSERT_TRUE(gen.EmitStatement(call)) << gen.error();
+  EXPECT_EQ(gen.result(), "  my_func(param1, param2);\n");
+}
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_case_test.cc

@@ -0,0 +1,109 @@
+// 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/ast/fallthrough_statement.h"
+#include "src/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using GlslGeneratorImplTest_Case = TestHelper;
+
+TEST_F(GlslGeneratorImplTest_Case, Emit_Case) {
+  auto* s = Switch(1, Case(Literal(5), Block(create<ast::BreakStatement>())),
+                   DefaultCase());
+  WrapInFunction(s);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitCase(s->body()[0])) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  case 5: {
+    break;
+  }
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Case, Emit_Case_BreaksByDefault) {
+  auto* s = Switch(1, Case(Literal(5), Block()), DefaultCase());
+  WrapInFunction(s);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitCase(s->body()[0])) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  case 5: {
+    break;
+  }
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Case, Emit_Case_WithFallthrough) {
+  auto* s =
+      Switch(1, Case(Literal(5), Block(create<ast::FallthroughStatement>())),
+             DefaultCase());
+  WrapInFunction(s);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitCase(s->body()[0])) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  case 5: {
+    /* fallthrough */
+  }
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Case, Emit_Case_MultipleSelectors) {
+  auto* s = Switch(
+      1, Case({Literal(5), Literal(6)}, Block(create<ast::BreakStatement>())),
+      DefaultCase());
+  WrapInFunction(s);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitCase(s->body()[0])) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  case 5:
+  case 6: {
+    break;
+  }
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Case, Emit_Case_Default) {
+  auto* s = Switch(1, DefaultCase(Block(create<ast::BreakStatement>())));
+  WrapInFunction(s);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitCase(s->body()[0])) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  default: {
+    break;
+  }
+)");
+}
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_cast_test.cc

@@ -0,0 +1,49 @@
+// 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/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using GlslGeneratorImplTest_Cast = TestHelper;
+
+TEST_F(GlslGeneratorImplTest_Cast, EmitExpression_Cast_Scalar) {
+  auto* cast = Construct<f32>(1);
+  WrapInFunction(cast);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, cast)) << gen.error();
+  EXPECT_EQ(out.str(), "float(1)");
+}
+
+TEST_F(GlslGeneratorImplTest_Cast, EmitExpression_Cast_Vector) {
+  auto* cast = vec3<f32>(vec3<i32>(1, 2, 3));
+  WrapInFunction(cast);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, cast)) << gen.error();
+  EXPECT_EQ(out.str(), "vec3(ivec3(1, 2, 3))");
+}
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_constructor_test.cc

@@ -0,0 +1,241 @@
+// 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 "gmock/gmock.h"
+#include "src/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using ::testing::HasSubstr;
+
+using GlslGeneratorImplTest_Constructor = TestHelper;
+
+TEST_F(GlslGeneratorImplTest_Constructor, EmitConstructor_Bool) {
+  WrapInFunction(Expr(false));
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("false"));
+}
+
+TEST_F(GlslGeneratorImplTest_Constructor, EmitConstructor_Int) {
+  WrapInFunction(Expr(-12345));
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("-12345"));
+}
+
+TEST_F(GlslGeneratorImplTest_Constructor, EmitConstructor_UInt) {
+  WrapInFunction(Expr(56779u));
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("56779u"));
+}
+
+TEST_F(GlslGeneratorImplTest_Constructor, EmitConstructor_Float) {
+  // Use a number close to 1<<30 but whose decimal representation ends in 0.
+  WrapInFunction(Expr(static_cast<float>((1 << 30) - 4)));
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("1073741824.0f"));
+}
+
+TEST_F(GlslGeneratorImplTest_Constructor, EmitConstructor_Type_Float) {
+  WrapInFunction(Construct<f32>(-1.2e-5f));
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("float(-0.000012f)"));
+}
+
+TEST_F(GlslGeneratorImplTest_Constructor, EmitConstructor_Type_Bool) {
+  WrapInFunction(Construct<bool>(true));
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("bool(true)"));
+}
+
+TEST_F(GlslGeneratorImplTest_Constructor, EmitConstructor_Type_Int) {
+  WrapInFunction(Construct<i32>(-12345));
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("int(-12345)"));
+}
+
+TEST_F(GlslGeneratorImplTest_Constructor, EmitConstructor_Type_Uint) {
+  WrapInFunction(Construct<u32>(12345u));
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("uint(12345u)"));
+}
+
+TEST_F(GlslGeneratorImplTest_Constructor, EmitConstructor_Type_Vec) {
+  WrapInFunction(vec3<f32>(1.f, 2.f, 3.f));
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("vec3(1.0f, 2.0f, 3.0f)"));
+}
+
+TEST_F(GlslGeneratorImplTest_Constructor, EmitConstructor_Type_Vec_Empty) {
+  WrapInFunction(vec3<f32>());
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("vec3(0.0f, 0.0f, 0.0f)"));
+}
+
+TEST_F(GlslGeneratorImplTest_Constructor,
+       EmitConstructor_Type_Vec_SingleScalar_Float) {
+  WrapInFunction(vec3<f32>(2.0f));
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("vec3((2.0f).xxx)"));
+}
+
+TEST_F(GlslGeneratorImplTest_Constructor,
+       EmitConstructor_Type_Vec_SingleScalar_Bool) {
+  WrapInFunction(vec3<bool>(true));
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("bvec3((true).xxx)"));
+}
+
+TEST_F(GlslGeneratorImplTest_Constructor,
+       EmitConstructor_Type_Vec_SingleScalar_Int) {
+  WrapInFunction(vec3<i32>(2));
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("ivec3((2).xxx)"));
+}
+
+TEST_F(GlslGeneratorImplTest_Constructor,
+       EmitConstructor_Type_Vec_SingleScalar_UInt) {
+  WrapInFunction(vec3<u32>(2u));
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("uvec3((2u).xxx)"));
+}
+
+TEST_F(GlslGeneratorImplTest_Constructor, EmitConstructor_Type_Mat) {
+  WrapInFunction(
+      mat2x3<f32>(vec3<f32>(1.f, 2.f, 3.f), vec3<f32>(3.f, 4.f, 5.f)));
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+
+  EXPECT_THAT(
+      gen.result(),
+      HasSubstr("mat2x3(vec3(1.0f, 2.0f, 3.0f), vec3(3.0f, 4.0f, 5.0f))"));
+}
+
+TEST_F(GlslGeneratorImplTest_Constructor, EmitConstructor_Type_Mat_Empty) {
+  WrapInFunction(mat2x3<f32>());
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+
+  EXPECT_THAT(gen.result(),
+              HasSubstr("mat2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)"));
+}
+
+TEST_F(GlslGeneratorImplTest_Constructor, EmitConstructor_Type_Array) {
+  WrapInFunction(Construct(ty.array(ty.vec3<f32>(), 3),
+                           vec3<f32>(1.f, 2.f, 3.f), vec3<f32>(4.f, 5.f, 6.f),
+                           vec3<f32>(7.f, 8.f, 9.f)));
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("vec3[3](vec3(1.0f, 2.0f, 3.0f), "
+                                      "vec3(4.0f, 5.0f, 6.0f), "
+                                      "vec3(7.0f, 8.0f, 9.0f))"));
+}
+
+// TODO(bclayton): Zero-init arrays
+TEST_F(GlslGeneratorImplTest_Constructor,
+       DISABLED_EmitConstructor_Type_Array_Empty) {
+  WrapInFunction(Construct(ty.array(ty.vec3<f32>(), 3)));
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(),
+              HasSubstr("{vec3(0.0f, 0.0f, 0.0f), vec3(0.0f, 0.0f, 0.0f),"
+                        " vec3(0.0f, 0.0f, 0.0f)}"));
+}
+
+TEST_F(GlslGeneratorImplTest_Constructor, EmitConstructor_Type_Struct) {
+  auto* str = Structure("S", {
+                                 Member("a", ty.i32()),
+                                 Member("b", ty.f32()),
+                                 Member("c", ty.vec3<i32>()),
+                             });
+
+  WrapInFunction(Construct(ty.Of(str), 1, 2.0f, vec3<i32>(3, 4, 5)));
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("S(1, 2.0f, ivec3(3, 4, 5))"));
+}
+
+TEST_F(GlslGeneratorImplTest_Constructor, EmitConstructor_Type_Struct_Empty) {
+  auto* str = Structure("S", {
+                                 Member("a", ty.i32()),
+                                 Member("b", ty.f32()),
+                                 Member("c", ty.vec3<i32>()),
+                             });
+
+  WrapInFunction(Construct(ty.Of(str)));
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("S(0"));
+}
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_continue_test.cc

@@ -0,0 +1,42 @@
+// 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/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using GlslGeneratorImplTest_Continue = TestHelper;
+
+TEST_F(GlslGeneratorImplTest_Continue, Emit_Continue) {
+  auto* loop = Loop(Block(create<ast::ContinueStatement>()));
+  WrapInFunction(loop);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(loop)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  while (true) {
+    continue;
+  }
+)");
+}
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_discard_test.cc

@@ -0,0 +1,39 @@
+// 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/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using GlslGeneratorImplTest_Discard = TestHelper;
+
+TEST_F(GlslGeneratorImplTest_Discard, Emit_Discard) {
+  auto* stmt = create<ast::DiscardStatement>();
+  WrapInFunction(stmt);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(stmt)) << gen.error();
+  EXPECT_EQ(gen.result(), "  discard;\n");
+}
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_identifier_test.cc

@@ -0,0 +1,40 @@
+// 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/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using GlslGeneratorImplTest_Identifier = TestHelper;
+
+TEST_F(GlslGeneratorImplTest_Identifier, EmitIdentifierExpression) {
+  Global("foo", ty.i32(), ast::StorageClass::kPrivate);
+
+  auto* i = Expr("foo");
+  WrapInFunction(i);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, i)) << gen.error();
+  EXPECT_EQ(out.str(), "foo");
+}
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_if_test.cc

@@ -0,0 +1,135 @@
+// 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/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using GlslGeneratorImplTest_If = TestHelper;
+
+TEST_F(GlslGeneratorImplTest_If, Emit_If) {
+  Global("cond", ty.bool_(), ast::StorageClass::kPrivate);
+
+  auto* cond = Expr("cond");
+  auto* body = Block(Return());
+  auto* i = If(cond, body);
+  WrapInFunction(i);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+  ASSERT_TRUE(gen.EmitStatement(i)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  if (cond) {
+    return;
+  }
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_If, Emit_IfWithElseIf) {
+  Global("cond", ty.bool_(), ast::StorageClass::kPrivate);
+  Global("else_cond", ty.bool_(), ast::StorageClass::kPrivate);
+
+  auto* else_cond = Expr("else_cond");
+  auto* else_body = Block(Return());
+
+  auto* cond = Expr("cond");
+  auto* body = Block(Return());
+  auto* i = If(
+      cond, body,
+      ast::ElseStatementList{create<ast::ElseStatement>(else_cond, else_body)});
+  WrapInFunction(i);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(i)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  if (cond) {
+    return;
+  } else {
+    if (else_cond) {
+      return;
+    }
+  }
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_If, Emit_IfWithElse) {
+  Global("cond", ty.bool_(), ast::StorageClass::kPrivate);
+
+  auto* else_body = Block(Return());
+
+  auto* cond = Expr("cond");
+  auto* body = Block(Return());
+  auto* i = If(
+      cond, body,
+      ast::ElseStatementList{create<ast::ElseStatement>(nullptr, else_body)});
+  WrapInFunction(i);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(i)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  if (cond) {
+    return;
+  } else {
+    return;
+  }
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_If, Emit_IfWithMultiple) {
+  Global("cond", ty.bool_(), ast::StorageClass::kPrivate);
+  Global("else_cond", ty.bool_(), ast::StorageClass::kPrivate);
+
+  auto* else_cond = Expr("else_cond");
+
+  auto* else_body = Block(Return());
+
+  auto* else_body_2 = Block(Return());
+
+  auto* cond = Expr("cond");
+  auto* body = Block(Return());
+  auto* i = If(cond, body,
+               ast::ElseStatementList{
+                   create<ast::ElseStatement>(else_cond, else_body),
+                   create<ast::ElseStatement>(nullptr, else_body_2),
+               });
+  WrapInFunction(i);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(i)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  if (cond) {
+    return;
+  } else {
+    if (else_cond) {
+      return;
+    } else {
+      return;
+    }
+  }
+)");
+}
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_import_test.cc

@@ -0,0 +1,282 @@
+// 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/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using GlslGeneratorImplTest_Import = TestHelper;
+
+struct GlslImportData {
+  const char* name;
+  const char* glsl_name;
+};
+inline std::ostream& operator<<(std::ostream& out, GlslImportData data) {
+  out << data.name;
+  return out;
+}
+
+using GlslImportData_SingleParamTest = TestParamHelper<GlslImportData>;
+TEST_P(GlslImportData_SingleParamTest, FloatScalar) {
+  auto param = GetParam();
+
+  auto* ident = Expr(param.name);
+  auto* expr = Call(ident, 1.f);
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitCall(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(), std::string(param.glsl_name) + "(1.0f)");
+}
+INSTANTIATE_TEST_SUITE_P(GlslGeneratorImplTest_Import,
+                         GlslImportData_SingleParamTest,
+                         testing::Values(GlslImportData{"abs", "abs"},
+                                         GlslImportData{"acos", "acos"},
+                                         GlslImportData{"asin", "asin"},
+                                         GlslImportData{"atan", "atan"},
+                                         GlslImportData{"cos", "cos"},
+                                         GlslImportData{"cosh", "cosh"},
+                                         GlslImportData{"ceil", "ceil"},
+                                         GlslImportData{"exp", "exp"},
+                                         GlslImportData{"exp2", "exp2"},
+                                         GlslImportData{"floor", "floor"},
+                                         GlslImportData{"fract", "frac"},
+                                         GlslImportData{"inverseSqrt", "rsqrt"},
+                                         GlslImportData{"length", "length"},
+                                         GlslImportData{"log", "log"},
+                                         GlslImportData{"log2", "log2"},
+                                         GlslImportData{"round", "round"},
+                                         GlslImportData{"sign", "sign"},
+                                         GlslImportData{"sin", "sin"},
+                                         GlslImportData{"sinh", "sinh"},
+                                         GlslImportData{"sqrt", "sqrt"},
+                                         GlslImportData{"tan", "tan"},
+                                         GlslImportData{"tanh", "tanh"},
+                                         GlslImportData{"trunc", "trunc"}));
+
+using GlslImportData_SingleIntParamTest = TestParamHelper<GlslImportData>;
+TEST_P(GlslImportData_SingleIntParamTest, IntScalar) {
+  auto param = GetParam();
+
+  auto* expr = Call(param.name, Expr(1));
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitCall(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(), std::string(param.glsl_name) + "(1)");
+}
+INSTANTIATE_TEST_SUITE_P(GlslGeneratorImplTest_Import,
+                         GlslImportData_SingleIntParamTest,
+                         testing::Values(GlslImportData{"abs", "abs"}));
+
+using GlslImportData_SingleVectorParamTest = TestParamHelper<GlslImportData>;
+TEST_P(GlslImportData_SingleVectorParamTest, FloatVector) {
+  auto param = GetParam();
+
+  auto* ident = Expr(param.name);
+  auto* expr = Call(ident, vec3<f32>(1.f, 2.f, 3.f));
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitCall(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(),
+            std::string(param.glsl_name) + "(vec3(1.0f, 2.0f, 3.0f))");
+}
+INSTANTIATE_TEST_SUITE_P(GlslGeneratorImplTest_Import,
+                         GlslImportData_SingleVectorParamTest,
+                         testing::Values(GlslImportData{"abs", "abs"},
+                                         GlslImportData{"acos", "acos"},
+                                         GlslImportData{"asin", "asin"},
+                                         GlslImportData{"atan", "atan"},
+                                         GlslImportData{"cos", "cos"},
+                                         GlslImportData{"cosh", "cosh"},
+                                         GlslImportData{"ceil", "ceil"},
+                                         GlslImportData{"exp", "exp"},
+                                         GlslImportData{"exp2", "exp2"},
+                                         GlslImportData{"floor", "floor"},
+                                         GlslImportData{"fract", "frac"},
+                                         GlslImportData{"inverseSqrt", "rsqrt"},
+                                         GlslImportData{"length", "length"},
+                                         GlslImportData{"log", "log"},
+                                         GlslImportData{"log2", "log2"},
+                                         GlslImportData{"normalize",
+                                                        "normalize"},
+                                         GlslImportData{"round", "round"},
+                                         GlslImportData{"sign", "sign"},
+                                         GlslImportData{"sin", "sin"},
+                                         GlslImportData{"sinh", "sinh"},
+                                         GlslImportData{"sqrt", "sqrt"},
+                                         GlslImportData{"tan", "tan"},
+                                         GlslImportData{"tanh", "tanh"},
+                                         GlslImportData{"trunc", "trunc"}));
+
+using GlslImportData_DualParam_ScalarTest = TestParamHelper<GlslImportData>;
+TEST_P(GlslImportData_DualParam_ScalarTest, Float) {
+  auto param = GetParam();
+
+  auto* expr = Call(param.name, 1.f, 2.f);
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitCall(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(), std::string(param.glsl_name) + "(1.0f, 2.0f)");
+}
+INSTANTIATE_TEST_SUITE_P(GlslGeneratorImplTest_Import,
+                         GlslImportData_DualParam_ScalarTest,
+                         testing::Values(GlslImportData{"atan2", "atan2"},
+                                         GlslImportData{"distance", "distance"},
+                                         GlslImportData{"max", "max"},
+                                         GlslImportData{"min", "min"},
+                                         GlslImportData{"pow", "pow"},
+                                         GlslImportData{"step", "step"}));
+
+using GlslImportData_DualParam_VectorTest = TestParamHelper<GlslImportData>;
+TEST_P(GlslImportData_DualParam_VectorTest, Float) {
+  auto param = GetParam();
+
+  auto* expr =
+      Call(param.name, vec3<f32>(1.f, 2.f, 3.f), vec3<f32>(4.f, 5.f, 6.f));
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitCall(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(), std::string(param.glsl_name) +
+                           "(vec3(1.0f, 2.0f, 3.0f), vec3(4.0f, 5.0f, 6.0f))");
+}
+INSTANTIATE_TEST_SUITE_P(GlslGeneratorImplTest_Import,
+                         GlslImportData_DualParam_VectorTest,
+                         testing::Values(GlslImportData{"atan2", "atan2"},
+                                         GlslImportData{"cross", "cross"},
+                                         GlslImportData{"distance", "distance"},
+                                         GlslImportData{"max", "max"},
+                                         GlslImportData{"min", "min"},
+                                         GlslImportData{"pow", "pow"},
+                                         GlslImportData{"reflect", "reflect"},
+                                         GlslImportData{"step", "step"}));
+
+using GlslImportData_DualParam_Int_Test = TestParamHelper<GlslImportData>;
+TEST_P(GlslImportData_DualParam_Int_Test, IntScalar) {
+  auto param = GetParam();
+
+  auto* expr = Call(param.name, 1, 2);
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitCall(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(), std::string(param.glsl_name) + "(1, 2)");
+}
+INSTANTIATE_TEST_SUITE_P(GlslGeneratorImplTest_Import,
+                         GlslImportData_DualParam_Int_Test,
+                         testing::Values(GlslImportData{"max", "max"},
+                                         GlslImportData{"min", "min"}));
+
+using GlslImportData_TripleParam_ScalarTest = TestParamHelper<GlslImportData>;
+TEST_P(GlslImportData_TripleParam_ScalarTest, Float) {
+  auto param = GetParam();
+
+  auto* expr = Call(param.name, 1.f, 2.f, 3.f);
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitCall(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(), std::string(param.glsl_name) + "(1.0f, 2.0f, 3.0f)");
+}
+INSTANTIATE_TEST_SUITE_P(GlslGeneratorImplTest_Import,
+                         GlslImportData_TripleParam_ScalarTest,
+                         testing::Values(GlslImportData{"fma", "mad"},
+                                         GlslImportData{"mix", "lerp"},
+                                         GlslImportData{"clamp", "clamp"},
+                                         GlslImportData{"smoothStep",
+                                                        "smoothstep"}));
+
+using GlslImportData_TripleParam_VectorTest = TestParamHelper<GlslImportData>;
+TEST_P(GlslImportData_TripleParam_VectorTest, Float) {
+  auto param = GetParam();
+
+  auto* expr = Call(param.name, vec3<f32>(1.f, 2.f, 3.f),
+                    vec3<f32>(4.f, 5.f, 6.f), vec3<f32>(7.f, 8.f, 9.f));
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitCall(out, expr)) << gen.error();
+  EXPECT_EQ(
+      out.str(),
+      std::string(param.glsl_name) +
+          R"((vec3(1.0f, 2.0f, 3.0f), vec3(4.0f, 5.0f, 6.0f), vec3(7.0f, 8.0f, 9.0f)))");
+}
+INSTANTIATE_TEST_SUITE_P(
+    GlslGeneratorImplTest_Import,
+    GlslImportData_TripleParam_VectorTest,
+    testing::Values(GlslImportData{"faceForward", "faceforward"},
+                    GlslImportData{"fma", "mad"},
+                    GlslImportData{"clamp", "clamp"},
+                    GlslImportData{"smoothStep", "smoothstep"}));
+
+TEST_F(GlslGeneratorImplTest_Import, DISABLED_GlslImportData_FMix) {
+  FAIL();
+}
+
+using GlslImportData_TripleParam_Int_Test = TestParamHelper<GlslImportData>;
+TEST_P(GlslImportData_TripleParam_Int_Test, IntScalar) {
+  auto param = GetParam();
+
+  auto* expr = Call(param.name, 1, 2, 3);
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitCall(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(), std::string(param.glsl_name) + "(1, 2, 3)");
+}
+INSTANTIATE_TEST_SUITE_P(GlslGeneratorImplTest_Import,
+                         GlslImportData_TripleParam_Int_Test,
+                         testing::Values(GlslImportData{"clamp", "clamp"}));
+
+TEST_F(GlslGeneratorImplTest_Import, GlslImportData_Determinant) {
+  Global("var", ty.mat3x3<f32>(), ast::StorageClass::kPrivate);
+
+  auto* expr = Call("determinant", "var");
+  WrapInFunction(expr);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitCall(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(), std::string("determinant(var)"));
+}
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_intrinsic_test.cc

@@ -0,0 +1,605 @@
+// 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 "gmock/gmock.h"
+#include "src/ast/call_statement.h"
+#include "src/ast/stage_decoration.h"
+#include "src/sem/call.h"
+#include "src/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using IntrinsicType = sem::IntrinsicType;
+
+using ::testing::HasSubstr;
+
+using GlslGeneratorImplTest_Intrinsic = TestHelper;
+
+enum class ParamType {
+  kF32,
+  kU32,
+  kBool,
+};
+
+struct IntrinsicData {
+  IntrinsicType intrinsic;
+  ParamType type;
+  const char* glsl_name;
+};
+inline std::ostream& operator<<(std::ostream& out, IntrinsicData data) {
+  out << data.glsl_name;
+  switch (data.type) {
+    case ParamType::kF32:
+      out << "f32";
+      break;
+    case ParamType::kU32:
+      out << "u32";
+      break;
+    case ParamType::kBool:
+      out << "bool";
+      break;
+  }
+  out << ">";
+  return out;
+}
+
+ast::CallExpression* GenerateCall(IntrinsicType intrinsic,
+                                  ParamType type,
+                                  ProgramBuilder* builder) {
+  std::string name;
+  std::ostringstream str(name);
+  str << intrinsic;
+  switch (intrinsic) {
+    case IntrinsicType::kAcos:
+    case IntrinsicType::kAsin:
+    case IntrinsicType::kAtan:
+    case IntrinsicType::kCeil:
+    case IntrinsicType::kCos:
+    case IntrinsicType::kCosh:
+    case IntrinsicType::kDpdx:
+    case IntrinsicType::kDpdxCoarse:
+    case IntrinsicType::kDpdxFine:
+    case IntrinsicType::kDpdy:
+    case IntrinsicType::kDpdyCoarse:
+    case IntrinsicType::kDpdyFine:
+    case IntrinsicType::kExp:
+    case IntrinsicType::kExp2:
+    case IntrinsicType::kFloor:
+    case IntrinsicType::kFract:
+    case IntrinsicType::kFwidth:
+    case IntrinsicType::kFwidthCoarse:
+    case IntrinsicType::kFwidthFine:
+    case IntrinsicType::kInverseSqrt:
+    case IntrinsicType::kIsFinite:
+    case IntrinsicType::kIsInf:
+    case IntrinsicType::kIsNan:
+    case IntrinsicType::kIsNormal:
+    case IntrinsicType::kLength:
+    case IntrinsicType::kLog:
+    case IntrinsicType::kLog2:
+    case IntrinsicType::kNormalize:
+    case IntrinsicType::kRound:
+    case IntrinsicType::kSin:
+    case IntrinsicType::kSinh:
+    case IntrinsicType::kSqrt:
+    case IntrinsicType::kTan:
+    case IntrinsicType::kTanh:
+    case IntrinsicType::kTrunc:
+    case IntrinsicType::kSign:
+      return builder->Call(str.str(), "f2");
+    case IntrinsicType::kLdexp:
+      return builder->Call(str.str(), "f2", "i2");
+    case IntrinsicType::kAtan2:
+    case IntrinsicType::kDot:
+    case IntrinsicType::kDistance:
+    case IntrinsicType::kPow:
+    case IntrinsicType::kReflect:
+    case IntrinsicType::kStep:
+      return builder->Call(str.str(), "f2", "f2");
+    case IntrinsicType::kCross:
+      return builder->Call(str.str(), "f3", "f3");
+    case IntrinsicType::kFma:
+    case IntrinsicType::kMix:
+    case IntrinsicType::kFaceForward:
+    case IntrinsicType::kSmoothStep:
+      return builder->Call(str.str(), "f2", "f2", "f2");
+    case IntrinsicType::kAll:
+    case IntrinsicType::kAny:
+      return builder->Call(str.str(), "b2");
+    case IntrinsicType::kAbs:
+      if (type == ParamType::kF32) {
+        return builder->Call(str.str(), "f2");
+      } else {
+        return builder->Call(str.str(), "u2");
+      }
+    case IntrinsicType::kCountOneBits:
+    case IntrinsicType::kReverseBits:
+      return builder->Call(str.str(), "u2");
+    case IntrinsicType::kMax:
+    case IntrinsicType::kMin:
+      if (type == ParamType::kF32) {
+        return builder->Call(str.str(), "f2", "f2");
+      } else {
+        return builder->Call(str.str(), "u2", "u2");
+      }
+    case IntrinsicType::kClamp:
+      if (type == ParamType::kF32) {
+        return builder->Call(str.str(), "f2", "f2", "f2");
+      } else {
+        return builder->Call(str.str(), "u2", "u2", "u2");
+      }
+    case IntrinsicType::kSelect:
+      return builder->Call(str.str(), "f2", "f2", "b2");
+    case IntrinsicType::kDeterminant:
+      return builder->Call(str.str(), "m2x2");
+    case IntrinsicType::kTranspose:
+      return builder->Call(str.str(), "m3x2");
+    default:
+      break;
+  }
+  return nullptr;
+}
+using GlslIntrinsicTest = TestParamHelper<IntrinsicData>;
+TEST_P(GlslIntrinsicTest, Emit) {
+  auto param = GetParam();
+
+  Global("f2", ty.vec2<f32>(), ast::StorageClass::kPrivate);
+  Global("f3", ty.vec3<f32>(), ast::StorageClass::kPrivate);
+  Global("u2", ty.vec2<u32>(), ast::StorageClass::kPrivate);
+  Global("i2", ty.vec2<i32>(), ast::StorageClass::kPrivate);
+  Global("b2", ty.vec2<bool>(), ast::StorageClass::kPrivate);
+  Global("m2x2", ty.mat2x2<f32>(), ast::StorageClass::kPrivate);
+  Global("m3x2", ty.mat3x2<f32>(), ast::StorageClass::kPrivate);
+
+  auto* call = GenerateCall(param.intrinsic, param.type, this);
+  ASSERT_NE(nullptr, call) << "Unhandled intrinsic";
+  Func("func", {}, ty.void_(), {Ignore(call)},
+       {create<ast::StageDecoration>(ast::PipelineStage::kFragment)});
+
+  GeneratorImpl& gen = Build();
+
+  auto* sem = program->Sem().Get(call);
+  ASSERT_NE(sem, nullptr);
+  auto* target = sem->Target();
+  ASSERT_NE(target, nullptr);
+  auto* intrinsic = target->As<sem::Intrinsic>();
+  ASSERT_NE(intrinsic, nullptr);
+
+  EXPECT_EQ(gen.generate_builtin_name(intrinsic), param.glsl_name);
+}
+INSTANTIATE_TEST_SUITE_P(
+    GlslGeneratorImplTest_Intrinsic,
+    GlslIntrinsicTest,
+    testing::Values(
+        IntrinsicData{IntrinsicType::kAbs, ParamType::kF32, "abs"},
+        IntrinsicData{IntrinsicType::kAbs, ParamType::kU32, "abs"},
+        IntrinsicData{IntrinsicType::kAcos, ParamType::kF32, "acos"},
+        IntrinsicData{IntrinsicType::kAll, ParamType::kBool, "all"},
+        IntrinsicData{IntrinsicType::kAny, ParamType::kBool, "any"},
+        IntrinsicData{IntrinsicType::kAsin, ParamType::kF32, "asin"},
+        IntrinsicData{IntrinsicType::kAtan, ParamType::kF32, "atan"},
+        IntrinsicData{IntrinsicType::kAtan2, ParamType::kF32, "atan2"},
+        IntrinsicData{IntrinsicType::kCeil, ParamType::kF32, "ceil"},
+        IntrinsicData{IntrinsicType::kClamp, ParamType::kF32, "clamp"},
+        IntrinsicData{IntrinsicType::kClamp, ParamType::kU32, "clamp"},
+        IntrinsicData{IntrinsicType::kCos, ParamType::kF32, "cos"},
+        IntrinsicData{IntrinsicType::kCosh, ParamType::kF32, "cosh"},
+        IntrinsicData{IntrinsicType::kCountOneBits, ParamType::kU32,
+                      "countbits"},
+        IntrinsicData{IntrinsicType::kCross, ParamType::kF32, "cross"},
+        IntrinsicData{IntrinsicType::kDeterminant, ParamType::kF32,
+                      "determinant"},
+        IntrinsicData{IntrinsicType::kDistance, ParamType::kF32, "distance"},
+        IntrinsicData{IntrinsicType::kDot, ParamType::kF32, "dot"},
+        IntrinsicData{IntrinsicType::kDpdx, ParamType::kF32, "ddx"},
+        IntrinsicData{IntrinsicType::kDpdxCoarse, ParamType::kF32,
+                      "ddx_coarse"},
+        IntrinsicData{IntrinsicType::kDpdxFine, ParamType::kF32, "ddx_fine"},
+        IntrinsicData{IntrinsicType::kDpdy, ParamType::kF32, "ddy"},
+        IntrinsicData{IntrinsicType::kDpdyCoarse, ParamType::kF32,
+                      "ddy_coarse"},
+        IntrinsicData{IntrinsicType::kDpdyFine, ParamType::kF32, "ddy_fine"},
+        IntrinsicData{IntrinsicType::kExp, ParamType::kF32, "exp"},
+        IntrinsicData{IntrinsicType::kExp2, ParamType::kF32, "exp2"},
+        IntrinsicData{IntrinsicType::kFaceForward, ParamType::kF32,
+                      "faceforward"},
+        IntrinsicData{IntrinsicType::kFloor, ParamType::kF32, "floor"},
+        IntrinsicData{IntrinsicType::kFma, ParamType::kF32, "mad"},
+        IntrinsicData{IntrinsicType::kFract, ParamType::kF32, "frac"},
+        IntrinsicData{IntrinsicType::kFwidth, ParamType::kF32, "fwidth"},
+        IntrinsicData{IntrinsicType::kFwidthCoarse, ParamType::kF32, "fwidth"},
+        IntrinsicData{IntrinsicType::kFwidthFine, ParamType::kF32, "fwidth"},
+        IntrinsicData{IntrinsicType::kInverseSqrt, ParamType::kF32, "rsqrt"},
+        IntrinsicData{IntrinsicType::kIsFinite, ParamType::kF32, "isfinite"},
+        IntrinsicData{IntrinsicType::kIsInf, ParamType::kF32, "isinf"},
+        IntrinsicData{IntrinsicType::kIsNan, ParamType::kF32, "isnan"},
+        IntrinsicData{IntrinsicType::kLdexp, ParamType::kF32, "ldexp"},
+        IntrinsicData{IntrinsicType::kLength, ParamType::kF32, "length"},
+        IntrinsicData{IntrinsicType::kLog, ParamType::kF32, "log"},
+        IntrinsicData{IntrinsicType::kLog2, ParamType::kF32, "log2"},
+        IntrinsicData{IntrinsicType::kMax, ParamType::kF32, "max"},
+        IntrinsicData{IntrinsicType::kMax, ParamType::kU32, "max"},
+        IntrinsicData{IntrinsicType::kMin, ParamType::kF32, "min"},
+        IntrinsicData{IntrinsicType::kMin, ParamType::kU32, "min"},
+        IntrinsicData{IntrinsicType::kMix, ParamType::kF32, "lerp"},
+        IntrinsicData{IntrinsicType::kNormalize, ParamType::kF32, "normalize"},
+        IntrinsicData{IntrinsicType::kPow, ParamType::kF32, "pow"},
+        IntrinsicData{IntrinsicType::kReflect, ParamType::kF32, "reflect"},
+        IntrinsicData{IntrinsicType::kReverseBits, ParamType::kU32,
+                      "reversebits"},
+        IntrinsicData{IntrinsicType::kRound, ParamType::kU32, "round"},
+        IntrinsicData{IntrinsicType::kSign, ParamType::kF32, "sign"},
+        IntrinsicData{IntrinsicType::kSin, ParamType::kF32, "sin"},
+        IntrinsicData{IntrinsicType::kSinh, ParamType::kF32, "sinh"},
+        IntrinsicData{IntrinsicType::kSmoothStep, ParamType::kF32,
+                      "smoothstep"},
+        IntrinsicData{IntrinsicType::kSqrt, ParamType::kF32, "sqrt"},
+        IntrinsicData{IntrinsicType::kStep, ParamType::kF32, "step"},
+        IntrinsicData{IntrinsicType::kTan, ParamType::kF32, "tan"},
+        IntrinsicData{IntrinsicType::kTanh, ParamType::kF32, "tanh"},
+        IntrinsicData{IntrinsicType::kTranspose, ParamType::kF32, "transpose"},
+        IntrinsicData{IntrinsicType::kTrunc, ParamType::kF32, "trunc"}));
+
+TEST_F(GlslGeneratorImplTest_Intrinsic, DISABLED_Intrinsic_IsNormal) {
+  FAIL();
+}
+
+TEST_F(GlslGeneratorImplTest_Intrinsic, Intrinsic_Call) {
+  auto* call = Call("dot", "param1", "param2");
+
+  Global("param1", ty.vec3<f32>(), ast::StorageClass::kPrivate);
+  Global("param2", ty.vec3<f32>(), ast::StorageClass::kPrivate);
+
+  WrapInFunction(call);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+  EXPECT_EQ(out.str(), "dot(param1, param2)");
+}
+
+TEST_F(GlslGeneratorImplTest_Intrinsic, Select_Scalar) {
+  auto* call = Call("select", 1.0f, 2.0f, true);
+  WrapInFunction(call);
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+  EXPECT_EQ(out.str(), "(true ? 2.0f : 1.0f)");
+}
+
+TEST_F(GlslGeneratorImplTest_Intrinsic, Select_Vector) {
+  auto* call =
+      Call("select", vec2<i32>(1, 2), vec2<i32>(3, 4), vec2<bool>(true, false));
+  WrapInFunction(call);
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+  EXPECT_EQ(out.str(), "(bvec2(true, false) ? ivec2(3, 4) : ivec2(1, 2))");
+}
+
+TEST_F(GlslGeneratorImplTest_Intrinsic, Modf_Scalar) {
+  auto* res = Var("res", ty.f32());
+  auto* call = Call("modf", 1.0f, AddressOf(res));
+  WrapInFunction(res, call);
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("modf(1.0f, res)"));
+}
+
+TEST_F(GlslGeneratorImplTest_Intrinsic, Modf_Vector) {
+  auto* res = Var("res", ty.vec3<f32>());
+  auto* call = Call("modf", vec3<f32>(), AddressOf(res));
+  WrapInFunction(res, call);
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("modf(vec3(0.0f, 0.0f, 0.0f), res)"));
+}
+
+#if 0
+TEST_F(GlslGeneratorImplTest_Intrinsic, Frexp_Scalar_i32) {
+  auto* exp = Var("exp", ty.i32());
+  auto* call = Call("frexp", 1.0f, AddressOf(exp));
+  WrapInFunction(exp, call);
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr(R"(
+  float tint_tmp;
+  float tint_tmp_1 = frexp(1.0f, tint_tmp);
+  exp = int(tint_tmp);
+  tint_tmp_1;
+)"));
+}
+
+TEST_F(GlslGeneratorImplTest_Intrinsic, Frexp_Vector_i32) {
+  auto* res = Var("res", ty.vec3<i32>());
+  auto* call = Call("frexp", vec3<f32>(), AddressOf(res));
+  WrapInFunction(res, call);
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr(R"(
+  vec3 tint_tmp;
+  vec3 tint_tmp_1 = frexp(vec3(0.0f, 0.0f, 0.0f), tint_tmp);
+  res = ivec3(tint_tmp);
+  tint_tmp_1;
+)"));
+}
+
+TEST_F(GlslGeneratorImplTest_Intrinsic, IsNormal_Scalar) {
+  auto* val = Var("val", ty.f32());
+  auto* call = Call("isNormal", val);
+  WrapInFunction(val, call);
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr(R"(
+  uint tint_isnormal_exponent = asuint(val) & 0x7f80000;
+  uint tint_isnormal_clamped = clamp(tint_isnormal_exponent, 0x0080000, 0x7f00000);
+  (tint_isnormal_clamped == tint_isnormal_exponent);
+)"));
+}
+
+TEST_F(GlslGeneratorImplTest_Intrinsic, IsNormal_Vector) {
+  auto* val = Var("val", ty.vec3<f32>());
+  auto* call = Call("isNormal", val);
+  WrapInFunction(val, call);
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr(R"(
+  uvec3 tint_isnormal_exponent = asuint(val) & 0x7f80000;
+  uvec3 tint_isnormal_clamped = clamp(tint_isnormal_exponent, 0x0080000, 0x7f00000);
+  (tint_isnormal_clamped == tint_isnormal_exponent);
+)"));
+}
+
+TEST_F(GlslGeneratorImplTest_Intrinsic, Pack4x8Snorm) {
+  auto* call = Call("pack4x8snorm", "p1");
+  Global("p1", ty.vec4<f32>(), ast::StorageClass::kPrivate);
+  WrapInFunction(call);
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("ivec4 tint_tmp = ivec4(round(clamp(p1, "
+                                      "-1.0, 1.0) * 127.0)) & 0xff;"));
+  EXPECT_THAT(out.str(), HasSubstr("asuint(tint_tmp.x | tint_tmp.y << 8 | "
+                                   "tint_tmp.z << 16 | tint_tmp.w << 24)"));
+}
+
+TEST_F(GlslGeneratorImplTest_Intrinsic, Pack4x8Unorm) {
+  auto* call = Call("pack4x8unorm", "p1");
+  Global("p1", ty.vec4<f32>(), ast::StorageClass::kPrivate);
+  WrapInFunction(call);
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("uvec4 tint_tmp = uvec4(round(clamp(p1, "
+                                      "0.0, 1.0) * 255.0));"));
+  EXPECT_THAT(out.str(), HasSubstr("(tint_tmp.x | tint_tmp.y << 8 | "
+                                   "tint_tmp.z << 16 | tint_tmp.w << 24)"));
+}
+
+TEST_F(GlslGeneratorImplTest_Intrinsic, Pack2x16Snorm) {
+  auto* call = Call("pack2x16snorm", "p1");
+  Global("p1", ty.vec2<f32>(), ast::StorageClass::kPrivate);
+  WrapInFunction(call);
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("int2 tint_tmp = int2(round(clamp(p1, "
+                                      "-1.0, 1.0) * 32767.0)) & 0xffff;"));
+  EXPECT_THAT(out.str(), HasSubstr("asuint(tint_tmp.x | tint_tmp.y << 16)"));
+}
+
+TEST_F(GlslGeneratorImplTest_Intrinsic, Pack2x16Unorm) {
+  auto* call = Call("pack2x16unorm", "p1");
+  Global("p1", ty.vec2<f32>(), ast::StorageClass::kPrivate);
+  WrapInFunction(call);
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("uint2 tint_tmp = uint2(round(clamp(p1, "
+                                      "0.0, 1.0) * 65535.0));"));
+  EXPECT_THAT(out.str(), HasSubstr("(tint_tmp.x | tint_tmp.y << 16)"));
+}
+
+TEST_F(GlslGeneratorImplTest_Intrinsic, Pack2x16Float) {
+  auto* call = Call("pack2x16float", "p1");
+  Global("p1", ty.vec2<f32>(), ast::StorageClass::kPrivate);
+  WrapInFunction(call);
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("uint2 tint_tmp = f32tof16(p1);"));
+  EXPECT_THAT(out.str(), HasSubstr("(tint_tmp.x | tint_tmp.y << 16)"));
+}
+
+TEST_F(GlslGeneratorImplTest_Intrinsic, Unpack4x8Snorm) {
+  auto* call = Call("unpack4x8snorm", "p1");
+  Global("p1", ty.u32(), ast::StorageClass::kPrivate);
+  WrapInFunction(call);
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("int tint_tmp_1 = int(p1);"));
+  EXPECT_THAT(gen.result(),
+              HasSubstr("ivec4 tint_tmp = ivec4(tint_tmp_1 << 24, tint_tmp_1 "
+                        "<< 16, tint_tmp_1 << 8, tint_tmp_1) >> 24;"));
+  EXPECT_THAT(out.str(),
+              HasSubstr("clamp(float4(tint_tmp) / 127.0, -1.0, 1.0)"));
+}
+
+TEST_F(GlslGeneratorImplTest_Intrinsic, Unpack4x8Unorm) {
+  auto* call = Call("unpack4x8unorm", "p1");
+  Global("p1", ty.u32(), ast::StorageClass::kPrivate);
+  WrapInFunction(call);
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("uint tint_tmp_1 = p1;"));
+  EXPECT_THAT(
+      gen.result(),
+      HasSubstr("uvec4 tint_tmp = uvec4(tint_tmp_1 & 0xff, (tint_tmp_1 >> "
+                "8) & 0xff, (tint_tmp_1 >> 16) & 0xff, tint_tmp_1 >> 24);"));
+  EXPECT_THAT(out.str(), HasSubstr("float4(tint_tmp) / 255.0"));
+}
+
+TEST_F(GlslGeneratorImplTest_Intrinsic, Unpack2x16Snorm) {
+  auto* call = Call("unpack2x16snorm", "p1");
+  Global("p1", ty.u32(), ast::StorageClass::kPrivate);
+  WrapInFunction(call);
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("int tint_tmp_1 = int(p1);"));
+  EXPECT_THAT(
+      gen.result(),
+      HasSubstr("int2 tint_tmp = int2(tint_tmp_1 << 16, tint_tmp_1) >> 16;"));
+  EXPECT_THAT(out.str(),
+              HasSubstr("clamp(float2(tint_tmp) / 32767.0, -1.0, 1.0)"));
+}
+
+TEST_F(GlslGeneratorImplTest_Intrinsic, Unpack2x16Unorm) {
+  auto* call = Call("unpack2x16unorm", "p1");
+  Global("p1", ty.u32(), ast::StorageClass::kPrivate);
+  WrapInFunction(call);
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("uint tint_tmp_1 = p1;"));
+  EXPECT_THAT(gen.result(),
+              HasSubstr("uint2 tint_tmp = uint2(tint_tmp_1 & 0xffff, "
+                        "tint_tmp_1 >> 16);"));
+  EXPECT_THAT(out.str(), HasSubstr("float2(tint_tmp) / 65535.0"));
+}
+
+TEST_F(GlslGeneratorImplTest_Intrinsic, Unpack2x16Float) {
+  auto* call = Call("unpack2x16float", "p1");
+  Global("p1", ty.u32(), ast::StorageClass::kPrivate);
+  WrapInFunction(call);
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, call)) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("uint tint_tmp = p1;"));
+  EXPECT_THAT(out.str(),
+              HasSubstr("f16tof32(uint2(tint_tmp & 0xffff, tint_tmp >> 16))"));
+}
+
+TEST_F(GlslGeneratorImplTest_Intrinsic, StorageBarrier) {
+  Func("main", {}, ty.void_(),
+       {create<ast::CallStatement>(Call("storageBarrier"))},
+       {
+           Stage(ast::PipelineStage::kCompute),
+           WorkgroupSize(1),
+       });
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_EQ(gen.result(), R"([numthreads(1, 1, 1)]
+void main() {
+  DeviceMemoryBarrierWithGroupSync();
+  return;
+}
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Intrinsic, WorkgroupBarrier) {
+  Func("main", {}, ty.void_(),
+       {create<ast::CallStatement>(Call("workgroupBarrier"))},
+       {
+           Stage(ast::PipelineStage::kCompute),
+           WorkgroupSize(1),
+       });
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_EQ(gen.result(), R"([numthreads(1, 1, 1)]
+void main() {
+  GroupMemoryBarrierWithGroupSync();
+  return;
+}
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Intrinsic, Ignore) {
+  Func("f", {Param("a", ty.i32()), Param("b", ty.i32()), Param("c", ty.i32())},
+       ty.i32(), {Return(Mul(Add("a", "b"), "c"))});
+
+  Func("main", {}, ty.void_(),
+       {create<ast::CallStatement>(Call("ignore", Call("f", 1, 2, 3)))},
+       {
+           Stage(ast::PipelineStage::kCompute),
+           WorkgroupSize(1),
+       });
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_EQ(gen.result(), R"(int f(int a, int b, int c) {
+  return ((a + b) * c);
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  f(1, 2, 3);
+  return;
+}
+)");
+}
+#endif
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_intrinsic_texture_test.cc

@@ -0,0 +1,293 @@
+// 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 "gmock/gmock.h"
+#include "src/ast/call_statement.h"
+#include "src/ast/intrinsic_texture_helper_test.h"
+#include "src/ast/stage_decoration.h"
+#include "src/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using ::testing::HasSubstr;
+
+struct ExpectedResult {
+  ExpectedResult(const char* o) : out(o) {}  // NOLINT
+
+  std::string pre;
+  std::string out;
+};
+
+ExpectedResult expected_texture_overload(
+    ast::intrinsic::test::ValidTextureOverload overload) {
+  using ValidTextureOverload = ast::intrinsic::test::ValidTextureOverload;
+  switch (overload) {
+    case ValidTextureOverload::kDimensions1d:
+    case ValidTextureOverload::kDimensionsStorageRO1d:
+    case ValidTextureOverload::kDimensionsStorageWO1d:
+    case ValidTextureOverload::kDimensions2d:
+    case ValidTextureOverload::kDimensionsDepth2d:
+    case ValidTextureOverload::kDimensionsStorageRO2d:
+    case ValidTextureOverload::kDimensionsStorageWO2d:
+    case ValidTextureOverload::kDimensionsDepthMultisampled2d:
+    case ValidTextureOverload::kDimensionsMultisampled2d:
+    case ValidTextureOverload::kDimensions2dArray:
+    case ValidTextureOverload::kDimensionsDepth2dArray:
+    case ValidTextureOverload::kDimensionsStorageRO2dArray:
+    case ValidTextureOverload::kDimensionsStorageWO2dArray:
+    case ValidTextureOverload::kDimensions3d:
+    case ValidTextureOverload::kDimensionsStorageRO3d:
+    case ValidTextureOverload::kDimensionsStorageWO3d:
+    case ValidTextureOverload::kDimensionsCube:
+    case ValidTextureOverload::kDimensionsDepthCube:
+    case ValidTextureOverload::kDimensionsCubeArray:
+    case ValidTextureOverload::kDimensionsDepthCubeArray:
+    case ValidTextureOverload::kDimensions2dLevel:
+    case ValidTextureOverload::kDimensionsDepth2dLevel:
+    case ValidTextureOverload::kDimensions2dArrayLevel:
+    case ValidTextureOverload::kDimensionsDepth2dArrayLevel:
+    case ValidTextureOverload::kDimensions3dLevel:
+    case ValidTextureOverload::kDimensionsCubeLevel:
+    case ValidTextureOverload::kDimensionsDepthCubeLevel:
+    case ValidTextureOverload::kDimensionsCubeArrayLevel:
+    case ValidTextureOverload::kDimensionsDepthCubeArrayLevel:
+      return {"textureSize"};
+    case ValidTextureOverload::kNumLayers2dArray:
+    case ValidTextureOverload::kNumLayersDepth2dArray:
+    case ValidTextureOverload::kNumLayersCubeArray:
+    case ValidTextureOverload::kNumLayersDepthCubeArray:
+    case ValidTextureOverload::kNumLayersStorageWO2dArray:
+    case ValidTextureOverload::kNumLevels2d:
+    case ValidTextureOverload::kNumLevelsCube:
+    case ValidTextureOverload::kNumLevelsDepth2d:
+    case ValidTextureOverload::kNumLevelsDepthCube:
+    case ValidTextureOverload::kNumLevels2dArray:
+    case ValidTextureOverload::kNumLevels3d:
+    case ValidTextureOverload::kNumLevelsCubeArray:
+    case ValidTextureOverload::kNumLevelsDepth2dArray:
+    case ValidTextureOverload::kNumLevelsDepthCubeArray:
+      return {"textureQueryLevels"};
+    case ValidTextureOverload::kNumSamplesDepthMultisampled2d:
+    case ValidTextureOverload::kNumSamplesMultisampled2d:
+      return {"textureSamples"};
+    case ValidTextureOverload::kSample1dF32:
+      return R"(texture.Sample(sampler, 1.0f);)";
+    case ValidTextureOverload::kSample2dF32:
+      return R"(texture.Sample(sampler, vec2(1.0f, 2.0f));)";
+    case ValidTextureOverload::kSample2dOffsetF32:
+      return R"(texture.Sample(sampler, vec2(1.0f, 2.0f), ivec2(3, 4));)";
+    case ValidTextureOverload::kSample2dArrayF32:
+      return R"(texture.Sample(sampler, vec3(1.0f, 2.0f, float(3)));)";
+    case ValidTextureOverload::kSample2dArrayOffsetF32:
+      return R"(texture.Sample(sampler, vec3(1.0f, 2.0f, float(3)), ivec2(4, 5));)";
+    case ValidTextureOverload::kSample3dF32:
+      return R"(texture.Sample(sampler, vec3(1.0f, 2.0f, 3.0f));)";
+    case ValidTextureOverload::kSample3dOffsetF32:
+      return R"(texture.Sample(sampler, vec3(1.0f, 2.0f, 3.0f), ivec3(4, 5, 6));)";
+    case ValidTextureOverload::kSampleCubeF32:
+      return R"(texture.Sample(sampler, vec3(1.0f, 2.0f, 3.0f));)";
+    case ValidTextureOverload::kSampleCubeArrayF32:
+      return R"(texture.Sample(sampler, vec4(1.0f, 2.0f, 3.0f, float(4)));)";
+    case ValidTextureOverload::kSampleDepth2dF32:
+      return R"(texture.Sample(sampler, vec2(1.0f, 2.0f)).x;)";
+    case ValidTextureOverload::kSampleDepth2dOffsetF32:
+      return R"(texture.Sample(sampler, vec2(1.0f, 2.0f), ivec2(3, 4)).x;)";
+    case ValidTextureOverload::kSampleDepth2dArrayF32:
+      return R"(texture.Sample(sampler, vec3(1.0f, 2.0f, float(3))).x;)";
+    case ValidTextureOverload::kSampleDepth2dArrayOffsetF32:
+      return R"(texture.Sample(sampler, vec3(1.0f, 2.0f, float(3)), ivec2(4, 5)).x;)";
+    case ValidTextureOverload::kSampleDepthCubeF32:
+      return R"(texture.Sample(sampler, vec3(1.0f, 2.0f, 3.0f)).x;)";
+    case ValidTextureOverload::kSampleDepthCubeArrayF32:
+      return R"(texture.Sample(sampler, vec4(1.0f, 2.0f, 3.0f, float(4))).x;)";
+    case ValidTextureOverload::kSampleBias2dF32:
+      return R"(texture.SampleBias(sampler, vec2(1.0f, 2.0f), 3.0f);)";
+    case ValidTextureOverload::kSampleBias2dOffsetF32:
+      return R"(texture.SampleBias(sampler, vec2(1.0f, 2.0f), 3.0f, ivec2(4, 5));)";
+    case ValidTextureOverload::kSampleBias2dArrayF32:
+      return R"(texture.SampleBias(sampler, vec3(1.0f, 2.0f, float(4)), 3.0f);)";
+    case ValidTextureOverload::kSampleBias2dArrayOffsetF32:
+      return R"(texture.SampleBias(sampler, vec3(1.0f, 2.0f, float(3)), 4.0f, ivec2(5, 6));)";
+    case ValidTextureOverload::kSampleBias3dF32:
+      return R"(texture.SampleBias(sampler, vec3(1.0f, 2.0f, 3.0f), 4.0f);)";
+    case ValidTextureOverload::kSampleBias3dOffsetF32:
+      return R"(texture.SampleBias(sampler, vec3(1.0f, 2.0f, 3.0f), 4.0f, ivec3(5, 6, 7));)";
+    case ValidTextureOverload::kSampleBiasCubeF32:
+      return R"(texture.SampleBias(sampler, vec3(1.0f, 2.0f, 3.0f), 4.0f);)";
+    case ValidTextureOverload::kSampleBiasCubeArrayF32:
+      return R"(texture.SampleBias(sampler, vec4(1.0f, 2.0f, 3.0f, float(3)), 4.0f);)";
+    case ValidTextureOverload::kSampleLevel2dF32:
+      return R"(texture.SampleLevel(sampler, vec2(1.0f, 2.0f), 3.0f);)";
+    case ValidTextureOverload::kSampleLevel2dOffsetF32:
+      return R"(texture.SampleLevel(sampler, vec2(1.0f, 2.0f), 3.0f, ivec2(4, 5));)";
+    case ValidTextureOverload::kSampleLevel2dArrayF32:
+      return R"(texture.SampleLevel(sampler, vec3(1.0f, 2.0f, float(3)), 4.0f);)";
+    case ValidTextureOverload::kSampleLevel2dArrayOffsetF32:
+      return R"(texture.SampleLevel(sampler, vec3(1.0f, 2.0f, float(3)), 4.0f, ivec2(5, 6));)";
+    case ValidTextureOverload::kSampleLevel3dF32:
+      return R"(texture.SampleLevel(sampler, vec3(1.0f, 2.0f, 3.0f), 4.0f);)";
+    case ValidTextureOverload::kSampleLevel3dOffsetF32:
+      return R"(texture.SampleLevel(sampler, vec3(1.0f, 2.0f, 3.0f), 4.0f, ivec3(5, 6, 7));)";
+    case ValidTextureOverload::kSampleLevelCubeF32:
+      return R"(texture.SampleLevel(sampler, vec3(1.0f, 2.0f, 3.0f), 4.0f);)";
+    case ValidTextureOverload::kSampleLevelCubeArrayF32:
+      return R"(texture.SampleLevel(sampler, vec4(1.0f, 2.0f, 3.0f, float(4)), 5.0f);)";
+    case ValidTextureOverload::kSampleLevelDepth2dF32:
+      return R"(texture.SampleLevel(sampler, vec2(1.0f, 2.0f), 3).x;)";
+    case ValidTextureOverload::kSampleLevelDepth2dOffsetF32:
+      return R"(texture.SampleLevel(sampler, vec2(1.0f, 2.0f), 3, ivec2(4, 5)).x;)";
+    case ValidTextureOverload::kSampleLevelDepth2dArrayF32:
+      return R"(texture.SampleLevel(sampler, vec3(1.0f, 2.0f, float(3)), 4).x;)";
+    case ValidTextureOverload::kSampleLevelDepth2dArrayOffsetF32:
+      return R"(texture.SampleLevel(sampler, vec3(1.0f, 2.0f, float(3)), 4, ivec2(5, 6)).x;)";
+    case ValidTextureOverload::kSampleLevelDepthCubeF32:
+      return R"(texture.SampleLevel(sampler, vec3(1.0f, 2.0f, 3.0f), 4).x;)";
+    case ValidTextureOverload::kSampleLevelDepthCubeArrayF32:
+      return R"(texture.SampleLevel(sampler, vec4(1.0f, 2.0f, 3.0f, float(4)), 5).x;)";
+    case ValidTextureOverload::kSampleGrad2dF32:
+      return R"(texture.SampleGrad(sampler, vec2(1.0f, 2.0f), vec2(3.0f, 4.0f), vec2(5.0f, 6.0f));)";
+    case ValidTextureOverload::kSampleGrad2dOffsetF32:
+      return R"(texture.SampleGrad(sampler, vec2(1.0f, 2.0f), vec2(3.0f, 4.0f), vec2(5.0f, 6.0f), ivec2(7, 7));)";
+    case ValidTextureOverload::kSampleGrad2dArrayF32:
+      return R"(texture.SampleGrad(sampler, vec3(1.0f, 2.0f, float(3)), vec2(4.0f, 5.0f), vec2(6.0f, 7.0f));)";
+    case ValidTextureOverload::kSampleGrad2dArrayOffsetF32:
+      return R"(texture.SampleGrad(sampler, vec3(1.0f, 2.0f, float(3)), vec2(4.0f, 5.0f), vec2(6.0f, 7.0f), ivec2(6, 7));)";
+    case ValidTextureOverload::kSampleGrad3dF32:
+      return R"(texture.SampleGrad(sampler, vec3(1.0f, 2.0f, 3.0f), vec3(4.0f, 5.0f, 6.0f), vec3(7.0f, 8.0f, 9.0f));)";
+    case ValidTextureOverload::kSampleGrad3dOffsetF32:
+      return R"(texture.SampleGrad(sampler, vec3(1.0f, 2.0f, 3.0f), vec3(4.0f, 5.0f, 6.0f), vec3(7.0f, 8.0f, 9.0f), ivec3(0, 1, 2));)";
+    case ValidTextureOverload::kSampleGradCubeF32:
+      return R"(texture.SampleGrad(sampler, vec3(1.0f, 2.0f, 3.0f), vec3(4.0f, 5.0f, 6.0f), vec3(7.0f, 8.0f, 9.0f));)";
+    case ValidTextureOverload::kSampleGradCubeArrayF32:
+      return R"(texture.SampleGrad(sampler, vec4(1.0f, 2.0f, 3.0f, float(4)), vec3(5.0f, 6.0f, 7.0f), vec3(8.0f, 9.0f, 10.0f));)";
+    case ValidTextureOverload::kSampleCompareDepth2dF32:
+      return R"(texture.SampleCmp(sampler, vec2(1.0f, 2.0f), 3.0f);)";
+    case ValidTextureOverload::kSampleCompareDepth2dOffsetF32:
+      return R"(texture.SampleCmp(sampler, vec2(1.0f, 2.0f), 3.0f, ivec2(4, 5));)";
+    case ValidTextureOverload::kSampleCompareDepth2dArrayF32:
+      return R"(texture.SampleCmp(sampler, vec3(1.0f, 2.0f, float(4)), 3.0f);)";
+    case ValidTextureOverload::kSampleCompareDepth2dArrayOffsetF32:
+      return R"(texture.SampleCmp(sampler, vec3(1.0f, 2.0f, float(4)), 3.0f, ivec2(5, 6));)";
+    case ValidTextureOverload::kSampleCompareDepthCubeF32:
+      return R"(texture.SampleCmp(sampler, vec3(1.0f, 2.0f, 3.0f), 4.0f);)";
+    case ValidTextureOverload::kSampleCompareDepthCubeArrayF32:
+      return R"(texture.SampleCmp(sampler, vec4(1.0f, 2.0f, 3.0f, float(4)), 5.0f);)";
+    case ValidTextureOverload::kSampleCompareLevelDepth2dF32:
+      return R"(texture.SampleCmpLevelZero(sampler, vec2(1.0f, 2.0f), 3.0f);)";
+    case ValidTextureOverload::kSampleCompareLevelDepth2dOffsetF32:
+      return R"(texture.SampleCmpLevelZero(sampler, vec2(1.0f, 2.0f), 3.0f, ivec2(4, 5));)";
+    case ValidTextureOverload::kSampleCompareLevelDepth2dArrayF32:
+      return R"(texture.SampleCmpLevelZero(sampler, vec3(1.0f, 2.0f, float(4)), 3.0f);)";
+    case ValidTextureOverload::kSampleCompareLevelDepth2dArrayOffsetF32:
+      return R"(texture.SampleCmpLevelZero(sampler, vec3(1.0f, 2.0f, float(4)), 3.0f, ivec2(5, 6));)";
+    case ValidTextureOverload::kSampleCompareLevelDepthCubeF32:
+      return R"(texture.SampleCmpLevelZero(sampler, vec3(1.0f, 2.0f, 3.0f), 4.0f);)";
+    case ValidTextureOverload::kSampleCompareLevelDepthCubeArrayF32:
+      return R"(texture.SampleCmpLevelZero(sampler, vec4(1.0f, 2.0f, 3.0f, float(4)), 5.0f);)";
+    case ValidTextureOverload::kLoad1dLevelF32:
+    case ValidTextureOverload::kLoad1dLevelU32:
+    case ValidTextureOverload::kLoad1dLevelI32:
+      return R"(texture.Load(ivec2(1, 3));)";
+    case ValidTextureOverload::kLoad2dLevelF32:
+    case ValidTextureOverload::kLoad2dLevelU32:
+    case ValidTextureOverload::kLoad2dLevelI32:
+      return R"(texture.Load(ivec3(1, 2, 3));)";
+    case ValidTextureOverload::kLoad2dArrayLevelF32:
+    case ValidTextureOverload::kLoad2dArrayLevelU32:
+    case ValidTextureOverload::kLoad2dArrayLevelI32:
+    case ValidTextureOverload::kLoad3dLevelF32:
+    case ValidTextureOverload::kLoad3dLevelU32:
+    case ValidTextureOverload::kLoad3dLevelI32:
+      return R"(texture.Load(ivec4(1, 2, 3, 4));)";
+    case ValidTextureOverload::kLoadDepthMultisampled2dF32:
+    case ValidTextureOverload::kLoadMultisampled2dF32:
+    case ValidTextureOverload::kLoadMultisampled2dU32:
+    case ValidTextureOverload::kLoadMultisampled2dI32:
+      return R"(texture.Load(ivec2(1, 2), 3);)";
+    case ValidTextureOverload::kLoadDepth2dLevelF32:
+      return R"(texture.Load(ivec3(1, 2, 3)).x;)";
+    case ValidTextureOverload::kLoadDepth2dArrayLevelF32:
+      return R"(texture.Load(ivec4(1, 2, 3, 4)).x;)";
+    case ValidTextureOverload::kLoadStorageRO1dRgba32float:
+      return R"(texture.Load(ivec2(1, 0));)";
+    case ValidTextureOverload::kLoadStorageRO2dRgba8unorm:
+    case ValidTextureOverload::kLoadStorageRO2dRgba8snorm:
+    case ValidTextureOverload::kLoadStorageRO2dRgba8uint:
+    case ValidTextureOverload::kLoadStorageRO2dRgba8sint:
+    case ValidTextureOverload::kLoadStorageRO2dRgba16uint:
+    case ValidTextureOverload::kLoadStorageRO2dRgba16sint:
+    case ValidTextureOverload::kLoadStorageRO2dRgba16float:
+    case ValidTextureOverload::kLoadStorageRO2dR32uint:
+    case ValidTextureOverload::kLoadStorageRO2dR32sint:
+    case ValidTextureOverload::kLoadStorageRO2dR32float:
+    case ValidTextureOverload::kLoadStorageRO2dRg32uint:
+    case ValidTextureOverload::kLoadStorageRO2dRg32sint:
+    case ValidTextureOverload::kLoadStorageRO2dRg32float:
+    case ValidTextureOverload::kLoadStorageRO2dRgba32uint:
+    case ValidTextureOverload::kLoadStorageRO2dRgba32sint:
+    case ValidTextureOverload::kLoadStorageRO2dRgba32float:
+      return R"(texture.Load(ivec3(1, 2, 0));)";
+    case ValidTextureOverload::kLoadStorageRO2dArrayRgba32float:
+    case ValidTextureOverload::kLoadStorageRO3dRgba32float:
+      return R"(texture.Load(ivec4(1, 2, 3, 0));)";
+    case ValidTextureOverload::kStoreWO1dRgba32float:
+      return R"(texture[1] = vec4(2.0f, 3.0f, 4.0f, 5.0f);)";
+    case ValidTextureOverload::kStoreWO2dRgba32float:
+      return R"(texture[ivec2(1, 2)] = vec4(3.0f, 4.0f, 5.0f, 6.0f);)";
+    case ValidTextureOverload::kStoreWO2dArrayRgba32float:
+      return R"(texture[ivec3(1, 2, 3)] = vec4(4.0f, 5.0f, 6.0f, 7.0f);)";
+    case ValidTextureOverload::kStoreWO3dRgba32float:
+      return R"(texture[ivec3(1, 2, 3)] = vec4(4.0f, 5.0f, 6.0f, 7.0f);)";
+  }
+  return "<unmatched texture overload>";
+}  // NOLINT - Ignore the length of this function
+
+class GlslGeneratorIntrinsicTextureTest
+    : public TestParamHelper<ast::intrinsic::test::TextureOverloadCase> {};
+
+TEST_P(GlslGeneratorIntrinsicTextureTest, Call) {
+  auto param = GetParam();
+
+  param.buildTextureVariable(this);
+  param.buildSamplerVariable(this);
+
+  auto* call = Call(param.function, param.args(this));
+  auto* stmt = ast::intrinsic::test::ReturnsVoid(param.overload)
+                   ? create<ast::CallStatement>(call)
+                   : Ignore(call);
+
+  Func("main", {}, ty.void_(), {stmt}, {Stage(ast::PipelineStage::kFragment)});
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+
+  auto expected = expected_texture_overload(param.overload);
+
+  EXPECT_THAT(gen.result(), HasSubstr(expected.pre));
+  EXPECT_THAT(gen.result(), HasSubstr(expected.out));
+}
+
+INSTANTIATE_TEST_SUITE_P(
+    GlslGeneratorIntrinsicTextureTest,
+    GlslGeneratorIntrinsicTextureTest,
+    testing::ValuesIn(ast::intrinsic::test::TextureOverloadCase::ValidCases()));
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_loop_test.cc

@@ -0,0 +1,418 @@
+// 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/ast/variable_decl_statement.h"
+#include "src/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using GlslGeneratorImplTest_Loop = TestHelper;
+
+TEST_F(GlslGeneratorImplTest_Loop, Emit_Loop) {
+  auto* body = Block(create<ast::DiscardStatement>());
+  auto* continuing = Block();
+  auto* l = Loop(body, continuing);
+
+  WrapInFunction(l);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(l)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  while (true) {
+    discard;
+  }
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Loop, Emit_LoopWithContinuing) {
+  Func("a_statement", {}, ty.void_(), {});
+
+  auto* body = Block(create<ast::DiscardStatement>());
+  auto* continuing = Block(create<ast::CallStatement>(Call("a_statement")));
+  auto* l = Loop(body, continuing);
+
+  WrapInFunction(l);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(l)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  while (true) {
+    discard;
+    {
+      a_statement();
+    }
+  }
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Loop, Emit_LoopNestedWithContinuing) {
+  Func("a_statement", {}, ty.void_(), {});
+
+  Global("lhs", ty.f32(), ast::StorageClass::kPrivate);
+  Global("rhs", ty.f32(), ast::StorageClass::kPrivate);
+
+  auto* body = Block(create<ast::DiscardStatement>());
+  auto* continuing = Block(create<ast::CallStatement>(Call("a_statement")));
+  auto* inner = Loop(body, continuing);
+
+  body = Block(inner);
+
+  auto* lhs = Expr("lhs");
+  auto* rhs = Expr("rhs");
+
+  continuing = Block(Assign(lhs, rhs));
+
+  auto* outer = Loop(body, continuing);
+  WrapInFunction(outer);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(outer)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  while (true) {
+    while (true) {
+      discard;
+      {
+        a_statement();
+      }
+    }
+    {
+      lhs = rhs;
+    }
+  }
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Loop, Emit_LoopWithVarUsedInContinuing) {
+  // loop {
+  //   var lhs : f32 = 2.4;
+  //   var other : f32;
+  //   continuing {
+  //     lhs = rhs
+  //   }
+  // }
+  //
+  // ->
+  // {
+  //   float lhs;
+  //   float other;
+  //   for (;;) {
+  //     if (continuing) {
+  //       lhs = rhs;
+  //     }
+  //     lhs = 2.4f;
+  //     other = 0.0f;
+  //   }
+  // }
+
+  Global("rhs", ty.f32(), ast::StorageClass::kPrivate);
+
+  auto* var = Var("lhs", ty.f32(), ast::StorageClass::kNone, Expr(2.4f));
+
+  auto* body = Block(Decl(var), Decl(Var("other", ty.f32())));
+
+  auto* lhs = Expr("lhs");
+  auto* rhs = Expr("rhs");
+
+  auto* continuing = Block(Assign(lhs, rhs));
+  auto* outer = Loop(body, continuing);
+  WrapInFunction(outer);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(outer)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  while (true) {
+    float lhs = 2.400000095f;
+    float other = 0.0f;
+    {
+      lhs = rhs;
+    }
+  }
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Loop, Emit_ForLoop) {
+  // for(; ; ) {
+  //   return;
+  // }
+
+  Func("a_statement", {}, ty.void_(), {});
+
+  auto* f = For(nullptr, nullptr, nullptr,
+                Block(create<ast::CallStatement>(Call("a_statement"))));
+  WrapInFunction(f);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  {
+    for(; ; ) {
+      a_statement();
+    }
+  }
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Loop, Emit_ForLoopWithSimpleInit) {
+  // for(var i : i32; ; ) {
+  //   return;
+  // }
+
+  Func("a_statement", {}, ty.void_(), {});
+
+  auto* f = For(Decl(Var("i", ty.i32())), nullptr, nullptr,
+                Block(create<ast::CallStatement>(Call("a_statement"))));
+  WrapInFunction(f);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  {
+    for(int i = 0; ; ) {
+      a_statement();
+    }
+  }
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Loop, Emit_ForLoopWithMultiStmtInit) {
+  // for(var b = true && false; ; ) {
+  //   return;
+  // }
+  Func("a_statement", {}, ty.void_(), {});
+
+  auto* multi_stmt = create<ast::BinaryExpression>(ast::BinaryOp::kLogicalAnd,
+                                                   Expr(true), Expr(false));
+  auto* f = For(Decl(Var("b", nullptr, multi_stmt)), nullptr, nullptr,
+                Block(create<ast::CallStatement>(Call("a_statement"))));
+  WrapInFunction(f);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  {
+    bool tint_tmp = true;
+    if (tint_tmp) {
+      tint_tmp = false;
+    }
+    bool b = (tint_tmp);
+    for(; ; ) {
+      a_statement();
+    }
+  }
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Loop, Emit_ForLoopWithSimpleCond) {
+  // for(; true; ) {
+  //   return;
+  // }
+
+  Func("a_statement", {}, ty.void_(), {});
+
+  auto* f = For(nullptr, true, nullptr,
+                Block(create<ast::CallStatement>(Call("a_statement"))));
+  WrapInFunction(f);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  {
+    for(; true; ) {
+      a_statement();
+    }
+  }
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Loop, Emit_ForLoopWithMultiStmtCond) {
+  // for(; true && false; ) {
+  //   return;
+  // }
+
+  Func("a_statement", {}, ty.void_(), {});
+
+  auto* multi_stmt = create<ast::BinaryExpression>(ast::BinaryOp::kLogicalAnd,
+                                                   Expr(true), Expr(false));
+  auto* f = For(nullptr, multi_stmt, nullptr,
+                Block(create<ast::CallStatement>(Call("a_statement"))));
+  WrapInFunction(f);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  {
+    while (true) {
+      bool tint_tmp = true;
+      if (tint_tmp) {
+        tint_tmp = false;
+      }
+      if (!((tint_tmp))) { break; }
+      a_statement();
+    }
+  }
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Loop, Emit_ForLoopWithSimpleCont) {
+  // for(; ; i = i + 1) {
+  //   return;
+  // }
+
+  Func("a_statement", {}, ty.void_(), {});
+
+  auto* v = Decl(Var("i", ty.i32()));
+  auto* f = For(nullptr, nullptr, Assign("i", Add("i", 1)),
+                Block(create<ast::CallStatement>(Call("a_statement"))));
+  WrapInFunction(v, f);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  {
+    for(; ; i = (i + 1)) {
+      a_statement();
+    }
+  }
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Loop, Emit_ForLoopWithMultiStmtCont) {
+  // for(; ; i = true && false) {
+  //   return;
+  // }
+
+  Func("a_statement", {}, ty.void_(), {});
+
+  auto* multi_stmt = create<ast::BinaryExpression>(ast::BinaryOp::kLogicalAnd,
+                                                   Expr(true), Expr(false));
+  auto* v = Decl(Var("i", ty.bool_()));
+  auto* f = For(nullptr, nullptr, Assign("i", multi_stmt),
+                Block(create<ast::CallStatement>(Call("a_statement"))));
+  WrapInFunction(v, f);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  {
+    while (true) {
+      a_statement();
+      bool tint_tmp = true;
+      if (tint_tmp) {
+        tint_tmp = false;
+      }
+      i = (tint_tmp);
+    }
+  }
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Loop, Emit_ForLoopWithSimpleInitCondCont) {
+  // for(var i : i32; true; i = i + 1) {
+  //   return;
+  // }
+
+  Func("a_statement", {}, ty.void_(), {});
+
+  auto* f = For(Decl(Var("i", ty.i32())), true, Assign("i", Add("i", 1)),
+                Block(create<ast::CallStatement>(Call("a_statement"))));
+  WrapInFunction(f);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  {
+    for(int i = 0; true; i = (i + 1)) {
+      a_statement();
+    }
+  }
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Loop, Emit_ForLoopWithMultiStmtInitCondCont) {
+  // for(var i = true && false; true && false; i = true && false) {
+  //   return;
+  // }
+  Func("a_statement", {}, ty.void_(), {});
+
+  auto* multi_stmt_a = create<ast::BinaryExpression>(ast::BinaryOp::kLogicalAnd,
+                                                     Expr(true), Expr(false));
+  auto* multi_stmt_b = create<ast::BinaryExpression>(ast::BinaryOp::kLogicalAnd,
+                                                     Expr(true), Expr(false));
+  auto* multi_stmt_c = create<ast::BinaryExpression>(ast::BinaryOp::kLogicalAnd,
+                                                     Expr(true), Expr(false));
+
+  auto* f = For(Decl(Var("i", nullptr, multi_stmt_a)), multi_stmt_b,
+                Assign("i", multi_stmt_c),
+                Block(create<ast::CallStatement>(Call("a_statement"))));
+  WrapInFunction(f);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  {
+    bool tint_tmp = true;
+    if (tint_tmp) {
+      tint_tmp = false;
+    }
+    bool i = (tint_tmp);
+    while (true) {
+      bool tint_tmp_1 = true;
+      if (tint_tmp_1) {
+        tint_tmp_1 = false;
+      }
+      if (!((tint_tmp_1))) { break; }
+      a_statement();
+      bool tint_tmp_2 = true;
+      if (tint_tmp_2) {
+        tint_tmp_2 = false;
+      }
+      i = (tint_tmp_2);
+    }
+  }
+)");
+}
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_member_accessor_test.cc

@@ -0,0 +1,816 @@
+// 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 "gmock/gmock.h"
+#include "src/ast/stage_decoration.h"
+#include "src/ast/struct_block_decoration.h"
+#include "src/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using ::testing::HasSubstr;
+
+using create_type_func_ptr =
+    ast::Type* (*)(const ProgramBuilder::TypesBuilder& ty);
+
+inline ast::Type* ty_i32(const ProgramBuilder::TypesBuilder& ty) {
+  return ty.i32();
+}
+inline ast::Type* ty_u32(const ProgramBuilder::TypesBuilder& ty) {
+  return ty.u32();
+}
+inline ast::Type* ty_f32(const ProgramBuilder::TypesBuilder& ty) {
+  return ty.f32();
+}
+template <typename T>
+inline ast::Type* ty_vec2(const ProgramBuilder::TypesBuilder& ty) {
+  return ty.vec2<T>();
+}
+template <typename T>
+inline ast::Type* ty_vec3(const ProgramBuilder::TypesBuilder& ty) {
+  return ty.vec3<T>();
+}
+template <typename T>
+inline ast::Type* ty_vec4(const ProgramBuilder::TypesBuilder& ty) {
+  return ty.vec4<T>();
+}
+template <typename T>
+inline ast::Type* ty_mat2x2(const ProgramBuilder::TypesBuilder& ty) {
+  return ty.mat2x2<T>();
+}
+template <typename T>
+inline ast::Type* ty_mat2x3(const ProgramBuilder::TypesBuilder& ty) {
+  return ty.mat2x3<T>();
+}
+template <typename T>
+inline ast::Type* ty_mat2x4(const ProgramBuilder::TypesBuilder& ty) {
+  return ty.mat2x4<T>();
+}
+template <typename T>
+inline ast::Type* ty_mat3x2(const ProgramBuilder::TypesBuilder& ty) {
+  return ty.mat3x2<T>();
+}
+template <typename T>
+inline ast::Type* ty_mat3x3(const ProgramBuilder::TypesBuilder& ty) {
+  return ty.mat3x3<T>();
+}
+template <typename T>
+inline ast::Type* ty_mat3x4(const ProgramBuilder::TypesBuilder& ty) {
+  return ty.mat3x4<T>();
+}
+template <typename T>
+inline ast::Type* ty_mat4x2(const ProgramBuilder::TypesBuilder& ty) {
+  return ty.mat4x2<T>();
+}
+template <typename T>
+inline ast::Type* ty_mat4x3(const ProgramBuilder::TypesBuilder& ty) {
+  return ty.mat4x3<T>();
+}
+template <typename T>
+inline ast::Type* ty_mat4x4(const ProgramBuilder::TypesBuilder& ty) {
+  return ty.mat4x4<T>();
+}
+
+using i32 = ProgramBuilder::i32;
+using u32 = ProgramBuilder::u32;
+using f32 = ProgramBuilder::f32;
+
+template <typename BASE>
+class GlslGeneratorImplTest_MemberAccessorBase : public BASE {
+ public:
+  void SetupStorageBuffer(ast::StructMemberList members) {
+    ProgramBuilder& b = *this;
+
+    auto* s =
+        b.Structure("Data", members, {b.create<ast::StructBlockDecoration>()});
+
+    b.Global("data", b.ty.Of(s), ast::StorageClass::kStorage,
+             ast::Access::kReadWrite,
+             ast::DecorationList{
+                 b.create<ast::BindingDecoration>(0),
+                 b.create<ast::GroupDecoration>(1),
+             });
+  }
+
+  void SetupFunction(ast::StatementList statements) {
+    ProgramBuilder& b = *this;
+    b.Func("main", ast::VariableList{}, b.ty.void_(), statements,
+           ast::DecorationList{
+               b.Stage(ast::PipelineStage::kFragment),
+           });
+  }
+};
+
+using GlslGeneratorImplTest_MemberAccessor =
+    GlslGeneratorImplTest_MemberAccessorBase<TestHelper>;
+
+template <typename T>
+using GlslGeneratorImplTest_MemberAccessorWithParam =
+    GlslGeneratorImplTest_MemberAccessorBase<TestParamHelper<T>>;
+
+TEST_F(GlslGeneratorImplTest_MemberAccessor, EmitExpression_MemberAccessor) {
+  auto* s = Structure("Data", {Member("mem", ty.f32())});
+  Global("str", ty.Of(s), ast::StorageClass::kPrivate);
+
+  auto* expr = MemberAccessor("str", "mem");
+  WrapInFunction(Var("expr", ty.f32(), ast::StorageClass::kNone, expr));
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_EQ(gen.result(), R"(#version 310 es
+precision mediump float;
+
+struct Data {
+  float mem;
+};
+
+static Data str = Data(0.0f);
+
+[numthreads(1, 1, 1)]
+void test_function() {
+  float expr = str.mem;
+  return;
+}
+void main() {
+  test_function();
+}
+
+
+)");
+}
+
+struct TypeCase {
+  create_type_func_ptr member_type;
+  std::string expected;
+};
+inline std::ostream& operator<<(std::ostream& out, TypeCase c) {
+  ProgramBuilder b;
+  auto* ty = c.member_type(b.ty);
+  out << ty->FriendlyName(b.Symbols());
+  return out;
+}
+
+using GlslGeneratorImplTest_MemberAccessor_StorageBufferLoad =
+    GlslGeneratorImplTest_MemberAccessorWithParam<TypeCase>;
+TEST_P(GlslGeneratorImplTest_MemberAccessor_StorageBufferLoad, Test) {
+  // struct Data {
+  //   a : i32;
+  //   b : <type>;
+  // };
+  // var<storage> data : Data;
+  // data.b;
+
+  auto p = GetParam();
+
+  SetupStorageBuffer({
+      Member("a", ty.i32()),
+      Member("b", p.member_type(ty)),
+  });
+
+  SetupFunction({
+      Decl(Var("x", nullptr, ast::StorageClass::kNone,
+               MemberAccessor("data", "b"))),
+  });
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr(p.expected));
+}
+
+INSTANTIATE_TEST_SUITE_P(GlslGeneratorImplTest_MemberAccessor,
+                         GlslGeneratorImplTest_MemberAccessor_StorageBufferLoad,
+                         testing::Values(TypeCase{ty_u32, "data.b"},
+                                         TypeCase{ty_f32, "data.b"},
+                                         TypeCase{ty_i32, "data.b"},
+                                         TypeCase{ty_vec2<u32>, "data.b"},
+                                         TypeCase{ty_vec2<f32>, "data.b"},
+                                         TypeCase{ty_vec2<i32>, "data.b"},
+                                         TypeCase{ty_vec3<u32>, "data.b"},
+                                         TypeCase{ty_vec3<f32>, "data.b"},
+                                         TypeCase{ty_vec3<i32>, "data.b"},
+                                         TypeCase{ty_vec4<u32>, "data.b"},
+                                         TypeCase{ty_vec4<f32>, "data.b"},
+                                         TypeCase{ty_vec4<i32>, "data.b"},
+                                         TypeCase{ty_mat2x2<f32>, "data.b"},
+                                         TypeCase{ty_mat2x3<f32>, "data.b"},
+                                         TypeCase{ty_mat2x4<f32>, "data.b"},
+                                         TypeCase{ty_mat3x2<f32>, "data.b"},
+                                         TypeCase{ty_mat3x3<f32>, "data.b"},
+                                         TypeCase{ty_mat3x4<f32>, "data.b"},
+                                         TypeCase{ty_mat4x2<f32>, "data.b"},
+                                         TypeCase{ty_mat4x3<f32>, "data.b"},
+                                         TypeCase{ty_mat4x4<f32>, "data.b"}));
+
+using GlslGeneratorImplTest_MemberAccessor_StorageBufferStore =
+    GlslGeneratorImplTest_MemberAccessorWithParam<TypeCase>;
+TEST_P(GlslGeneratorImplTest_MemberAccessor_StorageBufferStore, Test) {
+  // struct Data {
+  //   a : i32;
+  //   b : <type>;
+  // };
+  // var<storage> data : Data;
+  // data.b = <type>();
+
+  auto p = GetParam();
+
+  SetupStorageBuffer({
+      Member("a", ty.i32()),
+      Member("b", p.member_type(ty)),
+  });
+
+  SetupFunction({
+      Decl(Var("value", p.member_type(ty), ast::StorageClass::kNone,
+               Construct(p.member_type(ty)))),
+      Assign(MemberAccessor("data", "b"), Expr("value")),
+  });
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr(p.expected));
+}
+
+INSTANTIATE_TEST_SUITE_P(
+    GlslGeneratorImplTest_MemberAccessor,
+    GlslGeneratorImplTest_MemberAccessor_StorageBufferStore,
+    testing::Values(TypeCase{ty_u32, "data.b = value"},
+                    TypeCase{ty_f32, "data.b = value"},
+                    TypeCase{ty_i32, "data.b = value"},
+                    TypeCase{ty_vec2<u32>, "data.b = value"},
+                    TypeCase{ty_vec2<f32>, "data.b = value"},
+                    TypeCase{ty_vec2<i32>, "data.b = value"},
+                    TypeCase{ty_vec3<u32>, "data.b = value"},
+                    TypeCase{ty_vec3<f32>, "data.b = value"},
+                    TypeCase{ty_vec3<i32>, "data.b = value"},
+                    TypeCase{ty_vec4<u32>, "data.b = value"},
+                    TypeCase{ty_vec4<f32>, "data.b = value"},
+                    TypeCase{ty_vec4<i32>, "data.b = value"},
+                    TypeCase{ty_mat2x2<f32>, "data.b = value"},
+                    TypeCase{ty_mat2x3<f32>, "data.b = value"},
+                    TypeCase{ty_mat2x4<f32>, "data.b = value"},
+                    TypeCase{ty_mat3x2<f32>, "data.b = value"},
+                    TypeCase{ty_mat3x3<f32>, "data.b = value"},
+                    TypeCase{ty_mat3x4<f32>, "data.b = value"},
+                    TypeCase{ty_mat4x2<f32>, "data.b = value"},
+                    TypeCase{ty_mat4x3<f32>, "data.b = value"},
+                    TypeCase{ty_mat4x4<f32>, "data.b = value"}));
+
+TEST_F(GlslGeneratorImplTest_MemberAccessor, StorageBuffer_Store_Matrix_Empty) {
+  // struct Data {
+  //   z : f32;
+  //   a : mat2x3<f32>;
+  // };
+  // var<storage> data : Data;
+  // data.a = mat2x3<f32>();
+
+  SetupStorageBuffer({
+      Member("a", ty.i32()),
+      Member("b", ty.mat2x3<f32>()),
+  });
+
+  SetupFunction({
+      Assign(MemberAccessor("data", "b"),
+             Construct(ty.mat2x3<f32>(), ast::ExpressionList{})),
+  });
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  auto* expected =
+      R"(#version 310 es
+precision mediump float;
+
+
+Data data : register(u0, space1);
+
+void main() {
+  data.b = mat2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+  return;
+}
+void main() {
+  main();
+}
+
+
+)";
+  EXPECT_EQ(gen.result(), expected);
+}
+
+TEST_F(GlslGeneratorImplTest_MemberAccessor,
+       StorageBuffer_Load_Matrix_Single_Element) {
+  // struct Data {
+  //   z : f32;
+  //   a : mat4x3<f32>;
+  // };
+  // var<storage> data : Data;
+  // data.a[2][1];
+
+  SetupStorageBuffer({
+      Member("z", ty.f32()),
+      Member("a", ty.mat4x3<f32>()),
+  });
+
+  SetupFunction({
+      Decl(
+          Var("x", nullptr, ast::StorageClass::kNone,
+              IndexAccessor(IndexAccessor(MemberAccessor("data", "a"), 2), 1))),
+  });
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  auto* expected =
+      R"(#version 310 es
+precision mediump float;
+
+
+Data data : register(u0, space1);
+
+void main() {
+  float x = data.a[2][1];
+  return;
+}
+void main() {
+  main();
+}
+
+
+)";
+  EXPECT_EQ(gen.result(), expected);
+}
+
+TEST_F(GlslGeneratorImplTest_MemberAccessor,
+       EmitExpression_ArrayAccessor_StorageBuffer_Load_Int_FromArray) {
+  // struct Data {
+  //   a : [[stride(4)]] array<i32, 5>;
+  // };
+  // var<storage> data : Data;
+  // data.a[2];
+
+  SetupStorageBuffer({
+      Member("z", ty.f32()),
+      Member("a", ty.array<i32, 5>(4)),
+  });
+
+  SetupFunction({
+      Decl(Var("x", nullptr, ast::StorageClass::kNone,
+               IndexAccessor(MemberAccessor("data", "a"), 2))),
+  });
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  auto* expected =
+      R"(#version 310 es
+precision mediump float;
+
+
+Data data : register(u0, space1);
+
+void main() {
+  int x = data.a[2];
+  return;
+}
+void main() {
+  main();
+}
+
+
+)";
+  EXPECT_EQ(gen.result(), expected);
+}
+
+TEST_F(GlslGeneratorImplTest_MemberAccessor,
+       EmitExpression_ArrayAccessor_StorageBuffer_Load_Int_FromArray_ExprIdx) {
+  // struct Data {
+  //   a : [[stride(4)]] array<i32, 5>;
+  // };
+  // var<storage> data : Data;
+  // data.a[(2 + 4) - 3];
+
+  SetupStorageBuffer({
+      Member("z", ty.f32()),
+      Member("a", ty.array<i32, 5>(4)),
+  });
+
+  SetupFunction({
+      Decl(Var("x", nullptr, ast::StorageClass::kNone,
+               IndexAccessor(MemberAccessor("data", "a"),
+                             Sub(Add(2, Expr(4)), Expr(3))))),
+  });
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  auto* expected =
+      R"(#version 310 es
+precision mediump float;
+
+
+Data data : register(u0, space1);
+
+void main() {
+  int x = data.a[((2 + 4) - 3)];
+  return;
+}
+void main() {
+  main();
+}
+
+
+)";
+  EXPECT_EQ(gen.result(), expected);
+}
+
+TEST_F(GlslGeneratorImplTest_MemberAccessor, StorageBuffer_Store_ToArray) {
+  // struct Data {
+  //   a : [[stride(4)]] array<i32, 5>;
+  // };
+  // var<storage> data : Data;
+  // data.a[2] = 2;
+
+  SetupStorageBuffer({
+      Member("z", ty.f32()),
+      Member("a", ty.array<i32, 5>(4)),
+  });
+
+  SetupFunction({
+      Assign(IndexAccessor(MemberAccessor("data", "a"), 2), 2),
+  });
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  auto* expected =
+      R"(#version 310 es
+precision mediump float;
+
+
+Data data : register(u0, space1);
+
+void main() {
+  data.a[2] = 2;
+  return;
+}
+void main() {
+  main();
+}
+
+
+)";
+  EXPECT_EQ(gen.result(), expected);
+}
+
+TEST_F(GlslGeneratorImplTest_MemberAccessor, StorageBuffer_Load_MultiLevel) {
+  // struct Inner {
+  //   a : vec3<i32>;
+  //   b : vec3<f32>;
+  // };
+  // struct Data {
+  //   var c : [[stride(32)]] array<Inner, 4>;
+  // };
+  //
+  // var<storage> data : Pre;
+  // data.c[2].b
+
+  auto* inner = Structure("Inner", {
+                                       Member("a", ty.vec3<f32>()),
+                                       Member("b", ty.vec3<f32>()),
+                                   });
+
+  SetupStorageBuffer({
+      Member("c", ty.array(ty.Of(inner), 4, 32)),
+  });
+
+  SetupFunction({
+      Decl(Var(
+          "x", nullptr, ast::StorageClass::kNone,
+          MemberAccessor(IndexAccessor(MemberAccessor("data", "c"), 2), "b"))),
+  });
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  auto* expected =
+      R"(#version 310 es
+precision mediump float;
+
+
+Data data : register(u0, space1);
+
+void main() {
+  vec3 x = data.c[2].b;
+  return;
+}
+void main() {
+  main();
+}
+
+
+)";
+  EXPECT_EQ(gen.result(), expected);
+}
+
+TEST_F(GlslGeneratorImplTest_MemberAccessor,
+       StorageBuffer_Load_MultiLevel_Swizzle) {
+  // struct Inner {
+  //   a : vec3<i32>;
+  //   b : vec3<f32>;
+  // };
+  // struct Data {
+  //   var c : [[stride(32)]] array<Inner, 4>;
+  // };
+  //
+  // var<storage> data : Pre;
+  // data.c[2].b.xy
+
+  auto* inner = Structure("Inner", {
+                                       Member("a", ty.vec3<f32>()),
+                                       Member("b", ty.vec3<f32>()),
+                                   });
+
+  SetupStorageBuffer({
+      Member("c", ty.array(ty.Of(inner), 4, 32)),
+  });
+
+  SetupFunction({
+      Decl(Var("x", nullptr, ast::StorageClass::kNone,
+               MemberAccessor(
+                   MemberAccessor(IndexAccessor(MemberAccessor("data", "c"), 2),
+                                  "b"),
+                   "xy"))),
+  });
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  auto* expected =
+      R"(#version 310 es
+precision mediump float;
+
+
+Data data : register(u0, space1);
+
+void main() {
+  vec2 x = data.c[2].b.xy;
+  return;
+}
+void main() {
+  main();
+}
+
+
+)";
+  EXPECT_EQ(gen.result(), expected);
+}
+
+TEST_F(GlslGeneratorImplTest_MemberAccessor,
+       StorageBuffer_Load_MultiLevel_Swizzle_SingleLetter) {  // NOLINT
+  // struct Inner {
+  //   a : vec3<i32>;
+  //   b : vec3<f32>;
+  // };
+  // struct Data {
+  //   var c : [[stride(32)]] array<Inner, 4>;
+  // };
+  //
+  // var<storage> data : Pre;
+  // data.c[2].b.g
+
+  auto* inner = Structure("Inner", {
+                                       Member("a", ty.vec3<f32>()),
+                                       Member("b", ty.vec3<f32>()),
+                                   });
+
+  SetupStorageBuffer({
+      Member("c", ty.array(ty.Of(inner), 4, 32)),
+  });
+
+  SetupFunction({
+      Decl(Var("x", nullptr, ast::StorageClass::kNone,
+               MemberAccessor(
+                   MemberAccessor(IndexAccessor(MemberAccessor("data", "c"), 2),
+                                  "b"),
+                   "g"))),
+  });
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  auto* expected =
+      R"(#version 310 es
+precision mediump float;
+
+
+Data data : register(u0, space1);
+
+void main() {
+  float x = data.c[2].b.g;
+  return;
+}
+void main() {
+  main();
+}
+
+
+)";
+  EXPECT_EQ(gen.result(), expected);
+}
+
+TEST_F(GlslGeneratorImplTest_MemberAccessor,
+       StorageBuffer_Load_MultiLevel_Index) {
+  // struct Inner {
+  //   a : vec3<i32>;
+  //   b : vec3<f32>;
+  // };
+  // struct Data {
+  //   var c : [[stride(32)]] array<Inner, 4>;
+  // };
+  //
+  // var<storage> data : Pre;
+  // data.c[2].b[1]
+
+  auto* inner = Structure("Inner", {
+                                       Member("a", ty.vec3<f32>()),
+                                       Member("b", ty.vec3<f32>()),
+                                   });
+
+  SetupStorageBuffer({
+      Member("c", ty.array(ty.Of(inner), 4, 32)),
+  });
+
+  SetupFunction({
+      Decl(Var(
+          "x", nullptr, ast::StorageClass::kNone,
+          IndexAccessor(MemberAccessor(
+                            IndexAccessor(MemberAccessor("data", "c"), 2), "b"),
+                        1))),
+  });
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  auto* expected =
+      R"(#version 310 es
+precision mediump float;
+
+
+Data data : register(u0, space1);
+
+void main() {
+  float x = data.c[2].b[1];
+  return;
+}
+void main() {
+  main();
+}
+
+
+)";
+  EXPECT_EQ(gen.result(), expected);
+}
+
+TEST_F(GlslGeneratorImplTest_MemberAccessor, StorageBuffer_Store_MultiLevel) {
+  // struct Inner {
+  //   a : vec3<i32>;
+  //   b : vec3<f32>;
+  // };
+  // struct Data {
+  //   var c : [[stride(32)]] array<Inner, 4>;
+  // };
+  //
+  // var<storage> data : Pre;
+  // data.c[2].b = vec3<f32>(1.f, 2.f, 3.f);
+
+  auto* inner = Structure("Inner", {
+                                       Member("a", ty.vec3<f32>()),
+                                       Member("b", ty.vec3<f32>()),
+                                   });
+
+  SetupStorageBuffer({
+      Member("c", ty.array(ty.Of(inner), 4, 32)),
+  });
+
+  SetupFunction({
+      Assign(MemberAccessor(IndexAccessor(MemberAccessor("data", "c"), 2), "b"),
+             vec3<f32>(1.f, 2.f, 3.f)),
+  });
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  auto* expected =
+      R"(#version 310 es
+precision mediump float;
+
+
+Data data : register(u0, space1);
+
+void main() {
+  data.c[2].b = vec3(1.0f, 2.0f, 3.0f);
+  return;
+}
+void main() {
+  main();
+}
+
+
+)";
+  EXPECT_EQ(gen.result(), expected);
+}
+
+TEST_F(GlslGeneratorImplTest_MemberAccessor,
+       StorageBuffer_Store_Swizzle_SingleLetter) {
+  // struct Inner {
+  //   a : vec3<i32>;
+  //   b : vec3<f32>;
+  // };
+  // struct Data {
+  //   var c : [[stride(32)]] array<Inner, 4>;
+  // };
+  //
+  // var<storage> data : Pre;
+  // data.c[2].b.y = 1.f;
+
+  auto* inner = Structure("Inner", {
+                                       Member("a", ty.vec3<i32>()),
+                                       Member("b", ty.vec3<f32>()),
+                                   });
+
+  SetupStorageBuffer({
+      Member("c", ty.array(ty.Of(inner), 4, 32)),
+  });
+
+  SetupFunction({
+      Assign(MemberAccessor(
+                 MemberAccessor(IndexAccessor(MemberAccessor("data", "c"), 2),
+                                "b"),
+                 "y"),
+             Expr(1.f)),
+  });
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  auto* expected =
+      R"(#version 310 es
+precision mediump float;
+
+
+Data data : register(u0, space1);
+
+void main() {
+  data.c[2].b.y = 1.0f;
+  return;
+}
+void main() {
+  main();
+}
+
+
+)";
+  EXPECT_EQ(gen.result(), expected);
+}
+
+TEST_F(GlslGeneratorImplTest_MemberAccessor, Swizzle_xyz) {
+  auto* var = Var("my_vec", ty.vec4<f32>(), ast::StorageClass::kNone,
+                  vec4<f32>(1.f, 2.f, 3.f, 4.f));
+  auto* expr = MemberAccessor("my_vec", "xyz");
+  WrapInFunction(var, expr);
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("my_vec.xyz"));
+}
+
+TEST_F(GlslGeneratorImplTest_MemberAccessor, Swizzle_gbr) {
+  auto* var = Var("my_vec", ty.vec4<f32>(), ast::StorageClass::kNone,
+                  vec4<f32>(1.f, 2.f, 3.f, 4.f));
+  auto* expr = MemberAccessor("my_vec", "gbr");
+  WrapInFunction(var, expr);
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("my_vec.gbr"));
+}
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_module_constant_test.cc

@@ -0,0 +1,96 @@
+// 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/ast/override_decoration.h"
+#include "src/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using GlslGeneratorImplTest_ModuleConstant = TestHelper;
+
+TEST_F(GlslGeneratorImplTest_ModuleConstant, Emit_ModuleConstant) {
+  auto* var = Const("pos", ty.array<f32, 3>(), array<f32, 3>(1.f, 2.f, 3.f));
+  WrapInFunction(Decl(var));
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.EmitProgramConstVariable(var)) << gen.error();
+  EXPECT_EQ(gen.result(),
+            "static const float pos[3] = float[3](1.0f, 2.0f, 3.0f);\n");
+}
+
+TEST_F(GlslGeneratorImplTest_ModuleConstant, Emit_SpecConstant) {
+  auto* var = GlobalConst("pos", ty.f32(), Expr(3.0f),
+                          ast::DecorationList{
+                              create<ast::OverrideDecoration>(23),
+                          });
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.EmitProgramConstVariable(var)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(#ifndef WGSL_SPEC_CONSTANT_23
+#define WGSL_SPEC_CONSTANT_23 3.0f
+#endif
+static const float pos = WGSL_SPEC_CONSTANT_23;
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_ModuleConstant, Emit_SpecConstant_NoConstructor) {
+  auto* var = GlobalConst("pos", ty.f32(), nullptr,
+                          ast::DecorationList{
+                              create<ast::OverrideDecoration>(23),
+                          });
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.EmitProgramConstVariable(var)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(#ifndef WGSL_SPEC_CONSTANT_23
+#error spec constant required for constant id 23
+#endif
+static const float pos = WGSL_SPEC_CONSTANT_23;
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_ModuleConstant, Emit_SpecConstant_NoId) {
+  auto* a = GlobalConst("a", ty.f32(), Expr(3.0f),
+                        ast::DecorationList{
+                            create<ast::OverrideDecoration>(0),
+                        });
+  auto* b = GlobalConst("b", ty.f32(), Expr(2.0f),
+                        ast::DecorationList{
+                            create<ast::OverrideDecoration>(),
+                        });
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.EmitProgramConstVariable(a)) << gen.error();
+  ASSERT_TRUE(gen.EmitProgramConstVariable(b)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(#ifndef WGSL_SPEC_CONSTANT_0
+#define WGSL_SPEC_CONSTANT_0 3.0f
+#endif
+static const float a = WGSL_SPEC_CONSTANT_0;
+#ifndef WGSL_SPEC_CONSTANT_1
+#define WGSL_SPEC_CONSTANT_1 2.0f
+#endif
+static const float b = WGSL_SPEC_CONSTANT_1;
+)");
+}
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_return_test.cc

@@ -0,0 +1,51 @@
+// 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/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using GlslGeneratorImplTest_Return = TestHelper;
+
+TEST_F(GlslGeneratorImplTest_Return, Emit_Return) {
+  auto* r = Return();
+  WrapInFunction(r);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(r)) << gen.error();
+  EXPECT_EQ(gen.result(), "  return;\n");
+}
+
+TEST_F(GlslGeneratorImplTest_Return, Emit_ReturnWithValue) {
+  auto* r = Return(123);
+  Func("f", {}, ty.i32(), {r});
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(r)) << gen.error();
+  EXPECT_EQ(gen.result(), "  return 123;\n");
+}
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_sanitizer_test.cc

@@ -0,0 +1,351 @@
+// 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/ast/call_statement.h"
+#include "src/ast/stage_decoration.h"
+#include "src/ast/struct_block_decoration.h"
+#include "src/ast/variable_decl_statement.h"
+#include "src/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using GlslSanitizerTest = TestHelper;
+
+TEST_F(GlslSanitizerTest, Call_ArrayLength) {
+  auto* s = Structure("my_struct", {Member(0, "a", ty.array<f32>(4))},
+                      {create<ast::StructBlockDecoration>()});
+  Global("b", ty.Of(s), ast::StorageClass::kStorage, ast::Access::kRead,
+         ast::DecorationList{
+             create<ast::BindingDecoration>(1),
+             create<ast::GroupDecoration>(2),
+         });
+
+  Func("a_func", ast::VariableList{}, ty.void_(),
+       ast::StatementList{
+           Decl(Var("len", ty.u32(), ast::StorageClass::kNone,
+                    Call("arrayLength", AddressOf(MemberAccessor("b", "a"))))),
+       },
+       ast::DecorationList{
+           Stage(ast::PipelineStage::kFragment),
+       });
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+
+  auto got = gen.result();
+  auto* expect = R"(#version 310 es
+precision mediump float;
+
+
+my_struct b : register(t1, space2);
+
+void a_func() {
+  uint tint_symbol_1 = 0u;
+  b.GetDimensions(tint_symbol_1);
+  uint tint_symbol_2 = ((tint_symbol_1 - 0u) / 4u);
+  uint len = tint_symbol_2;
+  return;
+}
+void main() {
+  a_func();
+}
+
+
+)";
+  EXPECT_EQ(expect, got);
+}
+
+TEST_F(GlslSanitizerTest, Call_ArrayLength_OtherMembersInStruct) {
+  auto* s = Structure("my_struct",
+                      {
+                          Member(0, "z", ty.f32()),
+                          Member(4, "a", ty.array<f32>(4)),
+                      },
+                      {create<ast::StructBlockDecoration>()});
+  Global("b", ty.Of(s), ast::StorageClass::kStorage, ast::Access::kRead,
+         ast::DecorationList{
+             create<ast::BindingDecoration>(1),
+             create<ast::GroupDecoration>(2),
+         });
+
+  Func("a_func", ast::VariableList{}, ty.void_(),
+       ast::StatementList{
+           Decl(Var("len", ty.u32(), ast::StorageClass::kNone,
+                    Call("arrayLength", AddressOf(MemberAccessor("b", "a"))))),
+       },
+       ast::DecorationList{
+           Stage(ast::PipelineStage::kFragment),
+       });
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+
+  auto got = gen.result();
+  auto* expect = R"(#version 310 es
+precision mediump float;
+
+
+my_struct b : register(t1, space2);
+
+void a_func() {
+  uint tint_symbol_1 = 0u;
+  b.GetDimensions(tint_symbol_1);
+  uint tint_symbol_2 = ((tint_symbol_1 - 4u) / 4u);
+  uint len = tint_symbol_2;
+  return;
+}
+void main() {
+  a_func();
+}
+
+
+)";
+
+  EXPECT_EQ(expect, got);
+}
+
+TEST_F(GlslSanitizerTest, Call_ArrayLength_ViaLets) {
+  auto* s = Structure("my_struct", {Member(0, "a", ty.array<f32>(4))},
+                      {create<ast::StructBlockDecoration>()});
+  Global("b", ty.Of(s), ast::StorageClass::kStorage, ast::Access::kRead,
+         ast::DecorationList{
+             create<ast::BindingDecoration>(1),
+             create<ast::GroupDecoration>(2),
+         });
+
+  auto* p = Const("p", nullptr, AddressOf("b"));
+  auto* p2 = Const("p2", nullptr, AddressOf(MemberAccessor(Deref(p), "a")));
+
+  Func("a_func", ast::VariableList{}, ty.void_(),
+       ast::StatementList{
+           Decl(p),
+           Decl(p2),
+           Decl(Var("len", ty.u32(), ast::StorageClass::kNone,
+                    Call("arrayLength", p2))),
+       },
+       ast::DecorationList{
+           Stage(ast::PipelineStage::kFragment),
+       });
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+
+  auto got = gen.result();
+  auto* expect = R"(#version 310 es
+precision mediump float;
+
+
+my_struct b : register(t1, space2);
+
+void a_func() {
+  uint tint_symbol_1 = 0u;
+  b.GetDimensions(tint_symbol_1);
+  uint tint_symbol_2 = ((tint_symbol_1 - 0u) / 4u);
+  uint len = tint_symbol_2;
+  return;
+}
+void main() {
+  a_func();
+}
+
+
+)";
+
+  EXPECT_EQ(expect, got);
+}
+
+TEST_F(GlslSanitizerTest, PromoteArrayInitializerToConstVar) {
+  auto* array_init = array<i32, 4>(1, 2, 3, 4);
+  auto* array_index = IndexAccessor(array_init, 3);
+  auto* pos = Var("pos", ty.i32(), ast::StorageClass::kNone, array_index);
+
+  Func("main", ast::VariableList{}, ty.void_(),
+       {
+           Decl(pos),
+       },
+       {
+           Stage(ast::PipelineStage::kFragment),
+       });
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+
+  auto got = gen.result();
+  auto* expect = R"(#version 310 es
+precision mediump float;
+
+void main() {
+  int tint_symbol[4] = int[4](1, 2, 3, 4);
+  int pos = tint_symbol[3];
+  return;
+}
+void main() {
+  main();
+}
+
+
+)";
+  EXPECT_EQ(expect, got);
+}
+
+TEST_F(GlslSanitizerTest, PromoteStructInitializerToConstVar) {
+  auto* str = Structure("S", {
+                                 Member("a", ty.i32()),
+                                 Member("b", ty.vec3<f32>()),
+                                 Member("c", ty.i32()),
+                             });
+  auto* struct_init = Construct(ty.Of(str), 1, vec3<f32>(2.f, 3.f, 4.f), 4);
+  auto* struct_access = MemberAccessor(struct_init, "b");
+  auto* pos =
+      Var("pos", ty.vec3<f32>(), ast::StorageClass::kNone, struct_access);
+
+  Func("main", ast::VariableList{}, ty.void_(),
+       {
+           Decl(pos),
+       },
+       {
+           Stage(ast::PipelineStage::kFragment),
+       });
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+
+  auto got = gen.result();
+  auto* expect = R"(#version 310 es
+precision mediump float;
+
+struct S {
+  int a;
+  vec3 b;
+  int c;
+};
+
+void main() {
+  S tint_symbol = S(1, vec3(2.0f, 3.0f, 4.0f), 4);
+  vec3 pos = tint_symbol.b;
+  return;
+}
+void main() {
+  main();
+}
+
+
+)";
+  EXPECT_EQ(expect, got);
+}
+
+TEST_F(GlslSanitizerTest, InlinePtrLetsBasic) {
+  // var v : i32;
+  // let p : ptr<function, i32> = &v;
+  // let x : i32 = *p;
+  auto* v = Var("v", ty.i32());
+  auto* p =
+      Const("p", ty.pointer<i32>(ast::StorageClass::kFunction), AddressOf(v));
+  auto* x = Var("x", ty.i32(), ast::StorageClass::kNone, Deref(p));
+
+  Func("main", ast::VariableList{}, ty.void_(),
+       {
+           Decl(v),
+           Decl(p),
+           Decl(x),
+       },
+       {
+           Stage(ast::PipelineStage::kFragment),
+       });
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+
+  auto got = gen.result();
+  auto* expect = R"(#version 310 es
+precision mediump float;
+
+void main() {
+  int v = 0;
+  int x = v;
+  return;
+}
+void main() {
+  main();
+}
+
+
+)";
+  EXPECT_EQ(expect, got);
+}
+
+TEST_F(GlslSanitizerTest, InlinePtrLetsComplexChain) {
+  // var m : mat4x4<f32>;
+  // let mp : ptr<function, mat4x4<f32>> = &m;
+  // let vp : ptr<function, vec4<f32>> = &(*mp)[2];
+  // let fp : ptr<function, f32> = &(*vp)[1];
+  // let f : f32 = *fp;
+  auto* m = Var("m", ty.mat4x4<f32>());
+  auto* mp =
+      Const("mp", ty.pointer(ty.mat4x4<f32>(), ast::StorageClass::kFunction),
+            AddressOf(m));
+  auto* vp =
+      Const("vp", ty.pointer(ty.vec4<f32>(), ast::StorageClass::kFunction),
+            AddressOf(IndexAccessor(Deref(mp), 2)));
+  auto* fp = Const("fp", ty.pointer<f32>(ast::StorageClass::kFunction),
+                   AddressOf(IndexAccessor(Deref(vp), 1)));
+  auto* f = Var("f", ty.f32(), ast::StorageClass::kNone, Deref(fp));
+
+  Func("main", ast::VariableList{}, ty.void_(),
+       {
+           Decl(m),
+           Decl(mp),
+           Decl(vp),
+           Decl(fp),
+           Decl(f),
+       },
+       {
+           Stage(ast::PipelineStage::kFragment),
+       });
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+
+  auto got = gen.result();
+  auto* expect = R"(#version 310 es
+precision mediump float;
+
+void main() {
+  mat4 m = mat4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+  float f = m[2][1];
+  return;
+}
+void main() {
+  main();
+}
+
+
+)";
+  EXPECT_EQ(expect, got);
+}
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_switch_test.cc

@@ -0,0 +1,64 @@
+// 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/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using GlslGeneratorImplTest_Switch = TestHelper;
+
+TEST_F(GlslGeneratorImplTest_Switch, Emit_Switch) {
+  Global("cond", ty.i32(), ast::StorageClass::kPrivate);
+
+  auto* def_body = Block(create<ast::BreakStatement>());
+  auto* def = create<ast::CaseStatement>(ast::CaseSelectorList{}, def_body);
+
+  ast::CaseSelectorList case_val;
+  case_val.push_back(Literal(5));
+
+  auto* case_body = Block(create<ast::BreakStatement>());
+
+  auto* case_stmt = create<ast::CaseStatement>(case_val, case_body);
+
+  ast::CaseStatementList body;
+  body.push_back(case_stmt);
+  body.push_back(def);
+
+  auto* cond = Expr("cond");
+  auto* s = create<ast::SwitchStatement>(cond, body);
+  WrapInFunction(s);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(s)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  switch(cond) {
+    case 5: {
+      break;
+    }
+    default: {
+      break;
+    }
+  }
+)");
+}
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_test.cc

@@ -0,0 +1,87 @@
+// 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/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using GlslGeneratorImplTest = TestHelper;
+
+TEST_F(GlslGeneratorImplTest, ErrorIfSanitizerNotRun) {
+  auto program = std::make_unique<Program>(std::move(*this));
+  GeneratorImpl gen(program.get());
+  EXPECT_FALSE(gen.Generate());
+  EXPECT_EQ(
+      gen.error(),
+      "error: GLSL writer requires the transform::Glsl sanitizer to have been "
+      "applied to the input program");
+}
+
+TEST_F(GlslGeneratorImplTest, Generate) {
+  Func("my_func", ast::VariableList{}, ty.void_(), ast::StatementList{},
+       ast::DecorationList{});
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_EQ(gen.result(), R"(#version 310 es
+precision mediump float;
+
+void my_func() {
+}
+)");
+}
+
+struct GlslBuiltinData {
+  ast::Builtin builtin;
+  const char* attribute_name;
+};
+inline std::ostream& operator<<(std::ostream& out, GlslBuiltinData data) {
+  out << data.builtin;
+  return out;
+}
+using GlslBuiltinConversionTest = TestParamHelper<GlslBuiltinData>;
+TEST_P(GlslBuiltinConversionTest, Emit) {
+  auto params = GetParam();
+  GeneratorImpl& gen = Build();
+
+  EXPECT_EQ(gen.builtin_to_string(params.builtin),
+            std::string(params.attribute_name));
+}
+INSTANTIATE_TEST_SUITE_P(
+    GlslGeneratorImplTest,
+    GlslBuiltinConversionTest,
+    testing::Values(
+        GlslBuiltinData{ast::Builtin::kPosition, "gl_Position"},
+        GlslBuiltinData{ast::Builtin::kVertexIndex, "gl_VertexID"},
+        GlslBuiltinData{ast::Builtin::kInstanceIndex, "gl_InstanceID"},
+        GlslBuiltinData{ast::Builtin::kFrontFacing, "gl_FrontFacing"},
+        GlslBuiltinData{ast::Builtin::kFragDepth, "gl_FragDepth"},
+        GlslBuiltinData{ast::Builtin::kLocalInvocationId,
+                        "gl_LocalInvocationID"},
+        GlslBuiltinData{ast::Builtin::kLocalInvocationIndex,
+                        "gl_LocalInvocationIndex"},
+        GlslBuiltinData{ast::Builtin::kGlobalInvocationId,
+                        "gl_GlobalInvocationID"},
+        GlslBuiltinData{ast::Builtin::kWorkgroupId, "gl_WorkGroupID"},
+        GlslBuiltinData{ast::Builtin::kSampleIndex, "gl_SampleID"},
+        GlslBuiltinData{ast::Builtin::kSampleMask, "gl_SampleMask"}));
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_type_test.cc

@@ -0,0 +1,615 @@
+// 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 "gmock/gmock.h"
+#include "src/ast/call_statement.h"
+#include "src/ast/stage_decoration.h"
+#include "src/ast/struct_block_decoration.h"
+#include "src/sem/depth_texture_type.h"
+#include "src/sem/multisampled_texture_type.h"
+#include "src/sem/sampled_texture_type.h"
+#include "src/sem/sampler_type.h"
+#include "src/sem/storage_texture_type.h"
+#include "src/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using ::testing::HasSubstr;
+
+using GlslGeneratorImplTest_Type = TestHelper;
+
+TEST_F(GlslGeneratorImplTest_Type, EmitType_Array) {
+  auto* arr = ty.array<bool, 4>();
+  Global("G", arr, ast::StorageClass::kPrivate);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, program->TypeOf(arr), ast::StorageClass::kNone,
+                           ast::Access::kReadWrite, "ary"))
+      << gen.error();
+  EXPECT_EQ(out.str(), "bool ary[4]");
+}
+
+TEST_F(GlslGeneratorImplTest_Type, EmitType_ArrayOfArray) {
+  auto* arr = ty.array(ty.array<bool, 4>(), 5);
+  Global("G", arr, ast::StorageClass::kPrivate);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, program->TypeOf(arr), ast::StorageClass::kNone,
+                           ast::Access::kReadWrite, "ary"))
+      << gen.error();
+  EXPECT_EQ(out.str(), "bool ary[5][4]");
+}
+
+// TODO(dsinclair): Is this possible? What order should it output in?
+TEST_F(GlslGeneratorImplTest_Type,
+       DISABLED_EmitType_ArrayOfArrayOfRuntimeArray) {
+  auto* arr = ty.array(ty.array(ty.array<bool, 4>(), 5), 0);
+  Global("G", arr, ast::StorageClass::kPrivate);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, program->TypeOf(arr), ast::StorageClass::kNone,
+                           ast::Access::kReadWrite, "ary"))
+      << gen.error();
+  EXPECT_EQ(out.str(), "bool ary[5][4][1]");
+}
+
+TEST_F(GlslGeneratorImplTest_Type, EmitType_ArrayOfArrayOfArray) {
+  auto* arr = ty.array(ty.array(ty.array<bool, 4>(), 5), 6);
+  Global("G", arr, ast::StorageClass::kPrivate);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, program->TypeOf(arr), ast::StorageClass::kNone,
+                           ast::Access::kReadWrite, "ary"))
+      << gen.error();
+  EXPECT_EQ(out.str(), "bool ary[6][5][4]");
+}
+
+TEST_F(GlslGeneratorImplTest_Type, EmitType_Array_WithoutName) {
+  auto* arr = ty.array<bool, 4>();
+  Global("G", arr, ast::StorageClass::kPrivate);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, program->TypeOf(arr), ast::StorageClass::kNone,
+                           ast::Access::kReadWrite, ""))
+      << gen.error();
+  EXPECT_EQ(out.str(), "bool[4]");
+}
+
+TEST_F(GlslGeneratorImplTest_Type, EmitType_Bool) {
+  auto* bool_ = create<sem::Bool>();
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, bool_, ast::StorageClass::kNone,
+                           ast::Access::kReadWrite, ""))
+      << gen.error();
+  EXPECT_EQ(out.str(), "bool");
+}
+
+TEST_F(GlslGeneratorImplTest_Type, EmitType_F32) {
+  auto* f32 = create<sem::F32>();
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, f32, ast::StorageClass::kNone,
+                           ast::Access::kReadWrite, ""))
+      << gen.error();
+  EXPECT_EQ(out.str(), "float");
+}
+
+TEST_F(GlslGeneratorImplTest_Type, EmitType_I32) {
+  auto* i32 = create<sem::I32>();
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, i32, ast::StorageClass::kNone,
+                           ast::Access::kReadWrite, ""))
+      << gen.error();
+  EXPECT_EQ(out.str(), "int");
+}
+
+TEST_F(GlslGeneratorImplTest_Type, EmitType_Matrix) {
+  auto* f32 = create<sem::F32>();
+  auto* vec3 = create<sem::Vector>(f32, 3);
+  auto* mat2x3 = create<sem::Matrix>(vec3, 2);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, mat2x3, ast::StorageClass::kNone,
+                           ast::Access::kReadWrite, ""))
+      << gen.error();
+  EXPECT_EQ(out.str(), "mat2x3");
+}
+
+// TODO(dsinclair): How to annotate as workgroup?
+TEST_F(GlslGeneratorImplTest_Type, DISABLED_EmitType_Pointer) {
+  auto* f32 = create<sem::F32>();
+  auto* p = create<sem::Pointer>(f32, ast::StorageClass::kWorkgroup,
+                                 ast::Access::kReadWrite);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, p, ast::StorageClass::kNone,
+                           ast::Access::kReadWrite, ""))
+      << gen.error();
+  EXPECT_EQ(out.str(), "float*");
+}
+
+TEST_F(GlslGeneratorImplTest_Type, EmitType_StructDecl) {
+  auto* s = Structure("S", {
+                               Member("a", ty.i32()),
+                               Member("b", ty.f32()),
+                           });
+  Global("g", ty.Of(s), ast::StorageClass::kPrivate);
+
+  GeneratorImpl& gen = Build();
+
+  TextGenerator::TextBuffer buf;
+  auto* sem_s = program->TypeOf(s)->As<sem::Struct>();
+  ASSERT_TRUE(gen.EmitStructType(&buf, sem_s)) << gen.error();
+  EXPECT_EQ(buf.String(), R"(struct S {
+  int a;
+  float b;
+};
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Type, EmitType_StructDecl_OmittedIfStorageBuffer) {
+  auto* s = Structure("S",
+                      {
+                          Member("a", ty.i32()),
+                          Member("b", ty.f32()),
+                      },
+                      {create<ast::StructBlockDecoration>()});
+  Global("g", ty.Of(s), ast::StorageClass::kStorage, ast::Access::kReadWrite,
+         ast::DecorationList{
+             create<ast::BindingDecoration>(0),
+             create<ast::GroupDecoration>(0),
+         });
+
+  GeneratorImpl& gen = Build();
+
+  TextGenerator::TextBuffer buf;
+  auto* sem_s = program->TypeOf(s)->As<sem::Struct>();
+  ASSERT_TRUE(gen.EmitStructType(&buf, sem_s)) << gen.error();
+  EXPECT_EQ(buf.String(), "");
+}
+
+TEST_F(GlslGeneratorImplTest_Type, EmitType_Struct) {
+  auto* s = Structure("S", {
+                               Member("a", ty.i32()),
+                               Member("b", ty.f32()),
+                           });
+  Global("g", ty.Of(s), ast::StorageClass::kPrivate);
+
+  GeneratorImpl& gen = Build();
+
+  auto* sem_s = program->TypeOf(s)->As<sem::Struct>();
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, sem_s, ast::StorageClass::kNone,
+                           ast::Access::kReadWrite, ""))
+      << gen.error();
+  EXPECT_EQ(out.str(), "S");
+}
+
+TEST_F(GlslGeneratorImplTest_Type, EmitType_Struct_NameCollision) {
+  auto* s = Structure("S", {
+                               Member("double", ty.i32()),
+                               Member("float", ty.f32()),
+                           });
+  Global("g", ty.Of(s), ast::StorageClass::kPrivate);
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr(R"(struct S {
+  int tint_symbol;
+  float tint_symbol_1;
+};
+)"));
+}
+
+TEST_F(GlslGeneratorImplTest_Type, EmitType_Struct_WithOffsetAttributes) {
+  auto* s = Structure("S",
+                      {
+                          Member("a", ty.i32(), {MemberOffset(0)}),
+                          Member("b", ty.f32(), {MemberOffset(8)}),
+                      },
+                      {create<ast::StructBlockDecoration>()});
+  Global("g", ty.Of(s), ast::StorageClass::kPrivate);
+
+  GeneratorImpl& gen = Build();
+
+  TextGenerator::TextBuffer buf;
+  auto* sem_s = program->TypeOf(s)->As<sem::Struct>();
+  ASSERT_TRUE(gen.EmitStructType(&buf, sem_s)) << gen.error();
+  EXPECT_EQ(buf.String(), R"(struct S {
+  int a;
+  float b;
+};
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Type, EmitType_U32) {
+  auto* u32 = create<sem::U32>();
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, u32, ast::StorageClass::kNone,
+                           ast::Access::kReadWrite, ""))
+      << gen.error();
+  EXPECT_EQ(out.str(), "uint");
+}
+
+TEST_F(GlslGeneratorImplTest_Type, EmitType_Vector) {
+  auto* f32 = create<sem::F32>();
+  auto* vec3 = create<sem::Vector>(f32, 3);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, vec3, ast::StorageClass::kNone,
+                           ast::Access::kReadWrite, ""))
+      << gen.error();
+  EXPECT_EQ(out.str(), "vec3");
+}
+
+TEST_F(GlslGeneratorImplTest_Type, EmitType_Void) {
+  auto* void_ = create<sem::Void>();
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, void_, ast::StorageClass::kNone,
+                           ast::Access::kReadWrite, ""))
+      << gen.error();
+  EXPECT_EQ(out.str(), "void");
+}
+
+TEST_F(GlslGeneratorImplTest_Type, EmitSampler) {
+  auto* sampler = create<sem::Sampler>(ast::SamplerKind::kSampler);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, sampler, ast::StorageClass::kNone,
+                           ast::Access::kReadWrite, ""))
+      << gen.error();
+  EXPECT_EQ(out.str(), "SamplerState");
+}
+
+TEST_F(GlslGeneratorImplTest_Type, EmitSamplerComparison) {
+  auto* sampler = create<sem::Sampler>(ast::SamplerKind::kComparisonSampler);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, sampler, ast::StorageClass::kNone,
+                           ast::Access::kReadWrite, ""))
+      << gen.error();
+  EXPECT_EQ(out.str(), "SamplerComparisonState");
+}
+
+struct GlslDepthTextureData {
+  ast::TextureDimension dim;
+  std::string result;
+};
+inline std::ostream& operator<<(std::ostream& out, GlslDepthTextureData data) {
+  out << data.dim;
+  return out;
+}
+using GlslDepthTexturesTest = TestParamHelper<GlslDepthTextureData>;
+TEST_P(GlslDepthTexturesTest, Emit) {
+  auto params = GetParam();
+
+  auto* t = ty.depth_texture(params.dim);
+
+  Global("tex", t,
+         ast::DecorationList{
+             create<ast::BindingDecoration>(1),
+             create<ast::GroupDecoration>(2),
+         });
+
+  Func("main", {}, ty.void_(), {Ignore(Call("textureDimensions", "tex"))},
+       {Stage(ast::PipelineStage::kFragment)});
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr(params.result));
+}
+INSTANTIATE_TEST_SUITE_P(
+    GlslGeneratorImplTest_Type,
+    GlslDepthTexturesTest,
+    testing::Values(
+        GlslDepthTextureData{ast::TextureDimension::k2d,
+                             "Texture2D tex : register(t1, space2);"},
+        GlslDepthTextureData{ast::TextureDimension::k2dArray,
+                             "Texture2DArray tex : register(t1, space2);"},
+        GlslDepthTextureData{ast::TextureDimension::kCube,
+                             "TextureCube tex : register(t1, space2);"},
+        GlslDepthTextureData{ast::TextureDimension::kCubeArray,
+                             "TextureCubeArray tex : register(t1, space2);"}));
+
+using GlslDepthMultisampledTexturesTest = TestHelper;
+TEST_F(GlslDepthMultisampledTexturesTest, Emit) {
+  auto* t = ty.depth_multisampled_texture(ast::TextureDimension::k2d);
+
+  Global("tex", t,
+         ast::DecorationList{
+             create<ast::BindingDecoration>(1),
+             create<ast::GroupDecoration>(2),
+         });
+
+  Func("main", {}, ty.void_(), {Ignore(Call("textureDimensions", "tex"))},
+       {Stage(ast::PipelineStage::kFragment)});
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(),
+              HasSubstr("Texture2DMS<float4> tex : register(t1, space2);"));
+}
+
+enum class TextureDataType { F32, U32, I32 };
+struct GlslSampledTextureData {
+  ast::TextureDimension dim;
+  TextureDataType datatype;
+  std::string result;
+};
+inline std::ostream& operator<<(std::ostream& out,
+                                GlslSampledTextureData data) {
+  out << data.dim;
+  return out;
+}
+using GlslSampledTexturesTest = TestParamHelper<GlslSampledTextureData>;
+TEST_P(GlslSampledTexturesTest, Emit) {
+  auto params = GetParam();
+
+  ast::Type* datatype = nullptr;
+  switch (params.datatype) {
+    case TextureDataType::F32:
+      datatype = ty.f32();
+      break;
+    case TextureDataType::U32:
+      datatype = ty.u32();
+      break;
+    case TextureDataType::I32:
+      datatype = ty.i32();
+      break;
+  }
+  auto* t = ty.sampled_texture(params.dim, datatype);
+
+  Global("tex", t,
+         ast::DecorationList{
+             create<ast::BindingDecoration>(1),
+             create<ast::GroupDecoration>(2),
+         });
+
+  Func("main", {}, ty.void_(), {Ignore(Call("textureDimensions", "tex"))},
+       {Stage(ast::PipelineStage::kFragment)});
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr(params.result));
+}
+INSTANTIATE_TEST_SUITE_P(
+    GlslGeneratorImplTest_Type,
+    GlslSampledTexturesTest,
+    testing::Values(
+        GlslSampledTextureData{
+            ast::TextureDimension::k1d,
+            TextureDataType::F32,
+            "Texture1D<float4> tex : register(t1, space2);",
+        },
+        GlslSampledTextureData{
+            ast::TextureDimension::k2d,
+            TextureDataType::F32,
+            "Texture2D<float4> tex : register(t1, space2);",
+        },
+        GlslSampledTextureData{
+            ast::TextureDimension::k2dArray,
+            TextureDataType::F32,
+            "Texture2DArray<float4> tex : register(t1, space2);",
+        },
+        GlslSampledTextureData{
+            ast::TextureDimension::k3d,
+            TextureDataType::F32,
+            "Texture3D<float4> tex : register(t1, space2);",
+        },
+        GlslSampledTextureData{
+            ast::TextureDimension::kCube,
+            TextureDataType::F32,
+            "TextureCube<float4> tex : register(t1, space2);",
+        },
+        GlslSampledTextureData{
+            ast::TextureDimension::kCubeArray,
+            TextureDataType::F32,
+            "TextureCubeArray<float4> tex : register(t1, space2);",
+        },
+        GlslSampledTextureData{
+            ast::TextureDimension::k1d,
+            TextureDataType::U32,
+            "Texture1D<uint4> tex : register(t1, space2);",
+        },
+        GlslSampledTextureData{
+            ast::TextureDimension::k2d,
+            TextureDataType::U32,
+            "Texture2D<uint4> tex : register(t1, space2);",
+        },
+        GlslSampledTextureData{
+            ast::TextureDimension::k2dArray,
+            TextureDataType::U32,
+            "Texture2DArray<uint4> tex : register(t1, space2);",
+        },
+        GlslSampledTextureData{
+            ast::TextureDimension::k3d,
+            TextureDataType::U32,
+            "Texture3D<uint4> tex : register(t1, space2);",
+        },
+        GlslSampledTextureData{
+            ast::TextureDimension::kCube,
+            TextureDataType::U32,
+            "TextureCube<uint4> tex : register(t1, space2);",
+        },
+        GlslSampledTextureData{
+            ast::TextureDimension::kCubeArray,
+            TextureDataType::U32,
+            "TextureCubeArray<uint4> tex : register(t1, space2);",
+        },
+        GlslSampledTextureData{
+            ast::TextureDimension::k1d,
+            TextureDataType::I32,
+            "Texture1D<int4> tex : register(t1, space2);",
+        },
+        GlslSampledTextureData{
+            ast::TextureDimension::k2d,
+            TextureDataType::I32,
+            "Texture2D<int4> tex : register(t1, space2);",
+        },
+        GlslSampledTextureData{
+            ast::TextureDimension::k2dArray,
+            TextureDataType::I32,
+            "Texture2DArray<int4> tex : register(t1, space2);",
+        },
+        GlslSampledTextureData{
+            ast::TextureDimension::k3d,
+            TextureDataType::I32,
+            "Texture3D<int4> tex : register(t1, space2);",
+        },
+        GlslSampledTextureData{
+            ast::TextureDimension::kCube,
+            TextureDataType::I32,
+            "TextureCube<int4> tex : register(t1, space2);",
+        },
+        GlslSampledTextureData{
+            ast::TextureDimension::kCubeArray,
+            TextureDataType::I32,
+            "TextureCubeArray<int4> tex : register(t1, space2);",
+        }));
+
+TEST_F(GlslGeneratorImplTest_Type, EmitMultisampledTexture) {
+  auto* f32 = create<sem::F32>();
+  auto* s = create<sem::MultisampledTexture>(ast::TextureDimension::k2d, f32);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitType(out, s, ast::StorageClass::kNone,
+                           ast::Access::kReadWrite, ""))
+      << gen.error();
+  EXPECT_EQ(out.str(), "Texture2DMS<float4>");
+}
+
+struct GlslStorageTextureData {
+  ast::TextureDimension dim;
+  ast::ImageFormat imgfmt;
+  bool ro;
+  std::string result;
+};
+inline std::ostream& operator<<(std::ostream& out,
+                                GlslStorageTextureData data) {
+  out << data.dim << (data.ro ? "ReadOnly" : "WriteOnly");
+  return out;
+}
+using GlslStorageTexturesTest = TestParamHelper<GlslStorageTextureData>;
+TEST_P(GlslStorageTexturesTest, Emit) {
+  auto params = GetParam();
+
+  auto* t =
+      ty.storage_texture(params.dim, params.imgfmt,
+                         params.ro ? ast::Access::kRead : ast::Access::kWrite);
+
+  Global("tex", t,
+         ast::DecorationList{
+             create<ast::BindingDecoration>(1),
+             create<ast::GroupDecoration>(2),
+         });
+
+  Func("main", {}, ty.void_(), {Ignore(Call("textureDimensions", "tex"))},
+       {Stage(ast::PipelineStage::kFragment)});
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr(params.result));
+}
+INSTANTIATE_TEST_SUITE_P(
+    GlslGeneratorImplTest_Type,
+    GlslStorageTexturesTest,
+    testing::Values(
+        GlslStorageTextureData{ast::TextureDimension::k1d,
+                               ast::ImageFormat::kRgba8Unorm, true,
+                               "Texture1D<float4> tex : register(t1, space2);"},
+        GlslStorageTextureData{ast::TextureDimension::k2d,
+                               ast::ImageFormat::kRgba16Float, true,
+                               "Texture2D<float4> tex : register(t1, space2);"},
+        GlslStorageTextureData{
+            ast::TextureDimension::k2dArray, ast::ImageFormat::kR32Float, true,
+            "Texture2DArray<float4> tex : register(t1, space2);"},
+        GlslStorageTextureData{ast::TextureDimension::k3d,
+                               ast::ImageFormat::kRg32Float, true,
+                               "Texture3D<float4> tex : register(t1, space2);"},
+        GlslStorageTextureData{
+            ast::TextureDimension::k1d, ast::ImageFormat::kRgba32Float, false,
+            "RWTexture1D<float4> tex : register(u1, space2);"},
+        GlslStorageTextureData{
+            ast::TextureDimension::k2d, ast::ImageFormat::kRgba16Uint, false,
+            "RWTexture2D<uint4> tex : register(u1, space2);"},
+        GlslStorageTextureData{
+            ast::TextureDimension::k2dArray, ast::ImageFormat::kR32Uint, false,
+            "RWTexture2DArray<uint4> tex : register(u1, space2);"},
+        GlslStorageTextureData{
+            ast::TextureDimension::k3d, ast::ImageFormat::kRg32Uint, false,
+            "RWTexture3D<uint4> tex : register(u1, space2);"},
+        GlslStorageTextureData{ast::TextureDimension::k1d,
+                               ast::ImageFormat::kRgba32Uint, true,
+                               "Texture1D<uint4> tex : register(t1, space2);"},
+        GlslStorageTextureData{ast::TextureDimension::k2d,
+                               ast::ImageFormat::kRgba16Sint, true,
+                               "Texture2D<int4> tex : register(t1, space2);"},
+        GlslStorageTextureData{
+            ast::TextureDimension::k2dArray, ast::ImageFormat::kR32Sint, true,
+            "Texture2DArray<int4> tex : register(t1, space2);"},
+        GlslStorageTextureData{ast::TextureDimension::k3d,
+                               ast::ImageFormat::kRg32Sint, true,
+                               "Texture3D<int4> tex : register(t1, space2);"},
+        GlslStorageTextureData{
+            ast::TextureDimension::k1d, ast::ImageFormat::kRgba32Sint, false,
+            "RWTexture1D<int4> tex : register(u1, space2);"}));
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_unary_op_test.cc

@@ -0,0 +1,93 @@
+// 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/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using GlslUnaryOpTest = TestHelper;
+
+TEST_F(GlslUnaryOpTest, AddressOf) {
+  Global("expr", ty.f32(), ast::StorageClass::kPrivate);
+  auto* op =
+      create<ast::UnaryOpExpression>(ast::UnaryOp::kAddressOf, Expr("expr"));
+  WrapInFunction(op);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, op)) << gen.error();
+  EXPECT_EQ(out.str(), "expr");
+}
+
+TEST_F(GlslUnaryOpTest, Complement) {
+  Global("expr", ty.u32(), ast::StorageClass::kPrivate);
+  auto* op =
+      create<ast::UnaryOpExpression>(ast::UnaryOp::kComplement, Expr("expr"));
+  WrapInFunction(op);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, op)) << gen.error();
+  EXPECT_EQ(out.str(), "~(expr)");
+}
+
+TEST_F(GlslUnaryOpTest, Indirection) {
+  Global("G", ty.f32(), ast::StorageClass::kPrivate);
+  auto* p = Const(
+      "expr", nullptr,
+      create<ast::UnaryOpExpression>(ast::UnaryOp::kAddressOf, Expr("G")));
+  auto* op =
+      create<ast::UnaryOpExpression>(ast::UnaryOp::kIndirection, Expr("expr"));
+  WrapInFunction(p, op);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, op)) << gen.error();
+  EXPECT_EQ(out.str(), "expr");
+}
+
+TEST_F(GlslUnaryOpTest, Not) {
+  Global("expr", ty.bool_(), ast::StorageClass::kPrivate);
+  auto* op = create<ast::UnaryOpExpression>(ast::UnaryOp::kNot, Expr("expr"));
+  WrapInFunction(op);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, op)) << gen.error();
+  EXPECT_EQ(out.str(), "!(expr)");
+}
+
+TEST_F(GlslUnaryOpTest, Negation) {
+  Global("expr", ty.i32(), ast::StorageClass::kPrivate);
+  auto* op =
+      create<ast::UnaryOpExpression>(ast::UnaryOp::kNegation, Expr("expr"));
+  WrapInFunction(op);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, op)) << gen.error();
+  EXPECT_EQ(out.str(), "-(expr)");
+}
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_variable_decl_statement_test.cc

@@ -0,0 +1,129 @@
+// 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 "gmock/gmock.h"
+#include "src/ast/variable_decl_statement.h"
+#include "src/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+
+using ::testing::HasSubstr;
+
+using GlslGeneratorImplTest_VariableDecl = TestHelper;
+
+TEST_F(GlslGeneratorImplTest_VariableDecl, Emit_VariableDeclStatement) {
+  auto* var = Var("a", ty.f32());
+  auto* stmt = Decl(var);
+  WrapInFunction(stmt);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(stmt)) << gen.error();
+  EXPECT_EQ(gen.result(), "  float a = 0.0f;\n");
+}
+
+TEST_F(GlslGeneratorImplTest_VariableDecl, Emit_VariableDeclStatement_Const) {
+  auto* var = Const("a", ty.f32(), Construct(ty.f32()));
+  auto* stmt = Decl(var);
+  WrapInFunction(stmt);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(stmt)) << gen.error();
+  EXPECT_EQ(gen.result(), "  float a = 0.0f;\n");
+}
+
+TEST_F(GlslGeneratorImplTest_VariableDecl, Emit_VariableDeclStatement_Array) {
+  auto* var = Var("a", ty.array<f32, 5>());
+
+  WrapInFunction(var, Expr("a"));
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(
+      gen.result(),
+      HasSubstr("  float a[5] = float[5](0.0f, 0.0f, 0.0f, 0.0f, 0.0f);\n"));
+}
+
+TEST_F(GlslGeneratorImplTest_VariableDecl, Emit_VariableDeclStatement_Private) {
+  Global("a", ty.f32(), ast::StorageClass::kPrivate);
+
+  WrapInFunction(Expr("a"));
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("  static float a = 0.0f;\n"));
+}
+
+TEST_F(GlslGeneratorImplTest_VariableDecl,
+       Emit_VariableDeclStatement_Initializer_Private) {
+  Global("initializer", ty.f32(), ast::StorageClass::kPrivate);
+  Global("a", ty.f32(), ast::StorageClass::kPrivate, Expr("initializer"));
+
+  WrapInFunction(Expr("a"));
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr(R"(float a = initializer;
+)"));
+}
+
+TEST_F(GlslGeneratorImplTest_VariableDecl,
+       Emit_VariableDeclStatement_Initializer_ZeroVec) {
+  auto* var = Var("a", ty.vec3<f32>(), ast::StorageClass::kNone, vec3<f32>());
+
+  auto* stmt = Decl(var);
+  WrapInFunction(stmt);
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.EmitStatement(stmt)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(vec3 a = vec3(0.0f, 0.0f, 0.0f);
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_VariableDecl,
+       Emit_VariableDeclStatement_Initializer_ZeroMat) {
+  auto* var =
+      Var("a", ty.mat2x3<f32>(), ast::StorageClass::kNone, mat2x3<f32>());
+
+  auto* stmt = Decl(var);
+  WrapInFunction(stmt);
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.EmitStatement(stmt)) << gen.error();
+  EXPECT_EQ(gen.result(),
+            R"(mat2x3 a = mat2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+)");
+}
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/generator_impl_workgroup_var_test.cc

@@ -0,0 +1,61 @@
+// 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 "gmock/gmock.h"
+#include "src/ast/override_decoration.h"
+#include "src/ast/stage_decoration.h"
+#include "src/writer/glsl/test_helper.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+namespace {
+using ::testing::HasSubstr;
+
+using GlslGeneratorImplTest_WorkgroupVar = TestHelper;
+
+TEST_F(GlslGeneratorImplTest_WorkgroupVar, Basic) {
+  Global("wg", ty.f32(), ast::StorageClass::kWorkgroup);
+
+  Func("main", {}, ty.void_(), {Assign("wg", 1.2f)},
+       {
+           Stage(ast::PipelineStage::kCompute),
+           WorkgroupSize(1),
+       });
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("groupshared float wg;\n"));
+}
+
+TEST_F(GlslGeneratorImplTest_WorkgroupVar, Aliased) {
+  auto* alias = Alias("F32", ty.f32());
+
+  Global("wg", ty.Of(alias), ast::StorageClass::kWorkgroup);
+
+  Func("main", {}, ty.void_(), {Assign("wg", 1.2f)},
+       {
+           Stage(ast::PipelineStage::kCompute),
+           WorkgroupSize(1),
+       });
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_THAT(gen.result(), HasSubstr("groupshared float wg;\n"));
+}
+
+}  // namespace
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint

src/writer/glsl/test_helper.h

@@ -0,0 +1,114 @@
+// 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_WRITER_GLSL_TEST_HELPER_H_
+#define SRC_WRITER_GLSL_TEST_HELPER_H_
+
+#include <memory>
+#include <string>
+#include <utility>
+
+#include "gtest/gtest.h"
+#include "src/transform/glsl.h"
+#include "src/transform/manager.h"
+#include "src/transform/renamer.h"
+#include "src/writer/glsl/generator_impl.h"
+
+namespace tint {
+namespace writer {
+namespace glsl {
+
+/// Helper class for testing
+template <typename BODY>
+class TestHelperBase : public BODY, public ProgramBuilder {
+ public:
+  TestHelperBase() = default;
+  ~TestHelperBase() override = default;
+
+  /// Builds the program and returns a GeneratorImpl from the program.
+  /// @note The generator is only built once. Multiple calls to Build() will
+  /// return the same GeneratorImpl without rebuilding.
+  /// @return the built generator
+  GeneratorImpl& Build() {
+    if (gen_) {
+      return *gen_;
+    }
+    // Fake that the GLSL sanitizer has been applied, so that we can unit test
+    // the writer without it erroring.
+    SetTransformApplied<transform::Glsl>();
+    [&]() {
+      ASSERT_TRUE(IsValid()) << "Builder program is not valid\n"
+                             << diag::Formatter().format(Diagnostics());
+    }();
+    program = std::make_unique<Program>(std::move(*this));
+    [&]() {
+      ASSERT_TRUE(program->IsValid())
+          << diag::Formatter().format(program->Diagnostics());
+    }();
+    gen_ = std::make_unique<GeneratorImpl>(program.get());
+    return *gen_;
+  }
+
+  /// Builds the program, runs the program through the transform::Glsl sanitizer
+  /// and returns a GeneratorImpl from the sanitized program.
+  /// @note The generator is only built once. Multiple calls to Build() will
+  /// return the same GeneratorImpl without rebuilding.
+  /// @return the built generator
+  GeneratorImpl& SanitizeAndBuild() {
+    if (gen_) {
+      return *gen_;
+    }
+    diag::Formatter formatter;
+    [&]() {
+      ASSERT_TRUE(IsValid()) << "Builder program is not valid\n"
+                             << formatter.format(Diagnostics());
+    }();
+    program = std::make_unique<Program>(std::move(*this));
+    [&]() {
+      ASSERT_TRUE(program->IsValid())
+          << formatter.format(program->Diagnostics());
+    }();
+
+    transform::Manager transform_manager;
+    transform::DataMap transform_data;
+    transform_data.Add<transform::Renamer::Config>(
+        transform::Renamer::Target::kGlslKeywords);
+    transform_manager.Add<tint::transform::Renamer>();
+    transform_manager.Add<tint::transform::Glsl>();
+    auto result = transform_manager.Run(program.get(), transform_data);
+    [&]() {
+      ASSERT_TRUE(result.program.IsValid())
+          << formatter.format(result.program.Diagnostics());
+    }();
+    *program = std::move(result.program);
+    gen_ = std::make_unique<GeneratorImpl>(program.get());
+    return *gen_;
+  }
+
+  /// The program built with a call to Build()
+  std::unique_ptr<Program> program;
+
+ private:
+  std::unique_ptr<GeneratorImpl> gen_;
+};
+using TestHelper = TestHelperBase<testing::Test>;
+
+template <typename T>
+using TestParamHelper = TestHelperBase<testing::TestWithParam<T>>;
+
+}  // namespace glsl
+}  // namespace writer
+}  // namespace tint
+
+#endif  // SRC_WRITER_GLSL_TEST_HELPER_H_

test/BUILD.gn

@@ -624,6 +624,47 @@ tint_unittests_source_set("tint_unittests_hlsl_writer_src") {
   ]
 }
 
+tint_unittests_source_set("tint_unittests_glsl_writer_src") {
+  sources = [
+    "../src/transform/glsl_test.cc",
+    "../src/writer/glsl/generator_impl_array_accessor_test.cc",
+    "../src/writer/glsl/generator_impl_assign_test.cc",
+    "../src/writer/glsl/generator_impl_binary_test.cc",
+    "../src/writer/glsl/generator_impl_bitcast_test.cc",
+    "../src/writer/glsl/generator_impl_block_test.cc",
+    "../src/writer/glsl/generator_impl_break_test.cc",
+    "../src/writer/glsl/generator_impl_call_test.cc",
+    "../src/writer/glsl/generator_impl_case_test.cc",
+    "../src/writer/glsl/generator_impl_cast_test.cc",
+    "../src/writer/glsl/generator_impl_constructor_test.cc",
+    "../src/writer/glsl/generator_impl_continue_test.cc",
+    "../src/writer/glsl/generator_impl_discard_test.cc",
+    "../src/writer/glsl/generator_impl_function_test.cc",
+    "../src/writer/glsl/generator_impl_identifier_test.cc",
+    "../src/writer/glsl/generator_impl_if_test.cc",
+    "../src/writer/glsl/generator_impl_import_test.cc",
+    "../src/writer/glsl/generator_impl_intrinsic_test.cc",
+    "../src/writer/glsl/generator_impl_intrinsic_texture_test.cc",
+    "../src/writer/glsl/generator_impl_loop_test.cc",
+    "../src/writer/glsl/generator_impl_member_accessor_test.cc",
+    "../src/writer/glsl/generator_impl_module_constant_test.cc",
+    "../src/writer/glsl/generator_impl_return_test.cc",
+    "../src/writer/glsl/generator_impl_sanitizer_test.cc",
+    "../src/writer/glsl/generator_impl_switch_test.cc",
+    "../src/writer/glsl/generator_impl_test.cc",
+    "../src/writer/glsl/generator_impl_type_test.cc",
+    "../src/writer/glsl/generator_impl_unary_op_test.cc",
+    "../src/writer/glsl/generator_impl_variable_decl_statement_test.cc",
+    "../src/writer/glsl/generator_impl_workgroup_var_test.cc",
+    "../src/writer/glsl/test_helper.h",
+  ]
+
+  deps = [
+    ":tint_unittests_core_src",
+    "${tint_root_dir}/src:libtint_glsl_writer_src",
+  ]
+}
+
 source_set("tint_unittests_src") {
   testonly = true
 
@@ -653,6 +694,10 @@ source_set("tint_unittests_src") {
     deps += [ ":tint_unittests_hlsl_writer_src" ]
   }
 
+  if (tint_build_glsl_writer) {
+    deps += [ ":tint_unittests_glsl_writer_src" ]
+  }
+
   configs += [ ":tint_unittests_config" ]
 
   if (build_with_chromium) {

tint_overrides_with_defaults.gni

@@ -66,4 +66,9 @@ declare_args() {
   if (!defined(tint_build_hlsl_writer)) {
     tint_build_hlsl_writer = true
   }
+
+  # Build the GLSL output writer
+  if (!defined(tint_build_glsl_writer)) {
+    tint_build_glsl_writer = true
+  }
 }