Pre-parse token list.

This CL updates the parser and lexer to pre-parse the token list and
then the parser works off the list of tokens. This allows the parser to
work with references to the tokens and not have to worry about them
going stale.

For any splittable token a placeholder token is injected after the
token. If the token ends up getting split the placeholder is overwritten
with the new type.

Change-Id: I5a8ccca15d8c14b5027df7dd2734be6753e46fa9
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/97070
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Dan Sinclair <dsinclair@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/src/tint/reader/wgsl/lexer.cc b/src/tint/reader/wgsl/lexer.cc
index 48c3e5d..c9056e0 100644
--- a/src/tint/reader/wgsl/lexer.cc
+++ b/src/tint/reader/wgsl/lexer.cc
@@ -37,6 +37,8 @@
               "tint::reader::wgsl requires the size of a std::string element "
               "to be a single byte");
 
+static constexpr size_t kDefaultListSize = 512;
+
 bool read_blankspace(std::string_view str, size_t i, bool* is_blankspace, size_t* blankspace_size) {
     // See https://www.w3.org/TR/WGSL/#blankspace
 
@@ -88,6 +90,27 @@
 
 Lexer::~Lexer() = default;
 
+std::vector<Token> Lexer::Lex() {
+    std::vector<Token> tokens;
+    tokens.reserve(kDefaultListSize);
+    while (true) {
+        tokens.emplace_back(next());
+
+        // If the token can be split, we insert a placeholder element into
+        // the stream to hold the split character.
+        if (tokens.back().IsSplittable()) {
+            auto src = tokens.back().source();
+            src.range.begin.column++;
+            tokens.emplace_back(Token(Token::Type::kPlaceholder, src));
+        }
+
+        if (tokens.back().IsEof() || tokens.back().IsError()) {
+            break;
+        }
+    }
+    return tokens;
+}
+
 const std::string_view Lexer::line() const {
     if (file_->content.lines.size() == 0) {
         static const char* empty_string = "";
diff --git a/src/tint/reader/wgsl/lexer.h b/src/tint/reader/wgsl/lexer.h
index d51c26c..8e0306b 100644
--- a/src/tint/reader/wgsl/lexer.h
+++ b/src/tint/reader/wgsl/lexer.h
@@ -16,6 +16,7 @@
 #define SRC_TINT_READER_WGSL_LEXER_H_
 
 #include <string>
+#include <vector>
 
 #include "src/tint/reader/wgsl/token.h"
 
@@ -29,11 +30,14 @@
     explicit Lexer(const Source::File* file);
     ~Lexer();
 
+    /// @return the token list.
+    std::vector<Token> Lex();
+
+  private:
     /// Returns the next token in the input stream.
     /// @return Token
     Token next();
 
-  private:
     /// Advances past blankspace and comments, if present at the current position.
     /// @returns error token, EOF, or uninitialized
     Token skip_blankspace_and_comments();
diff --git a/src/tint/reader/wgsl/lexer_test.cc b/src/tint/reader/wgsl/lexer_test.cc
index 799f9c2..f82045a 100644
--- a/src/tint/reader/wgsl/lexer_test.cc
+++ b/src/tint/reader/wgsl/lexer_test.cc
@@ -46,24 +46,33 @@
 TEST_F(LexerTest, Empty) {
     Source::File file("", "");
     Lexer l(&file);
-    auto t = l.next();
-    EXPECT_TRUE(t.IsEof());
+
+    auto list = l.Lex();
+    ASSERT_EQ(1u, list.size());
+    EXPECT_TRUE(list[0].IsEof());
 }
 
 TEST_F(LexerTest, Skips_Blankspace_Basic) {
     Source::File file("", "\t\r\n\t    ident\t\n\t  \r ");
     Lexer l(&file);
 
-    auto t = l.next();
-    EXPECT_TRUE(t.IsIdentifier());
-    EXPECT_EQ(t.source().range.begin.line, 2u);
-    EXPECT_EQ(t.source().range.begin.column, 6u);
-    EXPECT_EQ(t.source().range.end.line, 2u);
-    EXPECT_EQ(t.source().range.end.column, 11u);
-    EXPECT_EQ(t.to_str(), "ident");
+    auto list = l.Lex();
+    ASSERT_EQ(2u, list.size());
 
-    t = l.next();
-    EXPECT_TRUE(t.IsEof());
+    {
+        auto& t = list[0];
+        EXPECT_TRUE(t.IsIdentifier());
+        EXPECT_EQ(t.source().range.begin.line, 2u);
+        EXPECT_EQ(t.source().range.begin.column, 6u);
+        EXPECT_EQ(t.source().range.end.line, 2u);
+        EXPECT_EQ(t.source().range.end.column, 11u);
+        EXPECT_EQ(t.to_str(), "ident");
+    }
+
+    {
+        auto& t = list[1];
+        EXPECT_TRUE(t.IsEof());
+    }
 }
 
 TEST_F(LexerTest, Skips_Blankspace_Exotic) {
@@ -73,16 +82,23 @@
                       kVTab kFF kNL kLS kPS kL2R kR2L);
     Lexer l(&file);
 
-    auto t = l.next();
-    EXPECT_TRUE(t.IsIdentifier());
-    EXPECT_EQ(t.source().range.begin.line, 6u);
-    EXPECT_EQ(t.source().range.begin.column, 7u);
-    EXPECT_EQ(t.source().range.end.line, 6u);
-    EXPECT_EQ(t.source().range.end.column, 12u);
-    EXPECT_EQ(t.to_str(), "ident");
+    auto list = l.Lex();
+    ASSERT_EQ(2u, list.size());
 
-    t = l.next();
-    EXPECT_TRUE(t.IsEof());
+    {
+        auto& t = list[0];
+        EXPECT_TRUE(t.IsIdentifier());
+        EXPECT_EQ(t.source().range.begin.line, 6u);
+        EXPECT_EQ(t.source().range.begin.column, 7u);
+        EXPECT_EQ(t.source().range.end.line, 6u);
+        EXPECT_EQ(t.source().range.end.column, 12u);
+        EXPECT_EQ(t.to_str(), "ident");
+    }
+
+    {
+        auto& t = list[1];
+        EXPECT_TRUE(t.IsEof());
+    }
 }
 
 TEST_F(LexerTest, Skips_Comments_Line) {
@@ -92,24 +108,33 @@
  ident2)");
     Lexer l(&file);
 
-    auto t = l.next();
-    EXPECT_TRUE(t.IsIdentifier());
-    EXPECT_EQ(t.source().range.begin.line, 2u);
-    EXPECT_EQ(t.source().range.begin.column, 1u);
-    EXPECT_EQ(t.source().range.end.line, 2u);
-    EXPECT_EQ(t.source().range.end.column, 7u);
-    EXPECT_EQ(t.to_str(), "ident1");
+    auto list = l.Lex();
+    ASSERT_EQ(3u, list.size());
 
-    t = l.next();
-    EXPECT_TRUE(t.IsIdentifier());
-    EXPECT_EQ(t.source().range.begin.line, 4u);
-    EXPECT_EQ(t.source().range.begin.column, 2u);
-    EXPECT_EQ(t.source().range.end.line, 4u);
-    EXPECT_EQ(t.source().range.end.column, 8u);
-    EXPECT_EQ(t.to_str(), "ident2");
+    {
+        auto& t = list[0];
+        EXPECT_TRUE(t.IsIdentifier());
+        EXPECT_EQ(t.source().range.begin.line, 2u);
+        EXPECT_EQ(t.source().range.begin.column, 1u);
+        EXPECT_EQ(t.source().range.end.line, 2u);
+        EXPECT_EQ(t.source().range.end.column, 7u);
+        EXPECT_EQ(t.to_str(), "ident1");
+    }
 
-    t = l.next();
-    EXPECT_TRUE(t.IsEof());
+    {
+        auto& t = list[1];
+        EXPECT_TRUE(t.IsIdentifier());
+        EXPECT_EQ(t.source().range.begin.line, 4u);
+        EXPECT_EQ(t.source().range.begin.column, 2u);
+        EXPECT_EQ(t.source().range.end.line, 4u);
+        EXPECT_EQ(t.source().range.end.column, 8u);
+        EXPECT_EQ(t.to_str(), "ident2");
+    }
+
+    {
+        auto& t = list[2];
+        EXPECT_TRUE(t.IsEof());
+    }
 }
 
 TEST_F(LexerTest, Skips_Comments_Unicode) {
@@ -119,24 +144,33 @@
  ident2)");
     Lexer l(&file);
 
-    auto t = l.next();
-    EXPECT_TRUE(t.IsIdentifier());
-    EXPECT_EQ(t.source().range.begin.line, 2u);
-    EXPECT_EQ(t.source().range.begin.column, 1u);
-    EXPECT_EQ(t.source().range.end.line, 2u);
-    EXPECT_EQ(t.source().range.end.column, 7u);
-    EXPECT_EQ(t.to_str(), "ident1");
+    auto list = l.Lex();
+    ASSERT_EQ(3u, list.size());
 
-    t = l.next();
-    EXPECT_TRUE(t.IsIdentifier());
-    EXPECT_EQ(t.source().range.begin.line, 4u);
-    EXPECT_EQ(t.source().range.begin.column, 2u);
-    EXPECT_EQ(t.source().range.end.line, 4u);
-    EXPECT_EQ(t.source().range.end.column, 8u);
-    EXPECT_EQ(t.to_str(), "ident2");
+    {
+        auto& t = list[0];
+        EXPECT_TRUE(t.IsIdentifier());
+        EXPECT_EQ(t.source().range.begin.line, 2u);
+        EXPECT_EQ(t.source().range.begin.column, 1u);
+        EXPECT_EQ(t.source().range.end.line, 2u);
+        EXPECT_EQ(t.source().range.end.column, 7u);
+        EXPECT_EQ(t.to_str(), "ident1");
+    }
 
-    t = l.next();
-    EXPECT_TRUE(t.IsEof());
+    {
+        auto& t = list[1];
+        EXPECT_TRUE(t.IsIdentifier());
+        EXPECT_EQ(t.source().range.begin.line, 4u);
+        EXPECT_EQ(t.source().range.begin.column, 2u);
+        EXPECT_EQ(t.source().range.end.line, 4u);
+        EXPECT_EQ(t.source().range.end.column, 8u);
+        EXPECT_EQ(t.to_str(), "ident2");
+    }
+
+    {
+        auto& t = list[2];
+        EXPECT_TRUE(t.IsEof());
+    }
 }
 
 using LineCommentTerminatorTest = testing::TestWithParam<const char*>;
@@ -150,21 +184,29 @@
     Source::File file("", src);
     Lexer l(&file);
 
-    auto t = l.next();
-    EXPECT_TRUE(t.Is(Token::Type::kConst));
-    EXPECT_EQ(t.source().range.begin.line, 1u);
-    EXPECT_EQ(t.source().range.begin.column, 1u);
-    EXPECT_EQ(t.source().range.end.line, 1u);
-    EXPECT_EQ(t.source().range.end.column, 6u);
-
     auto is_same_line = [](std::string_view v) {
         return v == kSpace || v == kHTab || v == kL2R || v == kR2L;
     };
 
+    auto list = l.Lex();
+    ASSERT_EQ(is_same_line(c) ? 2u : 3u, list.size());
+
+    size_t idx = 0;
+
+    {
+        auto& t = list[idx++];
+        EXPECT_TRUE(t.Is(Token::Type::kConst));
+        EXPECT_EQ(t.source().range.begin.line, 1u);
+        EXPECT_EQ(t.source().range.begin.column, 1u);
+        EXPECT_EQ(t.source().range.end.line, 1u);
+        EXPECT_EQ(t.source().range.end.column, 6u);
+    }
+
     if (!is_same_line(c)) {
         size_t line = is_same_line(c) ? 1u : 2u;
         size_t col = is_same_line(c) ? 25u : 1u;
-        t = l.next();
+
+        auto& t = list[idx++];
         EXPECT_TRUE(t.IsIdentifier());
         EXPECT_EQ(t.source().range.begin.line, line);
         EXPECT_EQ(t.source().range.begin.column, col);
@@ -173,8 +215,10 @@
         EXPECT_EQ(t.to_str(), "ident");
     }
 
-    t = l.next();
-    EXPECT_TRUE(t.IsEof());
+    {
+        auto& t = list[idx];
+        EXPECT_TRUE(t.IsEof());
+    }
 }
 INSTANTIATE_TEST_SUITE_P(LexerTest,
                          LineCommentTerminatorTest,
@@ -198,16 +242,23 @@
 text */ident)");
     Lexer l(&file);
 
-    auto t = l.next();
-    EXPECT_TRUE(t.IsIdentifier());
-    EXPECT_EQ(t.source().range.begin.line, 2u);
-    EXPECT_EQ(t.source().range.begin.column, 8u);
-    EXPECT_EQ(t.source().range.end.line, 2u);
-    EXPECT_EQ(t.source().range.end.column, 13u);
-    EXPECT_EQ(t.to_str(), "ident");
+    auto list = l.Lex();
+    ASSERT_EQ(2u, list.size());
 
-    t = l.next();
-    EXPECT_TRUE(t.IsEof());
+    {
+        auto& t = list[0];
+        EXPECT_TRUE(t.IsIdentifier());
+        EXPECT_EQ(t.source().range.begin.line, 2u);
+        EXPECT_EQ(t.source().range.begin.column, 8u);
+        EXPECT_EQ(t.source().range.end.line, 2u);
+        EXPECT_EQ(t.source().range.end.column, 13u);
+        EXPECT_EQ(t.to_str(), "ident");
+    }
+
+    {
+        auto& t = list[1];
+        EXPECT_TRUE(t.IsEof());
+    }
 }
 
 TEST_F(LexerTest, Skips_Comments_Block_Nested) {
@@ -216,16 +267,23 @@
 /////**/ */*/ident)");
     Lexer l(&file);
 
-    auto t = l.next();
-    EXPECT_TRUE(t.IsIdentifier());
-    EXPECT_EQ(t.source().range.begin.line, 3u);
-    EXPECT_EQ(t.source().range.begin.column, 14u);
-    EXPECT_EQ(t.source().range.end.line, 3u);
-    EXPECT_EQ(t.source().range.end.column, 19u);
-    EXPECT_EQ(t.to_str(), "ident");
+    auto list = l.Lex();
+    ASSERT_EQ(2u, list.size());
 
-    t = l.next();
-    EXPECT_TRUE(t.IsEof());
+    {
+        auto& t = list[0];
+        EXPECT_TRUE(t.IsIdentifier());
+        EXPECT_EQ(t.source().range.begin.line, 3u);
+        EXPECT_EQ(t.source().range.begin.column, 14u);
+        EXPECT_EQ(t.source().range.end.line, 3u);
+        EXPECT_EQ(t.source().range.end.column, 19u);
+        EXPECT_EQ(t.to_str(), "ident");
+    }
+
+    {
+        auto& t = list[1];
+        EXPECT_TRUE(t.IsEof());
+    }
 }
 
 TEST_F(LexerTest, Skips_Comments_Block_Unterminated) {
@@ -237,7 +295,10 @@
 abcd)");
     Lexer l(&file);
 
-    auto t = l.next();
+    auto list = l.Lex();
+    ASSERT_EQ(1u, list.size());
+
+    auto& t = list[0];
     ASSERT_TRUE(t.Is(Token::Type::kError));
     EXPECT_EQ(t.to_str(), "unterminated block comment");
     EXPECT_EQ(t.source().range.begin.line, 2u);
@@ -250,7 +311,10 @@
     Source::File file("", std::string{' ', 0, ' '});
     Lexer l(&file);
 
-    auto t = l.next();
+    auto list = l.Lex();
+    ASSERT_EQ(1u, list.size());
+
+    auto& t = list[0];
     EXPECT_TRUE(t.IsError());
     EXPECT_EQ(t.source().range.begin.line, 1u);
     EXPECT_EQ(t.source().range.begin.column, 2u);
@@ -263,7 +327,10 @@
     Source::File file("", std::string{'/', '/', ' ', 0, ' '});
     Lexer l(&file);
 
-    auto t = l.next();
+    auto list = l.Lex();
+    ASSERT_EQ(1u, list.size());
+
+    auto& t = list[0];
     EXPECT_TRUE(t.IsError());
     EXPECT_EQ(t.source().range.begin.line, 1u);
     EXPECT_EQ(t.source().range.begin.column, 4u);
@@ -276,7 +343,10 @@
     Source::File file("", std::string{'/', '*', ' ', 0, '*', '/'});
     Lexer l(&file);
 
-    auto t = l.next();
+    auto list = l.Lex();
+    ASSERT_EQ(1u, list.size());
+
+    auto& t = list[0];
     EXPECT_TRUE(t.IsError());
     EXPECT_EQ(t.source().range.begin.line, 1u);
     EXPECT_EQ(t.source().range.begin.column, 4u);
@@ -292,16 +362,24 @@
     Source::File file("", std::string{'a', 0, 'c'});
     Lexer l(&file);
 
-    auto t = l.next();
-    EXPECT_TRUE(t.IsIdentifier());
-    EXPECT_EQ(t.to_str(), "a");
-    t = l.next();
-    EXPECT_TRUE(t.IsError());
-    EXPECT_EQ(t.source().range.begin.line, 1u);
-    EXPECT_EQ(t.source().range.begin.column, 2u);
-    EXPECT_EQ(t.source().range.end.line, 1u);
-    EXPECT_EQ(t.source().range.end.column, 2u);
-    EXPECT_EQ(t.to_str(), "null character found");
+    auto list = l.Lex();
+    ASSERT_EQ(2u, list.size());
+
+    {
+        auto& t = list[0];
+        EXPECT_TRUE(t.IsIdentifier());
+        EXPECT_EQ(t.to_str(), "a");
+    }
+
+    {
+        auto& t = list[1];
+        EXPECT_TRUE(t.IsError());
+        EXPECT_EQ(t.source().range.begin.line, 1u);
+        EXPECT_EQ(t.source().range.begin.column, 2u);
+        EXPECT_EQ(t.source().range.end.line, 1u);
+        EXPECT_EQ(t.source().range.end.column, 2u);
+        EXPECT_EQ(t.to_str(), "null character found");
+    }
 }
 
 struct FloatData {
@@ -318,22 +396,29 @@
     Source::File file("", params.input);
     Lexer l(&file);
 
-    auto t = l.next();
-    if (std::string(params.input).back() == 'f') {
-        EXPECT_TRUE(t.Is(Token::Type::kFloatLiteral_F));
-    } else if (std::string(params.input).back() == 'h') {
-        EXPECT_TRUE(t.Is(Token::Type::kFloatLiteral_H));
-    } else {
-        EXPECT_TRUE(t.Is(Token::Type::kFloatLiteral));
-    }
-    EXPECT_EQ(t.to_f64(), params.result);
-    EXPECT_EQ(t.source().range.begin.line, 1u);
-    EXPECT_EQ(t.source().range.begin.column, 1u);
-    EXPECT_EQ(t.source().range.end.line, 1u);
-    EXPECT_EQ(t.source().range.end.column, 1u + strlen(params.input));
+    auto list = l.Lex();
+    ASSERT_EQ(2u, list.size());
 
-    t = l.next();
-    EXPECT_TRUE(t.IsEof());
+    {
+        auto& t = list[0];
+        if (std::string(params.input).back() == 'f') {
+            EXPECT_TRUE(t.Is(Token::Type::kFloatLiteral_F));
+        } else if (std::string(params.input).back() == 'h') {
+            EXPECT_TRUE(t.Is(Token::Type::kFloatLiteral_H));
+        } else {
+            EXPECT_TRUE(t.Is(Token::Type::kFloatLiteral));
+        }
+        EXPECT_EQ(t.to_f64(), params.result);
+        EXPECT_EQ(t.source().range.begin.line, 1u);
+        EXPECT_EQ(t.source().range.begin.column, 1u);
+        EXPECT_EQ(t.source().range.end.line, 1u);
+        EXPECT_EQ(t.source().range.end.column, 1u + strlen(params.input));
+    }
+
+    {
+        auto& t = list[1];
+        EXPECT_TRUE(t.IsEof());
+    }
 }
 INSTANTIATE_TEST_SUITE_P(LexerTest,
                          FloatTest,
@@ -437,7 +522,10 @@
     Source::File file("", GetParam());
     Lexer l(&file);
 
-    auto t = l.next();
+    auto list = l.Lex();
+    ASSERT_FALSE(list.empty());
+
+    auto& t = list[0];
     EXPECT_FALSE(t.Is(Token::Type::kFloatLiteral) || t.Is(Token::Type::kFloatLiteral_F) ||
                  t.Is(Token::Type::kFloatLiteral_H));
 }
@@ -478,13 +566,23 @@
     Source::File file("", GetParam());
     Lexer l(&file);
 
-    auto t = l.next();
-    EXPECT_TRUE(t.IsIdentifier());
-    EXPECT_EQ(t.source().range.begin.line, 1u);
-    EXPECT_EQ(t.source().range.begin.column, 1u);
-    EXPECT_EQ(t.source().range.end.line, 1u);
-    EXPECT_EQ(t.source().range.end.column, 1u + strlen(GetParam()));
-    EXPECT_EQ(t.to_str(), GetParam());
+    auto list = l.Lex();
+    ASSERT_EQ(2u, list.size());
+
+    {
+        auto& t = list[0];
+        EXPECT_TRUE(t.IsIdentifier());
+        EXPECT_EQ(t.source().range.begin.line, 1u);
+        EXPECT_EQ(t.source().range.begin.column, 1u);
+        EXPECT_EQ(t.source().range.end.line, 1u);
+        EXPECT_EQ(t.source().range.end.column, 1u + strlen(GetParam()));
+        EXPECT_EQ(t.to_str(), GetParam());
+    }
+
+    {
+        auto& t = list[1];
+        EXPECT_TRUE(t.IsEof());
+    }
 }
 INSTANTIATE_TEST_SUITE_P(LexerTest,
                          AsciiIdentifierTest,
@@ -510,13 +608,23 @@
     Source::File file("", GetParam().utf8);
     Lexer l(&file);
 
-    auto t = l.next();
-    EXPECT_TRUE(t.IsIdentifier());
-    EXPECT_EQ(t.source().range.begin.line, 1u);
-    EXPECT_EQ(t.source().range.begin.column, 1u);
-    EXPECT_EQ(t.source().range.end.line, 1u);
-    EXPECT_EQ(t.source().range.end.column, 1u + GetParam().count);
-    EXPECT_EQ(t.to_str(), GetParam().utf8);
+    auto list = l.Lex();
+    ASSERT_EQ(2u, list.size());
+
+    {
+        auto& t = list[0];
+        EXPECT_TRUE(t.IsIdentifier());
+        EXPECT_EQ(t.source().range.begin.line, 1u);
+        EXPECT_EQ(t.source().range.begin.column, 1u);
+        EXPECT_EQ(t.source().range.end.line, 1u);
+        EXPECT_EQ(t.source().range.end.column, 1u + GetParam().count);
+        EXPECT_EQ(t.to_str(), GetParam().utf8);
+    }
+
+    {
+        auto& t = list[1];
+        EXPECT_TRUE(t.IsEof());
+    }
 }
 INSTANTIATE_TEST_SUITE_P(
     LexerTest,
@@ -554,7 +662,10 @@
     Source::File file("", GetParam());
     Lexer l(&file);
 
-    auto t = l.next();
+    auto list = l.Lex();
+    ASSERT_FALSE(list.empty());
+
+    auto& t = list[0];
     EXPECT_TRUE(t.IsError());
     EXPECT_EQ(t.source().range.begin.line, 1u);
     EXPECT_EQ(t.source().range.begin.column, 1u);
@@ -602,7 +713,10 @@
     Source::File file("", "_");
     Lexer l(&file);
 
-    auto t = l.next();
+    auto list = l.Lex();
+    ASSERT_FALSE(list.empty());
+
+    auto& t = list[0];
     EXPECT_FALSE(t.IsIdentifier());
 }
 
@@ -610,7 +724,10 @@
     Source::File file("", "__test");
     Lexer l(&file);
 
-    auto t = l.next();
+    auto list = l.Lex();
+    ASSERT_FALSE(list.empty());
+
+    auto& t = list[0];
     EXPECT_FALSE(t.IsIdentifier());
 }
 
@@ -618,7 +735,10 @@
     Source::File file("", "01test");
     Lexer l(&file);
 
-    auto t = l.next();
+    auto list = l.Lex();
+    EXPECT_FALSE(list.empty());
+
+    auto& t = list[0];
     EXPECT_FALSE(t.IsIdentifier());
 }
 
@@ -641,7 +761,12 @@
     auto params = std::get<1>(GetParam());
     Source::File file("", params.input);
 
-    auto t = Lexer(&file).next();
+    Lexer l(&file);
+
+    auto list = l.Lex();
+    ASSERT_FALSE(list.empty());
+
+    auto& t = list[0];
     switch (suffix) {
         case 'i':
             EXPECT_TRUE(t.Is(Token::Type::kIntLiteral_I));
@@ -770,7 +895,12 @@
     auto type = std::get<0>(GetParam());
     auto source = std::get<1>(GetParam());
     Source::File file("", source);
-    auto t = Lexer(&file).next();
+
+    Lexer l(&file);
+    auto list = l.Lex();
+    ASSERT_FALSE(list.empty());
+
+    auto& t = list[0];
     EXPECT_TRUE(t.Is(Token::Type::kError));
     auto expect = "value cannot be represented as '" + std::string(type) + "'";
     EXPECT_EQ(t.to_str(), expect);
@@ -800,7 +930,12 @@
 using ParseIntegerTest_LeadingZeros = testing::TestWithParam<const char*>;
 TEST_P(ParseIntegerTest_LeadingZeros, Parse) {
     Source::File file("", GetParam());
-    auto t = Lexer(&file).next();
+
+    Lexer l(&file);
+    auto list = l.Lex();
+    ASSERT_FALSE(list.empty());
+
+    auto& t = list[0];
     EXPECT_TRUE(t.Is(Token::Type::kError));
     EXPECT_EQ(t.to_str(), "integer literal cannot have leading 0s");
 }
@@ -815,7 +950,12 @@
 using ParseIntegerTest_NoSignificantDigits = testing::TestWithParam<const char*>;
 TEST_P(ParseIntegerTest_NoSignificantDigits, Parse) {
     Source::File file("", GetParam());
-    auto t = Lexer(&file).next();
+
+    Lexer l(&file);
+    auto list = l.Lex();
+    ASSERT_FALSE(list.empty());
+
+    auto& t = list[0];
     EXPECT_TRUE(t.Is(Token::Type::kError));
     EXPECT_EQ(t.to_str(), "integer or float hex literal has no significant digits");
 }
@@ -849,15 +989,22 @@
     Source::File file("", params.input);
     Lexer l(&file);
 
-    auto t = l.next();
-    EXPECT_TRUE(t.Is(params.type));
-    EXPECT_EQ(t.source().range.begin.line, 1u);
-    EXPECT_EQ(t.source().range.begin.column, 1u);
-    EXPECT_EQ(t.source().range.end.line, 1u);
-    EXPECT_EQ(t.source().range.end.column, 1u + strlen(params.input));
+    auto list = l.Lex();
+    ASSERT_GE(list.size(), 2u);
 
-    t = l.next();
-    EXPECT_EQ(t.source().range.begin.column, 1 + std::string(params.input).size());
+    {
+        auto& t = list[0];
+        EXPECT_TRUE(t.Is(params.type));
+        EXPECT_EQ(t.source().range.begin.line, 1u);
+        EXPECT_EQ(t.source().range.begin.column, 1u);
+        EXPECT_EQ(t.source().range.end.line, 1u);
+        EXPECT_EQ(t.source().range.end.column, 1u + strlen(params.input));
+    }
+
+    {
+        auto& t = list[1];
+        EXPECT_EQ(t.source().range.begin.column, 1 + std::string(params.input).size());
+    }
 }
 INSTANTIATE_TEST_SUITE_P(LexerTest,
                          PunctuationTest,
@@ -876,8 +1023,6 @@
                                          TokenData{"=", Token::Type::kEqual},
                                          TokenData{"==", Token::Type::kEqualEqual},
                                          TokenData{">", Token::Type::kGreaterThan},
-                                         TokenData{">=", Token::Type::kGreaterThanEqual},
-                                         TokenData{">>", Token::Type::kShiftRight},
                                          TokenData{"<", Token::Type::kLessThan},
                                          TokenData{"<=", Token::Type::kLessThanEqual},
                                          TokenData{"<<", Token::Type::kShiftLeft},
@@ -906,21 +1051,60 @@
                                          TokenData{"|=", Token::Type::kOrEqual},
                                          TokenData{"^=", Token::Type::kXorEqual}));
 
+using SplittablePunctuationTest = testing::TestWithParam<TokenData>;
+TEST_P(SplittablePunctuationTest, Parses) {
+    auto params = GetParam();
+    Source::File file("", params.input);
+    Lexer l(&file);
+
+    auto list = l.Lex();
+    ASSERT_GE(list.size(), 3u);
+
+    {
+        auto& t = list[0];
+        EXPECT_TRUE(t.Is(params.type));
+        EXPECT_EQ(t.source().range.begin.line, 1u);
+        EXPECT_EQ(t.source().range.begin.column, 1u);
+        EXPECT_EQ(t.source().range.end.line, 1u);
+        EXPECT_EQ(t.source().range.end.column, 1u + strlen(params.input));
+    }
+
+    {
+        auto& t = list[1];
+        EXPECT_TRUE(t.Is(Token::Type::kPlaceholder));
+        EXPECT_EQ(t.source().range.begin.line, 1u);
+        EXPECT_EQ(t.source().range.begin.column, 2u);
+        EXPECT_EQ(t.source().range.end.line, 1u);
+        EXPECT_EQ(t.source().range.end.column, 1u + strlen(params.input));
+    }
+
+    {
+        auto& t = list[2];
+        EXPECT_EQ(t.source().range.begin.column, 1 + std::string(params.input).size());
+    }
+}
+INSTANTIATE_TEST_SUITE_P(LexerTest,
+                         SplittablePunctuationTest,
+                         testing::Values(TokenData{">=", Token::Type::kGreaterThanEqual},
+                                         TokenData{">>", Token::Type::kShiftRight}));
+
 using KeywordTest = testing::TestWithParam<TokenData>;
 TEST_P(KeywordTest, Parses) {
     auto params = GetParam();
     Source::File file("", params.input);
     Lexer l(&file);
 
-    auto t = l.next();
+    auto list = l.Lex();
+    ASSERT_GE(list.size(), 2u);
+
+    auto& t = list[0];
     EXPECT_TRUE(t.Is(params.type)) << params.input;
     EXPECT_EQ(t.source().range.begin.line, 1u);
     EXPECT_EQ(t.source().range.begin.column, 1u);
     EXPECT_EQ(t.source().range.end.line, 1u);
     EXPECT_EQ(t.source().range.end.column, 1u + strlen(params.input));
 
-    t = l.next();
-    EXPECT_EQ(t.source().range.begin.column, 1 + std::string(params.input).size());
+    EXPECT_EQ(list[1].source().range.begin.column, 1 + std::string(params.input).size());
 }
 INSTANTIATE_TEST_SUITE_P(
     LexerTest,
diff --git a/src/tint/reader/wgsl/parser_impl.cc b/src/tint/reader/wgsl/parser_impl.cc
index 87dc78e..e08f6cb 100644
--- a/src/tint/reader/wgsl/parser_impl.cc
+++ b/src/tint/reader/wgsl/parser_impl.cc
@@ -123,7 +123,7 @@
 const char kWorkgroupSizeAttribute[] = "workgroup_size";
 
 // https://gpuweb.github.io/gpuweb/wgsl.html#reserved-keywords
-bool is_reserved(Token t) {
+bool is_reserved(const Token& t) {
     return t == "asm" || t == "bf16" || t == "do" || t == "enum" || t == "f64" || t == "handle" ||
            t == "i8" || t == "i16" || t == "i64" || t == "mat" || t == "premerge" ||
            t == "regardless" || t == "typedef" || t == "u8" || t == "u16" || t == "u64" ||
@@ -181,7 +181,7 @@
     /// Implicit conversion to Source that returns the combined source from start
     /// to the current last token's source.
     operator Source() const {
-        Source end = parser_->last_token().source().End();
+        Source end = parser_->last_source().End();
         if (end < start_) {
             end = start_;
         }
@@ -241,7 +241,7 @@
 
 ParserImpl::VarDeclInfo::~VarDeclInfo() = default;
 
-ParserImpl::ParserImpl(Source::File const* file) : lexer_(std::make_unique<Lexer>(file)) {}
+ParserImpl::ParserImpl(Source::File const* file) : file_(file) {}
 
 ParserImpl::~ParserImpl() = default;
 
@@ -274,33 +274,67 @@
                                        "use of deprecated language feature: " + msg, source);
 }
 
-Token ParserImpl::next() {
-    if (!token_queue_.empty()) {
-        auto t = token_queue_.front();
-        token_queue_.pop_front();
-        last_token_ = t;
-        return last_token_;
+const Token& ParserImpl::next() {
+    if (!tokens_[next_token_idx_].IsEof() && !tokens_[next_token_idx_].IsError()) {
+        // Skip over any placeholder elements
+        while (true) {
+            if (!tokens_[next_token_idx_].IsPlaceholder()) {
+                break;
+            }
+            next_token_idx_++;
+        }
     }
-    last_token_ = lexer_->next();
-    return last_token_;
+    last_source_ = tokens_[next_token_idx_].source();
+    return tokens_[next_token_idx_++];
 }
 
-Token ParserImpl::peek(size_t idx) {
-    while (token_queue_.size() < (idx + 1)) {
-        token_queue_.push_back(lexer_->next());
+const Token& ParserImpl::peek(size_t idx) {
+    if (next_token_idx_ + idx >= tokens_.size()) {
+        return tokens_[tokens_.size() - 1];
     }
-    return token_queue_[idx];
+
+    // Skip over any placeholder elements
+    while (true) {
+        if (!tokens_[next_token_idx_ + idx].IsPlaceholder()) {
+            break;
+        }
+        idx++;
+    }
+
+    return tokens_[next_token_idx_ + idx];
 }
 
 bool ParserImpl::peek_is(Token::Type tok, size_t idx) {
     return peek(idx).Is(tok);
 }
 
-Token ParserImpl::last_token() const {
-    return last_token_;
+void ParserImpl::split_token(Token::Type lhs, Token::Type rhs) {
+    if (next_token_idx_ == 0) {
+        TINT_ICE(Reader, builder_.Diagnostics())
+            << "attempt to update placeholder at beginning of tokens";
+    }
+    if (next_token_idx_ >= tokens_.size()) {
+        TINT_ICE(Reader, builder_.Diagnostics())
+            << "attempt to update placeholder past end of tokens";
+    }
+    if (!tokens_[next_token_idx_].IsPlaceholder()) {
+        TINT_ICE(Reader, builder_.Diagnostics()) << "attempt to update non-placeholder token";
+    }
+    tokens_[next_token_idx_ - 1].SetType(lhs);
+    tokens_[next_token_idx_].SetType(rhs);
+}
+
+Source ParserImpl::last_source() const {
+    return last_source_;
+}
+
+void ParserImpl::InitializeLex() {
+    Lexer l{file_};
+    tokens_ = l.Lex();
 }
 
 bool ParserImpl::Parse() {
+    InitializeLex();
     translation_unit();
     return !has_error();
 }
@@ -310,7 +344,7 @@
 void ParserImpl::translation_unit() {
     bool after_global_decl = false;
     while (continue_parsing()) {
-        auto p = peek();
+        auto& p = peek();
         if (p.IsEof()) {
             break;
         }
@@ -353,7 +387,7 @@
 
         // Match the extension name.
         Expect<std::string> name = {""};
-        auto t = peek();
+        auto& t = peek();
         if (t.IsIdentifier()) {
             synchronized_ = true;
             next();
@@ -508,7 +542,7 @@
     }
 
     // We have a statement outside of a function?
-    auto t = peek();
+    auto& t = peek();
     auto stat = without_error([&] { return statement(); });
     if (stat.matched) {
         // Attempt to jump to the next '}' - the function might have just been
@@ -882,7 +916,7 @@
 //  | 'rgba32sint'
 //  | 'rgba32float'
 Expect<ast::TexelFormat> ParserImpl::expect_texel_format(std::string_view use) {
-    auto t = next();
+    auto& t = next();
     if (t == "rgba8unorm") {
         return ast::TexelFormat::kRgba8Unorm;
     }
@@ -951,7 +985,7 @@
         return Failure::kErrored;
     }
 
-    auto t = peek();
+    auto& t = peek();
     auto type = type_decl();
     if (type.errored) {
         return Failure::kErrored;
@@ -1021,7 +1055,7 @@
         return Failure::kNoMatch;
     }
 
-    auto t = next();
+    auto& t = next();
     const char* use = "type alias";
 
     auto name = expect_ident(use);
@@ -1069,7 +1103,7 @@
 //   | MAT4x4 LESS_THAN type_decl GREATER_THAN
 //   | texture_samplers
 Maybe<const ast::Type*> ParserImpl::type_decl() {
-    auto t = peek();
+    auto& t = peek();
     Source source;
     if (match(Token::Type::kIdentifier, &source)) {
         return builder_.create<ast::TypeName>(source, builder_.Symbols().Register(t.to_str()));
@@ -1139,7 +1173,7 @@
     return type.value;
 }
 
-Expect<const ast::Type*> ParserImpl::expect_type_decl_pointer(Token t) {
+Expect<const ast::Type*> ParserImpl::expect_type_decl_pointer(const Token& t) {
     const char* use = "ptr declaration";
 
     auto storage_class = ast::StorageClass::kNone;
@@ -1180,7 +1214,7 @@
                                access);
 }
 
-Expect<const ast::Type*> ParserImpl::expect_type_decl_atomic(Token t) {
+Expect<const ast::Type*> ParserImpl::expect_type_decl_atomic(const Token& t) {
     const char* use = "atomic declaration";
 
     auto subtype = expect_lt_gt_block(use, [&] { return expect_type(use); });
@@ -1191,7 +1225,7 @@
     return builder_.ty.atomic(make_source_range_from(t.source()), subtype.value);
 }
 
-Expect<const ast::Type*> ParserImpl::expect_type_decl_vector(Token t) {
+Expect<const ast::Type*> ParserImpl::expect_type_decl_vector(const Token& t) {
     uint32_t count = 2;
     if (t.Is(Token::Type::kVec3)) {
         count = 3;
@@ -1212,7 +1246,7 @@
     return builder_.ty.vec(make_source_range_from(t.source()), subtype, count);
 }
 
-Expect<const ast::Type*> ParserImpl::expect_type_decl_array(Token t) {
+Expect<const ast::Type*> ParserImpl::expect_type_decl_array(const Token& t) {
     const char* use = "array declaration";
 
     const ast::Expression* size = nullptr;
@@ -1244,7 +1278,7 @@
     return builder_.ty.array(make_source_range_from(t.source()), subtype.value, size);
 }
 
-Expect<const ast::Type*> ParserImpl::expect_type_decl_matrix(Token t) {
+Expect<const ast::Type*> ParserImpl::expect_type_decl_matrix(const Token& t) {
     uint32_t rows = 2;
     uint32_t columns = 2;
     if (t.IsMat3xN()) {
@@ -1305,7 +1339,7 @@
 // struct_decl
 //   : STRUCT IDENT struct_body_decl
 Maybe<const ast::Struct*> ParserImpl::struct_decl() {
-    auto t = peek();
+    auto& t = peek();
     auto source = t.source();
 
     if (!match(Token::Type::kStruct)) {
@@ -1334,7 +1368,7 @@
         bool errored = false;
         while (continue_parsing()) {
             // Check for the end of the list.
-            auto t = peek();
+            auto& t = peek();
             if (!t.IsIdentifier() && !t.Is(Token::Type::kAttr)) {
                 break;
             }
@@ -1481,7 +1515,7 @@
     ast::ParameterList ret;
     while (continue_parsing()) {
         // Check for the end of the list.
-        auto t = peek();
+        auto& t = peek();
         if (!t.IsIdentifier() && !t.Is(Token::Type::kAttr)) {
             break;
         }
@@ -1522,7 +1556,7 @@
 //   | FRAGMENT
 //   | COMPUTE
 Expect<ast::PipelineStage> ParserImpl::expect_pipeline_stage() {
-    auto t = peek();
+    auto& t = peek();
     if (t == kVertexStage) {
         next();  // Consume the peek
         return {ast::PipelineStage::kVertex, t.source()};
@@ -2010,7 +2044,7 @@
         return Failure::kNoMatch;
     }
 
-    auto t = next();
+    auto& t = next();
     auto source = t.source();
 
     ast::CaseSelectorList selector_list;
@@ -2263,8 +2297,8 @@
 // func_call_stmt
 //    : IDENT argument_expression_list
 Maybe<const ast::CallStatement*> ParserImpl::func_call_stmt() {
-    auto t = peek();
-    auto t2 = peek(1);
+    auto& t = peek();
+    auto& t2 = peek(1);
     if (!t.IsIdentifier() || !t2.Is(Token::Type::kParenLeft)) {
         return Failure::kNoMatch;
     }
@@ -2325,7 +2359,7 @@
 //   | paren_expression
 //   | BITCAST LESS_THAN type_decl GREATER_THAN paren_expression
 Maybe<const ast::Expression*> ParserImpl::primary_expression() {
-    auto t = peek();
+    auto& t = peek();
     auto source = t.source();
 
     auto lit = const_literal();
@@ -2491,7 +2525,7 @@
 //   | STAR unary_expression
 //   | AND unary_expression
 Maybe<const ast::Expression*> ParserImpl::unary_expression() {
-    auto t = peek();
+    auto& t = peek();
 
     if (match(Token::Type::kPlusPlus) || match(Token::Type::kMinusMinus)) {
         add_error(t.source(),
@@ -2556,7 +2590,7 @@
             return lhs;
         }
 
-        auto t = next();
+        auto& t = next();
         auto source = t.source();
         auto name = t.to_name();
 
@@ -2603,7 +2637,7 @@
             return lhs;
         }
 
-        auto t = next();
+        auto& t = next();
         auto source = t.source();
 
         auto rhs = multiplicative_expression();
@@ -2651,7 +2685,7 @@
             return lhs;
         }
 
-        auto t = next();
+        auto& t = next();
         auto source = t.source();
         auto rhs = additive_expression();
         if (rhs.errored) {
@@ -2702,7 +2736,7 @@
             return lhs;
         }
 
-        auto t = next();
+        auto& t = next();
         auto source = t.source();
         auto name = t.to_name();
 
@@ -2749,7 +2783,7 @@
             return lhs;
         }
 
-        auto t = next();
+        auto& t = next();
         auto source = t.source();
         auto name = t.to_name();
 
@@ -2790,7 +2824,7 @@
             return lhs;
         }
 
-        auto t = next();
+        auto& t = next();
         auto source = t.source();
 
         auto rhs = equality_expression();
@@ -2903,7 +2937,7 @@
             return lhs;
         }
 
-        auto t = next();
+        auto& t = next();
         auto source = t.source();
 
         auto rhs = inclusive_or_expression();
@@ -3013,7 +3047,7 @@
 // decrement_stmt
 // | lhs_expression MINUS_MINUS
 Maybe<const ast::Statement*> ParserImpl::assignment_stmt() {
-    auto t = peek();
+    auto& t = peek();
     auto source = t.source();
 
     // tint:295 - Test for `ident COLON` - this is invalid grammar, and without
@@ -3076,7 +3110,7 @@
 //   | TRUE
 //   | FALSE
 Maybe<const ast::LiteralExpression*> ParserImpl::const_literal() {
-    auto t = peek();
+    auto& t = peek();
     if (match(Token::Type::kIntLiteral)) {
         return create<ast::IntLiteralExpression>(t.source(), t.to_i64(),
                                                  ast::IntLiteralExpression::Suffix::kNone);
@@ -3141,7 +3175,7 @@
 }
 
 Expect<const ast::Attribute*> ParserImpl::expect_attribute() {
-    auto t = peek();
+    auto& t = peek();
     auto attr = attribute();
     if (attr.errored) {
         return Failure::kErrored;
@@ -3154,7 +3188,7 @@
 
 Maybe<const ast::Attribute*> ParserImpl::attribute() {
     using Result = Maybe<const ast::Attribute*>;
-    auto t = next();
+    auto& t = next();
 
     if (!t.IsIdentifier()) {
         return Failure::kNoMatch;
@@ -3204,7 +3238,7 @@
             ast::InterpolationType type;
             ast::InterpolationSampling sampling = ast::InterpolationSampling::kNone;
 
-            auto type_tok = next();
+            auto& type_tok = next();
             if (type_tok == "perspective") {
                 type = ast::InterpolationType::kPerspective;
             } else if (type_tok == "linear") {
@@ -3217,7 +3251,7 @@
 
             if (match(Token::Type::kComma)) {
                 if (!peek_is(Token::Type::kParenRight)) {
-                    auto sampling_tok = next();
+                    auto& sampling_tok = next();
                     if (sampling_tok == "center") {
                         sampling = ast::InterpolationSampling::kCenter;
                     } else if (sampling_tok == "centroid") {
@@ -3227,7 +3261,6 @@
                     } else {
                         return add_error(sampling_tok, "invalid interpolation sampling");
                     }
-
                     match(Token::Type::kComma);
                 }
             }
@@ -3384,7 +3417,7 @@
 }
 
 bool ParserImpl::match(Token::Type tok, Source* source /*= nullptr*/) {
-    auto t = peek();
+    auto& t = peek();
 
     if (source != nullptr) {
         *source = t.source();
@@ -3398,7 +3431,7 @@
 }
 
 bool ParserImpl::expect(std::string_view use, Token::Type tok) {
-    auto t = peek();
+    auto& t = peek();
     if (t.Is(tok)) {
         next();
         synchronized_ = true;
@@ -3414,9 +3447,9 @@
         auto source = t.source();
         source.range.begin.column++;
         if (t.Is(Token::Type::kShiftRight)) {
-            token_queue_.push_front(Token(Token::Type::kGreaterThan, source));
+            split_token(Token::Type::kGreaterThan, Token::Type::kGreaterThan);
         } else if (t.Is(Token::Type::kGreaterThanEqual)) {
-            token_queue_.push_front(Token(Token::Type::kEqual, source));
+            split_token(Token::Type::kGreaterThan, Token::Type::kEqual);
         }
 
         synchronized_ = true;
@@ -3439,7 +3472,7 @@
 }
 
 Expect<int32_t> ParserImpl::expect_sint(std::string_view use) {
-    auto t = peek();
+    auto& t = peek();
     if (!t.Is(Token::Type::kIntLiteral) && !t.Is(Token::Type::kIntLiteral_I)) {
         return add_error(t.source(), "expected signed integer literal", use);
     }
@@ -3482,7 +3515,7 @@
 }
 
 Expect<std::string> ParserImpl::expect_ident(std::string_view use) {
-    auto t = peek();
+    auto& t = peek();
     if (t.IsIdentifier()) {
         synchronized_ = true;
         next();
@@ -3578,7 +3611,7 @@
     BlockCounters counters;
 
     for (size_t i = 0; i < kMaxResynchronizeLookahead; i++) {
-        auto t = peek(i);
+        auto& t = peek(i);
         if (counters.consume(t) > 0) {
             continue;  // Nested block
         }
diff --git a/src/tint/reader/wgsl/parser_impl.h b/src/tint/reader/wgsl/parser_impl.h
index a9c0423..424fd2e 100644
--- a/src/tint/reader/wgsl/parser_impl.h
+++ b/src/tint/reader/wgsl/parser_impl.h
@@ -15,7 +15,6 @@
 #ifndef SRC_TINT_READER_WGSL_PARSER_IMPL_H_
 #define SRC_TINT_READER_WGSL_PARSER_IMPL_H_
 
-#include <deque>
 #include <memory>
 #include <string>
 #include <string_view>
@@ -300,6 +299,10 @@
     explicit ParserImpl(Source::File const* file);
     ~ParserImpl();
 
+    /// Reads tokens from the source file. This will be called automatically
+    /// by |parse|.
+    void InitializeLex();
+
     /// Run the parser
     /// @returns true if the parse was successful, false otherwise.
     bool Parse();
@@ -330,19 +333,19 @@
     ProgramBuilder& builder() { return builder_; }
 
     /// @returns the next token
-    Token next();
+    const Token& next();
     /// Peeks ahead and returns the token at `idx` ahead of the current position
     /// @param idx the index of the token to return
     /// @returns the token `idx` positions ahead without advancing
-    Token peek(size_t idx = 0);
+    const Token& peek(size_t idx = 0);
     /// Peeks ahead and returns true if the token at `idx` ahead of the current
     /// position is |tok|
     /// @param idx the index of the token to return
     /// @param tok the token to look for
     /// @returns true if the token `idx` positions ahead is |tok|
     bool peek_is(Token::Type tok, size_t idx = 0);
-    /// @returns the last token that was returned by `next()`
-    Token last_token() const;
+    /// @returns the last source location that was returned by `next()`
+    Source last_source() const;
     /// Appends an error at `t` with the message `msg`
     /// @param t the token to associate the error with
     /// @param msg the error message
@@ -812,11 +815,11 @@
     /// Used to ensure that all attributes are consumed.
     bool expect_attributes_consumed(ast::AttributeList& list);
 
-    Expect<const ast::Type*> expect_type_decl_pointer(Token t);
-    Expect<const ast::Type*> expect_type_decl_atomic(Token t);
-    Expect<const ast::Type*> expect_type_decl_vector(Token t);
-    Expect<const ast::Type*> expect_type_decl_array(Token t);
-    Expect<const ast::Type*> expect_type_decl_matrix(Token t);
+    Expect<const ast::Type*> expect_type_decl_pointer(const Token& t);
+    Expect<const ast::Type*> expect_type_decl_atomic(const Token& t);
+    Expect<const ast::Type*> expect_type_decl_vector(const Token& t);
+    Expect<const ast::Type*> expect_type_decl_array(const Token& t);
+    Expect<const ast::Type*> expect_type_decl_matrix(const Token& t);
 
     Expect<const ast::Type*> expect_type(std::string_view use);
 
@@ -824,6 +827,8 @@
     Maybe<const ast::Statement*> for_header_initializer();
     Maybe<const ast::Statement*> for_header_continuing();
 
+    void split_token(Token::Type lhs, Token::Type rhs);
+
     class MultiTokenSource;
     MultiTokenSource make_source_range();
     MultiTokenSource make_source_range_from(const Source& start);
@@ -837,9 +842,10 @@
         return builder_.create<T>(std::forward<ARGS>(args)...);
     }
 
-    std::unique_ptr<Lexer> lexer_;
-    std::deque<Token> token_queue_;
-    Token last_token_;
+    Source::File const* const file_;
+    std::vector<Token> tokens_;
+    size_t next_token_idx_ = 0;
+    Source last_source_;
     bool synchronized_ = true;
     uint32_t parse_depth_ = 0;
     std::vector<Token::Type> sync_tokens_;
diff --git a/src/tint/reader/wgsl/parser_impl_pipeline_stage_test.cc b/src/tint/reader/wgsl/parser_impl_pipeline_stage_test.cc
index 75a1ec2..3c730e2 100644
--- a/src/tint/reader/wgsl/parser_impl_pipeline_stage_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_pipeline_stage_test.cc
@@ -40,7 +40,7 @@
     EXPECT_EQ(stage.source.range.end.line, 1u);
     EXPECT_EQ(stage.source.range.end.column, 1u + params.input.size());
 
-    auto t = p->next();
+    auto& t = p->next();
     EXPECT_TRUE(t.IsEof());
 }
 INSTANTIATE_TEST_SUITE_P(
diff --git a/src/tint/reader/wgsl/parser_impl_storage_class_test.cc b/src/tint/reader/wgsl/parser_impl_storage_class_test.cc
index deb87fc..7144f2c 100644
--- a/src/tint/reader/wgsl/parser_impl_storage_class_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_storage_class_test.cc
@@ -37,7 +37,7 @@
     EXPECT_FALSE(p->has_error());
     EXPECT_EQ(sc.value, params.result);
 
-    auto t = p->next();
+    auto& t = p->next();
     EXPECT_TRUE(t.IsEof());
 }
 INSTANTIATE_TEST_SUITE_P(
diff --git a/src/tint/reader/wgsl/parser_impl_test_helper.h b/src/tint/reader/wgsl/parser_impl_test_helper.h
index f6ee8ea..1b14039 100644
--- a/src/tint/reader/wgsl/parser_impl_test_helper.h
+++ b/src/tint/reader/wgsl/parser_impl_test_helper.h
@@ -38,6 +38,7 @@
     std::unique_ptr<ParserImpl> parser(const std::string& str) {
         auto file = std::make_unique<Source::File>("test.wgsl", str);
         auto impl = std::make_unique<ParserImpl>(file.get());
+        impl->InitializeLex();
         files_.emplace_back(std::move(file));
         return impl;
     }
@@ -60,6 +61,7 @@
     std::unique_ptr<ParserImpl> parser(const std::string& str) {
         auto file = std::make_unique<Source::File>("test.wgsl", str);
         auto impl = std::make_unique<ParserImpl>(file.get());
+        impl->InitializeLex();
         files_.emplace_back(std::move(file));
         return impl;
     }
diff --git a/src/tint/reader/wgsl/parser_impl_variable_decl_test.cc b/src/tint/reader/wgsl/parser_impl_variable_decl_test.cc
index 39f8b24..3ed0948 100644
--- a/src/tint/reader/wgsl/parser_impl_variable_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_variable_decl_test.cc
@@ -68,7 +68,7 @@
     EXPECT_FALSE(v.errored);
     EXPECT_FALSE(p->has_error());
 
-    auto t = p->next();
+    auto& t = p->next();
     ASSERT_TRUE(t.IsIdentifier());
 }
 
diff --git a/src/tint/reader/wgsl/parser_impl_variable_qualifier_test.cc b/src/tint/reader/wgsl/parser_impl_variable_qualifier_test.cc
index aefd2b5..0dc0a1a 100644
--- a/src/tint/reader/wgsl/parser_impl_variable_qualifier_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_variable_qualifier_test.cc
@@ -40,7 +40,7 @@
     EXPECT_EQ(sc->storage_class, params.storage_class);
     EXPECT_EQ(sc->access, params.access);
 
-    auto t = p->next();
+    auto& t = p->next();
     EXPECT_TRUE(t.IsEof());
 }
 INSTANTIATE_TEST_SUITE_P(
@@ -83,7 +83,7 @@
     EXPECT_FALSE(sc.errored);
     EXPECT_FALSE(sc.matched);
 
-    auto t = p->next();
+    auto& t = p->next();
     ASSERT_TRUE(t.Is(Token::Type::kIdentifier));
 }
 
@@ -94,7 +94,7 @@
     EXPECT_FALSE(sc.errored);
     EXPECT_FALSE(sc.matched);
 
-    auto t = p->next();
+    auto& t = p->next();
     ASSERT_TRUE(t.Is(Token::Type::kIdentifier));
 }
 
diff --git a/src/tint/reader/wgsl/token.cc b/src/tint/reader/wgsl/token.cc
index 44255d3..a7b437f 100644
--- a/src/tint/reader/wgsl/token.cc
+++ b/src/tint/reader/wgsl/token.cc
@@ -37,6 +37,8 @@
             return "'i'-suffixed integer literal";
         case Token::Type::kIntLiteral_U:
             return "'u'-suffixed integer literal";
+        case Token::Type::kPlaceholder:
+            return "placeholder";
         case Token::Type::kUninitialized:
             return "uninitialized";
 
@@ -285,13 +287,9 @@
 
 Token::Token(Token&&) = default;
 
-Token::Token(const Token&) = default;
-
 Token::~Token() = default;
 
-Token& Token::operator=(const Token& rhs) = default;
-
-bool Token::operator==(std::string_view ident) {
+bool Token::operator==(std::string_view ident) const {
     if (type_ != Type::kIdentifier) {
         return false;
     }
diff --git a/src/tint/reader/wgsl/token.h b/src/tint/reader/wgsl/token.h
index 473588d..352c18f 100644
--- a/src/tint/reader/wgsl/token.h
+++ b/src/tint/reader/wgsl/token.h
@@ -32,6 +32,8 @@
         kError = -2,
         /// Uninitialized token
         kUninitialized = 0,
+        /// Placeholder token which maybe fillled in later
+        kPlaceholder = 1,
         /// End of input string reached
         kEOF,
 
@@ -310,19 +312,16 @@
     Token(Type type, const Source& source, double val);
     /// Move constructor
     Token(Token&&);
-    /// Copy constructor
-    Token(const Token&);
     ~Token();
 
-    /// Assignment operator
-    /// @param b the token to copy
-    /// @return Token
-    Token& operator=(const Token& b);
-
     /// Equality operator with an identifier
     /// @param ident the identifier string
     /// @return true if this token is an identifier and is equal to ident.
-    bool operator==(std::string_view ident);
+    bool operator==(std::string_view ident) const;
+
+    /// Sets the token to the given type
+    /// @param type the type to set
+    void SetType(Token::Type type) { type_ = type; }
 
     /// Returns true if the token is of the given type
     /// @param t the type to check against.
@@ -331,6 +330,8 @@
 
     /// @returns true if the token is uninitialized
     bool IsUninitialized() const { return type_ == Type::kUninitialized; }
+    /// @returns true if the token is a placeholder
+    bool IsPlaceholder() const { return type_ == Type::kPlaceholder; }
     /// @returns true if the token is EOF
     bool IsEof() const { return type_ == Type::kEOF; }
     /// @returns true if the token is Error
@@ -372,6 +373,11 @@
         return type_ == Type::kVec2 || type_ == Type::kVec3 || type_ == Type::kVec4;
     }
 
+    /// @returns true if the token can be split during parse into component tokens
+    bool IsSplittable() const {
+        return Is(Token::Type::kShiftRight) || Is(Token::Type::kGreaterThanEqual);
+    }
+
     /// @returns the source information for this token
     Source source() const { return source_; }