Tint: Add inputAttachmentLoad to core.def, wgsl.def, IR
Bug: 341117913
Change-Id: I2f67f72458deed4a6764e79b68181ce15f2c13ca
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/189860
Reviewed-by: dan sinclair <dsinclair@google.com>
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Auto-Submit: Quyen Le <lehoangquyen@chromium.org>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/cmd/fuzz/wgsl/dictionary.txt b/src/tint/cmd/fuzz/wgsl/dictionary.txt
index ffd821a..4f9d148 100644
--- a/src/tint/cmd/fuzz/wgsl/dictionary.txt
+++ b/src/tint/cmd/fuzz/wgsl/dictionary.txt
@@ -232,6 +232,7 @@
"id"
"if"
"info"
+"inputAttachmentLoad"
"input_attachment"
"input_attachment_index"
"insertBits"
diff --git a/src/tint/lang/core/builtin_fn.cc b/src/tint/lang/core/builtin_fn.cc
index 82e1e37..ea202f1 100644
--- a/src/tint/lang/core/builtin_fn.cc
+++ b/src/tint/lang/core/builtin_fn.cc
@@ -363,6 +363,9 @@
if (name == "textureLoad") {
return BuiltinFn::kTextureLoad;
}
+ if (name == "inputAttachmentLoad") {
+ return BuiltinFn::kInputAttachmentLoad;
+ }
if (name == "atomicLoad") {
return BuiltinFn::kAtomicLoad;
}
@@ -625,6 +628,8 @@
return "textureStore";
case BuiltinFn::kTextureLoad:
return "textureLoad";
+ case BuiltinFn::kInputAttachmentLoad:
+ return "inputAttachmentLoad";
case BuiltinFn::kAtomicLoad:
return "atomicLoad";
case BuiltinFn::kAtomicStore:
diff --git a/src/tint/lang/core/builtin_fn.h b/src/tint/lang/core/builtin_fn.h
index 7a87054..5c6cf01 100644
--- a/src/tint/lang/core/builtin_fn.h
+++ b/src/tint/lang/core/builtin_fn.h
@@ -155,6 +155,7 @@
kTextureSampleBaseClampToEdge,
kTextureStore,
kTextureLoad,
+ kInputAttachmentLoad,
kAtomicLoad,
kAtomicStore,
kAtomicAdd,
@@ -298,6 +299,7 @@
BuiltinFn::kTextureSampleBaseClampToEdge,
BuiltinFn::kTextureStore,
BuiltinFn::kTextureLoad,
+ BuiltinFn::kInputAttachmentLoad,
BuiltinFn::kAtomicLoad,
BuiltinFn::kAtomicStore,
BuiltinFn::kAtomicAdd,
@@ -423,6 +425,7 @@
"textureSampleBaseClampToEdge",
"textureStore",
"textureLoad",
+ "inputAttachmentLoad",
"atomicLoad",
"atomicStore",
"atomicAdd",
diff --git a/src/tint/lang/core/core.def b/src/tint/lang/core/core.def
index 175f61c..26907a3 100644
--- a/src/tint/lang/core/core.def
+++ b/src/tint/lang/core/core.def
@@ -270,6 +270,7 @@
type texture_storage_2d_array<F: texel_format, A: access>
type texture_storage_3d<F: texel_format, A: access>
type texture_external
+type input_attachment<T>
@display("__modf_result_{T}") type __modf_result<T>
@display("__modf_result_vec{N}_{T}") type __modf_result_vec<N: num, T>
@@ -668,6 +669,8 @@
@must_use fn textureLoad[C: iu32](texture: texture_storage_3d<i32_texel_format, readable>, coords: vec3<C>) -> vec4<i32>
@must_use fn textureLoad[C: iu32](texture: texture_storage_3d<u32_texel_format, readable>, coords: vec3<C>) -> vec4<u32>
+@stage("fragment") fn inputAttachmentLoad[T: fiu32](input_attachment: input_attachment<T>) -> vec4<T>
+
@stage("fragment", "compute") fn atomicLoad[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>) -> T
@stage("fragment", "compute") fn atomicStore[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, T)
@stage("fragment", "compute") fn atomicAdd[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T
diff --git a/src/tint/lang/core/intrinsic/data.cc b/src/tint/lang/core/intrinsic/data.cc
index 8165d76..eaafd91 100644
--- a/src/tint/lang/core/intrinsic/data.cc
+++ b/src/tint/lang/core/intrinsic/data.cc
@@ -868,6 +868,26 @@
};
+/// TypeMatcher for 'type input_attachment'
+constexpr TypeMatcher kInputAttachmentMatcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+ const Type* T = nullptr;
+ if (!MatchInputAttachment(state, ty, T)) {
+ return nullptr;
+ }
+ T = state.Type(T);
+ if (T == nullptr) {
+ return nullptr;
+ }
+ return BuildInputAttachment(state, ty, T);
+ },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
+ state->PrintType(T);
+ out << style::Type("input_attachment", "<", T, ">");
+ }
+};
+
+
/// TypeMatcher for 'type __modf_result'
constexpr TypeMatcher kModfResultMatcher {
/* match */ [](MatchState& state, const Type* ty) -> const Type* {
@@ -1461,23 +1481,24 @@
/* [42] */ kTextureStorage2DArrayMatcher,
/* [43] */ kTextureStorage3DMatcher,
/* [44] */ kTextureExternalMatcher,
- /* [45] */ kModfResultMatcher,
- /* [46] */ kModfResultVecMatcher,
- /* [47] */ kFrexpResultMatcher,
- /* [48] */ kFrexpResultVecMatcher,
- /* [49] */ kAtomicCompareExchangeResultMatcher,
- /* [50] */ kScalarMatcher,
- /* [51] */ kScalarNoF32Matcher,
- /* [52] */ kScalarNoF16Matcher,
- /* [53] */ kScalarNoI32Matcher,
- /* [54] */ kScalarNoU32Matcher,
- /* [55] */ kScalarNoBoolMatcher,
- /* [56] */ kFiu32F16Matcher,
- /* [57] */ kFiu32Matcher,
- /* [58] */ kFi32F16Matcher,
- /* [59] */ kFi32Matcher,
- /* [60] */ kF32F16Matcher,
- /* [61] */ kIu32Matcher,
+ /* [45] */ kInputAttachmentMatcher,
+ /* [46] */ kModfResultMatcher,
+ /* [47] */ kModfResultVecMatcher,
+ /* [48] */ kFrexpResultMatcher,
+ /* [49] */ kFrexpResultVecMatcher,
+ /* [50] */ kAtomicCompareExchangeResultMatcher,
+ /* [51] */ kScalarMatcher,
+ /* [52] */ kScalarNoF32Matcher,
+ /* [53] */ kScalarNoF16Matcher,
+ /* [54] */ kScalarNoI32Matcher,
+ /* [55] */ kScalarNoU32Matcher,
+ /* [56] */ kScalarNoBoolMatcher,
+ /* [57] */ kFiu32F16Matcher,
+ /* [58] */ kFiu32Matcher,
+ /* [59] */ kFi32F16Matcher,
+ /* [60] */ kFi32Matcher,
+ /* [61] */ kF32F16Matcher,
+ /* [62] */ kIu32Matcher,
};
/// The template numbers, and number matchers
@@ -1550,13 +1571,13 @@
/* [47] */ MatcherIndex(21),
/* [48] */ MatcherIndex(0),
/* [49] */ MatcherIndex(7),
- /* [50] */ MatcherIndex(48),
+ /* [50] */ MatcherIndex(49),
/* [51] */ MatcherIndex(0),
/* [52] */ MatcherIndex(1),
/* [53] */ MatcherIndex(21),
/* [54] */ MatcherIndex(0),
/* [55] */ MatcherIndex(2),
- /* [56] */ MatcherIndex(46),
+ /* [56] */ MatcherIndex(47),
/* [57] */ MatcherIndex(0),
/* [58] */ MatcherIndex(1),
/* [59] */ MatcherIndex(40),
@@ -1662,9 +1683,9 @@
/* [159] */ MatcherIndex(21),
/* [160] */ MatcherIndex(1),
/* [161] */ MatcherIndex(6),
- /* [162] */ MatcherIndex(47),
+ /* [162] */ MatcherIndex(48),
/* [163] */ MatcherIndex(0),
- /* [164] */ MatcherIndex(45),
+ /* [164] */ MatcherIndex(46),
/* [165] */ MatcherIndex(0),
/* [166] */ MatcherIndex(11),
/* [167] */ MatcherIndex(7),
@@ -1700,83 +1721,85 @@
/* [197] */ MatcherIndex(7),
/* [198] */ MatcherIndex(33),
/* [199] */ MatcherIndex(7),
- /* [200] */ MatcherIndex(49),
+ /* [200] */ MatcherIndex(45),
/* [201] */ MatcherIndex(0),
- /* [202] */ MatcherIndex(11),
- /* [203] */ MatcherIndex(1),
- /* [204] */ MatcherIndex(12),
- /* [205] */ MatcherIndex(0),
+ /* [202] */ MatcherIndex(50),
+ /* [203] */ MatcherIndex(0),
+ /* [204] */ MatcherIndex(11),
+ /* [205] */ MatcherIndex(1),
/* [206] */ MatcherIndex(12),
- /* [207] */ MatcherIndex(7),
+ /* [207] */ MatcherIndex(0),
/* [208] */ MatcherIndex(12),
- /* [209] */ MatcherIndex(8),
- /* [210] */ MatcherIndex(13),
- /* [211] */ MatcherIndex(0),
+ /* [209] */ MatcherIndex(7),
+ /* [210] */ MatcherIndex(12),
+ /* [211] */ MatcherIndex(8),
/* [212] */ MatcherIndex(13),
- /* [213] */ MatcherIndex(7),
+ /* [213] */ MatcherIndex(0),
/* [214] */ MatcherIndex(13),
- /* [215] */ MatcherIndex(8),
- /* [216] */ MatcherIndex(14),
- /* [217] */ MatcherIndex(0),
+ /* [215] */ MatcherIndex(7),
+ /* [216] */ MatcherIndex(13),
+ /* [217] */ MatcherIndex(8),
/* [218] */ MatcherIndex(14),
- /* [219] */ MatcherIndex(7),
+ /* [219] */ MatcherIndex(0),
/* [220] */ MatcherIndex(14),
- /* [221] */ MatcherIndex(8),
- /* [222] */ MatcherIndex(15),
- /* [223] */ MatcherIndex(0),
+ /* [221] */ MatcherIndex(7),
+ /* [222] */ MatcherIndex(14),
+ /* [223] */ MatcherIndex(8),
/* [224] */ MatcherIndex(15),
- /* [225] */ MatcherIndex(7),
+ /* [225] */ MatcherIndex(0),
/* [226] */ MatcherIndex(15),
- /* [227] */ MatcherIndex(8),
- /* [228] */ MatcherIndex(16),
- /* [229] */ MatcherIndex(0),
+ /* [227] */ MatcherIndex(7),
+ /* [228] */ MatcherIndex(15),
+ /* [229] */ MatcherIndex(8),
/* [230] */ MatcherIndex(16),
- /* [231] */ MatcherIndex(7),
+ /* [231] */ MatcherIndex(0),
/* [232] */ MatcherIndex(16),
- /* [233] */ MatcherIndex(8),
- /* [234] */ MatcherIndex(17),
- /* [235] */ MatcherIndex(0),
+ /* [233] */ MatcherIndex(7),
+ /* [234] */ MatcherIndex(16),
+ /* [235] */ MatcherIndex(8),
/* [236] */ MatcherIndex(17),
- /* [237] */ MatcherIndex(7),
+ /* [237] */ MatcherIndex(0),
/* [238] */ MatcherIndex(17),
- /* [239] */ MatcherIndex(8),
- /* [240] */ MatcherIndex(18),
- /* [241] */ MatcherIndex(0),
+ /* [239] */ MatcherIndex(7),
+ /* [240] */ MatcherIndex(17),
+ /* [241] */ MatcherIndex(8),
/* [242] */ MatcherIndex(18),
- /* [243] */ MatcherIndex(7),
+ /* [243] */ MatcherIndex(0),
/* [244] */ MatcherIndex(18),
- /* [245] */ MatcherIndex(8),
- /* [246] */ MatcherIndex(19),
- /* [247] */ MatcherIndex(0),
+ /* [245] */ MatcherIndex(7),
+ /* [246] */ MatcherIndex(18),
+ /* [247] */ MatcherIndex(8),
/* [248] */ MatcherIndex(19),
- /* [249] */ MatcherIndex(7),
+ /* [249] */ MatcherIndex(0),
/* [250] */ MatcherIndex(19),
- /* [251] */ MatcherIndex(8),
- /* [252] */ MatcherIndex(20),
- /* [253] */ MatcherIndex(0),
+ /* [251] */ MatcherIndex(7),
+ /* [252] */ MatcherIndex(19),
+ /* [253] */ MatcherIndex(8),
/* [254] */ MatcherIndex(20),
- /* [255] */ MatcherIndex(7),
+ /* [255] */ MatcherIndex(0),
/* [256] */ MatcherIndex(20),
- /* [257] */ MatcherIndex(8),
- /* [258] */ MatcherIndex(56),
- /* [259] */ MatcherIndex(60),
- /* [260] */ MatcherIndex(61),
- /* [261] */ MatcherIndex(50),
- /* [262] */ MatcherIndex(58),
- /* [263] */ MatcherIndex(57),
- /* [264] */ MatcherIndex(35),
- /* [265] */ MatcherIndex(36),
- /* [266] */ MatcherIndex(37),
- /* [267] */ MatcherIndex(38),
- /* [268] */ MatcherIndex(39),
- /* [269] */ MatcherIndex(44),
- /* [270] */ MatcherIndex(26),
- /* [271] */ MatcherIndex(27),
- /* [272] */ MatcherIndex(53),
- /* [273] */ MatcherIndex(54),
- /* [274] */ MatcherIndex(51),
- /* [275] */ MatcherIndex(52),
- /* [276] */ MatcherIndex(55),
+ /* [257] */ MatcherIndex(7),
+ /* [258] */ MatcherIndex(20),
+ /* [259] */ MatcherIndex(8),
+ /* [260] */ MatcherIndex(57),
+ /* [261] */ MatcherIndex(61),
+ /* [262] */ MatcherIndex(62),
+ /* [263] */ MatcherIndex(51),
+ /* [264] */ MatcherIndex(59),
+ /* [265] */ MatcherIndex(58),
+ /* [266] */ MatcherIndex(35),
+ /* [267] */ MatcherIndex(36),
+ /* [268] */ MatcherIndex(37),
+ /* [269] */ MatcherIndex(38),
+ /* [270] */ MatcherIndex(39),
+ /* [271] */ MatcherIndex(44),
+ /* [272] */ MatcherIndex(26),
+ /* [273] */ MatcherIndex(27),
+ /* [274] */ MatcherIndex(54),
+ /* [275] */ MatcherIndex(55),
+ /* [276] */ MatcherIndex(52),
+ /* [277] */ MatcherIndex(53),
+ /* [278] */ MatcherIndex(56),
};
static_assert(MatcherIndicesIndex::CanIndex(kMatcherIndices),
@@ -1886,7 +1909,7 @@
{
/* [20] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [21] */
@@ -1926,7 +1949,7 @@
{
/* [28] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [29] */
@@ -1946,12 +1969,12 @@
{
/* [32] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(265),
+ /* matcher_indices */ MatcherIndicesIndex(267),
},
{
/* [33] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(271),
+ /* matcher_indices */ MatcherIndicesIndex(273),
},
{
/* [34] */
@@ -1981,7 +2004,7 @@
{
/* [39] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [40] */
@@ -2011,7 +2034,7 @@
{
/* [45] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [46] */
@@ -2041,7 +2064,7 @@
{
/* [51] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [52] */
@@ -2071,7 +2094,7 @@
{
/* [57] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [58] */
@@ -2101,7 +2124,7 @@
{
/* [63] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [64] */
@@ -2126,12 +2149,12 @@
{
/* [68] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(265),
+ /* matcher_indices */ MatcherIndicesIndex(267),
},
{
/* [69] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [70] */
@@ -2166,7 +2189,7 @@
{
/* [76] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [77] */
@@ -2191,7 +2214,7 @@
{
/* [81] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [82] */
@@ -2206,12 +2229,12 @@
{
/* [84] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(265),
+ /* matcher_indices */ MatcherIndicesIndex(267),
},
{
/* [85] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [86] */
@@ -2231,12 +2254,12 @@
{
/* [89] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(264),
+ /* matcher_indices */ MatcherIndicesIndex(266),
},
{
/* [90] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(271),
+ /* matcher_indices */ MatcherIndicesIndex(273),
},
{
/* [91] */
@@ -2256,12 +2279,12 @@
{
/* [94] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(267),
+ /* matcher_indices */ MatcherIndicesIndex(269),
},
{
/* [95] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(271),
+ /* matcher_indices */ MatcherIndicesIndex(273),
},
{
/* [96] */
@@ -2286,7 +2309,7 @@
{
/* [100] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [101] */
@@ -2311,7 +2334,7 @@
{
/* [105] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [106] */
@@ -2336,7 +2359,7 @@
{
/* [110] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [111] */
@@ -2361,7 +2384,7 @@
{
/* [115] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [116] */
@@ -2386,7 +2409,7 @@
{
/* [120] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [121] */
@@ -2411,7 +2434,7 @@
{
/* [125] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [126] */
@@ -2436,7 +2459,7 @@
{
/* [130] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [131] */
@@ -2461,7 +2484,7 @@
{
/* [135] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [136] */
@@ -2481,12 +2504,12 @@
{
/* [139] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(264),
+ /* matcher_indices */ MatcherIndicesIndex(266),
},
{
/* [140] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [141] */
@@ -2506,12 +2529,12 @@
{
/* [144] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(267),
+ /* matcher_indices */ MatcherIndicesIndex(269),
},
{
/* [145] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [146] */
@@ -2561,7 +2584,7 @@
{
/* [155] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [156] */
@@ -2571,12 +2594,12 @@
{
/* [157] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(264),
+ /* matcher_indices */ MatcherIndicesIndex(266),
},
{
/* [158] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [159] */
@@ -2591,12 +2614,12 @@
{
/* [161] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(266),
+ /* matcher_indices */ MatcherIndicesIndex(268),
},
{
/* [162] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(271),
+ /* matcher_indices */ MatcherIndicesIndex(273),
},
{
/* [163] */
@@ -2616,7 +2639,7 @@
{
/* [166] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [167] */
@@ -2636,7 +2659,7 @@
{
/* [170] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [171] */
@@ -2656,7 +2679,7 @@
{
/* [174] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [175] */
@@ -2676,7 +2699,7 @@
{
/* [178] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [179] */
@@ -2691,12 +2714,12 @@
{
/* [181] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(266),
+ /* matcher_indices */ MatcherIndicesIndex(268),
},
{
/* [182] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [183] */
@@ -2791,7 +2814,7 @@
{
/* [201] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(265),
+ /* matcher_indices */ MatcherIndicesIndex(267),
},
{
/* [202] */
@@ -2976,7 +2999,7 @@
{
/* [238] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [239] */
@@ -2986,12 +3009,12 @@
{
/* [240] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(269),
+ /* matcher_indices */ MatcherIndicesIndex(271),
},
{
/* [241] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(270),
+ /* matcher_indices */ MatcherIndicesIndex(272),
},
{
/* [242] */
@@ -3196,7 +3219,7 @@
{
/* [282] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(264),
+ /* matcher_indices */ MatcherIndicesIndex(266),
},
{
/* [283] */
@@ -3211,7 +3234,7 @@
{
/* [285] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(268),
+ /* matcher_indices */ MatcherIndicesIndex(270),
},
{
/* [286] */
@@ -3396,7 +3419,7 @@
{
/* [322] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(264),
+ /* matcher_indices */ MatcherIndicesIndex(266),
},
{
/* [323] */
@@ -3406,7 +3429,7 @@
{
/* [324] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(265),
+ /* matcher_indices */ MatcherIndicesIndex(267),
},
{
/* [325] */
@@ -3416,7 +3439,7 @@
{
/* [326] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(266),
+ /* matcher_indices */ MatcherIndicesIndex(268),
},
{
/* [327] */
@@ -3426,7 +3449,7 @@
{
/* [328] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(267),
+ /* matcher_indices */ MatcherIndicesIndex(269),
},
{
/* [329] */
@@ -3436,7 +3459,7 @@
{
/* [330] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(269),
+ /* matcher_indices */ MatcherIndicesIndex(271),
},
{
/* [331] */
@@ -3715,28 +3738,28 @@
},
{
/* [386] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(86),
+ /* usage */ core::ParameterUsage::kInputAttachment,
+ /* matcher_indices */ MatcherIndicesIndex(200),
},
{
/* [387] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(9),
+ /* matcher_indices */ MatcherIndicesIndex(86),
},
{
/* [388] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(130),
+ /* matcher_indices */ MatcherIndicesIndex(9),
},
{
/* [389] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(89),
+ /* matcher_indices */ MatcherIndicesIndex(130),
},
{
/* [390] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(202),
+ /* matcher_indices */ MatcherIndicesIndex(89),
},
{
/* [391] */
@@ -3873,6 +3896,11 @@
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(256),
},
+ {
+ /* [418] */
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(258),
+ },
};
static_assert(ParameterIndex::CanIndex(kParameters),
@@ -3882,31 +3910,31 @@
{
/* [0] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(263),
+ /* matcher_indices */ MatcherIndicesIndex(265),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [1] */
/* name */ "C",
- /* matcher_indices */ MatcherIndicesIndex(260),
+ /* matcher_indices */ MatcherIndicesIndex(262),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [2] */
/* name */ "A",
- /* matcher_indices */ MatcherIndicesIndex(260),
+ /* matcher_indices */ MatcherIndicesIndex(262),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [3] */
/* name */ "L",
- /* matcher_indices */ MatcherIndicesIndex(260),
+ /* matcher_indices */ MatcherIndicesIndex(262),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [4] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(259),
+ /* matcher_indices */ MatcherIndicesIndex(261),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -3936,7 +3964,7 @@
{
/* [9] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(259),
+ /* matcher_indices */ MatcherIndicesIndex(261),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -3960,7 +3988,7 @@
{
/* [13] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(259),
+ /* matcher_indices */ MatcherIndicesIndex(261),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -3978,43 +4006,43 @@
{
/* [16] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(263),
+ /* matcher_indices */ MatcherIndicesIndex(265),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [17] */
/* name */ "C",
- /* matcher_indices */ MatcherIndicesIndex(260),
+ /* matcher_indices */ MatcherIndicesIndex(262),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [18] */
/* name */ "L",
- /* matcher_indices */ MatcherIndicesIndex(260),
+ /* matcher_indices */ MatcherIndicesIndex(262),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [19] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(263),
+ /* matcher_indices */ MatcherIndicesIndex(265),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [20] */
/* name */ "C",
- /* matcher_indices */ MatcherIndicesIndex(260),
+ /* matcher_indices */ MatcherIndicesIndex(262),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [21] */
/* name */ "S",
- /* matcher_indices */ MatcherIndicesIndex(260),
+ /* matcher_indices */ MatcherIndicesIndex(262),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [22] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(259),
+ /* matcher_indices */ MatcherIndicesIndex(261),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4038,7 +4066,7 @@
{
/* [26] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(258),
+ /* matcher_indices */ MatcherIndicesIndex(260),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4050,13 +4078,13 @@
{
/* [28] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(260),
+ /* matcher_indices */ MatcherIndicesIndex(262),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [29] */
/* name */ "S",
- /* matcher_indices */ MatcherIndicesIndex(204),
+ /* matcher_indices */ MatcherIndicesIndex(206),
/* kind */ TemplateInfo::Kind::kNumber,
},
{
@@ -4074,7 +4102,7 @@
{
/* [32] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(261),
+ /* matcher_indices */ MatcherIndicesIndex(263),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4086,7 +4114,7 @@
{
/* [34] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(261),
+ /* matcher_indices */ MatcherIndicesIndex(263),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4098,7 +4126,7 @@
{
/* [36] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(262),
+ /* matcher_indices */ MatcherIndicesIndex(264),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4110,13 +4138,13 @@
{
/* [38] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(263),
+ /* matcher_indices */ MatcherIndicesIndex(265),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [39] */
/* name */ "L",
- /* matcher_indices */ MatcherIndicesIndex(260),
+ /* matcher_indices */ MatcherIndicesIndex(262),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4134,7 +4162,7 @@
{
/* [42] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(260),
+ /* matcher_indices */ MatcherIndicesIndex(262),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4152,7 +4180,7 @@
{
/* [45] */
/* name */ "U",
- /* matcher_indices */ MatcherIndicesIndex(274),
+ /* matcher_indices */ MatcherIndicesIndex(276),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4164,7 +4192,7 @@
{
/* [47] */
/* name */ "U",
- /* matcher_indices */ MatcherIndicesIndex(275),
+ /* matcher_indices */ MatcherIndicesIndex(277),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4176,7 +4204,7 @@
{
/* [49] */
/* name */ "U",
- /* matcher_indices */ MatcherIndicesIndex(272),
+ /* matcher_indices */ MatcherIndicesIndex(274),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4188,7 +4216,7 @@
{
/* [51] */
/* name */ "U",
- /* matcher_indices */ MatcherIndicesIndex(273),
+ /* matcher_indices */ MatcherIndicesIndex(275),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4200,37 +4228,37 @@
{
/* [53] */
/* name */ "U",
- /* matcher_indices */ MatcherIndicesIndex(276),
+ /* matcher_indices */ MatcherIndicesIndex(278),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [54] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(272),
+ /* matcher_indices */ MatcherIndicesIndex(274),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [55] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(273),
+ /* matcher_indices */ MatcherIndicesIndex(275),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [56] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(274),
+ /* matcher_indices */ MatcherIndicesIndex(276),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [57] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(275),
+ /* matcher_indices */ MatcherIndicesIndex(277),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [58] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(276),
+ /* matcher_indices */ MatcherIndicesIndex(278),
/* kind */ TemplateInfo::Kind::kType,
},
};
@@ -5168,7 +5196,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 2,
/* templates */ TemplateIndex(44),
- /* parameters */ ParameterIndex(390),
+ /* parameters */ ParameterIndex(391),
/* return_matcher_indices */ MatcherIndicesIndex(186),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
@@ -5179,7 +5207,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 2,
/* templates */ TemplateIndex(46),
- /* parameters */ ParameterIndex(390),
+ /* parameters */ ParameterIndex(391),
/* return_matcher_indices */ MatcherIndicesIndex(186),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
@@ -5190,7 +5218,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 2,
/* templates */ TemplateIndex(48),
- /* parameters */ ParameterIndex(390),
+ /* parameters */ ParameterIndex(391),
/* return_matcher_indices */ MatcherIndicesIndex(186),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
@@ -5201,7 +5229,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 2,
/* templates */ TemplateIndex(50),
- /* parameters */ ParameterIndex(390),
+ /* parameters */ ParameterIndex(391),
/* return_matcher_indices */ MatcherIndicesIndex(186),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
@@ -5212,7 +5240,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 2,
/* templates */ TemplateIndex(52),
- /* parameters */ ParameterIndex(390),
+ /* parameters */ ParameterIndex(391),
/* return_matcher_indices */ MatcherIndicesIndex(186),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
@@ -5707,7 +5735,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 2,
/* templates */ TemplateIndex(44),
- /* parameters */ ParameterIndex(389),
+ /* parameters */ ParameterIndex(390),
/* return_matcher_indices */ MatcherIndicesIndex(73),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
@@ -5718,7 +5746,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 2,
/* templates */ TemplateIndex(46),
- /* parameters */ ParameterIndex(389),
+ /* parameters */ ParameterIndex(390),
/* return_matcher_indices */ MatcherIndicesIndex(73),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
@@ -5729,7 +5757,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 2,
/* templates */ TemplateIndex(48),
- /* parameters */ ParameterIndex(389),
+ /* parameters */ ParameterIndex(390),
/* return_matcher_indices */ MatcherIndicesIndex(73),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
@@ -5740,7 +5768,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 2,
/* templates */ TemplateIndex(50),
- /* parameters */ ParameterIndex(389),
+ /* parameters */ ParameterIndex(390),
/* return_matcher_indices */ MatcherIndicesIndex(73),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
@@ -5751,7 +5779,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 2,
/* templates */ TemplateIndex(52),
- /* parameters */ ParameterIndex(389),
+ /* parameters */ ParameterIndex(390),
/* return_matcher_indices */ MatcherIndicesIndex(73),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
@@ -6015,7 +6043,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 2,
/* templates */ TemplateIndex(44),
- /* parameters */ ParameterIndex(388),
+ /* parameters */ ParameterIndex(389),
/* return_matcher_indices */ MatcherIndicesIndex(126),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
@@ -6026,7 +6054,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 2,
/* templates */ TemplateIndex(46),
- /* parameters */ ParameterIndex(388),
+ /* parameters */ ParameterIndex(389),
/* return_matcher_indices */ MatcherIndicesIndex(126),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
@@ -6037,7 +6065,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 2,
/* templates */ TemplateIndex(48),
- /* parameters */ ParameterIndex(388),
+ /* parameters */ ParameterIndex(389),
/* return_matcher_indices */ MatcherIndicesIndex(126),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
@@ -6048,7 +6076,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 2,
/* templates */ TemplateIndex(50),
- /* parameters */ ParameterIndex(388),
+ /* parameters */ ParameterIndex(389),
/* return_matcher_indices */ MatcherIndicesIndex(126),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
@@ -6059,7 +6087,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 2,
/* templates */ TemplateIndex(52),
- /* parameters */ ParameterIndex(388),
+ /* parameters */ ParameterIndex(389),
/* return_matcher_indices */ MatcherIndicesIndex(126),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
@@ -6445,7 +6473,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(/* invalid */),
- /* return_matcher_indices */ MatcherIndicesIndex(204),
+ /* return_matcher_indices */ MatcherIndicesIndex(206),
/* const_eval_fn */ ConstEvalFunctionIndex(105),
},
{
@@ -6455,8 +6483,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
- /* parameters */ ParameterIndex(391),
- /* return_matcher_indices */ MatcherIndicesIndex(204),
+ /* parameters */ ParameterIndex(392),
+ /* return_matcher_indices */ MatcherIndicesIndex(206),
/* const_eval_fn */ ConstEvalFunctionIndex(106),
},
{
@@ -6467,7 +6495,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(1),
- /* return_matcher_indices */ MatcherIndicesIndex(204),
+ /* return_matcher_indices */ MatcherIndicesIndex(206),
/* const_eval_fn */ ConstEvalFunctionIndex(111),
},
{
@@ -6478,7 +6506,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(209),
- /* return_matcher_indices */ MatcherIndicesIndex(204),
+ /* return_matcher_indices */ MatcherIndicesIndex(206),
/* const_eval_fn */ ConstEvalFunctionIndex(112),
},
{
@@ -6488,8 +6516,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(46),
- /* parameters */ ParameterIndex(392),
- /* return_matcher_indices */ MatcherIndicesIndex(204),
+ /* parameters */ ParameterIndex(393),
+ /* return_matcher_indices */ MatcherIndicesIndex(206),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -6499,8 +6527,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(44),
- /* parameters */ ParameterIndex(393),
- /* return_matcher_indices */ MatcherIndicesIndex(204),
+ /* parameters */ ParameterIndex(394),
+ /* return_matcher_indices */ MatcherIndicesIndex(206),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -6511,7 +6539,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(/* invalid */),
- /* return_matcher_indices */ MatcherIndicesIndex(210),
+ /* return_matcher_indices */ MatcherIndicesIndex(212),
/* const_eval_fn */ ConstEvalFunctionIndex(105),
},
{
@@ -6521,8 +6549,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
- /* parameters */ ParameterIndex(394),
- /* return_matcher_indices */ MatcherIndicesIndex(210),
+ /* parameters */ ParameterIndex(395),
+ /* return_matcher_indices */ MatcherIndicesIndex(212),
/* const_eval_fn */ ConstEvalFunctionIndex(106),
},
{
@@ -6533,7 +6561,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(1),
- /* return_matcher_indices */ MatcherIndicesIndex(210),
+ /* return_matcher_indices */ MatcherIndicesIndex(212),
/* const_eval_fn */ ConstEvalFunctionIndex(111),
},
{
@@ -6544,7 +6572,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(213),
- /* return_matcher_indices */ MatcherIndicesIndex(210),
+ /* return_matcher_indices */ MatcherIndicesIndex(212),
/* const_eval_fn */ ConstEvalFunctionIndex(112),
},
{
@@ -6554,8 +6582,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(46),
- /* parameters */ ParameterIndex(395),
- /* return_matcher_indices */ MatcherIndicesIndex(210),
+ /* parameters */ ParameterIndex(396),
+ /* return_matcher_indices */ MatcherIndicesIndex(212),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -6565,8 +6593,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(44),
- /* parameters */ ParameterIndex(396),
- /* return_matcher_indices */ MatcherIndicesIndex(210),
+ /* parameters */ ParameterIndex(397),
+ /* return_matcher_indices */ MatcherIndicesIndex(212),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -6577,7 +6605,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(/* invalid */),
- /* return_matcher_indices */ MatcherIndicesIndex(216),
+ /* return_matcher_indices */ MatcherIndicesIndex(218),
/* const_eval_fn */ ConstEvalFunctionIndex(105),
},
{
@@ -6587,8 +6615,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
- /* parameters */ ParameterIndex(397),
- /* return_matcher_indices */ MatcherIndicesIndex(216),
+ /* parameters */ ParameterIndex(398),
+ /* return_matcher_indices */ MatcherIndicesIndex(218),
/* const_eval_fn */ ConstEvalFunctionIndex(106),
},
{
@@ -6599,7 +6627,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(1),
- /* return_matcher_indices */ MatcherIndicesIndex(216),
+ /* return_matcher_indices */ MatcherIndicesIndex(218),
/* const_eval_fn */ ConstEvalFunctionIndex(111),
},
{
@@ -6610,7 +6638,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(217),
- /* return_matcher_indices */ MatcherIndicesIndex(216),
+ /* return_matcher_indices */ MatcherIndicesIndex(218),
/* const_eval_fn */ ConstEvalFunctionIndex(112),
},
{
@@ -6620,8 +6648,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(46),
- /* parameters */ ParameterIndex(398),
- /* return_matcher_indices */ MatcherIndicesIndex(216),
+ /* parameters */ ParameterIndex(399),
+ /* return_matcher_indices */ MatcherIndicesIndex(218),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -6631,8 +6659,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(44),
- /* parameters */ ParameterIndex(399),
- /* return_matcher_indices */ MatcherIndicesIndex(216),
+ /* parameters */ ParameterIndex(400),
+ /* return_matcher_indices */ MatcherIndicesIndex(218),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -6643,7 +6671,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(/* invalid */),
- /* return_matcher_indices */ MatcherIndicesIndex(222),
+ /* return_matcher_indices */ MatcherIndicesIndex(224),
/* const_eval_fn */ ConstEvalFunctionIndex(105),
},
{
@@ -6653,8 +6681,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
- /* parameters */ ParameterIndex(400),
- /* return_matcher_indices */ MatcherIndicesIndex(222),
+ /* parameters */ ParameterIndex(401),
+ /* return_matcher_indices */ MatcherIndicesIndex(224),
/* const_eval_fn */ ConstEvalFunctionIndex(106),
},
{
@@ -6665,7 +6693,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(1),
- /* return_matcher_indices */ MatcherIndicesIndex(222),
+ /* return_matcher_indices */ MatcherIndicesIndex(224),
/* const_eval_fn */ ConstEvalFunctionIndex(111),
},
{
@@ -6676,7 +6704,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(209),
- /* return_matcher_indices */ MatcherIndicesIndex(222),
+ /* return_matcher_indices */ MatcherIndicesIndex(224),
/* const_eval_fn */ ConstEvalFunctionIndex(112),
},
{
@@ -6686,8 +6714,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(46),
- /* parameters */ ParameterIndex(401),
- /* return_matcher_indices */ MatcherIndicesIndex(222),
+ /* parameters */ ParameterIndex(402),
+ /* return_matcher_indices */ MatcherIndicesIndex(224),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -6697,8 +6725,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(44),
- /* parameters */ ParameterIndex(402),
- /* return_matcher_indices */ MatcherIndicesIndex(222),
+ /* parameters */ ParameterIndex(403),
+ /* return_matcher_indices */ MatcherIndicesIndex(224),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -6709,7 +6737,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(/* invalid */),
- /* return_matcher_indices */ MatcherIndicesIndex(228),
+ /* return_matcher_indices */ MatcherIndicesIndex(230),
/* const_eval_fn */ ConstEvalFunctionIndex(105),
},
{
@@ -6719,8 +6747,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
- /* parameters */ ParameterIndex(403),
- /* return_matcher_indices */ MatcherIndicesIndex(228),
+ /* parameters */ ParameterIndex(404),
+ /* return_matcher_indices */ MatcherIndicesIndex(230),
/* const_eval_fn */ ConstEvalFunctionIndex(106),
},
{
@@ -6731,7 +6759,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(1),
- /* return_matcher_indices */ MatcherIndicesIndex(228),
+ /* return_matcher_indices */ MatcherIndicesIndex(230),
/* const_eval_fn */ ConstEvalFunctionIndex(111),
},
{
@@ -6742,7 +6770,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(213),
- /* return_matcher_indices */ MatcherIndicesIndex(228),
+ /* return_matcher_indices */ MatcherIndicesIndex(230),
/* const_eval_fn */ ConstEvalFunctionIndex(112),
},
{
@@ -6752,8 +6780,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(46),
- /* parameters */ ParameterIndex(404),
- /* return_matcher_indices */ MatcherIndicesIndex(228),
+ /* parameters */ ParameterIndex(405),
+ /* return_matcher_indices */ MatcherIndicesIndex(230),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -6763,8 +6791,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(44),
- /* parameters */ ParameterIndex(405),
- /* return_matcher_indices */ MatcherIndicesIndex(228),
+ /* parameters */ ParameterIndex(406),
+ /* return_matcher_indices */ MatcherIndicesIndex(230),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -6775,7 +6803,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(/* invalid */),
- /* return_matcher_indices */ MatcherIndicesIndex(234),
+ /* return_matcher_indices */ MatcherIndicesIndex(236),
/* const_eval_fn */ ConstEvalFunctionIndex(105),
},
{
@@ -6785,8 +6813,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
- /* parameters */ ParameterIndex(406),
- /* return_matcher_indices */ MatcherIndicesIndex(234),
+ /* parameters */ ParameterIndex(407),
+ /* return_matcher_indices */ MatcherIndicesIndex(236),
/* const_eval_fn */ ConstEvalFunctionIndex(106),
},
{
@@ -6797,7 +6825,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(1),
- /* return_matcher_indices */ MatcherIndicesIndex(234),
+ /* return_matcher_indices */ MatcherIndicesIndex(236),
/* const_eval_fn */ ConstEvalFunctionIndex(111),
},
{
@@ -6808,7 +6836,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(217),
- /* return_matcher_indices */ MatcherIndicesIndex(234),
+ /* return_matcher_indices */ MatcherIndicesIndex(236),
/* const_eval_fn */ ConstEvalFunctionIndex(112),
},
{
@@ -6818,8 +6846,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(46),
- /* parameters */ ParameterIndex(407),
- /* return_matcher_indices */ MatcherIndicesIndex(234),
+ /* parameters */ ParameterIndex(408),
+ /* return_matcher_indices */ MatcherIndicesIndex(236),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -6829,8 +6857,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(44),
- /* parameters */ ParameterIndex(408),
- /* return_matcher_indices */ MatcherIndicesIndex(234),
+ /* parameters */ ParameterIndex(409),
+ /* return_matcher_indices */ MatcherIndicesIndex(236),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -6841,7 +6869,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(/* invalid */),
- /* return_matcher_indices */ MatcherIndicesIndex(240),
+ /* return_matcher_indices */ MatcherIndicesIndex(242),
/* const_eval_fn */ ConstEvalFunctionIndex(105),
},
{
@@ -6851,8 +6879,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
- /* parameters */ ParameterIndex(409),
- /* return_matcher_indices */ MatcherIndicesIndex(240),
+ /* parameters */ ParameterIndex(410),
+ /* return_matcher_indices */ MatcherIndicesIndex(242),
/* const_eval_fn */ ConstEvalFunctionIndex(106),
},
{
@@ -6863,7 +6891,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(1),
- /* return_matcher_indices */ MatcherIndicesIndex(240),
+ /* return_matcher_indices */ MatcherIndicesIndex(242),
/* const_eval_fn */ ConstEvalFunctionIndex(111),
},
{
@@ -6874,7 +6902,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(209),
- /* return_matcher_indices */ MatcherIndicesIndex(240),
+ /* return_matcher_indices */ MatcherIndicesIndex(242),
/* const_eval_fn */ ConstEvalFunctionIndex(112),
},
{
@@ -6884,8 +6912,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(46),
- /* parameters */ ParameterIndex(410),
- /* return_matcher_indices */ MatcherIndicesIndex(240),
+ /* parameters */ ParameterIndex(411),
+ /* return_matcher_indices */ MatcherIndicesIndex(242),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -6895,8 +6923,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(44),
- /* parameters */ ParameterIndex(411),
- /* return_matcher_indices */ MatcherIndicesIndex(240),
+ /* parameters */ ParameterIndex(412),
+ /* return_matcher_indices */ MatcherIndicesIndex(242),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -6907,7 +6935,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(/* invalid */),
- /* return_matcher_indices */ MatcherIndicesIndex(246),
+ /* return_matcher_indices */ MatcherIndicesIndex(248),
/* const_eval_fn */ ConstEvalFunctionIndex(105),
},
{
@@ -6917,8 +6945,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
- /* parameters */ ParameterIndex(412),
- /* return_matcher_indices */ MatcherIndicesIndex(246),
+ /* parameters */ ParameterIndex(413),
+ /* return_matcher_indices */ MatcherIndicesIndex(248),
/* const_eval_fn */ ConstEvalFunctionIndex(106),
},
{
@@ -6929,7 +6957,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(1),
- /* return_matcher_indices */ MatcherIndicesIndex(246),
+ /* return_matcher_indices */ MatcherIndicesIndex(248),
/* const_eval_fn */ ConstEvalFunctionIndex(111),
},
{
@@ -6940,7 +6968,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(213),
- /* return_matcher_indices */ MatcherIndicesIndex(246),
+ /* return_matcher_indices */ MatcherIndicesIndex(248),
/* const_eval_fn */ ConstEvalFunctionIndex(112),
},
{
@@ -6950,8 +6978,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(46),
- /* parameters */ ParameterIndex(413),
- /* return_matcher_indices */ MatcherIndicesIndex(246),
+ /* parameters */ ParameterIndex(414),
+ /* return_matcher_indices */ MatcherIndicesIndex(248),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -6961,8 +6989,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(44),
- /* parameters */ ParameterIndex(414),
- /* return_matcher_indices */ MatcherIndicesIndex(246),
+ /* parameters */ ParameterIndex(415),
+ /* return_matcher_indices */ MatcherIndicesIndex(248),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -6973,7 +7001,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(/* invalid */),
- /* return_matcher_indices */ MatcherIndicesIndex(252),
+ /* return_matcher_indices */ MatcherIndicesIndex(254),
/* const_eval_fn */ ConstEvalFunctionIndex(105),
},
{
@@ -6983,8 +7011,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
- /* parameters */ ParameterIndex(415),
- /* return_matcher_indices */ MatcherIndicesIndex(252),
+ /* parameters */ ParameterIndex(416),
+ /* return_matcher_indices */ MatcherIndicesIndex(254),
/* const_eval_fn */ ConstEvalFunctionIndex(106),
},
{
@@ -6995,7 +7023,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(1),
- /* return_matcher_indices */ MatcherIndicesIndex(252),
+ /* return_matcher_indices */ MatcherIndicesIndex(254),
/* const_eval_fn */ ConstEvalFunctionIndex(111),
},
{
@@ -7006,7 +7034,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(217),
- /* return_matcher_indices */ MatcherIndicesIndex(252),
+ /* return_matcher_indices */ MatcherIndicesIndex(254),
/* const_eval_fn */ ConstEvalFunctionIndex(112),
},
{
@@ -7016,8 +7044,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(46),
- /* parameters */ ParameterIndex(416),
- /* return_matcher_indices */ MatcherIndicesIndex(252),
+ /* parameters */ ParameterIndex(417),
+ /* return_matcher_indices */ MatcherIndicesIndex(254),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -7027,8 +7055,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(44),
- /* parameters */ ParameterIndex(417),
- /* return_matcher_indices */ MatcherIndicesIndex(252),
+ /* parameters */ ParameterIndex(418),
+ /* return_matcher_indices */ MatcherIndicesIndex(254),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -7456,7 +7484,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(386),
+ /* parameters */ ParameterIndex(387),
/* return_matcher_indices */ MatcherIndicesIndex(86),
/* const_eval_fn */ ConstEvalFunctionIndex(106),
},
@@ -7555,7 +7583,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(387),
+ /* parameters */ ParameterIndex(388),
/* return_matcher_indices */ MatcherIndicesIndex(9),
/* const_eval_fn */ ConstEvalFunctionIndex(106),
},
@@ -9409,6 +9437,17 @@
},
{
/* [459] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline),
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 1,
+ /* templates */ TemplateIndex(0),
+ /* parameters */ ParameterIndex(386),
+ /* return_matcher_indices */ MatcherIndicesIndex(186),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [460] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -9419,7 +9458,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [460] */
+ /* [461] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -9430,7 +9469,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [461] */
+ /* [462] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -9441,18 +9480,18 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [462] */
+ /* [463] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(28),
/* parameters */ ParameterIndex(0),
- /* return_matcher_indices */ MatcherIndicesIndex(200),
+ /* return_matcher_indices */ MatcherIndicesIndex(202),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [463] */
+ /* [464] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMustUse),
/* num_parameters */ 0,
/* num_explicit_templates */ 0,
@@ -9463,7 +9502,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [464] */
+ /* [465] */
/* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMustUse),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -9474,7 +9513,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(95),
},
{
- /* [465] */
+ /* [466] */
/* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMustUse),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -10344,78 +10383,84 @@
},
{
/* [108] */
- /* fn atomicLoad[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>) -> T */
+ /* fn inputAttachmentLoad[T : fiu32](input_attachment: input_attachment<T>) -> vec4<T> */
/* num overloads */ 1,
/* overloads */ OverloadIndex(459),
},
{
/* [109] */
- /* fn atomicStore[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) */
+ /* fn atomicLoad[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>) -> T */
/* num overloads */ 1,
/* overloads */ OverloadIndex(460),
},
{
/* [110] */
- /* fn atomicAdd[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
+ /* fn atomicStore[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) */
/* num overloads */ 1,
/* overloads */ OverloadIndex(461),
},
{
/* [111] */
- /* fn atomicSub[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
+ /* fn atomicAdd[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(461),
+ /* overloads */ OverloadIndex(462),
},
{
/* [112] */
- /* fn atomicMax[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
+ /* fn atomicSub[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(461),
+ /* overloads */ OverloadIndex(462),
},
{
/* [113] */
- /* fn atomicMin[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
+ /* fn atomicMax[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(461),
+ /* overloads */ OverloadIndex(462),
},
{
/* [114] */
- /* fn atomicAnd[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
+ /* fn atomicMin[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(461),
+ /* overloads */ OverloadIndex(462),
},
{
/* [115] */
- /* fn atomicOr[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
+ /* fn atomicAnd[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(461),
+ /* overloads */ OverloadIndex(462),
},
{
/* [116] */
- /* fn atomicXor[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
+ /* fn atomicOr[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(461),
+ /* overloads */ OverloadIndex(462),
},
{
/* [117] */
- /* fn atomicExchange[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
+ /* fn atomicXor[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(461),
+ /* overloads */ OverloadIndex(462),
},
{
/* [118] */
- /* fn atomicCompareExchangeWeak[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, T) -> __atomic_compare_exchange_result<T> */
+ /* fn atomicExchange[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
/* num overloads */ 1,
/* overloads */ OverloadIndex(462),
},
{
/* [119] */
- /* fn subgroupBallot() -> vec4<u32> */
+ /* fn atomicCompareExchangeWeak[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, T) -> __atomic_compare_exchange_result<T> */
/* num overloads */ 1,
/* overloads */ OverloadIndex(463),
},
{
/* [120] */
+ /* fn subgroupBallot() -> vec4<u32> */
+ /* num overloads */ 1,
+ /* overloads */ OverloadIndex(464),
+ },
+ {
+ /* [121] */
/* fn subgroupBroadcast[T : fiu32_f16](value: T, @const sourceLaneIndex: u32) -> T */
/* fn subgroupBroadcast[N : num, T : fiu32_f16](value: vec<N, T>, @const sourceLaneIndex: u32) -> vec<N, T> */
/* num overloads */ 2,
@@ -10532,13 +10577,13 @@
/* [8] */
/* op &&(bool, bool) -> bool */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(464),
+ /* overloads */ OverloadIndex(465),
},
{
/* [9] */
/* op ||(bool, bool) -> bool */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(465),
+ /* overloads */ OverloadIndex(466),
},
{
/* [10] */
diff --git a/src/tint/lang/core/ir/binary/decode.cc b/src/tint/lang/core/ir/binary/decode.cc
index ada52f9..770f89b 100644
--- a/src/tint/lang/core/ir/binary/decode.cc
+++ b/src/tint/lang/core/ir/binary/decode.cc
@@ -38,6 +38,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/invalid.h"
#include "src/tint/lang/core/type/multisampled_texture.h"
#include "src/tint/lang/core/type/sampled_texture.h"
@@ -624,6 +625,9 @@
auto& bp_in = var_in.binding_point();
var_out->SetBindingPoint(bp_in.group(), bp_in.binding());
}
+ if (var_in.has_input_attachment_index()) {
+ var_out->SetInputAttachmentIndex(var_in.input_attachment_index());
+ }
return var_out;
}
@@ -664,6 +668,8 @@
return CreateTypeExternalTexture(type_in.external_texture());
case pb::Type::KindCase::kSampler:
return CreateTypeSampler(type_in.sampler());
+ case pb::Type::KindCase::kInputAttachment:
+ return CreateTypeInputAttachment(type_in.input_attachment());
case pb::Type::KindCase::KIND_NOT_SET:
break;
}
@@ -863,6 +869,12 @@
return mod_out_.Types().Get<type::Sampler>(kind);
}
+ const type::InputAttachment* CreateTypeInputAttachment(
+ const pb::TypeInputAttachment& input_in) {
+ auto sub_type = Type(input_in.sub_type());
+ return mod_out_.Types().Get<type::InputAttachment>(sub_type);
+ }
+
const type::Type* Type(size_t id) {
if (TINT_UNLIKELY(id >= types_.Length())) {
Error() << "type id " << id << " out of range";
@@ -1595,6 +1607,8 @@
return core::BuiltinFn::kSubgroupBallot;
case pb::BuiltinFn::subgroup_broadcast:
return core::BuiltinFn::kSubgroupBroadcast;
+ case pb::BuiltinFn::input_attachment_load:
+ return core::BuiltinFn::kInputAttachmentLoad;
case pb::BuiltinFn::BuiltinFn_INT_MIN_SENTINEL_DO_NOT_USE_:
case pb::BuiltinFn::BuiltinFn_INT_MAX_SENTINEL_DO_NOT_USE_:
diff --git a/src/tint/lang/core/ir/binary/encode.cc b/src/tint/lang/core/ir/binary/encode.cc
index 05f06bd..bbd6290 100644
--- a/src/tint/lang/core/ir/binary/encode.cc
+++ b/src/tint/lang/core/ir/binary/encode.cc
@@ -74,6 +74,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"
@@ -343,6 +344,9 @@
auto& bp_out = *var_out.mutable_binding_point();
BindingPoint(bp_out, *bp_in);
}
+ if (auto iidx_in = var_in->InputAttachmentIndex()) {
+ var_out.set_input_attachment_index(iidx_in.value());
+ }
}
void InstructionUnreachable(pb::InstructionUnreachable&, const ir::Unreachable*) {}
@@ -387,6 +391,9 @@
TypeExternalTexture(*type_out.mutable_external_texture(), t);
},
[&](const core::type::Sampler* s) { TypeSampler(*type_out.mutable_sampler(), s); },
+ [&](const core::type::InputAttachment* i) {
+ TypeInputAttachment(*type_out.mutable_input_attachment(), i);
+ },
TINT_ICE_ON_NO_MATCH);
mod_out_.mutable_types()->Add(std::move(type_out));
@@ -487,6 +494,10 @@
}
void TypeExternalTexture(pb::TypeExternalTexture&, const core::type::ExternalTexture*) {}
+ void TypeInputAttachment(pb::TypeInputAttachment& input_attachment_out,
+ const core::type::InputAttachment* input_attachment_in) {
+ input_attachment_out.set_sub_type(Type(input_attachment_in->type()));
+ }
void TypeSampler(pb::TypeSampler& sampler_out, const core::type::Sampler* sampler_in) {
sampler_out.set_kind(SamplerKind(sampler_in->kind()));
@@ -1123,6 +1134,8 @@
return pb::BuiltinFn::subgroup_ballot;
case core::BuiltinFn::kSubgroupBroadcast:
return pb::BuiltinFn::subgroup_broadcast;
+ case core::BuiltinFn::kInputAttachmentLoad:
+ return pb::BuiltinFn::input_attachment_load;
case core::BuiltinFn::kNone:
break;
}
diff --git a/src/tint/lang/core/ir/binary/ir.proto b/src/tint/lang/core/ir/binary/ir.proto
index a5271c3..98ef951 100644
--- a/src/tint/lang/core/ir/binary/ir.proto
+++ b/src/tint/lang/core/ir/binary/ir.proto
@@ -57,6 +57,7 @@
TypeStorageTexture storage_texture = 12;
TypeExternalTexture external_texture = 13;
TypeSampler sampler = 14;
+ TypeInputAttachment input_attachment = 15;
}
}
@@ -140,6 +141,10 @@
SamplerKind kind = 1;
}
+message TypeInputAttachment {
+ uint32 sub_type = 1; // Module.types
+}
+
////////////////////////////////////////////////////////////////////////////////
// Values
////////////////////////////////////////////////////////////////////////////////
@@ -294,6 +299,7 @@
message InstructionVar {
optional BindingPoint binding_point = 1;
+ optional uint32 input_attachment_index = 2;
}
message InstructionConvert {}
@@ -625,4 +631,5 @@
atomic_compare_exchange_weak = 118;
subgroup_ballot = 119;
subgroup_broadcast = 120;
+ input_attachment_load = 121;
}
diff --git a/src/tint/lang/core/ir/binary/roundtrip_test.cc b/src/tint/lang/core/ir/binary/roundtrip_test.cc
index 0dee1d2..3cfe333 100644
--- a/src/tint/lang/core/ir/binary/roundtrip_test.cc
+++ b/src/tint/lang/core/ir/binary/roundtrip_test.cc
@@ -33,6 +33,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/multisampled_texture.h"
#include "src/tint/lang/core/type/sampled_texture.h"
#include "src/tint/lang/core/type/storage_texture.h"
@@ -732,5 +733,19 @@
RUN_TEST();
}
+TEST_F(IRBinaryRoundtripTest, InputAttachment) {
+ b.Append(b.ir.root_block, [&] {
+ auto* input_type = ty.Get<core::type::InputAttachment>(ty.i32());
+ auto* v = b.Var(ty.ptr(handle, input_type, read));
+ v->SetBindingPoint(10, 20);
+ v->SetInputAttachmentIndex(11);
+
+ auto* fn = b.Function("Function", ty.vec4<i32>());
+ b.Append(fn->Block(),
+ [&] { b.Return(fn, b.Call<i32>(core::BuiltinFn::kInputAttachmentLoad, v)); });
+ });
+ RUN_TEST();
+}
+
} // namespace
} // namespace tint::core::ir::binary
diff --git a/src/tint/lang/core/ir/var.h b/src/tint/lang/core/ir/var.h
index b3e06c4..251baeb 100644
--- a/src/tint/lang/core/ir/var.h
+++ b/src/tint/lang/core/ir/var.h
@@ -84,6 +84,12 @@
/// @returns the binding points if `Attributes` contains `kBindingPoint`
std::optional<struct BindingPoint> BindingPoint() const { return binding_point_; }
+ /// Sets the input attachment index
+ /// @param index the index
+ void SetInputAttachmentIndex(uint32_t index) { input_attachment_index_ = index; }
+ /// @returns the input attachment index if any
+ std::optional<uint32_t> InputAttachmentIndex() const { return input_attachment_index_; }
+
/// Sets the IO attributes
/// @param attrs the attributes
void SetAttributes(const IOAttributes& attrs) { attributes_ = attrs; }
@@ -98,6 +104,7 @@
private:
std::optional<struct BindingPoint> binding_point_;
+ std::optional<uint32_t> input_attachment_index_;
IOAttributes attributes_;
};
diff --git a/src/tint/lang/core/parameter_usage.cc b/src/tint/lang/core/parameter_usage.cc
index 3efbde9..431dec2 100644
--- a/src/tint/lang/core/parameter_usage.cc
+++ b/src/tint/lang/core/parameter_usage.cc
@@ -56,6 +56,8 @@
return "ddy";
case ParameterUsage::kDepthRef:
return "depth_ref";
+ case ParameterUsage::kInputAttachment:
+ return "input_attachment";
case ParameterUsage::kLevel:
return "level";
case ParameterUsage::kOffset:
diff --git a/src/tint/lang/core/parameter_usage.h b/src/tint/lang/core/parameter_usage.h
index aa1913b..10eecf6 100644
--- a/src/tint/lang/core/parameter_usage.h
+++ b/src/tint/lang/core/parameter_usage.h
@@ -54,6 +54,7 @@
kDdx,
kDdy,
kDepthRef,
+ kInputAttachment,
kLevel,
kOffset,
kSampleIndex,
diff --git a/src/tint/lang/spirv/writer/printer/printer.cc b/src/tint/lang/spirv/writer/printer/printer.cc
index e1a00e0..0695542 100644
--- a/src/tint/lang/spirv/writer/printer/printer.cc
+++ b/src/tint/lang/spirv/writer/printer/printer.cc
@@ -81,6 +81,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"
@@ -611,7 +612,8 @@
texture, //
[&](const core::type::SampledTexture* t) { return Type(t->type()); },
[&](const core::type::MultisampledTexture* t) { return Type(t->type()); },
- [&](const core::type::StorageTexture* t) { return Type(t->type()); }, //
+ [&](const core::type::StorageTexture* t) { return Type(t->type()); },
+ [&](const core::type::InputAttachment* t) { return Type(t->type()); }, //
TINT_ICE_ON_NO_MATCH);
uint32_t dim = SpvDimMax;
diff --git a/src/tint/lang/spirv/writer/raise/builtin_polyfill.cc b/src/tint/lang/spirv/writer/raise/builtin_polyfill.cc
index 70808f6..9c7e276 100644
--- a/src/tint/lang/spirv/writer/raise/builtin_polyfill.cc
+++ b/src/tint/lang/spirv/writer/raise/builtin_polyfill.cc
@@ -99,6 +99,7 @@
case core::BuiltinFn::kTextureSampleGrad:
case core::BuiltinFn::kTextureSampleLevel:
case core::BuiltinFn::kTextureStore:
+ case core::BuiltinFn::kInputAttachmentLoad:
worklist.Push(builtin);
break;
case core::BuiltinFn::kQuantizeToF16:
@@ -168,6 +169,9 @@
case core::BuiltinFn::kQuantizeToF16:
QuantizeToF16Vec(builtin);
break;
+ case core::BuiltinFn::kInputAttachmentLoad:
+ InputAttachmentLoad(builtin);
+ break;
default:
break;
}
@@ -853,6 +857,23 @@
construct->InsertBefore(builtin);
builtin->Destroy();
}
+
+ /// Handle an inputAttachmentLoad() builtin.
+ /// @param builtin the builtin call instruction
+ void InputAttachmentLoad(core::ir::CoreBuiltinCall* builtin) {
+ TINT_ASSERT(builtin->Args().Length() == 1);
+ // TODO(b/341117913): implement this
+
+ auto* result_ty = builtin->Result(0)->Type();
+
+ // Create placeholder result
+ auto* result = b.Construct(result_ty, b.Zero(result_ty));
+
+ result->InsertBefore(builtin);
+
+ result->SetResults(Vector{builtin->DetachResult()});
+ builtin->Destroy();
+ }
};
} // namespace
diff --git a/src/tint/lang/wgsl/ast/transform/transform.cc b/src/tint/lang/wgsl/ast/transform/transform.cc
index 9d9a22a..33a45c9 100644
--- a/src/tint/lang/wgsl/ast/transform/transform.cc
+++ b/src/tint/lang/wgsl/ast/transform/transform.cc
@@ -34,6 +34,7 @@
#include "src/tint/lang/core/fluent_types.h"
#include "src/tint/lang/core/type/atomic.h"
#include "src/tint/lang/core/type/depth_multisampled_texture.h"
+#include "src/tint/lang/core/type/input_attachment.h"
#include "src/tint/lang/core/type/reference.h"
#include "src/tint/lang/core/type/sampler.h"
#include "src/tint/lang/wgsl/program/clone_context.h"
@@ -186,6 +187,9 @@
address_space == core::AddressSpace::kStorage ? p->Access() : core::Access::kUndefined;
return ctx.dst->ty.ptr(address_space, CreateASTTypeFor(ctx, p->StoreType()), access);
}
+ if (auto* i = ty->As<core::type::InputAttachment>()) {
+ return ctx.dst->ty.input_attachment(CreateASTTypeFor(ctx, i->type()));
+ }
TINT_UNREACHABLE() << "Unhandled type: " << ty->TypeInfo().name;
}
diff --git a/src/tint/lang/wgsl/ast/variable.h b/src/tint/lang/wgsl/ast/variable.h
index a0d4be5..f348494 100644
--- a/src/tint/lang/wgsl/ast/variable.h
+++ b/src/tint/lang/wgsl/ast/variable.h
@@ -37,6 +37,7 @@
#include "src/tint/lang/wgsl/ast/binding_attribute.h"
#include "src/tint/lang/wgsl/ast/expression.h"
#include "src/tint/lang/wgsl/ast/group_attribute.h"
+#include "src/tint/lang/wgsl/ast/input_attachment_index_attribute.h"
#include "src/tint/lang/wgsl/ast/node.h"
#include "src/tint/lang/wgsl/ast/type.h"
@@ -81,6 +82,11 @@
HasAttribute<GroupAttribute>(attributes);
}
+ /// @returns true if the variable has an input_attachment_index attribute
+ bool HasInputAttachmentIndex() const {
+ return HasAttribute<InputAttachmentIndexAttribute>(attributes);
+ }
+
/// @returns the kind of the variable, which can be used in diagnostics
/// e.g. "var", "let", "const", etc
virtual const char* Kind() const = 0;
diff --git a/src/tint/lang/wgsl/builtin_fn.cc b/src/tint/lang/wgsl/builtin_fn.cc
index 78afd1b..286c8d1 100644
--- a/src/tint/lang/wgsl/builtin_fn.cc
+++ b/src/tint/lang/wgsl/builtin_fn.cc
@@ -369,6 +369,9 @@
if (name == "textureLoad") {
return BuiltinFn::kTextureLoad;
}
+ if (name == "inputAttachmentLoad") {
+ return BuiltinFn::kInputAttachmentLoad;
+ }
if (name == "atomicLoad") {
return BuiltinFn::kAtomicLoad;
}
@@ -638,6 +641,8 @@
return "textureStore";
case BuiltinFn::kTextureLoad:
return "textureLoad";
+ case BuiltinFn::kInputAttachmentLoad:
+ return "inputAttachmentLoad";
case BuiltinFn::kAtomicLoad:
return "atomicLoad";
case BuiltinFn::kAtomicStore:
diff --git a/src/tint/lang/wgsl/builtin_fn.h b/src/tint/lang/wgsl/builtin_fn.h
index 8f061ee..679e975 100644
--- a/src/tint/lang/wgsl/builtin_fn.h
+++ b/src/tint/lang/wgsl/builtin_fn.h
@@ -157,6 +157,7 @@
kTextureSampleBaseClampToEdge,
kTextureStore,
kTextureLoad,
+ kInputAttachmentLoad,
kAtomicLoad,
kAtomicStore,
kAtomicAdd,
@@ -303,6 +304,7 @@
BuiltinFn::kTextureSampleBaseClampToEdge,
BuiltinFn::kTextureStore,
BuiltinFn::kTextureLoad,
+ BuiltinFn::kInputAttachmentLoad,
BuiltinFn::kAtomicLoad,
BuiltinFn::kAtomicStore,
BuiltinFn::kAtomicAdd,
@@ -431,6 +433,7 @@
"textureSampleBaseClampToEdge",
"textureStore",
"textureLoad",
+ "inputAttachmentLoad",
"atomicLoad",
"atomicStore",
"atomicAdd",
diff --git a/src/tint/lang/wgsl/intrinsic/data.cc b/src/tint/lang/wgsl/intrinsic/data.cc
index 5c5b46d..73a8ff4 100644
--- a/src/tint/lang/wgsl/intrinsic/data.cc
+++ b/src/tint/lang/wgsl/intrinsic/data.cc
@@ -2035,94 +2035,96 @@
/* [177] */ MatcherIndex(0),
/* [178] */ MatcherIndex(12),
/* [179] */ MatcherIndex(1),
- /* [180] */ MatcherIndex(54),
+ /* [180] */ MatcherIndex(49),
/* [181] */ MatcherIndex(0),
- /* [182] */ MatcherIndex(11),
- /* [183] */ MatcherIndex(5),
- /* [184] */ MatcherIndex(12),
+ /* [182] */ MatcherIndex(54),
+ /* [183] */ MatcherIndex(0),
+ /* [184] */ MatcherIndex(11),
/* [185] */ MatcherIndex(5),
- /* [186] */ MatcherIndex(13),
+ /* [186] */ MatcherIndex(12),
/* [187] */ MatcherIndex(5),
/* [188] */ MatcherIndex(13),
- /* [189] */ MatcherIndex(1),
- /* [190] */ MatcherIndex(14),
- /* [191] */ MatcherIndex(9),
+ /* [189] */ MatcherIndex(5),
+ /* [190] */ MatcherIndex(13),
+ /* [191] */ MatcherIndex(1),
/* [192] */ MatcherIndex(14),
- /* [193] */ MatcherIndex(10),
- /* [194] */ MatcherIndex(15),
- /* [195] */ MatcherIndex(0),
+ /* [193] */ MatcherIndex(9),
+ /* [194] */ MatcherIndex(14),
+ /* [195] */ MatcherIndex(10),
/* [196] */ MatcherIndex(15),
- /* [197] */ MatcherIndex(9),
+ /* [197] */ MatcherIndex(0),
/* [198] */ MatcherIndex(15),
- /* [199] */ MatcherIndex(10),
- /* [200] */ MatcherIndex(16),
- /* [201] */ MatcherIndex(0),
+ /* [199] */ MatcherIndex(9),
+ /* [200] */ MatcherIndex(15),
+ /* [201] */ MatcherIndex(10),
/* [202] */ MatcherIndex(16),
- /* [203] */ MatcherIndex(9),
+ /* [203] */ MatcherIndex(0),
/* [204] */ MatcherIndex(16),
- /* [205] */ MatcherIndex(10),
- /* [206] */ MatcherIndex(17),
- /* [207] */ MatcherIndex(0),
+ /* [205] */ MatcherIndex(9),
+ /* [206] */ MatcherIndex(16),
+ /* [207] */ MatcherIndex(10),
/* [208] */ MatcherIndex(17),
- /* [209] */ MatcherIndex(9),
+ /* [209] */ MatcherIndex(0),
/* [210] */ MatcherIndex(17),
- /* [211] */ MatcherIndex(10),
- /* [212] */ MatcherIndex(18),
- /* [213] */ MatcherIndex(0),
+ /* [211] */ MatcherIndex(9),
+ /* [212] */ MatcherIndex(17),
+ /* [213] */ MatcherIndex(10),
/* [214] */ MatcherIndex(18),
- /* [215] */ MatcherIndex(9),
+ /* [215] */ MatcherIndex(0),
/* [216] */ MatcherIndex(18),
- /* [217] */ MatcherIndex(10),
- /* [218] */ MatcherIndex(19),
- /* [219] */ MatcherIndex(0),
+ /* [217] */ MatcherIndex(9),
+ /* [218] */ MatcherIndex(18),
+ /* [219] */ MatcherIndex(10),
/* [220] */ MatcherIndex(19),
- /* [221] */ MatcherIndex(9),
+ /* [221] */ MatcherIndex(0),
/* [222] */ MatcherIndex(19),
- /* [223] */ MatcherIndex(10),
- /* [224] */ MatcherIndex(20),
- /* [225] */ MatcherIndex(0),
+ /* [223] */ MatcherIndex(9),
+ /* [224] */ MatcherIndex(19),
+ /* [225] */ MatcherIndex(10),
/* [226] */ MatcherIndex(20),
- /* [227] */ MatcherIndex(9),
+ /* [227] */ MatcherIndex(0),
/* [228] */ MatcherIndex(20),
- /* [229] */ MatcherIndex(10),
- /* [230] */ MatcherIndex(21),
- /* [231] */ MatcherIndex(0),
+ /* [229] */ MatcherIndex(9),
+ /* [230] */ MatcherIndex(20),
+ /* [231] */ MatcherIndex(10),
/* [232] */ MatcherIndex(21),
- /* [233] */ MatcherIndex(9),
+ /* [233] */ MatcherIndex(0),
/* [234] */ MatcherIndex(21),
- /* [235] */ MatcherIndex(10),
- /* [236] */ MatcherIndex(22),
- /* [237] */ MatcherIndex(0),
+ /* [235] */ MatcherIndex(9),
+ /* [236] */ MatcherIndex(21),
+ /* [237] */ MatcherIndex(10),
/* [238] */ MatcherIndex(22),
- /* [239] */ MatcherIndex(9),
+ /* [239] */ MatcherIndex(0),
/* [240] */ MatcherIndex(22),
- /* [241] */ MatcherIndex(10),
- /* [242] */ MatcherIndex(48),
- /* [243] */ MatcherIndex(0),
- /* [244] */ MatcherIndex(62),
- /* [245] */ MatcherIndex(66),
- /* [246] */ MatcherIndex(74),
- /* [247] */ MatcherIndex(68),
- /* [248] */ MatcherIndex(55),
- /* [249] */ MatcherIndex(63),
- /* [250] */ MatcherIndex(38),
- /* [251] */ MatcherIndex(39),
- /* [252] */ MatcherIndex(40),
- /* [253] */ MatcherIndex(41),
- /* [254] */ MatcherIndex(42),
- /* [255] */ MatcherIndex(47),
- /* [256] */ MatcherIndex(29),
- /* [257] */ MatcherIndex(30),
- /* [258] */ MatcherIndex(6),
- /* [259] */ MatcherIndex(69),
- /* [260] */ MatcherIndex(67),
- /* [261] */ MatcherIndex(59),
- /* [262] */ MatcherIndex(60),
- /* [263] */ MatcherIndex(57),
- /* [264] */ MatcherIndex(58),
- /* [265] */ MatcherIndex(61),
- /* [266] */ MatcherIndex(56),
- /* [267] */ MatcherIndex(73),
+ /* [241] */ MatcherIndex(9),
+ /* [242] */ MatcherIndex(22),
+ /* [243] */ MatcherIndex(10),
+ /* [244] */ MatcherIndex(48),
+ /* [245] */ MatcherIndex(0),
+ /* [246] */ MatcherIndex(62),
+ /* [247] */ MatcherIndex(66),
+ /* [248] */ MatcherIndex(74),
+ /* [249] */ MatcherIndex(68),
+ /* [250] */ MatcherIndex(55),
+ /* [251] */ MatcherIndex(63),
+ /* [252] */ MatcherIndex(38),
+ /* [253] */ MatcherIndex(39),
+ /* [254] */ MatcherIndex(40),
+ /* [255] */ MatcherIndex(41),
+ /* [256] */ MatcherIndex(42),
+ /* [257] */ MatcherIndex(47),
+ /* [258] */ MatcherIndex(29),
+ /* [259] */ MatcherIndex(30),
+ /* [260] */ MatcherIndex(6),
+ /* [261] */ MatcherIndex(69),
+ /* [262] */ MatcherIndex(67),
+ /* [263] */ MatcherIndex(59),
+ /* [264] */ MatcherIndex(60),
+ /* [265] */ MatcherIndex(57),
+ /* [266] */ MatcherIndex(58),
+ /* [267] */ MatcherIndex(61),
+ /* [268] */ MatcherIndex(56),
+ /* [269] */ MatcherIndex(73),
};
static_assert(MatcherIndicesIndex::CanIndex(kMatcherIndices),
@@ -2232,7 +2234,7 @@
{
/* [20] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [21] */
@@ -2272,7 +2274,7 @@
{
/* [28] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [29] */
@@ -2292,12 +2294,12 @@
{
/* [32] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(251),
+ /* matcher_indices */ MatcherIndicesIndex(253),
},
{
/* [33] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(257),
+ /* matcher_indices */ MatcherIndicesIndex(259),
},
{
/* [34] */
@@ -2327,7 +2329,7 @@
{
/* [39] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [40] */
@@ -2357,7 +2359,7 @@
{
/* [45] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [46] */
@@ -2387,7 +2389,7 @@
{
/* [51] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [52] */
@@ -2417,7 +2419,7 @@
{
/* [57] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [58] */
@@ -2447,7 +2449,7 @@
{
/* [63] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [64] */
@@ -2472,12 +2474,12 @@
{
/* [68] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(251),
+ /* matcher_indices */ MatcherIndicesIndex(253),
},
{
/* [69] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [70] */
@@ -2512,7 +2514,7 @@
{
/* [76] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [77] */
@@ -2537,7 +2539,7 @@
{
/* [81] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [82] */
@@ -2552,12 +2554,12 @@
{
/* [84] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(251),
+ /* matcher_indices */ MatcherIndicesIndex(253),
},
{
/* [85] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [86] */
@@ -2577,12 +2579,12 @@
{
/* [89] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(250),
+ /* matcher_indices */ MatcherIndicesIndex(252),
},
{
/* [90] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(257),
+ /* matcher_indices */ MatcherIndicesIndex(259),
},
{
/* [91] */
@@ -2602,12 +2604,12 @@
{
/* [94] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(253),
+ /* matcher_indices */ MatcherIndicesIndex(255),
},
{
/* [95] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(257),
+ /* matcher_indices */ MatcherIndicesIndex(259),
},
{
/* [96] */
@@ -2632,7 +2634,7 @@
{
/* [100] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [101] */
@@ -2657,7 +2659,7 @@
{
/* [105] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [106] */
@@ -2682,7 +2684,7 @@
{
/* [110] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [111] */
@@ -2707,7 +2709,7 @@
{
/* [115] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [116] */
@@ -2732,7 +2734,7 @@
{
/* [120] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [121] */
@@ -2757,7 +2759,7 @@
{
/* [125] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [126] */
@@ -2782,7 +2784,7 @@
{
/* [130] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [131] */
@@ -2807,7 +2809,7 @@
{
/* [135] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [136] */
@@ -2827,12 +2829,12 @@
{
/* [139] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(250),
+ /* matcher_indices */ MatcherIndicesIndex(252),
},
{
/* [140] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [141] */
@@ -2852,12 +2854,12 @@
{
/* [144] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(253),
+ /* matcher_indices */ MatcherIndicesIndex(255),
},
{
/* [145] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [146] */
@@ -2907,7 +2909,7 @@
{
/* [155] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [156] */
@@ -2917,12 +2919,12 @@
{
/* [157] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(250),
+ /* matcher_indices */ MatcherIndicesIndex(252),
},
{
/* [158] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [159] */
@@ -2937,12 +2939,12 @@
{
/* [161] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(252),
+ /* matcher_indices */ MatcherIndicesIndex(254),
},
{
/* [162] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(257),
+ /* matcher_indices */ MatcherIndicesIndex(259),
},
{
/* [163] */
@@ -2962,7 +2964,7 @@
{
/* [166] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [167] */
@@ -2982,7 +2984,7 @@
{
/* [170] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [171] */
@@ -3002,7 +3004,7 @@
{
/* [174] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [175] */
@@ -3022,7 +3024,7 @@
{
/* [178] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [179] */
@@ -3037,12 +3039,12 @@
{
/* [181] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(252),
+ /* matcher_indices */ MatcherIndicesIndex(254),
},
{
/* [182] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [183] */
@@ -3137,7 +3139,7 @@
{
/* [201] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(251),
+ /* matcher_indices */ MatcherIndicesIndex(253),
},
{
/* [202] */
@@ -3322,7 +3324,7 @@
{
/* [238] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [239] */
@@ -3332,12 +3334,12 @@
{
/* [240] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(255),
+ /* matcher_indices */ MatcherIndicesIndex(257),
},
{
/* [241] */
/* usage */ core::ParameterUsage::kSampler,
- /* matcher_indices */ MatcherIndicesIndex(256),
+ /* matcher_indices */ MatcherIndicesIndex(258),
},
{
/* [242] */
@@ -3542,7 +3544,7 @@
{
/* [282] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(250),
+ /* matcher_indices */ MatcherIndicesIndex(252),
},
{
/* [283] */
@@ -3557,7 +3559,7 @@
{
/* [285] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(254),
+ /* matcher_indices */ MatcherIndicesIndex(256),
},
{
/* [286] */
@@ -3697,7 +3699,7 @@
{
/* [313] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(250),
+ /* matcher_indices */ MatcherIndicesIndex(252),
},
{
/* [314] */
@@ -3707,7 +3709,7 @@
{
/* [315] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(251),
+ /* matcher_indices */ MatcherIndicesIndex(253),
},
{
/* [316] */
@@ -3717,7 +3719,7 @@
{
/* [317] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(252),
+ /* matcher_indices */ MatcherIndicesIndex(254),
},
{
/* [318] */
@@ -3727,7 +3729,7 @@
{
/* [319] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(253),
+ /* matcher_indices */ MatcherIndicesIndex(255),
},
{
/* [320] */
@@ -3737,7 +3739,7 @@
{
/* [321] */
/* usage */ core::ParameterUsage::kTexture,
- /* matcher_indices */ MatcherIndicesIndex(255),
+ /* matcher_indices */ MatcherIndicesIndex(257),
},
{
/* [322] */
@@ -3971,38 +3973,38 @@
},
{
/* [368] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(30),
+ /* usage */ core::ParameterUsage::kInputAttachment,
+ /* matcher_indices */ MatcherIndicesIndex(180),
},
{
/* [369] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(26),
+ /* matcher_indices */ MatcherIndicesIndex(30),
},
{
/* [370] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(135),
+ /* matcher_indices */ MatcherIndicesIndex(26),
},
{
/* [371] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(178),
+ /* matcher_indices */ MatcherIndicesIndex(135),
},
{
/* [372] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(188),
+ /* matcher_indices */ MatcherIndicesIndex(178),
},
{
/* [373] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(23),
+ /* matcher_indices */ MatcherIndicesIndex(190),
},
{
/* [374] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(190),
+ /* matcher_indices */ MatcherIndicesIndex(23),
},
{
/* [375] */
@@ -4129,6 +4131,11 @@
/* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(240),
},
+ {
+ /* [400] */
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(242),
+ },
};
static_assert(ParameterIndex::CanIndex(kParameters),
@@ -4144,7 +4151,7 @@
{
/* [1] */
/* name */ "C",
- /* matcher_indices */ MatcherIndicesIndex(246),
+ /* matcher_indices */ MatcherIndicesIndex(248),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4162,19 +4169,19 @@
{
/* [4] */
/* name */ "A",
- /* matcher_indices */ MatcherIndicesIndex(246),
+ /* matcher_indices */ MatcherIndicesIndex(248),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [5] */
/* name */ "L",
- /* matcher_indices */ MatcherIndicesIndex(246),
+ /* matcher_indices */ MatcherIndicesIndex(248),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [6] */
/* name */ "C",
- /* matcher_indices */ MatcherIndicesIndex(246),
+ /* matcher_indices */ MatcherIndicesIndex(248),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4192,19 +4199,19 @@
{
/* [9] */
/* name */ "A",
- /* matcher_indices */ MatcherIndicesIndex(246),
+ /* matcher_indices */ MatcherIndicesIndex(248),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [10] */
/* name */ "C",
- /* matcher_indices */ MatcherIndicesIndex(246),
+ /* matcher_indices */ MatcherIndicesIndex(248),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [11] */
/* name */ "F",
- /* matcher_indices */ MatcherIndicesIndex(258),
+ /* matcher_indices */ MatcherIndicesIndex(260),
/* kind */ TemplateInfo::Kind::kNumber,
},
{
@@ -4216,7 +4223,7 @@
{
/* [13] */
/* name */ "A",
- /* matcher_indices */ MatcherIndicesIndex(246),
+ /* matcher_indices */ MatcherIndicesIndex(248),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4228,25 +4235,25 @@
{
/* [15] */
/* name */ "C",
- /* matcher_indices */ MatcherIndicesIndex(246),
+ /* matcher_indices */ MatcherIndicesIndex(248),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [16] */
/* name */ "A",
- /* matcher_indices */ MatcherIndicesIndex(246),
+ /* matcher_indices */ MatcherIndicesIndex(248),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [17] */
/* name */ "L",
- /* matcher_indices */ MatcherIndicesIndex(246),
+ /* matcher_indices */ MatcherIndicesIndex(248),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [18] */
/* name */ "C",
- /* matcher_indices */ MatcherIndicesIndex(246),
+ /* matcher_indices */ MatcherIndicesIndex(248),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4264,13 +4271,13 @@
{
/* [21] */
/* name */ "A",
- /* matcher_indices */ MatcherIndicesIndex(246),
+ /* matcher_indices */ MatcherIndicesIndex(248),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [22] */
/* name */ "C",
- /* matcher_indices */ MatcherIndicesIndex(246),
+ /* matcher_indices */ MatcherIndicesIndex(248),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4288,19 +4295,19 @@
{
/* [25] */
/* name */ "A",
- /* matcher_indices */ MatcherIndicesIndex(246),
+ /* matcher_indices */ MatcherIndicesIndex(248),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [26] */
/* name */ "C",
- /* matcher_indices */ MatcherIndicesIndex(246),
+ /* matcher_indices */ MatcherIndicesIndex(248),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [27] */
/* name */ "F",
- /* matcher_indices */ MatcherIndicesIndex(258),
+ /* matcher_indices */ MatcherIndicesIndex(260),
/* kind */ TemplateInfo::Kind::kNumber,
},
{
@@ -4312,13 +4319,13 @@
{
/* [29] */
/* name */ "A",
- /* matcher_indices */ MatcherIndicesIndex(246),
+ /* matcher_indices */ MatcherIndicesIndex(248),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [30] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(245),
+ /* matcher_indices */ MatcherIndicesIndex(247),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4366,13 +4373,13 @@
{
/* [38] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(245),
+ /* matcher_indices */ MatcherIndicesIndex(247),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [39] */
/* name */ "U",
- /* matcher_indices */ MatcherIndicesIndex(247),
+ /* matcher_indices */ MatcherIndicesIndex(249),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4390,7 +4397,7 @@
{
/* [42] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(245),
+ /* matcher_indices */ MatcherIndicesIndex(247),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4414,13 +4421,13 @@
{
/* [46] */
/* name */ "C",
- /* matcher_indices */ MatcherIndicesIndex(246),
+ /* matcher_indices */ MatcherIndicesIndex(248),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [47] */
/* name */ "L",
- /* matcher_indices */ MatcherIndicesIndex(246),
+ /* matcher_indices */ MatcherIndicesIndex(248),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4432,13 +4439,13 @@
{
/* [49] */
/* name */ "C",
- /* matcher_indices */ MatcherIndicesIndex(246),
+ /* matcher_indices */ MatcherIndicesIndex(248),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [50] */
/* name */ "S",
- /* matcher_indices */ MatcherIndicesIndex(246),
+ /* matcher_indices */ MatcherIndicesIndex(248),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4462,7 +4469,7 @@
{
/* [54] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(245),
+ /* matcher_indices */ MatcherIndicesIndex(247),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4486,7 +4493,7 @@
{
/* [58] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(244),
+ /* matcher_indices */ MatcherIndicesIndex(246),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4498,7 +4505,7 @@
{
/* [60] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(246),
+ /* matcher_indices */ MatcherIndicesIndex(248),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4534,7 +4541,7 @@
{
/* [66] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(248),
+ /* matcher_indices */ MatcherIndicesIndex(250),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4546,7 +4553,7 @@
{
/* [68] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(249),
+ /* matcher_indices */ MatcherIndicesIndex(251),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4558,7 +4565,7 @@
{
/* [70] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(259),
+ /* matcher_indices */ MatcherIndicesIndex(261),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4606,7 +4613,7 @@
{
/* [78] */
/* name */ "L",
- /* matcher_indices */ MatcherIndicesIndex(246),
+ /* matcher_indices */ MatcherIndicesIndex(248),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4624,7 +4631,7 @@
{
/* [81] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(260),
+ /* matcher_indices */ MatcherIndicesIndex(262),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4642,7 +4649,7 @@
{
/* [84] */
/* name */ "U",
- /* matcher_indices */ MatcherIndicesIndex(263),
+ /* matcher_indices */ MatcherIndicesIndex(265),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4654,7 +4661,7 @@
{
/* [86] */
/* name */ "U",
- /* matcher_indices */ MatcherIndicesIndex(264),
+ /* matcher_indices */ MatcherIndicesIndex(266),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4666,7 +4673,7 @@
{
/* [88] */
/* name */ "U",
- /* matcher_indices */ MatcherIndicesIndex(261),
+ /* matcher_indices */ MatcherIndicesIndex(263),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4678,7 +4685,7 @@
{
/* [90] */
/* name */ "U",
- /* matcher_indices */ MatcherIndicesIndex(262),
+ /* matcher_indices */ MatcherIndicesIndex(264),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4690,7 +4697,7 @@
{
/* [92] */
/* name */ "U",
- /* matcher_indices */ MatcherIndicesIndex(265),
+ /* matcher_indices */ MatcherIndicesIndex(267),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -4702,43 +4709,43 @@
{
/* [94] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(261),
+ /* matcher_indices */ MatcherIndicesIndex(263),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [95] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(262),
+ /* matcher_indices */ MatcherIndicesIndex(264),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [96] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(263),
+ /* matcher_indices */ MatcherIndicesIndex(265),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [97] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(264),
+ /* matcher_indices */ MatcherIndicesIndex(266),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [98] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(265),
+ /* matcher_indices */ MatcherIndicesIndex(267),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [99] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(266),
+ /* matcher_indices */ MatcherIndicesIndex(268),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [100] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(267),
+ /* matcher_indices */ MatcherIndicesIndex(269),
/* kind */ TemplateInfo::Kind::kType,
},
};
@@ -5172,7 +5179,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
- /* return_matcher_indices */ MatcherIndicesIndex(186),
+ /* return_matcher_indices */ MatcherIndicesIndex(188),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -5391,7 +5398,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 2,
/* templates */ TemplateIndex(83),
- /* parameters */ ParameterIndex(372),
+ /* parameters */ ParameterIndex(373),
/* return_matcher_indices */ MatcherIndicesIndex(156),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
@@ -5402,7 +5409,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 2,
/* templates */ TemplateIndex(85),
- /* parameters */ ParameterIndex(372),
+ /* parameters */ ParameterIndex(373),
/* return_matcher_indices */ MatcherIndicesIndex(156),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
@@ -5413,7 +5420,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 2,
/* templates */ TemplateIndex(87),
- /* parameters */ ParameterIndex(372),
+ /* parameters */ ParameterIndex(373),
/* return_matcher_indices */ MatcherIndicesIndex(156),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
@@ -5424,7 +5431,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 2,
/* templates */ TemplateIndex(89),
- /* parameters */ ParameterIndex(372),
+ /* parameters */ ParameterIndex(373),
/* return_matcher_indices */ MatcherIndicesIndex(156),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
@@ -5435,7 +5442,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 2,
/* templates */ TemplateIndex(91),
- /* parameters */ ParameterIndex(372),
+ /* parameters */ ParameterIndex(373),
/* return_matcher_indices */ MatcherIndicesIndex(156),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
@@ -5678,7 +5685,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
- /* return_matcher_indices */ MatcherIndicesIndex(184),
+ /* return_matcher_indices */ MatcherIndicesIndex(186),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -5809,7 +5816,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 2,
/* templates */ TemplateIndex(83),
- /* parameters */ ParameterIndex(371),
+ /* parameters */ ParameterIndex(372),
/* return_matcher_indices */ MatcherIndicesIndex(124),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
@@ -5820,7 +5827,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 2,
/* templates */ TemplateIndex(85),
- /* parameters */ ParameterIndex(371),
+ /* parameters */ ParameterIndex(372),
/* return_matcher_indices */ MatcherIndicesIndex(124),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
@@ -5831,7 +5838,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 2,
/* templates */ TemplateIndex(87),
- /* parameters */ ParameterIndex(371),
+ /* parameters */ ParameterIndex(372),
/* return_matcher_indices */ MatcherIndicesIndex(124),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
@@ -5842,7 +5849,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 2,
/* templates */ TemplateIndex(89),
- /* parameters */ ParameterIndex(371),
+ /* parameters */ ParameterIndex(372),
/* return_matcher_indices */ MatcherIndicesIndex(124),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
@@ -5853,7 +5860,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 2,
/* templates */ TemplateIndex(91),
- /* parameters */ ParameterIndex(371),
+ /* parameters */ ParameterIndex(372),
/* return_matcher_indices */ MatcherIndicesIndex(124),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
@@ -6184,7 +6191,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(/* invalid */),
- /* return_matcher_indices */ MatcherIndicesIndex(182),
+ /* return_matcher_indices */ MatcherIndicesIndex(184),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -6920,7 +6927,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(100),
- /* parameters */ ParameterIndex(373),
+ /* parameters */ ParameterIndex(374),
/* return_matcher_indices */ MatcherIndicesIndex(23),
/* const_eval_fn */ ConstEvalFunctionIndex(82),
},
@@ -6931,7 +6938,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
- /* parameters */ ParameterIndex(373),
+ /* parameters */ ParameterIndex(374),
/* return_matcher_indices */ MatcherIndicesIndex(23),
/* const_eval_fn */ ConstEvalFunctionIndex(82),
},
@@ -6986,7 +6993,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(85),
- /* parameters */ ParameterIndex(374),
+ /* parameters */ ParameterIndex(375),
/* return_matcher_indices */ MatcherIndicesIndex(23),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
@@ -6997,7 +7004,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(83),
- /* parameters */ ParameterIndex(375),
+ /* parameters */ ParameterIndex(376),
/* return_matcher_indices */ MatcherIndicesIndex(23),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
@@ -7009,7 +7016,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(100),
/* parameters */ ParameterIndex(/* invalid */),
- /* return_matcher_indices */ MatcherIndicesIndex(194),
+ /* return_matcher_indices */ MatcherIndicesIndex(196),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -7019,8 +7026,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(100),
- /* parameters */ ParameterIndex(376),
- /* return_matcher_indices */ MatcherIndicesIndex(194),
+ /* parameters */ ParameterIndex(377),
+ /* return_matcher_indices */ MatcherIndicesIndex(196),
/* const_eval_fn */ ConstEvalFunctionIndex(82),
},
{
@@ -7030,8 +7037,8 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
- /* parameters */ ParameterIndex(376),
- /* return_matcher_indices */ MatcherIndicesIndex(194),
+ /* parameters */ ParameterIndex(377),
+ /* return_matcher_indices */ MatcherIndicesIndex(196),
/* const_eval_fn */ ConstEvalFunctionIndex(82),
},
{
@@ -7042,7 +7049,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(1),
- /* return_matcher_indices */ MatcherIndicesIndex(194),
+ /* return_matcher_indices */ MatcherIndicesIndex(196),
/* const_eval_fn */ ConstEvalFunctionIndex(112),
},
{
@@ -7053,7 +7060,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(1),
- /* return_matcher_indices */ MatcherIndicesIndex(194),
+ /* return_matcher_indices */ MatcherIndicesIndex(196),
/* const_eval_fn */ ConstEvalFunctionIndex(112),
},
{
@@ -7064,7 +7071,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(213),
- /* return_matcher_indices */ MatcherIndicesIndex(194),
+ /* return_matcher_indices */ MatcherIndicesIndex(196),
/* const_eval_fn */ ConstEvalFunctionIndex(113),
},
{
@@ -7075,7 +7082,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(213),
- /* return_matcher_indices */ MatcherIndicesIndex(194),
+ /* return_matcher_indices */ MatcherIndicesIndex(196),
/* const_eval_fn */ ConstEvalFunctionIndex(113),
},
{
@@ -7085,8 +7092,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(85),
- /* parameters */ ParameterIndex(377),
- /* return_matcher_indices */ MatcherIndicesIndex(194),
+ /* parameters */ ParameterIndex(378),
+ /* return_matcher_indices */ MatcherIndicesIndex(196),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
{
@@ -7096,8 +7103,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(83),
- /* parameters */ ParameterIndex(378),
- /* return_matcher_indices */ MatcherIndicesIndex(194),
+ /* parameters */ ParameterIndex(379),
+ /* return_matcher_indices */ MatcherIndicesIndex(196),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
{
@@ -7108,7 +7115,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(100),
/* parameters */ ParameterIndex(/* invalid */),
- /* return_matcher_indices */ MatcherIndicesIndex(200),
+ /* return_matcher_indices */ MatcherIndicesIndex(202),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -7118,8 +7125,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(100),
- /* parameters */ ParameterIndex(379),
- /* return_matcher_indices */ MatcherIndicesIndex(200),
+ /* parameters */ ParameterIndex(380),
+ /* return_matcher_indices */ MatcherIndicesIndex(202),
/* const_eval_fn */ ConstEvalFunctionIndex(82),
},
{
@@ -7129,8 +7136,8 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
- /* parameters */ ParameterIndex(379),
- /* return_matcher_indices */ MatcherIndicesIndex(200),
+ /* parameters */ ParameterIndex(380),
+ /* return_matcher_indices */ MatcherIndicesIndex(202),
/* const_eval_fn */ ConstEvalFunctionIndex(82),
},
{
@@ -7141,7 +7148,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(1),
- /* return_matcher_indices */ MatcherIndicesIndex(200),
+ /* return_matcher_indices */ MatcherIndicesIndex(202),
/* const_eval_fn */ ConstEvalFunctionIndex(112),
},
{
@@ -7152,7 +7159,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(1),
- /* return_matcher_indices */ MatcherIndicesIndex(200),
+ /* return_matcher_indices */ MatcherIndicesIndex(202),
/* const_eval_fn */ ConstEvalFunctionIndex(112),
},
{
@@ -7163,7 +7170,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(217),
- /* return_matcher_indices */ MatcherIndicesIndex(200),
+ /* return_matcher_indices */ MatcherIndicesIndex(202),
/* const_eval_fn */ ConstEvalFunctionIndex(113),
},
{
@@ -7174,7 +7181,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(217),
- /* return_matcher_indices */ MatcherIndicesIndex(200),
+ /* return_matcher_indices */ MatcherIndicesIndex(202),
/* const_eval_fn */ ConstEvalFunctionIndex(113),
},
{
@@ -7184,8 +7191,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(85),
- /* parameters */ ParameterIndex(380),
- /* return_matcher_indices */ MatcherIndicesIndex(200),
+ /* parameters */ ParameterIndex(381),
+ /* return_matcher_indices */ MatcherIndicesIndex(202),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
{
@@ -7195,8 +7202,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(83),
- /* parameters */ ParameterIndex(381),
- /* return_matcher_indices */ MatcherIndicesIndex(200),
+ /* parameters */ ParameterIndex(382),
+ /* return_matcher_indices */ MatcherIndicesIndex(202),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
{
@@ -7207,7 +7214,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(100),
/* parameters */ ParameterIndex(/* invalid */),
- /* return_matcher_indices */ MatcherIndicesIndex(206),
+ /* return_matcher_indices */ MatcherIndicesIndex(208),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -7217,8 +7224,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(100),
- /* parameters */ ParameterIndex(382),
- /* return_matcher_indices */ MatcherIndicesIndex(206),
+ /* parameters */ ParameterIndex(383),
+ /* return_matcher_indices */ MatcherIndicesIndex(208),
/* const_eval_fn */ ConstEvalFunctionIndex(82),
},
{
@@ -7228,8 +7235,8 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
- /* parameters */ ParameterIndex(382),
- /* return_matcher_indices */ MatcherIndicesIndex(206),
+ /* parameters */ ParameterIndex(383),
+ /* return_matcher_indices */ MatcherIndicesIndex(208),
/* const_eval_fn */ ConstEvalFunctionIndex(82),
},
{
@@ -7240,7 +7247,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(1),
- /* return_matcher_indices */ MatcherIndicesIndex(206),
+ /* return_matcher_indices */ MatcherIndicesIndex(208),
/* const_eval_fn */ ConstEvalFunctionIndex(112),
},
{
@@ -7251,7 +7258,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(1),
- /* return_matcher_indices */ MatcherIndicesIndex(206),
+ /* return_matcher_indices */ MatcherIndicesIndex(208),
/* const_eval_fn */ ConstEvalFunctionIndex(112),
},
{
@@ -7262,7 +7269,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(209),
- /* return_matcher_indices */ MatcherIndicesIndex(206),
+ /* return_matcher_indices */ MatcherIndicesIndex(208),
/* const_eval_fn */ ConstEvalFunctionIndex(113),
},
{
@@ -7273,7 +7280,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(209),
- /* return_matcher_indices */ MatcherIndicesIndex(206),
+ /* return_matcher_indices */ MatcherIndicesIndex(208),
/* const_eval_fn */ ConstEvalFunctionIndex(113),
},
{
@@ -7283,8 +7290,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(85),
- /* parameters */ ParameterIndex(383),
- /* return_matcher_indices */ MatcherIndicesIndex(206),
+ /* parameters */ ParameterIndex(384),
+ /* return_matcher_indices */ MatcherIndicesIndex(208),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
{
@@ -7294,8 +7301,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(83),
- /* parameters */ ParameterIndex(384),
- /* return_matcher_indices */ MatcherIndicesIndex(206),
+ /* parameters */ ParameterIndex(385),
+ /* return_matcher_indices */ MatcherIndicesIndex(208),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
{
@@ -7306,7 +7313,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(100),
/* parameters */ ParameterIndex(/* invalid */),
- /* return_matcher_indices */ MatcherIndicesIndex(212),
+ /* return_matcher_indices */ MatcherIndicesIndex(214),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -7316,8 +7323,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(100),
- /* parameters */ ParameterIndex(385),
- /* return_matcher_indices */ MatcherIndicesIndex(212),
+ /* parameters */ ParameterIndex(386),
+ /* return_matcher_indices */ MatcherIndicesIndex(214),
/* const_eval_fn */ ConstEvalFunctionIndex(82),
},
{
@@ -7327,8 +7334,8 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
- /* parameters */ ParameterIndex(385),
- /* return_matcher_indices */ MatcherIndicesIndex(212),
+ /* parameters */ ParameterIndex(386),
+ /* return_matcher_indices */ MatcherIndicesIndex(214),
/* const_eval_fn */ ConstEvalFunctionIndex(82),
},
{
@@ -7339,7 +7346,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(1),
- /* return_matcher_indices */ MatcherIndicesIndex(212),
+ /* return_matcher_indices */ MatcherIndicesIndex(214),
/* const_eval_fn */ ConstEvalFunctionIndex(112),
},
{
@@ -7350,7 +7357,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(1),
- /* return_matcher_indices */ MatcherIndicesIndex(212),
+ /* return_matcher_indices */ MatcherIndicesIndex(214),
/* const_eval_fn */ ConstEvalFunctionIndex(112),
},
{
@@ -7361,7 +7368,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(213),
- /* return_matcher_indices */ MatcherIndicesIndex(212),
+ /* return_matcher_indices */ MatcherIndicesIndex(214),
/* const_eval_fn */ ConstEvalFunctionIndex(113),
},
{
@@ -7372,7 +7379,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(213),
- /* return_matcher_indices */ MatcherIndicesIndex(212),
+ /* return_matcher_indices */ MatcherIndicesIndex(214),
/* const_eval_fn */ ConstEvalFunctionIndex(113),
},
{
@@ -7382,8 +7389,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(85),
- /* parameters */ ParameterIndex(386),
- /* return_matcher_indices */ MatcherIndicesIndex(212),
+ /* parameters */ ParameterIndex(387),
+ /* return_matcher_indices */ MatcherIndicesIndex(214),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
{
@@ -7393,8 +7400,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(83),
- /* parameters */ ParameterIndex(387),
- /* return_matcher_indices */ MatcherIndicesIndex(212),
+ /* parameters */ ParameterIndex(388),
+ /* return_matcher_indices */ MatcherIndicesIndex(214),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
{
@@ -7405,7 +7412,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(100),
/* parameters */ ParameterIndex(/* invalid */),
- /* return_matcher_indices */ MatcherIndicesIndex(218),
+ /* return_matcher_indices */ MatcherIndicesIndex(220),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -7415,8 +7422,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(100),
- /* parameters */ ParameterIndex(388),
- /* return_matcher_indices */ MatcherIndicesIndex(218),
+ /* parameters */ ParameterIndex(389),
+ /* return_matcher_indices */ MatcherIndicesIndex(220),
/* const_eval_fn */ ConstEvalFunctionIndex(82),
},
{
@@ -7426,8 +7433,8 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
- /* parameters */ ParameterIndex(388),
- /* return_matcher_indices */ MatcherIndicesIndex(218),
+ /* parameters */ ParameterIndex(389),
+ /* return_matcher_indices */ MatcherIndicesIndex(220),
/* const_eval_fn */ ConstEvalFunctionIndex(82),
},
{
@@ -7438,7 +7445,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(1),
- /* return_matcher_indices */ MatcherIndicesIndex(218),
+ /* return_matcher_indices */ MatcherIndicesIndex(220),
/* const_eval_fn */ ConstEvalFunctionIndex(112),
},
{
@@ -7449,7 +7456,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(1),
- /* return_matcher_indices */ MatcherIndicesIndex(218),
+ /* return_matcher_indices */ MatcherIndicesIndex(220),
/* const_eval_fn */ ConstEvalFunctionIndex(112),
},
{
@@ -7460,7 +7467,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(217),
- /* return_matcher_indices */ MatcherIndicesIndex(218),
+ /* return_matcher_indices */ MatcherIndicesIndex(220),
/* const_eval_fn */ ConstEvalFunctionIndex(113),
},
{
@@ -7471,7 +7478,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(217),
- /* return_matcher_indices */ MatcherIndicesIndex(218),
+ /* return_matcher_indices */ MatcherIndicesIndex(220),
/* const_eval_fn */ ConstEvalFunctionIndex(113),
},
{
@@ -7481,8 +7488,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(85),
- /* parameters */ ParameterIndex(389),
- /* return_matcher_indices */ MatcherIndicesIndex(218),
+ /* parameters */ ParameterIndex(390),
+ /* return_matcher_indices */ MatcherIndicesIndex(220),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
{
@@ -7492,8 +7499,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(83),
- /* parameters */ ParameterIndex(390),
- /* return_matcher_indices */ MatcherIndicesIndex(218),
+ /* parameters */ ParameterIndex(391),
+ /* return_matcher_indices */ MatcherIndicesIndex(220),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
{
@@ -7504,7 +7511,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(100),
/* parameters */ ParameterIndex(/* invalid */),
- /* return_matcher_indices */ MatcherIndicesIndex(224),
+ /* return_matcher_indices */ MatcherIndicesIndex(226),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -7514,8 +7521,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(100),
- /* parameters */ ParameterIndex(391),
- /* return_matcher_indices */ MatcherIndicesIndex(224),
+ /* parameters */ ParameterIndex(392),
+ /* return_matcher_indices */ MatcherIndicesIndex(226),
/* const_eval_fn */ ConstEvalFunctionIndex(82),
},
{
@@ -7525,8 +7532,8 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
- /* parameters */ ParameterIndex(391),
- /* return_matcher_indices */ MatcherIndicesIndex(224),
+ /* parameters */ ParameterIndex(392),
+ /* return_matcher_indices */ MatcherIndicesIndex(226),
/* const_eval_fn */ ConstEvalFunctionIndex(82),
},
{
@@ -7537,7 +7544,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(1),
- /* return_matcher_indices */ MatcherIndicesIndex(224),
+ /* return_matcher_indices */ MatcherIndicesIndex(226),
/* const_eval_fn */ ConstEvalFunctionIndex(112),
},
{
@@ -7548,7 +7555,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(1),
- /* return_matcher_indices */ MatcherIndicesIndex(224),
+ /* return_matcher_indices */ MatcherIndicesIndex(226),
/* const_eval_fn */ ConstEvalFunctionIndex(112),
},
{
@@ -7559,7 +7566,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(209),
- /* return_matcher_indices */ MatcherIndicesIndex(224),
+ /* return_matcher_indices */ MatcherIndicesIndex(226),
/* const_eval_fn */ ConstEvalFunctionIndex(113),
},
{
@@ -7570,7 +7577,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(209),
- /* return_matcher_indices */ MatcherIndicesIndex(224),
+ /* return_matcher_indices */ MatcherIndicesIndex(226),
/* const_eval_fn */ ConstEvalFunctionIndex(113),
},
{
@@ -7580,8 +7587,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(85),
- /* parameters */ ParameterIndex(392),
- /* return_matcher_indices */ MatcherIndicesIndex(224),
+ /* parameters */ ParameterIndex(393),
+ /* return_matcher_indices */ MatcherIndicesIndex(226),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
{
@@ -7591,8 +7598,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(83),
- /* parameters */ ParameterIndex(393),
- /* return_matcher_indices */ MatcherIndicesIndex(224),
+ /* parameters */ ParameterIndex(394),
+ /* return_matcher_indices */ MatcherIndicesIndex(226),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
{
@@ -7603,7 +7610,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(100),
/* parameters */ ParameterIndex(/* invalid */),
- /* return_matcher_indices */ MatcherIndicesIndex(230),
+ /* return_matcher_indices */ MatcherIndicesIndex(232),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -7613,8 +7620,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(100),
- /* parameters */ ParameterIndex(394),
- /* return_matcher_indices */ MatcherIndicesIndex(230),
+ /* parameters */ ParameterIndex(395),
+ /* return_matcher_indices */ MatcherIndicesIndex(232),
/* const_eval_fn */ ConstEvalFunctionIndex(82),
},
{
@@ -7624,8 +7631,8 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
- /* parameters */ ParameterIndex(394),
- /* return_matcher_indices */ MatcherIndicesIndex(230),
+ /* parameters */ ParameterIndex(395),
+ /* return_matcher_indices */ MatcherIndicesIndex(232),
/* const_eval_fn */ ConstEvalFunctionIndex(82),
},
{
@@ -7636,7 +7643,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(1),
- /* return_matcher_indices */ MatcherIndicesIndex(230),
+ /* return_matcher_indices */ MatcherIndicesIndex(232),
/* const_eval_fn */ ConstEvalFunctionIndex(112),
},
{
@@ -7647,7 +7654,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(1),
- /* return_matcher_indices */ MatcherIndicesIndex(230),
+ /* return_matcher_indices */ MatcherIndicesIndex(232),
/* const_eval_fn */ ConstEvalFunctionIndex(112),
},
{
@@ -7658,7 +7665,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(213),
- /* return_matcher_indices */ MatcherIndicesIndex(230),
+ /* return_matcher_indices */ MatcherIndicesIndex(232),
/* const_eval_fn */ ConstEvalFunctionIndex(113),
},
{
@@ -7669,7 +7676,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(213),
- /* return_matcher_indices */ MatcherIndicesIndex(230),
+ /* return_matcher_indices */ MatcherIndicesIndex(232),
/* const_eval_fn */ ConstEvalFunctionIndex(113),
},
{
@@ -7679,8 +7686,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(85),
- /* parameters */ ParameterIndex(395),
- /* return_matcher_indices */ MatcherIndicesIndex(230),
+ /* parameters */ ParameterIndex(396),
+ /* return_matcher_indices */ MatcherIndicesIndex(232),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
{
@@ -7690,8 +7697,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(83),
- /* parameters */ ParameterIndex(396),
- /* return_matcher_indices */ MatcherIndicesIndex(230),
+ /* parameters */ ParameterIndex(397),
+ /* return_matcher_indices */ MatcherIndicesIndex(232),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
{
@@ -7702,7 +7709,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(100),
/* parameters */ ParameterIndex(/* invalid */),
- /* return_matcher_indices */ MatcherIndicesIndex(236),
+ /* return_matcher_indices */ MatcherIndicesIndex(238),
/* const_eval_fn */ ConstEvalFunctionIndex(107),
},
{
@@ -7712,8 +7719,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(100),
- /* parameters */ ParameterIndex(397),
- /* return_matcher_indices */ MatcherIndicesIndex(236),
+ /* parameters */ ParameterIndex(398),
+ /* return_matcher_indices */ MatcherIndicesIndex(238),
/* const_eval_fn */ ConstEvalFunctionIndex(82),
},
{
@@ -7723,8 +7730,8 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
- /* parameters */ ParameterIndex(397),
- /* return_matcher_indices */ MatcherIndicesIndex(236),
+ /* parameters */ ParameterIndex(398),
+ /* return_matcher_indices */ MatcherIndicesIndex(238),
/* const_eval_fn */ ConstEvalFunctionIndex(82),
},
{
@@ -7735,7 +7742,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(1),
- /* return_matcher_indices */ MatcherIndicesIndex(236),
+ /* return_matcher_indices */ MatcherIndicesIndex(238),
/* const_eval_fn */ ConstEvalFunctionIndex(112),
},
{
@@ -7746,7 +7753,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(1),
- /* return_matcher_indices */ MatcherIndicesIndex(236),
+ /* return_matcher_indices */ MatcherIndicesIndex(238),
/* const_eval_fn */ ConstEvalFunctionIndex(112),
},
{
@@ -7757,7 +7764,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(217),
- /* return_matcher_indices */ MatcherIndicesIndex(236),
+ /* return_matcher_indices */ MatcherIndicesIndex(238),
/* const_eval_fn */ ConstEvalFunctionIndex(113),
},
{
@@ -7768,7 +7775,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(30),
/* parameters */ ParameterIndex(217),
- /* return_matcher_indices */ MatcherIndicesIndex(236),
+ /* return_matcher_indices */ MatcherIndicesIndex(238),
/* const_eval_fn */ ConstEvalFunctionIndex(113),
},
{
@@ -7778,8 +7785,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(85),
- /* parameters */ ParameterIndex(398),
- /* return_matcher_indices */ MatcherIndicesIndex(236),
+ /* parameters */ ParameterIndex(399),
+ /* return_matcher_indices */ MatcherIndicesIndex(238),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
{
@@ -7789,8 +7796,8 @@
/* num_explicit_templates */ 1,
/* num_templates */ 1,
/* templates */ TemplateIndex(83),
- /* parameters */ ParameterIndex(399),
- /* return_matcher_indices */ MatcherIndicesIndex(236),
+ /* parameters */ ParameterIndex(400),
+ /* return_matcher_indices */ MatcherIndicesIndex(238),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
{
@@ -8592,7 +8599,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(370),
+ /* parameters */ ParameterIndex(371),
/* return_matcher_indices */ MatcherIndicesIndex(135),
/* const_eval_fn */ ConstEvalFunctionIndex(82),
},
@@ -10243,7 +10250,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(99),
/* parameters */ ParameterIndex(213),
- /* return_matcher_indices */ MatcherIndicesIndex(242),
+ /* return_matcher_indices */ MatcherIndicesIndex(244),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
{
@@ -10254,7 +10261,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(99),
/* parameters */ ParameterIndex(213),
- /* return_matcher_indices */ MatcherIndicesIndex(242),
+ /* return_matcher_indices */ MatcherIndicesIndex(244),
/* const_eval_fn */ ConstEvalFunctionIndex(108),
},
{
@@ -10578,6 +10585,17 @@
},
{
/* [519] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline),
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 1,
+ /* templates */ TemplateIndex(0),
+ /* parameters */ ParameterIndex(368),
+ /* return_matcher_indices */ MatcherIndicesIndex(156),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [520] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -10588,17 +10606,6 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [520] */
- /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 2,
- /* num_explicit_templates */ 0,
- /* num_templates */ 2,
- /* templates */ TemplateIndex(60),
- /* parameters */ ParameterIndex(0),
- /* return_matcher_indices */ MatcherIndicesIndex(/* invalid */),
- /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
- },
- {
/* [521] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
@@ -10606,22 +10613,33 @@
/* num_templates */ 2,
/* templates */ TemplateIndex(60),
/* parameters */ ParameterIndex(0),
- /* return_matcher_indices */ MatcherIndicesIndex(3),
+ /* return_matcher_indices */ MatcherIndicesIndex(/* invalid */),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
/* [522] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 2,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 2,
+ /* templates */ TemplateIndex(60),
+ /* parameters */ ParameterIndex(0),
+ /* return_matcher_indices */ MatcherIndicesIndex(3),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [523] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(60),
/* parameters */ ParameterIndex(0),
- /* return_matcher_indices */ MatcherIndicesIndex(180),
+ /* return_matcher_indices */ MatcherIndicesIndex(182),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [523] */
+ /* [524] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMustUse),
/* num_parameters */ 0,
/* num_explicit_templates */ 0,
@@ -10632,7 +10650,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [524] */
+ /* [525] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -10643,17 +10661,6 @@
/* const_eval_fn */ ConstEvalFunctionIndex(82),
},
{
- /* [525] */
- /* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMustUse),
- /* num_parameters */ 1,
- /* num_explicit_templates */ 0,
- /* num_templates */ 3,
- /* templates */ TemplateIndex(51),
- /* parameters */ ParameterIndex(368),
- /* return_matcher_indices */ MatcherIndicesIndex(26),
- /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
- },
- {
/* [526] */
/* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMustUse),
/* num_parameters */ 1,
@@ -10661,12 +10668,23 @@
/* num_templates */ 3,
/* templates */ TemplateIndex(51),
/* parameters */ ParameterIndex(369),
- /* return_matcher_indices */ MatcherIndicesIndex(30),
+ /* return_matcher_indices */ MatcherIndicesIndex(26),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
/* [527] */
/* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMustUse),
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 3,
+ /* templates */ TemplateIndex(51),
+ /* parameters */ ParameterIndex(370),
+ /* return_matcher_indices */ MatcherIndicesIndex(30),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [528] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMustUse),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 0,
@@ -10676,7 +10694,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(97),
},
{
- /* [528] */
+ /* [529] */
/* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMustUse),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -11567,88 +11585,94 @@
},
{
/* [110] */
- /* fn atomicLoad[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>) -> T */
+ /* fn inputAttachmentLoad[T : fiu32](input_attachment: input_attachment<T>) -> vec4<T> */
/* num overloads */ 1,
/* overloads */ OverloadIndex(519),
},
{
/* [111] */
- /* fn atomicStore[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) */
+ /* fn atomicLoad[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>) -> T */
/* num overloads */ 1,
/* overloads */ OverloadIndex(520),
},
{
/* [112] */
- /* fn atomicAdd[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
+ /* fn atomicStore[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) */
/* num overloads */ 1,
/* overloads */ OverloadIndex(521),
},
{
/* [113] */
- /* fn atomicSub[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
+ /* fn atomicAdd[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(521),
+ /* overloads */ OverloadIndex(522),
},
{
/* [114] */
- /* fn atomicMax[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
+ /* fn atomicSub[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(521),
+ /* overloads */ OverloadIndex(522),
},
{
/* [115] */
- /* fn atomicMin[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
+ /* fn atomicMax[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(521),
+ /* overloads */ OverloadIndex(522),
},
{
/* [116] */
- /* fn atomicAnd[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
+ /* fn atomicMin[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(521),
+ /* overloads */ OverloadIndex(522),
},
{
/* [117] */
- /* fn atomicOr[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
+ /* fn atomicAnd[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(521),
+ /* overloads */ OverloadIndex(522),
},
{
/* [118] */
- /* fn atomicXor[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
+ /* fn atomicOr[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(521),
+ /* overloads */ OverloadIndex(522),
},
{
/* [119] */
- /* fn atomicExchange[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
+ /* fn atomicXor[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(521),
+ /* overloads */ OverloadIndex(522),
},
{
/* [120] */
- /* fn atomicCompareExchangeWeak[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, T) -> __atomic_compare_exchange_result<T> */
+ /* fn atomicExchange[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
/* num overloads */ 1,
/* overloads */ OverloadIndex(522),
},
{
/* [121] */
- /* fn subgroupBallot() -> vec4<u32> */
+ /* fn atomicCompareExchangeWeak[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, T) -> __atomic_compare_exchange_result<T> */
/* num overloads */ 1,
/* overloads */ OverloadIndex(523),
},
{
/* [122] */
+ /* fn subgroupBallot() -> vec4<u32> */
+ /* num overloads */ 1,
+ /* overloads */ OverloadIndex(524),
+ },
+ {
+ /* [123] */
/* fn subgroupBroadcast[T : fiu32_f16](value: T, @const sourceLaneIndex: u32) -> T */
/* fn subgroupBroadcast[N : num, T : fiu32_f16](value: vec<N, T>, @const sourceLaneIndex: u32) -> vec<N, T> */
/* num overloads */ 2,
/* overloads */ OverloadIndex(462),
},
{
- /* [123] */
+ /* [124] */
/* fn __tint_materialize[T](T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(524),
+ /* overloads */ OverloadIndex(525),
},
};
@@ -11657,13 +11681,13 @@
/* [0] */
/* op &[S : address_space, T, A : access](ref<S, T, A>) -> ptr<S, T, A> */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(525),
+ /* overloads */ OverloadIndex(526),
},
{
/* [1] */
/* op *[S : address_space, T, A : access](ptr<S, T, A>) -> ref<S, T, A> */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(526),
+ /* overloads */ OverloadIndex(527),
},
{
/* [2] */
@@ -11775,13 +11799,13 @@
/* [8] */
/* op &&(bool, bool) -> bool */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(527),
+ /* overloads */ OverloadIndex(528),
},
{
/* [9] */
/* op ||(bool, bool) -> bool */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(528),
+ /* overloads */ OverloadIndex(529),
},
{
/* [10] */
diff --git a/src/tint/lang/wgsl/ir_roundtrip_test.cc b/src/tint/lang/wgsl/ir_roundtrip_test.cc
index 60345a5..df1f324 100644
--- a/src/tint/lang/wgsl/ir_roundtrip_test.cc
+++ b/src/tint/lang/wgsl/ir_roundtrip_test.cc
@@ -3314,6 +3314,22 @@
)");
}
+////////////////////////////////////////////////////////////////////////////////
+// chromium_internal_input_attachments
+////////////////////////////////////////////////////////////////////////////////
+TEST_F(IRToProgramRoundtripTest, Call_InputAttachmentLoad) {
+ RUN_TEST(R"(
+enable chromium_internal_input_attachments;
+
+@group(0u) @binding(0u) @input_attachment_index(3u) var input_tex : input_attachment<f32>;
+
+@fragment
+fn main() -> @location(0u) vec4<f32> {
+ return inputAttachmentLoad(input_tex);
+}
+)");
+}
+
TEST_F(IRToProgramRoundtripTest, WorkgroupSizeLargerThanI32) {
RUN_TEST(R"(
@compute @workgroup_size(4294967295u, 1u, 1u)
diff --git a/src/tint/lang/wgsl/reader/lower/lower.cc b/src/tint/lang/wgsl/reader/lower/lower.cc
index 33f1999..56d9fd3 100644
--- a/src/tint/lang/wgsl/reader/lower/lower.cc
+++ b/src/tint/lang/wgsl/reader/lower/lower.cc
@@ -167,6 +167,7 @@
CASE(kAtomicCompareExchangeWeak)
CASE(kSubgroupBallot)
CASE(kSubgroupBroadcast)
+ CASE(kInputAttachmentLoad)
case tint::wgsl::BuiltinFn::kBitcast: // should lower to ir::Bitcast
case tint::wgsl::BuiltinFn::kWorkgroupUniformLoad: // should be handled in Lower()
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
index 34d17d7..83139a4 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
@@ -1247,9 +1247,15 @@
}
current_block_->Append(val);
- if (auto* gv = sem->As<sem::GlobalVariable>(); gv && var->HasBindingPoint()) {
- val->SetBindingPoint(gv->Attributes().binding_point->group,
- gv->Attributes().binding_point->binding);
+ if (auto* gv = sem->As<sem::GlobalVariable>(); gv) {
+ if (var->HasBindingPoint()) {
+ val->SetBindingPoint(gv->Attributes().binding_point->group,
+ gv->Attributes().binding_point->binding);
+ }
+ if (var->HasInputAttachmentIndex()) {
+ val->SetInputAttachmentIndex(
+ gv->Attributes().input_attachment_index.value());
+ }
}
// Store the declaration so we can get the instruction to store too
diff --git a/src/tint/lang/wgsl/resolver/input_attachments_extension_test.cc b/src/tint/lang/wgsl/resolver/input_attachments_extension_test.cc
index 9f2a768..5387d09 100644
--- a/src/tint/lang/wgsl/resolver/input_attachments_extension_test.cc
+++ b/src/tint/lang/wgsl/resolver/input_attachments_extension_test.cc
@@ -165,5 +165,30 @@
note: '@input_attachment_index' must only be applied to declarations of 'input_attachment' type)");
}
+// Test that inputAttachmentLoad cannot be used in vertex shader
+TEST_F(InputAttachmenExtensionTest, InputAttachmentLoadInVertexStageError) {
+ // enable chromium_internal_input_attachments;
+ // @group(0) @binding(0) @input_attachment_index(3)
+ // var input_tex : input_attachment<f32>;
+ // @vertex fn f() -> @builtin(position) vec4f {
+ // return inputAttachmentLoad(input_tex);
+ // }
+
+ Enable(Source{{12, 34}}, wgsl::Extension::kChromiumInternalInputAttachments);
+
+ GlobalVar("input_tex", ty.input_attachment(ty.Of<f32>()),
+ Vector{Binding(0_u), Group(0_u), InputAttachmentIndex(3_u)});
+
+ Func("f", Empty, ty.vec4<f32>(),
+ Vector{
+ Return(Call("inputAttachmentLoad", "input_tex")),
+ },
+ Vector{Stage(ast::PipelineStage::kVertex)},
+ Vector{Builtin(core::BuiltinValue::kPosition)});
+
+ EXPECT_FALSE(r()->Resolve());
+ EXPECT_NE(r()->error().find(R"(cannot be used by vertex pipeline stage)"), std::string::npos);
+}
+
} // namespace
} // namespace tint::resolver
diff --git a/src/tint/lang/wgsl/wgsl.def b/src/tint/lang/wgsl/wgsl.def
index 2fbd877..b07d49c 100644
--- a/src/tint/lang/wgsl/wgsl.def
+++ b/src/tint/lang/wgsl/wgsl.def
@@ -581,6 +581,8 @@
@must_use fn textureLoad[C: iu32, F: i32_texel_format, R: readable](texture: texture_storage_3d<F, R>, coords: vec3<C>) -> vec4<i32>
@must_use fn textureLoad[C: iu32, F: u32_texel_format, R: readable](texture: texture_storage_3d<F, R>, coords: vec3<C>) -> vec4<u32>
+@stage("fragment") fn inputAttachmentLoad[T: fiu32](input_attachment: input_attachment<T>) -> vec4<T>
+
@stage("fragment", "compute") fn atomicLoad[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>) -> T
@stage("fragment", "compute") fn atomicStore[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, T)
@stage("fragment", "compute") fn atomicAdd[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T
diff --git a/src/tint/lang/wgsl/writer/ast_printer/ast_printer.cc b/src/tint/lang/wgsl/writer/ast_printer/ast_printer.cc
index 710c70a..44f8e4b 100644
--- a/src/tint/lang/wgsl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/wgsl/writer/ast_printer/ast_printer.cc
@@ -55,6 +55,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/internal_attribute.h"
#include "src/tint/lang/wgsl/ast/interpolate_attribute.h"
@@ -547,6 +548,11 @@
[&](const ast::StrideAttribute* stride) { out << "stride(" << stride->stride << ")"; },
[&](const ast::InternalAttribute* internal) {
out << "internal(" << internal->InternalName() << ")";
+ },
+ [&](const ast::InputAttachmentIndexAttribute* index) {
+ out << "input_attachment_index(";
+ EmitExpression(out, index->expr);
+ out << ")";
}, //
TINT_ICE_ON_NO_MATCH);
}
diff --git a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc
index 4482524..36a9ba3 100644
--- a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc
+++ b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc
@@ -72,6 +72,7 @@
#include "src/tint/lang/core/type/atomic.h"
#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/input_attachment.h"
#include "src/tint/lang/core/type/multisampled_texture.h"
#include "src/tint/lang/core/type/pointer.h"
#include "src/tint/lang/core/type/reference.h"
@@ -561,6 +562,10 @@
attrs.Push(b.Binding(u32(bp->binding)));
}
+ if (auto ii = var->InputAttachmentIndex()) {
+ attrs.Push(b.InputAttachmentIndex(u32(ii.value())));
+ }
+
const ast::Expression* init = nullptr;
if (var->Initializer()) {
init = Expr(var->Initializer());
@@ -984,6 +989,11 @@
},
[&](const core::type::Reference*) -> ast::Type {
TINT_ICE() << "reference types should never appear in the IR";
+ },
+ [&](const core::type::InputAttachment* i) {
+ Enable(wgsl::Extension::kChromiumInternalInputAttachments);
+ auto el = Type(i->type());
+ return b.ty.input_attachment(el);
}, //
TINT_ICE_ON_NO_MATCH);
}
diff --git a/src/tint/lang/wgsl/writer/raise/raise.cc b/src/tint/lang/wgsl/writer/raise/raise.cc
index d06cb72..6720392 100644
--- a/src/tint/lang/wgsl/writer/raise/raise.cc
+++ b/src/tint/lang/wgsl/writer/raise/raise.cc
@@ -170,6 +170,7 @@
CASE(kAtomicCompareExchangeWeak)
CASE(kSubgroupBallot)
CASE(kSubgroupBroadcast)
+ CASE(kInputAttachmentLoad)
case core::BuiltinFn::kNone:
break;
}
diff --git a/test/tint/builtins/gen/gen.wgsl.tmpl b/test/tint/builtins/gen/gen.wgsl.tmpl
index 7d3bd95..bc32f10 100644
--- a/test/tint/builtins/gen/gen.wgsl.tmpl
+++ b/test/tint/builtins/gen/gen.wgsl.tmpl
@@ -92,6 +92,9 @@
{{- if eq "ptr" $p.Type.Target.Name -}}
{{- $el_type := Eval "Type" (index $p.Type.TemplateArguments 1)}}
{{- if eq "handle" $class -}}
+{{- if HasPrefix $el_type "input_attachment" -}}
+ @input_attachment_index(3)
+{{ end -}}
@group(1) @binding({{$i}}) var arg_{{$i}}: {{$el_type}};
{{ $args.Put $i (printf "&arg_%v" $i) -}}
{{- else if eq "workgroup" $class -}}
@@ -104,6 +107,9 @@
{{- else -}}
{{- $type := Eval "Type" $p.Type}}
{{- if eq "handle" $class -}}
+{{- if HasPrefix $type "input_attachment" -}}
+ @input_attachment_index(3)
+{{ end -}}
@group(1) @binding({{$i}}) var arg_{{$i}}: {{$type}};
{{ $args.Put $i (printf "arg_%v" $i) -}}
{{- else if eq "workgroup" $class -}}
@@ -230,6 +236,12 @@
enable chromium_experimental_subgroups;
{{ end -}}
+
+{{- /* Emit 'enable chromium_internal_input_attachments' if required */ -}}
+{{- if eq "inputAttachmentLoad" $builtin_name}}
+enable chromium_internal_input_attachments;
+{{ end -}}
+
{{- /* Emit 'enable f16' if required */ -}}
{{- if OverloadUsesType $overload "f16"}}
enable f16;
@@ -283,6 +295,7 @@
{{- if eq $name "array" -}}storage
{{- else if HasPrefix $name "texture" -}}handle
{{- else if HasPrefix $name "sampler" -}}handle
+{{- else if HasPrefix $name "input_attachment" -}}handle
{{- else if eq $name "ptr" -}}{{(index .TemplateArguments 0).Target.Name}}
{{- else -}}function
{{- end -}}
diff --git a/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl b/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl
new file mode 100644
index 0000000..7b35e37
--- /dev/null
+++ b/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl
@@ -0,0 +1,52 @@
+// 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.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by 'tools/src/cmd/gen' using the template:
+// test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// To regenerate run: './tools/run gen'
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+enable chromium_internal_input_attachments;
+@input_attachment_index(3)
+@group(1) @binding(0) var arg_0: input_attachment<i32>;
+
+// fn inputAttachmentLoad(input_attachment: input_attachment<i32>) -> vec4<i32>
+fn inputAttachmentLoad_315bf5() {
+ var res: vec4<i32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<i32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_315bf5();
+}
diff --git a/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..a19367c
--- /dev/null
+++ b/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl.expected.dxc.hlsl
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<i32>;
+
+fn inputAttachmentLoad_315bf5() {
+ var res : vec4<i32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<i32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_315bf5();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl:38:8 error: HLSL backend does not support extension 'chromium_internal_input_attachments'
+enable chromium_internal_input_attachments;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..a19367c
--- /dev/null
+++ b/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl.expected.fxc.hlsl
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<i32>;
+
+fn inputAttachmentLoad_315bf5() {
+ var res : vec4<i32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<i32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_315bf5();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl:38:8 error: HLSL backend does not support extension 'chromium_internal_input_attachments'
+enable chromium_internal_input_attachments;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl.expected.glsl b/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl.expected.glsl
new file mode 100644
index 0000000..f9159de
--- /dev/null
+++ b/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl.expected.glsl
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<i32>;
+
+fn inputAttachmentLoad_315bf5() {
+ var res : vec4<i32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<i32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_315bf5();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl:44:44 error: unresolved value 'arg_0'
+ var res: vec4<i32> = inputAttachmentLoad(arg_0);
+ ^^^^^
+
diff --git a/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl.expected.ir.spvasm b/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl.expected.ir.spvasm
new file mode 100644
index 0000000..959770a
--- /dev/null
+++ b/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl.expected.ir.spvasm
@@ -0,0 +1,51 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 25
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %fragment_main "fragment_main"
+ OpExecutionMode %fragment_main OriginUpperLeft
+ OpName %arg_0 "arg_0"
+ OpMemberName %tint_symbol_1 0 "tint_symbol"
+ OpName %tint_symbol_1 "tint_symbol_1"
+ OpName %inputAttachmentLoad_315bf5 "inputAttachmentLoad_315bf5"
+ OpName %res "res"
+ OpName %fragment_main "fragment_main"
+ OpDecorate %arg_0 DescriptorSet 1
+ OpDecorate %arg_0 Binding 0
+ OpMemberDecorate %tint_symbol_1 0 Offset 0
+ OpDecorate %tint_symbol_1 Block
+ OpDecorate %5 DescriptorSet 2
+ OpDecorate %5 Binding 0
+ %int = OpTypeInt 32 1
+ %3 = OpTypeImage %int 2D 0 0 0 2 Unknown
+%_ptr_UniformConstant_3 = OpTypePointer UniformConstant %3
+ %arg_0 = OpVariable %_ptr_UniformConstant_3 UniformConstant
+ %v4int = OpTypeVector %int 4
+%tint_symbol_1 = OpTypeStruct %v4int
+%_ptr_StorageBuffer_tint_symbol_1 = OpTypePointer StorageBuffer %tint_symbol_1
+ %5 = OpVariable %_ptr_StorageBuffer_tint_symbol_1 StorageBuffer
+ %void = OpTypeVoid
+ %11 = OpTypeFunction %void
+ %14 = OpConstantNull %v4int
+%_ptr_Function_v4int = OpTypePointer Function %v4int
+%_ptr_StorageBuffer_v4int = OpTypePointer StorageBuffer %v4int
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%inputAttachmentLoad_315bf5 = OpFunction %void None %11
+ %12 = OpLabel
+ %res = OpVariable %_ptr_Function_v4int Function
+ %13 = OpLoad %3 %arg_0
+ OpStore %res %14
+ %17 = OpLoad %v4int %res
+ %18 = OpAccessChain %_ptr_StorageBuffer_v4int %5 %uint_0
+ OpStore %18 %17
+ OpReturn
+ OpFunctionEnd
+%fragment_main = OpFunction %void None %11
+ %23 = OpLabel
+ %24 = OpFunctionCall %void %inputAttachmentLoad_315bf5
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl.expected.msl b/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl.expected.msl
new file mode 100644
index 0000000..b780cbd
--- /dev/null
+++ b/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl.expected.msl
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<i32>;
+
+fn inputAttachmentLoad_315bf5() {
+ var res : vec4<i32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<i32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_315bf5();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl:39:2 error: '@input_attachment_index' is not valid for function parameters
+@input_attachment_index(3)
+ ^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl.expected.spvasm
new file mode 100644
index 0000000..cc98e07
--- /dev/null
+++ b/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl.expected.spvasm
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<i32>;
+
+fn inputAttachmentLoad_315bf5() {
+ var res : vec4<i32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<i32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_315bf5();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl:38:8 error: SPIR-V backend does not support extension 'chromium_internal_input_attachments'
+enable chromium_internal_input_attachments;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl.expected.wgsl
new file mode 100644
index 0000000..cce6bba
--- /dev/null
+++ b/test/tint/builtins/gen/literal/inputAttachmentLoad/315bf5.wgsl.expected.wgsl
@@ -0,0 +1,15 @@
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<i32>;
+
+fn inputAttachmentLoad_315bf5() {
+ var res : vec4<i32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<i32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_315bf5();
+}
diff --git a/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl b/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl
new file mode 100644
index 0000000..bf40179
--- /dev/null
+++ b/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl
@@ -0,0 +1,52 @@
+// 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.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by 'tools/src/cmd/gen' using the template:
+// test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// To regenerate run: './tools/run gen'
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+enable chromium_internal_input_attachments;
+@input_attachment_index(3)
+@group(1) @binding(0) var arg_0: input_attachment<f32>;
+
+// fn inputAttachmentLoad(input_attachment: input_attachment<f32>) -> vec4<f32>
+fn inputAttachmentLoad_c38b2f() {
+ var res: vec4<f32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<f32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_c38b2f();
+}
diff --git a/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..a87b277
--- /dev/null
+++ b/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl.expected.dxc.hlsl
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<f32>;
+
+fn inputAttachmentLoad_c38b2f() {
+ var res : vec4<f32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<f32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_c38b2f();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl:38:8 error: HLSL backend does not support extension 'chromium_internal_input_attachments'
+enable chromium_internal_input_attachments;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..a87b277
--- /dev/null
+++ b/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl.expected.fxc.hlsl
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<f32>;
+
+fn inputAttachmentLoad_c38b2f() {
+ var res : vec4<f32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<f32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_c38b2f();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl:38:8 error: HLSL backend does not support extension 'chromium_internal_input_attachments'
+enable chromium_internal_input_attachments;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl.expected.glsl b/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl.expected.glsl
new file mode 100644
index 0000000..33ac4ee
--- /dev/null
+++ b/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl.expected.glsl
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<f32>;
+
+fn inputAttachmentLoad_c38b2f() {
+ var res : vec4<f32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<f32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_c38b2f();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl:44:44 error: unresolved value 'arg_0'
+ var res: vec4<f32> = inputAttachmentLoad(arg_0);
+ ^^^^^
+
diff --git a/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl.expected.ir.spvasm b/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl.expected.ir.spvasm
new file mode 100644
index 0000000..85e6c37
--- /dev/null
+++ b/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl.expected.ir.spvasm
@@ -0,0 +1,51 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 25
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %fragment_main "fragment_main"
+ OpExecutionMode %fragment_main OriginUpperLeft
+ OpName %arg_0 "arg_0"
+ OpMemberName %tint_symbol_1 0 "tint_symbol"
+ OpName %tint_symbol_1 "tint_symbol_1"
+ OpName %inputAttachmentLoad_c38b2f "inputAttachmentLoad_c38b2f"
+ OpName %res "res"
+ OpName %fragment_main "fragment_main"
+ OpDecorate %arg_0 DescriptorSet 1
+ OpDecorate %arg_0 Binding 0
+ OpMemberDecorate %tint_symbol_1 0 Offset 0
+ OpDecorate %tint_symbol_1 Block
+ OpDecorate %5 DescriptorSet 2
+ OpDecorate %5 Binding 0
+ %float = OpTypeFloat 32
+ %3 = OpTypeImage %float 2D 0 0 0 2 Unknown
+%_ptr_UniformConstant_3 = OpTypePointer UniformConstant %3
+ %arg_0 = OpVariable %_ptr_UniformConstant_3 UniformConstant
+ %v4float = OpTypeVector %float 4
+%tint_symbol_1 = OpTypeStruct %v4float
+%_ptr_StorageBuffer_tint_symbol_1 = OpTypePointer StorageBuffer %tint_symbol_1
+ %5 = OpVariable %_ptr_StorageBuffer_tint_symbol_1 StorageBuffer
+ %void = OpTypeVoid
+ %11 = OpTypeFunction %void
+ %14 = OpConstantNull %v4float
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+%_ptr_StorageBuffer_v4float = OpTypePointer StorageBuffer %v4float
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%inputAttachmentLoad_c38b2f = OpFunction %void None %11
+ %12 = OpLabel
+ %res = OpVariable %_ptr_Function_v4float Function
+ %13 = OpLoad %3 %arg_0
+ OpStore %res %14
+ %17 = OpLoad %v4float %res
+ %18 = OpAccessChain %_ptr_StorageBuffer_v4float %5 %uint_0
+ OpStore %18 %17
+ OpReturn
+ OpFunctionEnd
+%fragment_main = OpFunction %void None %11
+ %23 = OpLabel
+ %24 = OpFunctionCall %void %inputAttachmentLoad_c38b2f
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl.expected.msl b/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl.expected.msl
new file mode 100644
index 0000000..24fc69f
--- /dev/null
+++ b/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl.expected.msl
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<f32>;
+
+fn inputAttachmentLoad_c38b2f() {
+ var res : vec4<f32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<f32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_c38b2f();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl:39:2 error: '@input_attachment_index' is not valid for function parameters
+@input_attachment_index(3)
+ ^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl.expected.spvasm
new file mode 100644
index 0000000..d1c2880
--- /dev/null
+++ b/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl.expected.spvasm
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<f32>;
+
+fn inputAttachmentLoad_c38b2f() {
+ var res : vec4<f32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<f32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_c38b2f();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl:38:8 error: SPIR-V backend does not support extension 'chromium_internal_input_attachments'
+enable chromium_internal_input_attachments;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl.expected.wgsl
new file mode 100644
index 0000000..d9b2d7f
--- /dev/null
+++ b/test/tint/builtins/gen/literal/inputAttachmentLoad/c38b2f.wgsl.expected.wgsl
@@ -0,0 +1,15 @@
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<f32>;
+
+fn inputAttachmentLoad_c38b2f() {
+ var res : vec4<f32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<f32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_c38b2f();
+}
diff --git a/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl b/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl
new file mode 100644
index 0000000..39cf4eb
--- /dev/null
+++ b/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl
@@ -0,0 +1,52 @@
+// 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.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by 'tools/src/cmd/gen' using the template:
+// test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// To regenerate run: './tools/run gen'
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+enable chromium_internal_input_attachments;
+@input_attachment_index(3)
+@group(1) @binding(0) var arg_0: input_attachment<u32>;
+
+// fn inputAttachmentLoad(input_attachment: input_attachment<u32>) -> vec4<u32>
+fn inputAttachmentLoad_fc4d97() {
+ var res: vec4<u32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_fc4d97();
+}
diff --git a/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..d0dc177
--- /dev/null
+++ b/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl.expected.dxc.hlsl
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<u32>;
+
+fn inputAttachmentLoad_fc4d97() {
+ var res : vec4<u32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_fc4d97();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl:38:8 error: HLSL backend does not support extension 'chromium_internal_input_attachments'
+enable chromium_internal_input_attachments;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..d0dc177
--- /dev/null
+++ b/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl.expected.fxc.hlsl
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<u32>;
+
+fn inputAttachmentLoad_fc4d97() {
+ var res : vec4<u32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_fc4d97();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl:38:8 error: HLSL backend does not support extension 'chromium_internal_input_attachments'
+enable chromium_internal_input_attachments;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl.expected.glsl b/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl.expected.glsl
new file mode 100644
index 0000000..2d23c1f
--- /dev/null
+++ b/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl.expected.glsl
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<u32>;
+
+fn inputAttachmentLoad_fc4d97() {
+ var res : vec4<u32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_fc4d97();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl:44:44 error: unresolved value 'arg_0'
+ var res: vec4<u32> = inputAttachmentLoad(arg_0);
+ ^^^^^
+
diff --git a/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl.expected.ir.spvasm b/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl.expected.ir.spvasm
new file mode 100644
index 0000000..ce0fda1
--- /dev/null
+++ b/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl.expected.ir.spvasm
@@ -0,0 +1,50 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 24
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %fragment_main "fragment_main"
+ OpExecutionMode %fragment_main OriginUpperLeft
+ OpName %arg_0 "arg_0"
+ OpMemberName %tint_symbol_1 0 "tint_symbol"
+ OpName %tint_symbol_1 "tint_symbol_1"
+ OpName %inputAttachmentLoad_fc4d97 "inputAttachmentLoad_fc4d97"
+ OpName %res "res"
+ OpName %fragment_main "fragment_main"
+ OpDecorate %arg_0 DescriptorSet 1
+ OpDecorate %arg_0 Binding 0
+ OpMemberDecorate %tint_symbol_1 0 Offset 0
+ OpDecorate %tint_symbol_1 Block
+ OpDecorate %5 DescriptorSet 2
+ OpDecorate %5 Binding 0
+ %uint = OpTypeInt 32 0
+ %3 = OpTypeImage %uint 2D 0 0 0 2 Unknown
+%_ptr_UniformConstant_3 = OpTypePointer UniformConstant %3
+ %arg_0 = OpVariable %_ptr_UniformConstant_3 UniformConstant
+ %v4uint = OpTypeVector %uint 4
+%tint_symbol_1 = OpTypeStruct %v4uint
+%_ptr_StorageBuffer_tint_symbol_1 = OpTypePointer StorageBuffer %tint_symbol_1
+ %5 = OpVariable %_ptr_StorageBuffer_tint_symbol_1 StorageBuffer
+ %void = OpTypeVoid
+ %11 = OpTypeFunction %void
+ %14 = OpConstantNull %v4uint
+%_ptr_Function_v4uint = OpTypePointer Function %v4uint
+%_ptr_StorageBuffer_v4uint = OpTypePointer StorageBuffer %v4uint
+ %uint_0 = OpConstant %uint 0
+%inputAttachmentLoad_fc4d97 = OpFunction %void None %11
+ %12 = OpLabel
+ %res = OpVariable %_ptr_Function_v4uint Function
+ %13 = OpLoad %3 %arg_0
+ OpStore %res %14
+ %17 = OpLoad %v4uint %res
+ %18 = OpAccessChain %_ptr_StorageBuffer_v4uint %5 %uint_0
+ OpStore %18 %17
+ OpReturn
+ OpFunctionEnd
+%fragment_main = OpFunction %void None %11
+ %22 = OpLabel
+ %23 = OpFunctionCall %void %inputAttachmentLoad_fc4d97
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl.expected.msl b/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl.expected.msl
new file mode 100644
index 0000000..13d004c
--- /dev/null
+++ b/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl.expected.msl
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<u32>;
+
+fn inputAttachmentLoad_fc4d97() {
+ var res : vec4<u32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_fc4d97();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl:39:2 error: '@input_attachment_index' is not valid for function parameters
+@input_attachment_index(3)
+ ^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl.expected.spvasm
new file mode 100644
index 0000000..42713c6
--- /dev/null
+++ b/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl.expected.spvasm
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<u32>;
+
+fn inputAttachmentLoad_fc4d97() {
+ var res : vec4<u32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_fc4d97();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl:38:8 error: SPIR-V backend does not support extension 'chromium_internal_input_attachments'
+enable chromium_internal_input_attachments;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl.expected.wgsl
new file mode 100644
index 0000000..a658bce
--- /dev/null
+++ b/test/tint/builtins/gen/literal/inputAttachmentLoad/fc4d97.wgsl.expected.wgsl
@@ -0,0 +1,15 @@
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<u32>;
+
+fn inputAttachmentLoad_fc4d97() {
+ var res : vec4<u32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_fc4d97();
+}
diff --git a/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl b/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl
new file mode 100644
index 0000000..7b35e37
--- /dev/null
+++ b/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl
@@ -0,0 +1,52 @@
+// 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.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by 'tools/src/cmd/gen' using the template:
+// test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// To regenerate run: './tools/run gen'
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+enable chromium_internal_input_attachments;
+@input_attachment_index(3)
+@group(1) @binding(0) var arg_0: input_attachment<i32>;
+
+// fn inputAttachmentLoad(input_attachment: input_attachment<i32>) -> vec4<i32>
+fn inputAttachmentLoad_315bf5() {
+ var res: vec4<i32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<i32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_315bf5();
+}
diff --git a/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..b1f9cfc
--- /dev/null
+++ b/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl.expected.dxc.hlsl
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<i32>;
+
+fn inputAttachmentLoad_315bf5() {
+ var res : vec4<i32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<i32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_315bf5();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl:38:8 error: HLSL backend does not support extension 'chromium_internal_input_attachments'
+enable chromium_internal_input_attachments;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..b1f9cfc
--- /dev/null
+++ b/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl.expected.fxc.hlsl
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<i32>;
+
+fn inputAttachmentLoad_315bf5() {
+ var res : vec4<i32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<i32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_315bf5();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl:38:8 error: HLSL backend does not support extension 'chromium_internal_input_attachments'
+enable chromium_internal_input_attachments;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl.expected.glsl b/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl.expected.glsl
new file mode 100644
index 0000000..d5e8b74
--- /dev/null
+++ b/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl.expected.glsl
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<i32>;
+
+fn inputAttachmentLoad_315bf5() {
+ var res : vec4<i32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<i32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_315bf5();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl:44:44 error: unresolved value 'arg_0'
+ var res: vec4<i32> = inputAttachmentLoad(arg_0);
+ ^^^^^
+
diff --git a/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl.expected.ir.spvasm b/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl.expected.ir.spvasm
new file mode 100644
index 0000000..959770a
--- /dev/null
+++ b/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl.expected.ir.spvasm
@@ -0,0 +1,51 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 25
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %fragment_main "fragment_main"
+ OpExecutionMode %fragment_main OriginUpperLeft
+ OpName %arg_0 "arg_0"
+ OpMemberName %tint_symbol_1 0 "tint_symbol"
+ OpName %tint_symbol_1 "tint_symbol_1"
+ OpName %inputAttachmentLoad_315bf5 "inputAttachmentLoad_315bf5"
+ OpName %res "res"
+ OpName %fragment_main "fragment_main"
+ OpDecorate %arg_0 DescriptorSet 1
+ OpDecorate %arg_0 Binding 0
+ OpMemberDecorate %tint_symbol_1 0 Offset 0
+ OpDecorate %tint_symbol_1 Block
+ OpDecorate %5 DescriptorSet 2
+ OpDecorate %5 Binding 0
+ %int = OpTypeInt 32 1
+ %3 = OpTypeImage %int 2D 0 0 0 2 Unknown
+%_ptr_UniformConstant_3 = OpTypePointer UniformConstant %3
+ %arg_0 = OpVariable %_ptr_UniformConstant_3 UniformConstant
+ %v4int = OpTypeVector %int 4
+%tint_symbol_1 = OpTypeStruct %v4int
+%_ptr_StorageBuffer_tint_symbol_1 = OpTypePointer StorageBuffer %tint_symbol_1
+ %5 = OpVariable %_ptr_StorageBuffer_tint_symbol_1 StorageBuffer
+ %void = OpTypeVoid
+ %11 = OpTypeFunction %void
+ %14 = OpConstantNull %v4int
+%_ptr_Function_v4int = OpTypePointer Function %v4int
+%_ptr_StorageBuffer_v4int = OpTypePointer StorageBuffer %v4int
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%inputAttachmentLoad_315bf5 = OpFunction %void None %11
+ %12 = OpLabel
+ %res = OpVariable %_ptr_Function_v4int Function
+ %13 = OpLoad %3 %arg_0
+ OpStore %res %14
+ %17 = OpLoad %v4int %res
+ %18 = OpAccessChain %_ptr_StorageBuffer_v4int %5 %uint_0
+ OpStore %18 %17
+ OpReturn
+ OpFunctionEnd
+%fragment_main = OpFunction %void None %11
+ %23 = OpLabel
+ %24 = OpFunctionCall %void %inputAttachmentLoad_315bf5
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl.expected.msl b/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl.expected.msl
new file mode 100644
index 0000000..392328f
--- /dev/null
+++ b/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl.expected.msl
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<i32>;
+
+fn inputAttachmentLoad_315bf5() {
+ var res : vec4<i32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<i32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_315bf5();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl:39:2 error: '@input_attachment_index' is not valid for function parameters
+@input_attachment_index(3)
+ ^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl.expected.spvasm b/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl.expected.spvasm
new file mode 100644
index 0000000..6eac7f9
--- /dev/null
+++ b/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl.expected.spvasm
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<i32>;
+
+fn inputAttachmentLoad_315bf5() {
+ var res : vec4<i32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<i32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_315bf5();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl:38:8 error: SPIR-V backend does not support extension 'chromium_internal_input_attachments'
+enable chromium_internal_input_attachments;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl.expected.wgsl b/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl.expected.wgsl
new file mode 100644
index 0000000..cce6bba
--- /dev/null
+++ b/test/tint/builtins/gen/var/inputAttachmentLoad/315bf5.wgsl.expected.wgsl
@@ -0,0 +1,15 @@
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<i32>;
+
+fn inputAttachmentLoad_315bf5() {
+ var res : vec4<i32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<i32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_315bf5();
+}
diff --git a/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl b/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl
new file mode 100644
index 0000000..bf40179
--- /dev/null
+++ b/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl
@@ -0,0 +1,52 @@
+// 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.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by 'tools/src/cmd/gen' using the template:
+// test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// To regenerate run: './tools/run gen'
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+enable chromium_internal_input_attachments;
+@input_attachment_index(3)
+@group(1) @binding(0) var arg_0: input_attachment<f32>;
+
+// fn inputAttachmentLoad(input_attachment: input_attachment<f32>) -> vec4<f32>
+fn inputAttachmentLoad_c38b2f() {
+ var res: vec4<f32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<f32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_c38b2f();
+}
diff --git a/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..c40fd88
--- /dev/null
+++ b/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl.expected.dxc.hlsl
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<f32>;
+
+fn inputAttachmentLoad_c38b2f() {
+ var res : vec4<f32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<f32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_c38b2f();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl:38:8 error: HLSL backend does not support extension 'chromium_internal_input_attachments'
+enable chromium_internal_input_attachments;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..c40fd88
--- /dev/null
+++ b/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl.expected.fxc.hlsl
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<f32>;
+
+fn inputAttachmentLoad_c38b2f() {
+ var res : vec4<f32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<f32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_c38b2f();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl:38:8 error: HLSL backend does not support extension 'chromium_internal_input_attachments'
+enable chromium_internal_input_attachments;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl.expected.glsl b/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl.expected.glsl
new file mode 100644
index 0000000..fabb939
--- /dev/null
+++ b/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl.expected.glsl
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<f32>;
+
+fn inputAttachmentLoad_c38b2f() {
+ var res : vec4<f32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<f32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_c38b2f();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl:44:44 error: unresolved value 'arg_0'
+ var res: vec4<f32> = inputAttachmentLoad(arg_0);
+ ^^^^^
+
diff --git a/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl.expected.ir.spvasm b/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl.expected.ir.spvasm
new file mode 100644
index 0000000..85e6c37
--- /dev/null
+++ b/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl.expected.ir.spvasm
@@ -0,0 +1,51 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 25
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %fragment_main "fragment_main"
+ OpExecutionMode %fragment_main OriginUpperLeft
+ OpName %arg_0 "arg_0"
+ OpMemberName %tint_symbol_1 0 "tint_symbol"
+ OpName %tint_symbol_1 "tint_symbol_1"
+ OpName %inputAttachmentLoad_c38b2f "inputAttachmentLoad_c38b2f"
+ OpName %res "res"
+ OpName %fragment_main "fragment_main"
+ OpDecorate %arg_0 DescriptorSet 1
+ OpDecorate %arg_0 Binding 0
+ OpMemberDecorate %tint_symbol_1 0 Offset 0
+ OpDecorate %tint_symbol_1 Block
+ OpDecorate %5 DescriptorSet 2
+ OpDecorate %5 Binding 0
+ %float = OpTypeFloat 32
+ %3 = OpTypeImage %float 2D 0 0 0 2 Unknown
+%_ptr_UniformConstant_3 = OpTypePointer UniformConstant %3
+ %arg_0 = OpVariable %_ptr_UniformConstant_3 UniformConstant
+ %v4float = OpTypeVector %float 4
+%tint_symbol_1 = OpTypeStruct %v4float
+%_ptr_StorageBuffer_tint_symbol_1 = OpTypePointer StorageBuffer %tint_symbol_1
+ %5 = OpVariable %_ptr_StorageBuffer_tint_symbol_1 StorageBuffer
+ %void = OpTypeVoid
+ %11 = OpTypeFunction %void
+ %14 = OpConstantNull %v4float
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+%_ptr_StorageBuffer_v4float = OpTypePointer StorageBuffer %v4float
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%inputAttachmentLoad_c38b2f = OpFunction %void None %11
+ %12 = OpLabel
+ %res = OpVariable %_ptr_Function_v4float Function
+ %13 = OpLoad %3 %arg_0
+ OpStore %res %14
+ %17 = OpLoad %v4float %res
+ %18 = OpAccessChain %_ptr_StorageBuffer_v4float %5 %uint_0
+ OpStore %18 %17
+ OpReturn
+ OpFunctionEnd
+%fragment_main = OpFunction %void None %11
+ %23 = OpLabel
+ %24 = OpFunctionCall %void %inputAttachmentLoad_c38b2f
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl.expected.msl b/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl.expected.msl
new file mode 100644
index 0000000..4619eae
--- /dev/null
+++ b/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl.expected.msl
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<f32>;
+
+fn inputAttachmentLoad_c38b2f() {
+ var res : vec4<f32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<f32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_c38b2f();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl:39:2 error: '@input_attachment_index' is not valid for function parameters
+@input_attachment_index(3)
+ ^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl.expected.spvasm b/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl.expected.spvasm
new file mode 100644
index 0000000..3de88ed
--- /dev/null
+++ b/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl.expected.spvasm
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<f32>;
+
+fn inputAttachmentLoad_c38b2f() {
+ var res : vec4<f32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<f32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_c38b2f();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl:38:8 error: SPIR-V backend does not support extension 'chromium_internal_input_attachments'
+enable chromium_internal_input_attachments;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl.expected.wgsl b/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl.expected.wgsl
new file mode 100644
index 0000000..d9b2d7f
--- /dev/null
+++ b/test/tint/builtins/gen/var/inputAttachmentLoad/c38b2f.wgsl.expected.wgsl
@@ -0,0 +1,15 @@
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<f32>;
+
+fn inputAttachmentLoad_c38b2f() {
+ var res : vec4<f32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<f32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_c38b2f();
+}
diff --git a/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl b/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl
new file mode 100644
index 0000000..39cf4eb
--- /dev/null
+++ b/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl
@@ -0,0 +1,52 @@
+// 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.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by 'tools/src/cmd/gen' using the template:
+// test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// To regenerate run: './tools/run gen'
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+enable chromium_internal_input_attachments;
+@input_attachment_index(3)
+@group(1) @binding(0) var arg_0: input_attachment<u32>;
+
+// fn inputAttachmentLoad(input_attachment: input_attachment<u32>) -> vec4<u32>
+fn inputAttachmentLoad_fc4d97() {
+ var res: vec4<u32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_fc4d97();
+}
diff --git a/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..ef2e2b3
--- /dev/null
+++ b/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl.expected.dxc.hlsl
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<u32>;
+
+fn inputAttachmentLoad_fc4d97() {
+ var res : vec4<u32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_fc4d97();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl:38:8 error: HLSL backend does not support extension 'chromium_internal_input_attachments'
+enable chromium_internal_input_attachments;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..ef2e2b3
--- /dev/null
+++ b/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl.expected.fxc.hlsl
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<u32>;
+
+fn inputAttachmentLoad_fc4d97() {
+ var res : vec4<u32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_fc4d97();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl:38:8 error: HLSL backend does not support extension 'chromium_internal_input_attachments'
+enable chromium_internal_input_attachments;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl.expected.glsl b/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl.expected.glsl
new file mode 100644
index 0000000..a277d54
--- /dev/null
+++ b/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl.expected.glsl
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<u32>;
+
+fn inputAttachmentLoad_fc4d97() {
+ var res : vec4<u32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_fc4d97();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl:44:44 error: unresolved value 'arg_0'
+ var res: vec4<u32> = inputAttachmentLoad(arg_0);
+ ^^^^^
+
diff --git a/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl.expected.ir.spvasm b/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl.expected.ir.spvasm
new file mode 100644
index 0000000..ce0fda1
--- /dev/null
+++ b/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl.expected.ir.spvasm
@@ -0,0 +1,50 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 24
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %fragment_main "fragment_main"
+ OpExecutionMode %fragment_main OriginUpperLeft
+ OpName %arg_0 "arg_0"
+ OpMemberName %tint_symbol_1 0 "tint_symbol"
+ OpName %tint_symbol_1 "tint_symbol_1"
+ OpName %inputAttachmentLoad_fc4d97 "inputAttachmentLoad_fc4d97"
+ OpName %res "res"
+ OpName %fragment_main "fragment_main"
+ OpDecorate %arg_0 DescriptorSet 1
+ OpDecorate %arg_0 Binding 0
+ OpMemberDecorate %tint_symbol_1 0 Offset 0
+ OpDecorate %tint_symbol_1 Block
+ OpDecorate %5 DescriptorSet 2
+ OpDecorate %5 Binding 0
+ %uint = OpTypeInt 32 0
+ %3 = OpTypeImage %uint 2D 0 0 0 2 Unknown
+%_ptr_UniformConstant_3 = OpTypePointer UniformConstant %3
+ %arg_0 = OpVariable %_ptr_UniformConstant_3 UniformConstant
+ %v4uint = OpTypeVector %uint 4
+%tint_symbol_1 = OpTypeStruct %v4uint
+%_ptr_StorageBuffer_tint_symbol_1 = OpTypePointer StorageBuffer %tint_symbol_1
+ %5 = OpVariable %_ptr_StorageBuffer_tint_symbol_1 StorageBuffer
+ %void = OpTypeVoid
+ %11 = OpTypeFunction %void
+ %14 = OpConstantNull %v4uint
+%_ptr_Function_v4uint = OpTypePointer Function %v4uint
+%_ptr_StorageBuffer_v4uint = OpTypePointer StorageBuffer %v4uint
+ %uint_0 = OpConstant %uint 0
+%inputAttachmentLoad_fc4d97 = OpFunction %void None %11
+ %12 = OpLabel
+ %res = OpVariable %_ptr_Function_v4uint Function
+ %13 = OpLoad %3 %arg_0
+ OpStore %res %14
+ %17 = OpLoad %v4uint %res
+ %18 = OpAccessChain %_ptr_StorageBuffer_v4uint %5 %uint_0
+ OpStore %18 %17
+ OpReturn
+ OpFunctionEnd
+%fragment_main = OpFunction %void None %11
+ %22 = OpLabel
+ %23 = OpFunctionCall %void %inputAttachmentLoad_fc4d97
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl.expected.msl b/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl.expected.msl
new file mode 100644
index 0000000..8917edb
--- /dev/null
+++ b/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl.expected.msl
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<u32>;
+
+fn inputAttachmentLoad_fc4d97() {
+ var res : vec4<u32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_fc4d97();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl:39:2 error: '@input_attachment_index' is not valid for function parameters
+@input_attachment_index(3)
+ ^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl.expected.spvasm b/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl.expected.spvasm
new file mode 100644
index 0000000..bd870e4
--- /dev/null
+++ b/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl.expected.spvasm
@@ -0,0 +1,23 @@
+SKIP: FAILED
+
+
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<u32>;
+
+fn inputAttachmentLoad_fc4d97() {
+ var res : vec4<u32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_fc4d97();
+}
+
+Failed to generate: <dawn>/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl:38:8 error: SPIR-V backend does not support extension 'chromium_internal_input_attachments'
+enable chromium_internal_input_attachments;
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
diff --git a/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl.expected.wgsl b/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl.expected.wgsl
new file mode 100644
index 0000000..a658bce
--- /dev/null
+++ b/test/tint/builtins/gen/var/inputAttachmentLoad/fc4d97.wgsl.expected.wgsl
@@ -0,0 +1,15 @@
+enable chromium_internal_input_attachments;
+
+@input_attachment_index(3) @group(1) @binding(0) var arg_0 : input_attachment<u32>;
+
+fn inputAttachmentLoad_fc4d97() {
+ var res : vec4<u32> = inputAttachmentLoad(arg_0);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
+
+@fragment
+fn fragment_main() {
+ inputAttachmentLoad_fc4d97();
+}
