diff --git a/src/ast/function.cc b/src/ast/function.cc
index 1c22b9a1e9..7e076b35e5 100644
--- a/src/ast/function.cc
+++ b/src/ast/function.cc
@@ -69,7 +69,7 @@ PipelineStage Function::pipeline_stage() const {
 
 void Function::add_referenced_module_variable(Variable* var) {
   for (const auto* v : referenced_module_vars_) {
-    if (v->name() == var->name()) {
+    if (v->symbol() == var->symbol()) {
       return;
     }
   }
@@ -78,7 +78,7 @@ void Function::add_referenced_module_variable(Variable* var) {
 
 void Function::add_local_referenced_module_variable(Variable* var) {
   for (const auto* v : local_referenced_module_vars_) {
-    if (v->name() == var->name()) {
+    if (v->symbol() == var->symbol()) {
       return;
     }
   }
diff --git a/src/ast/variable.cc b/src/ast/variable.cc
index a01a9dd6d2..777b6e33a0 100644
--- a/src/ast/variable.cc
+++ b/src/ast/variable.cc
@@ -85,7 +85,7 @@ uint32_t Variable::constant_id() const {
 
 Variable* Variable::Clone(CloneContext* ctx) const {
   return ctx->mod->create<Variable>(ctx->Clone(source()), ctx->Clone(symbol_),
-                                    name(), storage_class(), ctx->Clone(type()),
+                                    name_, storage_class(), ctx->Clone(type()),
                                     is_const_, ctx->Clone(constructor()),
                                     ctx->Clone(decorations_));
 }
diff --git a/src/ast/variable.h b/src/ast/variable.h
index 5403264322..a805485d6d 100644
--- a/src/ast/variable.h
+++ b/src/ast/variable.h
@@ -104,8 +104,6 @@ class Variable : public Castable<Variable, Node> {
 
   /// @returns the variable symbol
   const Symbol& symbol() const { return symbol_; }
-  /// @returns the variable name
-  const std::string& name() const { return name_; }
 
   /// @returns the variable's type.
   type::Type* type() const { return type_; }
diff --git a/src/ast/variable_test.cc b/src/ast/variable_test.cc
index c97442cd4d..925415d336 100644
--- a/src/ast/variable_test.cc
+++ b/src/ast/variable_test.cc
@@ -30,7 +30,6 @@ TEST_F(VariableTest, Creation) {
   auto* v = Var("my_var", StorageClass::kFunction, ty.i32);
 
   EXPECT_EQ(v->symbol(), Symbol(1));
-  EXPECT_EQ(v->name(), "my_var");
   EXPECT_EQ(v->storage_class(), StorageClass::kFunction);
   EXPECT_EQ(v->type(), ty.i32);
   EXPECT_EQ(v->source().range.begin.line, 0u);
@@ -45,7 +44,6 @@ TEST_F(VariableTest, CreationWithSource) {
       "i", StorageClass::kPrivate, ty.f32, nullptr, VariableDecorationList{});
 
   EXPECT_EQ(v->symbol(), Symbol(1));
-  EXPECT_EQ(v->name(), "i");
   EXPECT_EQ(v->storage_class(), StorageClass::kPrivate);
   EXPECT_EQ(v->type(), ty.f32);
   EXPECT_EQ(v->source().range.begin.line, 27u);
@@ -61,7 +59,6 @@ TEST_F(VariableTest, CreationEmpty) {
       VariableDecorationList{});
 
   EXPECT_EQ(v->symbol(), Symbol(1));
-  EXPECT_EQ(v->name(), "a_var");
   EXPECT_EQ(v->storage_class(), StorageClass::kWorkgroup);
   EXPECT_EQ(v->type(), ty.i32);
   EXPECT_EQ(v->source().range.begin.line, 27u);
diff --git a/src/inspector/inspector.cc b/src/inspector/inspector.cc
index 8d8c1fefaa..a06bd77518 100644
--- a/src/inspector/inspector.cc
+++ b/src/inspector/inspector.cc
@@ -64,10 +64,11 @@ std::vector<EntryPoint> Inspector::GetEntryPoints() {
              entry_point.workgroup_size_z) = func->workgroup_size();
 
     for (auto* var : func->referenced_module_variables()) {
+      auto name = namer_->NameFor(var->symbol());
       if (var->storage_class() == ast::StorageClass::kInput) {
-        entry_point.input_variables.push_back(var->name());
+        entry_point.input_variables.push_back(name);
       } else {
-        entry_point.output_variables.push_back(var->name());
+        entry_point.output_variables.push_back(name);
       }
     }
     result.push_back(std::move(entry_point));
diff --git a/src/reader/wgsl/parser_impl_function_decl_test.cc b/src/reader/wgsl/parser_impl_function_decl_test.cc
index e7e38ec929..c8de53554a 100644
--- a/src/reader/wgsl/parser_impl_function_decl_test.cc
+++ b/src/reader/wgsl/parser_impl_function_decl_test.cc
@@ -42,8 +42,8 @@ TEST_F(ParserImplTest, FunctionDecl) {
   EXPECT_TRUE(f->return_type()->Is<ast::type::Void>());
 
   ASSERT_EQ(f->params().size(), 2u);
-  EXPECT_EQ(f->params()[0]->name(), "a");
-  EXPECT_EQ(f->params()[1]->name(), "b");
+  EXPECT_EQ(f->params()[0]->symbol(), p->get_module().RegisterSymbol("a"));
+  EXPECT_EQ(f->params()[1]->symbol(), p->get_module().RegisterSymbol("b"));
 
   ASSERT_NE(f->return_type(), nullptr);
   EXPECT_TRUE(f->return_type()->Is<ast::type::Void>());
diff --git a/src/reader/wgsl/parser_impl_function_header_test.cc b/src/reader/wgsl/parser_impl_function_header_test.cc
index 9b53b21edf..529de06e1f 100644
--- a/src/reader/wgsl/parser_impl_function_header_test.cc
+++ b/src/reader/wgsl/parser_impl_function_header_test.cc
@@ -33,8 +33,8 @@ TEST_F(ParserImplTest, FunctionHeader) {
 
   EXPECT_EQ(f->name, "main");
   ASSERT_EQ(f->params.size(), 2u);
-  EXPECT_EQ(f->params[0]->name(), "a");
-  EXPECT_EQ(f->params[1]->name(), "b");
+  EXPECT_EQ(f->params[0]->symbol(), p->get_module().RegisterSymbol("a"));
+  EXPECT_EQ(f->params[1]->symbol(), p->get_module().RegisterSymbol("b"));
   EXPECT_TRUE(f->return_type->Is<ast::type::Void>());
 }
 
diff --git a/src/reader/wgsl/parser_impl_global_constant_decl_test.cc b/src/reader/wgsl/parser_impl_global_constant_decl_test.cc
index 5e2804b712..89864db510 100644
--- a/src/reader/wgsl/parser_impl_global_constant_decl_test.cc
+++ b/src/reader/wgsl/parser_impl_global_constant_decl_test.cc
@@ -33,7 +33,7 @@ TEST_F(ParserImplTest, GlobalConstantDecl) {
   ASSERT_NE(e.value, nullptr);
 
   EXPECT_TRUE(e->is_const());
-  EXPECT_EQ(e->name(), "a");
+  EXPECT_EQ(e->symbol(), p->get_module().RegisterSymbol("a"));
   ASSERT_NE(e->type(), nullptr);
   EXPECT_TRUE(e->type()->Is<ast::type::F32>());
 
diff --git a/src/reader/wgsl/parser_impl_global_decl_test.cc b/src/reader/wgsl/parser_impl_global_decl_test.cc
index 04adade435..6256f02a4e 100644
--- a/src/reader/wgsl/parser_impl_global_decl_test.cc
+++ b/src/reader/wgsl/parser_impl_global_decl_test.cc
@@ -38,7 +38,7 @@ TEST_F(ParserImplTest, GlobalDecl_GlobalVariable) {
   ASSERT_EQ(m.global_variables().size(), 1u);
 
   auto* v = m.global_variables()[0];
-  EXPECT_EQ(v->name(), "a");
+  EXPECT_EQ(v->symbol(), p->get_module().RegisterSymbol("a"));
 }
 
 TEST_F(ParserImplTest, GlobalDecl_GlobalVariable_Invalid) {
@@ -64,7 +64,7 @@ TEST_F(ParserImplTest, GlobalDecl_GlobalConstant) {
   ASSERT_EQ(m.global_variables().size(), 1u);
 
   auto* v = m.global_variables()[0];
-  EXPECT_EQ(v->name(), "a");
+  EXPECT_EQ(v->symbol(), p->get_module().RegisterSymbol("a"));
 }
 
 TEST_F(ParserImplTest, GlobalDecl_GlobalConstant_Invalid) {
diff --git a/src/reader/wgsl/parser_impl_global_variable_decl_test.cc b/src/reader/wgsl/parser_impl_global_variable_decl_test.cc
index a42c8d33d8..be1199adb4 100644
--- a/src/reader/wgsl/parser_impl_global_variable_decl_test.cc
+++ b/src/reader/wgsl/parser_impl_global_variable_decl_test.cc
@@ -36,7 +36,7 @@ TEST_F(ParserImplTest, GlobalVariableDecl_WithoutConstructor) {
   EXPECT_FALSE(e.errored);
   ASSERT_NE(e.value, nullptr);
 
-  EXPECT_EQ(e->name(), "a");
+  EXPECT_EQ(e->symbol(), p->get_module().RegisterSymbol("a"));
   EXPECT_TRUE(e->type()->Is<ast::type::F32>());
   EXPECT_EQ(e->storage_class(), ast::StorageClass::kOutput);
 
@@ -59,7 +59,7 @@ TEST_F(ParserImplTest, GlobalVariableDecl_WithConstructor) {
   EXPECT_FALSE(e.errored);
   ASSERT_NE(e.value, nullptr);
 
-  EXPECT_EQ(e->name(), "a");
+  EXPECT_EQ(e->symbol(), p->get_module().RegisterSymbol("a"));
   EXPECT_TRUE(e->type()->Is<ast::type::F32>());
   EXPECT_EQ(e->storage_class(), ast::StorageClass::kOutput);
 
@@ -84,7 +84,7 @@ TEST_F(ParserImplTest, GlobalVariableDecl_WithDecoration) {
   EXPECT_FALSE(e.errored);
   ASSERT_NE(e.value, nullptr);
 
-  EXPECT_EQ(e->name(), "a");
+  EXPECT_EQ(e->symbol(), p->get_module().RegisterSymbol("a"));
   ASSERT_NE(e->type(), nullptr);
   EXPECT_TRUE(e->type()->Is<ast::type::F32>());
   EXPECT_EQ(e->storage_class(), ast::StorageClass::kOutput);
@@ -114,7 +114,7 @@ TEST_F(ParserImplTest, GlobalVariableDecl_WithDecoration_MulitpleGroups) {
   EXPECT_FALSE(e.errored);
   ASSERT_NE(e.value, nullptr);
 
-  EXPECT_EQ(e->name(), "a");
+  EXPECT_EQ(e->symbol(), p->get_module().RegisterSymbol("a"));
   ASSERT_NE(e->type(), nullptr);
   EXPECT_TRUE(e->type()->Is<ast::type::F32>());
   EXPECT_EQ(e->storage_class(), ast::StorageClass::kOutput);
diff --git a/src/reader/wgsl/parser_impl_param_list_test.cc b/src/reader/wgsl/parser_impl_param_list_test.cc
index 622514d174..58b3cbd673 100644
--- a/src/reader/wgsl/parser_impl_param_list_test.cc
+++ b/src/reader/wgsl/parser_impl_param_list_test.cc
@@ -38,7 +38,7 @@ TEST_F(ParserImplTest, ParamList_Single) {
   ASSERT_FALSE(e.errored);
   EXPECT_EQ(e.value.size(), 1u);
 
-  EXPECT_EQ(e.value[0]->name(), "a");
+  EXPECT_EQ(e.value[0]->symbol(), p->get_module().RegisterSymbol("a"));
   EXPECT_EQ(e.value[0]->type(), i32);
   EXPECT_TRUE(e.value[0]->is_const());
 
@@ -61,7 +61,7 @@ TEST_F(ParserImplTest, ParamList_Multiple) {
   ASSERT_FALSE(e.errored);
   EXPECT_EQ(e.value.size(), 3u);
 
-  EXPECT_EQ(e.value[0]->name(), "a");
+  EXPECT_EQ(e.value[0]->symbol(), p->get_module().RegisterSymbol("a"));
   EXPECT_EQ(e.value[0]->type(), i32);
   EXPECT_TRUE(e.value[0]->is_const());
 
@@ -70,7 +70,7 @@ TEST_F(ParserImplTest, ParamList_Multiple) {
   ASSERT_EQ(e.value[0]->source().range.end.line, 1u);
   ASSERT_EQ(e.value[0]->source().range.end.column, 2u);
 
-  EXPECT_EQ(e.value[1]->name(), "b");
+  EXPECT_EQ(e.value[1]->symbol(), p->get_module().RegisterSymbol("b"));
   EXPECT_EQ(e.value[1]->type(), f32);
   EXPECT_TRUE(e.value[1]->is_const());
 
@@ -79,7 +79,7 @@ TEST_F(ParserImplTest, ParamList_Multiple) {
   ASSERT_EQ(e.value[1]->source().range.end.line, 1u);
   ASSERT_EQ(e.value[1]->source().range.end.column, 11u);
 
-  EXPECT_EQ(e.value[2]->name(), "c");
+  EXPECT_EQ(e.value[2]->symbol(), p->get_module().RegisterSymbol("c"));
   EXPECT_EQ(e.value[2]->type(), vec2);
   EXPECT_TRUE(e.value[2]->is_const());
 
diff --git a/src/reader/wgsl/parser_impl_variable_stmt_test.cc b/src/reader/wgsl/parser_impl_variable_stmt_test.cc
index d541af3f43..1a4aebaa73 100644
--- a/src/reader/wgsl/parser_impl_variable_stmt_test.cc
+++ b/src/reader/wgsl/parser_impl_variable_stmt_test.cc
@@ -32,7 +32,7 @@ TEST_F(ParserImplTest, VariableStmt_VariableDecl) {
   ASSERT_NE(e.value, nullptr);
   ASSERT_TRUE(e->Is<ast::VariableDeclStatement>());
   ASSERT_NE(e->variable(), nullptr);
-  EXPECT_EQ(e->variable()->name(), "a");
+  EXPECT_EQ(e->variable()->symbol(), p->get_module().RegisterSymbol("a"));
 
   ASSERT_EQ(e->source().range.begin.line, 1u);
   ASSERT_EQ(e->source().range.begin.column, 5u);
@@ -51,7 +51,7 @@ TEST_F(ParserImplTest, VariableStmt_VariableDecl_WithInit) {
   ASSERT_NE(e.value, nullptr);
   ASSERT_TRUE(e->Is<ast::VariableDeclStatement>());
   ASSERT_NE(e->variable(), nullptr);
-  EXPECT_EQ(e->variable()->name(), "a");
+  EXPECT_EQ(e->variable()->symbol(), p->get_module().RegisterSymbol("a"));
 
   ASSERT_EQ(e->source().range.begin.line, 1u);
   ASSERT_EQ(e->source().range.begin.column, 5u);
diff --git a/src/transform/first_index_offset.cc b/src/transform/first_index_offset.cc
index db1ddd4c8d..3cafccb748 100644
--- a/src/transform/first_index_offset.cc
+++ b/src/transform/first_index_offset.cc
@@ -86,7 +86,7 @@ Transform::Output FirstIndexOffset::Run(ast::Module* in) {
   // First do a quick check to see if the transform has already been applied.
   for (ast::Variable* var : in->global_variables()) {
     if (auto* dec_var = var->As<ast::Variable>()) {
-      if (dec_var->name() == kBufferName) {
+      if (dec_var->symbol() == in->RegisterSymbol(kBufferName)) {
         diag::Diagnostic err;
         err.message = "First index offset transform has already been applied.";
         err.severity = diag::Severity::Error;
@@ -110,8 +110,8 @@ Transform::Output FirstIndexOffset::Run(ast::Module* in) {
     return out;
   }
 
-  std::string vertex_index_name;
-  std::string instance_index_name;
+  Symbol vertex_index_sym;
+  Symbol instance_index_sym;
 
   // Lazilly construct the UniformBuffer on first call to
   // maybe_create_buffer_var()
@@ -134,15 +134,17 @@ Transform::Output FirstIndexOffset::Run(ast::Module* in) {
               if (auto* blt_dec = dec->As<ast::BuiltinDecoration>()) {
                 ast::Builtin blt_type = blt_dec->value();
                 if (blt_type == ast::Builtin::kVertexIdx) {
-                  vertex_index_name = var->name();
+                  vertex_index_sym = var->symbol();
                   has_vertex_index_ = true;
                   return clone_variable_with_new_name(
-                      ctx, var, kIndexOffsetPrefix + var->name());
+                      ctx, var,
+                      kIndexOffsetPrefix + in->SymbolToName(var->symbol()));
                 } else if (blt_type == ast::Builtin::kInstanceIdx) {
-                  instance_index_name = var->name();
+                  instance_index_sym = var->symbol();
                   has_instance_index_ = true;
                   return clone_variable_with_new_name(
-                      ctx, var, kIndexOffsetPrefix + var->name());
+                      ctx, var,
+                      kIndexOffsetPrefix + in->SymbolToName(var->symbol()));
                 }
               }
             }
@@ -161,11 +163,12 @@ Transform::Output FirstIndexOffset::Run(ast::Module* in) {
                  func->local_referenced_builtin_variables()) {
               if (data.second->value() == ast::Builtin::kVertexIdx) {
                 statements.emplace_back(CreateFirstIndexOffset(
-                    vertex_index_name, kFirstVertexName, buffer_var, ctx->mod));
+                    in->SymbolToName(vertex_index_sym), kFirstVertexName,
+                    buffer_var, ctx->mod));
               } else if (data.second->value() == ast::Builtin::kInstanceIdx) {
                 statements.emplace_back(CreateFirstIndexOffset(
-                    instance_index_name, kFirstInstanceName, buffer_var,
-                    ctx->mod));
+                    in->SymbolToName(instance_index_sym), kFirstInstanceName,
+                    buffer_var, ctx->mod));
               }
             }
             return CloneWithStatementsAtStart(ctx, func, statements);
@@ -252,12 +255,13 @@ ast::VariableDeclStatement* FirstIndexOffset::CreateFirstIndexOffset(
     ast::Variable* buffer_var,
     ast::Module* mod) {
   auto* buffer = mod->create<ast::IdentifierExpression>(
-      Source{}, mod->RegisterSymbol(buffer_var->name()), buffer_var->name());
+      Source{}, buffer_var->symbol(), mod->SymbolToName(buffer_var->symbol()));
+
+  auto lhs_name = kIndexOffsetPrefix + original_name;
   auto* constructor = mod->create<ast::BinaryExpression>(
       Source{}, ast::BinaryOp::kAdd,
       mod->create<ast::IdentifierExpression>(
-          Source{}, mod->RegisterSymbol(kIndexOffsetPrefix + original_name),
-          kIndexOffsetPrefix + original_name),
+          Source{}, mod->RegisterSymbol(lhs_name), lhs_name),
       mod->create<ast::MemberAccessorExpression>(
           Source{}, buffer,
           mod->create<ast::IdentifierExpression>(
diff --git a/src/transform/first_index_offset.h b/src/transform/first_index_offset.h
index 873ffc822b..64ad7a33ae 100644
--- a/src/transform/first_index_offset.h
+++ b/src/transform/first_index_offset.h
@@ -19,6 +19,7 @@
 
 #include "src/ast/module.h"
 #include "src/ast/variable_decl_statement.h"
+#include "src/symbol.h"
 #include "src/transform/transform.h"
 
 namespace tint {
diff --git a/src/transform/vertex_pulling.cc b/src/transform/vertex_pulling.cc
index 7b986faa0b..9709919d39 100644
--- a/src/transform/vertex_pulling.cc
+++ b/src/transform/vertex_pulling.cc
@@ -155,7 +155,7 @@ void VertexPulling::State::FindOrInsertVertexIndexIfUsed() {
     for (auto* d : v->decorations()) {
       if (auto* builtin = d->As<ast::BuiltinDecoration>()) {
         if (builtin->value() == ast::Builtin::kVertexIdx) {
-          vertex_index_name = v->name();
+          vertex_index_name = in->SymbolToName(v->symbol());
           return;
         }
       }
@@ -203,7 +203,7 @@ void VertexPulling::State::FindOrInsertInstanceIndexIfUsed() {
     for (auto* d : v->decorations()) {
       if (auto* builtin = d->As<ast::BuiltinDecoration>()) {
         if (builtin->value() == ast::Builtin::kInstanceIdx) {
-          instance_index_name = v->name();
+          instance_index_name = in->SymbolToName(v->symbol());
           return;
         }
       }
@@ -244,7 +244,7 @@ void VertexPulling::State::ConvertVertexInputVariablesToPrivate() {
         v = out->create<ast::Variable>(
             Source{},                        // source
             v->symbol(),                     // symbol
-            v->name(),                       // name
+            out->SymbolToName(v->symbol()),  // name
             ast::StorageClass::kPrivate,     // storage_class
             v->type(),                       // type
             false,                           // is_const
@@ -358,10 +358,11 @@ ast::BlockStatement* VertexPulling::State::CreateVertexPullingPreamble() const {
           Source{}, CreatePullingPositionIdent(), pos_value);
       stmts.emplace_back(set_pos_expr);
 
+      auto ident_name = in->SymbolToName(v->symbol());
       stmts.emplace_back(out->create<ast::AssignmentStatement>(
           Source{},
           out->create<ast::IdentifierExpression>(
-              Source{}, out->RegisterSymbol(v->name()), v->name()),
+              Source{}, out->RegisterSymbol(ident_name), ident_name),
           AccessByFormat(i, attribute_desc.format)));
     }
   }
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index 9dae4b847b..051e81f6f8 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -175,7 +175,7 @@ bool GeneratorImpl::Generate(std::ostream& out) {
     }
   }
 
-  std::unordered_set<std::string> emitted_globals;
+  std::unordered_set<uint32_t> emitted_globals;
   // Make sure all entry point data is emitted before the entry point functions
   for (auto* func : module_->functions()) {
     if (!func->IsEntryPoint()) {
@@ -1285,7 +1285,7 @@ bool GeneratorImpl::EmitFunctionInternal(std::ostream& out,
     }
     // Array name is output as part of the type
     if (!v->type()->Is<ast::type::Array>()) {
-      out << " " << v->name();
+      out << " " << namer_->NameFor(v->symbol());
     }
   }
 
@@ -1305,7 +1305,7 @@ bool GeneratorImpl::EmitFunctionInternal(std::ostream& out,
 bool GeneratorImpl::EmitEntryPointData(
     std::ostream& out,
     ast::Function* func,
-    std::unordered_set<std::string>& emitted_globals) {
+    std::unordered_set<uint32_t>& emitted_globals) {
   std::vector<std::pair<ast::Variable*, ast::VariableDecoration*>> in_variables;
   std::vector<std::pair<ast::Variable*, ast::VariableDecoration*>> outvariables;
   for (auto data : func->referenced_location_variables()) {
@@ -1338,22 +1338,24 @@ bool GeneratorImpl::EmitEntryPointData(
     // set. https://bugs.chromium.org/p/tint/issues/detail?id=104
     auto* binding = data.second.binding;
     if (binding == nullptr) {
-      error_ = "unable to find binding information for uniform: " + var->name();
+      error_ = "unable to find binding information for uniform: " +
+               module_->SymbolToName(var->symbol());
       return false;
     }
     // auto* set = data.second.set;
 
     // If the global has already been emitted we skip it, it's been emitted by
     // a previous entry point.
-    if (emitted_globals.count(var->name()) != 0) {
+    if (emitted_globals.count(var->symbol().value()) != 0) {
       continue;
     }
-    emitted_globals.insert(var->name());
+    emitted_globals.insert(var->symbol().value());
 
     auto* type = var->type()->UnwrapIfNeeded();
     if (auto* strct = type->As<ast::type::Struct>()) {
-      out << "ConstantBuffer<" << strct->name() << "> " << var->name()
-          << " : register(b" << binding->value() << ");" << std::endl;
+      out << "ConstantBuffer<" << strct->name() << "> "
+          << namer_->NameFor(var->symbol()) << " : register(b"
+          << binding->value() << ");" << std::endl;
     } else {
       // TODO(dsinclair): There is outstanding spec work to require all uniform
       // buffers to be [[block]] decorated, which means structs. This is
@@ -1361,7 +1363,7 @@ bool GeneratorImpl::EmitEntryPointData(
       // is not a block.
       // Relevant: https://github.com/gpuweb/gpuweb/issues/1004
       //           https://github.com/gpuweb/gpuweb/issues/1008
-      auto name = "cbuffer_" + var->name();
+      auto name = "cbuffer_" + namer_->NameFor(var->symbol());
       out << "cbuffer " << name << " : register(b" << binding->value() << ") {"
           << std::endl;
 
@@ -1370,7 +1372,7 @@ bool GeneratorImpl::EmitEntryPointData(
       if (!EmitType(out, type, Symbol())) {
         return false;
       }
-      out << " " << var->name() << ";" << std::endl;
+      out << " " << namer_->NameFor(var->symbol()) << ";" << std::endl;
       decrement_indent();
       out << "};" << std::endl;
     }
@@ -1388,10 +1390,10 @@ bool GeneratorImpl::EmitEntryPointData(
 
     // If the global has already been emitted we skip it, it's been emitted by
     // a previous entry point.
-    if (emitted_globals.count(var->name()) != 0) {
+    if (emitted_globals.count(var->symbol().value()) != 0) {
       continue;
     }
-    emitted_globals.insert(var->name());
+    emitted_globals.insert(var->symbol().value());
 
     auto* ac = var->type()->As<ast::type::AccessControl>();
     if (ac == nullptr) {
@@ -1402,8 +1404,8 @@ bool GeneratorImpl::EmitEntryPointData(
     if (ac->IsReadWrite()) {
       out << "RW";
     }
-    out << "ByteAddressBuffer " << var->name() << " : register(u"
-        << binding->value() << ");" << std::endl;
+    out << "ByteAddressBuffer " << namer_->NameFor(var->symbol())
+        << " : register(u" << binding->value() << ");" << std::endl;
     emitted_storagebuffer = true;
   }
   if (emitted_storagebuffer) {
@@ -1430,7 +1432,7 @@ bool GeneratorImpl::EmitEntryPointData(
         return false;
       }
 
-      out << " " << var->name() << " : ";
+      out << " " << namer_->NameFor(var->symbol()) << " : ";
       if (auto* location = deco->As<ast::LocationDecoration>()) {
         if (func->pipeline_stage() == ast::PipelineStage::kCompute) {
           error_ = "invalid location variable for pipeline stage";
@@ -1475,7 +1477,7 @@ bool GeneratorImpl::EmitEntryPointData(
         return false;
       }
 
-      out << " " << var->name() << " : ";
+      out << " " << namer_->NameFor(var->symbol()) << " : ";
 
       if (auto* location = deco->As<ast::LocationDecoration>()) {
         auto loc = location->value();
@@ -1697,7 +1699,7 @@ bool GeneratorImpl::EmitLoop(std::ostream& out, ast::LoopStatement* stmt) {
         }
         out << pre.str();
 
-        out << var->name() << " = ";
+        out << namer_->NameFor(var->symbol()) << " = ";
         if (var->constructor() != nullptr) {
           out << constructor_out.str();
         } else {
@@ -2289,7 +2291,7 @@ bool GeneratorImpl::EmitVariable(std::ostream& out,
     return false;
   }
   if (!var->type()->Is<ast::type::Array>()) {
-    out << " " << var->name();
+    out << " " << namer_->NameFor(var->symbol());
   }
   out << constructor_out.str() << ";" << std::endl;
 
@@ -2337,8 +2339,8 @@ bool GeneratorImpl::EmitProgramConstVariable(std::ostream& out,
     if (!EmitType(out, var->type(), var->symbol())) {
       return false;
     }
-    out << " " << var->name() << " = WGSL_SPEC_CONSTANT_" << const_id << ";"
-        << std::endl;
+    out << " " << namer_->NameFor(var->symbol()) << " = WGSL_SPEC_CONSTANT_"
+        << const_id << ";" << std::endl;
     out << "#undef WGSL_SPEC_CONSTANT_" << const_id << std::endl;
   } else {
     out << "static const ";
@@ -2346,7 +2348,7 @@ bool GeneratorImpl::EmitProgramConstVariable(std::ostream& out,
       return false;
     }
     if (!var->type()->Is<ast::type::Array>()) {
-      out << " " << var->name();
+      out << " " << namer_->NameFor(var->symbol());
     }
 
     if (var->constructor() != nullptr) {
diff --git a/src/writer/hlsl/generator_impl.h b/src/writer/hlsl/generator_impl.h
index 4a2a371971..9c1a1132cb 100644
--- a/src/writer/hlsl/generator_impl.h
+++ b/src/writer/hlsl/generator_impl.h
@@ -224,7 +224,7 @@ class GeneratorImpl {
   /// @returns true if the entry point data was emitted
   bool EmitEntryPointData(std::ostream& out,
                           ast::Function* func,
-                          std::unordered_set<std::string>& emitted_globals);
+                          std::unordered_set<uint32_t>& emitted_globals);
   /// Handles emitting the entry point function
   /// @param out the output stream
   /// @param func the entry point
diff --git a/src/writer/hlsl/generator_impl_function_entry_point_data_test.cc b/src/writer/hlsl/generator_impl_function_entry_point_data_test.cc
index 6bce585852..73bd4c43e7 100644
--- a/src/writer/hlsl/generator_impl_function_entry_point_data_test.cc
+++ b/src/writer/hlsl/generator_impl_function_entry_point_data_test.cc
@@ -72,7 +72,7 @@ TEST_F(HlslGeneratorImplTest_EntryPoint,
 
   mod->AddFunction(func);
 
-  std::unordered_set<std::string> globals;
+  std::unordered_set<uint32_t> globals;
 
   ASSERT_TRUE(td.Determine()) << td.error();
   ASSERT_TRUE(gen.EmitEntryPointData(out, func, globals)) << gen.error();
@@ -122,7 +122,7 @@ TEST_F(HlslGeneratorImplTest_EntryPoint,
 
   mod->AddFunction(func);
 
-  std::unordered_set<std::string> globals;
+  std::unordered_set<uint32_t> globals;
 
   ASSERT_TRUE(td.Determine()) << td.error();
   ASSERT_TRUE(gen.EmitEntryPointData(out, func, globals)) << gen.error();
@@ -172,7 +172,7 @@ TEST_F(HlslGeneratorImplTest_EntryPoint,
 
   mod->AddFunction(func);
 
-  std::unordered_set<std::string> globals;
+  std::unordered_set<uint32_t> globals;
 
   ASSERT_TRUE(td.Determine()) << td.error();
   ASSERT_TRUE(gen.EmitEntryPointData(out, func, globals)) << gen.error();
@@ -222,7 +222,7 @@ TEST_F(HlslGeneratorImplTest_EntryPoint,
 
   mod->AddFunction(func);
 
-  std::unordered_set<std::string> globals;
+  std::unordered_set<uint32_t> globals;
 
   ASSERT_TRUE(td.Determine()) << td.error();
   ASSERT_TRUE(gen.EmitEntryPointData(out, func, globals)) << gen.error();
@@ -269,7 +269,7 @@ TEST_F(HlslGeneratorImplTest_EntryPoint,
 
   mod->AddFunction(func);
 
-  std::unordered_set<std::string> globals;
+  std::unordered_set<uint32_t> globals;
 
   ASSERT_TRUE(td.Determine()) << td.error();
   ASSERT_FALSE(gen.EmitEntryPointData(out, func, globals)) << gen.error();
@@ -311,7 +311,7 @@ TEST_F(HlslGeneratorImplTest_EntryPoint,
 
   mod->AddFunction(func);
 
-  std::unordered_set<std::string> globals;
+  std::unordered_set<uint32_t> globals;
 
   ASSERT_TRUE(td.Determine()) << td.error();
   ASSERT_FALSE(gen.EmitEntryPointData(out, func, globals)) << gen.error();
@@ -361,7 +361,7 @@ TEST_F(HlslGeneratorImplTest_EntryPoint,
 
   mod->AddFunction(func);
 
-  std::unordered_set<std::string> globals;
+  std::unordered_set<uint32_t> globals;
 
   ASSERT_TRUE(td.Determine()) << td.error();
   ASSERT_TRUE(gen.EmitEntryPointData(out, func, globals)) << gen.error();
diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc
index 0ae66607e1..f9b425a025 100644
--- a/src/writer/msl/generator_impl.cc
+++ b/src/writer/msl/generator_impl.cc
@@ -607,7 +607,7 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) {
       out_ << ", ";
     }
     first = false;
-    out_ << var->name();
+    out_ << namer_->NameFor(var->symbol());
   }
 
   for (const auto& data : func->referenced_uniform_variables()) {
@@ -616,7 +616,7 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) {
       out_ << ", ";
     }
     first = false;
-    out_ << var->name();
+    out_ << namer_->NameFor(var->symbol());
   }
 
   for (const auto& data : func->referenced_storagebuffer_variables()) {
@@ -625,7 +625,7 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) {
       out_ << ", ";
     }
     first = false;
-    out_ << var->name();
+    out_ << namer_->NameFor(var->symbol());
   }
 
   const auto& params = expr->params();
@@ -1035,7 +1035,7 @@ bool GeneratorImpl::EmitEntryPointData(ast::Function* func) {
         return false;
       }
 
-      out_ << " " << var->name() << " [[";
+      out_ << " " << namer_->NameFor(var->symbol()) << " [[";
       if (func->pipeline_stage() == ast::PipelineStage::kVertex) {
         out_ << "attribute(" << loc << ")";
       } else if (func->pipeline_stage() == ast::PipelineStage::kFragment) {
@@ -1072,7 +1072,7 @@ bool GeneratorImpl::EmitEntryPointData(ast::Function* func) {
         return false;
       }
 
-      out_ << " " << var->name() << " [[";
+      out_ << " " << namer_->NameFor(var->symbol()) << " [[";
 
       if (auto* location = deco->As<ast::LocationDecoration>()) {
         auto loc = location->value();
@@ -1275,7 +1275,7 @@ bool GeneratorImpl::EmitFunctionInternal(ast::Function* func,
     if (!EmitType(var->type(), Symbol())) {
       return false;
     }
-    out_ << "& " << var->name();
+    out_ << "& " << namer_->NameFor(var->symbol());
   }
 
   for (const auto& data : func->referenced_uniform_variables()) {
@@ -1290,7 +1290,7 @@ bool GeneratorImpl::EmitFunctionInternal(ast::Function* func,
     if (!EmitType(var->type(), Symbol())) {
       return false;
     }
-    out_ << "& " << var->name();
+    out_ << "& " << namer_->NameFor(var->symbol());
   }
 
   for (const auto& data : func->referenced_storagebuffer_variables()) {
@@ -1313,7 +1313,7 @@ bool GeneratorImpl::EmitFunctionInternal(ast::Function* func,
     if (!EmitType(ac->type(), Symbol())) {
       return false;
     }
-    out_ << "& " << var->name();
+    out_ << "& " << namer_->NameFor(var->symbol());
   }
 
   for (auto* v : func->params()) {
@@ -1327,7 +1327,7 @@ bool GeneratorImpl::EmitFunctionInternal(ast::Function* func,
     }
     // Array name is output as part of the type
     if (!v->type()->Is<ast::type::Array>()) {
-      out_ << " " << v->name();
+      out_ << " " << namer_->NameFor(v->symbol());
     }
   }
 
@@ -1419,7 +1419,7 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) {
       error_ = "unknown builtin";
       return false;
     }
-    out_ << " " << var->name() << " [[" << attr << "]]";
+    out_ << " " << namer_->NameFor(var->symbol()) << " [[" << attr << "]]";
   }
 
   for (auto data : func->referenced_uniform_variables()) {
@@ -1434,7 +1434,8 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) {
     // set. https://bugs.chromium.org/p/tint/issues/detail?id=104
     auto* binding = data.second.binding;
     if (binding == nullptr) {
-      error_ = "unable to find binding information for uniform: " + var->name();
+      error_ = "unable to find binding information for uniform: " +
+               module_->SymbolToName(var->symbol());
       return false;
     }
     // auto* set = data.second.set;
@@ -1445,7 +1446,8 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) {
     if (!EmitType(var->type(), Symbol())) {
       return false;
     }
-    out_ << "& " << var->name() << " [[buffer(" << binding->value() << ")]]";
+    out_ << "& " << namer_->NameFor(var->symbol()) << " [[buffer("
+         << binding->value() << ")]]";
   }
 
   for (auto data : func->referenced_storagebuffer_variables()) {
@@ -1474,7 +1476,8 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) {
     if (!EmitType(ac->type(), Symbol())) {
       return false;
     }
-    out_ << "& " << var->name() << " [[buffer(" << binding->value() << ")]]";
+    out_ << "& " << namer_->NameFor(var->symbol()) << " [[buffer("
+         << binding->value() << ")]]";
   }
 
   out_ << ") {" << std::endl;
@@ -1589,7 +1592,7 @@ bool GeneratorImpl::EmitLoop(ast::LoopStatement* stmt) {
       make_indent();
 
       auto* var = decl->variable();
-      out_ << var->name() << " = ";
+      out_ << namer_->NameFor(var->symbol()) << " = ";
       if (var->constructor() != nullptr) {
         if (!EmitExpression(var->constructor())) {
           return false;
@@ -2020,7 +2023,7 @@ bool GeneratorImpl::EmitVariable(ast::Variable* var, bool skip_constructor) {
     return false;
   }
   if (!var->type()->Is<ast::type::Array>()) {
-    out_ << " " << var->name();
+    out_ << " " << namer_->NameFor(var->symbol());
   }
 
   if (!skip_constructor) {
@@ -2062,7 +2065,7 @@ bool GeneratorImpl::EmitProgramConstVariable(const ast::Variable* var) {
     return false;
   }
   if (!var->type()->Is<ast::type::Array>()) {
-    out_ << " " << var->name();
+    out_ << " " << namer_->NameFor(var->symbol());
   }
 
   if (var->HasConstantIdDecoration()) {
diff --git a/src/writer/wgsl/generator.cc b/src/writer/wgsl/generator.cc
index 3f01dc9e7f..e9d0a9b052 100644
--- a/src/writer/wgsl/generator.cc
+++ b/src/writer/wgsl/generator.cc
@@ -21,17 +21,18 @@ namespace writer {
 namespace wgsl {
 
 Generator::Generator(ast::Module module)
-    : Text(std::move(module)), impl_(std::make_unique<GeneratorImpl>()) {}
+    : Text(std::move(module)),
+      impl_(std::make_unique<GeneratorImpl>(&module_)) {}
 
 Generator::~Generator() = default;
 
 void Generator::Reset() {
   set_error("");
-  impl_ = std::make_unique<GeneratorImpl>();
+  impl_ = std::make_unique<GeneratorImpl>(&module_);
 }
 
 bool Generator::Generate() {
-  auto ret = impl_->Generate(module_);
+  auto ret = impl_->Generate();
   if (!ret) {
     error_ = impl_->error();
   }
@@ -40,7 +41,7 @@ bool Generator::Generate() {
 
 bool Generator::GenerateEntryPoint(ast::PipelineStage stage,
                                    const std::string& name) {
-  auto ret = impl_->GenerateEntryPoint(module_, stage, name);
+  auto ret = impl_->GenerateEntryPoint(stage, name);
   if (!ret) {
     error_ = impl_->error();
   }
diff --git a/src/writer/wgsl/generator_impl.cc b/src/writer/wgsl/generator_impl.cc
index 58982ddd2e..09509b447c 100644
--- a/src/writer/wgsl/generator_impl.cc
+++ b/src/writer/wgsl/generator_impl.cc
@@ -78,29 +78,30 @@ namespace tint {
 namespace writer {
 namespace wgsl {
 
-GeneratorImpl::GeneratorImpl() : TextGenerator() {}
+GeneratorImpl::GeneratorImpl(ast::Module* module)
+    : TextGenerator(), module_(*module) {}
 
 GeneratorImpl::~GeneratorImpl() = default;
 
-bool GeneratorImpl::Generate(const ast::Module& module) {
-  for (auto* const ty : module.constructed_types()) {
+bool GeneratorImpl::Generate() {
+  for (auto* const ty : module_.constructed_types()) {
     if (!EmitConstructedType(ty)) {
       return false;
     }
   }
-  if (!module.constructed_types().empty())
+  if (!module_.constructed_types().empty())
     out_ << std::endl;
 
-  for (auto* var : module.global_variables()) {
+  for (auto* var : module_.global_variables()) {
     if (!EmitVariable(var)) {
       return false;
     }
   }
-  if (!module.global_variables().empty()) {
+  if (!module_.global_variables().empty()) {
     out_ << std::endl;
   }
 
-  for (auto* func : module.functions()) {
+  for (auto* func : module_.functions()) {
     if (!EmitFunction(func)) {
       return false;
     }
@@ -110,11 +111,10 @@ bool GeneratorImpl::Generate(const ast::Module& module) {
   return true;
 }
 
-bool GeneratorImpl::GenerateEntryPoint(const ast::Module& module,
-                                       ast::PipelineStage stage,
+bool GeneratorImpl::GenerateEntryPoint(ast::PipelineStage stage,
                                        const std::string& name) {
   auto* func =
-      module.FindFunctionBySymbolAndStage(module.GetSymbol(name), stage);
+      module_.FindFunctionBySymbolAndStage(module_.GetSymbol(name), stage);
   if (func == nullptr) {
     error_ = "Unable to find requested entry point: " + name;
     return false;
@@ -122,18 +122,18 @@ bool GeneratorImpl::GenerateEntryPoint(const ast::Module& module,
 
   // TODO(dsinclair): We always emit constructed types even if they aren't
   // strictly needed
-  for (auto* const ty : module.constructed_types()) {
+  for (auto* const ty : module_.constructed_types()) {
     if (!EmitConstructedType(ty)) {
       return false;
     }
   }
-  if (!module.constructed_types().empty()) {
+  if (!module_.constructed_types().empty()) {
     out_ << std::endl;
   }
 
   // TODO(dsinclair): This should be smarter and only emit needed const
   // variables
-  for (auto* var : module.global_variables()) {
+  for (auto* var : module_.global_variables()) {
     if (!var->is_const()) {
       continue;
     }
@@ -153,8 +153,8 @@ bool GeneratorImpl::GenerateEntryPoint(const ast::Module& module,
     out_ << std::endl;
   }
 
-  for (auto* f : module.functions()) {
-    if (!f->HasAncestorEntryPoint(module.GetSymbol(name))) {
+  for (auto* f : module_.functions()) {
+    if (!f->HasAncestorEntryPoint(module_.GetSymbol(name))) {
       continue;
     }
 
@@ -370,7 +370,7 @@ bool GeneratorImpl::EmitFunction(ast::Function* func) {
     }
     first = false;
 
-    out_ << v->name() << " : ";
+    out_ << module_.SymbolToName(v->symbol()) << " : ";
 
     if (!EmitType(v->type())) {
       return false;
@@ -597,7 +597,7 @@ bool GeneratorImpl::EmitVariable(ast::Variable* var) {
     }
   }
 
-  out_ << " " << var->name() << " : ";
+  out_ << " " << module_.SymbolToName(var->symbol()) << " : ";
   if (!EmitType(var->type())) {
     return false;
   }
diff --git a/src/writer/wgsl/generator_impl.h b/src/writer/wgsl/generator_impl.h
index c4544e34b9..f21b602a15 100644
--- a/src/writer/wgsl/generator_impl.h
+++ b/src/writer/wgsl/generator_impl.h
@@ -53,22 +53,19 @@ namespace wgsl {
 class GeneratorImpl : public TextGenerator {
  public:
   /// Constructor
-  GeneratorImpl();
+  /// @param module the module to generate
+  explicit GeneratorImpl(ast::Module* mod);
   ~GeneratorImpl();
 
   /// Generates the result data
-  /// @param module the module to generate
   /// @returns true on successful generation; false otherwise
-  bool Generate(const ast::Module& module);
+  bool Generate();
 
   /// Generates a single entry point
-  /// @param module the module to generate from
   /// @param stage the pipeline stage
   /// @param name the entry point name
   /// @returns true on successful generation; false otherwise
-  bool GenerateEntryPoint(const ast::Module& module,
-                          ast::PipelineStage stage,
-                          const std::string& name);
+  bool GenerateEntryPoint(ast::PipelineStage stage, const std::string& name);
 
   /// Handles generating a constructed type
   /// @param ty the constructed to generate
@@ -206,6 +203,9 @@ class GeneratorImpl : public TextGenerator {
   /// @param var the decorated variable
   /// @returns true if the variable decoration was emitted
   bool EmitVariableDecorations(ast::Variable* var);
+
+ private:
+  ast::Module& module_;
 };
 
 }  // namespace wgsl
diff --git a/src/writer/wgsl/generator_impl_function_test.cc b/src/writer/wgsl/generator_impl_function_test.cc
index 2deb562379..37d61cb503 100644
--- a/src/writer/wgsl/generator_impl_function_test.cc
+++ b/src/writer/wgsl/generator_impl_function_test.cc
@@ -226,7 +226,7 @@ TEST_F(WgslGeneratorImplTest,
 
   ASSERT_TRUE(td.Determine()) << td.error();
 
-  ASSERT_TRUE(gen.Generate(*mod)) << gen.error();
+  ASSERT_TRUE(gen.Generate()) << gen.error();
   EXPECT_EQ(gen.result(), R"([[block]]
 struct Data {
   [[offset(0)]]
diff --git a/src/writer/wgsl/generator_impl_test.cc b/src/writer/wgsl/generator_impl_test.cc
index 99121cad5a..37f080a3f9 100644
--- a/src/writer/wgsl/generator_impl_test.cc
+++ b/src/writer/wgsl/generator_impl_test.cc
@@ -33,7 +33,7 @@ TEST_F(WgslGeneratorImplTest, Generate) {
   mod->AddFunction(Func("my_func", ast::VariableList{}, ty.void_,
                         ast::StatementList{}, ast::FunctionDecorationList{}));
 
-  ASSERT_TRUE(gen.Generate(*mod)) << gen.error();
+  ASSERT_TRUE(gen.Generate()) << gen.error();
   EXPECT_EQ(gen.result(), R"(fn my_func() -> void {
 }
 
diff --git a/src/writer/wgsl/test_helper.h b/src/writer/wgsl/test_helper.h
index d6c5a93393..8ef37bda96 100644
--- a/src/writer/wgsl/test_helper.h
+++ b/src/writer/wgsl/test_helper.h
@@ -31,7 +31,7 @@ namespace wgsl {
 template <typename BASE>
 class TestHelperBase : public BASE, public ast::BuilderWithModule {
  public:
-  TestHelperBase() : td(mod), gen() {}
+  TestHelperBase() : td(mod), gen(mod) {}
 
   ~TestHelperBase() = default;