diff --git a/src/tint/reader/wgsl/lexer.cc b/src/tint/reader/wgsl/lexer.cc
index 48c3e5d80b..c9056e0ab9 100644
--- a/src/tint/reader/wgsl/lexer.cc
+++ b/src/tint/reader/wgsl/lexer.cc
@@ -37,6 +37,8 @@ static_assert(sizeof(decltype(tint::Source::FileContent::data[0])) == sizeof(uin
               "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(const Source::File* file) : file_(file), location_{1, 1} {}
 
 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 d51c26c272..8e0306b6a8 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 @@ class Lexer {
     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 799f9c20b5..f82045ab80 100644
--- a/src/tint/reader/wgsl/lexer_test.cc
+++ b/src/tint/reader/wgsl/lexer_test.cc
@@ -46,24 +46,33 @@ using LexerTest = testing::Test;
 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 @@ TEST_F(LexerTest, Skips_Blankspace_Exotic) {
                       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 @@ ident1 //ends with comment
  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 @@ ident1 //ends with 🙂🙂🙂
  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 @@ TEST_P(LineCommentTerminatorTest, Terminators) {
     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 @@ TEST_P(LineCommentTerminatorTest, Terminators) {
         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 @@ TEST_F(LexerTest, Skips_Comments_Block) {
 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 @@ text // nested line comments are ignored /* more text
 /////**/ */*/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 @@ TEST_F(LexerTest, Skips_Comments_Block_Unterminated) {
 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 @@ TEST_F(LexerTest, Null_InBlankspace_IsError) {
     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 @@ TEST_F(LexerTest, Null_InLineComment_IsError) {
     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 @@ TEST_F(LexerTest, Null_InBlockComment_IsError) {
     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 @@ TEST_F(LexerTest, Null_InIdentifier_IsError) {
     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 @@ TEST_P(FloatTest, Parse) {
     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 @@ TEST_P(FloatTest_Invalid, Handles) {
     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 @@ TEST_P(AsciiIdentifierTest, Parse) {
     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 @@ TEST_P(ValidUnicodeIdentifierTest, Parse) {
     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 @@ TEST_P(InvalidUnicodeIdentifierTest, Parse) {
     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 @@ TEST_F(LexerTest, IdentifierTest_SingleUnderscoreDoesNotMatch) {
     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 @@ TEST_F(LexerTest, IdentifierTest_DoesNotStartWithDoubleUnderscore) {
     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 @@ TEST_F(LexerTest, IdentifierTest_DoesNotStartWithNumber) {
     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 @@ TEST_P(ParseIntegerTest, Parse) {
     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 @@ TEST_P(ParseIntegerTest_CannotBeRepresented, Parse) {
     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 @@ INSTANTIATE_TEST_SUITE_P(u32,
 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 @@ INSTANTIATE_TEST_SUITE_P(LeadingZero,
 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 @@ TEST_P(PunctuationTest, Parses) {
     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 @@ INSTANTIATE_TEST_SUITE_P(LexerTest,
                                          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 @@ INSTANTIATE_TEST_SUITE_P(LexerTest,
                                          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 87dc78e604..e08f6cb347 100644
--- a/src/tint/reader/wgsl/parser_impl.cc
+++ b/src/tint/reader/wgsl/parser_impl.cc
@@ -123,7 +123,7 @@ const char kStageAttribute[] = "stage";
 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 @@ class ParserImpl::MultiTokenSource {
     /// 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(Source source_in,
 
 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 @@ void ParserImpl::deprecated(const Source& source, const std::string& msg) {
                                        "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 @@ bool ParserImpl::Parse() {
 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 @@ Maybe<bool> ParserImpl::enable_directive() {
 
         // Match the extension name.
         Expect<std::string> name = {""};
-        auto t = peek();
+        auto& t = peek();
         if (t.IsIdentifier()) {
             synchronized_ = true;
             next();
@@ -508,7 +542,7 @@ Maybe<bool> ParserImpl::global_decl() {
     }
 
     // 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 @@ Maybe<const ast::Type*> ParserImpl::depth_texture() {
 //  | '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 @@ Expect<ParserImpl::TypedIdentifier> ParserImpl::expect_variable_ident_decl(std::
         return Failure::kErrored;
     }
 
-    auto t = peek();
+    auto& t = peek();
     auto type = type_decl();
     if (type.errored) {
         return Failure::kErrored;
@@ -1021,7 +1055,7 @@ Maybe<const ast::Alias*> ParserImpl::type_alias() {
         return Failure::kNoMatch;
     }
 
-    auto t = next();
+    auto& t = next();
     const char* use = "type alias";
 
     auto name = expect_ident(use);
@@ -1069,7 +1103,7 @@ Maybe<const ast::Alias*> ParserImpl::type_alias() {
 //   | 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 @@ Expect<const ast::Type*> ParserImpl::expect_type(std::string_view use) {
     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 @@ Expect<const ast::Type*> ParserImpl::expect_type_decl_pointer(Token t) {
                                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 @@ Expect<const ast::Type*> ParserImpl::expect_type_decl_atomic(Token t) {
     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 @@ Expect<const ast::Type*> ParserImpl::expect_type_decl_vector(Token t) {
     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 @@ Expect<const ast::Type*> ParserImpl::expect_type_decl_array(Token t) {
     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 @@ Expect<ast::StorageClass> ParserImpl::expect_storage_class(std::string_view use)
 // 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 @@ Expect<ast::StructMemberList> ParserImpl::expect_struct_body_decl() {
         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 @@ Expect<ast::ParameterList> ParserImpl::expect_param_list() {
     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 @@ Expect<ast::Parameter*> ParserImpl::expect_param() {
 //   | 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 @@ Maybe<const ast::CaseStatement*> ParserImpl::switch_body() {
         return Failure::kNoMatch;
     }
 
-    auto t = next();
+    auto& t = next();
     auto source = t.source();
 
     ast::CaseSelectorList selector_list;
@@ -2263,8 +2297,8 @@ Maybe<const ast::WhileStatement*> ParserImpl::while_stmt() {
 // 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 @@ Maybe<const ast::BlockStatement*> ParserImpl::continuing_stmt() {
 //   | 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 @@ Expect<ast::ExpressionList> ParserImpl::expect_argument_expression_list(std::str
 //   | 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 @@ Expect<const ast::Expression*> ParserImpl::expect_multiplicative_expr(const ast:
             return lhs;
         }
 
-        auto t = next();
+        auto& t = next();
         auto source = t.source();
         auto name = t.to_name();
 
@@ -2603,7 +2637,7 @@ Expect<const ast::Expression*> ParserImpl::expect_additive_expr(const ast::Expre
             return lhs;
         }
 
-        auto t = next();
+        auto& t = next();
         auto source = t.source();
 
         auto rhs = multiplicative_expression();
@@ -2651,7 +2685,7 @@ Expect<const ast::Expression*> ParserImpl::expect_shift_expr(const ast::Expressi
             return lhs;
         }
 
-        auto t = next();
+        auto& t = next();
         auto source = t.source();
         auto rhs = additive_expression();
         if (rhs.errored) {
@@ -2702,7 +2736,7 @@ Expect<const ast::Expression*> ParserImpl::expect_relational_expr(const ast::Exp
             return lhs;
         }
 
-        auto t = next();
+        auto& t = next();
         auto source = t.source();
         auto name = t.to_name();
 
@@ -2749,7 +2783,7 @@ Expect<const ast::Expression*> ParserImpl::expect_equality_expr(const ast::Expre
             return lhs;
         }
 
-        auto t = next();
+        auto& t = next();
         auto source = t.source();
         auto name = t.to_name();
 
@@ -2790,7 +2824,7 @@ Expect<const ast::Expression*> ParserImpl::expect_and_expr(const ast::Expression
             return lhs;
         }
 
-        auto t = next();
+        auto& t = next();
         auto source = t.source();
 
         auto rhs = equality_expression();
@@ -2903,7 +2937,7 @@ Expect<const ast::Expression*> ParserImpl::expect_logical_and_expr(const ast::Ex
             return lhs;
         }
 
-        auto t = next();
+        auto& t = next();
         auto source = t.source();
 
         auto rhs = inclusive_or_expression();
@@ -3013,7 +3047,7 @@ Maybe<ast::BinaryOp> ParserImpl::compound_assignment_operator() {
 // 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 @@ Maybe<const ast::Statement*> ParserImpl::assignment_stmt() {
 //   | 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 @@ Maybe<ast::AttributeList> ParserImpl::attribute_list() {
 }
 
 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 @@ Expect<const ast::Attribute*> ParserImpl::expect_attribute() {
 
 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 @@ Maybe<const ast::Attribute*> ParserImpl::attribute() {
             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 @@ Maybe<const ast::Attribute*> ParserImpl::attribute() {
 
             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 @@ Maybe<const ast::Attribute*> ParserImpl::attribute() {
                     } else {
                         return add_error(sampling_tok, "invalid interpolation sampling");
                     }
-
                     match(Token::Type::kComma);
                 }
             }
@@ -3384,7 +3417,7 @@ bool ParserImpl::expect_attributes_consumed(ast::AttributeList& in) {
 }
 
 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::match(Token::Type tok, Source* source /*= nullptr*/) {
 }
 
 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 @@ bool ParserImpl::expect(std::string_view use, Token::Type tok) {
         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 @@ bool ParserImpl::expect(std::string_view use, Token::Type tok) {
 }
 
 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<uint32_t> ParserImpl::expect_nonzero_positive_sint(std::string_view use)
 }
 
 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 @@ bool ParserImpl::sync_to(Token::Type tok, bool consume) {
     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 a9c0423d5a..424fd2eb42 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 @@ class ParserImpl {
     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 @@ class ParserImpl {
     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 @@ class ParserImpl {
     /// 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 @@ class ParserImpl {
     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 @@ class ParserImpl {
         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 75a1ec200a..3c730e2ab1 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 @@ TEST_P(PipelineStageTest, Parses) {
     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 deb87fc50b..7144f2cf9d 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 @@ TEST_P(StorageClassTest, Parses) {
     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 f6ee8ea3b7..1b140399fe 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 @@ class ParserImplTest : public testing::Test, public ProgramBuilder {
     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 @@ class ParserImplTestWithParam : public testing::TestWithParam<T>, public Program
     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 39f8b24dfb..3ed0948db3 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 @@ TEST_F(ParserImplTest, VariableDecl_MissingVar) {
     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 aefd2b5fb6..0dc0a1a9b8 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 @@ TEST_P(VariableQualifierTest, ParsesStorageClass) {
     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 @@ TEST_F(ParserImplTest, VariableQualifier_MissingLessThan) {
     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 @@ TEST_F(ParserImplTest, VariableQualifier_MissingLessThan_AfterSC) {
     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 44255d3239..a7b437fb28 100644
--- a/src/tint/reader/wgsl/token.cc
+++ b/src/tint/reader/wgsl/token.cc
@@ -37,6 +37,8 @@ std::string_view Token::TypeToName(Type type) {
             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(Type type, const Source& source) : type_(type), source_(source) {}
 
 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 473588dbce..352c18ff17 100644
--- a/src/tint/reader/wgsl/token.h
+++ b/src/tint/reader/wgsl/token.h
@@ -32,6 +32,8 @@ class Token {
         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 @@ class Token {
     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 @@ class Token {
 
     /// @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 @@ class Token {
         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_; }