diff --git a/src/tint/lang/wgsl/resolver/BUILD.bazel b/src/tint/lang/wgsl/resolver/BUILD.bazel
index 9893abb..e5ea130 100644
--- a/src/tint/lang/wgsl/resolver/BUILD.bazel
+++ b/src/tint/lang/wgsl/resolver/BUILD.bazel
@@ -125,6 +125,7 @@
     "host_shareable_validation_test.cc",
     "increment_decrement_validation_test.cc",
     "inferred_type_test.cc",
+    "input_attachments_extension_test.cc",
     "is_host_shareable_test.cc",
     "is_storeable_test.cc",
     "language_features_test.cc",
diff --git a/src/tint/lang/wgsl/resolver/BUILD.cmake b/src/tint/lang/wgsl/resolver/BUILD.cmake
index 429b9c0..326af91 100644
--- a/src/tint/lang/wgsl/resolver/BUILD.cmake
+++ b/src/tint/lang/wgsl/resolver/BUILD.cmake
@@ -123,6 +123,7 @@
   lang/wgsl/resolver/host_shareable_validation_test.cc
   lang/wgsl/resolver/increment_decrement_validation_test.cc
   lang/wgsl/resolver/inferred_type_test.cc
+  lang/wgsl/resolver/input_attachments_extension_test.cc
   lang/wgsl/resolver/is_host_shareable_test.cc
   lang/wgsl/resolver/is_storeable_test.cc
   lang/wgsl/resolver/language_features_test.cc
diff --git a/src/tint/lang/wgsl/resolver/BUILD.gn b/src/tint/lang/wgsl/resolver/BUILD.gn
index 9456dbb..6fb89d0 100644
--- a/src/tint/lang/wgsl/resolver/BUILD.gn
+++ b/src/tint/lang/wgsl/resolver/BUILD.gn
@@ -125,6 +125,7 @@
       "host_shareable_validation_test.cc",
       "increment_decrement_validation_test.cc",
       "inferred_type_test.cc",
+      "input_attachments_extension_test.cc",
       "is_host_shareable_test.cc",
       "is_storeable_test.cc",
       "language_features_test.cc",
diff --git a/src/tint/lang/wgsl/resolver/attribute_validation_test.cc b/src/tint/lang/wgsl/resolver/attribute_validation_test.cc
index e94ebbe..809c8e3 100644
--- a/src/tint/lang/wgsl/resolver/attribute_validation_test.cc
+++ b/src/tint/lang/wgsl/resolver/attribute_validation_test.cc
@@ -65,6 +65,7 @@
     kDiagnostic,
     kGroup,
     kId,
+    kInputAttachmentIndex,
     kInterpolate,
     kInvariant,
     kLocation,
@@ -93,6 +94,8 @@
             return o << "@group";
         case AttributeKind::kId:
             return o << "@id";
+        case AttributeKind::kInputAttachmentIndex:
+            return o << "@input_attachment_index";
         case AttributeKind::kInterpolate:
             return o << "@interpolate";
         case AttributeKind::kInvariant:
@@ -166,6 +169,10 @@
                 "1:2 error: '@id' is not valid for " + thing,
             },
             TestParams{
+                {AttributeKind::kInputAttachmentIndex},
+                "1:2 error: '@input_attachment_index' is not valid for " + thing,
+            },
+            TestParams{
                 {AttributeKind::kInterpolate},
                 "1:2 error: '@interpolate' is not valid for " + thing,
             },
@@ -230,6 +237,8 @@
             return builder.Group(source, 1_a);
         case AttributeKind::kId:
             return builder.Id(source, 0_a);
+        case AttributeKind::kInputAttachmentIndex:
+            return builder.InputAttachmentIndex(source, 2_a);
         case AttributeKind::kBlendSrc:
             return builder.BlendSrc(source, 0_a);
         case AttributeKind::kInterpolate:
@@ -264,6 +273,9 @@
             case AttributeKind::kBlendSrc:
                 Enable(wgsl::Extension::kChromiumInternalDualSourceBlending);
                 break;
+            case AttributeKind::kInputAttachmentIndex:
+                Enable(wgsl::Extension::kChromiumInternalInputAttachments);
+                break;
             default:
                 break;
         }
@@ -345,6 +357,10 @@
             R"(1:2 error: '@id' is not valid for functions)",
         },
         TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for functions)",
+        },
+        TestParams{
             {AttributeKind::kInterpolate},
             R"(1:2 error: '@interpolate' is not valid for functions)",
         },
@@ -393,81 +409,86 @@
 
     CHECK();
 }
-INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
-                         NonVoidFunctionAttributeTest,
-                         testing::Values(
-                             TestParams{
-                                 {AttributeKind::kAlign},
-                                 R"(1:2 error: '@align' is not valid for functions)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kBinding},
-                                 R"(1:2 error: '@binding' is not valid for functions)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kBlendSrc},
-                                 R"(1:2 error: '@blend_src' is not valid for functions)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kBuiltinPosition},
-                                 R"(1:2 error: '@builtin' is not valid for functions)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kColor},
-                                 R"(1:2 error: '@color' is not valid for functions)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kDiagnostic},
-                                 Pass,
-                             },
-                             TestParams{
-                                 {AttributeKind::kGroup},
-                                 R"(1:2 error: '@group' is not valid for functions)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kId},
-                                 R"(1:2 error: '@id' is not valid for functions)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kInterpolate},
-                                 R"(1:2 error: '@interpolate' is not valid for functions)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kInvariant},
-                                 R"(1:2 error: '@invariant' is not valid for functions)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kLocation},
-                                 R"(1:2 error: '@location' is not valid for functions)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kMustUse},
-                                 Pass,
-                             },
-                             TestParams{
-                                 {AttributeKind::kOffset},
-                                 R"(1:2 error: '@offset' is not valid for functions)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kSize},
-                                 R"(1:2 error: '@size' is not valid for functions)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kStageCompute},
-                                 R"(9:9 error: missing entry point IO attribute on return type)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kStageCompute, AttributeKind::kWorkgroupSize},
-                                 R"(9:9 error: missing entry point IO attribute on return type)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kStride},
-                                 R"(1:2 error: '@stride' is not valid for functions)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kWorkgroupSize},
-                                 R"(1:2 error: '@workgroup_size' is only valid for compute stages)",
-                             }));
+INSTANTIATE_TEST_SUITE_P(
+    ResolverAttributeValidationTest,
+    NonVoidFunctionAttributeTest,
+    testing::Values(
+        TestParams{
+            {AttributeKind::kAlign},
+            R"(1:2 error: '@align' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kBinding},
+            R"(1:2 error: '@binding' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kBlendSrc},
+            R"(1:2 error: '@blend_src' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kBuiltinPosition},
+            R"(1:2 error: '@builtin' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kColor},
+            R"(1:2 error: '@color' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kDiagnostic},
+            Pass,
+        },
+        TestParams{
+            {AttributeKind::kGroup},
+            R"(1:2 error: '@group' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kId},
+            R"(1:2 error: '@id' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kInterpolate},
+            R"(1:2 error: '@interpolate' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kInvariant},
+            R"(1:2 error: '@invariant' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kLocation},
+            R"(1:2 error: '@location' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kMustUse},
+            Pass,
+        },
+        TestParams{
+            {AttributeKind::kOffset},
+            R"(1:2 error: '@offset' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kSize},
+            R"(1:2 error: '@size' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kStageCompute},
+            R"(9:9 error: missing entry point IO attribute on return type)",
+        },
+        TestParams{
+            {AttributeKind::kStageCompute, AttributeKind::kWorkgroupSize},
+            R"(9:9 error: missing entry point IO attribute on return type)",
+        },
+        TestParams{
+            {AttributeKind::kStride},
+            R"(1:2 error: '@stride' is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kWorkgroupSize},
+            R"(1:2 error: '@workgroup_size' is only valid for compute stages)",
+        }));
 }  // namespace FunctionTests
 
 namespace FunctionInputAndOutputTests {
@@ -520,6 +541,10 @@
             R"(1:2 error: '@id' is not valid for function parameters)",
         },
         TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for function parameters)",
+        },
+        TestParams{
             {AttributeKind::kInterpolate},
             R"(1:2 error: '@interpolate' is not valid for non-entry point function parameters)",
         },
@@ -605,6 +630,10 @@
             R"(1:2 error: '@id' is not valid for non-entry point function return types)",
         },
         TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for non-entry point function return types)",
+        },
+        TestParams{
             {AttributeKind::kInterpolate},
             R"(1:2 error: '@interpolate' is not valid for non-entry point function return types)",
         },
@@ -695,6 +724,10 @@
             R"(1:2 error: '@id' is not valid for function parameters)",
         },
         TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for function parameters)",
+        },
+        TestParams{
             {AttributeKind::kInterpolate},
             R"(1:2 error: '@interpolate' cannot be used by compute shaders)",
         },
@@ -784,6 +817,10 @@
             R"(1:2 error: '@id' is not valid for function parameters)",
         },
         TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for function parameters)",
+        },
+        TestParams{
             {AttributeKind::kInterpolate},
             R"(9:9 error: missing entry point IO attribute on parameter)",
         },
@@ -887,6 +924,10 @@
             R"(1:2 error: '@id' is not valid for function parameters)",
         },
         TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for function parameters)",
+        },
+        TestParams{
             {AttributeKind::kInterpolate},
             R"(9:9 error: missing entry point IO attribute on parameter)",
         },
@@ -992,6 +1033,10 @@
             R"(1:2 error: '@id' is not valid for entry point return types)",
         },
         TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for entry point return types)",
+        },
+        TestParams{
             {AttributeKind::kInterpolate},
             R"(1:2 error: '@interpolate' cannot be used by compute shaders)",
         },
@@ -1082,6 +1127,10 @@
             R"(1:2 error: '@id' is not valid for entry point return types)",
         },
         TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for entry point return types)",
+        },
+        TestParams{
             {AttributeKind::kInterpolate},
             R"(9:9 error: missing entry point IO attribute on return type)",
         },
@@ -1190,6 +1239,10 @@
             R"(1:2 error: '@id' is not valid for entry point return types)",
         },
         TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for entry point return types)",
+        },
+        TestParams{
             {AttributeKind::kInterpolate},
             R"(1:2 error: '@interpolate' can only be used with '@location')",
         },
@@ -1315,6 +1368,10 @@
             R"(1:2 error: '@id' is not valid for 'struct' declarations)",
         },
         TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for 'struct' declarations)",
+        },
+        TestParams{
             {AttributeKind::kInterpolate},
             R"(1:2 error: '@interpolate' is not valid for 'struct' declarations)",
         },
@@ -1399,6 +1456,10 @@
             R"(1:2 error: '@id' is not valid for 'struct' members)",
         },
         TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for 'struct' members)",
+        },
+        TestParams{
             {AttributeKind::kInterpolate},
             R"(1:2 error: '@interpolate' can only be used with '@location')",
         },
@@ -1635,82 +1696,87 @@
 
     CHECK();
 }
-INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
-                         ArrayAttributeTest,
-                         testing::Values(
-                             TestParams{
-                                 {AttributeKind::kAlign},
-                                 R"(1:2 error: '@align' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kBinding},
-                                 R"(1:2 error: '@binding' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kBlendSrc},
-                                 R"(1:2 error: '@blend_src' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kBuiltinPosition},
-                                 R"(1:2 error: '@builtin' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kDiagnostic},
-                                 R"(1:2 error: '@diagnostic' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kGroup},
-                                 R"(1:2 error: '@group' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kId},
-                                 R"(1:2 error: '@id' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kInterpolate},
-                                 R"(1:2 error: '@interpolate' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kInvariant},
-                                 R"(1:2 error: '@invariant' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kLocation},
-                                 R"(1:2 error: '@location' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kMustUse},
-                                 R"(1:2 error: '@must_use' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kOffset},
-                                 R"(1:2 error: '@offset' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kSize},
-                                 R"(1:2 error: '@size' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kStageCompute},
-                                 R"(1:2 error: '@compute' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kStride},
-                                 Pass,
-                             },
-                             TestParams{
-                                 {AttributeKind::kWorkgroupSize},
-                                 R"(1:2 error: '@workgroup_size' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kBinding, AttributeKind::kGroup},
-                                 R"(1:2 error: '@binding' is not valid for 'array' types)",
-                             },
-                             TestParams{
-                                 {AttributeKind::kStride, AttributeKind::kStride},
-                                 R"(3:4 error: duplicate stride attribute
+INSTANTIATE_TEST_SUITE_P(
+    ResolverAttributeValidationTest,
+    ArrayAttributeTest,
+    testing::Values(
+        TestParams{
+            {AttributeKind::kAlign},
+            R"(1:2 error: '@align' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kBinding},
+            R"(1:2 error: '@binding' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kBlendSrc},
+            R"(1:2 error: '@blend_src' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kBuiltinPosition},
+            R"(1:2 error: '@builtin' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kDiagnostic},
+            R"(1:2 error: '@diagnostic' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kGroup},
+            R"(1:2 error: '@group' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kId},
+            R"(1:2 error: '@id' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kInterpolate},
+            R"(1:2 error: '@interpolate' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kInvariant},
+            R"(1:2 error: '@invariant' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kLocation},
+            R"(1:2 error: '@location' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kMustUse},
+            R"(1:2 error: '@must_use' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kOffset},
+            R"(1:2 error: '@offset' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kSize},
+            R"(1:2 error: '@size' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kStageCompute},
+            R"(1:2 error: '@compute' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kStride},
+            Pass,
+        },
+        TestParams{
+            {AttributeKind::kWorkgroupSize},
+            R"(1:2 error: '@workgroup_size' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kBinding, AttributeKind::kGroup},
+            R"(1:2 error: '@binding' is not valid for 'array' types)",
+        },
+        TestParams{
+            {AttributeKind::kStride, AttributeKind::kStride},
+            R"(3:4 error: duplicate stride attribute
 1:2 note: first attribute declared here)",
-                             }));
+        }));
 
 using VariableAttributeTest = TestWithParams;
 TEST_P(VariableAttributeTest, IsValid) {
@@ -1862,6 +1928,10 @@
             R"(1:2 error: '@id' is not valid for 'const' declaration)",
         },
         TestParams{
+            {AttributeKind::kInputAttachmentIndex},
+            R"(1:2 error: '@input_attachment_index' is not valid for 'const' declaration)",
+        },
+        TestParams{
             {AttributeKind::kInterpolate},
             R"(1:2 error: '@interpolate' is not valid for 'const' declaration)",
         },
diff --git a/src/tint/lang/wgsl/resolver/dependency_graph.cc b/src/tint/lang/wgsl/resolver/dependency_graph.cc
index 4838cde..bffa92e 100644
--- a/src/tint/lang/wgsl/resolver/dependency_graph.cc
+++ b/src/tint/lang/wgsl/resolver/dependency_graph.cc
@@ -52,6 +52,7 @@
 #include "src/tint/lang/wgsl/ast/identifier.h"
 #include "src/tint/lang/wgsl/ast/if_statement.h"
 #include "src/tint/lang/wgsl/ast/increment_decrement_statement.h"
+#include "src/tint/lang/wgsl/ast/input_attachment_index_attribute.h"
 #include "src/tint/lang/wgsl/ast/internal_attribute.h"
 #include "src/tint/lang/wgsl/ast/interpolate_attribute.h"
 #include "src/tint/lang/wgsl/ast/invariant_attribute.h"
@@ -379,6 +380,7 @@
             [&](const ast::ColorAttribute* color) { TraverseExpression(color->expr); },
             [&](const ast::GroupAttribute* group) { TraverseExpression(group->expr); },
             [&](const ast::IdAttribute* id) { TraverseExpression(id->expr); },
+            [&](const ast::InputAttachmentIndexAttribute* idx) { TraverseExpression(idx->expr); },
             [&](const ast::BlendSrcAttribute* index) { TraverseExpression(index->expr); },
             [&](const ast::InterpolateAttribute* interpolate) {
                 TraverseExpression(interpolate->type);
diff --git a/src/tint/lang/wgsl/resolver/dependency_graph_test.cc b/src/tint/lang/wgsl/resolver/dependency_graph_test.cc
index e435301..f64306a 100644
--- a/src/tint/lang/wgsl/resolver/dependency_graph_test.cc
+++ b/src/tint/lang/wgsl/resolver/dependency_graph_test.cc
@@ -1761,6 +1761,7 @@
     GlobalVar(Sym(), ty.sampler(core::type::SamplerKind::kSampler));
 
     GlobalVar(Sym(), ty.i32(), Vector{Binding(V), Group(V)});
+    GlobalVar(Sym(), ty.input_attachment(T), Vector{Binding(V), Group(V), InputAttachmentIndex(V)});
     GlobalVar(Sym(), ty.i32(), Vector{Location(V)});
     Override(Sym(), ty.i32(), Vector{Id(V)});
 
diff --git a/src/tint/lang/wgsl/resolver/input_attachments_extension_test.cc b/src/tint/lang/wgsl/resolver/input_attachments_extension_test.cc
new file mode 100644
index 0000000..2b0449d
--- /dev/null
+++ b/src/tint/lang/wgsl/resolver/input_attachments_extension_test.cc
@@ -0,0 +1,166 @@
+// Copyright 2024 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/resolver/resolver.h"
+#include "src/tint/lang/wgsl/resolver/resolver_helper_test.h"
+
+namespace tint::resolver {
+namespace {
+
+using namespace tint::core::fluent_types;     // NOLINT
+using namespace tint::core::number_suffixes;  // NOLINT
+
+using InputAttachmenExtensionTest = ResolverTest;
+
+// Test that input_attachment cannot be used without extension.
+TEST_F(InputAttachmenExtensionTest, InputAttachmentWithoutExtension) {
+    // @group(0) @binding(0) @input_attachment_index(3)
+    // var input_tex : input_attachment<f32>;
+
+    GlobalVar("input_tex", ty.input_attachment(ty.Of<f32>()),
+              Vector{Binding(0_u), Group(0_u), InputAttachmentIndex(3_u)});
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(
+        r()->error(),
+        R"(error: use of 'input_attachment' requires enabling extension 'chromium_internal_input_attachments')");
+}
+
+// Test that input_attachment cannot be declared locally.
+TEST_F(InputAttachmenExtensionTest, InputAttachmentLocalDecl) {
+    // enable chromium_internal_input_attachments;
+    // @fragment fn f() {
+    //    var input_tex : input_attachment<f32>;
+    // }
+
+    Enable(Source{{12, 34}}, wgsl::Extension::kChromiumInternalInputAttachments);
+
+    Func("f", Empty, ty.void_(),
+         Vector{
+             Decl(Var("input_tex", ty.input_attachment(ty.Of<f32>()))),
+         },
+         Vector{Stage(ast::PipelineStage::kFragment)});
+    EXPECT_FALSE(r()->Resolve());
+}
+
+// Test that input_attachment cannot be declared without index.
+TEST_F(InputAttachmenExtensionTest, InputAttachmentWithoutIndex) {
+    // enable chromium_internal_input_attachments;
+    // @group(0) @binding(0)
+    // var input_tex : input_attachment<f32>;
+
+    Enable(Source{{12, 34}}, wgsl::Extension::kChromiumInternalInputAttachments);
+
+    GlobalVar("input_tex", ty.input_attachment(ty.Of<f32>()), Vector{Binding(0_u), Group(0_u)});
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(),
+              R"(error: 'input_attachment' variables require '@input_attachment_index' attribute)");
+}
+
+// Test that Resolver can get input_attachment_index value.
+TEST_F(InputAttachmenExtensionTest, InputAttachmentIndexValue) {
+    // enable chromium_internal_input_attachments;
+    // @group(0) @binding(0) @input_attachment_index(3)
+    // var input_tex : input_attachment<f32>;
+
+    Enable(Source{{12, 34}}, wgsl::Extension::kChromiumInternalInputAttachments);
+
+    auto* ast_var = GlobalVar("input_tex", ty.input_attachment(ty.Of<f32>()),
+                              Vector{Binding(0_u), Group(0_u), InputAttachmentIndex(3_u)});
+
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* sem_var = Sem().Get<sem::GlobalVariable>(ast_var);
+    ASSERT_NE(sem_var, nullptr);
+    EXPECT_EQ(sem_var->Attributes().input_attachment_index, 3u);
+}
+
+// Test that @input_attachment_index cannot be used without extension.
+TEST_F(InputAttachmenExtensionTest, InputAttachmentIndexWithoutExtension) {
+    // @group(0) @binding(0) @input_attachment_index(3)
+    // var input_tex : texture_2d<f32>;
+
+    GlobalVar("input_tex", ty.sampled_texture(core::type::TextureDimension::k2d, ty.Of<f32>()),
+              Vector{Binding(0_u), Group(0_u), InputAttachmentIndex(3_u)});
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(
+        r()->error(),
+        R"(error: use of '@input_attachment_index' requires enabling extension 'chromium_internal_input_attachments')");
+}
+
+// Test that input_attachment_index's value cannot be float.
+TEST_F(InputAttachmenExtensionTest, InputAttachmentIndexInvalidValueType) {
+    // enable chromium_internal_input_attachments;
+    // @group(0) @binding(0) @input_attachment_index(3.0)
+    // var input_tex : input_attachment<f32>;
+
+    Enable(Source{{12, 34}}, wgsl::Extension::kChromiumInternalInputAttachments);
+
+    GlobalVar("input_tex", ty.input_attachment(ty.Of<f32>()),
+              Vector{Binding(0_u), Group(0_u), InputAttachmentIndex(3_f)});
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(), R"(error: '@input_attachment_index' must be an 'i32' or 'u32' value)");
+}
+
+// Test that input_attachment_index's value cannot be negative.
+TEST_F(InputAttachmenExtensionTest, InputAttachmentIndexNegative) {
+    // enable chromium_internal_input_attachments;
+    // @group(0) @binding(0) @input_attachment_index(-2)
+    // var input_tex : input_attachment<f32>;
+
+    Enable(Source{{12, 34}}, wgsl::Extension::kChromiumInternalInputAttachments);
+
+    GlobalVar("input_tex", ty.input_attachment(ty.Of<f32>()),
+              Vector{Binding(0_u), Group(0_u), InputAttachmentIndex(core::i32(-2))});
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(), R"(error: '@input_attachment_index' value must be non-negative)");
+}
+
+// Test that input_attachment_index cannot be used on non input_attachment variable.
+TEST_F(InputAttachmenExtensionTest, InputAttachmentIndexInvalidType) {
+    // enable chromium_internal_input_attachments;
+    // @group(0) @binding(0) @input_attachment_index(3)
+    // var input_tex : texture_2d<f32>;
+
+    Enable(Source{{12, 34}}, wgsl::Extension::kChromiumInternalInputAttachments);
+
+    GlobalVar("input_tex", ty.sampled_texture(core::type::TextureDimension::k2d, ty.Of<f32>()),
+              Vector{Binding(0_u), Group(0_u), InputAttachmentIndex(3_u)});
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(
+        r()->error(),
+        R"(error: cannot apply '@input_attachment_index' to declaration of type 'texture_2d<f32>'
+note: '@input_attachment_index' must only be applied to declarations of 'input_attachment' type)");
+}
+
+}  // namespace
+}  // namespace tint::resolver
diff --git a/src/tint/lang/wgsl/resolver/resolver.cc b/src/tint/lang/wgsl/resolver/resolver.cc
index 0a17841..fe24ffc 100644
--- a/src/tint/lang/wgsl/resolver/resolver.cc
+++ b/src/tint/lang/wgsl/resolver/resolver.cc
@@ -65,6 +65,7 @@
 #include "src/tint/lang/wgsl/ast/for_loop_statement.h"
 #include "src/tint/lang/wgsl/ast/id_attribute.h"
 #include "src/tint/lang/wgsl/ast/if_statement.h"
+#include "src/tint/lang/wgsl/ast/input_attachment_index_attribute.h"
 #include "src/tint/lang/wgsl/ast/internal_attribute.h"
 #include "src/tint/lang/wgsl/ast/interpolate_attribute.h"
 #include "src/tint/lang/wgsl/ast/loop_statement.h"
@@ -614,7 +615,7 @@
         bool has_io_address_space = sem->AddressSpace() == core::AddressSpace::kIn ||
                                     sem->AddressSpace() == core::AddressSpace::kOut;
 
-        std::optional<uint32_t> group, binding;
+        std::optional<uint32_t> group, binding, input_attachment_index;
         for (auto* attribute : var->attributes) {
             Mark(attribute);
             enum Status { kSuccess, kErrored, kInvalid };
@@ -636,6 +637,14 @@
                     group = value.Get();
                     return kSuccess;
                 },
+                [&](const ast::InputAttachmentIndexAttribute* attr) {
+                    auto value = InputAttachmentIndexAttribute(attr);
+                    if (value != Success) {
+                        return kErrored;
+                    }
+                    input_attachment_index = value.Get();
+                    return kSuccess;
+                },
                 [&](const ast::LocationAttribute* attr) {
                     if (!has_io_address_space) {
                         return kInvalid;
@@ -708,6 +717,10 @@
             global->Attributes().binding_point = BindingPoint{group.value(), binding.value()};
         }
 
+        if (input_attachment_index) {
+            global->Attributes().input_attachment_index = input_attachment_index;
+        }
+
     } else {
         for (auto* attribute : var->attributes) {
             Mark(attribute);
@@ -3876,6 +3889,31 @@
     return static_cast<uint32_t>(value);
 }
 
+tint::Result<uint32_t> Resolver::InputAttachmentIndexAttribute(
+    const ast::InputAttachmentIndexAttribute* attr) {
+    ExprEvalStageConstraint constraint{core::EvaluationStage::kConstant, "@input_attachment_index"};
+    TINT_SCOPED_ASSIGNMENT(expr_eval_stage_constraint_, constraint);
+
+    auto* materialized = Materialize(ValueExpression(attr->expr));
+    if (!materialized) {
+        return Failure{};
+    }
+    if (!materialized->Type()->IsAnyOf<core::type::I32, core::type::U32>()) {
+        AddError(attr->source) << style::Attribute("@input_attachment_index") << " must be an "
+                               << style::Type("i32") << " or " << style::Type("u32") << " value";
+        return Failure{};
+    }
+
+    auto const_value = materialized->ConstantValue();
+    auto value = const_value->ValueAs<AInt>();
+    if (value < 0) {
+        AddError(attr->source) << style::Attribute("@input_attachment_index")
+                               << " value must be non-negative";
+        return Failure{};
+    }
+    return static_cast<uint32_t>(value);
+}
+
 tint::Result<sem::WorkgroupSize> Resolver::WorkgroupAttribute(const ast::WorkgroupAttribute* attr) {
     // Set work-group size defaults.
     sem::WorkgroupSize ws;
diff --git a/src/tint/lang/wgsl/resolver/resolver.h b/src/tint/lang/wgsl/resolver/resolver.h
index df18f91..1781493 100644
--- a/src/tint/lang/wgsl/resolver/resolver.h
+++ b/src/tint/lang/wgsl/resolver/resolver.h
@@ -442,6 +442,11 @@
     /// @returns the group value on success.
     tint::Result<uint32_t> GroupAttribute(const ast::GroupAttribute* attr);
 
+    /// Resolves the `@input_attachment_index` attribute @p attr
+    /// @returns the index value on success.
+    tint::Result<uint32_t> InputAttachmentIndexAttribute(
+        const ast::InputAttachmentIndexAttribute* attr);
+
     /// Resolves the `@workgroup_size` attribute @p attr
     /// @returns the workgroup size on success.
     tint::Result<sem::WorkgroupSize> WorkgroupAttribute(const ast::WorkgroupAttribute* attr);
diff --git a/src/tint/lang/wgsl/resolver/validator.cc b/src/tint/lang/wgsl/resolver/validator.cc
index f482497..559aa50 100644
--- a/src/tint/lang/wgsl/resolver/validator.cc
+++ b/src/tint/lang/wgsl/resolver/validator.cc
@@ -442,6 +442,12 @@
 }
 
 bool Validator::InputAttachment(const core::type::InputAttachment* t, const Source& source) const {
+    if (!enabled_extensions_.Contains(wgsl::Extension::kChromiumInternalInputAttachments)) {
+        AddError(source) << "use of " << style::Type("input_attachment")
+                         << " requires enabling extension "
+                         << style::Code("chromium_internal_input_attachments");
+        return false;
+    }
     if (!t->type()->UnwrapRef()->IsAnyOf<core::type::F32, core::type::I32, core::type::U32>()) {
         AddError(source) << "input_attachment<type>: type must be f32, i32 or u32";
         return false;
@@ -450,6 +456,29 @@
     return true;
 }
 
+bool Validator::InputAttachmentIndexAttribute(const ast::InputAttachmentIndexAttribute* attr,
+                                              const core::type::Type* type,
+                                              const Source& source) const {
+    if (!enabled_extensions_.Contains(wgsl::Extension::kChromiumInternalInputAttachments)) {
+        AddError(source) << "use of " << style::Attribute("@input_attachment_index")
+                         << " requires enabling extension "
+                         << style::Code("chromium_internal_input_attachments");
+        return false;
+    }
+
+    if (!type->Is<core::type::InputAttachment>()) {
+        std::string invalid_type = sem_.TypeNameOf(type);
+        AddError(source) << "cannot apply " << style::Attribute("@input_attachment_index")
+                         << " to declaration of type " << style::Type(invalid_type);
+        AddNote(attr->source) << style::Attribute("@input_attachment_index")
+                              << " must only be applied to declarations of "
+                              << style::Type("input_attachment") << " type";
+        return false;
+    }
+
+    return true;
+}
+
 bool Validator::Materialize(const core::type::Type* to,
                             const core::type::Type* from,
                             const Source& source) const {
@@ -708,6 +737,13 @@
         return false;
     }
 
+    auto* input_attachment_index_attr =
+        ast::GetAttribute<ast::InputAttachmentIndexAttribute>(decl->attributes);
+    if (input_attachment_index_attr &&
+        !InputAttachmentIndexAttribute(input_attachment_index_attr, global->Type()->UnwrapRef(),
+                                       decl->source)) {
+        return false;
+    }
     switch (global->AddressSpace()) {
         case core::AddressSpace::kUniform:
         case core::AddressSpace::kStorage:
@@ -720,6 +756,13 @@
                     << style::Attribute("@binding") << " attributes";
                 return false;
             }
+            if (global->Type()->UnwrapRef()->Is<core::type::InputAttachment>() &&
+                !input_attachment_index_attr) {
+                AddError(decl->source)
+                    << style::Type("input_attachment") << " variables require "
+                    << style::Attribute("@input_attachment_index") << " attribute";
+                return false;
+            }
             break;
         }
         default: {
diff --git a/src/tint/lang/wgsl/resolver/validator.h b/src/tint/lang/wgsl/resolver/validator.h
index 6cdbfaf..fb54048 100644
--- a/src/tint/lang/wgsl/resolver/validator.h
+++ b/src/tint/lang/wgsl/resolver/validator.h
@@ -434,6 +434,15 @@
     /// @returns true on success, false otherwise
     bool InputAttachment(const core::type::InputAttachment* t, const Source& source) const;
 
+    /// Validates a input attachment index attribute
+    /// @param attr the input attachment index attribute to validate
+    /// @param type the variable type
+    /// @param source the source of declaration using the attribute
+    /// @returns true on success, false otherwise.
+    bool InputAttachmentIndexAttribute(const ast::InputAttachmentIndexAttribute* attr,
+                                       const core::type::Type* type,
+                                       const Source& source) const;
+
     /// Validates a structure
     /// @param str the structure to validate
     /// @param stage the current pipeline stage
diff --git a/src/tint/lang/wgsl/sem/variable.h b/src/tint/lang/wgsl/sem/variable.h
index bf97929..51cac2a 100644
--- a/src/tint/lang/wgsl/sem/variable.h
+++ b/src/tint/lang/wgsl/sem/variable.h
@@ -169,6 +169,8 @@
     /// @note a GlobalVariable generally doesn't have a `color` in WGSL, as it isn't allowed by
     /// the spec. The location maybe attached by transforms such as CanonicalizeEntryPointIO.
     std::optional<uint32_t> color;
+    /// The `input_attachment_index` attribute value for the variable, if set
+    std::optional<uint32_t> input_attachment_index;
 };
 
 /// GlobalVariable is a module-scope variable