Import Tint changes from Dawn
Changes:
- 1b0266e46488b89b0a98ccb9c538d80cf7436c91 [tint][ir] Add Decode() overload that takes a protobuf st... by Ben Clayton <bclayton@google.com>
- efa2c60efb27f761a0c6e265c39d84d8371ecf28 [tint][type] Move private validity checks to static membe... by Ben Clayton <bclayton@google.com>
- f403412d25f8dd5d8f5e7c17232e53232c78ea78 [tint] Add constant::Invalid by Ben Clayton <bclayton@google.com>
- a6788bdfee90a790c2e4991ba795e2a8a61635eb [tint][core] Test that OOB Splat::Index() returns nullptr by Ben Clayton <bclayton@google.com>
- e4e461377ef4d02328f0f20b1ef79897f14d75b0 Tint: Add input_attachment as builtin type. by Le Hoang Quyen <lehoangquyen@chromium.org>
- 59eb7f31f5507fbddcfead963581523c78e69b8f Tint: Add input_attachment_index attribute class and def ... by Le Hoang Quyen <lehoangquyen@chromium.org>
- 81da61165530deb3049ad38a54309278c3565413 Make textureDimensions(texture_external) use the visible ... by Corentin Wallez <cwallez@chromium.org>
- 4e043e9741d726a213b588a369d0c5623370c910 [msl] Add ArrayLengthFromUniform transform by James Price <jrprice@google.com>
GitOrigin-RevId: 1b0266e46488b89b0a98ccb9c538d80cf7436c91
Change-Id: Idf601440a3444bcda88f60ebc28becf157d7d013
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/189661
Commit-Queue: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
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
