diff --git a/src/tint/api/options/BUILD.bazel b/src/tint/api/options/BUILD.bazel
index 5adc20f..3e1de60 100644
--- a/src/tint/api/options/BUILD.bazel
+++ b/src/tint/api/options/BUILD.bazel
@@ -50,7 +50,6 @@
   ],
   deps = [
     "//src/tint/api/common",
-    "//src/tint/lang/core",
     "//src/tint/utils/macros",
     "//src/tint/utils/math",
     "//src/tint/utils/reflection",
diff --git a/src/tint/api/options/BUILD.cmake b/src/tint/api/options/BUILD.cmake
index d6fb7a5..9cf5e92 100644
--- a/src/tint/api/options/BUILD.cmake
+++ b/src/tint/api/options/BUILD.cmake
@@ -49,7 +49,6 @@
 
 tint_target_add_dependencies(tint_api_options lib
   tint_api_common
-  tint_lang_core
   tint_utils_macros
   tint_utils_math
   tint_utils_reflection
diff --git a/src/tint/api/options/BUILD.gn b/src/tint/api/options/BUILD.gn
index 3c016b5..3b567ca 100644
--- a/src/tint/api/options/BUILD.gn
+++ b/src/tint/api/options/BUILD.gn
@@ -49,7 +49,6 @@
   ]
   deps = [
     "${tint_src_dir}/api/common",
-    "${tint_src_dir}/lang/core",
     "${tint_src_dir}/utils/macros",
     "${tint_src_dir}/utils/math",
     "${tint_src_dir}/utils/reflection",
diff --git a/src/tint/api/options/binding_remapper.h b/src/tint/api/options/binding_remapper.h
index d560493..493c854 100644
--- a/src/tint/api/options/binding_remapper.h
+++ b/src/tint/api/options/binding_remapper.h
@@ -31,7 +31,6 @@
 #include <unordered_map>
 
 #include "src/tint/api/common/binding_point.h"
-#include "src/tint/lang/core/access.h"
 
 namespace tint {
 
diff --git a/src/tint/cmd/fuzz/wgsl/BUILD.bazel b/src/tint/cmd/fuzz/wgsl/BUILD.bazel
index 7c71462..4c77253 100644
--- a/src/tint/cmd/fuzz/wgsl/BUILD.bazel
+++ b/src/tint/cmd/fuzz/wgsl/BUILD.bazel
@@ -42,3 +42,8 @@
   actual = "//src/tint:tint_build_wgsl_reader_true",
 )
 
+alias(
+  name = "tint_build_wgsl_writer",
+  actual = "//src/tint:tint_build_wgsl_writer_true",
+)
+
diff --git a/src/tint/cmd/fuzz/wgsl/BUILD.cmake b/src/tint/cmd/fuzz/wgsl/BUILD.cmake
index 9ba6530..9648f95 100644
--- a/src/tint/cmd/fuzz/wgsl/BUILD.cmake
+++ b/src/tint/cmd/fuzz/wgsl/BUILD.cmake
@@ -51,6 +51,7 @@
   tint_lang_wgsl
   tint_lang_wgsl_ast
   tint_lang_wgsl_program
+  tint_lang_wgsl_program_fuzz
   tint_lang_wgsl_sem
   tint_lang_wgsl_fuzz
   tint_utils_cli
@@ -76,6 +77,12 @@
   )
 endif(TINT_BUILD_WGSL_READER)
 
+if(TINT_BUILD_WGSL_WRITER)
+  tint_target_add_dependencies(tint_cmd_fuzz_wgsl_fuzz_cmd fuzz_cmd
+    tint_lang_wgsl_writer_ast_printer_fuzz
+  )
+endif(TINT_BUILD_WGSL_WRITER)
+
 tint_target_set_output_name(tint_cmd_fuzz_wgsl_fuzz_cmd fuzz_cmd "tint_wgsl_fuzzer")
 
 endif(TINT_BUILD_WGSL_READER)
diff --git a/src/tint/cmd/fuzz/wgsl/BUILD.gn b/src/tint/cmd/fuzz/wgsl/BUILD.gn
index 5c4ccc6..5b86127 100644
--- a/src/tint/cmd/fuzz/wgsl/BUILD.gn
+++ b/src/tint/cmd/fuzz/wgsl/BUILD.gn
@@ -86,6 +86,7 @@
       "${tint_src_dir}/lang/wgsl:fuzz",
       "${tint_src_dir}/lang/wgsl/ast",
       "${tint_src_dir}/lang/wgsl/program",
+      "${tint_src_dir}/lang/wgsl/program:fuzz",
       "${tint_src_dir}/lang/wgsl/sem",
       "${tint_src_dir}/utils/cli",
       "${tint_src_dir}/utils/containers",
@@ -109,5 +110,9 @@
         "${tint_src_dir}/lang/wgsl/ast/transform:fuzz",
       ]
     }
+
+    if (tint_build_wgsl_writer) {
+      deps += [ "${tint_src_dir}/lang/wgsl/writer/ast_printer:fuzz" ]
+    }
   }
 }
diff --git a/src/tint/fuzzers/BUILD.gn b/src/tint/fuzzers/BUILD.gn
index 7f5b0cd..375d00d 100644
--- a/src/tint/fuzzers/BUILD.gn
+++ b/src/tint/fuzzers/BUILD.gn
@@ -134,20 +134,6 @@
   }
 
   if (tint_build_wgsl_reader && tint_build_wgsl_writer) {
-    fuzzer_test("tint_ast_clone_fuzzer") {
-      sources = [ "tint_ast_clone_fuzzer.cc" ]
-      deps = [ ":tint_fuzzer_common_with_init_src" ]
-      deps += [
-        "${tint_src_dir}/lang/wgsl/reader",
-        "${tint_src_dir}/lang/wgsl/reader/parser",
-      ]
-
-      dict = "dictionary.txt"
-      libfuzzer_options = tint_fuzzer_common_libfuzzer_options
-      seed_corpus = fuzzer_corpus_wgsl_dir
-      seed_corpus_deps = [ ":tint_generate_wgsl_corpus" ]
-    }
-
     if (build_with_chromium) {
       fuzzer_test("tint_ast_wgsl_writer_fuzzer") {
         sources = [ "tint_ast_fuzzer/tint_ast_wgsl_writer_fuzzer.cc" ]
@@ -165,15 +151,6 @@
       seed_corpus = fuzzer_corpus_wgsl_dir
       seed_corpus_deps = [ ":tint_generate_wgsl_corpus" ]
     }
-
-    fuzzer_test("tint_wgsl_reader_wgsl_writer_fuzzer") {
-      sources = [ "tint_wgsl_reader_wgsl_writer_fuzzer.cc" ]
-      deps = [ ":tint_fuzzer_common_with_init_src" ]
-      dict = "dictionary.txt"
-      libfuzzer_options = tint_fuzzer_common_libfuzzer_options
-      seed_corpus = fuzzer_corpus_wgsl_dir
-      seed_corpus_deps = [ ":tint_generate_wgsl_corpus" ]
-    }
   }
 
   if (tint_build_wgsl_reader && tint_build_spv_writer) {
@@ -343,11 +320,7 @@
     deps = []
 
     if (tint_build_wgsl_reader && tint_build_wgsl_writer) {
-      deps += [
-        ":tint_ast_clone_fuzzer",
-        ":tint_regex_wgsl_writer_fuzzer",
-        ":tint_wgsl_reader_wgsl_writer_fuzzer",
-      ]
+      deps += [ ":tint_regex_wgsl_writer_fuzzer" ]
       if (build_with_chromium) {
         deps += [ ":tint_ast_wgsl_writer_fuzzer" ]
       }
diff --git a/src/tint/fuzzers/CMakeLists.txt b/src/tint/fuzzers/CMakeLists.txt
index cd5a057..3801e7f 100644
--- a/src/tint/fuzzers/CMakeLists.txt
+++ b/src/tint/fuzzers/CMakeLists.txt
@@ -55,10 +55,6 @@
   target_compile_options(${NAME} PRIVATE -Wno-missing-prototypes)
 endfunction()
 
-if (${TINT_BUILD_WGSL_READER} AND ${TINT_BUILD_WGSL_WRITER})
-  add_tint_fuzzer(tint_wgsl_reader_wgsl_writer_fuzzer)
-endif()
-
 if (${TINT_BUILD_WGSL_READER} AND ${TINT_BUILD_SPV_WRITER})
   add_tint_fuzzer(tint_all_transforms_fuzzer)
   add_tint_fuzzer(tint_binding_remapper_fuzzer)
@@ -95,10 +91,6 @@
   add_tint_fuzzer(tint_spv_reader_msl_writer_fuzzer)
 endif()
 
-if (${TINT_BUILD_WGSL_READER} AND ${TINT_BUILD_WGSL_WRITER})
-  add_tint_fuzzer(tint_ast_clone_fuzzer)
-endif()
-
 if (${TINT_BUILD_AST_FUZZER})
   add_subdirectory(tint_ast_fuzzer)
 endif()
diff --git a/src/tint/fuzzers/tint_ast_clone_fuzzer.cc b/src/tint/fuzzers/tint_ast_clone_fuzzer.cc
deleted file mode 100644
index e33d3f9..0000000
--- a/src/tint/fuzzers/tint_ast_clone_fuzzer.cc
+++ /dev/null
@@ -1,126 +0,0 @@
-// Copyright 2020 The Dawn & Tint Authors
-//
-// Redistribution and use in source and binary forms, with or without
-// modification, are permitted provided that the following conditions are met:
-//
-// 1. Redistributions of source code must retain the above copyright notice, this
-//    list of conditions and the following disclaimer.
-//
-// 2. Redistributions in binary form must reproduce the above copyright notice,
-//    this list of conditions and the following disclaimer in the documentation
-//    and/or other materials provided with the distribution.
-//
-// 3. Neither the name of the copyright holder nor the names of its
-//    contributors may be used to endorse or promote products derived from
-//    this software without specific prior written permission.
-//
-// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
-// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
-// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
-// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
-// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
-// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
-// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
-// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
-// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
-// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-
-#include <iostream>
-#include <string>
-#include <unordered_set>
-
-#include "src/tint/lang/wgsl/reader/parser/parser.h"
-#include "src/tint/lang/wgsl/writer/writer.h"
-
-#define ASSERT_EQ(A, B)                                        \
-    do {                                                       \
-        decltype(A) assert_a = (A);                            \
-        decltype(B) assert_b = (B);                            \
-        if (assert_a != assert_b) {                            \
-            std::cerr << "ASSERT_EQ(" #A ", " #B ") failed:\n" \
-                      << #A << " was: " << assert_a << "\n"    \
-                      << #B << " was: " << assert_b << "\n";   \
-            __builtin_trap();                                  \
-        }                                                      \
-    } while (false)
-
-#define ASSERT_TRUE(A)                                                                          \
-    do {                                                                                        \
-        decltype(A) assert_a = (A);                                                             \
-        if (!assert_a) {                                                                        \
-            std::cerr << "ASSERT_TRUE(" #A ") failed:\n" << #A << " was: " << assert_a << "\n"; \
-            __builtin_trap();                                                                   \
-        }                                                                                       \
-    } while (false)
-
-[[noreturn]] void TintInternalCompilerErrorReporter(const tint::InternalCompilerError& err) {
-    std::cerr << err.Error() << std::endl;
-    __builtin_trap();
-}
-
-extern "C" int LLVMFuzzerTestOneInput(const uint8_t* data, size_t size) {
-    std::string str(reinterpret_cast<const char*>(data), size);
-
-    tint::SetInternalCompilerErrorReporter(&TintInternalCompilerErrorReporter);
-
-    tint::Source::File file("test.wgsl", str);
-
-    // Parse the wgsl, create the src program
-    tint::wgsl::reader::Parser parser(&file);
-    parser.set_max_errors(1);
-    if (!parser.Parse()) {
-        return 0;
-    }
-    auto src = parser.program();
-    if (!src.IsValid()) {
-        return 0;
-    }
-
-    // Clone the src program to dst
-    tint::Program dst(src.Clone());
-
-    // Expect the printed strings to match
-    ASSERT_EQ(tint::Program::printer(src), tint::Program::printer(dst));
-
-    // Check that none of the AST nodes or type pointers in dst are found in src
-    std::unordered_set<const tint::ast::Node*> src_nodes;
-    for (auto* src_node : src.ASTNodes().Objects()) {
-        src_nodes.emplace(src_node);
-    }
-    std::unordered_set<const tint::core::type::Type*> src_types;
-    for (auto* src_type : src.Types()) {
-        src_types.emplace(src_type);
-    }
-    for (auto* dst_node : dst.ASTNodes().Objects()) {
-        ASSERT_EQ(src_nodes.count(dst_node), 0u);
-    }
-    for (auto* dst_type : dst.Types()) {
-        ASSERT_EQ(src_types.count(dst_type), 0u);
-    }
-
-    // Regenerate the wgsl for the src program. We use this instead of the
-    // original source so that reformatting doesn't impact the final wgsl
-    // comparison.
-    std::string src_wgsl;
-    tint::wgsl::writer::Options wgsl_options;
-    {
-        auto result = tint::wgsl::writer::Generate(src, wgsl_options);
-        ASSERT_TRUE(result == true);
-        src_wgsl = result->wgsl;
-
-        // Move the src program to a temporary that'll be dropped, so that the src
-        // program is released before we attempt to print the dst program. This
-        // guarantee that all the source program nodes and types are destructed and
-        // freed. ASAN should error if there's any remaining references in dst when
-        // we try to reconstruct the WGSL.
-        auto tmp = std::move(src);
-    }
-
-    // Print the dst program, check it matches the original source
-    auto result = tint::wgsl::writer::Generate(dst, wgsl_options);
-    ASSERT_TRUE(result == true);
-    auto dst_wgsl = result->wgsl;
-    ASSERT_EQ(src_wgsl, dst_wgsl);
-
-    return 0;
-}
diff --git a/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc b/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
index cec5d7e..b0c7219 100644
--- a/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
@@ -319,6 +319,9 @@
                 // Record the required extension for generating extension directive later
                 RecordExtension(enable);
             },  //
+            [&](const ast::Requires*) {
+                // Do nothing for requiring language features in GLSL.
+            },  //
             TINT_ICE_ON_NO_MATCH);
     }
 
diff --git a/src/tint/lang/glsl/writer/ast_printer/ast_printer_test.cc b/src/tint/lang/glsl/writer/ast_printer/ast_printer_test.cc
index 7d70089..9834a54 100644
--- a/src/tint/lang/glsl/writer/ast_printer/ast_printer_test.cc
+++ b/src/tint/lang/glsl/writer/ast_printer/ast_printer_test.cc
@@ -131,5 +131,16 @@
               R"(12:34 error: GLSL backend does not support extension 'undefined')");
 }
 
+TEST_F(GlslASTPrinterTest, RequiresDirective) {
+    Require(wgsl::LanguageFeature::kReadonlyAndReadwriteStorageTextures);
+
+    ASTPrinter& gen = Build();
+
+    ASSERT_TRUE(gen.Generate()) << gen.Diagnostics();
+    EXPECT_EQ(gen.Result(), R"(#version 310 es
+
+)");
+}
+
 }  // namespace
 }  // namespace tint::glsl::writer
diff --git a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
index d84e98a..b23c19f 100644
--- a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
@@ -370,7 +370,8 @@
 
     auto* mod = builder_.Sem().Module();
     for (auto* decl : mod->DependencyOrderedDeclarations()) {
-        if (decl->IsAnyOf<ast::Alias, ast::DiagnosticDirective, ast::Enable, ast::ConstAssert>()) {
+        if (decl->IsAnyOf<ast::Alias, ast::DiagnosticDirective, ast::Enable, ast::Requires,
+                          ast::ConstAssert>()) {
             continue;  // These are not emitted.
         }
 
diff --git a/src/tint/lang/hlsl/writer/ast_printer/ast_printer_test.cc b/src/tint/lang/hlsl/writer/ast_printer/ast_printer_test.cc
index a42a651..23c9713 100644
--- a/src/tint/lang/hlsl/writer/ast_printer/ast_printer_test.cc
+++ b/src/tint/lang/hlsl/writer/ast_printer/ast_printer_test.cc
@@ -54,6 +54,15 @@
               R"(12:34 error: HLSL backend does not support extension 'undefined')");
 }
 
+TEST_F(HlslASTPrinterTest, RequiresDirective) {
+    Require(wgsl::LanguageFeature::kReadonlyAndReadwriteStorageTextures);
+
+    ASTPrinter& gen = Build();
+
+    ASSERT_TRUE(gen.Generate()) << gen.Diagnostics();
+    EXPECT_EQ(gen.Result(), "");
+}
+
 TEST_F(HlslASTPrinterTest, Generate) {
     Func("my_func", {}, ty.void_(), {});
 
diff --git a/src/tint/lang/hlsl/writer/common/options.h b/src/tint/lang/hlsl/writer/common/options.h
index b71fe89..e8d5cc0 100644
--- a/src/tint/lang/hlsl/writer/common/options.h
+++ b/src/tint/lang/hlsl/writer/common/options.h
@@ -37,6 +37,7 @@
 #include "src/tint/api/options/array_length_from_uniform.h"
 #include "src/tint/api/options/binding_remapper.h"
 #include "src/tint/api/options/external_texture.h"
+#include "src/tint/lang/core/access.h"
 #include "src/tint/utils/reflection/reflection.h"
 
 namespace tint::hlsl::writer {
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
index 69a4f31..3d49d28 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
@@ -325,6 +325,10 @@
                 // Do nothing for enabling extension in MSL
                 return true;
             },
+            [&](const ast::Requires*) {
+                // Do nothing for requiring language features in MSL.
+                return true;
+            },
             [&](const ast::ConstAssert*) {
                 return true;  // Not emitted
             },                //
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc b/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc
index 3b738f5..49fa3c1 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc
@@ -58,6 +58,18 @@
               R"(12:34 error: MSL backend does not support extension 'undefined')");
 }
 
+TEST_F(MslASTPrinterTest, RequiresDirective) {
+    Require(wgsl::LanguageFeature::kReadonlyAndReadwriteStorageTextures);
+
+    ASTPrinter& gen = Build();
+
+    ASSERT_TRUE(gen.Generate()) << gen.Diagnostics();
+    EXPECT_EQ(gen.Result(), R"(#include <metal_stdlib>
+
+using namespace metal;
+)");
+}
+
 TEST_F(MslASTPrinterTest, Generate) {
     Func("my_func", tint::Empty, ty.void_(), tint::Empty,
          Vector{
diff --git a/src/tint/lang/spirv/reader/ast_parser/ast_parser.cc b/src/tint/lang/spirv/reader/ast_parser/ast_parser.cc
index 0bf7f61..9370817 100644
--- a/src/tint/lang/spirv/reader/ast_parser/ast_parser.cc
+++ b/src/tint/lang/spirv/reader/ast_parser/ast_parser.cc
@@ -2583,6 +2583,9 @@
         } else {
             const auto access =
                 usage.IsStorageReadWriteTexture() ? core::Access::kReadWrite : core::Access::kWrite;
+            if (access == core::Access::kReadWrite) {
+                Require(wgsl::LanguageFeature::kReadonlyAndReadwriteStorageTextures);
+            }
             const auto format = enum_converter_.ToTexelFormat(image_type->format());
             if (format == core::TexelFormat::kUndefined) {
                 return nullptr;
diff --git a/src/tint/lang/spirv/reader/ast_parser/ast_parser.h b/src/tint/lang/spirv/reader/ast_parser/ast_parser.h
index 59d3c02..ddd8ded 100644
--- a/src/tint/lang/spirv/reader/ast_parser/ast_parser.h
+++ b/src/tint/lang/spirv/reader/ast_parser/ast_parser.h
@@ -759,6 +759,14 @@
         }
     }
 
+    /// Require a WGSL language feature, if not already required.
+    /// @param feature the language feature to require
+    void Require(wgsl::LanguageFeature feature) {
+        if (required_features_.Add(feature)) {
+            builder_.Require(feature);
+        }
+    }
+
   private:
     /// Converts a specific SPIR-V type to a Tint type. Integer case
     const Type* ConvertType(const spvtools::opt::analysis::Integer* int_ty);
@@ -948,6 +956,8 @@
 
     /// Set of WGSL extensions that have been enabled.
     Hashset<wgsl::Extension, 4> enabled_extensions_;
+    /// Set of WGSL language features that have been required.
+    Hashset<wgsl::LanguageFeature, 4> required_features_;
 };
 
 }  // namespace tint::spirv::reader::ast_parser
diff --git a/src/tint/lang/spirv/reader/ast_parser/function.cc b/src/tint/lang/spirv/reader/ast_parser/function.cc
index 198f58f..2ba5673 100644
--- a/src/tint/lang/spirv/reader/ast_parser/function.cc
+++ b/src/tint/lang/spirv/reader/ast_parser/function.cc
@@ -5325,6 +5325,7 @@
         if (memory != uint32_t(spv::Scope::Workgroup)) {
             return Fail() << "textureBarrier requires workgroup memory scope";
         }
+        parser_impl_.Require(wgsl::LanguageFeature::kReadonlyAndReadwriteStorageTextures);
         AddStatement(builder_.CallStmt(builder_.Call("textureBarrier")));
         semantics &= ~static_cast<uint32_t>(spv::MemorySemanticsMask::ImageMemory);
     }
diff --git a/src/tint/lang/spirv/reader/ast_parser/handle_test.cc b/src/tint/lang/spirv/reader/ast_parser/handle_test.cc
index 154ccc2..dd05ded 100644
--- a/src/tint/lang/spirv/reader/ast_parser/handle_test.cc
+++ b/src/tint/lang/spirv/reader/ast_parser/handle_test.cc
@@ -4234,7 +4234,9 @@
     EXPECT_TRUE(p->error().empty()) << p->error();
     const auto got = test::ToString(p->program());
     auto* expect =
-        R"(@group(0) @binding(0) var RWTexture2D : texture_storage_2d<rgba32float, read_write>;
+        R"(requires readonly_and_readwrite_storage_textures;
+
+@group(0) @binding(0) var RWTexture2D : texture_storage_2d<rgba32float, read_write>;
 
 const x_9 = vec2u(1u);
 
diff --git a/src/tint/lang/spirv/writer/ast_printer/ast_printer_test.cc b/src/tint/lang/spirv/writer/ast_printer/ast_printer_test.cc
index 932ab4d..9fda078 100644
--- a/src/tint/lang/spirv/writer/ast_printer/ast_printer_test.cc
+++ b/src/tint/lang/spirv/writer/ast_printer/ast_printer_test.cc
@@ -53,5 +53,13 @@
               R"(12:34 error: SPIR-V backend does not support extension 'undefined')");
 }
 
+TEST_F(SpirvASTPrinterTest, RequiresDirective) {
+    Require(wgsl::LanguageFeature::kReadonlyAndReadwriteStorageTextures);
+
+    auto program = resolver::Resolve(*this);
+    auto result = Generate(program, Options{});
+    EXPECT_TRUE(result);
+}
+
 }  // namespace
 }  // namespace tint::spirv::writer
diff --git a/src/tint/lang/spirv/writer/common/option_builder.cc b/src/tint/lang/spirv/writer/common/option_builder.cc
index 6b66447..55181fc 100644
--- a/src/tint/lang/spirv/writer/common/option_builder.cc
+++ b/src/tint/lang/spirv/writer/common/option_builder.cc
@@ -32,31 +32,41 @@
 namespace tint::spirv::writer {
 
 bool ValidateBindingOptions(const Options& options, diag::List& diagnostics) {
-    tint::Hashset<tint::BindingPoint, 8> seen_wgsl_bindings{};
-    tint::Hashset<binding::BindingInfo, 8> seen_spirv_bindings{};
+    tint::Hashmap<tint::BindingPoint, binding::BindingInfo, 8> seen_wgsl_bindings{};
+    tint::Hashmap<binding::BindingInfo, tint::BindingPoint, 8> seen_spirv_bindings{};
 
-    auto wgsl_seen = [&diagnostics, &seen_wgsl_bindings](const tint::BindingPoint& info) -> bool {
-        if (seen_wgsl_bindings.Contains(info)) {
-            std::stringstream str;
-            str << "found duplicate WGSL binding point: " << info;
+    // Both wgsl_seen and spirv_seen check to see if the pair of [src, dst] are unique. If we have
+    // multiple entries that map the same [src, dst] pair, that's fine. We treat it as valid as it's
+    // possible for multiple entry points to use the remapper at the same time. If the pair doesn't
+    // match, then we report an error about a duplicate binding point.
 
-            diagnostics.add_error(diag::System::Writer, str.str());
-            return true;
+    auto wgsl_seen = [&diagnostics, &seen_wgsl_bindings](const tint::BindingPoint& src,
+                                                         const binding::BindingInfo& dst) -> bool {
+        if (auto binding = seen_wgsl_bindings.Find(src)) {
+            if (*binding != dst) {
+                std::stringstream str;
+                str << "found duplicate WGSL binding point: " << src;
+
+                diagnostics.add_error(diag::System::Writer, str.str());
+                return true;
+            }
         }
-        seen_wgsl_bindings.Add(info);
+        seen_wgsl_bindings.Add(src, dst);
         return false;
     };
 
-    auto spirv_seen = [&diagnostics,
-                       &seen_spirv_bindings](const binding::BindingInfo& info) -> bool {
-        if (seen_spirv_bindings.Contains(info)) {
-            std::stringstream str;
-            str << "found duplicate SPIR-V binding point: [group: " << info.group
-                << ", binding: " << info.binding << "]";
-            diagnostics.add_error(diag::System::Writer, str.str());
-            return true;
+    auto spirv_seen = [&diagnostics, &seen_spirv_bindings](const binding::BindingInfo& src,
+                                                           const tint::BindingPoint& dst) -> bool {
+        if (auto binding = seen_spirv_bindings.Find(src)) {
+            if (*binding != dst) {
+                std::stringstream str;
+                str << "found duplicate SPIR-V binding point: [group: " << src.group
+                    << ", binding: " << src.binding << "]";
+                diagnostics.add_error(diag::System::Writer, str.str());
+                return true;
+            }
         }
-        seen_spirv_bindings.Add(info);
+        seen_spirv_bindings.Add(src, dst);
         return false;
     };
 
@@ -65,11 +75,11 @@
             const auto& src_binding = it.first;
             const auto& dst_binding = it.second;
 
-            if (wgsl_seen(src_binding)) {
+            if (wgsl_seen(src_binding, dst_binding)) {
                 return false;
             }
 
-            if (spirv_seen(dst_binding)) {
+            if (spirv_seen(dst_binding, src_binding)) {
                 return false;
             }
         }
@@ -104,20 +114,20 @@
         const auto& metadata = it.second.metadata;
 
         // Validate with the actual source regardless of what the remapper will do
-        if (wgsl_seen(src_binding)) {
+        if (wgsl_seen(src_binding, plane0)) {
             diagnostics.add_note(diag::System::Writer, "when processing external_texture", {});
             return false;
         }
 
-        if (spirv_seen(plane0)) {
+        if (spirv_seen(plane0, src_binding)) {
             diagnostics.add_note(diag::System::Writer, "when processing external_texture", {});
             return false;
         }
-        if (spirv_seen(plane1)) {
+        if (spirv_seen(plane1, src_binding)) {
             diagnostics.add_note(diag::System::Writer, "when processing external_texture", {});
             return false;
         }
-        if (spirv_seen(metadata)) {
+        if (spirv_seen(metadata, src_binding)) {
             diagnostics.add_note(diag::System::Writer, "when processing external_texture", {});
             return false;
         }
diff --git a/src/tint/lang/spirv/writer/common/options.h b/src/tint/lang/spirv/writer/common/options.h
index c0f5e28..96b585e 100644
--- a/src/tint/lang/spirv/writer/common/options.h
+++ b/src/tint/lang/spirv/writer/common/options.h
@@ -49,6 +49,10 @@
     inline bool operator==(const BindingInfo& rhs) const {
         return group == rhs.group && binding == rhs.binding;
     }
+    /// Inequality operator
+    /// @param rhs the BindingInfo to compare against
+    /// @returns true if this BindingInfo is not equal to `rhs`
+    inline bool operator!=(const BindingInfo& rhs) const { return !(*this == rhs); }
 
     /// Reflect the fields of this class so that it can be used by tint::ForeachField()
     TINT_REFLECT(group, binding);
diff --git a/src/tint/lang/spirv/writer/helpers/generate_bindings.cc b/src/tint/lang/spirv/writer/helpers/generate_bindings.cc
index 1f429ba..21ef88b 100644
--- a/src/tint/lang/spirv/writer/helpers/generate_bindings.cc
+++ b/src/tint/lang/spirv/writer/helpers/generate_bindings.cc
@@ -50,28 +50,12 @@
 
     Bindings bindings{};
 
-    std::unordered_set<tint::BindingPoint> seen_binding_points;
-
     // Collect next valid binding number per group
     Hashmap<uint32_t, uint32_t, 4> group_to_next_binding_number;
     Vector<tint::BindingPoint, 4> ext_tex_bps;
     for (auto* var : program.AST().GlobalVariables()) {
         if (auto* sem_var = program.Sem().Get(var)->As<sem::GlobalVariable>()) {
             if (auto bp = sem_var->BindingPoint()) {
-                // This is a bit of a hack. The binding points must be unique over all the `binding`
-                // entries. But, this is looking at _all_ entry points where bindings can overlap.
-                // In the case where both entry points used the same type (uniform, sampler, etc)
-                // then it would be fine as it just overwrites with itself. But, if one entry point
-                // has a uniform and the other a sampler at the same (group,binding) pair then we'll
-                // get a validation error due to duplicate WGSL bindings.
-                //
-                // For generating bindings we don't really care as we always map to itself, so if it
-                // exists anywhere, we just pretend that's the only one.
-                if (seen_binding_points.find(*bp) != seen_binding_points.end()) {
-                    continue;
-                }
-                seen_binding_points.emplace(*bp);
-
                 if (auto val = group_to_next_binding_number.Find(bp->group)) {
                     *val = std::max(*val, bp->binding + 1);
                 } else {
diff --git a/src/tint/lang/wgsl/BUILD.bazel b/src/tint/lang/wgsl/BUILD.bazel
index 9b7655d..0dd8726 100644
--- a/src/tint/lang/wgsl/BUILD.bazel
+++ b/src/tint/lang/wgsl/BUILD.bazel
@@ -43,12 +43,14 @@
     "diagnostic_rule.cc",
     "diagnostic_severity.cc",
     "extension.cc",
+    "language_feature.cc",
   ],
   hdrs = [
     "builtin_fn.h",
     "diagnostic_rule.h",
     "diagnostic_severity.h",
     "extension.h",
+    "language_feature.h",
   ],
   deps = [
     "//src/tint/utils/containers",
@@ -71,6 +73,7 @@
     "diagnostic_rule_test.cc",
     "diagnostic_severity_test.cc",
     "extension_test.cc",
+    "language_feature_test.cc",
     "wgsl_test.cc",
   ] + select({
     "//conditions:default": [],
diff --git a/src/tint/lang/wgsl/BUILD.cmake b/src/tint/lang/wgsl/BUILD.cmake
index cea7aa0..7a1599d 100644
--- a/src/tint/lang/wgsl/BUILD.cmake
+++ b/src/tint/lang/wgsl/BUILD.cmake
@@ -58,6 +58,8 @@
   lang/wgsl/diagnostic_severity.h
   lang/wgsl/extension.cc
   lang/wgsl/extension.h
+  lang/wgsl/language_feature.cc
+  lang/wgsl/language_feature.h
 )
 
 tint_target_add_dependencies(tint_lang_wgsl lib
@@ -80,6 +82,7 @@
   lang/wgsl/diagnostic_rule_test.cc
   lang/wgsl/diagnostic_severity_test.cc
   lang/wgsl/extension_test.cc
+  lang/wgsl/language_feature_test.cc
   lang/wgsl/wgsl_test.cc
 )
 
diff --git a/src/tint/lang/wgsl/BUILD.gn b/src/tint/lang/wgsl/BUILD.gn
index 10d334a..619e1f8 100644
--- a/src/tint/lang/wgsl/BUILD.gn
+++ b/src/tint/lang/wgsl/BUILD.gn
@@ -52,6 +52,8 @@
     "diagnostic_severity.h",
     "extension.cc",
     "extension.h",
+    "language_feature.cc",
+    "language_feature.h",
   ]
   deps = [
     "${tint_src_dir}/utils/containers",
@@ -71,6 +73,7 @@
       "diagnostic_rule_test.cc",
       "diagnostic_severity_test.cc",
       "extension_test.cc",
+      "language_feature_test.cc",
       "wgsl_test.cc",
     ]
     deps = [
diff --git a/src/tint/lang/wgsl/ast/BUILD.bazel b/src/tint/lang/wgsl/ast/BUILD.bazel
index eabb9ca..bc6462d 100644
--- a/src/tint/lang/wgsl/ast/BUILD.bazel
+++ b/src/tint/lang/wgsl/ast/BUILD.bazel
@@ -97,6 +97,7 @@
     "parameter.cc",
     "phony_expression.cc",
     "pipeline_stage.cc",
+    "requires.cc",
     "return_statement.cc",
     "stage_attribute.cc",
     "statement.cc",
@@ -177,6 +178,7 @@
     "parameter.h",
     "phony_expression.h",
     "pipeline_stage.h",
+    "requires.h",
     "return_statement.h",
     "stage_attribute.h",
     "statement.h",
@@ -271,6 +273,7 @@
     "member_accessor_expression_test.cc",
     "module_test.cc",
     "phony_expression_test.cc",
+    "requires_test.cc",
     "return_statement_test.cc",
     "stage_attribute_test.cc",
     "stride_attribute_test.cc",
diff --git a/src/tint/lang/wgsl/ast/BUILD.cmake b/src/tint/lang/wgsl/ast/BUILD.cmake
index e72ea46..e35a8b1 100644
--- a/src/tint/lang/wgsl/ast/BUILD.cmake
+++ b/src/tint/lang/wgsl/ast/BUILD.cmake
@@ -158,6 +158,8 @@
   lang/wgsl/ast/phony_expression.h
   lang/wgsl/ast/pipeline_stage.cc
   lang/wgsl/ast/pipeline_stage.h
+  lang/wgsl/ast/requires.cc
+  lang/wgsl/ast/requires.h
   lang/wgsl/ast/return_statement.cc
   lang/wgsl/ast/return_statement.h
   lang/wgsl/ast/stage_attribute.cc
@@ -271,6 +273,7 @@
   lang/wgsl/ast/member_accessor_expression_test.cc
   lang/wgsl/ast/module_test.cc
   lang/wgsl/ast/phony_expression_test.cc
+  lang/wgsl/ast/requires_test.cc
   lang/wgsl/ast/return_statement_test.cc
   lang/wgsl/ast/stage_attribute_test.cc
   lang/wgsl/ast/stride_attribute_test.cc
diff --git a/src/tint/lang/wgsl/ast/BUILD.gn b/src/tint/lang/wgsl/ast/BUILD.gn
index e440921..9c36fd0 100644
--- a/src/tint/lang/wgsl/ast/BUILD.gn
+++ b/src/tint/lang/wgsl/ast/BUILD.gn
@@ -161,6 +161,8 @@
     "phony_expression.h",
     "pipeline_stage.cc",
     "pipeline_stage.h",
+    "requires.cc",
+    "requires.h",
     "return_statement.cc",
     "return_statement.h",
     "stage_attribute.cc",
@@ -271,6 +273,7 @@
       "member_accessor_expression_test.cc",
       "module_test.cc",
       "phony_expression_test.cc",
+      "requires_test.cc",
       "return_statement_test.cc",
       "stage_attribute_test.cc",
       "stride_attribute_test.cc",
diff --git a/src/tint/lang/wgsl/ast/builder.h b/src/tint/lang/wgsl/ast/builder.h
index bab8fb9..1cbdfba 100644
--- a/src/tint/lang/wgsl/ast/builder.h
+++ b/src/tint/lang/wgsl/ast/builder.h
@@ -97,6 +97,7 @@
 #include "src/tint/lang/wgsl/ast/override.h"
 #include "src/tint/lang/wgsl/ast/parameter.h"
 #include "src/tint/lang/wgsl/ast/phony_expression.h"
+#include "src/tint/lang/wgsl/ast/requires.h"
 #include "src/tint/lang/wgsl/ast/return_statement.h"
 #include "src/tint/lang/wgsl/ast/stage_attribute.h"
 #include "src/tint/lang/wgsl/ast/stride_attribute.h"
@@ -1617,6 +1618,25 @@
         return enable;
     }
 
+    /// Adds the language feature to the list of requires directives at the top of the module.
+    /// @param feature the feature to require
+    /// @return a `ast::Requires` requiring the given language feature.
+    const ast::Requires* Require(wgsl::LanguageFeature feature) {
+        auto* req = create<ast::Requires>(wgsl::LanguageFeatures({feature}));
+        AST().AddRequires(req);
+        return req;
+    }
+
+    /// Adds the language feature to the list of requires directives at the top of the module.
+    /// @param source the requires source
+    /// @param feature the feature to require
+    /// @return a `ast::Requires` requiring the given language feature.
+    const ast::Requires* Require(const Source& source, wgsl::LanguageFeature feature) {
+        auto* req = create<ast::Requires>(source, wgsl::LanguageFeatures({feature}));
+        AST().AddRequires(req);
+        return req;
+    }
+
     /// @param name the variable name
     /// @param options the extra options passed to the ast::Var initializer
     /// Can be any of the following, in any order:
diff --git a/src/tint/lang/wgsl/ast/module.cc b/src/tint/lang/wgsl/ast/module.cc
index 3ee8cfd..974abad 100644
--- a/src/tint/lang/wgsl/ast/module.cc
+++ b/src/tint/lang/wgsl/ast/module.cc
@@ -89,6 +89,10 @@
             TINT_ASSERT_GENERATION_IDS_EQUAL_IF_VALID(enable, generation_id);
             enables_.Push(enable);
         },
+        [&](const ast::Requires* req) {
+            TINT_ASSERT_GENERATION_IDS_EQUAL_IF_VALID(req, generation_id);
+            requires_.Push(req);
+        },
         [&](const ConstAssert* assertion) {
             TINT_ASSERT_GENERATION_IDS_EQUAL_IF_VALID(assertion, generation_id);
             const_asserts_.Push(assertion);
@@ -110,6 +114,13 @@
     enables_.Push(enable);
 }
 
+void Module::AddRequires(const ast::Requires* req) {
+    TINT_ASSERT(req);
+    TINT_ASSERT_GENERATION_IDS_EQUAL_IF_VALID(req, generation_id);
+    global_declarations_.Push(req);
+    requires_.Push(req);
+}
+
 void Module::AddGlobalVariable(const Variable* var) {
     TINT_ASSERT(var);
     TINT_ASSERT_GENERATION_IDS_EQUAL_IF_VALID(var, generation_id);
diff --git a/src/tint/lang/wgsl/ast/module.h b/src/tint/lang/wgsl/ast/module.h
index 7214aec..8513348 100644
--- a/src/tint/lang/wgsl/ast/module.h
+++ b/src/tint/lang/wgsl/ast/module.h
@@ -34,6 +34,7 @@
 #include "src/tint/lang/wgsl/ast/diagnostic_directive.h"
 #include "src/tint/lang/wgsl/ast/enable.h"
 #include "src/tint/lang/wgsl/ast/function.h"
+#include "src/tint/lang/wgsl/ast/requires.h"
 #include "src/tint/utils/containers/vector.h"
 
 namespace tint::ast {
@@ -110,12 +111,19 @@
     /// @param ext the enable directive to add
     void AddEnable(const Enable* ext);
 
+    /// Add a requires directive to the module
+    /// @param req the requires directive to add
+    void AddRequires(const Requires* req);
+
     /// @returns the diagnostic directives for the module
     const auto& DiagnosticDirectives() const { return diagnostic_directives_; }
 
     /// @returns the extension set for the module
     const auto& Enables() const { return enables_; }
 
+    /// @returns the requires directives for the module
+    const auto& Requires() const { return requires_; }
+
     /// Add a global const assertion to the module
     /// @param assertion the const assert to add
     void AddConstAssert(const ConstAssert* assertion);
@@ -168,6 +176,7 @@
     tint::Vector<const Variable*, 32> global_variables_;
     tint::Vector<const DiagnosticDirective*, 8> diagnostic_directives_;
     tint::Vector<const Enable*, 8> enables_;
+    tint::Vector<const ast::Requires*, 8> requires_;
     tint::Vector<const ConstAssert*, 8> const_asserts_;
 };
 
diff --git a/src/tint/lang/wgsl/ast/module_clone_test.cc b/src/tint/lang/wgsl/ast/module_clone_test.cc
index 6e64f65..3444805 100644
--- a/src/tint/lang/wgsl/ast/module_clone_test.cc
+++ b/src/tint/lang/wgsl/ast/module_clone_test.cc
@@ -38,7 +38,6 @@
 
 TEST(ModuleCloneTest, Clone) {
     // Shader that exercises the bulk of the AST nodes and types.
-    // See also fuzzers/tint_ast_clone_fuzzer.cc for further coverage of cloning.
     Source::File file("test.wgsl", R"(enable f16;
 diagnostic(off, chromium.unreachable_code);
 
diff --git a/src/tint/fuzzers/tint_wgsl_reader_wgsl_writer_fuzzer.cc b/src/tint/lang/wgsl/ast/requires.cc
similarity index 70%
copy from src/tint/fuzzers/tint_wgsl_reader_wgsl_writer_fuzzer.cc
copy to src/tint/lang/wgsl/ast/requires.cc
index 7ea3510..059c35f 100644
--- a/src/tint/fuzzers/tint_wgsl_reader_wgsl_writer_fuzzer.cc
+++ b/src/tint/lang/wgsl/ast/requires.cc
@@ -1,4 +1,4 @@
-// Copyright 2021 The Dawn & Tint Authors
+// Copyright 2023 The Dawn & Tint Authors
 //
 // Redistribution and use in source and binary forms, with or without
 // modification, are permitted provided that the following conditions are met:
@@ -25,19 +25,23 @@
 // OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
 // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
-#include <string>
+#include "src/tint/lang/wgsl/ast/requires.h"
 
-#include "src/tint/fuzzers/fuzzer_init.h"
-#include "src/tint/fuzzers/tint_reader_writer_fuzzer.h"
+#include "src/tint/lang/wgsl/ast/builder.h"
+#include "src/tint/lang/wgsl/ast/clone_context.h"
 
-namespace tint::fuzzers {
+TINT_INSTANTIATE_TYPEINFO(tint::ast::Requires);
 
-extern "C" int LLVMFuzzerTestOneInput(const uint8_t* data, size_t size) {
-    tint::fuzzers::ReaderWriterFuzzer fuzzer(InputFormat::kWGSL, OutputFormat::kWGSL);
-    fuzzer.SetDumpInput(GetCliParams().dump_input);
-    fuzzer.SetEnforceValidity(GetCliParams().enforce_validity);
+namespace tint::ast {
 
-    return fuzzer.Run(data, size);
+Requires::Requires(GenerationID pid, NodeID nid, const Source& src, wgsl::LanguageFeatures feats)
+    : Base(pid, nid, src), features(std::move(feats)) {}
+
+Requires::~Requires() = default;
+
+const Requires* Requires::Clone(CloneContext& ctx) const {
+    auto src = ctx.Clone(source);
+    return ctx.dst->create<Requires>(src, features);
 }
 
-}  // namespace tint::fuzzers
+}  // namespace tint::ast
diff --git a/src/tint/lang/wgsl/ast/requires.h b/src/tint/lang/wgsl/ast/requires.h
new file mode 100644
index 0000000..c94d6b0
--- /dev/null
+++ b/src/tint/lang/wgsl/ast/requires.h
@@ -0,0 +1,68 @@
+// Copyright 2023 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+//    list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+//    this list of conditions and the following disclaimer in the documentation
+//    and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+//    contributors may be used to endorse or promote products derived from
+//    this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#ifndef SRC_TINT_LANG_WGSL_AST_REQUIRES_H_
+#define SRC_TINT_LANG_WGSL_AST_REQUIRES_H_
+
+#include <string>
+#include <utility>
+#include <vector>
+
+#include "src/tint/lang/wgsl/ast/node.h"
+#include "src/tint/lang/wgsl/language_feature.h"
+
+namespace tint::ast {
+
+/// A "requires" directive. Example:
+/// ```
+///   // Require a language feature named "foo"
+///   requires foo;
+/// ```
+class Requires final : public Castable<Requires, Node> {
+  public:
+    /// Create a requires directive
+    /// @param pid the identifier of the program that owns this node
+    /// @param nid the unique node identifier
+    /// @param src the source of this node
+    /// @param feats the language features being required by this directive
+    Requires(GenerationID pid, NodeID nid, const Source& src, wgsl::LanguageFeatures feats);
+
+    /// Destructor
+    ~Requires() override;
+
+    /// Clones this node and all transitive child nodes using the `CloneContext` `ctx`.
+    /// @param ctx the clone context
+    /// @return the newly cloned node
+    const Requires* Clone(CloneContext& ctx) const override;
+
+    /// The features being required by this directive.
+    const wgsl::LanguageFeatures features;
+};
+
+}  // namespace tint::ast
+
+#endif  // SRC_TINT_LANG_WGSL_AST_REQUIRES_H_
diff --git a/src/tint/fuzzers/tint_wgsl_reader_wgsl_writer_fuzzer.cc b/src/tint/lang/wgsl/ast/requires_test.cc
similarity index 64%
copy from src/tint/fuzzers/tint_wgsl_reader_wgsl_writer_fuzzer.cc
copy to src/tint/lang/wgsl/ast/requires_test.cc
index 7ea3510..8650f3f 100644
--- a/src/tint/fuzzers/tint_wgsl_reader_wgsl_writer_fuzzer.cc
+++ b/src/tint/lang/wgsl/ast/requires_test.cc
@@ -1,4 +1,4 @@
-// Copyright 2021 The Dawn & Tint Authors
+// Copyright 2023 The Dawn & Tint Authors
 //
 // Redistribution and use in source and binary forms, with or without
 // modification, are permitted provided that the following conditions are met:
@@ -25,19 +25,27 @@
 // OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
 // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
-#include <string>
+#include "src/tint/lang/wgsl/ast/requires.h"
 
-#include "src/tint/fuzzers/fuzzer_init.h"
-#include "src/tint/fuzzers/tint_reader_writer_fuzzer.h"
+#include "src/tint/lang/wgsl/ast/helper_test.h"
 
-namespace tint::fuzzers {
+namespace tint::ast {
+namespace {
 
-extern "C" int LLVMFuzzerTestOneInput(const uint8_t* data, size_t size) {
-    tint::fuzzers::ReaderWriterFuzzer fuzzer(InputFormat::kWGSL, OutputFormat::kWGSL);
-    fuzzer.SetDumpInput(GetCliParams().dump_input);
-    fuzzer.SetEnforceValidity(GetCliParams().enforce_validity);
+using RequiresTest = TestHelper;
 
-    return fuzzer.Run(data, size);
+TEST_F(RequiresTest, Creation) {
+    auto* req = Require(Source{{{20, 2}, {20, 5}}},
+                        wgsl::LanguageFeature::kReadonlyAndReadwriteStorageTextures);
+    EXPECT_EQ(req->source.range.begin.line, 20u);
+    EXPECT_EQ(req->source.range.begin.column, 2u);
+    EXPECT_EQ(req->source.range.end.line, 20u);
+    EXPECT_EQ(req->source.range.end.column, 5u);
+    ASSERT_EQ(req->features.Length(), 1u);
+    EXPECT_EQ(req->features[0], wgsl::LanguageFeature::kReadonlyAndReadwriteStorageTextures);
+    ASSERT_EQ(AST().Requires().Length(), 1u);
+    EXPECT_EQ(AST().Requires()[0], req);
 }
 
-}  // namespace tint::fuzzers
+}  // namespace
+}  // namespace tint::ast
diff --git a/src/tint/lang/wgsl/ast/transform/single_entry_point.cc b/src/tint/lang/wgsl/ast/transform/single_entry_point.cc
index 44e4b86..daae6f8 100644
--- a/src/tint/lang/wgsl/ast/transform/single_entry_point.cc
+++ b/src/tint/lang/wgsl/ast/transform/single_entry_point.cc
@@ -128,6 +128,10 @@
                 }
             },
             [&](const Enable* ext) { b.AST().AddEnable(ctx.Clone(ext)); },
+            [&](const Requires*) {
+                // Drop requires directives as they are optional, and it's non-trivial to determine
+                // which features are needed for which entry points.
+            },
             [&](const DiagnosticDirective* d) { b.AST().AddDiagnosticDirective(ctx.Clone(d)); },  //
             TINT_ICE_ON_NO_MATCH);
     }
diff --git a/src/tint/lang/wgsl/ast/transform/single_entry_point_test.cc b/src/tint/lang/wgsl/ast/transform/single_entry_point_test.cc
index 609f97b..f22ddb5 100644
--- a/src/tint/lang/wgsl/ast/transform/single_entry_point_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/single_entry_point_test.cc
@@ -644,5 +644,30 @@
     EXPECT_EQ(src, str(got));
 }
 
+TEST_F(SingleEntryPointTest, Requires) {
+    // Make sure that requires are handled (and dropped).
+    auto* src = R"(
+requires readonly_and_readwrite_storage_textures;
+
+@compute @workgroup_size(1)
+fn main() {
+}
+)";
+
+    auto* expect = R"(
+@compute @workgroup_size(1)
+fn main() {
+}
+)";
+
+    SingleEntryPoint::Config cfg("main");
+
+    DataMap data;
+    data.Add<SingleEntryPoint::Config>(cfg);
+    auto got = Run<SingleEntryPoint>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
 }  // namespace
 }  // namespace tint::ast::transform
diff --git a/src/tint/lang/wgsl/language_feature.cc b/src/tint/lang/wgsl/language_feature.cc
new file mode 100644
index 0000000..a5e0430
--- /dev/null
+++ b/src/tint/lang/wgsl/language_feature.cc
@@ -0,0 +1,61 @@
+// Copyright 2023 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+//    list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+//    this list of conditions and the following disclaimer in the documentation
+//    and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+//    contributors may be used to endorse or promote products derived from
+//    this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by 'tools/src/cmd/gen' using the template:
+//   src/tint/lang/wgsl/language_feature.cc.tmpl
+//
+// To regenerate run: './tools/run gen'
+//
+//                       Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+#include "src/tint/lang/wgsl/language_feature.h"
+
+namespace tint::wgsl {
+
+/// ParseLanguageFeature parses a LanguageFeature from a string.
+/// @param str the string to parse
+/// @returns the parsed enum, or LanguageFeature::kUndefined if the string could not be parsed.
+LanguageFeature ParseLanguageFeature(std::string_view str) {
+    if (str == "readonly_and_readwrite_storage_textures") {
+        return LanguageFeature::kReadonlyAndReadwriteStorageTextures;
+    }
+    return LanguageFeature::kUndefined;
+}
+
+std::string_view ToString(LanguageFeature value) {
+    switch (value) {
+        case LanguageFeature::kUndefined:
+            return "undefined";
+        case LanguageFeature::kReadonlyAndReadwriteStorageTextures:
+            return "readonly_and_readwrite_storage_textures";
+    }
+    return "<unknown>";
+}
+
+}  // namespace tint::wgsl
diff --git a/src/tint/lang/wgsl/language_feature.cc.tmpl b/src/tint/lang/wgsl/language_feature.cc.tmpl
new file mode 100644
index 0000000..782b515
--- /dev/null
+++ b/src/tint/lang/wgsl/language_feature.cc.tmpl
@@ -0,0 +1,23 @@
+{{- /*
+--------------------------------------------------------------------------------
+Template file for use with tools/src/cmd/gen to generate language_feature.cc
+
+See:
+* tools/src/cmd/gen for structures used by this template
+* https://golang.org/pkg/text/template/ for documentation on the template syntax
+--------------------------------------------------------------------------------
+*/ -}}
+
+{{- $I := LoadIntrinsics "src/tint/lang/wgsl/wgsl.def" -}}
+{{- Import "src/tint/utils/templates/enums.tmpl.inc" -}}
+{{- $enum := ($I.Sem.Enum "language_feature") -}}
+
+#include "src/tint/lang/wgsl/language_feature.h"
+
+namespace tint::wgsl {
+
+{{ Eval "ParseEnum" $enum}}
+
+{{ Eval "EnumOStream" $enum}}
+
+}  // namespace tint::wgsl
diff --git a/src/tint/lang/wgsl/language_feature.h b/src/tint/lang/wgsl/language_feature.h
new file mode 100644
index 0000000..0a5573a
--- /dev/null
+++ b/src/tint/lang/wgsl/language_feature.h
@@ -0,0 +1,78 @@
+// Copyright 2023 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+//    list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+//    this list of conditions and the following disclaimer in the documentation
+//    and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+//    contributors may be used to endorse or promote products derived from
+//    this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by 'tools/src/cmd/gen' using the template:
+//   src/tint/lang/wgsl/language_feature.h.tmpl
+//
+// To regenerate run: './tools/run gen'
+//
+//                       Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+#ifndef SRC_TINT_LANG_WGSL_LANGUAGE_FEATURE_H_
+#define SRC_TINT_LANG_WGSL_LANGUAGE_FEATURE_H_
+
+#include "src/tint/utils/containers/unique_vector.h"
+#include "src/tint/utils/traits/traits.h"
+
+namespace tint::wgsl {
+
+/// An enumerator of WGSL language features
+/// @see src/tint/lang/wgsl/intrinsics.def for language feature descriptions
+enum class LanguageFeature : uint8_t {
+    kUndefined,
+    kReadonlyAndReadwriteStorageTextures,
+};
+
+/// @param value the enum value
+/// @returns the string for the given enum value
+std::string_view ToString(LanguageFeature value);
+
+/// @param out the stream to write to
+/// @param value the LanguageFeature
+/// @returns @p out so calls can be chained
+template <typename STREAM, typename = traits::EnableIfIsOStream<STREAM>>
+auto& operator<<(STREAM& out, LanguageFeature value) {
+    return out << ToString(value);
+}
+
+/// ParseLanguageFeature parses a LanguageFeature from a string.
+/// @param str the string to parse
+/// @returns the parsed enum, or LanguageFeature::kUndefined if the string could not be parsed.
+LanguageFeature ParseLanguageFeature(std::string_view str);
+
+constexpr const char* kLanguageFeatureStrings[] = {
+    "readonly_and_readwrite_storage_textures",
+};
+
+// A unique vector of language features
+using LanguageFeatures = UniqueVector<LanguageFeature, 4>;
+
+}  // namespace tint::wgsl
+
+#endif  // SRC_TINT_LANG_WGSL_LANGUAGE_FEATURE_H_
diff --git a/src/tint/lang/wgsl/language_feature.h.tmpl b/src/tint/lang/wgsl/language_feature.h.tmpl
new file mode 100644
index 0000000..cbe9f78
--- /dev/null
+++ b/src/tint/lang/wgsl/language_feature.h.tmpl
@@ -0,0 +1,32 @@
+{{- /*
+--------------------------------------------------------------------------------
+Template file for use with tools/src/cmd/gen to generate language_feature.h
+
+See:
+* tools/src/cmd/gen for structures used by this template
+* https://golang.org/pkg/text/template/ for documentation on the template syntax
+--------------------------------------------------------------------------------
+*/ -}}
+
+{{- $I := LoadIntrinsics "src/tint/lang/wgsl/wgsl.def" -}}
+{{- Import "src/tint/utils/templates/enums.tmpl.inc" -}}
+{{- $enum := ($I.Sem.Enum "language_feature") -}}
+
+#ifndef SRC_TINT_LANG_WGSL_LANGUAGE_FEATURE_H_
+#define SRC_TINT_LANG_WGSL_LANGUAGE_FEATURE_H_
+
+#include "src/tint/utils/traits/traits.h"
+#include "src/tint/utils/containers/unique_vector.h"
+
+namespace tint::wgsl {
+
+/// An enumerator of WGSL language features
+/// @see src/tint/lang/wgsl/intrinsics.def for language feature descriptions
+{{ Eval "DeclareEnum" $enum}}
+
+// A unique vector of language features
+using LanguageFeatures = UniqueVector<LanguageFeature, 4>;
+
+}  // namespace tint::wgsl
+
+#endif  // SRC_TINT_LANG_WGSL_LANGUAGE_FEATURE_H_
diff --git a/src/tint/lang/wgsl/language_feature_test.cc b/src/tint/lang/wgsl/language_feature_test.cc
new file mode 100644
index 0000000..4597ce9
--- /dev/null
+++ b/src/tint/lang/wgsl/language_feature_test.cc
@@ -0,0 +1,94 @@
+// Copyright 2023 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+//    list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+//    this list of conditions and the following disclaimer in the documentation
+//    and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+//    contributors may be used to endorse or promote products derived from
+//    this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by 'tools/src/cmd/gen' using the template:
+//   src/tint/lang/wgsl/language_feature_test.cc.tmpl
+//
+// To regenerate run: './tools/run gen'
+//
+//                       Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+#include "src/tint/lang/wgsl/language_feature.h"
+
+#include <gtest/gtest.h>
+
+#include <string>
+
+#include "src/tint/utils/text/string.h"
+
+namespace tint::wgsl {
+namespace {
+
+namespace parse_print_tests {
+
+struct Case {
+    const char* string;
+    LanguageFeature value;
+};
+
+inline std::ostream& operator<<(std::ostream& out, Case c) {
+    return out << "'" << std::string(c.string) << "'";
+}
+
+static constexpr Case kValidCases[] = {
+    {"readonly_and_readwrite_storage_textures",
+     LanguageFeature::kReadonlyAndReadwriteStorageTextures},
+};
+
+static constexpr Case kInvalidCases[] = {
+    {"eadonly_and_readwrite_stccrage_textures", LanguageFeature::kUndefined},
+    {"rladonly_a3readrite_storage_textures", LanguageFeature::kUndefined},
+    {"readonly_and_readwriVe_storage_textures", LanguageFeature::kUndefined},
+};
+
+using LanguageFeatureParseTest = testing::TestWithParam<Case>;
+
+TEST_P(LanguageFeatureParseTest, Parse) {
+    const char* string = GetParam().string;
+    LanguageFeature expect = GetParam().value;
+    EXPECT_EQ(expect, ParseLanguageFeature(string));
+}
+
+INSTANTIATE_TEST_SUITE_P(ValidCases, LanguageFeatureParseTest, testing::ValuesIn(kValidCases));
+INSTANTIATE_TEST_SUITE_P(InvalidCases, LanguageFeatureParseTest, testing::ValuesIn(kInvalidCases));
+
+using LanguageFeaturePrintTest = testing::TestWithParam<Case>;
+
+TEST_P(LanguageFeaturePrintTest, Print) {
+    LanguageFeature value = GetParam().value;
+    const char* expect = GetParam().string;
+    EXPECT_EQ(expect, tint::ToString(value));
+}
+
+INSTANTIATE_TEST_SUITE_P(ValidCases, LanguageFeaturePrintTest, testing::ValuesIn(kValidCases));
+
+}  // namespace parse_print_tests
+
+}  // namespace
+}  // namespace tint::wgsl
diff --git a/src/tint/lang/wgsl/language_feature_test.cc.tmpl b/src/tint/lang/wgsl/language_feature_test.cc.tmpl
new file mode 100644
index 0000000..fe98798
--- /dev/null
+++ b/src/tint/lang/wgsl/language_feature_test.cc.tmpl
@@ -0,0 +1,29 @@
+{{- /*
+--------------------------------------------------------------------------------
+Template file for use with tools/src/cmd/gen to generate language_feature_test.cc
+
+See:
+* tools/src/cmd/gen for structures used by this template
+* https://golang.org/pkg/text/template/ for documentation on the template syntax
+--------------------------------------------------------------------------------
+*/ -}}
+
+{{- $I := LoadIntrinsics "src/tint/lang/wgsl/wgsl.def" -}}
+{{- Import "src/tint/utils/templates/enums.tmpl.inc" -}}
+{{- $enum := ($I.Sem.Enum "language_feature") -}}
+
+#include "src/tint/lang/wgsl/language_feature.h"
+
+#include <gtest/gtest.h>
+
+#include <string>
+
+#include "src/tint/utils/text/string.h"
+
+namespace tint::wgsl {
+namespace {
+
+{{ Eval "TestParsePrintEnum" $enum}}
+
+}  // namespace
+}  // namespace tint::wgsl
diff --git a/src/tint/lang/wgsl/program/BUILD.bazel b/src/tint/lang/wgsl/program/BUILD.bazel
index 0cc90fa..1906aad 100644
--- a/src/tint/lang/wgsl/program/BUILD.bazel
+++ b/src/tint/lang/wgsl/program/BUILD.bazel
@@ -111,3 +111,21 @@
   visibility = ["//visibility:public"],
 )
 
+alias(
+  name = "tint_build_wgsl_reader",
+  actual = "//src/tint:tint_build_wgsl_reader_true",
+)
+
+alias(
+  name = "tint_build_wgsl_writer",
+  actual = "//src/tint:tint_build_wgsl_writer_true",
+)
+
+selects.config_setting_group(
+    name = "tint_build_wgsl_reader_and_tint_build_wgsl_writer",
+    match_all = [
+        ":tint_build_wgsl_reader",
+        ":tint_build_wgsl_writer",
+    ],
+)
+
diff --git a/src/tint/lang/wgsl/program/BUILD.cmake b/src/tint/lang/wgsl/program/BUILD.cmake
index 07c09e0..eea928d 100644
--- a/src/tint/lang/wgsl/program/BUILD.cmake
+++ b/src/tint/lang/wgsl/program/BUILD.cmake
@@ -109,3 +109,54 @@
 tint_target_add_external_dependencies(tint_lang_wgsl_program_test test
   "gtest"
 )
+
+################################################################################
+# Target:    tint_lang_wgsl_program_fuzz
+# Kind:      fuzz
+################################################################################
+tint_add_target(tint_lang_wgsl_program_fuzz fuzz
+)
+
+tint_target_add_dependencies(tint_lang_wgsl_program_fuzz fuzz
+  tint_api_common
+  tint_lang_core
+  tint_lang_core_constant
+  tint_lang_core_type
+  tint_lang_wgsl
+  tint_lang_wgsl_ast
+  tint_lang_wgsl_program
+  tint_lang_wgsl_resolver
+  tint_lang_wgsl_sem
+  tint_utils_containers
+  tint_utils_diagnostic
+  tint_utils_ice
+  tint_utils_id
+  tint_utils_macros
+  tint_utils_math
+  tint_utils_memory
+  tint_utils_reflection
+  tint_utils_result
+  tint_utils_rtti
+  tint_utils_symbol
+  tint_utils_text
+  tint_utils_traits
+)
+
+if(TINT_BUILD_WGSL_READER)
+  tint_target_add_dependencies(tint_lang_wgsl_program_fuzz fuzz
+    tint_cmd_fuzz_wgsl_fuzz
+    tint_lang_wgsl_reader_parser
+  )
+endif(TINT_BUILD_WGSL_READER)
+
+if(TINT_BUILD_WGSL_READER AND TINT_BUILD_WGSL_WRITER)
+  tint_target_add_sources(tint_lang_wgsl_program_fuzz fuzz
+    "lang/wgsl/program/clone_context_fuzz.cc"
+  )
+endif(TINT_BUILD_WGSL_READER AND TINT_BUILD_WGSL_WRITER)
+
+if(TINT_BUILD_WGSL_WRITER)
+  tint_target_add_dependencies(tint_lang_wgsl_program_fuzz fuzz
+    tint_lang_wgsl_writer
+  )
+endif(TINT_BUILD_WGSL_WRITER)
diff --git a/src/tint/lang/wgsl/program/BUILD.gn b/src/tint/lang/wgsl/program/BUILD.gn
index 6cf7d8b..5641e61 100644
--- a/src/tint/lang/wgsl/program/BUILD.gn
+++ b/src/tint/lang/wgsl/program/BUILD.gn
@@ -109,3 +109,46 @@
     ]
   }
 }
+
+tint_fuzz_source_set("fuzz") {
+  sources = []
+  deps = [
+    "${tint_src_dir}/api/common",
+    "${tint_src_dir}/lang/core",
+    "${tint_src_dir}/lang/core/constant",
+    "${tint_src_dir}/lang/core/type",
+    "${tint_src_dir}/lang/wgsl",
+    "${tint_src_dir}/lang/wgsl/ast",
+    "${tint_src_dir}/lang/wgsl/program",
+    "${tint_src_dir}/lang/wgsl/resolver",
+    "${tint_src_dir}/lang/wgsl/sem",
+    "${tint_src_dir}/utils/containers",
+    "${tint_src_dir}/utils/diagnostic",
+    "${tint_src_dir}/utils/ice",
+    "${tint_src_dir}/utils/id",
+    "${tint_src_dir}/utils/macros",
+    "${tint_src_dir}/utils/math",
+    "${tint_src_dir}/utils/memory",
+    "${tint_src_dir}/utils/reflection",
+    "${tint_src_dir}/utils/result",
+    "${tint_src_dir}/utils/rtti",
+    "${tint_src_dir}/utils/symbol",
+    "${tint_src_dir}/utils/text",
+    "${tint_src_dir}/utils/traits",
+  ]
+
+  if (tint_build_wgsl_reader) {
+    deps += [
+      "${tint_src_dir}/cmd/fuzz/wgsl:fuzz",
+      "${tint_src_dir}/lang/wgsl/reader/parser",
+    ]
+  }
+
+  if (tint_build_wgsl_reader && tint_build_wgsl_writer) {
+    sources += [ "clone_context_fuzz.cc" ]
+  }
+
+  if (tint_build_wgsl_writer) {
+    deps += [ "${tint_src_dir}/lang/wgsl/writer" ]
+  }
+}
diff --git a/src/tint/lang/wgsl/program/clone_context_fuzz.cc b/src/tint/lang/wgsl/program/clone_context_fuzz.cc
new file mode 100644
index 0000000..e4cfc2f
--- /dev/null
+++ b/src/tint/lang/wgsl/program/clone_context_fuzz.cc
@@ -0,0 +1,94 @@
+// Copyright 2020 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+//    list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+//    this list of conditions and the following disclaimer in the documentation
+//    and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+//    contributors may be used to endorse or promote products derived from
+//    this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+// GEN_BUILD:CONDITION(tint_build_wgsl_reader && tint_build_wgsl_writer)
+
+#include <string>
+#include <unordered_set>
+
+#include "src/tint/cmd/fuzz/wgsl/wgsl_fuzz.h"
+#include "src/tint/lang/wgsl/reader/parser/parser.h"
+#include "src/tint/lang/wgsl/writer/writer.h"
+
+namespace tint::program {
+
+#define ASSERT_EQ(A, B)                                         \
+    do {                                                        \
+        decltype(A) assert_a = (A);                             \
+        decltype(B) assert_b = (B);                             \
+        if (assert_a != assert_b) {                             \
+            TINT_ICE() << "ASSERT_EQ(" #A ", " #B ") failed:\n" \
+                       << #A << " was: " << assert_a << "\n"    \
+                       << #B << " was: " << assert_b << "\n";   \
+        }                                                       \
+    } while (false)
+
+#define ASSERT_TRUE(A)                                                                           \
+    do {                                                                                         \
+        decltype(A) assert_a = (A);                                                              \
+        if (!assert_a) {                                                                         \
+            TINT_ICE() << "ASSERT_TRUE(" #A ") failed:\n" << #A << " was: " << assert_a << "\n"; \
+        }                                                                                        \
+    } while (false)
+
+void CloneContextFuzzer(const tint::Program& src) {
+    // Clone the src program to dst
+    tint::Program dst(src.Clone());
+
+    // Expect the printed strings to match
+    ASSERT_EQ(tint::Program::printer(src), tint::Program::printer(dst));
+
+    // Check that none of the AST nodes or type pointers in dst are found in src
+    std::unordered_set<const tint::ast::Node*> src_nodes;
+    for (auto* src_node : src.ASTNodes().Objects()) {
+        src_nodes.emplace(src_node);
+    }
+    std::unordered_set<const tint::core::type::Type*> src_types;
+    for (auto* src_type : src.Types()) {
+        src_types.emplace(src_type);
+    }
+    for (auto* dst_node : dst.ASTNodes().Objects()) {
+        ASSERT_EQ(src_nodes.count(dst_node), 0u);
+    }
+    for (auto* dst_type : dst.Types()) {
+        ASSERT_EQ(src_types.count(dst_type), 0u);
+    }
+
+    tint::wgsl::writer::Options wgsl_options;
+
+    auto src_wgsl = tint::wgsl::writer::Generate(src, wgsl_options);
+    ASSERT_TRUE(src_wgsl);
+
+    auto dst_wgsl = tint::wgsl::writer::Generate(dst, wgsl_options);
+    ASSERT_TRUE(dst_wgsl);
+
+    ASSERT_EQ(src_wgsl->wgsl, dst_wgsl->wgsl);
+}
+
+}  // namespace tint::program
+
+TINT_WGSL_PROGRAM_FUZZER(tint::program::CloneContextFuzzer);
diff --git a/src/tint/lang/wgsl/reader/parser/parser.cc b/src/tint/lang/wgsl/reader/parser/parser.cc
index 40d1484..54aeb27 100644
--- a/src/tint/lang/wgsl/reader/parser/parser.cc
+++ b/src/tint/lang/wgsl/reader/parser/parser.cc
@@ -456,6 +456,7 @@
 //  : require identifier (COMMA identifier)* COMMA? SEMICOLON
 Maybe<Void> Parser::requires_directive() {
     return sync(Token::Type::kSemicolon, [&]() -> Maybe<Void> {
+        MultiTokenSource decl_source(this);
         if (!match(Token::Type::kRequires)) {
             return Failure::kNoMatch;
         }
@@ -473,34 +474,43 @@
             return add_error(t.source(), "requires directives don't take parenthesis");
         }
 
+        wgsl::LanguageFeatures features;
         while (continue_parsing()) {
-            auto& t2 = peek();
-
-            // Match the require name.
+            auto& t2 = next();
             if (handle_error(t2)) {
                 // The token might itself be an error.
                 return Failure::kErrored;
             }
 
+            // Match the require name.
             if (t2.IsIdentifier()) {
-                // TODO(dsinclair): When there are actual values for a requires directive they
-                // should be checked here.
-
-                // Any identifer is a valid feature name, so we correctly handle new feature
-                // names getting added in the future, they just all get flagged as not supported.
-                return add_error(t2.source(), "feature '" + t2.to_str() + "' is not supported");
-            }
-            if (t2.Is(Token::Type::kSemicolon)) {
-                break;
-            }
-            if (!match(Token::Type::kComma)) {
+                auto feature = wgsl::ParseLanguageFeature(t2.to_str_view());
+                if (feature == LanguageFeature::kUndefined) {
+                    // Any identifier is a valid feature name, so we correctly handle new feature
+                    // names getting added in the future, they just all get flagged as not
+                    // supported.
+                    return add_error(t2.source(), "feature '" + t2.to_str() + "' is not supported");
+                }
+                features.Add(feature);
+            } else {
                 return add_error(t2.source(), "invalid feature name for requires");
             }
+
+            if (!match(Token::Type::kComma)) {
+                break;
+            }
+            if (peek_is(Token::Type::kSemicolon)) {
+                break;
+            }
         }
-        // TODO(dsinclair): When there are actual values for a requires directive then the
-        // `while` will need to keep track if any were seen, and this needs to become
-        // conditional.
-        return add_error(t.source(), "missing feature names in requires directive");
+
+        if (!expect("requires directive", Token::Type::kSemicolon)) {
+            return Failure::kErrored;
+        }
+
+        builder_.AST().AddRequires(
+            create<ast::Requires>(decl_source.Source(), std::move(features)));
+        return kSuccess;
     });
 }
 
diff --git a/src/tint/lang/wgsl/reader/parser/require_directive_test.cc b/src/tint/lang/wgsl/reader/parser/require_directive_test.cc
index 43ea6be..8e7ff03 100644
--- a/src/tint/lang/wgsl/reader/parser/require_directive_test.cc
+++ b/src/tint/lang/wgsl/reader/parser/require_directive_test.cc
@@ -33,11 +33,64 @@
 using RequiresDirectiveTest = WGSLParserTest;
 
 // Test a valid require directive.
-// There currently are no valid require directives
-TEST_F(RequiresDirectiveTest, DISABLED_Valid) {
-    auto p = parser("requires <sometime>;");
+TEST_F(RequiresDirectiveTest, Single) {
+    auto p = parser("requires readonly_and_readwrite_storage_textures;");
     p->requires_directive();
     EXPECT_FALSE(p->has_error()) << p->error();
+
+    auto program = p->program();
+    auto& ast = program.AST();
+    ASSERT_EQ(ast.Requires().Length(), 1u);
+    auto* req = ast.Requires()[0];
+    EXPECT_EQ(req->source.range.begin.line, 1u);
+    EXPECT_EQ(req->source.range.begin.column, 1u);
+    EXPECT_EQ(req->source.range.end.line, 1u);
+    EXPECT_EQ(req->source.range.end.column, 50u);
+    ASSERT_EQ(req->features.Length(), 1u);
+    EXPECT_EQ(req->features[0], wgsl::LanguageFeature::kReadonlyAndReadwriteStorageTextures);
+    ASSERT_EQ(ast.GlobalDeclarations().Length(), 1u);
+    EXPECT_EQ(ast.GlobalDeclarations()[0], req);
+}
+
+// Test a valid require directive with a trailing comma.
+TEST_F(RequiresDirectiveTest, Single_TrailingComma) {
+    auto p = parser("requires readonly_and_readwrite_storage_textures,;");
+    p->requires_directive();
+    EXPECT_FALSE(p->has_error()) << p->error();
+
+    auto program = p->program();
+    auto& ast = program.AST();
+    ASSERT_EQ(ast.Requires().Length(), 1u);
+    auto* req = ast.Requires()[0];
+    EXPECT_EQ(req->source.range.begin.line, 1u);
+    EXPECT_EQ(req->source.range.begin.column, 1u);
+    EXPECT_EQ(req->source.range.end.line, 1u);
+    EXPECT_EQ(req->source.range.end.column, 51u);
+    ASSERT_EQ(req->features.Length(), 1u);
+    EXPECT_EQ(req->features[0], wgsl::LanguageFeature::kReadonlyAndReadwriteStorageTextures);
+    ASSERT_EQ(ast.GlobalDeclarations().Length(), 1u);
+    EXPECT_EQ(ast.GlobalDeclarations()[0], req);
+}
+
+TEST_F(RequiresDirectiveTest, Multiple_Repeated) {
+    auto p = parser(
+        "requires readonly_and_readwrite_storage_textures, "
+        "readonly_and_readwrite_storage_textures;");
+    p->requires_directive();
+    EXPECT_FALSE(p->has_error()) << p->error();
+
+    auto program = p->program();
+    auto& ast = program.AST();
+    ASSERT_EQ(ast.Requires().Length(), 1u);
+    auto* req = ast.Requires()[0];
+    EXPECT_EQ(req->source.range.begin.line, 1u);
+    EXPECT_EQ(req->source.range.begin.column, 1u);
+    EXPECT_EQ(req->source.range.end.line, 1u);
+    EXPECT_EQ(req->source.range.end.column, 91u);
+    ASSERT_EQ(req->features.Length(), 1u);
+    EXPECT_EQ(req->features[0], wgsl::LanguageFeature::kReadonlyAndReadwriteStorageTextures);
+    ASSERT_EQ(ast.GlobalDeclarations().Length(), 1u);
+    EXPECT_EQ(ast.GlobalDeclarations()[0], req);
 }
 
 // Test an unknown require identifier.
@@ -76,7 +129,7 @@
         auto p = parser("requires;");
         p->translation_unit();
         EXPECT_TRUE(p->has_error());
-        EXPECT_EQ(p->error(), R"(1:9: missing feature names in requires directive)");
+        EXPECT_EQ(p->error(), R"(1:9: invalid feature name for requires)");
     }
 }
 
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
index cf38225..a0c10a8 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
@@ -86,6 +86,7 @@
 #include "src/tint/lang/wgsl/ast/member_accessor_expression.h"
 #include "src/tint/lang/wgsl/ast/override.h"
 #include "src/tint/lang/wgsl/ast/phony_expression.h"
+#include "src/tint/lang/wgsl/ast/requires.h"
 #include "src/tint/lang/wgsl/ast/return_statement.h"
 #include "src/tint/lang/wgsl/ast/statement.h"
 #include "src/tint/lang/wgsl/ast/struct.h"
@@ -259,6 +260,9 @@
                 [&](const ast::DiagnosticDirective*) {
                     // Ignored for now.
                 },  //
+                [&](const ast::Requires*) {
+                    // Ignored for now.
+                },  //
                 TINT_ICE_ON_NO_MATCH);
         }
 
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc
index 19d448b..f704fa5 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc
@@ -1142,5 +1142,27 @@
 )");
 }
 
+TEST_F(IR_FromProgramTest, Requires) {
+    Require(wgsl::LanguageFeature::kReadonlyAndReadwriteStorageTextures);
+    Func("f", tint::Empty, ty.void_(), tint::Empty);
+
+    auto m = Build();
+    ASSERT_TRUE(m) << m;
+
+    ASSERT_EQ(1u, m->functions.Length());
+
+    auto* f = m->functions[0];
+    ASSERT_NE(f->Block(), nullptr);
+
+    EXPECT_EQ(m->functions[0]->Stage(), core::ir::Function::PipelineStage::kUndefined);
+
+    EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():void -> %b1 {
+  %b1 = block {
+    ret
+  }
+}
+)");
+}
+
 }  // namespace
 }  // namespace tint::wgsl::reader
diff --git a/src/tint/lang/wgsl/resolver/BUILD.bazel b/src/tint/lang/wgsl/resolver/BUILD.bazel
index 719635c..578ead6 100644
--- a/src/tint/lang/wgsl/resolver/BUILD.bazel
+++ b/src/tint/lang/wgsl/resolver/BUILD.bazel
@@ -123,6 +123,7 @@
     "inferred_type_test.cc",
     "is_host_shareable_test.cc",
     "is_storeable_test.cc",
+    "language_features_test.cc",
     "load_test.cc",
     "materialize_test.cc",
     "override_test.cc",
diff --git a/src/tint/lang/wgsl/resolver/BUILD.cmake b/src/tint/lang/wgsl/resolver/BUILD.cmake
index 75280ee..63026bf 100644
--- a/src/tint/lang/wgsl/resolver/BUILD.cmake
+++ b/src/tint/lang/wgsl/resolver/BUILD.cmake
@@ -121,6 +121,7 @@
   lang/wgsl/resolver/inferred_type_test.cc
   lang/wgsl/resolver/is_host_shareable_test.cc
   lang/wgsl/resolver/is_storeable_test.cc
+  lang/wgsl/resolver/language_features_test.cc
   lang/wgsl/resolver/load_test.cc
   lang/wgsl/resolver/materialize_test.cc
   lang/wgsl/resolver/override_test.cc
diff --git a/src/tint/lang/wgsl/resolver/BUILD.gn b/src/tint/lang/wgsl/resolver/BUILD.gn
index b4b367f..86dc084 100644
--- a/src/tint/lang/wgsl/resolver/BUILD.gn
+++ b/src/tint/lang/wgsl/resolver/BUILD.gn
@@ -123,6 +123,7 @@
       "inferred_type_test.cc",
       "is_host_shareable_test.cc",
       "is_storeable_test.cc",
+      "language_features_test.cc",
       "load_test.cc",
       "materialize_test.cc",
       "override_test.cc",
diff --git a/src/tint/lang/wgsl/resolver/dependency_graph.cc b/src/tint/lang/wgsl/resolver/dependency_graph.cc
index a261629..dda9251 100644
--- a/src/tint/lang/wgsl/resolver/dependency_graph.cc
+++ b/src/tint/lang/wgsl/resolver/dependency_graph.cc
@@ -201,6 +201,9 @@
             [&](const ast::Enable*) {
                 // Enable directives do not affect the dependency graph.
             },
+            [&](const ast::Requires*) {
+                // Requires directives do not affect the dependency graph.
+            },
             [&](const ast::ConstAssert* assertion) {
                 TraverseExpression(assertion->condition);
             },  //
@@ -611,6 +614,7 @@
             [&](const ast::Variable* var) { return var->name->symbol; },
             [&](const ast::DiagnosticDirective*) { return Symbol(); },
             [&](const ast::Enable*) { return Symbol(); },
+            [&](const ast::Requires*) { return Symbol(); },
             [&](const ast::ConstAssert*) { return Symbol(); },  //
             TINT_ICE_ON_NO_MATCH);
     }
@@ -714,13 +718,13 @@
 
         // Make sure all directives go before any other global declarations.
         for (auto* global : declaration_order_) {
-            if (global->node->IsAnyOf<ast::DiagnosticDirective, ast::Enable>()) {
+            if (global->node->IsAnyOf<ast::DiagnosticDirective, ast::Enable, ast::Requires>()) {
                 sorted_.Add(global->node);
             }
         }
 
         for (auto* global : declaration_order_) {
-            if (global->node->IsAnyOf<ast::DiagnosticDirective, ast::Enable>()) {
+            if (global->node->IsAnyOf<ast::DiagnosticDirective, ast::Enable, ast::Requires>()) {
                 // Skip directives here, as they are already added.
                 continue;
             }
diff --git a/src/tint/lang/wgsl/resolver/dependency_graph_test.cc b/src/tint/lang/wgsl/resolver/dependency_graph_test.cc
index 160df59..8a1f136 100644
--- a/src/tint/lang/wgsl/resolver/dependency_graph_test.cc
+++ b/src/tint/lang/wgsl/resolver/dependency_graph_test.cc
@@ -1122,9 +1122,12 @@
     auto* enable = Enable(wgsl::Extension::kF16);
     auto* var_2 = GlobalVar("SYMBOL2", ty.f32());
     auto* diagnostic = DiagnosticDirective(wgsl::DiagnosticSeverity::kWarning, "foo");
+    auto* var_3 = GlobalVar("SYMBOL3", ty.u32());
+    auto* req = Require(wgsl::LanguageFeature::kReadonlyAndReadwriteStorageTextures);
 
-    EXPECT_THAT(AST().GlobalDeclarations(), ElementsAre(var_1, enable, var_2, diagnostic));
-    EXPECT_THAT(Build().ordered_globals, ElementsAre(enable, diagnostic, var_1, var_2));
+    EXPECT_THAT(AST().GlobalDeclarations(),
+                ElementsAre(var_1, enable, var_2, diagnostic, var_3, req));
+    EXPECT_THAT(Build().ordered_globals, ElementsAre(enable, diagnostic, req, var_1, var_2, var_3));
 }
 }  // namespace ordered_globals
 
diff --git a/src/tint/fuzzers/tint_wgsl_reader_wgsl_writer_fuzzer.cc b/src/tint/lang/wgsl/resolver/language_features_test.cc
similarity index 69%
copy from src/tint/fuzzers/tint_wgsl_reader_wgsl_writer_fuzzer.cc
copy to src/tint/lang/wgsl/resolver/language_features_test.cc
index 7ea3510..47a1990 100644
--- a/src/tint/fuzzers/tint_wgsl_reader_wgsl_writer_fuzzer.cc
+++ b/src/tint/lang/wgsl/resolver/language_features_test.cc
@@ -1,4 +1,4 @@
-// Copyright 2021 The Dawn & Tint Authors
+// Copyright 2023 The Dawn & Tint Authors
 //
 // Redistribution and use in source and binary forms, with or without
 // modification, are permitted provided that the following conditions are met:
@@ -25,19 +25,24 @@
 // OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
 // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
-#include <string>
+#include "src/tint/lang/wgsl/resolver/resolver.h"
 
-#include "src/tint/fuzzers/fuzzer_init.h"
-#include "src/tint/fuzzers/tint_reader_writer_fuzzer.h"
+#include "gmock/gmock.h"
+#include "src/tint/lang/core/fluent_types.h"
+#include "src/tint/lang/wgsl/resolver/resolver_helper_test.h"
 
-namespace tint::fuzzers {
+using namespace tint::core::number_suffixes;  // NOLINT
+using namespace tint::core::fluent_types;     // NOLINT
 
-extern "C" int LLVMFuzzerTestOneInput(const uint8_t* data, size_t size) {
-    tint::fuzzers::ReaderWriterFuzzer fuzzer(InputFormat::kWGSL, OutputFormat::kWGSL);
-    fuzzer.SetDumpInput(GetCliParams().dump_input);
-    fuzzer.SetEnforceValidity(GetCliParams().enforce_validity);
+namespace tint::resolver {
+namespace {
 
-    return fuzzer.Run(data, size);
+using ResolverLanguageFeaturesTest = ResolverTest;
+
+TEST_F(ResolverLanguageFeaturesTest, Requires) {
+    Require(wgsl::LanguageFeature::kReadonlyAndReadwriteStorageTextures);
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
-}  // namespace tint::fuzzers
+}  // namespace
+}  // namespace tint::resolver
diff --git a/src/tint/lang/wgsl/resolver/resolver.cc b/src/tint/lang/wgsl/resolver/resolver.cc
index 8ea7220..33b4e81 100644
--- a/src/tint/lang/wgsl/resolver/resolver.cc
+++ b/src/tint/lang/wgsl/resolver/resolver.cc
@@ -189,6 +189,7 @@
                     return DiagnosticControl(d->control);
                 },
                 [&](const ast::Enable* e) { return Enable(e); },
+                [&](const ast::Requires* r) { return Requires(r); },
                 [&](const ast::TypeDecl* td) { return TypeDecl(td); },
                 [&](const ast::Function* func) { return Function(func); },
                 [&](const ast::Variable* var) { return GlobalVariable(var); },
@@ -3958,6 +3959,11 @@
     return true;
 }
 
+bool Resolver::Requires(const ast::Requires*) {
+    // TODO(crbug.com/tint/2081): Check that all features are allowed.
+    return true;
+}
+
 core::type::Type* Resolver::TypeDecl(const ast::TypeDecl* named_type) {
     Mark(named_type->name);
 
diff --git a/src/tint/lang/wgsl/resolver/resolver.h b/src/tint/lang/wgsl/resolver/resolver.h
index fa353d4..887b117 100644
--- a/src/tint/lang/wgsl/resolver/resolver.h
+++ b/src/tint/lang/wgsl/resolver/resolver.h
@@ -466,9 +466,13 @@
     bool DiagnosticControl(const ast::DiagnosticControl& control);
 
     /// @param enable the enable declaration
-    /// @returns the resolved extension
+    /// @returns true on success, false on failure
     bool Enable(const ast::Enable* enable);
 
+    /// @param req the requires declaration
+    /// @returns true on success, false on failure
+    bool Requires(const ast::Requires* req);
+
     /// @param named_type the named type to resolve
     /// @returns the resolved semantic type
     core::type::Type* TypeDecl(const ast::TypeDecl* named_type);
diff --git a/src/tint/lang/wgsl/wgsl.def b/src/tint/lang/wgsl/wgsl.def
index 06039b2..a66c2ba 100644
--- a/src/tint/lang/wgsl/wgsl.def
+++ b/src/tint/lang/wgsl/wgsl.def
@@ -90,6 +90,11 @@
   chromium_experimental_pixel_local
 }
 
+// https://gpuweb.github.io/gpuweb/wgsl/#language-extensions-sec
+enum language_feature {
+  readonly_and_readwrite_storage_textures
+}
+
 ////////////////////////////////////////////////////////////////////////////////
 // WGSL primitive types                                                       //
 // Types may be decorated with @precedence(N) to prioritize which type        //
diff --git a/src/tint/lang/wgsl/writer/ast_printer/BUILD.bazel b/src/tint/lang/wgsl/writer/ast_printer/BUILD.bazel
index 16f6447..92a7f93 100644
--- a/src/tint/lang/wgsl/writer/ast_printer/BUILD.bazel
+++ b/src/tint/lang/wgsl/writer/ast_printer/BUILD.bazel
@@ -99,6 +99,7 @@
     "literal_test.cc",
     "loop_test.cc",
     "member_accessor_test.cc",
+    "requires_test.cc",
     "return_test.cc",
     "switch_test.cc",
     "type_test.cc",
@@ -142,6 +143,11 @@
 )
 
 alias(
+  name = "tint_build_wgsl_reader",
+  actual = "//src/tint:tint_build_wgsl_reader_true",
+)
+
+alias(
   name = "tint_build_wgsl_writer",
   actual = "//src/tint:tint_build_wgsl_writer_true",
 )
diff --git a/src/tint/lang/wgsl/writer/ast_printer/BUILD.cmake b/src/tint/lang/wgsl/writer/ast_printer/BUILD.cmake
index b65e2a0..a3abc50 100644
--- a/src/tint/lang/wgsl/writer/ast_printer/BUILD.cmake
+++ b/src/tint/lang/wgsl/writer/ast_printer/BUILD.cmake
@@ -102,6 +102,7 @@
   lang/wgsl/writer/ast_printer/literal_test.cc
   lang/wgsl/writer/ast_printer/loop_test.cc
   lang/wgsl/writer/ast_printer/member_accessor_test.cc
+  lang/wgsl/writer/ast_printer/requires_test.cc
   lang/wgsl/writer/ast_printer/return_test.cc
   lang/wgsl/writer/ast_printer/switch_test.cc
   lang/wgsl/writer/ast_printer/type_test.cc
@@ -146,4 +147,52 @@
   )
 endif(TINT_BUILD_WGSL_WRITER)
 
+endif(TINT_BUILD_WGSL_WRITER)
+if(TINT_BUILD_WGSL_WRITER)
+################################################################################
+# Target:    tint_lang_wgsl_writer_ast_printer_fuzz
+# Kind:      fuzz
+# Condition: TINT_BUILD_WGSL_WRITER
+################################################################################
+tint_add_target(tint_lang_wgsl_writer_ast_printer_fuzz fuzz
+)
+
+tint_target_add_dependencies(tint_lang_wgsl_writer_ast_printer_fuzz fuzz
+  tint_lang_core
+  tint_lang_core_constant
+  tint_lang_core_type
+  tint_lang_wgsl
+  tint_lang_wgsl_ast
+  tint_lang_wgsl_program
+  tint_lang_wgsl_sem
+  tint_utils_containers
+  tint_utils_diagnostic
+  tint_utils_generator
+  tint_utils_ice
+  tint_utils_id
+  tint_utils_macros
+  tint_utils_math
+  tint_utils_memory
+  tint_utils_result
+  tint_utils_rtti
+  tint_utils_symbol
+  tint_utils_text
+  tint_utils_traits
+)
+
+if(TINT_BUILD_WGSL_READER)
+  tint_target_add_sources(tint_lang_wgsl_writer_ast_printer_fuzz fuzz
+    "lang/wgsl/writer/ast_printer/ast_printer_fuzz.cc"
+  )
+  tint_target_add_dependencies(tint_lang_wgsl_writer_ast_printer_fuzz fuzz
+    tint_cmd_fuzz_wgsl_fuzz
+  )
+endif(TINT_BUILD_WGSL_READER)
+
+if(TINT_BUILD_WGSL_WRITER)
+  tint_target_add_dependencies(tint_lang_wgsl_writer_ast_printer_fuzz fuzz
+    tint_lang_wgsl_writer_ast_printer
+  )
+endif(TINT_BUILD_WGSL_WRITER)
+
 endif(TINT_BUILD_WGSL_WRITER)
\ No newline at end of file
diff --git a/src/tint/lang/wgsl/writer/ast_printer/BUILD.gn b/src/tint/lang/wgsl/writer/ast_printer/BUILD.gn
index 0f53128..ab5e562 100644
--- a/src/tint/lang/wgsl/writer/ast_printer/BUILD.gn
+++ b/src/tint/lang/wgsl/writer/ast_printer/BUILD.gn
@@ -101,6 +101,7 @@
         "literal_test.cc",
         "loop_test.cc",
         "member_accessor_test.cc",
+        "requires_test.cc",
         "return_test.cc",
         "switch_test.cc",
         "type_test.cc",
@@ -141,3 +142,39 @@
     }
   }
 }
+if (tint_build_wgsl_writer) {
+  tint_fuzz_source_set("fuzz") {
+    sources = []
+    deps = [
+      "${tint_src_dir}/lang/core",
+      "${tint_src_dir}/lang/core/constant",
+      "${tint_src_dir}/lang/core/type",
+      "${tint_src_dir}/lang/wgsl",
+      "${tint_src_dir}/lang/wgsl/ast",
+      "${tint_src_dir}/lang/wgsl/program",
+      "${tint_src_dir}/lang/wgsl/sem",
+      "${tint_src_dir}/utils/containers",
+      "${tint_src_dir}/utils/diagnostic",
+      "${tint_src_dir}/utils/generator",
+      "${tint_src_dir}/utils/ice",
+      "${tint_src_dir}/utils/id",
+      "${tint_src_dir}/utils/macros",
+      "${tint_src_dir}/utils/math",
+      "${tint_src_dir}/utils/memory",
+      "${tint_src_dir}/utils/result",
+      "${tint_src_dir}/utils/rtti",
+      "${tint_src_dir}/utils/symbol",
+      "${tint_src_dir}/utils/text",
+      "${tint_src_dir}/utils/traits",
+    ]
+
+    if (tint_build_wgsl_reader) {
+      sources += [ "ast_printer_fuzz.cc" ]
+      deps += [ "${tint_src_dir}/cmd/fuzz/wgsl:fuzz" ]
+    }
+
+    if (tint_build_wgsl_writer) {
+      deps += [ "${tint_src_dir}/lang/wgsl/writer/ast_printer" ]
+    }
+  }
+}
diff --git a/src/tint/lang/wgsl/writer/ast_printer/ast_printer.cc b/src/tint/lang/wgsl/writer/ast_printer/ast_printer.cc
index 7ff0c0d..ec38e3d 100644
--- a/src/tint/lang/wgsl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/wgsl/writer/ast_printer/ast_printer.cc
@@ -102,6 +102,10 @@
         EmitEnable(enable);
         has_directives = true;
     }
+    for (auto req : program_.AST().Requires()) {
+        EmitRequires(req);
+        has_directives = true;
+    }
     for (auto diagnostic : program_.AST().DiagnosticDirectives()) {
         auto out = Line();
         EmitDiagnosticControl(out, diagnostic->control);
@@ -113,7 +117,7 @@
     }
     // Generate global declarations in the order they appear in the module.
     for (auto* decl : program_.AST().GlobalDeclarations()) {
-        if (decl->IsAnyOf<ast::DiagnosticDirective, ast::Enable>()) {
+        if (decl->IsAnyOf<ast::DiagnosticDirective, ast::Enable, ast::Requires>()) {
             continue;
         }
         Switch(
@@ -148,6 +152,20 @@
     out << ";";
 }
 
+void ASTPrinter::EmitRequires(const ast::Requires* req) {
+    auto out = Line();
+    out << "requires ";
+    bool first = true;
+    for (auto feature : req->features) {
+        if (!first) {
+            out << ", ";
+        }
+        out << feature;
+        first = false;
+    }
+    out << ";";
+}
+
 void ASTPrinter::EmitTypeDecl(const ast::TypeDecl* ty) {
     Switch(
         ty,  //
diff --git a/src/tint/lang/wgsl/writer/ast_printer/ast_printer.h b/src/tint/lang/wgsl/writer/ast_printer/ast_printer.h
index 9fe6ca3..71bb09a 100644
--- a/src/tint/lang/wgsl/writer/ast_printer/ast_printer.h
+++ b/src/tint/lang/wgsl/writer/ast_printer/ast_printer.h
@@ -67,6 +67,7 @@
 class LiteralExpression;
 class LoopStatement;
 class MemberAccessorExpression;
+class Requires;
 class ReturnStatement;
 class Statement;
 class Statement;
@@ -104,6 +105,9 @@
     /// Handles generating an enable directive
     /// @param enable the enable node
     void EmitEnable(const ast::Enable* enable);
+    /// Handles generating a requires directive
+    /// @param req the requires node
+    void EmitRequires(const ast::Requires* req);
     /// Handles generating a declared type
     /// @param ty the declared type to generate
     void EmitTypeDecl(const ast::TypeDecl* ty);
diff --git a/src/tint/fuzzers/tint_wgsl_reader_wgsl_writer_fuzzer.cc b/src/tint/lang/wgsl/writer/ast_printer/ast_printer_fuzz.cc
similarity index 75%
rename from src/tint/fuzzers/tint_wgsl_reader_wgsl_writer_fuzzer.cc
rename to src/tint/lang/wgsl/writer/ast_printer/ast_printer_fuzz.cc
index 7ea3510..77022be 100644
--- a/src/tint/fuzzers/tint_wgsl_reader_wgsl_writer_fuzzer.cc
+++ b/src/tint/lang/wgsl/writer/ast_printer/ast_printer_fuzz.cc
@@ -25,19 +25,20 @@
 // OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
 // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
-#include <string>
+// GEN_BUILD:CONDITION(tint_build_wgsl_reader)
 
-#include "src/tint/fuzzers/fuzzer_init.h"
-#include "src/tint/fuzzers/tint_reader_writer_fuzzer.h"
+#include "src/tint/lang/wgsl/writer/ast_printer/ast_printer.h"
 
-namespace tint::fuzzers {
+#include "src/tint/cmd/fuzz/wgsl/wgsl_fuzz.h"
 
-extern "C" int LLVMFuzzerTestOneInput(const uint8_t* data, size_t size) {
-    tint::fuzzers::ReaderWriterFuzzer fuzzer(InputFormat::kWGSL, OutputFormat::kWGSL);
-    fuzzer.SetDumpInput(GetCliParams().dump_input);
-    fuzzer.SetEnforceValidity(GetCliParams().enforce_validity);
+namespace tint::wgsl::writer {
+namespace {
 
-    return fuzzer.Run(data, size);
+void ASTPrinterFuzzer(const tint::Program& program) {
+    ASTPrinter{program}.Generate();
 }
 
-}  // namespace tint::fuzzers
+}  // namespace
+}  // namespace tint::wgsl::writer
+
+TINT_WGSL_PROGRAM_FUZZER(tint::wgsl::writer::ASTPrinterFuzzer);
diff --git a/src/tint/lang/wgsl/writer/ast_printer/requires_test.cc b/src/tint/lang/wgsl/writer/ast_printer/requires_test.cc
new file mode 100644
index 0000000..1526fbb
--- /dev/null
+++ b/src/tint/lang/wgsl/writer/ast_printer/requires_test.cc
@@ -0,0 +1,66 @@
+// Copyright 2023 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+//    list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+//    this list of conditions and the following disclaimer in the documentation
+//    and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+//    contributors may be used to endorse or promote products derived from
+//    this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/wgsl/writer/ast_printer/helper_test.h"
+
+#include "gmock/gmock.h"
+
+namespace tint::wgsl::writer {
+namespace {
+
+using WgslASTPrinterTest = TestHelper;
+
+TEST_F(WgslASTPrinterTest, Emit_Requires) {
+    auto* req = Require(wgsl::LanguageFeature::kReadonlyAndReadwriteStorageTextures);
+
+    ASTPrinter& gen = Build();
+
+    gen.EmitRequires(req);
+    EXPECT_THAT(gen.Diagnostics(), testing::IsEmpty());
+    EXPECT_EQ(gen.Result(), R"(requires readonly_and_readwrite_storage_textures;
+)");
+}
+
+// TODO(jrprice): Enable this once we have multiple language features.
+TEST_F(WgslASTPrinterTest, DISABLED_Emit_Requires_Multiple) {
+    auto* req = create<ast::Requires>(
+        wgsl::LanguageFeatures({wgsl::LanguageFeature::kReadonlyAndReadwriteStorageTextures,
+                                wgsl::LanguageFeature::kReadonlyAndReadwriteStorageTextures}));
+    AST().AddRequires(req);
+
+    ASTPrinter& gen = Build();
+
+    gen.EmitRequires(req);
+    EXPECT_THAT(gen.Diagnostics(), testing::IsEmpty());
+    EXPECT_EQ(
+        gen.Result(),
+        R"(requires readonly_and_readwrite_storage_textures, readonly_and_readwrite_storage_textures;
+)");
+}
+
+}  // namespace
+}  // namespace tint::wgsl::writer