diff --git a/src/tint/cmd/fuzz/wgsl/dictionary.txt b/src/tint/cmd/fuzz/wgsl/dictionary.txt
index 39c4691..ffd821a 100644
--- a/src/tint/cmd/fuzz/wgsl/dictionary.txt
+++ b/src/tint/cmd/fuzz/wgsl/dictionary.txt
@@ -233,6 +233,7 @@
 "if"
 "info"
 "input_attachment"
+"input_attachment_index"
 "insertBits"
 "instance_index"
 "interpolate"
diff --git a/src/tint/lang/core/attribute.cc b/src/tint/lang/core/attribute.cc
index c6fa693..ead3dae 100644
--- a/src/tint/lang/core/attribute.cc
+++ b/src/tint/lang/core/attribute.cc
@@ -72,6 +72,9 @@
     if (str == "id") {
         return Attribute::kId;
     }
+    if (str == "input_attachment_index") {
+        return Attribute::kInputAttachmentIndex;
+    }
     if (str == "interpolate") {
         return Attribute::kInterpolate;
     }
@@ -120,6 +123,8 @@
             return "group";
         case Attribute::kId:
             return "id";
+        case Attribute::kInputAttachmentIndex:
+            return "input_attachment_index";
         case Attribute::kInterpolate:
             return "interpolate";
         case Attribute::kInvariant:
diff --git a/src/tint/lang/core/attribute.h b/src/tint/lang/core/attribute.h
index 28dd3d4..ecee8bc 100644
--- a/src/tint/lang/core/attribute.h
+++ b/src/tint/lang/core/attribute.h
@@ -59,6 +59,7 @@
     kFragment,
     kGroup,
     kId,
+    kInputAttachmentIndex,
     kInterpolate,
     kInvariant,
     kLocation,
@@ -86,9 +87,24 @@
 Attribute ParseAttribute(std::string_view str);
 
 constexpr std::string_view kAttributeStrings[] = {
-    "align",      "binding",  "blend_src", "builtin", "color",          "compute",
-    "diagnostic", "fragment", "group",     "id",      "interpolate",    "invariant",
-    "location",   "must_use", "size",      "vertex",  "workgroup_size",
+    "align",
+    "binding",
+    "blend_src",
+    "builtin",
+    "color",
+    "compute",
+    "diagnostic",
+    "fragment",
+    "group",
+    "id",
+    "input_attachment_index",
+    "interpolate",
+    "invariant",
+    "location",
+    "must_use",
+    "size",
+    "vertex",
+    "workgroup_size",
 };
 
 }  // namespace tint::core
diff --git a/src/tint/lang/core/attribute_bench.cc b/src/tint/lang/core/attribute_bench.cc
index 924456c..171390e 100644
--- a/src/tint/lang/core/attribute_bench.cc
+++ b/src/tint/lang/core/attribute_bench.cc
@@ -115,55 +115,62 @@
         "5O8",
         "",
         "DDbbB",
-        "interpolaKKe",
-        "33terpolate",
-        "intoott6QQlate",
+        "input_attachment_indKKx",
+        "33nput_attacOmentindex",
+        "ittp6t_attaQQhmeoot_9ndex",
+        "input_attachment_index",
+        "nput_attac66ment_index",
+        "inpt_attachmOnx_inde6zz",
+        "input_attacyyment_index",
+        "nterZZHate",
+        "WWnerpolaq4e",
+        "interOOoate",
         "interpolate",
-        "66terpolate",
-        "intxp66latOz",
-        "yynterpolate",
-        "HHnariZt",
-        "iWW44rianq",
-        "iOOvaiant",
+        "iYohepolte",
+        "inteolat",
+        "interoFate",
+        "invarwnt",
+        "GvrKanff",
+        "KKnqariant",
         "invariant",
-        "ivariYnt",
-        "nvaria",
-        "invaranF",
-        "wocatio",
-        "Kcatoff",
-        "qocKKtion",
+        "invammia3t",
+        "invarint",
+        "inqriant",
+        "obbatibn",
+        "lcaiiio",
+        "loqatiOn",
         "location",
-        "lFcmmt3on",
-        "locaion",
-        "locaton",
-        "ubbt_ube",
-        "mstiius",
-        "muqt_uOe",
+        "loTTatiovv",
+        "FFocation",
+        "fo00QPio",
+        "must_uPe",
+        "ust_uss77",
+        "mustCubbRe",
         "must_use",
-        "muTTt_usvv",
-        "FFust_use",
-        "fu00QPus",
-        "siPe",
-        "sis77",
-        "CiRbbe",
+        "must_usXX",
+        "muOOt_CCse",
+        "sst_usL",
+        "sizX",
+        "sze",
+        "qq",
         "size",
-        "sizXX",
-        "qCCOOO",
-        "szL",
-        "verteX",
-        "verte",
-        "qqrx",
+        "s22ze",
+        "iz0yz",
+        "VPzi",
+        "vrtCnn",
+        "ArqqHHx",
+        "ertex",
         "vertex",
-        "verte22",
-        "vzzyrte",
-        "eriVPx",
-        "wokgroupnnsCze",
-        "workgrouq_sizHA",
-        "workgrup_size",
+        "vefKK",
+        "Pegge",
+        "verte",
+        "4orkgrouTT_Nizc",
+        "worgrouppp7lize",
+        "wgzrkgrup_siNe",
         "workgroup_size",
-        "forroupKKsize",
-        "workrop_Pigge",
-        "workgoup_size",
+        "uuorkgbop_sXXze",
+        "worgroup_size",
+        "worKgroQ_88ze",
     };
     for (auto _ : state) {
         for (auto* str : kStrings) {
diff --git a/src/tint/lang/core/attribute_test.cc b/src/tint/lang/core/attribute_test.cc
index 4d4ed0d..c8eb487 100644
--- a/src/tint/lang/core/attribute_test.cc
+++ b/src/tint/lang/core/attribute_test.cc
@@ -67,6 +67,7 @@
     {"fragment", Attribute::kFragment},
     {"group", Attribute::kGroup},
     {"id", Attribute::kId},
+    {"input_attachment_index", Attribute::kInputAttachmentIndex},
     {"interpolate", Attribute::kInterpolate},
     {"invariant", Attribute::kInvariant},
     {"location", Attribute::kLocation},
@@ -107,27 +108,30 @@
     {"HH", Attribute::kUndefined},
     {"p", Attribute::kUndefined},
     {"1ii", Attribute::kUndefined},
-    {"interpoXXate", Attribute::kUndefined},
-    {"intII99r55olate", Attribute::kUndefined},
-    {"intaarpoSSrHHYe", Attribute::kUndefined},
-    {"kkvHant", Attribute::kUndefined},
-    {"jgaianRR", Attribute::kUndefined},
-    {"inaianb", Attribute::kUndefined},
-    {"locajion", Attribute::kUndefined},
-    {"locaton", Attribute::kUndefined},
-    {"loqion", Attribute::kUndefined},
-    {"mustusNN", Attribute::kUndefined},
-    {"usvvuse", Attribute::kUndefined},
-    {"mut_QQse", Attribute::kUndefined},
-    {"srf", Attribute::kUndefined},
-    {"sije", Attribute::kUndefined},
-    {"NNz2w", Attribute::kUndefined},
-    {"vrtex", Attribute::kUndefined},
-    {"rrertex", Attribute::kUndefined},
-    {"vGrtex", Attribute::kUndefined},
-    {"workgroup_sizFF", Attribute::kUndefined},
-    {"wErkrp_size", Attribute::kUndefined},
-    {"worgrroup_size", Attribute::kUndefined},
+    {"input_attachment_iXXdex", Attribute::kUndefined},
+    {"input_attnnIIhme99t55index", Attribute::kUndefined},
+    {"inputYarrHHacaamentSSindex", Attribute::kUndefined},
+    {"nHkkolate", Attribute::kUndefined},
+    {"jRRterogte", Attribute::kUndefined},
+    {"iterpolbe", Attribute::kUndefined},
+    {"invariajt", Attribute::kUndefined},
+    {"ivariant", Attribute::kUndefined},
+    {"ivariqt", Attribute::kUndefined},
+    {"locaioNN", Attribute::kUndefined},
+    {"ocvvion", Attribute::kUndefined},
+    {"loatQQon", Attribute::kUndefined},
+    {"musf_re", Attribute::kUndefined},
+    {"mujt_use", Attribute::kUndefined},
+    {"muswNN82e", Attribute::kUndefined},
+    {"sze", Attribute::kUndefined},
+    {"sirre", Attribute::kUndefined},
+    {"sGze", Attribute::kUndefined},
+    {"verFFex", Attribute::kUndefined},
+    {"vre", Attribute::kUndefined},
+    {"vrrrte", Attribute::kUndefined},
+    {"orkgroup_sze", Attribute::kUndefined},
+    {"DXkgroJJp_size", Attribute::kUndefined},
+    {"wo8kroup_sz", Attribute::kUndefined},
 };
 
 using AttributeParseTest = testing::TestWithParam<Case>;
diff --git a/src/tint/lang/core/builtin_type.cc b/src/tint/lang/core/builtin_type.cc
index 190cf7f..4fddb22 100644
--- a/src/tint/lang/core/builtin_type.cc
+++ b/src/tint/lang/core/builtin_type.cc
@@ -141,6 +141,9 @@
     if (str == "i32") {
         return BuiltinType::kI32;
     }
+    if (str == "input_attachment") {
+        return BuiltinType::kInputAttachment;
+    }
     if (str == "mat2x2") {
         return BuiltinType::kMat2X2;
     }
@@ -403,6 +406,8 @@
             return "f32";
         case BuiltinType::kI32:
             return "i32";
+        case BuiltinType::kInputAttachment:
+            return "input_attachment";
         case BuiltinType::kMat2X2:
             return "mat2x2";
         case BuiltinType::kMat2X2F:
diff --git a/src/tint/lang/core/builtin_type.h b/src/tint/lang/core/builtin_type.h
index ac0a22e..0f7a9f0 100644
--- a/src/tint/lang/core/builtin_type.h
+++ b/src/tint/lang/core/builtin_type.h
@@ -79,6 +79,7 @@
     kF16,
     kF32,
     kI32,
+    kInputAttachment,
     kMat2X2,
     kMat2X2F,
     kMat2X2H,
@@ -195,6 +196,7 @@
     "f16",
     "f32",
     "i32",
+    "input_attachment",
     "mat2x2",
     "mat2x2f",
     "mat2x2h",
diff --git a/src/tint/lang/core/builtin_type_bench.cc b/src/tint/lang/core/builtin_type_bench.cc
index f963bc7..1e5b673 100644
--- a/src/tint/lang/core/builtin_type_bench.cc
+++ b/src/tint/lang/core/builtin_type_bench.cc
@@ -276,447 +276,454 @@
         "iJJ",
         "UfCDD",
         "i3g",
-        "CCtx",
-        "mt2x2",
-        "mat2x__",
+        "inCCt_atachmen",
+        "iput_attachment",
+        "input__attachmenI",
+        "input_attachment",
+        "inpt_PttttaNment",
+        "input_attddc3ment",
+        "inKutyyattahment",
+        "mtx",
+        "mat03nn",
+        "mvtnCuuK",
         "mat2x2",
-        "attxPP",
-        "mdd32x2",
-        "yyK2x2",
-        "m2uu",
-        "ma0nnx2i",
-        "KanuuCC2f",
-        "mat2x2f",
-        "mlX2x2f",
-        "oat2pp2f",
-        "wwat22f",
+        "lXatx2",
+        "matppxo",
+        "maw2x2",
         "matguum",
         "mt2ma2",
-        "Tat2xZRRh",
+        "Tat2xZRRf",
+        "mat2x2f",
+        "ma8T2xOf",
+        "m0at2x2f",
+        "mBBt2x2f",
+        "at2ppM",
+        "matOO2h",
+        "mGG2x2G",
         "mat2x2h",
-        "ma8T2xOh",
-        "m0at2x2h",
-        "mBBt2x2h",
-        "Matpp",
-        "Oat2x3",
-        "GGG2x3",
+        "mHHt2x11h",
+        "eat2x6Fh",
+        "atx2h",
+        "mKl2ii3",
+        "at2x3",
+        "9IIat2x3",
         "mat2x3",
-        "mHHt2113",
-        "mateF63",
-        "matx",
-        "mat2ii3l",
-        "mt2x3f",
-        "IIvvt2x39",
-        "mat2x3f",
-        "mat23f",
-        "mat2h3f",
-        "mllt2xPzz",
-        "t3h",
-        "mtffxqqh",
+        "ma2x3",
+        "hat2x3",
+        "zzPt2xll",
+        "t3f",
+        "mtffxqqf",
         "mtJJx3dd",
-        "mat2x3h",
-        "mzz2X3h",
+        "mat2x3f",
+        "mzz2X3f",
         "matx32",
-        "maN2yy3h",
-        "atxO",
-        "rauExP",
-        "meet22dd",
+        "maN2yy3f",
+        "mt2O3",
+        "uZtrx3h",
+        "latdde23h",
+        "mat2x3h",
+        "matVVh",
+        "1IIt2x3",
+        "mat2xbh",
+        "t7i4",
+        "maio2x4",
+        "a5x4",
         "mat2x4",
-        "maV92",
-        "maI2x1",
-        "mab2x4",
-        "matzf",
-        "mao2ii4f",
-        "mat45",
+        "St2x4",
+        "m22tx",
+        "maC7G4",
+        "maff284f",
+        "t2x4f",
+        "SSatJJx4f",
         "mat2x4f",
-        "at2xSf",
-        "mat22f",
-        "maG1C4f",
-        "maff284h",
-        "t2x4h",
-        "SSatJJx4h",
+        "atx9f",
+        "maJJbbTT4f",
+        "66a2xf",
+        "ut2x466",
+        "aW2x4h",
+        "mtx4h",
         "mat2x4h",
-        "atx9h",
-        "maJJbbTT4h",
-        "66a2xh",
-        "ma663u",
-        "Wa3x2",
-        "ma32",
+        "mOt2x4",
+        "THat2xBB",
+        "maRR2xh",
+        "VLLa30",
+        "at3KOO",
+        "awg3x2",
         "mat3x2",
-        "ma3x2",
-        "rat3x2",
-        "m2t3xB",
-        "matxBBf",
-        "maRR3xf",
-        "maVV0Lf",
-        "mat3x2f",
-        "a3OOK2f",
-        "magw3xf",
-        "hht3L2f",
-        "aKii3xh",
-        "ma3x2h",
+        "maLhhp",
+        "iitEx2",
+        "at3x2",
         "UUa3882",
-        "mat3x2h",
-        "rrvvt3x2h",
+        "rrvvt3x2f",
         "m3xwmm",
-        "j443x2h",
-        "matXx3",
-        "m8t3x3",
-        "mat3vEE",
+        "mat3x2f",
+        "j443x2f",
+        "matXx2f",
+        "mat8x2f",
+        "mvEE3x2h",
+        "mai93zz",
+        "mQQJJnxGGh",
+        "mat3x2h",
+        "mass3x2h",
+        "matKxPh",
+        "mat3tth",
+        "atx3",
+        "maMM3x3",
+        "mJ0t3x3",
         "mat3x3",
-        "mzzi3x",
-        "maGGQ3JJ3",
-        "mat3ss3",
-        "matKxPf",
-        "mat3ttf",
-        "mt3x3",
-        "mat3x3f",
-        "mMMt3x3f",
-        "maJ03x3f",
-        "V8x3",
-        "maKggx3hh",
-        "maf3x3h",
-        "matQ7x3h",
-        "mat3x3h",
-        "mat3YYh",
+        "V38",
+        "hggat3x3",
+        "maf3x3",
+        "matQ7x3f",
+        "mat3YYf",
         "mak3x3",
+        "mat3x3f",
         "man3x2",
-        "mFFx4",
-        "GGatPPuUU",
-        "mEEFa4",
+        "ma3FFf",
+        "uGGtIUPP3f",
+        "EEvFx3h",
+        "ddtBBDDeh",
+        "m55tMccE3",
+        "mat3x3h",
+        "aKKx3",
+        "mat3x3R",
+        "maDx39",
+        "mt3x4",
+        "aIt34",
+        "mat3x77",
         "mat3x4",
-        "mBet3dd4",
-        "55atExcc",
-        "txKK",
-        "mat3x4R",
-        "maDx49",
-        "mt3x4f",
+        "matIx4",
+        "matd4",
+        "mt3x4",
+        "mtt4f",
+        "ma3XX3x4f",
+        "Eat34f",
         "mat3x4f",
-        "aaat3I",
-        "m77t3x4f",
-        "matIx4f",
-        "md3x4h",
-        "mat34h",
-        "mtt4h",
-        "mat3x4h",
-        "ma3XX3x4h",
-        "Eat34h",
         "maXX3x4",
-        "mxBt4x2",
-        "Wt4x",
-        "mat66x2",
+        "matxBx4f",
+        "m3x4f",
+        "mat3x466",
+        "matxv0",
+        "txh",
+        "mat3x4h",
+        "mpt3xh",
+        "at114h",
+        "EaJ3yBBh",
+        "mqIm4x2",
+        "ma4F2",
+        "aY4x2",
         "mat4x2",
-        "atTv0",
-        "kt",
-        "mpt4x",
-        "at112f",
-        "EaJ4yBBf",
-        "mqIm4x2f",
+        "matDHh2",
+        "24H2",
+        "m4x2",
+        "matx2f",
+        "matx2f",
+        "matddx2f",
         "mat4x2f",
-        "ma4xFf",
-        "Yt4x2f",
-        "mHHtDh2f",
-        "Ht22h",
-        "matx2",
-        "matx2h",
+        "Oat4x2f",
+        "atBbb2",
+        "m004x2f",
+        "mat4xhh",
+        "mgYtx2h",
+        "mat4O2h",
         "mat4x2h",
-        "matx2h",
-        "matddx2h",
-        "Oat4x2h",
-        "bbtB3",
-        "m00tx3",
-        "hat4x3",
+        "ah4xh",
+        "fpaEEx2h",
+        "mavx2h",
+        "at4zz3",
+        "mat4x",
+        "OiiJt4",
         "mat4x3",
-        "matgYx",
-        "Oat4x3",
-        "mhx3",
-        "fpaEEx3f",
-        "mavx3f",
-        "mzztx3f",
+        "mGt4xf",
+        "mTT224x3",
+        "lat4x3",
+        "bat4x3f",
+        "BBatx3f",
+        "PPIXt4S3f",
         "mat4x3f",
-        "ma4x3f",
-        "OOaJxii",
-        "mft4G3f",
-        "mat4x322T",
-        "datlx3h",
-        "bat4x3h",
+        "mjjt4x3f",
+        "_at4xccf",
+        "ax6z4xSSf",
+        "a4xG",
+        "mx44N3v",
+        "atAAx00p",
         "mat4x3h",
-        "BBatx3h",
-        "PPIXt4S3h",
-        "mjjt4x3h",
-        "macc4_4",
-        "SS6zz4xx",
-        "mtx",
+        "eeytyh",
+        "mabWWx0h",
+        "ttatMMxmh",
+        "4d4",
+        "mav_x",
+        "mVt4xEE",
         "mat4x4",
-        "mxxtvN",
-        "AA00t44",
-        "tyexy",
-        "mabWWx0f",
-        "ttatMMxmf",
-        "madf",
+        "m4t4x4",
+        "mVVtgxX",
+        "oat4V4",
+        "ma4x4f",
+        "KKattxf",
+        "G4xf",
         "mat4x4f",
-        "mat_4f",
-        "Vat4EE4f",
-        "mat44f",
-        "mRIxah",
-        "ma4mmh",
-        "at4x4p",
+        "ma4x4",
+        "mad4x4f",
+        "CCat4xPtf",
+        "matx4",
+        "m5t4x4h",
+        "m9jtffx4h",
         "mat4x4h",
-        "mat4xh",
-        "aaxh",
-        "mad4x4h",
-        "pCPtd",
-        "p",
-        "5tr",
+        "matRvxXXh",
+        "ctx4h",
+        "XX8t5x4",
+        "t",
+        "ppcc",
+        "vtr",
         "ptr",
-        "ff99j",
-        "YYvXR",
-        "r",
-        "XX8m5le",
-        "mpler",
-        "sccmlppr",
+        "EESS",
+        "p",
+        "",
+        "saplwwr",
+        "sa99plddr",
+        "99amplP",
         "sampler",
-        "sampver",
-        "EESSmplr",
-        "smplr",
-        "samplecomp_risa",
-        "sampler_co_prwwson",
-        "samplerdd99omparison",
+        "saKKler",
+        "smDDooeM",
+        "smlei",
+        "samler_comparqon",
+        "samplercmparis0LLn",
+        "snvvpler_com66arison",
         "sampler_comparison",
-        "ampler_o99paPPison",
-        "saplerKKcomparison",
-        "saMpler_oomDDarison",
-        "teiie_1B",
-        "txureq1d",
-        "txt00rLL_d",
+        "samplen_comparison",
+        "samlercompeerixxon",
+        "sONNpler_comparCCson",
+        "txture_1d",
+        "tex4uae_1d",
+        "extuNNe_1NN",
         "texture_1d",
-        "tnxture_16vv",
-        "trrxtur_nd",
-        "xxture_eed",
-        "CCNOxture_2d",
-        "txture_2d",
-        "tex4uae_2d",
+        "texture1d",
+        "tuxtre1d",
+        "textErYYS1A",
+        "tex0ure_2d",
+        "texaaure_2d",
+        "tqqmmtur_d",
         "texture_2d",
-        "extuNNe_2NN",
-        "texture2d",
-        "tuxtre2d",
-        "teYYtuAe_2d_arESy",
-        "texture_2d_0rray",
-        "texture_2d_aarray",
-        "texture_2d_array",
-        "texturmm_2d_arra",
-        "texture_2d_aray",
-        "teEuUUe_2darray",
+        "textue_2d",
+        "tUUxtureE",
         "tKKture_Dd",
-        "text__r0_3d",
-        "tAtuel3p",
+        "0exture_2d_a__rat",
+        "tpturl_2_Array",
+        "txture_2M_array",
+        "texture_2d_array",
+        "BBexture_2d_array",
+        "nnxtbbre_2_a99ray",
+        "texturEE2d_aAAray",
+        "tex66ur5_3d",
+        "tHxture_3d",
+        "teHHuxe_3d",
         "texture_3d",
-        "textue_3d",
-        "texturBB_3d",
-        "nnbb99re_3d",
-        "AAEExture_cub",
-        "t66Ttu5e_cube",
-        "textuHe_cube",
+        "tzxturny_0d",
+        "teture_3d",
+        "Hexture3k",
+        "textur_cubc",
+        "trrxtoorecbe",
+        "textreJJcub",
         "texture_cube",
-        "textrexxHcub",
-        "tzx0uryy_cnbe",
-        "texture_cue",
-        "texurH_kube_array",
-        "exture_cube_array",
-        "ooexrrre_cbe_array",
+        "0Cxture_cube",
+        "99eFtureAAcbx",
+        "textcre_cube",
+        "Sexture_cube_array",
+        "txtureBBcobe_arras",
+        "teture_cube_array",
         "texture_cube_array",
-        "textre_cubJJarray",
-        "tCCxtu0e_cube_arry",
-        "texturAAcxbe99aFray",
-        "textcre_depth_2d",
-        "texture_Septh_2d",
-        "textureodpthBB2d",
+        "texturemmebe_array",
+        "tQQturPP_cube_arragg",
+        "textr_cubBrray",
+        "teKKNllre_eth_2d",
+        "texture_ephrr2d",
+        "rppture_depth_2d",
         "texture_depth_2d",
-        "texture_dept_2d",
-        "textummedepth_2d",
-        "toxture_ggeQQtPP2d",
-        "tetur_dptB_2d_rray",
-        "texNure_deKKh2d_arrlly",
-        "teture_dpth_2d_arrray",
+        "texture_deyPth_2d",
+        "extue_dZZpth_ccd",
+        "texure_depth_2",
+        "textu00e_depth_2d_array",
+        "tPJxture_BBsspth_2d_arry",
+        "teffture_wwepth_2dJJarra",
         "texture_depth_2d_array",
-        "texture_epth_ppd_array",
-        "teyturPP_depth_2d_array",
-        "texture_ZZptcc_2d_arry",
-        "txtue_depth_cube",
-        "texture00depth_cube",
-        "texPPuBB_deJth_cusse",
+        "Iextu1e_dpth_2d_arraXX",
+        "extur_depth2_arry",
+        "texture_depth_2_array",
+        "tuxtreKKdepth_cube",
+        "44edture_depmh_ube",
+        "pexture_deoth_cube",
         "texture_depth_cube",
-        "texJJre_dffpwwh_fube",
-        "textIre_depXXh_cub",
-        "textur_ph_cue",
-        "textue_depth_cube_array",
-        "tKKxtue_depth_cube_array",
-        "teture_d44ptm_cube_adray",
+        "textre_djphhhNNcHHbe",
+        "textu33e_depth_EuwwUU",
+        "texture_dethuucbe",
+        "excurrr_depddh_cube_array",
+        "2exturePPttpth_cube_array",
+        "texture_depthwwcub1sarray",
         "texture_depth_cube_array",
-        "pexture_deoth_cube_array",
-        "thhHxtureNdepth_cubejarray",
-        "texwwuUUe_depthEc33be_array",
-        "texture_dept_multiuuampled_2",
-        "ddextKre_depth_ultisampcerr_2d",
-        "textuPPe_depr_multttsample2_2d",
+        "nnexture_depthcc11be_array",
+        "texture_depthcube_array",
+        "texture_depth_cube_arry",
+        "texture_deIth_multsa66pleaaSS2d",
+        "texture_depth_mDltisamEEld_2d",
+        "texture_Iepth_multiccamRleV_2d",
         "texture_depth_multisampled_2d",
-        "1exture_depth_wwsltisampled_2d",
-        "textuce_depth_mnnltisamp11ed_2d",
-        "texture_depth_multisapled_2d",
-        "texture_externl",
-        "teSS66ue_exaaeInal",
-        "textuEEe_extenal",
+        "texture_dephmultisampled92d",
+        "texthraa_depth_multsampled_2d",
+        "texture_depth_SultisaLLped_2d",
+        "txfure_ertermma",
+        "teVturem4xqerna",
+        "textu___extenal",
         "texture_external",
-        "ccexture_exVerIRl",
-        "te9tue_extrnal",
-        "taaxture_exterha",
-        "texture_multisamLLeS_2d",
-        "tefurmm_mutisampled_2d",
-        "texture_mul4isampld_qV",
+        "texurQ_eternal",
+        "textureddetRErnal",
+        "text9re_extenal",
+        "t0xtre_multisamCCe_2d",
+        "texure_zulisampled_2d",
+        "texccure_multisampled_2d",
         "texture_multisampled_2d",
-        "texture_multisa_pled_d",
-        "texure_multisampledQd",
-        "texRRuremultisampledEd2d",
-        "textur_st9rage_1d",
-        "tCCx0ure_strag_1",
-        "textuezstorae_1d",
+        "tOxture_mu_tisampled2QQ",
+        "exture_multsampledtt2d",
+        "textCCEe_mult33samzzled_ppd",
+        "textudde_storaghh_1d",
+        "_etur77_66torage_1d",
+        "texture_storaPe_1d",
         "texture_storage_1d",
-        "texccure_storage_1d",
-        "textureOQQ_orge_1d",
-        "teturettstrage_1d",
-        "textCCrepzzstEr33ge_2d",
-        "textudde_storaghh_2d",
-        "_etur77_66torage_2d",
+        "twxture_storage_1d",
+        "textur_straguu_1",
+        "texture_storaXXe_6d",
+        "textRRr_sorag_2d",
+        "textrestVVrag12d",
+        "texture_HHtorGGge_2d",
         "texture_storage_2d",
-        "texture_storaPe_2d",
-        "twxture_storage_2d",
-        "textur_straguu_2",
-        "textureXXstorage_6d_array",
-        "extuRRestorage_2d_aray",
-        "txtre_storage_2dVVarr1",
+        "tFFxture_storMge_7d",
+        "texture_storage_d",
+        "3xTugge_stoage_2d",
+        "text_rP_Qtorage12d_arrKKy",
+        "tExture_strage_2d_array",
+        "extMre_storage_d_array",
         "texture_storage_2d_array",
-        "GGexture_storHHge_2d_array",
-        "MFFxt7re_storage_2d_array",
-        "texture_storage_d_array",
-        "3xTugge_stoage_3d",
-        "texturP_QtKKrag1__d",
-        "textre_storageE3d",
+        "texturGGst77ragX_2d_SSrray",
+        "textturFF_storae_2d_arraK",
+        "textuoe_storage_2dssZarrUUy",
+        "texturestorage_3d",
+        "telture_storage_3d",
+        "texture_htorage_3d",
         "texture_storage_3d",
-        "tMture_storage_d",
-        "t77xture_sGGorSSe_3d",
-        "txtttre_storage_3FF",
-        "uZZss2",
-        "u2",
-        "u3l",
+        "texturTTKstorage_d",
+        "texturww_storagee3",
+        "textureKsjjragvv_3d",
+        "Y2",
+        "EEI",
+        "u3QQ",
         "u32",
-        "u3h",
-        "uTT",
-        "ww2",
-        "vKvjj",
-        "vYY",
-        "EcI2",
+        "P",
+        "H3ff",
+        "u3n",
+        "Fc66",
+        "ssech",
+        "llFec",
         "vec2",
-        "vecQQ",
-        "Pc",
-        "veffH",
-        "vec2n",
-        "g6F2f",
-        "vssh8f",
-        "vec2f",
-        "veFllf",
-        "00e2j",
-        "gec2f",
-        "vece",
-        "ffc2h",
+        "jj20",
+        "veg2",
+        "vce",
+        "ffc2f",
         "ve",
-        "vec2h",
-        "ve2h",
-        "vqc2h",
+        "ve2f",
+        "vec2f",
+        "vqc2f",
         "AAe",
-        "ec2i",
+        "ec2f",
         "vec2j",
-        "ZZec2i",
-        "vec2i",
+        "ZZec2h",
         "PPecII2",
-        "ZZec2i",
-        "vnnc2i",
+        "vec2h",
+        "ZZec2h",
+        "vnnc2h",
         "HekkZ222",
         "ec2",
         "RcNQQ",
-        "vec2u",
-        "eDu",
+        "eDi",
+        "vec2i",
         "s3c2cu",
-        "vRR2u",
-        "vlJJ3",
-        "MM",
-        "vT63",
-        "vec3",
-        "QQec3",
-        "vuA",
-        "e3",
+        "vRR2i",
+        "vc2JJl",
+        "MMu",
+        "v66T7",
+        "vQQcJu",
+        "vec2u",
+        "Auc2",
+        "ve2",
         "yeq3",
-        "vec3xx",
-        "crr",
-        "vec3f",
-        "v99cf",
-        "vecf",
-        "ecHl",
-        "e_h",
+        "vxxc3",
+        "rr",
+        "e993",
+        "vec3",
+        "ec3",
+        "vlH",
+        "ec",
         "uec3",
-        "vc3h",
-        "vec3h",
-        "EEtmec3h",
+        "vc3f",
+        "EEtmec3f",
+        "vec3f",
         "vec",
         "ec3rr",
-        "xc3i",
+        "xc3f",
         "vezz",
         "vec3e",
-        "vec3i",
         "uc3Zp",
+        "vec3h",
         "00uc7TT",
         "vvJJ",
-        "vecQu",
+        "vecQh",
         "ve3R",
         "e",
-        "vec3u",
         "veprPP",
-        "xxeDD88u",
-        "lldmYYqqu",
-        "FFec4",
-        "rGecNN",
-        "Mecl",
+        "vec3i",
+        "xxeDD88i",
+        "lldmYYqqi",
+        "vFFc__",
+        "rrNNc3u",
+        "leM3u",
+        "c3",
+        "vec3u",
+        "xlc3u",
+        "ec3u",
+        "ae44u",
+        "WWeG4",
+        "vjjc",
+        "vjjc4",
         "vec4",
-        "c",
-        "qxl4",
+        "vj4",
         "ve4",
-        "ae44f",
-        "vec4WW",
-        "vecjj",
-        "vec4f",
-        "vjjc4f",
-        "vj1f",
-        "vc4f",
-        "vec499",
-        "vyVV4h",
+        "99ec4",
+        "vyVV4f",
         "ec4xZ",
-        "vec4h",
-        "v33vvh",
+        "v33vvf",
+        "vec4f",
         "vecs9",
         "veF4",
-        "uec4i",
+        "uec4f",
         "eIKK",
         "ve4J",
-        "vec4i",
-        "vSSCCXXi",
+        "vSSCCXXh",
+        "vec4h",
         "JecWW6ZZ",
         "ecd5",
         "vBBcBU",
         "JJ0c411",
-        "vecttu",
-        "vec4u",
+        "vectti",
         "vttc",
-        "veL4u",
-        "v1c4u",
+        "vec4i",
+        "veL4i",
+        "v1c4i",
+        "veww4i",
+        "ece",
+        "vc4u",
+        "vec4NN",
+        "vec4u",
+        "vUlRR__u",
+        "vecHu",
+        "vrCC4u",
     };
     for (auto _ : state) {
         for (auto* str : kStrings) {
diff --git a/src/tint/lang/core/builtin_type_test.cc b/src/tint/lang/core/builtin_type_test.cc
index 8fd8d6a..886a7e3 100644
--- a/src/tint/lang/core/builtin_type_test.cc
+++ b/src/tint/lang/core/builtin_type_test.cc
@@ -90,6 +90,7 @@
     {"f16", BuiltinType::kF16},
     {"f32", BuiltinType::kF32},
     {"i32", BuiltinType::kI32},
+    {"input_attachment", BuiltinType::kInputAttachment},
     {"mat2x2", BuiltinType::kMat2X2},
     {"mat2x2f", BuiltinType::kMat2X2F},
     {"mat2x2h", BuiltinType::kMat2X2H},
@@ -255,195 +256,198 @@
     {"0yz2", BuiltinType::kUndefined},
     {"iVP", BuiltinType::kUndefined},
     {"Cnn", BuiltinType::kUndefined},
-    {"AtqqHH2", BuiltinType::kUndefined},
-    {"at2x2", BuiltinType::kUndefined},
-    {"mafKK", BuiltinType::kUndefined},
-    {"ltgg2f", BuiltinType::kUndefined},
-    {"mat2xf", BuiltinType::kUndefined},
-    {"NTTtcx4f", BuiltinType::kUndefined},
-    {"ma7ppl2h", BuiltinType::kUndefined},
+    {"nput_attacAAmeHHt", BuiltinType::kUndefined},
+    {"nput_attachment", BuiltinType::kUndefined},
+    {"input_attKKfmen", BuiltinType::kUndefined},
+    {"Paggx", BuiltinType::kUndefined},
+    {"mat2x", BuiltinType::kUndefined},
+    {"maN2c42", BuiltinType::kUndefined},
+    {"ma7ppl2f", BuiltinType::kUndefined},
     {"mNNt2xg", BuiltinType::kUndefined},
-    {"uub2XX2h", BuiltinType::kUndefined},
-    {"mt2x3", BuiltinType::kUndefined},
-    {"m88xK", BuiltinType::kUndefined},
-    {"maqx3", BuiltinType::kUndefined},
-    {"m11t2x3f", BuiltinType::kUndefined},
-    {"22at2iif", BuiltinType::kUndefined},
-    {"at2x377", BuiltinType::kUndefined},
-    {"m2t2xNh", BuiltinType::kUndefined},
-    {"mVVt2x3h", BuiltinType::kUndefined},
-    {"FaWW2w11h", BuiltinType::kUndefined},
-    {"matww4", BuiltinType::kUndefined},
-    {"mat2D4", BuiltinType::kUndefined},
-    {"maKx4", BuiltinType::kUndefined},
-    {"mat21PPhf", BuiltinType::kUndefined},
-    {"mat24f", BuiltinType::kUndefined},
-    {"mYYt2x4f", BuiltinType::kUndefined},
+    {"uub2XX2f", BuiltinType::kUndefined},
+    {"matx2h", BuiltinType::kUndefined},
+    {"Qt882h", BuiltinType::kUndefined},
+    {"mt9q2h", BuiltinType::kUndefined},
+    {"mat2113", BuiltinType::kUndefined},
+    {"Ft2xi22", BuiltinType::kUndefined},
+    {"m7t2x3", BuiltinType::kUndefined},
+    {"m2t2xNf", BuiltinType::kUndefined},
+    {"mVVt2x3f", BuiltinType::kUndefined},
+    {"FaWW2w11f", BuiltinType::kUndefined},
+    {"mawwx3h", BuiltinType::kUndefined},
+    {"Dat2x3h", BuiltinType::kUndefined},
+    {"mt2x3K", BuiltinType::kUndefined},
+    {"11at2xPP", BuiltinType::kUndefined},
+    {"mat2x", BuiltinType::kUndefined},
+    {"mat2xYY", BuiltinType::kUndefined},
     {"mttHH4kk", BuiltinType::kUndefined},
-    {"mat2rr4h", BuiltinType::kUndefined},
-    {"WWas2x4h", BuiltinType::kUndefined},
-    {"maYx2", BuiltinType::kUndefined},
-    {"mq3f2", BuiltinType::kUndefined},
-    {"vvafu222", BuiltinType::kUndefined},
-    {"t3x2f", BuiltinType::kUndefined},
-    {"YYat3f", BuiltinType::kUndefined},
-    {"may3x2EYY", BuiltinType::kUndefined},
-    {"da3xMoh", BuiltinType::kUndefined},
+    {"mat2rr4f", BuiltinType::kUndefined},
+    {"WWas2x4f", BuiltinType::kUndefined},
+    {"Yt2x4h", BuiltinType::kUndefined},
+    {"mt2qfh", BuiltinType::kUndefined},
+    {"mav222xuh", BuiltinType::kUndefined},
+    {"mt32", BuiltinType::kUndefined},
+    {"maY32", BuiltinType::kUndefined},
+    {"YYa7y3E2", BuiltinType::kUndefined},
+    {"da3xMof", BuiltinType::kUndefined},
     {"matMMx2", BuiltinType::kUndefined},
-    {"mat3x55h", BuiltinType::kUndefined},
-    {"maN3", BuiltinType::kUndefined},
-    {"ma33x3", BuiltinType::kUndefined},
-    {"mt3x3", BuiltinType::kUndefined},
-    {"mm66Issf", BuiltinType::kUndefined},
-    {"mat3x1f", BuiltinType::kUndefined},
-    {"Xt3x3", BuiltinType::kUndefined},
-    {"LatIx3h", BuiltinType::kUndefined},
-    {"at3fh", BuiltinType::kUndefined},
+    {"mat3x55f", BuiltinType::kUndefined},
+    {"maN32", BuiltinType::kUndefined},
+    {"ma3Ox33", BuiltinType::kUndefined},
+    {"m3t3x2h", BuiltinType::kUndefined},
+    {"mt3I3", BuiltinType::kUndefined},
+    {"mnnt3xr", BuiltinType::kUndefined},
+    {"mXX", BuiltinType::kUndefined},
+    {"LatIx3f", BuiltinType::kUndefined},
+    {"at3ff", BuiltinType::kUndefined},
     {"mYtURD3", BuiltinType::kUndefined},
-    {"mah3x4", BuiltinType::kUndefined},
-    {"muqII4", BuiltinType::kUndefined},
-    {"mat3xH", BuiltinType::kUndefined},
-    {"at3QQvv", BuiltinType::kUndefined},
-    {"at66ef", BuiltinType::kUndefined},
-    {"ma7O4f", BuiltinType::kUndefined},
+    {"mah3x3h", BuiltinType::kUndefined},
+    {"uuIqt3x", BuiltinType::kUndefined},
+    {"maH3x3h", BuiltinType::kUndefined},
+    {"at3Qvv", BuiltinType::kUndefined},
+    {"66ate", BuiltinType::kUndefined},
+    {"mat7x", BuiltinType::kUndefined},
     {"m55t3x0DD", BuiltinType::kUndefined},
     {"maH3x4II", BuiltinType::kUndefined},
     {"at3x4", BuiltinType::kUndefined},
-    {"ma994x2", BuiltinType::kUndefined},
-    {"mWWt4Gt2", BuiltinType::kUndefined},
-    {"ay42", BuiltinType::kUndefined},
-    {"mt4x2f", BuiltinType::kUndefined},
-    {"IIaBB4x2f", BuiltinType::kUndefined},
-    {"TTat4x833", BuiltinType::kUndefined},
-    {"ddUUnntYYx2h", BuiltinType::kUndefined},
+    {"rat3x499", BuiltinType::kUndefined},
+    {"mGtt31W4h", BuiltinType::kUndefined},
+    {"yatx4", BuiltinType::kUndefined},
+    {"mt4x2", BuiltinType::kUndefined},
+    {"maBBI4x2", BuiltinType::kUndefined},
+    {"mat8TTx2", BuiltinType::kUndefined},
+    {"ddUUnntYYx2f", BuiltinType::kUndefined},
     {"m5CCxxdZ", BuiltinType::kUndefined},
-    {"matkkq2h", BuiltinType::kUndefined},
-    {"5iitp00", BuiltinType::kUndefined},
-    {"mnntIIx3", BuiltinType::kUndefined},
-    {"ccaKx", BuiltinType::kUndefined},
-    {"m43KK", BuiltinType::kUndefined},
-    {"mat66x3f", BuiltinType::kUndefined},
-    {"Et4PP3K", BuiltinType::kUndefined},
-    {"xxatx3h", BuiltinType::kUndefined},
-    {"qat4x3h", BuiltinType::kUndefined},
-    {"MMayySrxh", BuiltinType::kUndefined},
-    {"uat4", BuiltinType::kUndefined},
-    {"tx4", BuiltinType::kUndefined},
-    {"ma54FF4", BuiltinType::kUndefined},
-    {"rra444z4f", BuiltinType::kUndefined},
-    {"matWW", BuiltinType::kUndefined},
-    {"CatZJXx4f", BuiltinType::kUndefined},
-    {"maPPx4h", BuiltinType::kUndefined},
-    {"mat4c4h", BuiltinType::kUndefined},
-    {"matPPll6h", BuiltinType::kUndefined},
-    {"9tyy", BuiltinType::kUndefined},
-    {"ptKK", BuiltinType::kUndefined},
-    {"x_", BuiltinType::kUndefined},
-    {"ayKer", BuiltinType::kUndefined},
-    {"szmpVek", BuiltinType::kUndefined},
-    {"sampqeK", BuiltinType::kUndefined},
-    {"sampler_comparisn", BuiltinType::kUndefined},
-    {"sapler_comparisVVn", BuiltinType::kUndefined},
-    {"samplerIcompaAUison", BuiltinType::kUndefined},
-    {"jexurbRd", BuiltinType::kUndefined},
-    {"exure_YYd", BuiltinType::kUndefined},
-    {"exture_1d", BuiltinType::kUndefined},
+    {"matkkq2f", BuiltinType::kUndefined},
+    {"005itpxh", BuiltinType::kUndefined},
+    {"maIInnx2h", BuiltinType::kUndefined},
+    {"Ka4Wcc", BuiltinType::kUndefined},
+    {"mtKK", BuiltinType::kUndefined},
+    {"ma664x3", BuiltinType::kUndefined},
+    {"mKKtPx", BuiltinType::kUndefined},
+    {"xxatx3f", BuiltinType::kUndefined},
+    {"qat4x3f", BuiltinType::kUndefined},
+    {"MMayySrxf", BuiltinType::kUndefined},
+    {"mat3h", BuiltinType::kUndefined},
+    {"tx3h", BuiltinType::kUndefined},
+    {"ma5F4x3h", BuiltinType::kUndefined},
+    {"rraz44x4", BuiltinType::kUndefined},
+    {"aWWx", BuiltinType::kUndefined},
+    {"ZZJJtCxX", BuiltinType::kUndefined},
+    {"maPPx4f", BuiltinType::kUndefined},
+    {"mat4c4f", BuiltinType::kUndefined},
+    {"matPPll6f", BuiltinType::kUndefined},
+    {"mat994yy", BuiltinType::kUndefined},
+    {"mat4JKKh", BuiltinType::kUndefined},
+    {"mat4_h", BuiltinType::kUndefined},
+    {"K", BuiltinType::kUndefined},
+    {"kVz", BuiltinType::kUndefined},
+    {"KtS", BuiltinType::kUndefined},
+    {"ampler", BuiltinType::kUndefined},
+    {"aVVpler", BuiltinType::kUndefined},
+    {"AAamIlUr", BuiltinType::kUndefined},
+    {"jamper_compaRson", BuiltinType::kUndefined},
+    {"amplr44compYYriso", BuiltinType::kUndefined},
+    {"samper_comparison", BuiltinType::kUndefined},
     {"texxxur_1d", BuiltinType::kUndefined},
-    {"tJxucce_2d", BuiltinType::kUndefined},
+    {"tJxucce_1d", BuiltinType::kUndefined},
     {"texure_JJd", BuiltinType::kUndefined},
-    {"lDexture_fCC_arraU", BuiltinType::kUndefined},
-    {"tegture_2d_array", BuiltinType::kUndefined},
-    {"teCCure2d_arra", BuiltinType::kUndefined},
-    {"textue_3d", BuiltinType::kUndefined},
-    {"tIx__ure_3d", BuiltinType::kUndefined},
-    {"texurettPP", BuiltinType::kUndefined},
-    {"tddx3ure_cube", BuiltinType::kUndefined},
-    {"teKyyur_cube", BuiltinType::kUndefined},
-    {"teturecub", BuiltinType::kUndefined},
-    {"textinne_c03e_array", BuiltinType::kUndefined},
-    {"nextCCruuvcubK_array", BuiltinType::kUndefined},
-    {"tXxturellcbe_array", BuiltinType::kUndefined},
-    {"tppxture_depth_2d", BuiltinType::kUndefined},
-    {"txture_deptww_2d", BuiltinType::kUndefined},
-    {"gexturedemmthuu2", BuiltinType::kUndefined},
-    {"texmmre_depthaa2daray", BuiltinType::kUndefined},
-    {"texture_RRepth_Td_ccZray", BuiltinType::kUndefined},
-    {"text88re_depthTOd_array", BuiltinType::kUndefined},
-    {"texture_depth_cm00e", BuiltinType::kUndefined},
-    {"texture_Bmepth_cube", BuiltinType::kUndefined},
-    {"Mextre_ppeph_cube", BuiltinType::kUndefined},
-    {"texturOO_depth_cub_array", BuiltinType::kUndefined},
-    {"GeGGture_depthcube_array", BuiltinType::kUndefined},
-    {"texture11Hdepth_cube_array", BuiltinType::kUndefined},
-    {"textu6e_FFepth_multiameeled_2d", BuiltinType::kUndefined},
-    {"texture_epth_mltisampled_2d", BuiltinType::kUndefined},
-    {"texture_depth_mullKsaiipled_2d", BuiltinType::kUndefined},
-    {"texture_extenal", BuiltinType::kUndefined},
-    {"IIext99reexvvernal", BuiltinType::kUndefined},
-    {"texture_externl", BuiltinType::kUndefined},
-    {"texture_mhltisampled_2d", BuiltinType::kUndefined},
-    {"texturemPllltisampzzed_2d", BuiltinType::kUndefined},
-    {"exture_mltisamed_2d", BuiltinType::kUndefined},
-    {"texture_qqtoragff_1", BuiltinType::kUndefined},
-    {"textre_JJddorage_1W", BuiltinType::kUndefined},
-    {"XXrxture_storae1zz", BuiltinType::kUndefined},
-    {"texturestorag2_2d", BuiltinType::kUndefined},
-    {"yyNxture_storage_2d", BuiltinType::kUndefined},
-    {"etue_storage_2OO", BuiltinType::kUndefined},
-    {"reutuPe_storZgeE2d_array", BuiltinType::kUndefined},
-    {"texlure_storddeee_d_22rray", BuiltinType::kUndefined},
-    {"texture_mtorage_2V_a9ra", BuiltinType::kUndefined},
-    {"teII1re_storage_3d", BuiltinType::kUndefined},
-    {"texture_storagb_3d", BuiltinType::kUndefined},
-    {"texizrestorge73d", BuiltinType::kUndefined},
-    {"u3oi", BuiltinType::kUndefined},
-    {"3", BuiltinType::kUndefined},
-    {"S2", BuiltinType::kUndefined},
-    {"e22", BuiltinType::kUndefined},
-    {"1eC2", BuiltinType::kUndefined},
-    {"vf8c2", BuiltinType::kUndefined},
-    {"c2f", BuiltinType::kUndefined},
-    {"JJecSSf", BuiltinType::kUndefined},
-    {"92f", BuiltinType::kUndefined},
+    {"lextuUe_2fDC", BuiltinType::kUndefined},
+    {"textgre_2d", BuiltinType::kUndefined},
+    {"CCxtue_2", BuiltinType::kUndefined},
+    {"txture_2d_array", BuiltinType::kUndefined},
+    {"textu__e_2d_arraI", BuiltinType::kUndefined},
+    {"texurPtt_2dNrray", BuiltinType::kUndefined},
+    {"tddx3ure_3d", BuiltinType::kUndefined},
+    {"teyytur_Kd", BuiltinType::kUndefined},
+    {"tetuuud", BuiltinType::kUndefined},
+    {"tex0unne_cu3e", BuiltinType::kUndefined},
+    {"uuexnure_cuvCK", BuiltinType::kUndefined},
+    {"tXllure_cube", BuiltinType::kUndefined},
+    {"texture_cppbo_array", BuiltinType::kUndefined},
+    {"teture_wwube_array", BuiltinType::kUndefined},
+    {"texturgm_cuube_aay", BuiltinType::kUndefined},
+    {"texmmre_deaath_", BuiltinType::kUndefined},
+    {"textRRcceTdeptZ_2d", BuiltinType::kUndefined},
+    {"te88tureOdTpth_2d", BuiltinType::kUndefined},
+    {"textume_00epth_2d_array", BuiltinType::kUndefined},
+    {"texturm_depth_BBd_array", BuiltinType::kUndefined},
+    {"texppur_depth_M_array", BuiltinType::kUndefined},
+    {"textre_deptOO_cube", BuiltinType::kUndefined},
+    {"texture_deptG_GGbe", BuiltinType::kUndefined},
+    {"tHHx11ure_depth_cube", BuiltinType::kUndefined},
+    {"texeere_depth_c6bFF_array", BuiltinType::kUndefined},
+    {"texure_depth_cub_array", BuiltinType::kUndefined},
+    {"tKxtiire_depth_cule_array", BuiltinType::kUndefined},
+    {"texture_depth_mulisampled_2d", BuiltinType::kUndefined},
+    {"texture99vdpth_multisIImpled_2d", BuiltinType::kUndefined},
+    {"texture_depth_multisampled_2", BuiltinType::kUndefined},
+    {"texture_hxternal", BuiltinType::kUndefined},
+    {"textullPP_extzzrna", BuiltinType::kUndefined},
+    {"texure_teral", BuiltinType::kUndefined},
+    {"texture_mufftiqqamled_2d", BuiltinType::kUndefined},
+    {"tddtuJJe_multisampled_2d", BuiltinType::kUndefined},
+    {"texuXXe_multisampldzz2d", BuiltinType::kUndefined},
+    {"texturestorag2_1d", BuiltinType::kUndefined},
+    {"yyNxture_storage_1d", BuiltinType::kUndefined},
+    {"etue_storage_1OO", BuiltinType::kUndefined},
+    {"reuEure_storZge_2P", BuiltinType::kUndefined},
+    {"teexture_sl2ddrage_d", BuiltinType::kUndefined},
+    {"tVVxture_9tagm_2d", BuiltinType::kUndefined},
+    {"tIIture_storage_2d_arr1y", BuiltinType::kUndefined},
+    {"texture_storagb_2d_array", BuiltinType::kUndefined},
+    {"texture_stzrage_2dai7a", BuiltinType::kUndefined},
+    {"texoure_storagii_3d", BuiltinType::kUndefined},
+    {"exture_ntor5ge3d", BuiltinType::kUndefined},
+    {"exture_storiSe_3d", BuiltinType::kUndefined},
+    {"22", BuiltinType::kUndefined},
+    {"uGC", BuiltinType::kUndefined},
+    {"f832", BuiltinType::kUndefined},
+    {"e2", BuiltinType::kUndefined},
+    {"SJJec2", BuiltinType::kUndefined},
+    {"9c", BuiltinType::kUndefined},
     {"vbbJJ2TT", BuiltinType::kUndefined},
-    {"e66h", BuiltinType::kUndefined},
-    {"u662h", BuiltinType::kUndefined},
-    {"vW2i", BuiltinType::kUndefined},
-    {"v2i", BuiltinType::kUndefined},
-    {"veci", BuiltinType::kUndefined},
-    {"rec2u", BuiltinType::kUndefined},
+    {"e66f", BuiltinType::kUndefined},
+    {"u662f", BuiltinType::kUndefined},
+    {"vW2h", BuiltinType::kUndefined},
+    {"v2h", BuiltinType::kUndefined},
+    {"vech", BuiltinType::kUndefined},
+    {"rec2i", BuiltinType::kUndefined},
     {"2ec2B", BuiltinType::kUndefined},
-    {"vcBBu", BuiltinType::kUndefined},
-    {"veRR", BuiltinType::kUndefined},
-    {"VLL0", BuiltinType::kUndefined},
-    {"KOe3", BuiltinType::kUndefined},
-    {"vgwcf", BuiltinType::kUndefined},
-    {"vLphf", BuiltinType::kUndefined},
-    {"eiiEf", BuiltinType::kUndefined},
-    {"ec3h", BuiltinType::kUndefined},
+    {"vcBBi", BuiltinType::kUndefined},
+    {"vRc2u", BuiltinType::kUndefined},
+    {"v2LL0", BuiltinType::kUndefined},
+    {"vKOOu", BuiltinType::kUndefined},
+    {"gwc3", BuiltinType::kUndefined},
+    {"hpLc", BuiltinType::kUndefined},
+    {"Eiic", BuiltinType::kUndefined},
+    {"ec3f", BuiltinType::kUndefined},
     {"UU883", BuiltinType::kUndefined},
-    {"rrecvvh", BuiltinType::kUndefined},
+    {"rrecvvf", BuiltinType::kUndefined},
     {"ecmm", BuiltinType::kUndefined},
     {"vec4j", BuiltinType::kUndefined},
     {"vec3X", BuiltinType::kUndefined},
     {"vec38", BuiltinType::kUndefined},
     {"vecvEE", BuiltinType::kUndefined},
     {"z99ci", BuiltinType::kUndefined},
-    {"JJGeQQ4", BuiltinType::kUndefined},
-    {"ssec4", BuiltinType::kUndefined},
-    {"PecK", BuiltinType::kUndefined},
-    {"tpc4f", BuiltinType::kUndefined},
-    {"vec", BuiltinType::kUndefined},
-    {"MMec4f", BuiltinType::kUndefined},
+    {"GGeJJA3u", BuiltinType::kUndefined},
+    {"vess3u", BuiltinType::kUndefined},
+    {"vPcKu", BuiltinType::kUndefined},
+    {"vtc4", BuiltinType::kUndefined},
+    {"e4", BuiltinType::kUndefined},
+    {"veMM4", BuiltinType::kUndefined},
     {"vJJc40", BuiltinType::kUndefined},
     {"8c", BuiltinType::kUndefined},
     {"vecggKh", BuiltinType::kUndefined},
-    {"vecfi", BuiltinType::kUndefined},
+    {"vecfh", BuiltinType::kUndefined},
     {"vec47Q", BuiltinType::kUndefined},
-    {"veY4i", BuiltinType::kUndefined},
-    {"keSu", BuiltinType::kUndefined},
+    {"veY4h", BuiltinType::kUndefined},
+    {"keSi", BuiltinType::kUndefined},
     {"n422", BuiltinType::kUndefined},
-    {"vFFu", BuiltinType::kUndefined},
+    {"vFFi", BuiltinType::kUndefined},
+    {"GGIec4PP", BuiltinType::kUndefined},
+    {"aeEE4", BuiltinType::kUndefined},
+    {"ddBee4u", BuiltinType::kUndefined},
 };
 
 using BuiltinTypeParseTest = testing::TestWithParam<Case>;
diff --git a/src/tint/lang/core/constant/BUILD.bazel b/src/tint/lang/core/constant/BUILD.bazel
index b092c56..c58c65c 100644
--- a/src/tint/lang/core/constant/BUILD.bazel
+++ b/src/tint/lang/core/constant/BUILD.bazel
@@ -41,6 +41,7 @@
   srcs = [
     "composite.cc",
     "eval.cc",
+    "invalid.cc",
     "manager.cc",
     "node.cc",
     "scalar.cc",
@@ -51,6 +52,7 @@
     "clone_context.h",
     "composite.h",
     "eval.h",
+    "invalid.h",
     "manager.h",
     "node.h",
     "scalar.h",
@@ -93,6 +95,7 @@
     "eval_test.h",
     "eval_unary_op_test.cc",
     "helper_test.h",
+    "invalid_test.cc",
     "manager_test.cc",
     "scalar_test.cc",
     "splat_test.cc",
diff --git a/src/tint/lang/core/constant/BUILD.cmake b/src/tint/lang/core/constant/BUILD.cmake
index b7dbca1..c4f4558 100644
--- a/src/tint/lang/core/constant/BUILD.cmake
+++ b/src/tint/lang/core/constant/BUILD.cmake
@@ -44,6 +44,8 @@
   lang/core/constant/composite.h
   lang/core/constant/eval.cc
   lang/core/constant/eval.h
+  lang/core/constant/invalid.cc
+  lang/core/constant/invalid.h
   lang/core/constant/manager.cc
   lang/core/constant/manager.h
   lang/core/constant/node.cc
@@ -91,6 +93,7 @@
   lang/core/constant/eval_test.h
   lang/core/constant/eval_unary_op_test.cc
   lang/core/constant/helper_test.h
+  lang/core/constant/invalid_test.cc
   lang/core/constant/manager_test.cc
   lang/core/constant/scalar_test.cc
   lang/core/constant/splat_test.cc
diff --git a/src/tint/lang/core/constant/BUILD.gn b/src/tint/lang/core/constant/BUILD.gn
index 8549c0b..2a343e5 100644
--- a/src/tint/lang/core/constant/BUILD.gn
+++ b/src/tint/lang/core/constant/BUILD.gn
@@ -49,6 +49,8 @@
     "composite.h",
     "eval.cc",
     "eval.h",
+    "invalid.cc",
+    "invalid.h",
     "manager.cc",
     "manager.h",
     "node.cc",
@@ -93,6 +95,7 @@
       "eval_test.h",
       "eval_unary_op_test.cc",
       "helper_test.h",
+      "invalid_test.cc",
       "manager_test.cc",
       "scalar_test.cc",
       "splat_test.cc",
diff --git a/src/tint/lang/core/constant/invalid.cc b/src/tint/lang/core/constant/invalid.cc
new file mode 100644
index 0000000..ecf434a
--- /dev/null
+++ b/src/tint/lang/core/constant/invalid.cc
@@ -0,0 +1,45 @@
+// 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/core/constant/invalid.h"
+
+#include "src/tint/lang/core/constant/manager.h"
+#include "src/tint/lang/core/type/invalid.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::core::constant::Invalid);
+
+namespace tint::core::constant {
+
+Invalid::Invalid(const core::type::Invalid* ty) : type(ty) {}
+
+Invalid::~Invalid() = default;
+
+const Invalid* Invalid::Clone(CloneContext& ctx) const {
+    return ctx.dst.Invalid();
+}
+
+}  // namespace tint::core::constant
diff --git a/src/tint/lang/core/constant/invalid.h b/src/tint/lang/core/constant/invalid.h
new file mode 100644
index 0000000..e767172
--- /dev/null
+++ b/src/tint/lang/core/constant/invalid.h
@@ -0,0 +1,79 @@
+// 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.
+
+#ifndef SRC_TINT_LANG_CORE_CONSTANT_INVALID_H_
+#define SRC_TINT_LANG_CORE_CONSTANT_INVALID_H_
+
+#include "src/tint/lang/core/constant/value.h"
+#include "src/tint/lang/core/type/invalid.h"
+#include "src/tint/utils/rtti/castable.h"
+
+namespace tint::core::constant {
+
+/// Invalid represents an invalid constant, used as a placeholder in a failed parse / resolve.
+class Invalid : public Castable<Invalid, Value> {
+  public:
+    /// Constructor
+    /// @param ty the Invalid type
+    explicit Invalid(const core::type::Invalid* ty);
+    ~Invalid() override;
+
+    /// @returns the type of the Invalid
+    const core::type::Type* Type() const override { return type; }
+
+    /// Retrieve item at index @p i
+    /// @param i the index to retrieve
+    /// @returns the element, or nullptr if out of bounds
+    const Value* Index([[maybe_unused]] size_t i) const override { return nullptr; }
+
+    /// @copydoc Value::NumElements()
+    size_t NumElements() const override { return 0; }
+
+    /// @returns true if the element is zero
+    bool AllZero() const override { return false; }
+    /// @returns true if the element is zero
+    bool AnyZero() const override { return false; }
+
+    /// @returns the hash for the Invalid
+    HashCode Hash() const override { return tint::Hash(type); }
+
+    /// Clones the constant into the provided context
+    /// @param ctx the clone context
+    /// @returns the cloned node
+    const Invalid* Clone(CloneContext& ctx) const override;
+
+    /// The Invalid type
+    core::type::Invalid const* const type;
+
+  protected:
+    /// @returns a monostate variant.
+    std::variant<std::monostate, AInt, AFloat> InternalValue() const override { return {}; }
+};
+
+}  // namespace tint::core::constant
+
+#endif  // SRC_TINT_LANG_CORE_CONSTANT_INVALID_H_
diff --git a/src/tint/lang/core/constant/invalid_test.cc b/src/tint/lang/core/constant/invalid_test.cc
new file mode 100644
index 0000000..bcbcb7a
--- /dev/null
+++ b/src/tint/lang/core/constant/invalid_test.cc
@@ -0,0 +1,73 @@
+// 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/core/constant/invalid.h"
+
+#include "src/tint/lang/core/constant/helper_test.h"
+#include "src/tint/lang/core/constant/scalar.h"
+#include "src/tint/lang/core/fluent_types.h"
+
+using namespace tint::core::number_suffixes;  // NOLINT
+using namespace tint::core::fluent_types;     // NOLINT
+
+namespace tint::core::constant {
+namespace {
+
+using ConstantTest_Invalid = TestHelper;
+
+TEST_F(ConstantTest_Invalid, AllZero) {
+    auto* invalid = constants.Invalid();
+    EXPECT_FALSE(invalid->AllZero());
+}
+
+TEST_F(ConstantTest_Invalid, AnyZero) {
+    auto* invalid = constants.Invalid();
+    EXPECT_FALSE(invalid->AnyZero());
+}
+
+TEST_F(ConstantTest_Invalid, Index) {
+    auto* invalid = constants.Invalid();
+    EXPECT_EQ(invalid->Index(0), nullptr);
+    EXPECT_EQ(invalid->Index(1), nullptr);
+    EXPECT_EQ(invalid->Index(2), nullptr);
+}
+
+TEST_F(ConstantTest_Invalid, Clone) {
+    auto* invalid = constants.Invalid();
+
+    constant::Manager mgr;
+    constant::CloneContext ctx{core::type::CloneContext{{nullptr}, {nullptr, &mgr.types}}, mgr};
+
+    auto* cloned = invalid->Clone(ctx);
+    EXPECT_NE(cloned, invalid);
+    ASSERT_NE(cloned, nullptr);
+    EXPECT_TRUE(cloned->type->Is<core::type::Invalid>());
+    EXPECT_TRUE(cloned->Is<core::constant::Invalid>());
+}
+
+}  // namespace
+}  // namespace tint::core::constant
diff --git a/src/tint/lang/core/constant/manager.cc b/src/tint/lang/core/constant/manager.cc
index ecd4a38..bf8a822 100644
--- a/src/tint/lang/core/constant/manager.cc
+++ b/src/tint/lang/core/constant/manager.cc
@@ -28,6 +28,7 @@
 #include "src/tint/lang/core/constant/manager.h"
 
 #include "src/tint/lang/core/constant/composite.h"
+#include "src/tint/lang/core/constant/invalid.h"
 #include "src/tint/lang/core/constant/scalar.h"
 #include "src/tint/lang/core/constant/splat.h"
 #include "src/tint/lang/core/type/abstract_float.h"
@@ -162,7 +163,12 @@
         [&](const core::type::F32*) { return Get(f32(0)); },               //
         [&](const core::type::F16*) { return Get(f16(0)); },               //
         [&](const core::type::Bool*) { return Get(false); },               //
+        [&](const core::type::Invalid*) { return Invalid(); },             //
         TINT_ICE_ON_NO_MATCH);
 }
 
+const constant::Invalid* Manager::Invalid() {
+    return values_.Get<constant::Invalid>(types.invalid());
+}
+
 }  // namespace tint::core::constant
diff --git a/src/tint/lang/core/constant/manager.h b/src/tint/lang/core/constant/manager.h
index 43497db..7105996 100644
--- a/src/tint/lang/core/constant/manager.h
+++ b/src/tint/lang/core/constant/manager.h
@@ -30,6 +30,7 @@
 
 #include <utility>
 
+#include "src/tint/lang/core/constant/invalid.h"
 #include "src/tint/lang/core/constant/value.h"
 #include "src/tint/lang/core/number.h"
 #include "src/tint/lang/core/type/manager.h"
@@ -144,6 +145,10 @@
     /// @returns a constant zero-value for the type
     const Value* Zero(const core::type::Type* type);
 
+    /// Constructs an invalid constant
+    /// @returns an invalid constant
+    const constant::Invalid* Invalid();
+
     /// The type manager
     core::type::Manager types;
 
diff --git a/src/tint/lang/core/constant/splat_test.cc b/src/tint/lang/core/constant/splat_test.cc
index 553cf14..10c67f2 100644
--- a/src/tint/lang/core/constant/splat_test.cc
+++ b/src/tint/lang/core/constant/splat_test.cc
@@ -80,6 +80,7 @@
     ASSERT_NE(sp->Index(0), nullptr);
     ASSERT_NE(sp->Index(1), nullptr);
     ASSERT_NE(sp->Index(2), nullptr);
+    EXPECT_EQ(sp->Index(3), nullptr);
 
     EXPECT_EQ(sp->Index(0)->As<Scalar<f32>>()->ValueOf(), 1.f);
     EXPECT_EQ(sp->Index(1)->As<Scalar<f32>>()->ValueOf(), 1.f);
@@ -95,6 +96,7 @@
     constant::CloneContext ctx{core::type::CloneContext{{nullptr}, {nullptr, &mgr.types}}, mgr};
 
     auto* r = sp->Clone(ctx);
+    EXPECT_NE(r, sp);
     ASSERT_NE(r, nullptr);
     EXPECT_TRUE(r->type->Is<core::type::Vector>());
     EXPECT_TRUE(r->el->Is<Scalar<i32>>());
diff --git a/src/tint/lang/core/constant/value.cc b/src/tint/lang/core/constant/value.cc
index f426010..97a64f1 100644
--- a/src/tint/lang/core/constant/value.cc
+++ b/src/tint/lang/core/constant/value.cc
@@ -29,6 +29,7 @@
 
 #include "src/tint/lang/core/constant/splat.h"
 #include "src/tint/lang/core/type/array.h"
+#include "src/tint/lang/core/type/invalid.h"
 #include "src/tint/lang/core/type/matrix.h"
 #include "src/tint/lang/core/type/struct.h"
 #include "src/tint/lang/core/type/vector.h"
@@ -106,13 +107,8 @@
             }
             return false;
         },
-        [&](Default) {
-            auto va = InternalValue();
-            auto vb = b->InternalValue();
-            TINT_ASSERT(!std::holds_alternative<std::monostate>(va));
-            TINT_ASSERT(!std::holds_alternative<std::monostate>(vb));
-            return va == vb;
-        });
+        [&](const core::type::Invalid*) { return true; },
+        [&](Default) { return InternalValue() == b->InternalValue(); });
 }
 
 }  // namespace tint::core::constant
diff --git a/src/tint/lang/core/core.def b/src/tint/lang/core/core.def
index 0f2d0d8..175f61c 100644
--- a/src/tint/lang/core/core.def
+++ b/src/tint/lang/core/core.def
@@ -161,6 +161,8 @@
   texture_storage_3d
   // https://www.w3.org/TR/WGSL/#external-texture-type
   texture_external
+  // chromium_internal_input_attachments
+  input_attachment
 
   // Internal types.
   __packed_vec3  // note: not core type, but used by legacy MSL writer
@@ -217,6 +219,9 @@
 
   // framebuffer-fetch input
   color
+
+  // chromium_internal_input_attachments
+  input_attachment_index
 }
 
 ////////////////////////////////////////////////////////////////////////////////
diff --git a/src/tint/lang/core/ir/binary/decode.cc b/src/tint/lang/core/ir/binary/decode.cc
index 780afd2..1960c60 100644
--- a/src/tint/lang/core/ir/binary/decode.cc
+++ b/src/tint/lang/core/ir/binary/decode.cc
@@ -50,8 +50,9 @@
 namespace {
 
 struct Decoder {
-    pb::Module& mod_in_;
-    Module& mod_out_;
+    const pb::Module& mod_in_;
+
+    Module mod_out_{};
     Vector<ir::Block*, 32> blocks_{};
     Vector<const type::Type*, 32> types_{};
     Vector<const core::constant::Value*, 32> constant_values_{};
@@ -65,7 +66,7 @@
     Vector<ir::BreakIf*, 32> break_ifs_{};
     Vector<ir::Continue*, 32> continues_{};
 
-    void Decode() {
+    Result<Module> Decode() {
         {
             const size_t n = static_cast<size_t>(mod_in_.types().size());
             types_.Reserve(n);
@@ -132,6 +133,8 @@
         for (auto* cont : continues_) {
             InferControlInstruction(cont, &Continue::SetLoop);
         }
+
+        return std::move(mod_out_);
     }
 
     template <typename EXIT, typename CTRL_INST>
@@ -1395,10 +1398,11 @@
         return Failure{"failed to deserialize protobuf"};
     }
 
-    Module mod_out;
-    Decoder{mod_in, mod_out}.Decode();
+    return Decode(mod_in);
+}
 
-    return mod_out;
+Result<Module> Decode(const pb::Module& mod_in) {
+    return Decoder{mod_in}.Decode();
 }
 
 }  // namespace tint::core::ir::binary
diff --git a/src/tint/lang/core/ir/binary/decode.h b/src/tint/lang/core/ir/binary/decode.h
index d4810b3..3e1a3ed 100644
--- a/src/tint/lang/core/ir/binary/decode.h
+++ b/src/tint/lang/core/ir/binary/decode.h
@@ -30,15 +30,22 @@
 
 #include "src/tint/utils/result/result.h"
 
-// Forward declarartion
+// Forward declarations
 namespace tint::core::ir {
 class Module;
 }  // namespace tint::core::ir
+namespace tint::core::ir::binary::pb {
+class Module;
+}  // namespace tint::core::ir::binary::pb
 
 namespace tint::core::ir::binary {
 
+/// @returns the decoded Module from the serialized protobuf.
 Result<Module> Decode(Slice<const std::byte> encoded);
 
+/// @returns the decoded Module from the protobuf.
+Result<Module> Decode(const pb::Module& module);
+
 }  // namespace tint::core::ir::binary
 
 #endif  // SRC_TINT_LANG_CORE_IR_BINARY_DECODE_H_
diff --git a/src/tint/lang/core/ir/builder.h b/src/tint/lang/core/ir/builder.h
index e3b797b..48cad54 100644
--- a/src/tint/lang/core/ir/builder.h
+++ b/src/tint/lang/core/ir/builder.h
@@ -361,6 +361,10 @@
         return Constant(ConstantValue(v));
     }
 
+    /// Creates a new invalid ir::Constant
+    /// @returns the new constant
+    ir::Constant* InvalidConstant() { return Constant(ir.constant_values.Invalid()); }
+
     /// Retrieves the inner constant from an ir::Constant
     /// @param constant the ir constant
     /// @returns the core::constant::Value inside the constant
diff --git a/src/tint/lang/core/ir/transform/BUILD.bazel b/src/tint/lang/core/ir/transform/BUILD.bazel
index 3ea71c4..6444ea9 100644
--- a/src/tint/lang/core/ir/transform/BUILD.bazel
+++ b/src/tint/lang/core/ir/transform/BUILD.bazel
@@ -40,6 +40,7 @@
   name = "transform",
   srcs = [
     "add_empty_entry_point.cc",
+    "array_length_from_uniform.cc",
     "bgra8unorm_polyfill.cc",
     "binary_polyfill.cc",
     "binding_remapper.cc",
@@ -60,6 +61,7 @@
   ],
   hdrs = [
     "add_empty_entry_point.h",
+    "array_length_from_uniform.h",
     "bgra8unorm_polyfill.h",
     "binary_polyfill.h",
     "binding_remapper.h",
@@ -109,6 +111,7 @@
   alwayslink = True,
   srcs = [
     "add_empty_entry_point_test.cc",
+    "array_length_from_uniform_test.cc",
     "bgra8unorm_polyfill_test.cc",
     "binary_polyfill_test.cc",
     "binding_remapper_test.cc",
diff --git a/src/tint/lang/core/ir/transform/BUILD.cmake b/src/tint/lang/core/ir/transform/BUILD.cmake
index f3a8b32..d8c3431 100644
--- a/src/tint/lang/core/ir/transform/BUILD.cmake
+++ b/src/tint/lang/core/ir/transform/BUILD.cmake
@@ -43,6 +43,8 @@
 tint_add_target(tint_lang_core_ir_transform lib
   lang/core/ir/transform/add_empty_entry_point.cc
   lang/core/ir/transform/add_empty_entry_point.h
+  lang/core/ir/transform/array_length_from_uniform.cc
+  lang/core/ir/transform/array_length_from_uniform.h
   lang/core/ir/transform/bgra8unorm_polyfill.cc
   lang/core/ir/transform/bgra8unorm_polyfill.h
   lang/core/ir/transform/binary_polyfill.cc
@@ -109,6 +111,7 @@
 ################################################################################
 tint_add_target(tint_lang_core_ir_transform_test test
   lang/core/ir/transform/add_empty_entry_point_test.cc
+  lang/core/ir/transform/array_length_from_uniform_test.cc
   lang/core/ir/transform/bgra8unorm_polyfill_test.cc
   lang/core/ir/transform/binary_polyfill_test.cc
   lang/core/ir/transform/binding_remapper_test.cc
@@ -189,6 +192,7 @@
 ################################################################################
 tint_add_target(tint_lang_core_ir_transform_fuzz fuzz
   lang/core/ir/transform/add_empty_entry_point_fuzz.cc
+  lang/core/ir/transform/array_length_from_uniform_fuzz.cc
   lang/core/ir/transform/bgra8unorm_polyfill_fuzz.cc
   lang/core/ir/transform/binary_polyfill_fuzz.cc
   lang/core/ir/transform/binding_remapper_fuzz.cc
diff --git a/src/tint/lang/core/ir/transform/BUILD.gn b/src/tint/lang/core/ir/transform/BUILD.gn
index db194b7..16765ad 100644
--- a/src/tint/lang/core/ir/transform/BUILD.gn
+++ b/src/tint/lang/core/ir/transform/BUILD.gn
@@ -46,6 +46,8 @@
   sources = [
     "add_empty_entry_point.cc",
     "add_empty_entry_point.h",
+    "array_length_from_uniform.cc",
+    "array_length_from_uniform.h",
     "bgra8unorm_polyfill.cc",
     "bgra8unorm_polyfill.h",
     "binary_polyfill.cc",
@@ -109,6 +111,7 @@
   tint_unittests_source_set("unittests") {
     sources = [
       "add_empty_entry_point_test.cc",
+      "array_length_from_uniform_test.cc",
       "bgra8unorm_polyfill_test.cc",
       "binary_polyfill_test.cc",
       "binding_remapper_test.cc",
@@ -180,6 +183,7 @@
 tint_fuzz_source_set("fuzz") {
   sources = [
     "add_empty_entry_point_fuzz.cc",
+    "array_length_from_uniform_fuzz.cc",
     "bgra8unorm_polyfill_fuzz.cc",
     "binary_polyfill_fuzz.cc",
     "binding_remapper_fuzz.cc",
diff --git a/src/tint/lang/core/ir/transform/array_length_from_uniform.cc b/src/tint/lang/core/ir/transform/array_length_from_uniform.cc
new file mode 100644
index 0000000..79dfb0a
--- /dev/null
+++ b/src/tint/lang/core/ir/transform/array_length_from_uniform.cc
@@ -0,0 +1,235 @@
+// 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/core/ir/transform/array_length_from_uniform.h"
+
+#include <algorithm>
+#include <utility>
+
+#include "src/tint/lang/core/ir/builder.h"
+#include "src/tint/lang/core/ir/module.h"
+#include "src/tint/lang/core/ir/validator.h"
+
+using namespace tint::core::fluent_types;     // NOLINT
+using namespace tint::core::number_suffixes;  // NOLINT
+
+namespace tint::core::ir::transform {
+
+namespace {
+
+/// PIMPL state for the transform.
+struct State {
+    /// The IR module.
+    Module& ir;
+
+    /// The binding point to use for the uniform buffer.
+    BindingPoint ubo_binding;
+
+    /// The map from binding point to the element index which holds the size of that buffer.
+    const std::unordered_map<BindingPoint, uint32_t>& bindpoint_to_size_index;
+
+    /// The IR builder.
+    core::ir::Builder b{ir};
+
+    /// The type manager.
+    core::type::Manager& ty{ir.Types()};
+
+    /// The uniform buffer variable that holds the total size of each storage buffer.
+    Var* buffer_sizes_var = nullptr;
+
+    /// A map from an array function parameter to the function parameter that holds its length.
+    Hashmap<FunctionParam*, FunctionParam*, 8> array_param_to_length_param{};
+
+    /// Process the module.
+    void Process() {
+        // Look for and replace calls to the array length builtin.
+        for (auto* inst : ir.Instructions()) {
+            if (auto* call = inst->As<CoreBuiltinCall>()) {
+                if (call->Func() == BuiltinFn::kArrayLength) {
+                    MaybeReplace(call);
+                }
+            }
+        }
+    }
+
+    /// Replace a call to an array length builtin, if the variable appears in the bindpoint map.
+    /// @param call the arrayLength call to replace
+    void MaybeReplace(CoreBuiltinCall* call) {
+        if (auto* length = GetComputedLength(call->Args()[0], call)) {
+            call->Result(0)->ReplaceAllUsesWith(length);
+            call->Destroy();
+        }
+    }
+
+    /// Get the computed length value for a runtime-sized array pointer.
+    /// @param ptr the pointer to the runtime-sized array
+    /// @param insertion_point the insertion point for new instructions
+    /// @returns the computed length, or nullptr if the original builtin should be used
+    Value* GetComputedLength(Value* ptr, Instruction* insertion_point) {
+        // Trace back from the value until we reach the originating variable.
+        while (true) {
+            if (auto* param = ptr->As<FunctionParam>()) {
+                // The length of an array pointer passed as a function parameter will be passed as
+                // an additional parameter to the function.
+                return GetArrayLengthParam(param);
+            }
+
+            if (auto* result = ptr->As<InstructionResult>()) {
+                if (auto* var = result->Instruction()->As<Var>()) {
+                    // We found the originating variable, so compute its array length.
+                    return ComputeArrayLength(var, insertion_point);
+                }
+                if (auto* access = result->Instruction()->As<Access>()) {
+                    ptr = access->Object();
+                    continue;
+                }
+                if (auto* let = result->Instruction()->As<Let>()) {
+                    ptr = let->Value();
+                    continue;
+                }
+                TINT_UNREACHABLE() << "unhandled source of a storage buffer pointer: "
+                                   << result->Instruction()->TypeInfo().name;
+            }
+            TINT_UNREACHABLE() << "unhandled source of a storage buffer pointer: "
+                               << ptr->TypeInfo().name;
+        }
+    }
+
+    /// Get (or create) the array length parameter that corresponds to an array parameter.
+    /// @param array_param the array parameter
+    /// @returns the array length parameter
+    FunctionParam* GetArrayLengthParam(FunctionParam* array_param) {
+        return array_param_to_length_param.GetOrAdd(array_param, [&] {
+            // Add a new parameter to receive the array length.
+            auto* length = b.FunctionParam<u32>("tint_array_length");
+            array_param->Function()->AppendParam(length);
+
+            // Update callsites of this function to pass the array length to it.
+            array_param->Function()->ForEachUse([&](core::ir::Usage use) {
+                if (auto* call = use.instruction->As<core::ir::UserCall>()) {
+                    // Get the length of the array in the calling function and pass that.
+                    auto* arg = call->Args()[array_param->Index()];
+                    call->AppendArg(GetComputedLength(arg, call));
+                }
+            });
+
+            return length;
+        });
+    }
+
+    /// Compute the array length of the runtime-sized array that is inside a storage buffer
+    /// variable. If the variable's binding point is not found in the bindpoint map, returns nullptr
+    /// to indicate that the original arrayLength builtin should be used instead.
+    ///
+    /// @param var the storage buffer variable that contains the runtime-sized array
+    /// @param insertion_point the insertion point for new instructions
+    /// @returns the length of the array, or nullptr if the original builtin should be used
+    Value* ComputeArrayLength(Var* var, Instruction* insertion_point) {
+        auto binding = var->BindingPoint();
+        TINT_ASSERT(binding);
+
+        auto idx_it = bindpoint_to_size_index.find(*binding);
+        if (idx_it == bindpoint_to_size_index.end()) {
+            // If the bindpoint_to_size_index map does not contain an entry for the storage buffer,
+            // then we preserve the arrayLength() call.
+            return nullptr;
+        }
+
+        Value* result = nullptr;
+        b.InsertBefore(insertion_point, [&] {
+            // Load the total storage buffer size from the uniform buffer.
+            // The sizes are packed into vec4s to satisfy the 16-byte alignment requirement for
+            // array elements in uniform buffers, so we have to find the vector and element that
+            // correspond to the index that we want.
+            const uint32_t size_index = idx_it->second;
+            const uint32_t array_index = size_index / 4;
+            const uint32_t vec_index = size_index % 4;
+            auto* vec_ptr = b.Access<ptr<uniform, vec4<u32>>>(BufferSizes(), u32(array_index));
+            auto* total_buffer_size = b.LoadVectorElement(vec_ptr, u32(vec_index))->Result(0);
+
+            // Calculate actual array length:
+            //                total_buffer_size - array_offset
+            // array_length = --------------------------------
+            //                             array_stride
+            auto* array_size = total_buffer_size;
+            auto* storage_buffer_type = var->Result(0)->Type()->UnwrapPtr();
+            const type::Array* array_type = nullptr;
+            if (auto* str = storage_buffer_type->As<core::type::Struct>()) {
+                // The variable is a struct, so subtract the byte offset of the array member.
+                auto* member = str->Members().Back();
+                array_type = member->Type()->As<core::type::Array>();
+                array_size = b.Subtract<u32>(total_buffer_size, u32(member->Offset()))->Result(0);
+            } else {
+                array_type = storage_buffer_type->As<core::type::Array>();
+            }
+            TINT_ASSERT(array_type);
+            result = b.Divide<u32>(array_size, u32(array_type->Stride()))->Result(0);
+        });
+        return result;
+    }
+
+    /// Get (or create, on first call) the uniform buffer that contains the storage buffer sizes.
+    /// @returns the uniform buffer pointer
+    Value* BufferSizes() {
+        if (buffer_sizes_var) {
+            return buffer_sizes_var->Result(0);
+        }
+
+        // Find the largest index declared in the map, in order to determine the number of elements
+        // needed in the array of buffer sizes.
+        // The buffer sizes will be packed into vec4s to satisfy the 16-byte alignment requirement
+        // for array elements in uniform buffers.
+        uint32_t max_index = 0;
+        for (auto& entry : bindpoint_to_size_index) {
+            max_index = std::max(max_index, entry.second);
+        }
+        uint32_t num_elements = (max_index / 4) + 1;
+        b.Append(ir.root_block, [&] {
+            buffer_sizes_var = b.Var("tint_storage_buffer_sizes",
+                                     ty.ptr<uniform>(ty.array(ty.vec4<u32>(), num_elements)));
+        });
+        return buffer_sizes_var->Result(0);
+    }
+};
+
+}  // namespace
+
+Result<SuccessType> ArrayLengthFromUniform(
+    Module& ir,
+    BindingPoint ubo_binding,
+    const std::unordered_map<BindingPoint, uint32_t>& bindpoint_to_size_index) {
+    auto result = ValidateAndDumpIfNeeded(ir, "ArrayLengthFromUniform transform");
+    if (result != Success) {
+        return result;
+    }
+
+    State{ir, ubo_binding, bindpoint_to_size_index}.Process();
+
+    return Success;
+}
+
+}  // namespace tint::core::ir::transform
diff --git a/src/tint/lang/core/ir/transform/array_length_from_uniform.h b/src/tint/lang/core/ir/transform/array_length_from_uniform.h
new file mode 100644
index 0000000..1128aa8
--- /dev/null
+++ b/src/tint/lang/core/ir/transform/array_length_from_uniform.h
@@ -0,0 +1,67 @@
+// 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.
+
+#ifndef SRC_TINT_LANG_CORE_IR_TRANSFORM_ARRAY_LENGTH_FROM_UNIFORM_H_
+#define SRC_TINT_LANG_CORE_IR_TRANSFORM_ARRAY_LENGTH_FROM_UNIFORM_H_
+
+#include <unordered_map>
+
+#include "src/tint/api/common/binding_point.h"
+#include "src/tint/utils/result/result.h"
+
+// Forward declarations.
+namespace tint::core::ir {
+class Module;
+}
+
+namespace tint::core::ir::transform {
+
+/// ArrayLengthFromUniform is a transform that replaces calls to the arrayLength() builtin by
+/// calculating the array length from the total size of the storage buffer, which is received via a
+/// uniform buffer.
+///
+/// The generated uniform buffer will have the form:
+/// ```
+/// @group(0) @binding(30)
+/// var<uniform> buffer_size_ubo : array<vec4<u32>, 8>;
+/// ```
+/// The binding group and number used for this uniform buffer is provided via the transform config.
+/// The transform config also defines the mapping from a storage buffer's `BindingPoint` to the
+/// element index that will be used to get the size of that buffer.
+///
+/// @param module the module to transform
+/// @param ubo_binding the binding point to use for the uniform buffer
+/// @param bindpoint_to_size_index the map from binding point to an index which holds the size
+/// @returns success or failure
+Result<SuccessType> ArrayLengthFromUniform(
+    Module& module,
+    BindingPoint ubo_binding,
+    const std::unordered_map<BindingPoint, uint32_t>& bindpoint_to_size_index);
+
+}  // namespace tint::core::ir::transform
+
+#endif  // SRC_TINT_LANG_CORE_IR_TRANSFORM_ARRAY_LENGTH_FROM_UNIFORM_H_
diff --git a/src/tint/lang/core/ir/transform/array_length_from_uniform_fuzz.cc b/src/tint/lang/core/ir/transform/array_length_from_uniform_fuzz.cc
new file mode 100644
index 0000000..8228b05
--- /dev/null
+++ b/src/tint/lang/core/ir/transform/array_length_from_uniform_fuzz.cc
@@ -0,0 +1,54 @@
+// 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/core/ir/transform/array_length_from_uniform.h"
+
+#include "src/tint/cmd/fuzz/ir/fuzz.h"
+#include "src/tint/lang/core/ir/validator.h"
+
+namespace tint::core::ir::transform {
+namespace {
+
+void ArrayLengthFromUniformFuzzer(
+    Module& module,
+    BindingPoint ubo_binding,
+    const std::unordered_map<BindingPoint, uint32_t>& bindpoint_to_size_index) {
+    if (auto res = ArrayLengthFromUniform(module, ubo_binding, bindpoint_to_size_index);
+        res != Success) {
+        return;
+    }
+
+    Capabilities capabilities;
+    if (auto res = Validate(module, capabilities); res != Success) {
+        TINT_ICE() << "result of ArrayLengthFromUniform failed IR validation\n" << res.Failure();
+    }
+}
+
+}  // namespace
+}  // namespace tint::core::ir::transform
+
+TINT_IR_MODULE_FUZZER(tint::core::ir::transform::ArrayLengthFromUniformFuzzer);
diff --git a/src/tint/lang/core/ir/transform/array_length_from_uniform_test.cc b/src/tint/lang/core/ir/transform/array_length_from_uniform_test.cc
new file mode 100644
index 0000000..b418de9
--- /dev/null
+++ b/src/tint/lang/core/ir/transform/array_length_from_uniform_test.cc
@@ -0,0 +1,1069 @@
+// 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/core/ir/transform/array_length_from_uniform.h"
+
+#include <utility>
+
+#include "src/tint/lang/core/ir/transform/helper_test.h"
+
+namespace tint::core::ir::transform {
+namespace {
+
+using namespace tint::core::fluent_types;     // NOLINT
+using namespace tint::core::number_suffixes;  // NOLINT
+
+using IR_ArrayLengthFromUniformTest = TransformTest;
+
+TEST_F(IR_ArrayLengthFromUniformTest, NoModify_UserFunction) {
+    auto* arr = ty.array<i32>();
+    auto* arr_ptr = ty.ptr<storage>(arr);
+
+    auto* buffer = b.Var("buffer", arr_ptr);
+    buffer->SetBindingPoint(0, 0);
+    mod.root_block->Append(buffer);
+
+    auto* user_func = b.Function("arrayLength", ty.u32());
+    auto* param = b.FunctionParam("arr", arr_ptr);
+    user_func->SetParams({param});
+    b.Append(user_func->Block(), [&] {  //
+        b.Return(user_func, 42_u);
+    });
+
+    auto* func = b.Function("foo", ty.void_());
+    b.Append(func->Block(), [&] {
+        b.Call(user_func, buffer);
+        b.Return(func);
+    });
+
+    auto* src = R"(
+$B1: {  # root
+  %buffer:ptr<storage, array<i32>, read_write> = var @binding_point(0, 0)
+}
+
+%arrayLength = func(%arr:ptr<storage, array<i32>, read_write>):u32 {
+  $B2: {
+    ret 42u
+  }
+}
+%foo = func():void {
+  $B3: {
+    %5:u32 = call %arrayLength, %buffer
+    ret
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = src;
+
+    std::unordered_map<BindingPoint, uint32_t> bindpoint_to_index;
+    bindpoint_to_index[{0, 0}] = 0;
+    Run(ArrayLengthFromUniform, BindingPoint{1, 2}, bindpoint_to_index);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_ArrayLengthFromUniformTest, DirectUse) {
+    auto* arr = ty.array<i32>();
+    auto* arr_ptr = ty.ptr<storage>(arr);
+
+    auto* buffer = b.Var("buffer", arr_ptr);
+    buffer->SetBindingPoint(0, 0);
+    mod.root_block->Append(buffer);
+
+    auto* func = b.Function("foo", ty.u32());
+    b.Append(func->Block(), [&] {
+        auto* len = b.Call<u32>(BuiltinFn::kArrayLength, buffer);
+        b.Return(func, len);
+    });
+
+    auto* src = R"(
+$B1: {  # root
+  %buffer:ptr<storage, array<i32>, read_write> = var @binding_point(0, 0)
+}
+
+%foo = func():u32 {
+  $B2: {
+    %3:u32 = arrayLength %buffer
+    ret %3
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+$B1: {  # root
+  %buffer:ptr<storage, array<i32>, read_write> = var @binding_point(0, 0)
+  %tint_storage_buffer_sizes:ptr<uniform, array<vec4<u32>, 1>, read> = var
+}
+
+%foo = func():u32 {
+  $B2: {
+    %4:ptr<uniform, vec4<u32>, read> = access %tint_storage_buffer_sizes, 0u
+    %5:u32 = load_vector_element %4, 0u
+    %6:u32 = div %5, 4u
+    ret %6
+  }
+}
+)";
+
+    std::unordered_map<BindingPoint, uint32_t> bindpoint_to_index;
+    bindpoint_to_index[{0, 0}] = 0;
+    Run(ArrayLengthFromUniform, BindingPoint{1, 2}, bindpoint_to_index);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_ArrayLengthFromUniformTest, DirectUse_NonZeroIndex) {
+    auto* arr = ty.array<i32>();
+    auto* arr_ptr = ty.ptr<storage>(arr);
+
+    auto* buffer = b.Var("buffer", arr_ptr);
+    buffer->SetBindingPoint(0, 0);
+    mod.root_block->Append(buffer);
+
+    auto* func = b.Function("foo", ty.u32());
+    b.Append(func->Block(), [&] {
+        auto* len = b.Call<u32>(BuiltinFn::kArrayLength, buffer);
+        b.Return(func, len);
+    });
+
+    auto* src = R"(
+$B1: {  # root
+  %buffer:ptr<storage, array<i32>, read_write> = var @binding_point(0, 0)
+}
+
+%foo = func():u32 {
+  $B2: {
+    %3:u32 = arrayLength %buffer
+    ret %3
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+$B1: {  # root
+  %buffer:ptr<storage, array<i32>, read_write> = var @binding_point(0, 0)
+  %tint_storage_buffer_sizes:ptr<uniform, array<vec4<u32>, 2>, read> = var
+}
+
+%foo = func():u32 {
+  $B2: {
+    %4:ptr<uniform, vec4<u32>, read> = access %tint_storage_buffer_sizes, 1u
+    %5:u32 = load_vector_element %4, 3u
+    %6:u32 = div %5, 4u
+    ret %6
+  }
+}
+)";
+
+    std::unordered_map<BindingPoint, uint32_t> bindpoint_to_index;
+    bindpoint_to_index[{0, 0}] = 7;
+    Run(ArrayLengthFromUniform, BindingPoint{1, 2}, bindpoint_to_index);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_ArrayLengthFromUniformTest, DirectUse_NotInMap) {
+    auto* arr = ty.array<i32>();
+    auto* arr_ptr = ty.ptr<storage>(arr);
+
+    auto* buffer = b.Var("buffer", arr_ptr);
+    buffer->SetBindingPoint(0, 1);
+    mod.root_block->Append(buffer);
+
+    auto* func = b.Function("foo", ty.u32());
+    b.Append(func->Block(), [&] {
+        auto* len = b.Call<u32>(BuiltinFn::kArrayLength, buffer);
+        b.Return(func, len);
+    });
+
+    auto* src = R"(
+$B1: {  # root
+  %buffer:ptr<storage, array<i32>, read_write> = var @binding_point(0, 1)
+}
+
+%foo = func():u32 {
+  $B2: {
+    %3:u32 = arrayLength %buffer
+    ret %3
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+$B1: {  # root
+  %buffer:ptr<storage, array<i32>, read_write> = var @binding_point(0, 1)
+}
+
+%foo = func():u32 {
+  $B2: {
+    %3:u32 = arrayLength %buffer
+    ret %3
+  }
+}
+)";
+
+    std::unordered_map<BindingPoint, uint32_t> bindpoint_to_index;
+    bindpoint_to_index[{0, 0}] = 0;
+    Run(ArrayLengthFromUniform, BindingPoint{1, 2}, bindpoint_to_index);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_ArrayLengthFromUniformTest, ViaAccess) {
+    auto* arr = ty.array<i32>();
+    auto* arr_ptr = ty.ptr<storage>(arr);
+
+    auto* buffer = b.Var("buffer", arr_ptr);
+    buffer->SetBindingPoint(0, 0);
+    mod.root_block->Append(buffer);
+
+    auto* func = b.Function("foo", ty.u32());
+    b.Append(func->Block(), [&] {
+        auto* len = b.Call<u32>(BuiltinFn::kArrayLength, b.Access(arr_ptr, buffer));
+        b.Return(func, len);
+    });
+
+    auto* src = R"(
+$B1: {  # root
+  %buffer:ptr<storage, array<i32>, read_write> = var @binding_point(0, 0)
+}
+
+%foo = func():u32 {
+  $B2: {
+    %3:ptr<storage, array<i32>, read_write> = access %buffer
+    %4:u32 = arrayLength %3
+    ret %4
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+$B1: {  # root
+  %buffer:ptr<storage, array<i32>, read_write> = var @binding_point(0, 0)
+  %tint_storage_buffer_sizes:ptr<uniform, array<vec4<u32>, 1>, read> = var
+}
+
+%foo = func():u32 {
+  $B2: {
+    %4:ptr<storage, array<i32>, read_write> = access %buffer
+    %5:ptr<uniform, vec4<u32>, read> = access %tint_storage_buffer_sizes, 0u
+    %6:u32 = load_vector_element %5, 0u
+    %7:u32 = div %6, 4u
+    ret %7
+  }
+}
+)";
+
+    std::unordered_map<BindingPoint, uint32_t> bindpoint_to_index;
+    bindpoint_to_index[{0, 0}] = 0;
+    Run(ArrayLengthFromUniform, BindingPoint{1, 2}, bindpoint_to_index);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_ArrayLengthFromUniformTest, ViaAccess_StructMember) {
+    auto* arr = ty.array<i32>();
+    auto* structure = ty.Struct(mod.symbols.New("MyStruct"), {
+                                                                 {mod.symbols.New("a"), arr},
+                                                             });
+    auto* arr_ptr = ty.ptr<storage>(arr);
+    auto* structure_ptr = ty.ptr<storage>(structure);
+
+    auto* buffer = b.Var("buffer", structure_ptr);
+    buffer->SetBindingPoint(0, 0);
+    mod.root_block->Append(buffer);
+
+    auto* func = b.Function("foo", ty.u32());
+    b.Append(func->Block(), [&] {
+        auto* len = b.Call<u32>(BuiltinFn::kArrayLength, b.Access(arr_ptr, buffer, 0_u));
+        b.Return(func, len);
+    });
+
+    auto* src = R"(
+MyStruct = struct @align(4) {
+  a:array<i32> @offset(0)
+}
+
+$B1: {  # root
+  %buffer:ptr<storage, MyStruct, read_write> = var @binding_point(0, 0)
+}
+
+%foo = func():u32 {
+  $B2: {
+    %3:ptr<storage, array<i32>, read_write> = access %buffer, 0u
+    %4:u32 = arrayLength %3
+    ret %4
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+MyStruct = struct @align(4) {
+  a:array<i32> @offset(0)
+}
+
+$B1: {  # root
+  %buffer:ptr<storage, MyStruct, read_write> = var @binding_point(0, 0)
+  %tint_storage_buffer_sizes:ptr<uniform, array<vec4<u32>, 1>, read> = var
+}
+
+%foo = func():u32 {
+  $B2: {
+    %4:ptr<storage, array<i32>, read_write> = access %buffer, 0u
+    %5:ptr<uniform, vec4<u32>, read> = access %tint_storage_buffer_sizes, 0u
+    %6:u32 = load_vector_element %5, 0u
+    %7:u32 = sub %6, 0u
+    %8:u32 = div %7, 4u
+    ret %8
+  }
+}
+)";
+
+    std::unordered_map<BindingPoint, uint32_t> bindpoint_to_index;
+    bindpoint_to_index[{0, 0}] = 0;
+    Run(ArrayLengthFromUniform, BindingPoint{1, 2}, bindpoint_to_index);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_ArrayLengthFromUniformTest, ViaAccess_StructMember_NonZeroOffset) {
+    auto* arr = ty.array<i32>();
+    auto* structure = ty.Struct(mod.symbols.New("MyStruct"), {
+                                                                 {mod.symbols.New("u1"), ty.u32()},
+                                                                 {mod.symbols.New("u2"), ty.u32()},
+                                                                 {mod.symbols.New("u3"), ty.u32()},
+                                                                 {mod.symbols.New("u4"), ty.u32()},
+                                                                 {mod.symbols.New("u5"), ty.u32()},
+                                                                 {mod.symbols.New("a"), arr},
+                                                             });
+    auto* arr_ptr = ty.ptr<storage>(arr);
+    auto* structure_ptr = ty.ptr<storage>(structure);
+
+    auto* buffer = b.Var("buffer", structure_ptr);
+    buffer->SetBindingPoint(0, 0);
+    mod.root_block->Append(buffer);
+
+    auto* func = b.Function("foo", ty.u32());
+    b.Append(func->Block(), [&] {
+        auto* len = b.Call<u32>(BuiltinFn::kArrayLength, b.Access(arr_ptr, buffer, 5_u));
+        b.Return(func, len);
+    });
+
+    auto* src = R"(
+MyStruct = struct @align(4) {
+  u1:u32 @offset(0)
+  u2:u32 @offset(4)
+  u3:u32 @offset(8)
+  u4:u32 @offset(12)
+  u5:u32 @offset(16)
+  a:array<i32> @offset(20)
+}
+
+$B1: {  # root
+  %buffer:ptr<storage, MyStruct, read_write> = var @binding_point(0, 0)
+}
+
+%foo = func():u32 {
+  $B2: {
+    %3:ptr<storage, array<i32>, read_write> = access %buffer, 5u
+    %4:u32 = arrayLength %3
+    ret %4
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+MyStruct = struct @align(4) {
+  u1:u32 @offset(0)
+  u2:u32 @offset(4)
+  u3:u32 @offset(8)
+  u4:u32 @offset(12)
+  u5:u32 @offset(16)
+  a:array<i32> @offset(20)
+}
+
+$B1: {  # root
+  %buffer:ptr<storage, MyStruct, read_write> = var @binding_point(0, 0)
+  %tint_storage_buffer_sizes:ptr<uniform, array<vec4<u32>, 1>, read> = var
+}
+
+%foo = func():u32 {
+  $B2: {
+    %4:ptr<storage, array<i32>, read_write> = access %buffer, 5u
+    %5:ptr<uniform, vec4<u32>, read> = access %tint_storage_buffer_sizes, 0u
+    %6:u32 = load_vector_element %5, 0u
+    %7:u32 = sub %6, 20u
+    %8:u32 = div %7, 4u
+    ret %8
+  }
+}
+)";
+
+    std::unordered_map<BindingPoint, uint32_t> bindpoint_to_index;
+    bindpoint_to_index[{0, 0}] = 0;
+    Run(ArrayLengthFromUniform, BindingPoint{1, 2}, bindpoint_to_index);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_ArrayLengthFromUniformTest, ViaLet) {
+    auto* arr = ty.array<i32>();
+    auto* arr_ptr = ty.ptr<storage>(arr);
+
+    auto* buffer = b.Var("buffer", arr_ptr);
+    buffer->SetBindingPoint(0, 0);
+    mod.root_block->Append(buffer);
+
+    auto* func = b.Function("foo", ty.u32());
+    b.Append(func->Block(), [&] {
+        auto* len = b.Call<u32>(BuiltinFn::kArrayLength, b.Let("let", buffer));
+        b.Return(func, len);
+    });
+
+    auto* src = R"(
+$B1: {  # root
+  %buffer:ptr<storage, array<i32>, read_write> = var @binding_point(0, 0)
+}
+
+%foo = func():u32 {
+  $B2: {
+    %let:ptr<storage, array<i32>, read_write> = let %buffer
+    %4:u32 = arrayLength %let
+    ret %4
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+$B1: {  # root
+  %buffer:ptr<storage, array<i32>, read_write> = var @binding_point(0, 0)
+  %tint_storage_buffer_sizes:ptr<uniform, array<vec4<u32>, 1>, read> = var
+}
+
+%foo = func():u32 {
+  $B2: {
+    %let:ptr<storage, array<i32>, read_write> = let %buffer
+    %5:ptr<uniform, vec4<u32>, read> = access %tint_storage_buffer_sizes, 0u
+    %6:u32 = load_vector_element %5, 0u
+    %7:u32 = div %6, 4u
+    ret %7
+  }
+}
+)";
+
+    std::unordered_map<BindingPoint, uint32_t> bindpoint_to_index;
+    bindpoint_to_index[{0, 0}] = 0;
+    Run(ArrayLengthFromUniform, BindingPoint{1, 2}, bindpoint_to_index);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_ArrayLengthFromUniformTest, ViaParameter) {
+    auto* arr = ty.array<i32>();
+    auto* arr_ptr = ty.ptr<storage>(arr);
+
+    auto* buffer = b.Var("buffer", arr_ptr);
+    buffer->SetBindingPoint(0, 0);
+    mod.root_block->Append(buffer);
+
+    auto* bar = b.Function("bar", ty.u32());
+    auto* param = b.FunctionParam("param", arr_ptr);
+    bar->SetParams({param});
+    b.Append(bar->Block(), [&] {
+        auto* len = b.Call<u32>(BuiltinFn::kArrayLength, param);
+        b.Return(bar, len);
+    });
+
+    auto* foo = b.Function("foo", ty.u32());
+    b.Append(foo->Block(), [&] {
+        auto* len = b.Call<u32>(bar, buffer);
+        b.Return(foo, len);
+    });
+
+    auto* src = R"(
+$B1: {  # root
+  %buffer:ptr<storage, array<i32>, read_write> = var @binding_point(0, 0)
+}
+
+%bar = func(%param:ptr<storage, array<i32>, read_write>):u32 {
+  $B2: {
+    %4:u32 = arrayLength %param
+    ret %4
+  }
+}
+%foo = func():u32 {
+  $B3: {
+    %6:u32 = call %bar, %buffer
+    ret %6
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+$B1: {  # root
+  %buffer:ptr<storage, array<i32>, read_write> = var @binding_point(0, 0)
+  %tint_storage_buffer_sizes:ptr<uniform, array<vec4<u32>, 1>, read> = var
+}
+
+%bar = func(%param:ptr<storage, array<i32>, read_write>, %tint_array_length:u32):u32 {
+  $B2: {
+    ret %tint_array_length
+  }
+}
+%foo = func():u32 {
+  $B3: {
+    %7:ptr<uniform, vec4<u32>, read> = access %tint_storage_buffer_sizes, 0u
+    %8:u32 = load_vector_element %7, 0u
+    %9:u32 = div %8, 4u
+    %10:u32 = call %bar, %buffer, %9
+    ret %10
+  }
+}
+)";
+
+    std::unordered_map<BindingPoint, uint32_t> bindpoint_to_index;
+    bindpoint_to_index[{0, 0}] = 0;
+    Run(ArrayLengthFromUniform, BindingPoint{1, 2}, bindpoint_to_index);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_ArrayLengthFromUniformTest, ViaParameterChain) {
+    auto* arr = ty.array<i32>();
+    auto* arr_ptr = ty.ptr<storage>(arr);
+
+    auto* buffer = b.Var("buffer", arr_ptr);
+    buffer->SetBindingPoint(0, 0);
+    mod.root_block->Append(buffer);
+
+    auto* zoo = b.Function("foo", ty.u32());
+    auto* param_zoo = b.FunctionParam("param_zoo", arr_ptr);
+    zoo->SetParams({param_zoo});
+    b.Append(zoo->Block(), [&] {
+        auto* len = b.Call<u32>(BuiltinFn::kArrayLength, param_zoo);
+        b.Return(zoo, len);
+    });
+
+    auto* bar = b.Function("foo", ty.u32());
+    auto* param_bar = b.FunctionParam("param_bar", arr_ptr);
+    bar->SetParams({param_bar});
+    b.Append(bar->Block(), [&] {
+        auto* len = b.Call<u32>(zoo, param_bar);
+        b.Return(bar, len);
+    });
+
+    auto* foo = b.Function("foo", ty.u32());
+    b.Append(foo->Block(), [&] {
+        auto* len = b.Call<u32>(bar, buffer);
+        b.Return(foo, len);
+    });
+
+    auto* src = R"(
+$B1: {  # root
+  %buffer:ptr<storage, array<i32>, read_write> = var @binding_point(0, 0)
+}
+
+%foo = func(%param_zoo:ptr<storage, array<i32>, read_write>):u32 {
+  $B2: {
+    %4:u32 = arrayLength %param_zoo
+    ret %4
+  }
+}
+%foo_1 = func(%param_bar:ptr<storage, array<i32>, read_write>):u32 {  # %foo_1: 'foo'
+  $B3: {
+    %7:u32 = call %foo, %param_bar
+    ret %7
+  }
+}
+%foo_2 = func():u32 {  # %foo_2: 'foo'
+  $B4: {
+    %9:u32 = call %foo_1, %buffer
+    ret %9
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+$B1: {  # root
+  %buffer:ptr<storage, array<i32>, read_write> = var @binding_point(0, 0)
+  %tint_storage_buffer_sizes:ptr<uniform, array<vec4<u32>, 1>, read> = var
+}
+
+%foo = func(%param_zoo:ptr<storage, array<i32>, read_write>, %tint_array_length:u32):u32 {
+  $B2: {
+    ret %tint_array_length
+  }
+}
+%foo_1 = func(%param_bar:ptr<storage, array<i32>, read_write>, %tint_array_length_1:u32):u32 {  # %foo_1: 'foo', %tint_array_length_1: 'tint_array_length'
+  $B3: {
+    %9:u32 = call %foo, %param_bar, %tint_array_length_1
+    ret %9
+  }
+}
+%foo_2 = func():u32 {  # %foo_2: 'foo'
+  $B4: {
+    %11:ptr<uniform, vec4<u32>, read> = access %tint_storage_buffer_sizes, 0u
+    %12:u32 = load_vector_element %11, 0u
+    %13:u32 = div %12, 4u
+    %14:u32 = call %foo_1, %buffer, %13
+    ret %14
+  }
+}
+)";
+
+    std::unordered_map<BindingPoint, uint32_t> bindpoint_to_index;
+    bindpoint_to_index[{0, 0}] = 0;
+    Run(ArrayLengthFromUniform, BindingPoint{1, 2}, bindpoint_to_index);
+
+    EXPECT_EQ(expect, str());
+}
+
+// Test that we re-use the length parameter for multiple arrayLength calls on the same parameter.
+TEST_F(IR_ArrayLengthFromUniformTest, ViaParameter_MultipleCallsSameParameter) {
+    auto* arr = ty.array<i32>();
+    auto* arr_ptr = ty.ptr<storage>(arr);
+
+    auto* buffer = b.Var("buffer", arr_ptr);
+    buffer->SetBindingPoint(0, 0);
+    mod.root_block->Append(buffer);
+
+    auto* bar = b.Function("bar", ty.u32());
+    auto* param = b.FunctionParam("param", arr_ptr);
+    bar->SetParams({param});
+    b.Append(bar->Block(), [&] {
+        auto* len_a = b.Call<u32>(BuiltinFn::kArrayLength, param);
+        auto* len_b = b.Call<u32>(BuiltinFn::kArrayLength, param);
+        auto* len_c = b.Call<u32>(BuiltinFn::kArrayLength, param);
+        b.Return(bar, b.Add<u32>(len_a, b.Add<u32>(len_b, len_c)));
+    });
+
+    auto* foo = b.Function("foo", ty.u32());
+    b.Append(foo->Block(), [&] {
+        auto* len = b.Call<u32>(bar, buffer);
+        b.Return(foo, len);
+    });
+
+    auto* src = R"(
+$B1: {  # root
+  %buffer:ptr<storage, array<i32>, read_write> = var @binding_point(0, 0)
+}
+
+%bar = func(%param:ptr<storage, array<i32>, read_write>):u32 {
+  $B2: {
+    %4:u32 = arrayLength %param
+    %5:u32 = arrayLength %param
+    %6:u32 = arrayLength %param
+    %7:u32 = add %5, %6
+    %8:u32 = add %4, %7
+    ret %8
+  }
+}
+%foo = func():u32 {
+  $B3: {
+    %10:u32 = call %bar, %buffer
+    ret %10
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+$B1: {  # root
+  %buffer:ptr<storage, array<i32>, read_write> = var @binding_point(0, 0)
+  %tint_storage_buffer_sizes:ptr<uniform, array<vec4<u32>, 1>, read> = var
+}
+
+%bar = func(%param:ptr<storage, array<i32>, read_write>, %tint_array_length:u32):u32 {
+  $B2: {
+    %6:u32 = add %tint_array_length, %tint_array_length
+    %7:u32 = add %tint_array_length, %6
+    ret %7
+  }
+}
+%foo = func():u32 {
+  $B3: {
+    %9:ptr<uniform, vec4<u32>, read> = access %tint_storage_buffer_sizes, 0u
+    %10:u32 = load_vector_element %9, 0u
+    %11:u32 = div %10, 4u
+    %12:u32 = call %bar, %buffer, %11
+    ret %12
+  }
+}
+)";
+
+    std::unordered_map<BindingPoint, uint32_t> bindpoint_to_index;
+    bindpoint_to_index[{0, 0}] = 0;
+    Run(ArrayLengthFromUniform, BindingPoint{1, 2}, bindpoint_to_index);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_ArrayLengthFromUniformTest, ViaParameter_MultipleCallsDifferentParameters) {
+    auto* arr = ty.array<i32>();
+    auto* arr_ptr = ty.ptr<storage>(arr);
+
+    auto* buffer = b.Var("buffer", arr_ptr);
+    buffer->SetBindingPoint(0, 0);
+    mod.root_block->Append(buffer);
+
+    auto* bar = b.Function("bar", ty.u32());
+    auto* param_a = b.FunctionParam("param_a", arr_ptr);
+    auto* param_b = b.FunctionParam("param_b", arr_ptr);
+    auto* param_c = b.FunctionParam("param_c", arr_ptr);
+    bar->SetParams({param_a, param_b, param_c});
+    b.Append(bar->Block(), [&] {
+        auto* len_a = b.Call<u32>(BuiltinFn::kArrayLength, param_a);
+        auto* len_b = b.Call<u32>(BuiltinFn::kArrayLength, param_b);
+        auto* len_c = b.Call<u32>(BuiltinFn::kArrayLength, param_c);
+        b.Return(bar, b.Add<u32>(len_a, b.Add<u32>(len_b, len_c)));
+    });
+
+    auto* foo = b.Function("foo", ty.u32());
+    b.Append(foo->Block(), [&] {
+        auto* len = b.Call<u32>(bar, buffer, buffer, buffer);
+        b.Return(foo, len);
+    });
+
+    auto* src = R"(
+$B1: {  # root
+  %buffer:ptr<storage, array<i32>, read_write> = var @binding_point(0, 0)
+}
+
+%bar = func(%param_a:ptr<storage, array<i32>, read_write>, %param_b:ptr<storage, array<i32>, read_write>, %param_c:ptr<storage, array<i32>, read_write>):u32 {
+  $B2: {
+    %6:u32 = arrayLength %param_a
+    %7:u32 = arrayLength %param_b
+    %8:u32 = arrayLength %param_c
+    %9:u32 = add %7, %8
+    %10:u32 = add %6, %9
+    ret %10
+  }
+}
+%foo = func():u32 {
+  $B3: {
+    %12:u32 = call %bar, %buffer, %buffer, %buffer
+    ret %12
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+$B1: {  # root
+  %buffer:ptr<storage, array<i32>, read_write> = var @binding_point(0, 0)
+  %tint_storage_buffer_sizes:ptr<uniform, array<vec4<u32>, 1>, read> = var
+}
+
+%bar = func(%param_a:ptr<storage, array<i32>, read_write>, %param_b:ptr<storage, array<i32>, read_write>, %param_c:ptr<storage, array<i32>, read_write>, %tint_array_length:u32, %tint_array_length_1:u32, %tint_array_length_2:u32):u32 {  # %tint_array_length_1: 'tint_array_length', %tint_array_length_2: 'tint_array_length'
+  $B2: {
+    %10:u32 = add %tint_array_length_1, %tint_array_length_2
+    %11:u32 = add %tint_array_length, %10
+    ret %11
+  }
+}
+%foo = func():u32 {
+  $B3: {
+    %13:ptr<uniform, vec4<u32>, read> = access %tint_storage_buffer_sizes, 0u
+    %14:u32 = load_vector_element %13, 0u
+    %15:u32 = div %14, 4u
+    %16:ptr<uniform, vec4<u32>, read> = access %tint_storage_buffer_sizes, 0u
+    %17:u32 = load_vector_element %16, 0u
+    %18:u32 = div %17, 4u
+    %19:ptr<uniform, vec4<u32>, read> = access %tint_storage_buffer_sizes, 0u
+    %20:u32 = load_vector_element %19, 0u
+    %21:u32 = div %20, 4u
+    %22:u32 = call %bar, %buffer, %buffer, %buffer, %15, %18, %21
+    ret %22
+  }
+}
+)";
+
+    std::unordered_map<BindingPoint, uint32_t> bindpoint_to_index;
+    bindpoint_to_index[{0, 0}] = 0;
+    Run(ArrayLengthFromUniform, BindingPoint{1, 2}, bindpoint_to_index);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_ArrayLengthFromUniformTest, ViaComplexChain) {
+    auto* arr = ty.array<i32>();
+    auto* structure = ty.Struct(mod.symbols.New("MyStruct"), {
+                                                                 {mod.symbols.New("u1"), ty.u32()},
+                                                                 {mod.symbols.New("a"), arr},
+                                                             });
+    auto* arr_ptr = ty.ptr<storage>(arr);
+    auto* structure_ptr = ty.ptr<storage>(structure);
+
+    auto* buffer = b.Var("buffer", structure_ptr);
+    buffer->SetBindingPoint(0, 0);
+    mod.root_block->Append(buffer);
+
+    auto* bar = b.Function("bar", ty.u32());
+    auto* param = b.FunctionParam("param", arr_ptr);
+    bar->SetParams({param});
+    b.Append(bar->Block(), [&] {
+        auto* access = b.Access(arr_ptr, param);
+        auto* let = b.Let("let", access);
+        auto* len = b.Call<u32>(BuiltinFn::kArrayLength, let);
+        b.Return(bar, len);
+    });
+
+    auto* foo = b.Function("foo", ty.u32());
+    b.Append(foo->Block(), [&] {
+        auto* access = b.Access(arr_ptr, buffer, 1_u);
+        auto* let = b.Let("let", access);
+        auto* len = b.Call<u32>(bar, let);
+        b.Return(foo, len);
+    });
+
+    auto* src = R"(
+MyStruct = struct @align(4) {
+  u1:u32 @offset(0)
+  a:array<i32> @offset(4)
+}
+
+$B1: {  # root
+  %buffer:ptr<storage, MyStruct, read_write> = var @binding_point(0, 0)
+}
+
+%bar = func(%param:ptr<storage, array<i32>, read_write>):u32 {
+  $B2: {
+    %4:ptr<storage, array<i32>, read_write> = access %param
+    %let:ptr<storage, array<i32>, read_write> = let %4
+    %6:u32 = arrayLength %let
+    ret %6
+  }
+}
+%foo = func():u32 {
+  $B3: {
+    %8:ptr<storage, array<i32>, read_write> = access %buffer, 1u
+    %let_1:ptr<storage, array<i32>, read_write> = let %8  # %let_1: 'let'
+    %10:u32 = call %bar, %let_1
+    ret %10
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+MyStruct = struct @align(4) {
+  u1:u32 @offset(0)
+  a:array<i32> @offset(4)
+}
+
+$B1: {  # root
+  %buffer:ptr<storage, MyStruct, read_write> = var @binding_point(0, 0)
+  %tint_storage_buffer_sizes:ptr<uniform, array<vec4<u32>, 1>, read> = var
+}
+
+%bar = func(%param:ptr<storage, array<i32>, read_write>, %tint_array_length:u32):u32 {
+  $B2: {
+    %6:ptr<storage, array<i32>, read_write> = access %param
+    %let:ptr<storage, array<i32>, read_write> = let %6
+    ret %tint_array_length
+  }
+}
+%foo = func():u32 {
+  $B3: {
+    %9:ptr<storage, array<i32>, read_write> = access %buffer, 1u
+    %let_1:ptr<storage, array<i32>, read_write> = let %9  # %let_1: 'let'
+    %11:ptr<uniform, vec4<u32>, read> = access %tint_storage_buffer_sizes, 0u
+    %12:u32 = load_vector_element %11, 0u
+    %13:u32 = sub %12, 4u
+    %14:u32 = div %13, 4u
+    %15:u32 = call %bar, %let_1, %14
+    ret %15
+  }
+}
+)";
+
+    std::unordered_map<BindingPoint, uint32_t> bindpoint_to_index;
+    bindpoint_to_index[{0, 0}] = 0;
+    Run(ArrayLengthFromUniform, BindingPoint{1, 2}, bindpoint_to_index);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_ArrayLengthFromUniformTest, ElementStrideLargerThanSize) {
+    auto* arr = ty.array<vec3<i32>>();
+    auto* arr_ptr = ty.ptr<storage>(arr);
+
+    auto* buffer = b.Var("buffer", arr_ptr);
+    buffer->SetBindingPoint(0, 0);
+    mod.root_block->Append(buffer);
+
+    auto* func = b.Function("foo", ty.u32());
+    b.Append(func->Block(), [&] {
+        auto* len = b.Call<u32>(BuiltinFn::kArrayLength, buffer);
+        b.Return(func, len);
+    });
+
+    auto* src = R"(
+$B1: {  # root
+  %buffer:ptr<storage, array<vec3<i32>>, read_write> = var @binding_point(0, 0)
+}
+
+%foo = func():u32 {
+  $B2: {
+    %3:u32 = arrayLength %buffer
+    ret %3
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+$B1: {  # root
+  %buffer:ptr<storage, array<vec3<i32>>, read_write> = var @binding_point(0, 0)
+  %tint_storage_buffer_sizes:ptr<uniform, array<vec4<u32>, 1>, read> = var
+}
+
+%foo = func():u32 {
+  $B2: {
+    %4:ptr<uniform, vec4<u32>, read> = access %tint_storage_buffer_sizes, 0u
+    %5:u32 = load_vector_element %4, 0u
+    %6:u32 = div %5, 16u
+    ret %6
+  }
+}
+)";
+
+    std::unordered_map<BindingPoint, uint32_t> bindpoint_to_index;
+    bindpoint_to_index[{0, 0}] = 0;
+    Run(ArrayLengthFromUniform, BindingPoint{1, 2}, bindpoint_to_index);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_ArrayLengthFromUniformTest, MultipleVars) {
+    auto* arr = ty.array<i32>();
+    auto* arr_ptr = ty.ptr<storage>(arr);
+
+    auto* buffer_a = b.Var("buffer_a", arr_ptr);
+    auto* buffer_b = b.Var("buffer_b", arr_ptr);
+    auto* buffer_c = b.Var("buffer_c", arr_ptr);
+    auto* buffer_d = b.Var("buffer_d", arr_ptr);
+    auto* buffer_e = b.Var("buffer_e", arr_ptr);
+    buffer_a->SetBindingPoint(0, 0);
+    buffer_b->SetBindingPoint(0, 1);
+    buffer_c->SetBindingPoint(1, 0);
+    buffer_d->SetBindingPoint(1, 1);
+    buffer_e->SetBindingPoint(2, 3);
+    mod.root_block->Append(buffer_a);
+    mod.root_block->Append(buffer_b);
+    mod.root_block->Append(buffer_c);
+    mod.root_block->Append(buffer_d);
+    mod.root_block->Append(buffer_e);
+
+    auto* func = b.Function("foo", ty.void_());
+    b.Append(func->Block(), [&] {
+        b.Call<u32>(BuiltinFn::kArrayLength, buffer_a);
+        b.Call<u32>(BuiltinFn::kArrayLength, buffer_b);
+        b.Call<u32>(BuiltinFn::kArrayLength, buffer_c);
+        b.Call<u32>(BuiltinFn::kArrayLength, buffer_d);
+        b.Call<u32>(BuiltinFn::kArrayLength, buffer_e);
+        b.Return(func);
+    });
+
+    auto* src = R"(
+$B1: {  # root
+  %buffer_a:ptr<storage, array<i32>, read_write> = var @binding_point(0, 0)
+  %buffer_b:ptr<storage, array<i32>, read_write> = var @binding_point(0, 1)
+  %buffer_c:ptr<storage, array<i32>, read_write> = var @binding_point(1, 0)
+  %buffer_d:ptr<storage, array<i32>, read_write> = var @binding_point(1, 1)
+  %buffer_e:ptr<storage, array<i32>, read_write> = var @binding_point(2, 3)
+}
+
+%foo = func():void {
+  $B2: {
+    %7:u32 = arrayLength %buffer_a
+    %8:u32 = arrayLength %buffer_b
+    %9:u32 = arrayLength %buffer_c
+    %10:u32 = arrayLength %buffer_d
+    %11:u32 = arrayLength %buffer_e
+    ret
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+$B1: {  # root
+  %buffer_a:ptr<storage, array<i32>, read_write> = var @binding_point(0, 0)
+  %buffer_b:ptr<storage, array<i32>, read_write> = var @binding_point(0, 1)
+  %buffer_c:ptr<storage, array<i32>, read_write> = var @binding_point(1, 0)
+  %buffer_d:ptr<storage, array<i32>, read_write> = var @binding_point(1, 1)
+  %buffer_e:ptr<storage, array<i32>, read_write> = var @binding_point(2, 3)
+  %tint_storage_buffer_sizes:ptr<uniform, array<vec4<u32>, 2>, read> = var
+}
+
+%foo = func():void {
+  $B2: {
+    %8:ptr<uniform, vec4<u32>, read> = access %tint_storage_buffer_sizes, 0u
+    %9:u32 = load_vector_element %8, 0u
+    %10:u32 = div %9, 4u
+    %11:ptr<uniform, vec4<u32>, read> = access %tint_storage_buffer_sizes, 1u
+    %12:u32 = load_vector_element %11, 1u
+    %13:u32 = div %12, 4u
+    %14:ptr<uniform, vec4<u32>, read> = access %tint_storage_buffer_sizes, 0u
+    %15:u32 = load_vector_element %14, 3u
+    %16:u32 = div %15, 4u
+    %17:ptr<uniform, vec4<u32>, read> = access %tint_storage_buffer_sizes, 0u
+    %18:u32 = load_vector_element %17, 2u
+    %19:u32 = div %18, 4u
+    %20:ptr<uniform, vec4<u32>, read> = access %tint_storage_buffer_sizes, 1u
+    %21:u32 = load_vector_element %20, 0u
+    %22:u32 = div %21, 4u
+    ret
+  }
+}
+)";
+
+    std::unordered_map<BindingPoint, uint32_t> bindpoint_to_index;
+    bindpoint_to_index[{0, 0}] = 0;
+    bindpoint_to_index[{0, 1}] = 5;
+    bindpoint_to_index[{1, 0}] = 3;
+    bindpoint_to_index[{1, 1}] = 2;
+    bindpoint_to_index[{2, 3}] = 4;
+    Run(ArrayLengthFromUniform, BindingPoint{1, 2}, bindpoint_to_index);
+
+    EXPECT_EQ(expect, str());
+}
+
+}  // namespace
+}  // namespace tint::core::ir::transform
diff --git a/src/tint/lang/core/ir/transform/multiplanar_external_texture.cc b/src/tint/lang/core/ir/transform/multiplanar_external_texture.cc
index 8637697..4a0de1f 100644
--- a/src/tint/lang/core/ir/transform/multiplanar_external_texture.cc
+++ b/src/tint/lang/core/ir/transform/multiplanar_external_texture.cc
@@ -228,8 +228,14 @@
                 },
                 [&](CoreBuiltinCall* call) {
                     if (call->Func() == core::BuiltinFn::kTextureDimensions) {
-                        // Use the first plane for the `textureDimensions()` call.
-                        call->SetOperand(use.operand_index, plane_0);
+                        // Use params.visibleSize + vec2u(1, 1) instead of the textureDimensions.
+                        b.InsertBefore(call, [&] {
+                            auto* visible_size = b.Access<vec2<u32>>(params, 12_u);
+                            auto* vec2u_1_1 = b.Splat<vec2<u32>>(1_u);
+                            auto* dimensions = b.Add<vec2<u32>>(visible_size, vec2u_1_1);
+                            dimensions->SetResults(Vector{call->DetachResult()});
+                        });
+                        call->Destroy();
                     } else if (call->Func() == core::BuiltinFn::kTextureLoad) {
                         // Convert the coordinates to unsigned integers if necessary.
                         auto* coords = call->Args()[1];
@@ -410,11 +416,10 @@
             auto* yuv_to_rgb_conversion_only = b.Access(ty.u32(), params, 1_u);
             auto* yuv_to_rgb_conversion = b.Access(ty.mat3x4<f32>(), params, 2_u);
             auto* load_transform_matrix = b.Access(ty.mat3x2<f32>(), params, 7_u);
-            auto* display_visible_rect_max = b.Access(ty.vec2<u32>(), params, 12_u);
+            auto* visible_size = b.Access(ty.vec2<u32>(), params, 12_u);
             auto* plane1_coord_factor = b.Access(ty.vec2<f32>(), params, 13_u);
 
-            auto* clamped_coords =
-                b.Call(vec2u, core::BuiltinFn::kMin, coords, display_visible_rect_max);
+            auto* clamped_coords = b.Call(vec2u, core::BuiltinFn::kMin, coords, visible_size);
             auto* clamped_coords_f = b.Convert(vec2f, clamped_coords);
             auto* modified_coords =
                 b.Multiply(vec2f, load_transform_matrix, b.Construct(vec3f, clamped_coords_f, 1_f));
diff --git a/src/tint/lang/core/ir/transform/multiplanar_external_texture_test.cc b/src/tint/lang/core/ir/transform/multiplanar_external_texture_test.cc
index 39b06e8..c8dc646 100644
--- a/src/tint/lang/core/ir/transform/multiplanar_external_texture_test.cc
+++ b/src/tint/lang/core/ir/transform/multiplanar_external_texture_test.cc
@@ -270,7 +270,8 @@
     %5:texture_2d<f32> = load %texture_plane0
     %6:texture_2d<f32> = load %texture_plane1
     %7:tint_ExternalTextureParams = load %texture_params
-    %result:vec2<u32> = textureDimensions %5
+    %8:vec2<u32> = access %7, 12u
+    %result:vec2<u32> = add %8, vec2<u32>(1u)
     ret %result
   }
 }
@@ -1071,99 +1072,100 @@
     %15:texture_2d<f32> = load %texture_plane0
     %16:texture_2d<f32> = load %texture_plane1
     %17:tint_ExternalTextureParams = load %texture_params
-    %18:vec2<u32> = textureDimensions %15
-    %19:texture_2d<f32> = load %texture_plane0
-    %20:texture_2d<f32> = load %texture_plane1
-    %21:tint_ExternalTextureParams = load %texture_params
-    %22:vec4<f32> = call %tint_TextureSampleExternal, %19, %20, %21, %sampler_1, %coords_1
-    %23:texture_2d<f32> = load %texture_plane0
-    %24:texture_2d<f32> = load %texture_plane1
-    %25:tint_ExternalTextureParams = load %texture_params
-    %26:vec4<f32> = call %tint_TextureSampleExternal, %23, %24, %25, %sampler_1, %coords_1
-    %27:texture_2d<f32> = load %texture_plane0
-    %28:texture_2d<f32> = load %texture_plane1
-    %29:tint_ExternalTextureParams = load %texture_params
-    %result_a:vec4<f32> = call %foo, %27, %28, %29, %sampler_1, %coords_1
-    %result_b:vec4<f32> = call %foo, %27, %28, %29, %sampler_1, %coords_1
-    %32:vec4<f32> = add %result_a, %result_b
-    ret %32
+    %18:vec2<u32> = access %17, 12u
+    %19:vec2<u32> = add %18, vec2<u32>(1u)
+    %20:texture_2d<f32> = load %texture_plane0
+    %21:texture_2d<f32> = load %texture_plane1
+    %22:tint_ExternalTextureParams = load %texture_params
+    %23:vec4<f32> = call %tint_TextureSampleExternal, %20, %21, %22, %sampler_1, %coords_1
+    %24:texture_2d<f32> = load %texture_plane0
+    %25:texture_2d<f32> = load %texture_plane1
+    %26:tint_ExternalTextureParams = load %texture_params
+    %27:vec4<f32> = call %tint_TextureSampleExternal, %24, %25, %26, %sampler_1, %coords_1
+    %28:texture_2d<f32> = load %texture_plane0
+    %29:texture_2d<f32> = load %texture_plane1
+    %30:tint_ExternalTextureParams = load %texture_params
+    %result_a:vec4<f32> = call %foo, %28, %29, %30, %sampler_1, %coords_1
+    %result_b:vec4<f32> = call %foo, %28, %29, %30, %sampler_1, %coords_1
+    %33:vec4<f32> = add %result_a, %result_b
+    ret %33
   }
 }
 %tint_TextureSampleExternal = func(%plane_0:texture_2d<f32>, %plane_1:texture_2d<f32>, %params:tint_ExternalTextureParams, %sampler_2:sampler, %coords_2:vec2<f32>):vec4<f32> {  # %sampler_2: 'sampler', %coords_2: 'coords'
   $B4: {
-    %38:u32 = access %params, 1u
-    %39:mat3x4<f32> = access %params, 2u
-    %40:mat3x2<f32> = access %params, 6u
-    %41:vec2<f32> = access %params, 8u
-    %42:vec2<f32> = access %params, 9u
-    %43:vec2<f32> = access %params, 10u
-    %44:vec2<f32> = access %params, 11u
-    %45:vec3<f32> = construct %coords_2, 1.0f
-    %46:vec2<f32> = mul %40, %45
-    %47:vec2<f32> = clamp %46, %41, %42
-    %48:u32 = access %params, 0u
-    %49:bool = eq %48, 1u
-    %50:vec3<f32>, %51:f32 = if %49 [t: $B5, f: $B6] {  # if_1
+    %39:u32 = access %params, 1u
+    %40:mat3x4<f32> = access %params, 2u
+    %41:mat3x2<f32> = access %params, 6u
+    %42:vec2<f32> = access %params, 8u
+    %43:vec2<f32> = access %params, 9u
+    %44:vec2<f32> = access %params, 10u
+    %45:vec2<f32> = access %params, 11u
+    %46:vec3<f32> = construct %coords_2, 1.0f
+    %47:vec2<f32> = mul %41, %46
+    %48:vec2<f32> = clamp %47, %42, %43
+    %49:u32 = access %params, 0u
+    %50:bool = eq %49, 1u
+    %51:vec3<f32>, %52:f32 = if %50 [t: $B5, f: $B6] {  # if_1
       $B5: {  # true
-        %52:vec4<f32> = textureSampleLevel %plane_0, %sampler_2, %47, 0.0f
-        %53:vec3<f32> = swizzle %52, xyz
-        %54:f32 = access %52, 3u
-        exit_if %53, %54  # if_1
+        %53:vec4<f32> = textureSampleLevel %plane_0, %sampler_2, %48, 0.0f
+        %54:vec3<f32> = swizzle %53, xyz
+        %55:f32 = access %53, 3u
+        exit_if %54, %55  # if_1
       }
       $B6: {  # false
-        %55:vec4<f32> = textureSampleLevel %plane_0, %sampler_2, %47, 0.0f
-        %56:f32 = access %55, 0u
-        %57:vec2<f32> = clamp %46, %43, %44
-        %58:vec4<f32> = textureSampleLevel %plane_1, %sampler_2, %57, 0.0f
-        %59:vec2<f32> = swizzle %58, xy
-        %60:vec4<f32> = construct %56, %59, 1.0f
-        %61:vec3<f32> = mul %60, %39
-        exit_if %61, 1.0f  # if_1
+        %56:vec4<f32> = textureSampleLevel %plane_0, %sampler_2, %48, 0.0f
+        %57:f32 = access %56, 0u
+        %58:vec2<f32> = clamp %47, %44, %45
+        %59:vec4<f32> = textureSampleLevel %plane_1, %sampler_2, %58, 0.0f
+        %60:vec2<f32> = swizzle %59, xy
+        %61:vec4<f32> = construct %57, %60, 1.0f
+        %62:vec3<f32> = mul %61, %40
+        exit_if %62, 1.0f  # if_1
       }
     }
-    %62:bool = eq %38, 0u
-    %63:vec3<f32> = if %62 [t: $B7, f: $B8] {  # if_2
+    %63:bool = eq %39, 0u
+    %64:vec3<f32> = if %63 [t: $B7, f: $B8] {  # if_2
       $B7: {  # true
-        %64:tint_GammaTransferParams = access %params, 3u
-        %65:tint_GammaTransferParams = access %params, 4u
-        %66:mat3x3<f32> = access %params, 5u
-        %67:vec3<f32> = call %tint_GammaCorrection, %50, %64
-        %69:vec3<f32> = mul %66, %67
-        %70:vec3<f32> = call %tint_GammaCorrection, %69, %65
-        exit_if %70  # if_2
+        %65:tint_GammaTransferParams = access %params, 3u
+        %66:tint_GammaTransferParams = access %params, 4u
+        %67:mat3x3<f32> = access %params, 5u
+        %68:vec3<f32> = call %tint_GammaCorrection, %51, %65
+        %70:vec3<f32> = mul %67, %68
+        %71:vec3<f32> = call %tint_GammaCorrection, %70, %66
+        exit_if %71  # if_2
       }
       $B8: {  # false
-        exit_if %50  # if_2
+        exit_if %51  # if_2
       }
     }
-    %71:vec4<f32> = construct %63, %51
-    ret %71
+    %72:vec4<f32> = construct %64, %52
+    ret %72
   }
 }
 %tint_GammaCorrection = func(%v:vec3<f32>, %params_1:tint_GammaTransferParams):vec3<f32> {  # %params_1: 'params'
   $B9: {
-    %74:f32 = access %params_1, 0u
-    %75:f32 = access %params_1, 1u
-    %76:f32 = access %params_1, 2u
-    %77:f32 = access %params_1, 3u
-    %78:f32 = access %params_1, 4u
-    %79:f32 = access %params_1, 5u
-    %80:f32 = access %params_1, 6u
-    %81:vec3<f32> = construct %74
-    %82:vec3<f32> = construct %78
-    %83:vec3<f32> = abs %v
-    %84:vec3<f32> = sign %v
-    %85:vec3<bool> = lt %83, %82
-    %86:vec3<f32> = mul %77, %83
-    %87:vec3<f32> = add %86, %80
-    %88:vec3<f32> = mul %84, %87
-    %89:vec3<f32> = mul %75, %83
-    %90:vec3<f32> = add %89, %76
-    %91:vec3<f32> = pow %90, %81
-    %92:vec3<f32> = add %91, %79
-    %93:vec3<f32> = mul %84, %92
-    %94:vec3<f32> = select %93, %88, %85
-    ret %94
+    %75:f32 = access %params_1, 0u
+    %76:f32 = access %params_1, 1u
+    %77:f32 = access %params_1, 2u
+    %78:f32 = access %params_1, 3u
+    %79:f32 = access %params_1, 4u
+    %80:f32 = access %params_1, 5u
+    %81:f32 = access %params_1, 6u
+    %82:vec3<f32> = construct %75
+    %83:vec3<f32> = construct %79
+    %84:vec3<f32> = abs %v
+    %85:vec3<f32> = sign %v
+    %86:vec3<bool> = lt %84, %83
+    %87:vec3<f32> = mul %78, %84
+    %88:vec3<f32> = add %87, %81
+    %89:vec3<f32> = mul %85, %88
+    %90:vec3<f32> = mul %76, %84
+    %91:vec3<f32> = add %90, %77
+    %92:vec3<f32> = pow %91, %82
+    %93:vec3<f32> = add %92, %80
+    %94:vec3<f32> = mul %85, %93
+    %95:vec3<f32> = select %94, %89, %86
+    ret %95
   }
 }
 )";
diff --git a/src/tint/lang/core/type/depth_multisampled_texture.cc b/src/tint/lang/core/type/depth_multisampled_texture.cc
index 5daed0e..81bf7de 100644
--- a/src/tint/lang/core/type/depth_multisampled_texture.cc
+++ b/src/tint/lang/core/type/depth_multisampled_texture.cc
@@ -37,17 +37,10 @@
 TINT_INSTANTIATE_TYPEINFO(tint::core::type::DepthMultisampledTexture);
 
 namespace tint::core::type {
-namespace {
-
-bool IsValidDepthDimension(TextureDimension dim) {
-    return dim == TextureDimension::k2d;
-}
-
-}  // namespace
 
 DepthMultisampledTexture::DepthMultisampledTexture(TextureDimension dim)
     : Base(Hash(TypeCode::Of<DepthMultisampledTexture>().bits, dim), dim) {
-    TINT_ASSERT(IsValidDepthDimension(dim));
+    TINT_ASSERT(IsValidDimension(dim));
 }
 
 DepthMultisampledTexture::~DepthMultisampledTexture() = default;
@@ -69,4 +62,7 @@
     return ctx.dst.mgr->Get<DepthMultisampledTexture>(dim());
 }
 
+bool DepthMultisampledTexture::IsValidDimension(TextureDimension dim) {
+    return dim == TextureDimension::k2d;
+}
 }  // namespace tint::core::type
diff --git a/src/tint/lang/core/type/depth_multisampled_texture.h b/src/tint/lang/core/type/depth_multisampled_texture.h
index 4774ead..0771f8f 100644
--- a/src/tint/lang/core/type/depth_multisampled_texture.h
+++ b/src/tint/lang/core/type/depth_multisampled_texture.h
@@ -56,6 +56,9 @@
     /// @param ctx the clone context
     /// @returns a clone of this type
     DepthMultisampledTexture* Clone(CloneContext& ctx) const override;
+
+    /// @returns true if @p dim is a valid TextureDimension for a DepthMultisampledTexture
+    static bool IsValidDimension(TextureDimension dim);
 };
 
 }  // namespace tint::core::type
diff --git a/src/tint/lang/core/type/depth_texture.cc b/src/tint/lang/core/type/depth_texture.cc
index a246ed7..b6a2e2b 100644
--- a/src/tint/lang/core/type/depth_texture.cc
+++ b/src/tint/lang/core/type/depth_texture.cc
@@ -37,18 +37,10 @@
 TINT_INSTANTIATE_TYPEINFO(tint::core::type::DepthTexture);
 
 namespace tint::core::type {
-namespace {
-
-bool IsValidDepthDimension(TextureDimension dim) {
-    return dim == TextureDimension::k2d || dim == TextureDimension::k2dArray ||
-           dim == TextureDimension::kCube || dim == TextureDimension::kCubeArray;
-}
-
-}  // namespace
 
 DepthTexture::DepthTexture(TextureDimension dim)
     : Base(Hash(TypeCode::Of<DepthTexture>().bits, dim), dim) {
-    TINT_ASSERT(IsValidDepthDimension(dim));
+    TINT_ASSERT(IsValidDimension(dim));
 }
 
 DepthTexture::~DepthTexture() = default;
@@ -70,4 +62,9 @@
     return ctx.dst.mgr->Get<DepthTexture>(dim());
 }
 
+bool DepthTexture::IsValidDimension(TextureDimension dim) {
+    return dim == TextureDimension::k2d || dim == TextureDimension::k2dArray ||
+           dim == TextureDimension::kCube || dim == TextureDimension::kCubeArray;
+}
+
 }  // namespace tint::core::type
diff --git a/src/tint/lang/core/type/depth_texture.h b/src/tint/lang/core/type/depth_texture.h
index 2f1e738..62171a5 100644
--- a/src/tint/lang/core/type/depth_texture.h
+++ b/src/tint/lang/core/type/depth_texture.h
@@ -56,6 +56,9 @@
     /// @param ctx the clone context
     /// @returns a clone of this type
     DepthTexture* Clone(CloneContext& ctx) const override;
+
+    /// @returns true iff @p dim is a valid TextureDimension for a DepthTexture
+    static bool IsValidDimension(TextureDimension dim);
 };
 
 }  // namespace tint::core::type
diff --git a/src/tint/lang/msl/writer/raise/raise.cc b/src/tint/lang/msl/writer/raise/raise.cc
index fea8afd..09caa3c 100644
--- a/src/tint/lang/msl/writer/raise/raise.cc
+++ b/src/tint/lang/msl/writer/raise/raise.cc
@@ -29,6 +29,7 @@
 
 #include <utility>
 
+#include "src/tint/lang/core/ir/transform/array_length_from_uniform.h"
 #include "src/tint/lang/core/ir/transform/binary_polyfill.h"
 #include "src/tint/lang/core/ir/transform/binding_remapper.h"
 #include "src/tint/lang/core/ir/transform/builtin_polyfill.h"
@@ -93,6 +94,9 @@
     }
 
     RUN_TRANSFORM(core::ir::transform::MultiplanarExternalTexture, external_texture_options);
+    RUN_TRANSFORM(core::ir::transform::ArrayLengthFromUniform,
+                  BindingPoint{0u, array_length_from_uniform_options.ubo_binding},
+                  array_length_from_uniform_options.bindpoint_to_size_index);
 
     if (!options.disable_workgroup_init) {
         RUN_TRANSFORM(core::ir::transform::ZeroInitWorkgroupMemory);
diff --git a/src/tint/lang/wgsl/ast/BUILD.bazel b/src/tint/lang/wgsl/ast/BUILD.bazel
index b45a457..be215b2 100644
--- a/src/tint/lang/wgsl/ast/BUILD.bazel
+++ b/src/tint/lang/wgsl/ast/BUILD.bazel
@@ -81,6 +81,7 @@
     "if_statement.cc",
     "increment_decrement_statement.cc",
     "index_accessor_expression.cc",
+    "input_attachment_index_attribute.cc",
     "int_literal_expression.cc",
     "internal_attribute.cc",
     "interpolate_attribute.cc",
@@ -161,6 +162,7 @@
     "if_statement.h",
     "increment_decrement_statement.h",
     "index_accessor_expression.h",
+    "input_attachment_index_attribute.h",
     "int_literal_expression.h",
     "internal_attribute.h",
     "interpolate_attribute.h",
@@ -267,6 +269,7 @@
     "if_statement_test.cc",
     "increment_decrement_statement_test.cc",
     "index_accessor_expression_test.cc",
+    "input_attachment_index_attribute_test.cc",
     "int_literal_expression_test.cc",
     "interpolate_attribute_test.cc",
     "location_attribute_test.cc",
diff --git a/src/tint/lang/wgsl/ast/BUILD.cmake b/src/tint/lang/wgsl/ast/BUILD.cmake
index c144385..3d36157 100644
--- a/src/tint/lang/wgsl/ast/BUILD.cmake
+++ b/src/tint/lang/wgsl/ast/BUILD.cmake
@@ -125,6 +125,8 @@
   lang/wgsl/ast/increment_decrement_statement.h
   lang/wgsl/ast/index_accessor_expression.cc
   lang/wgsl/ast/index_accessor_expression.h
+  lang/wgsl/ast/input_attachment_index_attribute.cc
+  lang/wgsl/ast/input_attachment_index_attribute.h
   lang/wgsl/ast/int_literal_expression.cc
   lang/wgsl/ast/int_literal_expression.h
   lang/wgsl/ast/internal_attribute.cc
@@ -267,6 +269,7 @@
   lang/wgsl/ast/if_statement_test.cc
   lang/wgsl/ast/increment_decrement_statement_test.cc
   lang/wgsl/ast/index_accessor_expression_test.cc
+  lang/wgsl/ast/input_attachment_index_attribute_test.cc
   lang/wgsl/ast/int_literal_expression_test.cc
   lang/wgsl/ast/interpolate_attribute_test.cc
   lang/wgsl/ast/location_attribute_test.cc
diff --git a/src/tint/lang/wgsl/ast/BUILD.gn b/src/tint/lang/wgsl/ast/BUILD.gn
index 249bb50..165f2a7 100644
--- a/src/tint/lang/wgsl/ast/BUILD.gn
+++ b/src/tint/lang/wgsl/ast/BUILD.gn
@@ -128,6 +128,8 @@
     "increment_decrement_statement.h",
     "index_accessor_expression.cc",
     "index_accessor_expression.h",
+    "input_attachment_index_attribute.cc",
+    "input_attachment_index_attribute.h",
     "int_literal_expression.cc",
     "int_literal_expression.h",
     "internal_attribute.cc",
@@ -267,6 +269,7 @@
       "if_statement_test.cc",
       "increment_decrement_statement_test.cc",
       "index_accessor_expression_test.cc",
+      "input_attachment_index_attribute_test.cc",
       "int_literal_expression_test.cc",
       "interpolate_attribute_test.cc",
       "location_attribute_test.cc",
diff --git a/src/tint/lang/wgsl/ast/builder.h b/src/tint/lang/wgsl/ast/builder.h
index c4eb4f1..b643fc4 100644
--- a/src/tint/lang/wgsl/ast/builder.h
+++ b/src/tint/lang/wgsl/ast/builder.h
@@ -46,6 +46,7 @@
 #include "src/tint/lang/core/type/f16.h"
 #include "src/tint/lang/core/type/f32.h"
 #include "src/tint/lang/core/type/i32.h"
+#include "src/tint/lang/core/type/input_attachment.h"
 #include "src/tint/lang/core/type/matrix.h"
 #include "src/tint/lang/core/type/multisampled_texture.h"
 #include "src/tint/lang/core/type/pointer.h"
@@ -86,6 +87,7 @@
 #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/index_accessor_expression.h"
+#include "src/tint/lang/wgsl/ast/input_attachment_index_attribute.h"
 #include "src/tint/lang/wgsl/ast/int_literal_expression.h"
 #include "src/tint/lang/wgsl/ast/interpolate_attribute.h"
 #include "src/tint/lang/wgsl/ast/invariant_attribute.h"
@@ -1297,6 +1299,12 @@
         /// @returns the external texture
         ast::Type external_texture() const { return (*this)("texture_external"); }
 
+        /// @param subtype the texture subtype.
+        /// @returns the input attachment
+        ast::Type input_attachment(ast::Type subtype) const {
+            return (*this)("input_attachment", subtype);
+        }
+
         /// @param source the Source of the node
         /// @returns the external texture
         ast::Type external_texture(const Source& source) const {
@@ -3240,6 +3248,24 @@
         return create<ast::IdAttribute>(Expr(std::forward<EXPR>(id)));
     }
 
+    /// Creates an ast::InputAttachmentIndexAttribute
+    /// @param index the index value expression
+    /// @returns the index attribute pointer
+    template <typename EXPR>
+    const ast::InputAttachmentIndexAttribute* InputAttachmentIndex(EXPR&& index) {
+        return create<ast::InputAttachmentIndexAttribute>(source_, Expr(std::forward<EXPR>(index)));
+    }
+
+    /// Creates an ast::InputAttachmentIndexAttribute
+    /// @param source the source information
+    /// @param index the index value expression
+    /// @returns the index attribute pointer
+    template <typename EXPR>
+    const ast::InputAttachmentIndexAttribute* InputAttachmentIndex(const Source& source,
+                                                                   EXPR&& index) {
+        return create<ast::InputAttachmentIndexAttribute>(source, Expr(std::forward<EXPR>(index)));
+    }
+
     /// Creates an ast::StageAttribute
     /// @param source the source information
     /// @param stage the pipeline stage
diff --git a/src/tint/lang/wgsl/ast/input_attachment_index_attribute.cc b/src/tint/lang/wgsl/ast/input_attachment_index_attribute.cc
new file mode 100644
index 0000000..5441ec4
--- /dev/null
+++ b/src/tint/lang/wgsl/ast/input_attachment_index_attribute.cc
@@ -0,0 +1,58 @@
+// 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/ast/input_attachment_index_attribute.h"
+
+#include <string>
+
+#include "src/tint/lang/wgsl/ast/builder.h"
+#include "src/tint/lang/wgsl/ast/clone_context.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ast::InputAttachmentIndexAttribute);
+
+namespace tint::ast {
+
+InputAttachmentIndexAttribute::InputAttachmentIndexAttribute(GenerationID pid,
+                                                             NodeID nid,
+                                                             const Source& src,
+                                                             const Expression* exp)
+    : Base(pid, nid, src), expr(exp) {}
+
+InputAttachmentIndexAttribute::~InputAttachmentIndexAttribute() = default;
+
+std::string InputAttachmentIndexAttribute::Name() const {
+    return "input_attachment_index";
+}
+
+const InputAttachmentIndexAttribute* InputAttachmentIndexAttribute::Clone(CloneContext& ctx) const {
+    // Clone arguments outside of create() call to have deterministic ordering
+    auto src = ctx.Clone(source);
+    auto* expr_ = ctx.Clone(expr);
+    return ctx.dst->create<InputAttachmentIndexAttribute>(src, expr_);
+}
+
+}  // namespace tint::ast
diff --git a/src/tint/lang/wgsl/ast/input_attachment_index_attribute.h b/src/tint/lang/wgsl/ast/input_attachment_index_attribute.h
new file mode 100644
index 0000000..3477070
--- /dev/null
+++ b/src/tint/lang/wgsl/ast/input_attachment_index_attribute.h
@@ -0,0 +1,72 @@
+// 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.
+
+#ifndef SRC_TINT_LANG_WGSL_AST_INPUT_ATTACHMENT_INDEX_ATTRIBUTE_H_
+#define SRC_TINT_LANG_WGSL_AST_INPUT_ATTACHMENT_INDEX_ATTRIBUTE_H_
+
+#include <string>
+
+#include "src/tint/lang/wgsl/ast/attribute.h"
+
+// Forward declarations
+namespace tint::ast {
+class Expression;
+}
+
+namespace tint::ast {
+
+/// An input attachment attribute (enabled with the input attachments extension)
+class InputAttachmentIndexAttribute final
+    : public Castable<InputAttachmentIndexAttribute, Attribute> {
+  public:
+    /// constructor
+    /// @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 expr the input attachment index value
+    InputAttachmentIndexAttribute(GenerationID pid,
+                                  NodeID nid,
+                                  const Source& src,
+                                  const Expression* expr);
+    ~InputAttachmentIndexAttribute() override;
+
+    /// @returns the WGSL name for the attribute
+    std::string Name() const override;
+
+    /// Clones this node and all transitive child nodes using the `CloneContext`
+    /// `ctx`.
+    /// @param ctx the clone context
+    /// @return the newly cloned node
+    const InputAttachmentIndexAttribute* Clone(CloneContext& ctx) const override;
+
+    /// The index value expression
+    const Expression* const expr;
+};
+
+}  // namespace tint::ast
+
+#endif  // SRC_TINT_LANG_WGSL_AST_INPUT_ATTACHMENT_INDEX_ATTRIBUTE_H_
diff --git a/src/tint/lang/wgsl/ast/input_attachment_index_attribute_test.cc b/src/tint/lang/wgsl/ast/input_attachment_index_attribute_test.cc
new file mode 100644
index 0000000..28ce19d
--- /dev/null
+++ b/src/tint/lang/wgsl/ast/input_attachment_index_attribute_test.cc
@@ -0,0 +1,45 @@
+// 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/ast/input_attachment_index_attribute.h"
+#include "src/tint/lang/wgsl/ast/helper_test.h"
+
+using namespace tint::core::number_suffixes;  // NOLINT
+
+namespace tint::ast {
+namespace {
+
+using InputAttachmentIndexAttributeTest = TestHelper;
+
+TEST_F(InputAttachmentIndexAttributeTest, Creation) {
+    auto* expr = Expr(1_u);
+    auto* c = InputAttachmentIndex(expr);
+    EXPECT_EQ(c->expr, expr);
+}
+
+}  // namespace
+}  // namespace tint::ast
diff --git a/src/tint/lang/wgsl/ast/transform/multiplanar_external_texture.cc b/src/tint/lang/wgsl/ast/transform/multiplanar_external_texture.cc
index 0b59b7d..c9b1313 100644
--- a/src/tint/lang/wgsl/ast/transform/multiplanar_external_texture.cc
+++ b/src/tint/lang/wgsl/ast/transform/multiplanar_external_texture.cc
@@ -221,13 +221,12 @@
 
         // Transform the external texture builtin calls into calls to the external texture
         // functions.
-        ctx.ReplaceAll([&](const CallExpression* expr) -> const CallExpression* {
+        ctx.ReplaceAll([&](const CallExpression* expr) -> const Expression* {
             auto* call = sem.Get(expr)->UnwrapMaterialize()->As<sem::Call>();
             auto* builtin = call->Target()->As<sem::BuiltinFn>();
 
             if (builtin && !builtin->Parameters().IsEmpty() &&
-                builtin->Parameters()[0]->Type()->Is<core::type::ExternalTexture>() &&
-                builtin->Fn() != wgsl::BuiltinFn::kTextureDimensions) {
+                builtin->Parameters()[0]->Type()->Is<core::type::ExternalTexture>()) {
                 if (auto* var_user =
                         sem.GetVal(expr->args[0])->UnwrapLoad()->As<sem::VariableUser>()) {
                     auto it = new_binding_symbols.find(var_user->Variable());
@@ -245,6 +244,8 @@
                             return createTextureLoad(call, syms);
                         case wgsl::BuiltinFn::kTextureSampleBaseClampToEdge:
                             return createTextureSampleBaseClampToEdge(expr, syms);
+                        case wgsl::BuiltinFn::kTextureDimensions:
+                            return createTextureDimensions(call, syms);
                         default:
                             break;
                     }
@@ -553,6 +554,19 @@
         return b.Call(texture_load_external_sym, plane_0_binding_arg, syms.plane_1,
                       ctx.Clone(args[1]->Declaration()), syms.params);
     }
+
+    /// Returns the expression used to replace a textureDimensions call.
+    /// @param call the call expression being transformed
+    /// @param syms the expanded symbols to be used in the new call
+    /// @returns a load of params.visibleSize
+    const Expression* createTextureDimensions(const sem::Call* call, NewBindingSymbols syms) {
+        if (TINT_UNLIKELY(call->Arguments().Length() != 1)) {
+            TINT_ICE() << "expected textureDimensions call with a texture_external to have 1 "
+                          "arguments, found "
+                       << call->Arguments().Length() << " arguments";
+        }
+        return b.Add(b.MemberAccessor(syms.params, "visibleSize"), b.Call<vec2<u32>>(1_a));
+    }
 };
 
 MultiplanarExternalTexture::NewBindingPoints::NewBindingPoints() = default;
diff --git a/src/tint/lang/wgsl/ast/transform/multiplanar_external_texture_test.cc b/src/tint/lang/wgsl/ast/transform/multiplanar_external_texture_test.cc
index 65c7b41..a16eb09 100644
--- a/src/tint/lang/wgsl/ast/transform/multiplanar_external_texture_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/multiplanar_external_texture_test.cc
@@ -171,7 +171,7 @@
 @fragment
 fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
   var dim : vec2<u32>;
-  dim = textureDimensions(ext_tex);
+  dim = (ext_tex_params.visibleSize + vec2<u32>(1));
   return vec4<f32>(0.0, 0.0, 0.0, 0.0);
 }
 )";
@@ -318,7 +318,7 @@
 @fragment
 fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
   var dim : vec2<u32>;
-  dim = textureDimensions(ext_tex);
+  dim = (ext_tex_params.visibleSize + vec2<u32>(1));
   return vec4<f32>(0.0, 0.0, 0.0, 0.0);
 }
 
@@ -1791,7 +1791,7 @@
 }
 
 fn f(ext_tex : texture_2d<f32>, ext_tex_plane_1 : texture_2d<f32>, ext_tex_params : ExternalTextureParams) -> vec2<u32> {
-  return textureDimensions(ext_tex);
+  return (ext_tex_params.visibleSize + vec2<u32>(1));
 }
 )";
 
diff --git a/src/tint/lang/wgsl/ast/transform/renamer_test.cc b/src/tint/lang/wgsl/ast/transform/renamer_test.cc
index b3aa2ac..6912b2e 100644
--- a/src/tint/lang/wgsl/ast/transform/renamer_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/renamer_test.cc
@@ -1866,7 +1866,8 @@
     std::vector<std::string_view> out;
     for (auto type : core::kBuiltinTypeStrings) {
         if (type != "ptr" && type != "atomic" && !tint::HasPrefix(type, "sampler") &&
-            !tint::HasPrefix(type, "texture") && !tint::HasPrefix(type, "__")) {
+            !tint::HasPrefix(type, "texture") && !tint::HasPrefix(type, "__") &&
+            !tint::HasPrefix(type, "input_attachment")) {
             out.push_back(type);
         }
     }
diff --git a/src/tint/lang/wgsl/reader/parser/error_resync_test.cc b/src/tint/lang/wgsl/reader/parser/error_resync_test.cc
index 54818ea..f5d39eb 100644
--- a/src/tint/lang/wgsl/reader/parser/error_resync_test.cc
+++ b/src/tint/lang/wgsl/reader/parser/error_resync_test.cc
@@ -67,7 +67,7 @@
      ^
 
 test.wgsl:4:2 error: expected attribute
-Possible values: 'align', 'binding', 'blend_src', 'builtin', 'color', 'compute', 'diagnostic', 'fragment', 'group', 'id', 'interpolate', 'invariant', 'location', 'must_use', 'size', 'vertex', 'workgroup_size'
+Possible values: 'align', 'binding', 'blend_src', 'builtin', 'color', 'compute', 'diagnostic', 'fragment', 'group', 'id', 'input_attachment_index', 'interpolate', 'invariant', 'location', 'must_use', 'size', 'vertex', 'workgroup_size'
 @_ fn -> {}
  ^
 )");
@@ -135,7 +135,7 @@
          ^^^^
 
 test.wgsl:7:6 error: expected attribute
-Possible values: 'align', 'binding', 'blend_src', 'builtin', 'color', 'compute', 'diagnostic', 'fragment', 'group', 'id', 'interpolate', 'invariant', 'location', 'must_use', 'size', 'vertex', 'workgroup_size'
+Possible values: 'align', 'binding', 'blend_src', 'builtin', 'color', 'compute', 'diagnostic', 'fragment', 'group', 'id', 'input_attachment_index', 'interpolate', 'invariant', 'location', 'must_use', 'size', 'vertex', 'workgroup_size'
     @- x : i32,
      ^
 )");
diff --git a/src/tint/lang/wgsl/reader/parser/function_attribute_list_test.cc b/src/tint/lang/wgsl/reader/parser/function_attribute_list_test.cc
index 3e66246..99e5cb6 100644
--- a/src/tint/lang/wgsl/reader/parser/function_attribute_list_test.cc
+++ b/src/tint/lang/wgsl/reader/parser/function_attribute_list_test.cc
@@ -66,7 +66,7 @@
     EXPECT_TRUE(attrs.value.IsEmpty());
     EXPECT_EQ(p->error(), R"(1:2: expected attribute
 Did you mean 'invariant'?
-Possible values: 'align', 'binding', 'blend_src', 'builtin', 'color', 'compute', 'diagnostic', 'fragment', 'group', 'id', 'interpolate', 'invariant', 'location', 'must_use', 'size', 'vertex', 'workgroup_size')");
+Possible values: 'align', 'binding', 'blend_src', 'builtin', 'color', 'compute', 'diagnostic', 'fragment', 'group', 'id', 'input_attachment_index', 'interpolate', 'invariant', 'location', 'must_use', 'size', 'vertex', 'workgroup_size')");
 }
 
 }  // namespace
diff --git a/src/tint/lang/wgsl/reader/parser/variable_attribute_list_test.cc b/src/tint/lang/wgsl/reader/parser/variable_attribute_list_test.cc
index 7b82086..c1361b7 100644
--- a/src/tint/lang/wgsl/reader/parser/variable_attribute_list_test.cc
+++ b/src/tint/lang/wgsl/reader/parser/variable_attribute_list_test.cc
@@ -64,7 +64,7 @@
     EXPECT_TRUE(attrs.value.IsEmpty());
     EXPECT_EQ(p->error(), R"(1:2: expected attribute
 Did you mean 'invariant'?
-Possible values: 'align', 'binding', 'blend_src', 'builtin', 'color', 'compute', 'diagnostic', 'fragment', 'group', 'id', 'interpolate', 'invariant', 'location', 'must_use', 'size', 'vertex', 'workgroup_size')");
+Possible values: 'align', 'binding', 'blend_src', 'builtin', 'color', 'compute', 'diagnostic', 'fragment', 'group', 'id', 'input_attachment_index', 'interpolate', 'invariant', 'location', 'must_use', 'size', 'vertex', 'workgroup_size')");
 }
 
 }  // namespace
diff --git a/src/tint/lang/wgsl/resolver/resolver.cc b/src/tint/lang/wgsl/resolver/resolver.cc
index 34d3af3..0a17841 100644
--- a/src/tint/lang/wgsl/resolver/resolver.cc
+++ b/src/tint/lang/wgsl/resolver/resolver.cc
@@ -46,6 +46,7 @@
 #include "src/tint/lang/core/type/depth_multisampled_texture.h"
 #include "src/tint/lang/core/type/depth_texture.h"
 #include "src/tint/lang/core/type/external_texture.h"
+#include "src/tint/lang/core/type/input_attachment.h"
 #include "src/tint/lang/core/type/memory_view.h"
 #include "src/tint/lang/core/type/multisampled_texture.h"
 #include "src/tint/lang/core/type/pointer.h"
@@ -2654,6 +2655,8 @@
             return StorageTexture(ident, core::type::TextureDimension::k2dArray);
         case core::BuiltinType::kTextureStorage3D:
             return StorageTexture(ident, core::type::TextureDimension::k3d);
+        case core::BuiltinType::kInputAttachment:
+            return InputAttachment(ident);
         case core::BuiltinType::kPackedVec3:
             return PackedVec3T(ident);
         case core::BuiltinType::kAtomicCompareExchangeResultI32:
@@ -2968,6 +2971,21 @@
     return tex;
 }
 
+core::type::InputAttachment* Resolver::InputAttachment(const ast::Identifier* ident) {
+    auto* tmpl_ident = TemplatedIdentifier(ident, 1);
+    if (TINT_UNLIKELY(!tmpl_ident)) {
+        return nullptr;
+    }
+
+    auto* ty_expr = sem_.GetType(tmpl_ident->arguments[0]);
+    if (TINT_UNLIKELY(!ty_expr)) {
+        return nullptr;
+    }
+
+    auto* out = b.create<core::type::InputAttachment>(ty_expr);
+    return validator_.InputAttachment(out, ident->source) ? out : nullptr;
+}
+
 core::type::Vector* Resolver::PackedVec3T(const ast::Identifier* ident) {
     auto* tmpl_ident = TemplatedIdentifier(ident, 1);
     if (TINT_UNLIKELY(!tmpl_ident)) {
diff --git a/src/tint/lang/wgsl/resolver/resolver.h b/src/tint/lang/wgsl/resolver/resolver.h
index 715d62a..df18f91 100644
--- a/src/tint/lang/wgsl/resolver/resolver.h
+++ b/src/tint/lang/wgsl/resolver/resolver.h
@@ -223,6 +223,9 @@
     core::type::StorageTexture* StorageTexture(const ast::Identifier* ident,
                                                core::type::TextureDimension dim);
 
+    /// @returns an input attachment resolved from the templated identifier @p ident
+    core::type::InputAttachment* InputAttachment(const ast::Identifier* ident);
+
     /// @returns a packed vec3 resolved from the templated identifier @p ident.
     core::type::Vector* PackedVec3T(const ast::Identifier* ident);
 
diff --git a/src/tint/lang/wgsl/resolver/validator.cc b/src/tint/lang/wgsl/resolver/validator.cc
index 155aaf1..f482497 100644
--- a/src/tint/lang/wgsl/resolver/validator.cc
+++ b/src/tint/lang/wgsl/resolver/validator.cc
@@ -441,6 +441,15 @@
     return true;
 }
 
+bool Validator::InputAttachment(const core::type::InputAttachment* t, const Source& source) const {
+    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;
+    }
+
+    return true;
+}
+
 bool Validator::Materialize(const core::type::Type* to,
                             const core::type::Type* from,
                             const Source& source) const {
diff --git a/src/tint/lang/wgsl/resolver/validator.h b/src/tint/lang/wgsl/resolver/validator.h
index 0fa3bd9..6cdbfaf 100644
--- a/src/tint/lang/wgsl/resolver/validator.h
+++ b/src/tint/lang/wgsl/resolver/validator.h
@@ -428,6 +428,12 @@
     /// @returns true on success, false otherwise
     bool MultisampledTexture(const core::type::MultisampledTexture* t, const Source& source) const;
 
+    /// Validates a input attachment
+    /// @param t the input attachment to validate
+    /// @param source the source of the input attachment
+    /// @returns true on success, false otherwise
+    bool InputAttachment(const core::type::InputAttachment* t, const Source& source) const;
+
     /// Validates a structure
     /// @param str the structure to validate
     /// @param stage the current pipeline stage