Import Tint changes from Dawn Changes: - 548ee54a9742b6368325414911deef9e21d986ee Add missing bits of build support for libprotobuf-mutator by Ryan Harrison <rharrison@chromium.org> - cc43f15ddc5e0a25cbdc141610483b2be2b941d6 Tint: Add `dual_source_blending` as a valid WGSL extension by Jiawei Shao <jiawei.shao@intel.com> - d7dfc6c14ee7b97e4ff6dd85fd0aa7bbe708f233 Tint: Add input attachments support to spirv AST printer. by Le Hoang Quyen <lehoangquyen@chromium.org> - 10d0b65f6e0d3d7b689217fbe3cc60d08e94480b Tint: Add input attachments support to spirv IR printer. by Le Hoang Quyen <lehoangquyen@chromium.org> - c224131e83b98f9525b511fdc79a78d478b67146 Tint: Add inputAttachmentLoad to core.def, wgsl.def, IR by Le Hoang Quyen <lehoangquyen@chromium.org> - 985b4e8921edd65074abc577eb80bb88639eb415 [msl] Populate Output::workgroup_allocations by James Price <jrprice@google.com> - 7796bc02b45e2f9f54cb6bd67c6d36d6ef125d9d [tint][ast][msl] Remove the need for SingleEntryPoint to ... by Ben Clayton <bclayton@google.com> - 0be0f1bda4660667da6f5f30ac79d603c359102a Fix misc-include-cleaner warnings from clang-tidy by David Neto <dneto@google.com> - d8bde2c4c5d124ac9997ad431aac172b1963751c [tint][ir] Validate compute entry point has workgroup_size by Ben Clayton <bclayton@google.com> - b4bd417d7c423b95f05de5c3b5f1887c5226a7b4 [tint][ast][msl] Skip the MSL fuzzer if the program has m... by Ben Clayton <bclayton@google.com> GitOrigin-RevId: 548ee54a9742b6368325414911deef9e21d986ee Change-Id: I9031fdf1b7d60ec2baf3bee208950e9ee29a5df5 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/190640 Kokoro: Kokoro <noreply+kokoro@google.com> Commit-Queue: James Price <jrprice@google.com> Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/tint/cmd/fuzz/wgsl/dictionary.txt b/src/tint/cmd/fuzz/wgsl/dictionary.txt index ffd821a..a9118b3 100644 --- a/src/tint/cmd/fuzz/wgsl/dictionary.txt +++ b/src/tint/cmd/fuzz/wgsl/dictionary.txt
@@ -195,6 +195,7 @@ "dpdy" "dpdyCoarse" "dpdyFine" +"dual_source_blending" "else" "enable" "error" @@ -232,6 +233,7 @@ "id" "if" "info" +"inputAttachmentLoad" "input_attachment" "input_attachment_index" "insertBits"
diff --git a/src/tint/cmd/fuzz/wgsl/fuzz.cc b/src/tint/cmd/fuzz/wgsl/fuzz.cc index 20de4f5..b5467bc 100644 --- a/src/tint/cmd/fuzz/wgsl/fuzz.cc +++ b/src/tint/cmd/fuzz/wgsl/fuzz.cc
@@ -36,6 +36,7 @@ #include "src/tint/lang/wgsl/ast/alias.h" #include "src/tint/lang/wgsl/ast/function.h" #include "src/tint/lang/wgsl/ast/identifier.h" +#include "src/tint/lang/wgsl/ast/module.h" #include "src/tint/lang/wgsl/ast/struct.h" #include "src/tint/lang/wgsl/ast/variable.h" #include "src/tint/lang/wgsl/builtin_fn.h" @@ -104,6 +105,19 @@ break; // Early exit - nothing more to find. } } + + // Check for multiple entry points + bool entry_point_found = false; + for (auto* fn : program.AST().Functions()) { + if (fn->IsEntryPoint()) { + if (entry_point_found) { + out.Add(ProgramProperties::kMultipleEntryPoints); + break; + } + entry_point_found = true; + } + } + return out; }
diff --git a/src/tint/cmd/fuzz/wgsl/fuzz.h b/src/tint/cmd/fuzz/wgsl/fuzz.h index a24514b..aa3cbc6 100644 --- a/src/tint/cmd/fuzz/wgsl/fuzz.h +++ b/src/tint/cmd/fuzz/wgsl/fuzz.h
@@ -61,6 +61,8 @@ kBuiltinFnsShadowed, /// The program has builtin types which have been shadowed kBuiltinTypesShadowed, + /// The program has multiple entry points + kMultipleEntryPoints, }; /// Context holds information about the fuzzer options and the input program.
diff --git a/src/tint/externals.json b/src/tint/externals.json index 8b8bd96..1745609 100644 --- a/src/tint/externals.json +++ b/src/tint/externals.json
@@ -89,5 +89,10 @@ "winsock2.h" ], "Condition": "tint_build_is_win" + }, + "libprotobuf-mutator": { + "IncludePatterns": [ + "libprotobuf-mutator/**" + ] } }
diff --git a/src/tint/lang/core/builtin_fn.cc b/src/tint/lang/core/builtin_fn.cc index 82e1e37..36b542f 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: @@ -681,7 +686,7 @@ f == BuiltinFn::kTextureSampleCompareLevel || // f == BuiltinFn::kTextureSampleGrad || // f == BuiltinFn::kTextureSampleLevel || // - f == BuiltinFn::kTextureStore; + f == BuiltinFn::kTextureStore || f == BuiltinFn::kInputAttachmentLoad; } bool IsImageQuery(BuiltinFn f) {
diff --git a/src/tint/lang/core/builtin_fn.cc.tmpl b/src/tint/lang/core/builtin_fn.cc.tmpl index 7b47108..3b922b0 100644 --- a/src/tint/lang/core/builtin_fn.cc.tmpl +++ b/src/tint/lang/core/builtin_fn.cc.tmpl
@@ -65,7 +65,7 @@ f == BuiltinFn::kTextureSampleCompareLevel || // f == BuiltinFn::kTextureSampleGrad || // f == BuiltinFn::kTextureSampleLevel || // - f == BuiltinFn::kTextureStore; + f == BuiltinFn::kTextureStore || f == BuiltinFn::kInputAttachmentLoad; } bool IsImageQuery(BuiltinFn f) {
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/constant/invalid.h b/src/tint/lang/core/constant/invalid.h index e767172..df4b68a 100644 --- a/src/tint/lang/core/constant/invalid.h +++ b/src/tint/lang/core/constant/invalid.h
@@ -28,7 +28,9 @@ #ifndef SRC_TINT_LANG_CORE_CONSTANT_INVALID_H_ #define SRC_TINT_LANG_CORE_CONSTANT_INVALID_H_ +#include <variant> #include "src/tint/lang/core/constant/value.h" +#include "src/tint/lang/core/number.h" #include "src/tint/lang/core/type/invalid.h" #include "src/tint/utils/rtti/castable.h"
diff --git a/src/tint/lang/core/constant/invalid_test.cc b/src/tint/lang/core/constant/invalid_test.cc index bcbcb7a..b4f29f2 100644 --- a/src/tint/lang/core/constant/invalid_test.cc +++ b/src/tint/lang/core/constant/invalid_test.cc
@@ -30,6 +30,7 @@ #include "src/tint/lang/core/constant/helper_test.h" #include "src/tint/lang/core/constant/scalar.h" #include "src/tint/lang/core/fluent_types.h" +#include "src/tint/lang/core/type/clone_context.h" using namespace tint::core::number_suffixes; // NOLINT using namespace tint::core::fluent_types; // NOLINT
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/function.h b/src/tint/lang/core/ir/function.h index f7c7ebd..7b2976d 100644 --- a/src/tint/lang/core/ir/function.h +++ b/src/tint/lang/core/ir/function.h
@@ -90,6 +90,10 @@ /// @param z the z size void SetWorkgroupSize(uint32_t x, uint32_t y, uint32_t z) { workgroup_size_ = {x, y, z}; } + /// Sets the workgroup size + /// @param size the new size + void SetWorkgroupSize(std::array<uint32_t, 3> size) { workgroup_size_ = size; } + /// Clears the workgroup size. void ClearWorkgroupSize() { workgroup_size_ = {}; }
diff --git a/src/tint/lang/core/ir/function_param.cc b/src/tint/lang/core/ir/function_param.cc index 00d72e5..fb3acfd 100644 --- a/src/tint/lang/core/ir/function_param.cc +++ b/src/tint/lang/core/ir/function_param.cc
@@ -29,6 +29,7 @@ #include "src/tint/lang/core/ir/clone_context.h" #include "src/tint/lang/core/ir/module.h" +#include "src/tint/lang/core/type/type.h" #include "src/tint/utils/ice/ice.h" TINT_INSTANTIATE_TYPEINFO(tint::core::ir::FunctionParam);
diff --git a/src/tint/lang/core/ir/validator.cc b/src/tint/lang/core/ir/validator.cc index 2912315..f42eeb5 100644 --- a/src/tint/lang/core/ir/validator.cc +++ b/src/tint/lang/core/ir/validator.cc
@@ -781,6 +781,13 @@ scope_stack_.Add(param); } + + if (func->Stage() == Function::PipelineStage::kCompute) { + if (TINT_UNLIKELY(!func->WorkgroupSize().has_value())) { + AddError(func) << "compute entry point requires workgroup size attribute"; + } + } + if (HoldsType<type::Reference>(func->ReturnType())) { AddError(func) << "references are not permitted as return types"; }
diff --git a/src/tint/lang/core/ir/validator_test.cc b/src/tint/lang/core/ir/validator_test.cc index 9d5b0f1..39029fb 100644 --- a/src/tint/lang/core/ir/validator_test.cc +++ b/src/tint/lang/core/ir/validator_test.cc
@@ -261,6 +261,26 @@ )"); } +TEST_F(IR_ValidatorTest, Function_MissingWorkgroupSize) { + auto* f = b.Function("f", ty.void_(), Function::PipelineStage::kCompute); + b.Append(f->Block(), [&] { b.Return(f); }); + + auto res = ir::Validate(mod); + ASSERT_NE(res, Success); + EXPECT_EQ(res.Failure().reason.Str(), + R"(:1:1 error: compute entry point requires workgroup size attribute +%f = @compute func():void { +^^ + +note: # Disassembly +%f = @compute func():void { + $B1: { + ret + } +} +)"); +} + TEST_F(IR_ValidatorTest, CallToFunctionOutsideModule) { auto* f = b.Function("f", ty.void_()); auto* g = b.Function("g", ty.void_()); @@ -296,6 +316,7 @@ TEST_F(IR_ValidatorTest, CallToEntryPointFunction) { auto* f = b.Function("f", ty.void_()); auto* g = b.Function("g", ty.void_(), Function::PipelineStage::kCompute); + g->SetWorkgroupSize(1, 1, 1); b.Append(f->Block(), [&] { b.Call(g); @@ -321,7 +342,7 @@ ret } } -%g = @compute func():void { +%g = @compute @workgroup_size(1, 1, 1) func():void { $B2: { ret }
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/glsl/writer/ast_printer/ast_printer.cc b/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc index 4046eb6..1cbcee4 100644 --- a/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc +++ b/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
@@ -281,6 +281,7 @@ wgsl::Extension::kChromiumInternalDualSourceBlending, wgsl::Extension::kChromiumInternalGraphite, wgsl::Extension::kF16, + wgsl::Extension::kDualSourceBlending, })) { return false; } @@ -379,7 +380,8 @@ requires_f16_extension_ = true; } - if (enable->HasExtension(wgsl::Extension::kChromiumInternalDualSourceBlending)) { + if (enable->HasExtension(wgsl::Extension::kChromiumInternalDualSourceBlending) || + enable->HasExtension(wgsl::Extension::kDualSourceBlending)) { requires_dual_source_blending_extension_ = true; } }
diff --git a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc index 28a5f6f..2345c72 100644 --- a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc +++ b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
@@ -399,6 +399,7 @@ wgsl::Extension::kChromiumInternalDualSourceBlending, wgsl::Extension::kChromiumInternalGraphite, wgsl::Extension::kF16, + wgsl::Extension::kDualSourceBlending, })) { return false; }
diff --git a/src/tint/lang/msl/writer/BUILD.bazel b/src/tint/lang/msl/writer/BUILD.bazel index f7ccfaf..c578d9d 100644 --- a/src/tint/lang/msl/writer/BUILD.bazel +++ b/src/tint/lang/msl/writer/BUILD.bazel
@@ -104,6 +104,7 @@ "return_test.cc", "type_test.cc", "var_test.cc", + "writer_test.cc", ], deps = [ "//src/tint/api/common",
diff --git a/src/tint/lang/msl/writer/BUILD.cmake b/src/tint/lang/msl/writer/BUILD.cmake index 0236187..4aeb2ed 100644 --- a/src/tint/lang/msl/writer/BUILD.cmake +++ b/src/tint/lang/msl/writer/BUILD.cmake
@@ -116,6 +116,7 @@ lang/msl/writer/return_test.cc lang/msl/writer/type_test.cc lang/msl/writer/var_test.cc + lang/msl/writer/writer_test.cc ) tint_target_add_dependencies(tint_lang_msl_writer_test test
diff --git a/src/tint/lang/msl/writer/BUILD.gn b/src/tint/lang/msl/writer/BUILD.gn index 4fdf454..5635792 100644 --- a/src/tint/lang/msl/writer/BUILD.gn +++ b/src/tint/lang/msl/writer/BUILD.gn
@@ -106,6 +106,7 @@ "return_test.cc", "type_test.cc", "var_test.cc", + "writer_test.cc", ] deps = [ "${tint_src_dir}:gmock_and_gtest",
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc index 7a514a6..6fb4080 100644 --- a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc +++ b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
@@ -288,6 +288,7 @@ wgsl::Extension::kChromiumInternalGraphite, wgsl::Extension::kChromiumInternalRelaxedUniformLayout, wgsl::Extension::kF16, + wgsl::Extension::kDualSourceBlending, })) { return false; }
diff --git a/src/tint/lang/msl/writer/ast_raise/pixel_local.cc b/src/tint/lang/msl/writer/ast_raise/pixel_local.cc index 16d849b..63c58b2 100644 --- a/src/tint/lang/msl/writer/ast_raise/pixel_local.cc +++ b/src/tint/lang/msl/writer/ast_raise/pixel_local.cc
@@ -89,25 +89,26 @@ } } - // Find the single entry point - const sem::Function* entry_point = nullptr; + /// The pixel local struct + Hashset<const sem::Struct*, 1> pixel_local_structs; + + // Find the entry points for (auto* fn : src.AST().Functions()) { - if (fn->IsEntryPoint()) { - if (entry_point != nullptr) { - TINT_ICE() << "PixelLocal transform requires that the SingleEntryPoint " - "transform has already been run"; + if (!fn->IsEntryPoint()) { + continue; + } + + auto* entry_point = sem.Get(fn); + + // Look for a `var<pixel_local>` used by the entry point... + for (auto* global : entry_point->TransitivelyReferencedGlobals()) { + if (global->AddressSpace() != core::AddressSpace::kPixelLocal) { + continue; } - entry_point = sem.Get(fn); - // Look for a `var<pixel_local>` used by the entry point... - for (auto* global : entry_point->TransitivelyReferencedGlobals()) { - if (global->AddressSpace() != core::AddressSpace::kPixelLocal) { - continue; - } - - // Obtain struct of the pixel local. - auto* pixel_local_str = global->Type()->UnwrapRef()->As<sem::Struct>(); - + // Obtain struct of the pixel local. + auto* pixel_local_str = global->Type()->UnwrapRef()->As<sem::Struct>(); + if (pixel_local_structs.Add(pixel_local_str)) { // Add an Color attribute to each member of the pixel_local structure. for (auto* member : pixel_local_str->Members()) { ctx.InsertBack(member->Declaration()->attributes, @@ -115,12 +116,11 @@ ctx.InsertBack(member->Declaration()->attributes, b.Disable(ast::DisabledValidation::kEntryPointParameter)); } - - TransformEntryPoint(entry_point, global, pixel_local_str); - made_changes = true; - - break; // Only a single `var<pixel_local>` can be used by an entry point. } + + TransformEntryPoint(entry_point, global, pixel_local_str); + made_changes = true; + break; // Only a single `var<pixel_local>` can be used by an entry point. } } @@ -217,14 +217,6 @@ auto& member_attrs = member->Declaration()->attributes; add_member(member->Type(), ctx.Clone(member_attrs)); return_args.Push(b.MemberAccessor(call_result, ctx.Clone(member->Name()))); - if (auto* location = ast::GetAttribute<ast::LocationAttribute>(member_attrs)) { - // Remove the @location attribute from the member of the inner function's - // output structure. - // Note: This will break other entry points that share the same output - // structure, however this transform assumes that the SingleEntryPoint - // transform will have already been run. - ctx.Remove(member_attrs, location); - } } } else { // The entry point returned a non-structure
diff --git a/src/tint/lang/msl/writer/ast_raise/pixel_local.h b/src/tint/lang/msl/writer/ast_raise/pixel_local.h index 4459e59..2fa06d5 100644 --- a/src/tint/lang/msl/writer/ast_raise/pixel_local.h +++ b/src/tint/lang/msl/writer/ast_raise/pixel_local.h
@@ -48,7 +48,6 @@ /// copied to the module-scope var before calling the 'inner' function. /// * The outer function will have a new struct return type which holds both the pixel local members /// and the returned value(s) of the 'inner' function. -/// @note PixelLocal requires that the SingleEntryPoint transform has already been run class PixelLocal final : public Castable<PixelLocal, ast::transform::Transform> { public: /// Transform configuration options
diff --git a/src/tint/lang/msl/writer/ast_raise/pixel_local_test.cc b/src/tint/lang/msl/writer/ast_raise/pixel_local_test.cc index fba0af3..09776c1 100644 --- a/src/tint/lang/msl/writer/ast_raise/pixel_local_test.cc +++ b/src/tint/lang/msl/writer/ast_raise/pixel_local_test.cc
@@ -84,7 +84,7 @@ EXPECT_EQ(expect, str(got)); } -TEST_F(PixelLocalTest, UseInEntryPoint) { +TEST_F(PixelLocalTest, UsedInSingleEntryPoint) { auto* src = R"( enable chromium_experimental_pixel_local; @@ -133,6 +133,231 @@ EXPECT_EQ(expect, str(got)); } +TEST_F(PixelLocalTest, SameVarUsedInMultipleEntryPoints) { + auto* src = R"( +enable chromium_experimental_pixel_local; + +struct PixelLocal { + a : u32, +} + +var<pixel_local> P : PixelLocal; + +@fragment +fn F1() { + P.a += 10; +} + +@fragment +fn F2() { + P.a += 20; +} +)"; + + auto* expect = + R"( +enable chromium_experimental_framebuffer_fetch; + +struct F1_res { + @location(1) + output_0 : u32, +} + +@fragment +fn F1(pixel_local_1 : PixelLocal) -> F1_res { + P = pixel_local_1; + F1_inner(); + return F1_res(P.a); +} + +struct F2_res { + @location(1) + output_0 : u32, +} + +@fragment +fn F2(pixel_local_2 : PixelLocal) -> F2_res { + P = pixel_local_2; + F2_inner(); + return F2_res(P.a); +} + +struct PixelLocal { + @color(1u) @internal(disable_validation__entry_point_parameter) + a : u32, +} + +var<private> P : PixelLocal; + +fn F1_inner() { + P.a += 10; +} + +fn F2_inner() { + P.a += 20; +} +)"; + + auto got = Run<PixelLocal>(src, Bindings({{0, 1}})); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(PixelLocalTest, DifferentVarUsedInMultipleEntryPoints) { + auto* src = R"( +enable chromium_experimental_pixel_local; + +struct PixelLocal { + a : u32, +} + +var<pixel_local> P1 : PixelLocal; +var<pixel_local> P2 : PixelLocal; + +@fragment +fn F1() { + P1.a += 10; +} + +@fragment +fn F2() { + P2.a += 20; +} +)"; + + auto* expect = + R"( +enable chromium_experimental_framebuffer_fetch; + +struct F1_res { + @location(1) + output_0 : u32, +} + +@fragment +fn F1(pixel_local_1 : PixelLocal) -> F1_res { + P1 = pixel_local_1; + F1_inner(); + return F1_res(P1.a); +} + +struct F2_res { + @location(1) + output_0 : u32, +} + +@fragment +fn F2(pixel_local_2 : PixelLocal) -> F2_res { + P2 = pixel_local_2; + F2_inner(); + return F2_res(P2.a); +} + +struct PixelLocal { + @color(1u) @internal(disable_validation__entry_point_parameter) + a : u32, +} + +var<private> P1 : PixelLocal; + +var<private> P2 : PixelLocal; + +fn F1_inner() { + P1.a += 10; +} + +fn F2_inner() { + P2.a += 20; +} +)"; + + auto got = Run<PixelLocal>(src, Bindings({{0, 1}})); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(PixelLocalTest, DifferentStructUsedInMultipleEntryPoints) { + auto* src = R"( +enable chromium_experimental_pixel_local; + +struct PixelLocal1 { + a : u32, +} + +struct PixelLocal2 { + a : u32, +} + +var<pixel_local> P1 : PixelLocal1; +var<pixel_local> P2 : PixelLocal2; + +@fragment +fn F1() { + P1.a += 10; +} + +@fragment +fn F2() { + P2.a += 20; +} +)"; + + auto* expect = + R"( +enable chromium_experimental_framebuffer_fetch; + +struct F1_res { + @location(1) + output_0 : u32, +} + +@fragment +fn F1(pixel_local_1 : PixelLocal1) -> F1_res { + P1 = pixel_local_1; + F1_inner(); + return F1_res(P1.a); +} + +struct F2_res { + @location(1) + output_0 : u32, +} + +@fragment +fn F2(pixel_local_2 : PixelLocal2) -> F2_res { + P2 = pixel_local_2; + F2_inner(); + return F2_res(P2.a); +} + +struct PixelLocal1 { + @color(1u) @internal(disable_validation__entry_point_parameter) + a : u32, +} + +struct PixelLocal2 { + @color(1u) @internal(disable_validation__entry_point_parameter) + a : u32, +} + +var<private> P1 : PixelLocal1; + +var<private> P2 : PixelLocal2; + +fn F1_inner() { + P1.a += 10; +} + +fn F2_inner() { + P2.a += 20; +} +)"; + + auto got = Run<PixelLocal>(src, Bindings({{0, 1}})); + + EXPECT_EQ(expect, str(got)); +} + TEST_F(PixelLocalTest, UseInCallee) { auto* src = R"( enable chromium_experimental_pixel_local; @@ -793,7 +1018,9 @@ var<private> P : PixelLocal; struct Output { + @location(0) x : vec4f, + @location(2) y : vec4f, } @@ -808,5 +1035,87 @@ EXPECT_EQ(expect, str(got)); } +TEST_F(PixelLocalTest, OutputStructUsedByMultipleEntryPoints) { + auto* src = R"( +enable chromium_experimental_pixel_local; + +struct PixelLocal { + a : u32, + b : u32, +} + +var<pixel_local> P : PixelLocal; + +struct Output { + @location(0) x : vec4f, + @location(2) y : vec4f, +} + +@fragment +fn F1() -> Output { + P.a += 42; + return Output(vec4f(1), vec4f(9)); +} + +@fragment +fn F2() -> Output { + return Output(vec4f(1), vec4f(9)); +} +)"; + + auto* expect = + R"( +enable chromium_experimental_framebuffer_fetch; + +struct F1_res { + @location(1) + output_0 : u32, + @location(5) + output_1 : u32, + @location(0) + output_2 : vec4<f32>, + @location(2) + output_3 : vec4<f32>, +} + +@fragment +fn F1(pixel_local_1 : PixelLocal) -> F1_res { + P = pixel_local_1; + let result = F1_inner(); + return F1_res(P.a, P.b, result.x, result.y); +} + +struct PixelLocal { + @color(1u) @internal(disable_validation__entry_point_parameter) + a : u32, + @color(5u) @internal(disable_validation__entry_point_parameter) + b : u32, +} + +var<private> P : PixelLocal; + +struct Output { + @location(0) + x : vec4f, + @location(2) + y : vec4f, +} + +fn F1_inner() -> Output { + P.a += 42; + return Output(vec4f(1), vec4f(9)); +} + +@fragment +fn F2() -> Output { + return Output(vec4f(1), vec4f(9)); +} +)"; + + auto got = Run<PixelLocal>(src, Bindings({{0, 1}, {1, 5}})); + + EXPECT_EQ(expect, str(got)); +} + } // namespace } // namespace tint::msl::writer
diff --git a/src/tint/lang/msl/writer/printer/printer.cc b/src/tint/lang/msl/writer/printer/printer.cc index 0ae4046..b2de07b 100644 --- a/src/tint/lang/msl/writer/printer/printer.cc +++ b/src/tint/lang/msl/writer/printer/printer.cc
@@ -106,7 +106,7 @@ explicit Printer(core::ir::Module& module) : ir_(module) {} /// @returns the generated MSL shader - tint::Result<std::string> Generate() { + tint::Result<PrintResult> Generate() { auto valid = core::ir::ValidateAndDumpIfNeeded(ir_, "MSL writer"); if (valid != Success) { return std::move(valid.Failure()); @@ -128,10 +128,15 @@ StringStream ss; ss << preamble_buffer_.String() << std::endl << main_buffer_.String(); - return ss.str(); + result_.msl = ss.str(); + + return std::move(result_); } private: + /// The result of printing the module. + PrintResult result_; + /// Map of builtin structure to unique generated name std::unordered_map<const core::type::Struct*, std::string> builtin_struct_names_; @@ -253,6 +258,8 @@ { auto out = Line(); + auto func_name = NameOf(func); + switch (func->Stage()) { case core::ir::Function::PipelineStage::kCompute: out << "kernel "; @@ -266,11 +273,14 @@ case core::ir::Function::PipelineStage::kUndefined: break; } + if (func->Stage() != core::ir::Function::PipelineStage::kUndefined) { + result_.workgroup_allocations.insert({func_name, {}}); + } // TODO(dsinclair): Handle return type attributes EmitType(out, func->ReturnType()); - out << " " << NameOf(func) << "("; + out << " " << func_name << "("; size_t i = 0; for (auto* param : func->Params()) { @@ -328,9 +338,10 @@ out << "]]"; } + auto ptr = param->Type()->As<core::type::Pointer>(); if (auto binding_point = param->BindingPoint()) { TINT_ASSERT(binding_point->group == 0); - if (auto ptr = param->Type()->As<core::type::Pointer>()) { + if (ptr) { switch (ptr->AddressSpace()) { case core::AddressSpace::kStorage: case core::AddressSpace::kUniform: @@ -353,6 +364,12 @@ TINT_ICE_ON_NO_MATCH); } } + if (ptr && ptr->AddressSpace() == core::AddressSpace::kWorkgroup && + func->Stage() == core::ir::Function::PipelineStage::kCompute) { + auto& allocations = result_.workgroup_allocations.at(func_name); + out << " [[threadgroup(" << allocations.size() << ")]]"; + allocations.push_back(ptr->StoreType()->Size()); + } } out << ") {"; @@ -1502,8 +1519,16 @@ } // namespace -Result<std::string> Print(core::ir::Module& module) { +Result<PrintResult> Print(core::ir::Module& module) { return Printer{module}.Generate(); } +PrintResult::PrintResult() = default; + +PrintResult::~PrintResult() = default; + +PrintResult::PrintResult(const PrintResult&) = default; + +PrintResult& PrintResult::operator=(const PrintResult&) = default; + } // namespace tint::msl::writer
diff --git a/src/tint/lang/msl/writer/printer/printer.h b/src/tint/lang/msl/writer/printer/printer.h index a5218c3..bba5c8c 100644 --- a/src/tint/lang/msl/writer/printer/printer.h +++ b/src/tint/lang/msl/writer/printer/printer.h
@@ -29,6 +29,8 @@ #define SRC_TINT_LANG_MSL_WRITER_PRINTER_PRINTER_H_ #include <string> +#include <unordered_map> +#include <vector> #include "src/tint/utils/result/result.h" @@ -39,9 +41,33 @@ namespace tint::msl::writer { -/// @returns the generated MSL shader on success, or failure +/// The output produced when printing MSL. +struct PrintResult { + /// Constructor + PrintResult(); + + /// Destructor + ~PrintResult(); + + /// Copy constructor + PrintResult(const PrintResult&); + + /// Copy assignment + /// @returns this + PrintResult& operator=(const PrintResult&); + + /// The generated MSL. + std::string msl = ""; + + /// A map from entry point name to a list of dynamic workgroup allocations. + /// Each element of the vector is the size of the workgroup allocation that should be created + /// for that index. + std::unordered_map<std::string, std::vector<uint32_t>> workgroup_allocations; +}; + /// @param module the Tint IR module to generate -Result<std::string> Print(core::ir::Module& module); +/// @returns the result of printing the MSL shader on success, or failure +Result<PrintResult> Print(core::ir::Module& module); } // namespace tint::msl::writer
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc index 915307b..bf8a720 100644 --- a/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc +++ b/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
@@ -46,13 +46,14 @@ TEST_F(MslWriter_BuiltinPolyfillTest, WorkgroupBarrier) { auto* func = b.Function("foo", ty.void_()); func->SetStage(core::ir::Function::PipelineStage::kCompute); + func->SetWorkgroupSize(1, 1, 1); b.Append(func->Block(), [&] { b.Call(ty.void_(), core::BuiltinFn::kWorkgroupBarrier); b.Return(func); }); auto* src = R"( -%foo = @compute func():void { +%foo = @compute @workgroup_size(1, 1, 1) func():void { $B1: { %2:void = workgroupBarrier ret @@ -62,7 +63,7 @@ EXPECT_EQ(src, str()); auto* expect = R"( -%foo = @compute func():void { +%foo = @compute @workgroup_size(1, 1, 1) func():void { $B1: { %2:void = msl.threadgroup_barrier 4u ret @@ -78,13 +79,14 @@ TEST_F(MslWriter_BuiltinPolyfillTest, StorageBarrier) { auto* func = b.Function("foo", ty.void_()); func->SetStage(core::ir::Function::PipelineStage::kCompute); + func->SetWorkgroupSize(1, 1, 1); b.Append(func->Block(), [&] { b.Call(ty.void_(), core::BuiltinFn::kStorageBarrier); b.Return(func); }); auto* src = R"( -%foo = @compute func():void { +%foo = @compute @workgroup_size(1, 1, 1) func():void { $B1: { %2:void = storageBarrier ret @@ -94,7 +96,7 @@ EXPECT_EQ(src, str()); auto* expect = R"( -%foo = @compute func():void { +%foo = @compute @workgroup_size(1, 1, 1) func():void { $B1: { %2:void = msl.threadgroup_barrier 1u ret @@ -110,13 +112,14 @@ TEST_F(MslWriter_BuiltinPolyfillTest, TextureBarrier) { auto* func = b.Function("foo", ty.void_()); func->SetStage(core::ir::Function::PipelineStage::kCompute); + func->SetWorkgroupSize(1, 1, 1); b.Append(func->Block(), [&] { b.Call(ty.void_(), core::BuiltinFn::kTextureBarrier); b.Return(func); }); auto* src = R"( -%foo = @compute func():void { +%foo = @compute @workgroup_size(1, 1, 1) func():void { $B1: { %2:void = textureBarrier ret @@ -126,7 +129,7 @@ EXPECT_EQ(src, str()); auto* expect = R"( -%foo = @compute func():void { +%foo = @compute @workgroup_size(1, 1, 1) func():void { $B1: { %2:void = msl.threadgroup_barrier 2u ret
diff --git a/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc b/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc index cf3fe76..2892d60 100644 --- a/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc +++ b/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc
@@ -29,6 +29,7 @@ #include <utility> +#include "src/tint/lang/core/fluent_types.h" #include "src/tint/lang/core/ir/transform/helper_test.h" #include "src/tint/lang/core/type/sampled_texture.h"
diff --git a/src/tint/lang/msl/writer/writer.cc b/src/tint/lang/msl/writer/writer.cc index c0f9d41..af54c6f 100644 --- a/src/tint/lang/msl/writer/writer.cc +++ b/src/tint/lang/msl/writer/writer.cc
@@ -62,7 +62,11 @@ if (result != Success) { return result.Failure(); } - output.msl = result.Get(); + output.msl = result->msl; + output.workgroup_allocations = std::move(result->workgroup_allocations); + // TODO(crbug.com/42251016): Set has_invariant. + // TODO(crbug.com/42251016): Set needs_storage_buffer_sizes. + // TODO(crbug.com/42251016): Set used_array_length_from_uniform_indices. return output; }
diff --git a/src/tint/lang/msl/writer/writer_ast_fuzz.cc b/src/tint/lang/msl/writer/writer_ast_fuzz.cc index 937a4a8..1c6e09a 100644 --- a/src/tint/lang/msl/writer/writer_ast_fuzz.cc +++ b/src/tint/lang/msl/writer/writer_ast_fuzz.cc
@@ -35,8 +35,9 @@ namespace tint::msl::writer { namespace { -void ASTFuzzer(const tint::Program& program, Options options) { +void ASTFuzzer(const tint::Program& program, const fuzz::wgsl::Context& context, Options options) { if (program.AST().HasOverrides()) { + // MSL writer assumes the SubstituteOverride and SingleEntryPoint transforms have been run return; }
diff --git a/src/tint/lang/msl/writer/writer_test.cc b/src/tint/lang/msl/writer/writer_test.cc new file mode 100644 index 0000000..52896f7 --- /dev/null +++ b/src/tint/lang/msl/writer/writer_test.cc
@@ -0,0 +1,90 @@ +// Copyright 2024 The Dawn & Tint Authors +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// +// 1. Redistributions of source code must retain the above copyright notice, this +// list of conditions and the following disclaimer. +// +// 2. Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// 3. Neither the name of the copyright holder nor the names of its +// contributors may be used to endorse or promote products derived from +// this software without specific prior written permission. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE +// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE +// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL +// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR +// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER +// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, +// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + +#include "src/tint/lang/msl/writer/helper_test.h" + +#include "gmock/gmock.h" + +namespace tint::msl::writer { +namespace { + +using namespace tint::core::fluent_types; // NOLINT +using namespace tint::core::number_suffixes; // NOLINT + +TEST_F(MslWriterTest, WorkgroupAllocations) { + auto* var_a = b.Var("a", ty.ptr<workgroup, i32>()); + auto* var_b = b.Var("b", ty.ptr<workgroup, i32>()); + mod.root_block->Append(var_a); + mod.root_block->Append(var_b); + + auto* foo = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kCompute, + std::array<uint32_t, 3>{1u, 1u, 1u}); + b.Append(foo->Block(), [&] { + auto* load_a = b.Load(var_a); + auto* load_b = b.Load(var_b); + b.Store(var_a, b.Add<i32>(load_a, load_b)); + b.Return(foo); + }); + + // No allocations, but still needs an entry in the map. + auto* bar = b.Function("bar", ty.void_(), core::ir::Function::PipelineStage::kCompute, + std::array<uint32_t, 3>{1u, 1u, 1u}); + b.Append(bar->Block(), [&] { b.Return(bar); }); + + ASSERT_TRUE(Generate()) << err_ << output_.msl; + EXPECT_EQ(output_.msl, R"(#include <metal_stdlib> +using namespace metal; +struct tint_symbol_2 { + int tint_symbol; + int tint_symbol_1; +}; +struct tint_module_vars_struct { + threadgroup int* a; + threadgroup int* b; +}; + +kernel void foo(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v [[threadgroup(0)]]) { + tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.a=(&(*v).tint_symbol), .b=(&(*v).tint_symbol_1)}; + if ((tint_local_index == 0u)) { + (*tint_module_vars.a) = 0; + (*tint_module_vars.b) = 0; + } + threadgroup_barrier(mem_flags::mem_threadgroup); + (*tint_module_vars.a) = ((*tint_module_vars.a) + (*tint_module_vars.b)); +} +kernel void bar() { +} +)"); + ASSERT_EQ(output_.workgroup_allocations.size(), 2u); + ASSERT_EQ(output_.workgroup_allocations.count("foo"), 1u); + ASSERT_EQ(output_.workgroup_allocations.count("bar"), 1u); + EXPECT_THAT(output_.workgroup_allocations.at("foo"), testing::ElementsAre(8u)); + EXPECT_THAT(output_.workgroup_allocations.at("bar"), testing::ElementsAre()); +} + +} // namespace +} // namespace tint::msl::writer
diff --git a/src/tint/lang/spirv/intrinsic/data.cc b/src/tint/lang/spirv/intrinsic/data.cc index b00c77c..5b40302 100644 --- a/src/tint/lang/spirv/intrinsic/data.cc +++ b/src/tint/lang/spirv/intrinsic/data.cc
@@ -837,6 +837,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 struct_with_runtime_array' constexpr TypeMatcher kStructWithRuntimeArrayMatcher { /* match */ [](MatchState& state, const Type* ty) -> const Type* { @@ -1149,13 +1169,14 @@ /* [40] */ kTextureStorage2DArrayMatcher, /* [41] */ kTextureStorage3DMatcher, /* [42] */ kPtrMatcher, - /* [43] */ kStructWithRuntimeArrayMatcher, - /* [44] */ kSampledImageMatcher, - /* [45] */ kF32F16Matcher, - /* [46] */ kIu32Matcher, - /* [47] */ kFiu32Matcher, - /* [48] */ kScalarMatcher, - /* [49] */ kSamplersMatcher, + /* [43] */ kInputAttachmentMatcher, + /* [44] */ kStructWithRuntimeArrayMatcher, + /* [45] */ kSampledImageMatcher, + /* [46] */ kF32F16Matcher, + /* [47] */ kIu32Matcher, + /* [48] */ kFiu32Matcher, + /* [49] */ kScalarMatcher, + /* [50] */ kSamplersMatcher, }; /// The template numbers, and number matchers @@ -1182,7 +1203,7 @@ /* [4] */ MatcherIndex(4), /* [5] */ MatcherIndex(42), /* [6] */ MatcherIndex(5), - /* [7] */ MatcherIndex(43), + /* [7] */ MatcherIndex(44), /* [8] */ MatcherIndex(1), /* [9] */ MatcherIndex(22), /* [10] */ MatcherIndex(2), @@ -1203,22 +1224,22 @@ /* [25] */ MatcherIndex(21), /* [26] */ MatcherIndex(0), /* [27] */ MatcherIndex(1), - /* [28] */ MatcherIndex(44), + /* [28] */ MatcherIndex(45), /* [29] */ MatcherIndex(26), /* [30] */ MatcherIndex(0), - /* [31] */ MatcherIndex(44), + /* [31] */ MatcherIndex(45), /* [32] */ MatcherIndex(27), /* [33] */ MatcherIndex(0), - /* [34] */ MatcherIndex(44), + /* [34] */ MatcherIndex(45), /* [35] */ MatcherIndex(28), /* [36] */ MatcherIndex(0), - /* [37] */ MatcherIndex(44), + /* [37] */ MatcherIndex(45), /* [38] */ MatcherIndex(29), /* [39] */ MatcherIndex(0), - /* [40] */ MatcherIndex(44), + /* [40] */ MatcherIndex(45), /* [41] */ MatcherIndex(30), /* [42] */ MatcherIndex(0), - /* [43] */ MatcherIndex(44), + /* [43] */ MatcherIndex(45), /* [44] */ MatcherIndex(31), /* [45] */ MatcherIndex(0), /* [46] */ MatcherIndex(38), @@ -1282,19 +1303,19 @@ /* [104] */ MatcherIndex(21), /* [105] */ MatcherIndex(0), /* [106] */ MatcherIndex(4), - /* [107] */ MatcherIndex(44), + /* [107] */ MatcherIndex(45), /* [108] */ MatcherIndex(33), /* [109] */ MatcherIndex(9), /* [110] */ MatcherIndex(5), /* [111] */ MatcherIndex(9), /* [112] */ MatcherIndex(2), - /* [113] */ MatcherIndex(44), + /* [113] */ MatcherIndex(45), /* [114] */ MatcherIndex(34), /* [115] */ MatcherIndex(10), /* [116] */ MatcherIndex(5), - /* [117] */ MatcherIndex(44), + /* [117] */ MatcherIndex(45), /* [118] */ MatcherIndex(35), - /* [119] */ MatcherIndex(44), + /* [119] */ MatcherIndex(45), /* [120] */ MatcherIndex(36), /* [121] */ MatcherIndex(9), /* [122] */ MatcherIndex(1), @@ -1312,17 +1333,19 @@ /* [134] */ MatcherIndex(8), /* [135] */ MatcherIndex(10), /* [136] */ MatcherIndex(2), - /* [137] */ MatcherIndex(9), + /* [137] */ MatcherIndex(43), /* [138] */ MatcherIndex(0), - /* [139] */ MatcherIndex(10), + /* [139] */ MatcherIndex(9), /* [140] */ MatcherIndex(0), - /* [141] */ MatcherIndex(46), - /* [142] */ MatcherIndex(6), - /* [143] */ MatcherIndex(45), - /* [144] */ MatcherIndex(47), - /* [145] */ MatcherIndex(37), - /* [146] */ MatcherIndex(49), - /* [147] */ MatcherIndex(48), + /* [141] */ MatcherIndex(10), + /* [142] */ MatcherIndex(0), + /* [143] */ MatcherIndex(47), + /* [144] */ MatcherIndex(6), + /* [145] */ MatcherIndex(46), + /* [146] */ MatcherIndex(48), + /* [147] */ MatcherIndex(37), + /* [148] */ MatcherIndex(50), + /* [149] */ MatcherIndex(49), }; static_assert(MatcherIndicesIndex::CanIndex(kMatcherIndices), @@ -2112,7 +2135,7 @@ { /* [156] */ /* usage */ core::ParameterUsage::kNone, - /* matcher_indices */ MatcherIndicesIndex(145), + /* matcher_indices */ MatcherIndicesIndex(147), }, { /* [157] */ @@ -2457,7 +2480,7 @@ { /* [225] */ /* usage */ core::ParameterUsage::kNone, - /* matcher_indices */ MatcherIndicesIndex(137), + /* matcher_indices */ MatcherIndicesIndex(139), }, { /* [226] */ @@ -2477,7 +2500,7 @@ { /* [229] */ /* usage */ core::ParameterUsage::kNone, - /* matcher_indices */ MatcherIndicesIndex(137), + /* matcher_indices */ MatcherIndicesIndex(139), }, { /* [230] */ @@ -2497,7 +2520,7 @@ { /* [233] */ /* usage */ core::ParameterUsage::kNone, - /* matcher_indices */ MatcherIndicesIndex(137), + /* matcher_indices */ MatcherIndicesIndex(139), }, { /* [234] */ @@ -2517,7 +2540,7 @@ { /* [237] */ /* usage */ core::ParameterUsage::kNone, - /* matcher_indices */ MatcherIndicesIndex(139), + /* matcher_indices */ MatcherIndicesIndex(141), }, { /* [238] */ @@ -2537,7 +2560,7 @@ { /* [241] */ /* usage */ core::ParameterUsage::kNone, - /* matcher_indices */ MatcherIndicesIndex(139), + /* matcher_indices */ MatcherIndicesIndex(141), }, { /* [242] */ @@ -2557,7 +2580,7 @@ { /* [245] */ /* usage */ core::ParameterUsage::kNone, - /* matcher_indices */ MatcherIndicesIndex(139), + /* matcher_indices */ MatcherIndicesIndex(141), }, { /* [246] */ @@ -2577,7 +2600,7 @@ { /* [249] */ /* usage */ core::ParameterUsage::kNone, - /* matcher_indices */ MatcherIndicesIndex(139), + /* matcher_indices */ MatcherIndicesIndex(141), }, { /* [250] */ @@ -2597,7 +2620,7 @@ { /* [253] */ /* usage */ core::ParameterUsage::kNone, - /* matcher_indices */ MatcherIndicesIndex(139), + /* matcher_indices */ MatcherIndicesIndex(141), }, { /* [254] */ @@ -2617,7 +2640,7 @@ { /* [257] */ /* usage */ core::ParameterUsage::kNone, - /* matcher_indices */ MatcherIndicesIndex(139), + /* matcher_indices */ MatcherIndicesIndex(141), }, { /* [258] */ @@ -2851,23 +2874,23 @@ }, { /* [304] */ - /* usage */ core::ParameterUsage::kNone, - /* matcher_indices */ MatcherIndicesIndex(13), + /* usage */ core::ParameterUsage::kInputAttachment, + /* matcher_indices */ MatcherIndicesIndex(137), }, { /* [305] */ /* usage */ core::ParameterUsage::kNone, - /* matcher_indices */ MatcherIndicesIndex(17), + /* matcher_indices */ MatcherIndicesIndex(121), }, { /* [306] */ /* usage */ core::ParameterUsage::kNone, - /* matcher_indices */ MatcherIndicesIndex(21), + /* matcher_indices */ MatcherIndicesIndex(13), }, { /* [307] */ /* usage */ core::ParameterUsage::kNone, - /* matcher_indices */ MatcherIndicesIndex(3), + /* matcher_indices */ MatcherIndicesIndex(17), }, { /* [308] */ @@ -2877,12 +2900,12 @@ { /* [309] */ /* usage */ core::ParameterUsage::kNone, - /* matcher_indices */ MatcherIndicesIndex(101), + /* matcher_indices */ MatcherIndicesIndex(3), }, { /* [310] */ /* usage */ core::ParameterUsage::kNone, - /* matcher_indices */ MatcherIndicesIndex(17), + /* matcher_indices */ MatcherIndicesIndex(21), }, { /* [311] */ @@ -2892,6 +2915,16 @@ { /* [312] */ /* usage */ core::ParameterUsage::kNone, + /* matcher_indices */ MatcherIndicesIndex(17), + }, + { + /* [313] */ + /* usage */ core::ParameterUsage::kNone, + /* matcher_indices */ MatcherIndicesIndex(101), + }, + { + /* [314] */ + /* usage */ core::ParameterUsage::kNone, /* matcher_indices */ MatcherIndicesIndex(3), }, }; @@ -2903,55 +2936,55 @@ { /* [0] */ /* name */ "T", - /* matcher_indices */ MatcherIndicesIndex(144), + /* matcher_indices */ MatcherIndicesIndex(146), /* kind */ TemplateInfo::Kind::kType, }, { /* [1] */ /* name */ "C", - /* matcher_indices */ MatcherIndicesIndex(141), + /* matcher_indices */ MatcherIndicesIndex(143), /* kind */ TemplateInfo::Kind::kType, }, { /* [2] */ /* name */ "I", - /* matcher_indices */ MatcherIndicesIndex(141), + /* matcher_indices */ MatcherIndicesIndex(143), /* kind */ TemplateInfo::Kind::kType, }, { /* [3] */ /* name */ "S", - /* matcher_indices */ MatcherIndicesIndex(141), + /* matcher_indices */ MatcherIndicesIndex(143), /* kind */ TemplateInfo::Kind::kType, }, { /* [4] */ /* name */ "T", - /* matcher_indices */ MatcherIndicesIndex(144), + /* matcher_indices */ MatcherIndicesIndex(146), /* kind */ TemplateInfo::Kind::kType, }, { /* [5] */ /* name */ "A", - /* matcher_indices */ MatcherIndicesIndex(141), + /* matcher_indices */ MatcherIndicesIndex(143), /* kind */ TemplateInfo::Kind::kType, }, { /* [6] */ /* name */ "B", - /* matcher_indices */ MatcherIndicesIndex(141), + /* matcher_indices */ MatcherIndicesIndex(143), /* kind */ TemplateInfo::Kind::kType, }, { /* [7] */ /* name */ "C", - /* matcher_indices */ MatcherIndicesIndex(141), + /* matcher_indices */ MatcherIndicesIndex(143), /* kind */ TemplateInfo::Kind::kType, }, { /* [8] */ /* name */ "D", - /* matcher_indices */ MatcherIndicesIndex(141), + /* matcher_indices */ MatcherIndicesIndex(143), /* kind */ TemplateInfo::Kind::kType, }, { @@ -2969,13 +3002,13 @@ { /* [11] */ /* name */ "C", - /* matcher_indices */ MatcherIndicesIndex(141), + /* matcher_indices */ MatcherIndicesIndex(143), /* kind */ TemplateInfo::Kind::kType, }, { /* [12] */ /* name */ "S", - /* matcher_indices */ MatcherIndicesIndex(141), + /* matcher_indices */ MatcherIndicesIndex(143), /* kind */ TemplateInfo::Kind::kType, }, { @@ -2993,13 +3026,13 @@ { /* [15] */ /* name */ "C", - /* matcher_indices */ MatcherIndicesIndex(141), + /* matcher_indices */ MatcherIndicesIndex(143), /* kind */ TemplateInfo::Kind::kType, }, { /* [16] */ /* name */ "S", - /* matcher_indices */ MatcherIndicesIndex(141), + /* matcher_indices */ MatcherIndicesIndex(143), /* kind */ TemplateInfo::Kind::kType, }, { @@ -3017,19 +3050,19 @@ { /* [19] */ /* name */ "C", - /* matcher_indices */ MatcherIndicesIndex(141), + /* matcher_indices */ MatcherIndicesIndex(143), /* kind */ TemplateInfo::Kind::kType, }, { /* [20] */ /* name */ "S", - /* matcher_indices */ MatcherIndicesIndex(141), + /* matcher_indices */ MatcherIndicesIndex(143), /* kind */ TemplateInfo::Kind::kType, }, { /* [21] */ /* name */ "T", - /* matcher_indices */ MatcherIndicesIndex(143), + /* matcher_indices */ MatcherIndicesIndex(145), /* kind */ TemplateInfo::Kind::kType, }, { @@ -3053,7 +3086,7 @@ { /* [25] */ /* name */ "T", - /* matcher_indices */ MatcherIndicesIndex(141), + /* matcher_indices */ MatcherIndicesIndex(143), /* kind */ TemplateInfo::Kind::kType, }, { @@ -3065,7 +3098,7 @@ { /* [27] */ /* name */ "S", - /* matcher_indices */ MatcherIndicesIndex(142), + /* matcher_indices */ MatcherIndicesIndex(144), /* kind */ TemplateInfo::Kind::kNumber, }, { @@ -3077,55 +3110,55 @@ { /* [29] */ /* name */ "B", - /* matcher_indices */ MatcherIndicesIndex(141), + /* matcher_indices */ MatcherIndicesIndex(143), /* kind */ TemplateInfo::Kind::kType, }, { /* [30] */ /* name */ "C", - /* matcher_indices */ MatcherIndicesIndex(141), + /* matcher_indices */ MatcherIndicesIndex(143), /* kind */ TemplateInfo::Kind::kType, }, { /* [31] */ /* name */ "I", - /* matcher_indices */ MatcherIndicesIndex(141), + /* matcher_indices */ MatcherIndicesIndex(143), /* kind */ TemplateInfo::Kind::kType, }, { /* [32] */ /* name */ "C", - /* matcher_indices */ MatcherIndicesIndex(141), + /* matcher_indices */ MatcherIndicesIndex(143), /* kind */ TemplateInfo::Kind::kType, }, { /* [33] */ /* name */ "S", - /* matcher_indices */ MatcherIndicesIndex(141), + /* matcher_indices */ MatcherIndicesIndex(143), /* kind */ TemplateInfo::Kind::kType, }, { /* [34] */ /* name */ "T", - /* matcher_indices */ MatcherIndicesIndex(144), + /* matcher_indices */ MatcherIndicesIndex(146), /* kind */ TemplateInfo::Kind::kType, }, { /* [35] */ /* name */ "C", - /* matcher_indices */ MatcherIndicesIndex(141), + /* matcher_indices */ MatcherIndicesIndex(143), /* kind */ TemplateInfo::Kind::kType, }, { /* [36] */ /* name */ "D", - /* matcher_indices */ MatcherIndicesIndex(141), + /* matcher_indices */ MatcherIndicesIndex(143), /* kind */ TemplateInfo::Kind::kType, }, { /* [37] */ /* name */ "T", - /* matcher_indices */ MatcherIndicesIndex(143), + /* matcher_indices */ MatcherIndicesIndex(145), /* kind */ TemplateInfo::Kind::kType, }, { @@ -3161,7 +3194,7 @@ { /* [43] */ /* name */ "T", - /* matcher_indices */ MatcherIndicesIndex(143), + /* matcher_indices */ MatcherIndicesIndex(145), /* kind */ TemplateInfo::Kind::kType, }, { @@ -3179,13 +3212,13 @@ { /* [46] */ /* name */ "T", - /* matcher_indices */ MatcherIndicesIndex(144), + /* matcher_indices */ MatcherIndicesIndex(146), /* kind */ TemplateInfo::Kind::kType, }, { /* [47] */ /* name */ "S", - /* matcher_indices */ MatcherIndicesIndex(146), + /* matcher_indices */ MatcherIndicesIndex(148), /* kind */ TemplateInfo::Kind::kType, }, { @@ -3197,7 +3230,7 @@ { /* [49] */ /* name */ "T", - /* matcher_indices */ MatcherIndicesIndex(147), + /* matcher_indices */ MatcherIndicesIndex(149), /* kind */ TemplateInfo::Kind::kType, }, }; @@ -4177,12 +4210,12 @@ { /* [88] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), - /* num_parameters */ 4, + /* num_parameters */ 2, /* num_explicit_templates */ 0, /* num_templates */ 2, - /* templates */ TemplateIndex(7), - /* parameters */ ParameterIndex(212), - /* return_matcher_indices */ MatcherIndicesIndex(/* invalid */), + /* templates */ TemplateIndex(0), + /* parameters */ ParameterIndex(304), + /* return_matcher_indices */ MatcherIndicesIndex(64), /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { @@ -4192,7 +4225,7 @@ /* num_explicit_templates */ 0, /* num_templates */ 2, /* templates */ TemplateIndex(7), - /* parameters */ ParameterIndex(216), + /* parameters */ ParameterIndex(212), /* return_matcher_indices */ MatcherIndicesIndex(/* invalid */), /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, @@ -4203,7 +4236,7 @@ /* num_explicit_templates */ 0, /* num_templates */ 2, /* templates */ TemplateIndex(7), - /* parameters */ ParameterIndex(220), + /* parameters */ ParameterIndex(216), /* return_matcher_indices */ MatcherIndicesIndex(/* invalid */), /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, @@ -4214,7 +4247,7 @@ /* num_explicit_templates */ 0, /* num_templates */ 2, /* templates */ TemplateIndex(7), - /* parameters */ ParameterIndex(224), + /* parameters */ ParameterIndex(220), /* return_matcher_indices */ MatcherIndicesIndex(/* invalid */), /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, @@ -4225,7 +4258,7 @@ /* num_explicit_templates */ 0, /* num_templates */ 2, /* templates */ TemplateIndex(7), - /* parameters */ ParameterIndex(228), + /* parameters */ ParameterIndex(224), /* return_matcher_indices */ MatcherIndicesIndex(/* invalid */), /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, @@ -4236,7 +4269,7 @@ /* num_explicit_templates */ 0, /* num_templates */ 2, /* templates */ TemplateIndex(7), - /* parameters */ ParameterIndex(232), + /* parameters */ ParameterIndex(228), /* return_matcher_indices */ MatcherIndicesIndex(/* invalid */), /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, @@ -4247,7 +4280,7 @@ /* num_explicit_templates */ 0, /* num_templates */ 2, /* templates */ TemplateIndex(7), - /* parameters */ ParameterIndex(236), + /* parameters */ ParameterIndex(232), /* return_matcher_indices */ MatcherIndicesIndex(/* invalid */), /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, @@ -4258,7 +4291,7 @@ /* num_explicit_templates */ 0, /* num_templates */ 2, /* templates */ TemplateIndex(7), - /* parameters */ ParameterIndex(240), + /* parameters */ ParameterIndex(236), /* return_matcher_indices */ MatcherIndicesIndex(/* invalid */), /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, @@ -4269,7 +4302,7 @@ /* num_explicit_templates */ 0, /* num_templates */ 2, /* templates */ TemplateIndex(7), - /* parameters */ ParameterIndex(244), + /* parameters */ ParameterIndex(240), /* return_matcher_indices */ MatcherIndicesIndex(/* invalid */), /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, @@ -4280,7 +4313,7 @@ /* num_explicit_templates */ 0, /* num_templates */ 2, /* templates */ TemplateIndex(7), - /* parameters */ ParameterIndex(248), + /* parameters */ ParameterIndex(244), /* return_matcher_indices */ MatcherIndicesIndex(/* invalid */), /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, @@ -4291,7 +4324,7 @@ /* num_explicit_templates */ 0, /* num_templates */ 2, /* templates */ TemplateIndex(7), - /* parameters */ ParameterIndex(252), + /* parameters */ ParameterIndex(248), /* return_matcher_indices */ MatcherIndicesIndex(/* invalid */), /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, @@ -4302,13 +4335,24 @@ /* num_explicit_templates */ 0, /* num_templates */ 2, /* templates */ TemplateIndex(7), - /* parameters */ ParameterIndex(256), + /* parameters */ ParameterIndex(252), /* return_matcher_indices */ MatcherIndicesIndex(/* invalid */), /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { /* [100] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), + /* num_parameters */ 4, + /* num_explicit_templates */ 0, + /* num_templates */ 2, + /* templates */ TemplateIndex(7), + /* parameters */ ParameterIndex(256), + /* return_matcher_indices */ MatcherIndicesIndex(/* invalid */), + /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), + }, + { + /* [101] */ + /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 2, /* num_explicit_templates */ 0, /* num_templates */ 2, @@ -4318,7 +4362,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [101] */ + /* [102] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 2, /* num_explicit_templates */ 0, @@ -4329,24 +4373,13 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [102] */ - /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), - /* num_parameters */ 2, - /* num_explicit_templates */ 0, - /* num_templates */ 2, - /* templates */ TemplateIndex(4), - /* parameters */ ParameterIndex(288), - /* return_matcher_indices */ MatcherIndicesIndex(133), - /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), - }, - { /* [103] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 2, /* num_explicit_templates */ 0, /* num_templates */ 2, /* templates */ TemplateIndex(4), - /* parameters */ ParameterIndex(290), + /* parameters */ ParameterIndex(288), /* return_matcher_indices */ MatcherIndicesIndex(133), /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, @@ -4357,12 +4390,23 @@ /* num_explicit_templates */ 0, /* num_templates */ 2, /* templates */ TemplateIndex(4), + /* parameters */ ParameterIndex(290), + /* return_matcher_indices */ MatcherIndicesIndex(133), + /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), + }, + { + /* [105] */ + /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), + /* num_parameters */ 2, + /* num_explicit_templates */ 0, + /* num_templates */ 2, + /* templates */ TemplateIndex(4), /* parameters */ ParameterIndex(292), /* return_matcher_indices */ MatcherIndicesIndex(131), /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [105] */ + /* [106] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 2, /* num_explicit_templates */ 0, @@ -4373,7 +4417,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [106] */ + /* [107] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 2, /* num_explicit_templates */ 0, @@ -4384,7 +4428,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [107] */ + /* [108] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 2, /* num_explicit_templates */ 0, @@ -4395,7 +4439,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [108] */ + /* [109] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 2, /* num_explicit_templates */ 0, @@ -4406,7 +4450,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [109] */ + /* [110] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 2, /* num_explicit_templates */ 0, @@ -4417,7 +4461,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [110] */ + /* [111] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 2, /* num_explicit_templates */ 0, @@ -4428,7 +4472,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [111] */ + /* [112] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 2, /* num_explicit_templates */ 0, @@ -4439,7 +4483,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [112] */ + /* [113] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 2, /* num_explicit_templates */ 0, @@ -4450,7 +4494,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [113] */ + /* [114] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 2, /* num_explicit_templates */ 0, @@ -4461,7 +4505,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [114] */ + /* [115] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 2, /* num_explicit_templates */ 0, @@ -4472,7 +4516,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [115] */ + /* [116] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 2, /* num_explicit_templates */ 0, @@ -4483,7 +4527,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [116] */ + /* [117] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 2, /* num_explicit_templates */ 0, @@ -4494,7 +4538,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [117] */ + /* [118] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 2, /* num_explicit_templates */ 0, @@ -4505,7 +4549,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [118] */ + /* [119] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 2, /* num_explicit_templates */ 0, @@ -4516,7 +4560,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [119] */ + /* [120] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 2, /* num_explicit_templates */ 0, @@ -4527,7 +4571,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [120] */ + /* [121] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 4, /* num_explicit_templates */ 0, @@ -4538,7 +4582,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [121] */ + /* [122] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 4, /* num_explicit_templates */ 0, @@ -4549,7 +4593,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [122] */ + /* [123] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 4, /* num_explicit_templates */ 0, @@ -4560,7 +4604,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [123] */ + /* [124] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 4, /* num_explicit_templates */ 0, @@ -4571,7 +4615,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [124] */ + /* [125] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 4, /* num_explicit_templates */ 0, @@ -4582,7 +4626,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [125] */ + /* [126] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 4, /* num_explicit_templates */ 0, @@ -4593,7 +4637,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [126] */ + /* [127] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 4, /* num_explicit_templates */ 0, @@ -4604,7 +4648,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [127] */ + /* [128] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 4, /* num_explicit_templates */ 0, @@ -4615,7 +4659,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [128] */ + /* [129] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 4, /* num_explicit_templates */ 0, @@ -4626,33 +4670,22 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [129] */ - /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), - /* num_parameters */ 5, - /* num_explicit_templates */ 0, - /* num_templates */ 3, - /* templates */ TemplateIndex(28), - /* parameters */ ParameterIndex(36), - /* return_matcher_indices */ MatcherIndicesIndex(60), - /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), - }, - { /* [130] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), - /* num_parameters */ 4, + /* num_parameters */ 5, /* num_explicit_templates */ 0, - /* num_templates */ 2, + /* num_templates */ 3, /* templates */ TemplateIndex(28), - /* parameters */ ParameterIndex(41), + /* parameters */ ParameterIndex(36), /* return_matcher_indices */ MatcherIndicesIndex(60), /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { /* [131] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), - /* num_parameters */ 5, + /* num_parameters */ 4, /* num_explicit_templates */ 0, - /* num_templates */ 3, + /* num_templates */ 2, /* templates */ TemplateIndex(28), /* parameters */ ParameterIndex(41), /* return_matcher_indices */ MatcherIndicesIndex(60), @@ -4661,11 +4694,11 @@ { /* [132] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), - /* num_parameters */ 4, + /* num_parameters */ 5, /* num_explicit_templates */ 0, - /* num_templates */ 2, + /* num_templates */ 3, /* templates */ TemplateIndex(28), - /* parameters */ ParameterIndex(120), + /* parameters */ ParameterIndex(41), /* return_matcher_indices */ MatcherIndicesIndex(60), /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, @@ -4676,7 +4709,7 @@ /* num_explicit_templates */ 0, /* num_templates */ 2, /* templates */ TemplateIndex(28), - /* parameters */ ParameterIndex(124), + /* parameters */ ParameterIndex(120), /* return_matcher_indices */ MatcherIndicesIndex(60), /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, @@ -4685,6 +4718,17 @@ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 4, /* num_explicit_templates */ 0, + /* num_templates */ 2, + /* templates */ TemplateIndex(28), + /* parameters */ ParameterIndex(124), + /* return_matcher_indices */ MatcherIndicesIndex(60), + /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), + }, + { + /* [135] */ + /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), + /* num_parameters */ 4, + /* num_explicit_templates */ 0, /* num_templates */ 1, /* templates */ TemplateIndex(1), /* parameters */ ParameterIndex(24), @@ -4692,7 +4736,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [135] */ + /* [136] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 5, /* num_explicit_templates */ 0, @@ -4703,7 +4747,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [136] */ + /* [137] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 4, /* num_explicit_templates */ 0, @@ -4714,7 +4758,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [137] */ + /* [138] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 5, /* num_explicit_templates */ 0, @@ -4725,7 +4769,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [138] */ + /* [139] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 4, /* num_explicit_templates */ 0, @@ -4736,7 +4780,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [139] */ + /* [140] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 4, /* num_explicit_templates */ 0, @@ -4747,7 +4791,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [140] */ + /* [141] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 5, /* num_explicit_templates */ 0, @@ -4758,34 +4802,23 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [141] */ - /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), - /* num_parameters */ 6, - /* num_explicit_templates */ 0, - /* num_templates */ 2, - /* templates */ TemplateIndex(7), - /* parameters */ ParameterIndex(24), - /* return_matcher_indices */ MatcherIndicesIndex(6), - /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), - }, - { /* [142] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), - /* num_parameters */ 5, + /* num_parameters */ 6, /* num_explicit_templates */ 0, - /* num_templates */ 1, - /* templates */ TemplateIndex(1), - /* parameters */ ParameterIndex(30), + /* num_templates */ 2, + /* templates */ TemplateIndex(7), + /* parameters */ ParameterIndex(24), /* return_matcher_indices */ MatcherIndicesIndex(6), /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { /* [143] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), - /* num_parameters */ 6, + /* num_parameters */ 5, /* num_explicit_templates */ 0, - /* num_templates */ 2, - /* templates */ TemplateIndex(7), + /* num_templates */ 1, + /* templates */ TemplateIndex(1), /* parameters */ ParameterIndex(30), /* return_matcher_indices */ MatcherIndicesIndex(6), /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), @@ -4793,11 +4826,11 @@ { /* [144] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), - /* num_parameters */ 5, + /* num_parameters */ 6, /* num_explicit_templates */ 0, - /* num_templates */ 1, - /* templates */ TemplateIndex(1), - /* parameters */ ParameterIndex(106), + /* num_templates */ 2, + /* templates */ TemplateIndex(7), + /* parameters */ ParameterIndex(30), /* return_matcher_indices */ MatcherIndicesIndex(6), /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, @@ -4808,13 +4841,24 @@ /* num_explicit_templates */ 0, /* num_templates */ 1, /* templates */ TemplateIndex(1), - /* parameters */ ParameterIndex(111), + /* parameters */ ParameterIndex(106), /* return_matcher_indices */ MatcherIndicesIndex(6), /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { /* [146] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), + /* num_parameters */ 5, + /* num_explicit_templates */ 0, + /* num_templates */ 1, + /* templates */ TemplateIndex(1), + /* parameters */ ParameterIndex(111), + /* return_matcher_indices */ MatcherIndicesIndex(6), + /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), + }, + { + /* [147] */ + /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 3, /* num_explicit_templates */ 0, /* num_templates */ 1, @@ -4824,7 +4868,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [147] */ + /* [148] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 3, /* num_explicit_templates */ 0, @@ -4835,7 +4879,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [148] */ + /* [149] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 2, /* num_explicit_templates */ 0, @@ -4846,7 +4890,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [149] */ + /* [150] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 4, /* num_explicit_templates */ 0, @@ -4857,7 +4901,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [150] */ + /* [151] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 6, /* num_explicit_templates */ 0, @@ -4868,7 +4912,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [151] */ + /* [152] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 3, /* num_explicit_templates */ 0, @@ -4879,7 +4923,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [152] */ + /* [153] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 4, /* num_explicit_templates */ 0, @@ -4890,7 +4934,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [153] */ + /* [154] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 2, /* num_explicit_templates */ 0, @@ -4901,25 +4945,14 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [154] */ + /* [155] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 2, /* num_explicit_templates */ 0, /* num_templates */ 4, /* templates */ TemplateIndex(21), - /* parameters */ ParameterIndex(304), - /* return_matcher_indices */ MatcherIndicesIndex(9), - /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), - }, - { - /* [155] */ - /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), - /* num_parameters */ 2, - /* num_explicit_templates */ 0, - /* num_templates */ 3, - /* templates */ TemplateIndex(37), /* parameters */ ParameterIndex(306), - /* return_matcher_indices */ MatcherIndicesIndex(21), + /* return_matcher_indices */ MatcherIndicesIndex(9), /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { @@ -4930,7 +4963,7 @@ /* num_templates */ 3, /* templates */ TemplateIndex(37), /* parameters */ ParameterIndex(308), - /* return_matcher_indices */ MatcherIndicesIndex(98), + /* return_matcher_indices */ MatcherIndicesIndex(21), /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { @@ -4940,7 +4973,7 @@ /* num_explicit_templates */ 0, /* num_templates */ 3, /* templates */ TemplateIndex(37), - /* parameters */ ParameterIndex(309), + /* parameters */ ParameterIndex(310), /* return_matcher_indices */ MatcherIndicesIndex(98), /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, @@ -4949,15 +4982,26 @@ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 2, /* num_explicit_templates */ 0, - /* num_templates */ 2, + /* num_templates */ 3, /* templates */ TemplateIndex(37), /* parameters */ ParameterIndex(311), - /* return_matcher_indices */ MatcherIndicesIndex(101), + /* return_matcher_indices */ MatcherIndicesIndex(98), /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { /* [159] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), + /* num_parameters */ 2, + /* num_explicit_templates */ 0, + /* num_templates */ 2, + /* templates */ TemplateIndex(37), + /* parameters */ ParameterIndex(313), + /* return_matcher_indices */ MatcherIndicesIndex(101), + /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), + }, + { + /* [160] */ + /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 3, /* num_explicit_templates */ 0, /* num_templates */ 0, @@ -4967,7 +5011,7 @@ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */), }, { - /* [160] */ + /* [161] */ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), /* num_parameters */ 3, /* num_explicit_templates */ 0, @@ -4987,91 +5031,91 @@ /* [0] */ /* fn array_length[I : u32, A : access](ptr<storage, struct_with_runtime_array, A>, I) -> u32 */ /* num overloads */ 1, - /* overloads */ OverloadIndex(148), + /* overloads */ OverloadIndex(149), }, { /* [1] */ /* fn atomic_and[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */ /* num overloads */ 1, - /* overloads */ OverloadIndex(149), + /* overloads */ OverloadIndex(150), }, { /* [2] */ /* fn atomic_compare_exchange[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, U, T, T) -> T */ /* num overloads */ 1, - /* overloads */ OverloadIndex(150), + /* overloads */ OverloadIndex(151), }, { /* [3] */ /* fn atomic_exchange[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */ /* num overloads */ 1, - /* overloads */ OverloadIndex(149), + /* overloads */ OverloadIndex(150), }, { /* [4] */ /* fn atomic_iadd[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */ /* num overloads */ 1, - /* overloads */ OverloadIndex(149), + /* overloads */ OverloadIndex(150), }, { /* [5] */ /* fn atomic_isub[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */ /* num overloads */ 1, - /* overloads */ OverloadIndex(149), + /* overloads */ OverloadIndex(150), }, { /* [6] */ /* fn atomic_load[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U) -> T */ /* num overloads */ 1, - /* overloads */ OverloadIndex(151), + /* overloads */ OverloadIndex(152), }, { /* [7] */ /* fn atomic_or[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */ /* num overloads */ 1, - /* overloads */ OverloadIndex(149), + /* overloads */ OverloadIndex(150), }, { /* [8] */ /* fn atomic_smax[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */ /* num overloads */ 1, - /* overloads */ OverloadIndex(149), + /* overloads */ OverloadIndex(150), }, { /* [9] */ /* fn atomic_smin[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */ /* num overloads */ 1, - /* overloads */ OverloadIndex(149), + /* overloads */ OverloadIndex(150), }, { /* [10] */ /* fn atomic_store[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) */ /* num overloads */ 1, - /* overloads */ OverloadIndex(152), + /* overloads */ OverloadIndex(153), }, { /* [11] */ /* fn atomic_umax[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */ /* num overloads */ 1, - /* overloads */ OverloadIndex(149), + /* overloads */ OverloadIndex(150), }, { /* [12] */ /* fn atomic_umin[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */ /* num overloads */ 1, - /* overloads */ OverloadIndex(149), + /* overloads */ OverloadIndex(150), }, { /* [13] */ /* fn atomic_xor[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */ /* num overloads */ 1, - /* overloads */ OverloadIndex(149), + /* overloads */ OverloadIndex(150), }, { /* [14] */ /* fn dot[N : num, T : f32_f16](vec<N, T>, vec<N, T>) -> T */ /* num overloads */ 1, - /* overloads */ OverloadIndex(153), + /* overloads */ OverloadIndex(154), }, { /* [15] */ @@ -5082,7 +5126,7 @@ /* fn image_dref_gather[A : f32, B : iu32](sampled_image<texture_depth_cube>, vec3<f32>, A, B) -> vec4<f32> */ /* fn image_dref_gather[A : f32, B : iu32](sampled_image<texture_depth_cube_array>, vec4<f32>, A, B) -> vec4<f32> */ /* num overloads */ 6, - /* overloads */ OverloadIndex(128), + /* overloads */ OverloadIndex(129), }, { /* [16] */ @@ -5095,7 +5139,7 @@ /* fn image_fetch[I : iu32, C : iu32, S : iu32](texture_depth_2d_array, vec3<C>, I, S) -> vec4<f32> */ /* fn image_fetch[I : iu32, C : iu32, S : iu32](texture_depth_multisampled_2d, vec2<C>, I, S) -> vec4<f32> */ /* num overloads */ 8, - /* overloads */ OverloadIndex(120), + /* overloads */ OverloadIndex(121), }, { /* [17] */ @@ -5150,7 +5194,7 @@ /* fn image_query_size_lod[A : iu32](texture_depth_cube, A) -> vec2<u32> */ /* fn image_query_size_lod[A : iu32](texture_depth_cube_array, A) -> vec3<u32> */ /* num overloads */ 10, - /* overloads */ OverloadIndex(100), + /* overloads */ OverloadIndex(101), }, { /* [20] */ @@ -5166,7 +5210,8 @@ /* fn image_read[F : f32_texel_format, A : readable, C : iu32, S : iu32](texture_storage_3d<F, A>, vec3<C>, S) -> vec4<f32> */ /* fn image_read[F : i32_texel_format, A : readable, C : iu32, S : iu32](texture_storage_3d<F, A>, vec3<C>, S) -> vec4<i32> */ /* fn image_read[F : u32_texel_format, A : readable, C : iu32, S : iu32](texture_storage_3d<F, A>, vec3<C>, S) -> vec4<u32> */ - /* num overloads */ 12, + /* fn image_read[T : fiu32, C : iu32](input_attachment: input_attachment<T>, vec2<C>) -> vec4<T> */ + /* num overloads */ 13, /* overloads */ OverloadIndex(76), }, { @@ -5234,7 +5279,7 @@ /* fn image_sample_dref_implicit_lod[C : iu32](sampled_image<texture_depth_cube>, vec3<f32>, f32, C) -> f32 */ /* fn image_sample_dref_implicit_lod[C : iu32](sampled_image<texture_depth_cube_array>, vec4<f32>, f32, C) -> f32 */ /* num overloads */ 6, - /* overloads */ OverloadIndex(134), + /* overloads */ OverloadIndex(135), }, { /* [24] */ @@ -5245,7 +5290,7 @@ /* fn image_sample_dref_explicit_lod[C : iu32](sampled_image<texture_depth_cube>, vec3<f32>, f32, C, f32) -> f32 */ /* fn image_sample_dref_explicit_lod[C : iu32](sampled_image<texture_depth_cube_array>, vec4<f32>, f32, C, f32) -> f32 */ /* num overloads */ 6, - /* overloads */ OverloadIndex(140), + /* overloads */ OverloadIndex(141), }, { /* [25] */ @@ -5262,25 +5307,25 @@ /* fn image_write[C : iu32, D : iu32](texture_storage_3d<i32_texel_format, writable>, vec3<C>, vec4<i32>, D) */ /* fn image_write[C : iu32, D : iu32](texture_storage_3d<u32_texel_format, writable>, vec3<C>, vec4<u32>, D) */ /* num overloads */ 12, - /* overloads */ OverloadIndex(88), + /* overloads */ OverloadIndex(89), }, { /* [26] */ /* fn matrix_times_matrix[T : f32_f16, K : num, C : num, R : num](mat<K, R, T>, mat<C, K, T>) -> mat<C, R, T> */ /* num overloads */ 1, - /* overloads */ OverloadIndex(154), + /* overloads */ OverloadIndex(155), }, { /* [27] */ /* fn matrix_times_scalar[T : f32_f16, N : num, M : num](mat<N, M, T>, T) -> mat<N, M, T> */ /* num overloads */ 1, - /* overloads */ OverloadIndex(155), + /* overloads */ OverloadIndex(156), }, { /* [28] */ /* fn matrix_times_vector[T : f32_f16, N : num, M : num](mat<N, M, T>, vec<N, T>) -> vec<M, T> */ /* num overloads */ 1, - /* overloads */ OverloadIndex(156), + /* overloads */ OverloadIndex(157), }, { /* [29] */ @@ -5295,38 +5340,38 @@ /* fn sampled_image[S : samplers](texture_depth_cube, S) -> sampled_image<texture_depth_cube> */ /* fn sampled_image[S : samplers](texture_depth_cube_array, S) -> sampled_image<texture_depth_cube_array> */ /* num overloads */ 10, - /* overloads */ OverloadIndex(110), + /* overloads */ OverloadIndex(111), }, { /* [30] */ /* fn select[T : scalar](bool, T, T) -> T */ /* fn select[N : num, T : scalar](vec<N, bool>, vec<N, T>, vec<N, T>) -> vec<N, T> */ /* num overloads */ 2, - /* overloads */ OverloadIndex(146), + /* overloads */ OverloadIndex(147), }, { /* [31] */ /* fn vector_times_matrix[T : f32_f16, N : num, M : num](vec<N, T>, mat<M, N, T>) -> vec<M, T> */ /* num overloads */ 1, - /* overloads */ OverloadIndex(157), + /* overloads */ OverloadIndex(158), }, { /* [32] */ /* fn vector_times_scalar[T : f32_f16, N : num](vec<N, T>, T) -> vec<N, T> */ /* num overloads */ 1, - /* overloads */ OverloadIndex(158), + /* overloads */ OverloadIndex(159), }, { /* [33] */ /* fn sdot(u32, u32, u32) -> i32 */ /* num overloads */ 1, - /* overloads */ OverloadIndex(159), + /* overloads */ OverloadIndex(160), }, { /* [34] */ /* fn udot(u32, u32, u32) -> u32 */ /* num overloads */ 1, - /* overloads */ OverloadIndex(160), + /* overloads */ OverloadIndex(161), }, };
diff --git a/src/tint/lang/spirv/reader/lower/shader_io_test.cc b/src/tint/lang/spirv/reader/lower/shader_io_test.cc index 79beb46..349db27 100644 --- a/src/tint/lang/spirv/reader/lower/shader_io_test.cc +++ b/src/tint/lang/spirv/reader/lower/shader_io_test.cc
@@ -57,13 +57,14 @@ TEST_F(SpirvReader_ShaderIOTest, NoInputsOrOutputs) { auto* ep = b.Function("foo", ty.void_()); ep->SetStage(core::ir::Function::PipelineStage::kCompute); + ep->SetWorkgroupSize(1, 1, 1); b.Append(ep->Block(), [&] { // b.Return(ep); }); auto* src = R"( -%foo = @compute func():void { +%foo = @compute @workgroup_size(1, 1, 1) func():void { $B1: { ret } @@ -342,6 +343,7 @@ // Use a different subset of the inputs in the entry point. auto* ep = b.Function("main1", ty.void_(), core::ir::Function::PipelineStage::kCompute); + ep->SetWorkgroupSize(1, 1, 1); b.Append(ep->Block(), [&] { auto* group_value = b.Load(group_id); auto* gid_value = b.Load(gid); @@ -365,7 +367,7 @@ ret } } -%main1 = @compute func():void { +%main1 = @compute @workgroup_size(1, 1, 1) func():void { $B3: { %9:vec3<u32> = load %group_id %10:vec3<u32> = load %gid @@ -384,7 +386,7 @@ ret } } -%main1 = @compute func(%gid_1:vec3<u32> [@global_invocation_id], %lid_1:vec3<u32> [@local_invocation_id], %group_id:vec3<u32> [@workgroup_id]):void { # %gid_1: 'gid', %lid_1: 'lid' +%main1 = @compute @workgroup_size(1, 1, 1) func(%gid_1:vec3<u32> [@global_invocation_id], %lid_1:vec3<u32> [@local_invocation_id], %group_id:vec3<u32> [@workgroup_id]):void { # %gid_1: 'gid', %lid_1: 'lid' $B2: { %9:vec3<u32> = add %group_id, %gid_1 %10:void = call %foo, %gid_1, %lid_1 @@ -422,6 +424,7 @@ mod.root_block->Append(group_id); auto* ep = b.Function("main1", ty.void_(), core::ir::Function::PipelineStage::kCompute); + ep->SetWorkgroupSize(1, 1, 1); auto* foo = b.Function("foo", ty.void_()); // Use a subset of the inputs in the entry point. @@ -448,7 +451,7 @@ %group_id:ptr<__in, vec3<u32>, read> = var @builtin(workgroup_id) } -%main1 = @compute func():void { +%main1 = @compute @workgroup_size(1, 1, 1) func():void { $B2: { %5:vec3<u32> = load %group_id %6:vec3<u32> = load %gid @@ -469,7 +472,7 @@ EXPECT_EQ(src, str()); auto* expect = R"( -%main1 = @compute func(%gid:vec3<u32> [@global_invocation_id], %lid:vec3<u32> [@local_invocation_id], %group_id:vec3<u32> [@workgroup_id]):void { +%main1 = @compute @workgroup_size(1, 1, 1) func(%gid:vec3<u32> [@global_invocation_id], %lid:vec3<u32> [@local_invocation_id], %group_id:vec3<u32> [@workgroup_id]):void { $B1: { %5:vec3<u32> = add %group_id, %gid %6:void = call %foo, %gid, %lid @@ -523,6 +526,7 @@ // Call the helper without directly referencing any inputs. auto* ep1 = b.Function("main1", ty.void_(), core::ir::Function::PipelineStage::kCompute); + ep1->SetWorkgroupSize(1, 1, 1); b.Append(ep1->Block(), [&] { b.Call(foo); b.Return(ep1); @@ -530,6 +534,7 @@ // Reference another input and then call the helper. auto* ep2 = b.Function("main2", ty.void_(), core::ir::Function::PipelineStage::kCompute); + ep2->SetWorkgroupSize(1, 1, 1); b.Append(ep2->Block(), [&] { auto* group_value = b.Load(group_id); b.Add(ty.vec3<u32>(), group_value, group_value); @@ -552,13 +557,13 @@ ret } } -%main1 = @compute func():void { +%main1 = @compute @workgroup_size(1, 1, 1) func():void { $B3: { %9:void = call %foo ret } } -%main2 = @compute func():void { +%main2 = @compute @workgroup_size(1, 1, 1) func():void { $B4: { %11:vec3<u32> = load %group_id %12:vec3<u32> = add %11, %11 @@ -576,13 +581,13 @@ ret } } -%main1 = @compute func(%gid_1:vec3<u32> [@global_invocation_id], %lid_1:vec3<u32> [@local_invocation_id]):void { # %gid_1: 'gid', %lid_1: 'lid' +%main1 = @compute @workgroup_size(1, 1, 1) func(%gid_1:vec3<u32> [@global_invocation_id], %lid_1:vec3<u32> [@local_invocation_id]):void { # %gid_1: 'gid', %lid_1: 'lid' $B2: { %8:void = call %foo, %gid_1, %lid_1 ret } } -%main2 = @compute func(%gid_2:vec3<u32> [@global_invocation_id], %lid_2:vec3<u32> [@local_invocation_id], %group_id:vec3<u32> [@workgroup_id]):void { # %gid_2: 'gid', %lid_2: 'lid' +%main2 = @compute @workgroup_size(1, 1, 1) func(%gid_2:vec3<u32> [@global_invocation_id], %lid_2:vec3<u32> [@local_invocation_id], %group_id:vec3<u32> [@workgroup_id]):void { # %gid_2: 'gid', %lid_2: 'lid' $B3: { %13:vec3<u32> = add %group_id, %group_id %14:void = call %foo, %gid_2, %lid_2 @@ -606,6 +611,7 @@ mod.root_block->Append(lid); auto* ep = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kCompute); + ep->SetWorkgroupSize(1, 1, 1); b.Append(ep->Block(), [&] { b.LoadVectorElement(lid, 2_u); b.Return(ep); @@ -616,7 +622,7 @@ %lid:ptr<__in, vec3<u32>, read> = var @builtin(local_invocation_id) } -%foo = @compute func():void { +%foo = @compute @workgroup_size(1, 1, 1) func():void { $B2: { %3:u32 = load_vector_element %lid, 2u ret @@ -626,7 +632,7 @@ EXPECT_EQ(src, str()); auto* expect = R"( -%foo = @compute func(%lid:vec3<u32> [@local_invocation_id]):void { +%foo = @compute @workgroup_size(1, 1, 1) func(%lid:vec3<u32> [@local_invocation_id]):void { $B1: { %3:u32 = access %lid, 2u ret @@ -649,6 +655,7 @@ mod.root_block->Append(lid); auto* ep = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kCompute); + ep->SetWorkgroupSize(1, 1, 1); b.Append(ep->Block(), [&] { auto* access_1 = b.Access(ty.ptr(core::AddressSpace::kIn, ty.vec3<u32>()), lid); auto* access_2 = b.Access(ty.ptr(core::AddressSpace::kIn, ty.vec3<u32>()), access_1); @@ -663,7 +670,7 @@ %lid:ptr<__in, vec3<u32>, read> = var @builtin(local_invocation_id) } -%foo = @compute func():void { +%foo = @compute @workgroup_size(1, 1, 1) func():void { $B2: { %3:ptr<__in, vec3<u32>, read> = access %lid %4:ptr<__in, vec3<u32>, read> = access %3 @@ -677,7 +684,7 @@ EXPECT_EQ(src, str()); auto* expect = R"( -%foo = @compute func(%lid:vec3<u32> [@local_invocation_id]):void { +%foo = @compute @workgroup_size(1, 1, 1) func(%lid:vec3<u32> [@local_invocation_id]):void { $B1: { %3:u32 = access %lid, 2u %4:vec3<u32> = mul %lid, %3
diff --git a/src/tint/lang/spirv/spirv.def b/src/tint/lang/spirv/spirv.def index 9dabb13..4860dc5 100644 --- a/src/tint/lang/spirv/spirv.def +++ b/src/tint/lang/spirv/spirv.def
@@ -78,6 +78,7 @@ type texture_storage_2d_array<F: texel_format, A: access> type texture_storage_3d<F: texel_format, A: access> type ptr<S: address_space, T, A: access> +type input_attachment<T> type struct_with_runtime_array type sampled_image<T> @@ -219,6 +220,7 @@ fn image_read[F: f32_texel_format, A: readable, C: iu32, S: iu32](texture_storage_3d<F, A>, vec3<C>, S) -> vec4<f32> fn image_read[F: i32_texel_format, A: readable, C: iu32, S: iu32](texture_storage_3d<F, A>, vec3<C>, S) -> vec4<i32> fn image_read[F: u32_texel_format, A: readable, C: iu32, S: iu32](texture_storage_3d<F, A>, vec3<C>, S) -> vec4<u32> +fn image_read[T: fiu32, C: iu32](input_attachment: input_attachment<T>, vec2<C>) -> vec4<T> fn image_sample_implicit_lod[T: fiu32, C: iu32](sampled_image<texture_1d<T> >, f32, C) -> vec4<f32> fn image_sample_implicit_lod[T: fiu32, C: iu32](sampled_image<texture_2d<T> >, vec2<f32>, C) -> vec4<f32>
diff --git a/src/tint/lang/spirv/writer/ast_printer/ast_builtin_test.cc b/src/tint/lang/spirv/writer/ast_printer/ast_builtin_test.cc index 9c58491..299c444 100644 --- a/src/tint/lang/spirv/writer/ast_printer/ast_builtin_test.cc +++ b/src/tint/lang/spirv/writer/ast_printer/ast_builtin_test.cc
@@ -26,6 +26,7 @@ // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include "src/tint/lang/core/type/depth_texture.h" +#include "src/tint/lang/core/type/input_attachment.h" #include "src/tint/lang/core/type/texture_dimension.h" #include "src/tint/lang/spirv/writer/ast_printer/helper_test.h" #include "src/tint/lang/spirv/writer/common/spv_dump_test.h" @@ -4248,5 +4249,60 @@ } // namespace Packed_4x8_integer_dot_product_builtin_tests +namespace InputAttachments_builtin_tests { + +TEST_F(BuiltinSpirvASTPrinterTest, Call_InputAttachmentLoad) { + // enable chromium_internal_input_attachments; + // @group(0) @binding(0) @input_attachment_index(3) + // var input_tex : input_attachment<f32>; + // fn f() -> vec4f { + // return inputAttachmentLoad(input_tex); + // } + + Enable(wgsl::Extension::kChromiumInternalInputAttachments); + + auto* input_tex = GlobalVar("input_tex", ty.input_attachment(ty.Of<f32>()), + Vector{Binding(0_u), Group(0_u), InputAttachmentIndex(3_u)}); + + auto* func = Func("f", Empty, ty.vec4<f32>(), + Vector{ + Return(Call("inputAttachmentLoad", "input_tex")), + }); + + Builder& b = Build(); + + ASSERT_TRUE(b.GenerateExtension(wgsl::Extension::kChromiumInternalInputAttachments)) + << b.Diagnostics(); + ASSERT_TRUE(b.GenerateGlobalVariable(input_tex)) << b.Diagnostics(); + ASSERT_TRUE(b.GenerateFunction(func)) << b.Diagnostics(); + + auto got = DumpModule(b.Module()); + auto expect = R"(OpCapability InputAttachment +OpName %1 "input_tex" +OpName %7 "f" +OpDecorate %1 Binding 0 +OpDecorate %1 DescriptorSet 0 +OpDecorate %1 InputAttachmentIndex 3 +%4 = OpTypeFloat 32 +%3 = OpTypeImage %4 SubpassData 0 0 0 2 Unknown +%2 = OpTypePointer UniformConstant %3 +%1 = OpVariable %2 UniformConstant +%6 = OpTypeVector %4 4 +%5 = OpTypeFunction %6 +%12 = OpTypeInt 32 1 +%11 = OpTypeVector %12 2 +%13 = OpConstantNull %11 +%7 = OpFunction %6 None %5 +%8 = OpLabel +%10 = OpLoad %3 %1 +%9 = OpImageRead %6 %10 %13 +OpReturnValue %9 +OpFunctionEnd +)"; + EXPECT_EQ(got, expect); +} + +} // namespace InputAttachments_builtin_tests + } // namespace } // namespace tint::spirv::writer
diff --git a/src/tint/lang/spirv/writer/ast_printer/builder.cc b/src/tint/lang/spirv/writer/ast_printer/builder.cc index da7014c..563ac13 100644 --- a/src/tint/lang/spirv/writer/ast_printer/builder.cc +++ b/src/tint/lang/spirv/writer/ast_printer/builder.cc
@@ -37,6 +37,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/reference.h" #include "src/tint/lang/core/type/sampled_texture.h" @@ -281,7 +282,9 @@ wgsl::Extension::kChromiumExperimentalSubgroups, wgsl::Extension::kChromiumInternalDualSourceBlending, wgsl::Extension::kChromiumInternalGraphite, + wgsl::Extension::kChromiumInternalInputAttachments, wgsl::Extension::kF16, + wgsl::Extension::kDualSourceBlending, })) { return false; } @@ -351,6 +354,9 @@ module_.PushCapability(SpvCapabilityUniformAndStorageBuffer16BitAccess); module_.PushCapability(SpvCapabilityStorageBuffer16BitAccess); break; + case wgsl::Extension::kChromiumInternalInputAttachments: + module_.PushCapability(SpvCapabilityInputAttachment); + break; default: return false; } @@ -843,7 +849,14 @@ }, [&](const ast::InternalAttribute*) { return true; // ignored - }, // + }, + [&](const ast::InputAttachmentIndexAttribute*) { + auto iidx = sem->Attributes().input_attachment_index; + module_.PushAnnot(spv::Op::OpDecorate, + {Operand(var_id), U32Operand(SpvDecorationInputAttachmentIndex), + Operand(iidx.value())}); + return true; + }, // TINT_ICE_ON_NO_MATCH); if (!ok) { return false; @@ -2617,7 +2630,13 @@ return gen(argument); }; - auto* texture = arg(Usage::kTexture); + Usage textureUsage; + if (builtin->Fn() == wgsl::BuiltinFn::kInputAttachmentLoad) { + textureUsage = Usage::kInputAttachment; + } else { + textureUsage = Usage::kTexture; + } + auto* texture = arg(textureUsage); if (TINT_UNLIKELY(!texture)) { TINT_ICE() << "missing texture argument"; } @@ -2967,6 +2986,18 @@ Operand(GenerateConstantIfNeeded(ScalarConstant::F32(0.0)))}); break; } + case wgsl::BuiltinFn::kInputAttachmentLoad: { + op = spv::Op::OpImageRead; + append_result_type_and_id_to_spirv_params_for_read(); + spirv_params.emplace_back(gen_arg(Usage::kInputAttachment)); + + // coords for input_attachment are always (0, 0) + auto* vec2i = + builder_.create<core::type::Vector>(builder_.create<core::type::I32>(), 2u); + spirv_params.emplace_back(Operand(GenerateConstantNullIfNeeded(vec2i))); + + break; + } default: TINT_UNREACHABLE(); } @@ -3759,6 +3790,10 @@ dim_literal = SpvDimCube; } + if (texture->Is<core::type::InputAttachment>()) { + dim_literal = SpvDimSubpassData; + } + uint32_t ms_literal = 0u; if (texture->IsAnyOf<core::type::MultisampledTexture, core::type::DepthMultisampledTexture>()) { ms_literal = 1u; @@ -3791,7 +3826,8 @@ }, [&](const core::type::SampledTexture* t) { return GenerateTypeIfNeeded(t->type()); }, [&](const core::type::MultisampledTexture* t) { return GenerateTypeIfNeeded(t->type()); }, - [&](const core::type::StorageTexture* t) { return GenerateTypeIfNeeded(t->type()); }, // + [&](const core::type::StorageTexture* t) { return GenerateTypeIfNeeded(t->type()); }, + [&](const core::type::InputAttachment* t) { return GenerateTypeIfNeeded(t->type()); }, // TINT_ICE_ON_NO_MATCH); if (type_id == 0u) { return false;
diff --git a/src/tint/lang/spirv/writer/printer/printer.cc b/src/tint/lang/spirv/writer/printer/printer.cc index e1a00e0..c2e7104 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; @@ -630,7 +632,12 @@ break; } case core::type::TextureDimension::k2d: { - dim = SpvDim2D; + if (texture->Is<core::type::InputAttachment>()) { + module_.PushCapability(SpvCapabilityInputAttachment); + dim = SpvDimSubpassData; + } else { + dim = SpvDim2D; + } break; } case core::type::TextureDimension::k2dArray: { @@ -2118,6 +2125,14 @@ {id, U32Operand(SpvDecorationNonReadable)}); } } + + auto iidx = var->InputAttachmentIndex(); + if (iidx) { + TINT_ASSERT(store_ty->Is<core::type::InputAttachment>()); + module_.PushAnnot( + spv::Op::OpDecorate, + {id, U32Operand(SpvDecorationInputAttachmentIndex), iidx.value()}); + } break; } case core::AddressSpace::kWorkgroup: {
diff --git a/src/tint/lang/spirv/writer/raise/builtin_polyfill.cc b/src/tint/lang/spirv/writer/raise/builtin_polyfill.cc index 70808f6..9ef3192 100644 --- a/src/tint/lang/spirv/writer/raise/builtin_polyfill.cc +++ b/src/tint/lang/spirv/writer/raise/builtin_polyfill.cc
@@ -37,6 +37,7 @@ #include "src/tint/lang/core/type/builtin_structs.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/sampled_texture.h" #include "src/tint/lang/core/type/storage_texture.h" @@ -99,6 +100,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 +170,9 @@ case core::BuiltinFn::kQuantizeToF16: QuantizeToF16Vec(builtin); break; + case core::BuiltinFn::kInputAttachmentLoad: + InputAttachmentLoad(builtin); + break; default: break; } @@ -853,6 +858,34 @@ 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); + + auto* texture = builtin->Args()[0]; + // coords for input_attachment are always (0, 0) + auto* coords = b.Composite(ty.vec2<i32>(), 0_i, 0_i); + + // Start building the argument list for the builtin. + // The first two operands are always the texture and then the coordinates. + Vector<core::ir::Value*, 8> builtin_args; + builtin_args.Push(texture); + builtin_args.Push(coords); + + // Call the builtin. + // The result is always a vec4 in SPIR-V. + auto* result_ty = builtin->Result(0)->Type(); + TINT_ASSERT(result_ty->Is<core::type::Vector>()); + + core::ir::Instruction* result = b.Call<spirv::ir::BuiltinCall>( + result_ty, spirv::BuiltinFn::kImageRead, std::move(builtin_args)); + result->InsertBefore(builtin); + + result->SetResults(Vector{builtin->DetachResult()}); + builtin->Destroy(); + } }; } // namespace
diff --git a/src/tint/lang/spirv/writer/raise/builtin_polyfill_test.cc b/src/tint/lang/spirv/writer/raise/builtin_polyfill_test.cc index a215e73..9d845a2 100644 --- a/src/tint/lang/spirv/writer/raise/builtin_polyfill_test.cc +++ b/src/tint/lang/spirv/writer/raise/builtin_polyfill_test.cc
@@ -34,6 +34,7 @@ #include "src/tint/lang/core/type/atomic.h" #include "src/tint/lang/core/type/builtin_structs.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/sampled_texture.h" #include "src/tint/lang/core/type/storage_texture.h" @@ -2957,5 +2958,39 @@ EXPECT_EQ(expect, str()); } +TEST_F(SpirvWriter_BuiltinPolyfillTest, InputAttachmentLoad) { + auto* t = b.FunctionParam("t", ty.Get<core::type::InputAttachment>(ty.f32())); + auto* func = b.Function("foo", ty.vec4<f32>()); + func->SetParams({t}); + + b.Append(func->Block(), [&] { + auto* result = b.Call(ty.vec4<f32>(), core::BuiltinFn::kInputAttachmentLoad, t); + b.Return(func, result); + }); + + auto* src = R"( +%foo = func(%t:input_attachment<f32>):vec4<f32> { + $B1: { + %3:vec4<f32> = inputAttachmentLoad %t + ret %3 + } +} +)"; + EXPECT_EQ(src, str()); + + auto* expect = R"( +%foo = func(%t:input_attachment<f32>):vec4<f32> { + $B1: { + %3:vec4<f32> = spirv.image_read %t, vec2<i32>(0i) + ret %3 + } +} +)"; + + Run(BuiltinPolyfill); + + EXPECT_EQ(expect, str()); +} + } // namespace } // namespace tint::spirv::writer::raise
diff --git a/src/tint/lang/spirv/writer/raise/shader_io_test.cc b/src/tint/lang/spirv/writer/raise/shader_io_test.cc index 4faaf5d..d01528f 100644 --- a/src/tint/lang/spirv/writer/raise/shader_io_test.cc +++ b/src/tint/lang/spirv/writer/raise/shader_io_test.cc
@@ -42,13 +42,14 @@ TEST_F(SpirvWriter_ShaderIOTest, NoInputsOrOutputs) { auto* ep = b.Function("foo", ty.void_()); ep->SetStage(core::ir::Function::PipelineStage::kCompute); + ep->SetWorkgroupSize(1, 1, 1); b.Append(ep->Block(), [&] { // b.Return(ep); }); auto* src = R"( -%foo = @compute func():void { +%foo = @compute @workgroup_size(1, 1, 1) func():void { $B1: { ret }
diff --git a/src/tint/lang/wgsl/ast/builder.h b/src/tint/lang/wgsl/ast/builder.h index b643fc4..03cc5b0 100644 --- a/src/tint/lang/wgsl/ast/builder.h +++ b/src/tint/lang/wgsl/ast/builder.h
@@ -46,7 +46,6 @@ #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"
diff --git a/src/tint/lang/wgsl/ast/transform/canonicalize_entry_point_io_test.cc b/src/tint/lang/wgsl/ast/transform/canonicalize_entry_point_io_test.cc index ec83e9e..1bd2d64 100644 --- a/src/tint/lang/wgsl/ast/transform/canonicalize_entry_point_io_test.cc +++ b/src/tint/lang/wgsl/ast/transform/canonicalize_entry_point_io_test.cc
@@ -4131,7 +4131,7 @@ TEST_F(CanonicalizeEntryPointIOTest, Return_Struct_Index_Attribute_Spirv) { auto* src = R"( -enable chromium_internal_dual_source_blending; +enable dual_source_blending; struct FragOutput { @location(0) @blend_src(0) color : vec4<f32>, @@ -4152,7 +4152,7 @@ )"; auto* expect = R"( -enable chromium_internal_dual_source_blending; +enable dual_source_blending; @location(0) @blend_src(0) @internal(disable_validation__ignore_address_space) var<__out> color_1 : vec4<f32>; @@ -4197,7 +4197,7 @@ TEST_F(CanonicalizeEntryPointIOTest, Return_Struct_Index_Attribute_Msl) { auto* src = R"( -enable chromium_internal_dual_source_blending; +enable dual_source_blending; struct FragOutput { @location(0) @blend_src(0) color : vec4<f32>, @@ -4218,7 +4218,7 @@ )"; auto* expect = R"( -enable chromium_internal_dual_source_blending; +enable dual_source_blending; struct FragOutput { color : vec4<f32>, @@ -4268,7 +4268,7 @@ TEST_F(CanonicalizeEntryPointIOTest, Return_Struct_Index_Attribute_Hlsl) { auto* src = R"( -enable chromium_internal_dual_source_blending; +enable dual_source_blending; struct FragOutput { @location(0) @blend_src(0) color : vec4<f32>, @@ -4289,7 +4289,7 @@ )"; auto* expect = R"( -enable chromium_internal_dual_source_blending; +enable dual_source_blending; struct FragOutput { color : vec4<f32>,
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..f03fada 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: @@ -696,7 +701,7 @@ f == BuiltinFn::kTextureSampleCompareLevel || // f == BuiltinFn::kTextureSampleGrad || // f == BuiltinFn::kTextureSampleLevel || // - f == BuiltinFn::kTextureStore; + f == BuiltinFn::kTextureStore || f == BuiltinFn::kInputAttachmentLoad; } bool IsImageQuery(BuiltinFn f) {
diff --git a/src/tint/lang/wgsl/builtin_fn.cc.tmpl b/src/tint/lang/wgsl/builtin_fn.cc.tmpl index 613a862..288bdac 100644 --- a/src/tint/lang/wgsl/builtin_fn.cc.tmpl +++ b/src/tint/lang/wgsl/builtin_fn.cc.tmpl
@@ -65,7 +65,7 @@ f == BuiltinFn::kTextureSampleCompareLevel || // f == BuiltinFn::kTextureSampleGrad || // f == BuiltinFn::kTextureSampleLevel || // - f == BuiltinFn::kTextureStore; + f == BuiltinFn::kTextureStore || f == BuiltinFn::kInputAttachmentLoad; } bool IsImageQuery(BuiltinFn f) {
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/extension.cc b/src/tint/lang/wgsl/extension.cc index e6790ec..395a0a4 100644 --- a/src/tint/lang/wgsl/extension.cc +++ b/src/tint/lang/wgsl/extension.cc
@@ -69,6 +69,9 @@ if (str == "chromium_internal_relaxed_uniform_layout") { return Extension::kChromiumInternalRelaxedUniformLayout; } + if (str == "dual_source_blending") { + return Extension::kDualSourceBlending; + } if (str == "f16") { return Extension::kF16; } @@ -97,6 +100,8 @@ return "chromium_internal_input_attachments"; case Extension::kChromiumInternalRelaxedUniformLayout: return "chromium_internal_relaxed_uniform_layout"; + case Extension::kDualSourceBlending: + return "dual_source_blending"; case Extension::kF16: return "f16"; }
diff --git a/src/tint/lang/wgsl/extension.h b/src/tint/lang/wgsl/extension.h index ed1ee13..dbf3ff7 100644 --- a/src/tint/lang/wgsl/extension.h +++ b/src/tint/lang/wgsl/extension.h
@@ -55,6 +55,7 @@ kChromiumInternalGraphite, kChromiumInternalInputAttachments, kChromiumInternalRelaxedUniformLayout, + kDualSourceBlending, kF16, }; @@ -85,6 +86,7 @@ "chromium_internal_graphite", "chromium_internal_input_attachments", "chromium_internal_relaxed_uniform_layout", + "dual_source_blending", "f16", }; @@ -99,6 +101,7 @@ Extension::kChromiumInternalGraphite, Extension::kChromiumInternalInputAttachments, Extension::kChromiumInternalRelaxedUniformLayout, + Extension::kDualSourceBlending, Extension::kF16, };
diff --git a/src/tint/lang/wgsl/extension_bench.cc b/src/tint/lang/wgsl/extension_bench.cc index 7e209ff..18731c7 100644 --- a/src/tint/lang/wgsl/extension_bench.cc +++ b/src/tint/lang/wgsl/extension_bench.cc
@@ -108,13 +108,20 @@ "chromiuminternal_relaxed_uniform_layut", "cXroDium_internal_rJJlaed_uniform_layout", "chromium_int8nal_relaed_uniform_layut", - "k", - "16", - "J1", + "dul_okrc_blen11ing", + "dua_source_blending", + "duJl_source_blendig", + "dual_source_blending", + "dual_source_clending", + "dual_sOurce_blending", + "dualKKs__urce_blttvnding", + "xx8", + "__F", + "f1q", "f16", - "c16", - "fO6", - "_KKttvv", + "331O", + "ftt6QQ", + "666", }; for (auto _ : state) { for (auto* str : kStrings) {
diff --git a/src/tint/lang/wgsl/extension_test.cc b/src/tint/lang/wgsl/extension_test.cc index 9e6c182..0237bd5 100644 --- a/src/tint/lang/wgsl/extension_test.cc +++ b/src/tint/lang/wgsl/extension_test.cc
@@ -66,6 +66,7 @@ {"chromium_internal_graphite", Extension::kChromiumInternalGraphite}, {"chromium_internal_input_attachments", Extension::kChromiumInternalInputAttachments}, {"chromium_internal_relaxed_uniform_layout", Extension::kChromiumInternalRelaxedUniformLayout}, + {"dual_source_blending", Extension::kDualSourceBlending}, {"f16", Extension::kF16}, }; @@ -97,9 +98,12 @@ {"chromium_internl_relaxyd_uniform_layout", Extension::kUndefined}, {"chromnnum_internrr77_Gelaxell_uniform_layout", Extension::kUndefined}, {"chromium_intern4l_relaxe00_uniform_layout", Extension::kUndefined}, - {"5", Extension::kUndefined}, - {"u16", Extension::kUndefined}, - {"f", Extension::kUndefined}, + {"dua_ource_bledoong", Extension::kUndefined}, + {"dualsorce_blendzzng", Extension::kUndefined}, + {"ua_sopiirce_bl11nding", Extension::kUndefined}, + {"f1XX", Extension::kUndefined}, + {"55199II", Extension::kUndefined}, + {"frSSHHa", Extension::kUndefined}, }; using ExtensionParseTest = testing::TestWithParam<Case>;
diff --git a/src/tint/lang/wgsl/inspector/inspector_test.cc b/src/tint/lang/wgsl/inspector/inspector_test.cc index 1a7c557..3dd5661 100644 --- a/src/tint/lang/wgsl/inspector/inspector_test.cc +++ b/src/tint/lang/wgsl/inspector/inspector_test.cc
@@ -4049,7 +4049,7 @@ } TEST_F(InspectorGetBlendSrcTest, Basic) { - Enable(wgsl::Extension::kChromiumInternalDualSourceBlending); + Enable(wgsl::Extension::kDualSourceBlending); Structure("out_struct", Vector{
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..35308c6 100644 --- a/src/tint/lang/wgsl/ir_roundtrip_test.cc +++ b/src/tint/lang/wgsl/ir_roundtrip_test.cc
@@ -260,7 +260,7 @@ TEST_F(IRToProgramRoundtripTest, StructDecl_MemberIndex) { RUN_TEST(R"( -enable chromium_internal_dual_source_blending; +enable dual_source_blending; struct S { a : i32, @@ -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/parser/enable_directive_test.cc b/src/tint/lang/wgsl/reader/parser/enable_directive_test.cc index 34d0d00..3051677 100644 --- a/src/tint/lang/wgsl/reader/parser/enable_directive_test.cc +++ b/src/tint/lang/wgsl/reader/parser/enable_directive_test.cc
@@ -177,7 +177,7 @@ // Error when unknown extension found EXPECT_TRUE(p->has_error()); EXPECT_EQ(p->error(), R"(1:8: expected extension -Possible values: 'f16')"); +Possible values: 'dual_source_blending', 'f16')"); auto program = p->program(); auto& ast = program.AST(); EXPECT_EQ(ast.Enables().Length(), 0u); @@ -191,7 +191,7 @@ EXPECT_TRUE(p->has_error()); EXPECT_EQ(p->error(), R"(1:8: expected extension Did you mean 'f16'? -Possible values: 'f16')"); +Possible values: 'dual_source_blending', 'f16')"); auto program = p->program(); auto& ast = program.AST(); EXPECT_EQ(ast.Enables().Length(), 0u); @@ -205,7 +205,7 @@ // Error when unknown extension found EXPECT_TRUE(p->has_error()); EXPECT_EQ(p->error(), R"(1:8: expected extension -Possible values: 'chromium_disable_uniformity_analysis', 'chromium_experimental_framebuffer_fetch', 'chromium_experimental_pixel_local', 'chromium_experimental_push_constant', 'chromium_experimental_subgroups', 'chromium_internal_dual_source_blending', 'chromium_internal_graphite', 'chromium_internal_input_attachments', 'chromium_internal_relaxed_uniform_layout', 'f16')"); +Possible values: 'chromium_disable_uniformity_analysis', 'chromium_experimental_framebuffer_fetch', 'chromium_experimental_pixel_local', 'chromium_experimental_push_constant', 'chromium_experimental_subgroups', 'chromium_internal_dual_source_blending', 'chromium_internal_graphite', 'chromium_internal_input_attachments', 'chromium_internal_relaxed_uniform_layout', 'dual_source_blending', 'f16')"); auto program = p->program(); auto& ast = program.AST(); EXPECT_EQ(ast.Enables().Length(), 0u); @@ -253,7 +253,7 @@ p->translation_unit(); EXPECT_TRUE(p->has_error()); EXPECT_EQ(p->error(), R"(1:8: expected extension -Possible values: 'f16')"); +Possible values: 'dual_source_blending', 'f16')"); auto program = p->program(); auto& ast = program.AST(); EXPECT_EQ(ast.Enables().Length(), 0u); @@ -264,7 +264,7 @@ p->translation_unit(); EXPECT_TRUE(p->has_error()); EXPECT_EQ(p->error(), R"(1:8: expected extension -Possible values: 'f16')"); +Possible values: 'dual_source_blending', 'f16')"); auto program = p->program(); auto& ast = program.AST(); EXPECT_EQ(ast.Enables().Length(), 0u); @@ -276,7 +276,7 @@ EXPECT_TRUE(p->has_error()); EXPECT_EQ(p->error(), R"(1:8: expected extension Did you mean 'f16'? -Possible values: 'f16')"); +Possible values: 'dual_source_blending', 'f16')"); auto program = p->program(); auto& ast = program.AST(); EXPECT_EQ(ast.Enables().Length(), 0u);
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/attribute_validation_test.cc b/src/tint/lang/wgsl/resolver/attribute_validation_test.cc index 809c8e3..f7b37a4 100644 --- a/src/tint/lang/wgsl/resolver/attribute_validation_test.cc +++ b/src/tint/lang/wgsl/resolver/attribute_validation_test.cc
@@ -271,7 +271,7 @@ Enable(wgsl::Extension::kChromiumExperimentalFramebufferFetch); break; case AttributeKind::kBlendSrc: - Enable(wgsl::Extension::kChromiumInternalDualSourceBlending); + Enable(wgsl::Extension::kDualSourceBlending); break; case AttributeKind::kInputAttachmentIndex: Enable(wgsl::Extension::kChromiumInternalInputAttachments);
diff --git a/src/tint/lang/wgsl/resolver/dual_source_blending_extension_test.cc b/src/tint/lang/wgsl/resolver/dual_source_blending_extension_test.cc index 55c6704..e747333 100644 --- a/src/tint/lang/wgsl/resolver/dual_source_blending_extension_test.cc +++ b/src/tint/lang/wgsl/resolver/dual_source_blending_extension_test.cc
@@ -39,8 +39,7 @@ using DualSourceBlendingExtensionTest = ResolverTest; -// Using the @blend_src attribute without chromium_internal_dual_source_blending enabled should -// fail. +// Using the @blend_src attribute without dual_source_blending enabled should fail. TEST_F(DualSourceBlendingExtensionTest, UseBlendSrcAttribWithoutExtensionError) { Structure("Output", Vector{ Member("a", ty.vec4<f32>(), @@ -50,14 +49,12 @@ EXPECT_FALSE(r()->Resolve()); EXPECT_EQ( r()->error(), - R"(12:34 error: use of '@blend_src' requires enabling extension 'chromium_internal_dual_source_blending')"); + R"(12:34 error: use of '@blend_src' requires enabling extension 'dual_source_blending')"); } class DualSourceBlendingExtensionTests : public ResolverTest { public: - DualSourceBlendingExtensionTests() { - Enable(wgsl::Extension::kChromiumInternalDualSourceBlending); - } + DualSourceBlendingExtensionTests() { Enable(wgsl::Extension::kDualSourceBlending); } }; // Using an F32 as an index value should fail. @@ -260,9 +257,7 @@ class DualSourceBlendingExtensionTestWithParams : public ResolverTestWithParam<int> { public: - DualSourceBlendingExtensionTestWithParams() { - Enable(wgsl::Extension::kChromiumInternalDualSourceBlending); - } + DualSourceBlendingExtensionTestWithParams() { Enable(wgsl::Extension::kDualSourceBlending); } }; // Rendering to multiple render targets while using dual source blending should fail.
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 2b0449d..5387d09 100644 --- a/src/tint/lang/wgsl/resolver/input_attachments_extension_test.cc +++ b/src/tint/lang/wgsl/resolver/input_attachments_extension_test.cc
@@ -25,6 +25,9 @@ // OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +#include "gmock/gmock.h" +#include "src/tint/lang/core/number.h" +#include "src/tint/lang/wgsl/extension.h" #include "src/tint/lang/wgsl/resolver/resolver.h" #include "src/tint/lang/wgsl/resolver/resolver_helper_test.h" @@ -162,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/resolver/resolver.cc b/src/tint/lang/wgsl/resolver/resolver.cc index fe24ffc..56b70e5 100644 --- a/src/tint/lang/wgsl/resolver/resolver.cc +++ b/src/tint/lang/wgsl/resolver/resolver.cc
@@ -3072,6 +3072,12 @@ void Resolver::CollectTextureSamplerPairs(const sem::BuiltinFn* builtin, VectorRef<const sem::ValueExpression*> args) const { + if (builtin->Fn() == wgsl::BuiltinFn::kInputAttachmentLoad) { + // inputAttachmentLoad() is considered a texture function, however it doesn't need sampler, + // and its parameter has ParameterUsage::kInputAttachment, so return early. + return; + } + // Collect a texture/sampler pair for this builtin. const auto& signature = builtin->Signature(); int texture_index = signature.IndexOf(core::ParameterUsage::kTexture);
diff --git a/src/tint/lang/wgsl/resolver/resolver.h b/src/tint/lang/wgsl/resolver/resolver.h index 1781493..a1abde7 100644 --- a/src/tint/lang/wgsl/resolver/resolver.h +++ b/src/tint/lang/wgsl/resolver/resolver.h
@@ -40,6 +40,7 @@ #include "src/tint/lang/core/constant/eval.h" #include "src/tint/lang/core/constant/value.h" #include "src/tint/lang/core/intrinsic/table.h" +#include "src/tint/lang/core/type/input_attachment.h" #include "src/tint/lang/wgsl/common/allowed_features.h" #include "src/tint/lang/wgsl/intrinsic/dialect.h" #include "src/tint/lang/wgsl/program/program_builder.h"
diff --git a/src/tint/lang/wgsl/resolver/validator.cc b/src/tint/lang/wgsl/resolver/validator.cc index 559aa50..45fbe2b 100644 --- a/src/tint/lang/wgsl/resolver/validator.cc +++ b/src/tint/lang/wgsl/resolver/validator.cc
@@ -38,6 +38,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" @@ -2482,10 +2483,11 @@ bool Validator::BlendSrcAttribute(const ast::BlendSrcAttribute* attr, ast::PipelineStage stage, const std::optional<bool> is_input) const { - if (!enabled_extensions_.Contains(wgsl::Extension::kChromiumInternalDualSourceBlending)) { + if (!enabled_extensions_.Contains(wgsl::Extension::kChromiumInternalDualSourceBlending) && + !enabled_extensions_.Contains(wgsl::Extension::kDualSourceBlending)) { AddError(attr->source) << "use of " << style::Attribute("@blend_src") << " requires enabling extension " - << style::Code("chromium_internal_dual_source_blending"); + << style::Code("dual_source_blending"); return false; }
diff --git a/src/tint/lang/wgsl/resolver/validator.h b/src/tint/lang/wgsl/resolver/validator.h index fb54048..eff9ce6 100644 --- a/src/tint/lang/wgsl/resolver/validator.h +++ b/src/tint/lang/wgsl/resolver/validator.h
@@ -33,6 +33,8 @@ #include <utility> #include "src/tint/lang/core/evaluation_stage.h" +#include "src/tint/lang/core/type/input_attachment.h" +#include "src/tint/lang/wgsl/ast/input_attachment_index_attribute.h" #include "src/tint/lang/wgsl/ast/pipeline_stage.h" #include "src/tint/lang/wgsl/common/allowed_features.h" #include "src/tint/lang/wgsl/program/program_builder.h"
diff --git a/src/tint/lang/wgsl/wgsl.def b/src/tint/lang/wgsl/wgsl.def index 2fbd877..770c6c7 100644 --- a/src/tint/lang/wgsl/wgsl.def +++ b/src/tint/lang/wgsl/wgsl.def
@@ -70,6 +70,8 @@ enum extension { // WGSL Extension "f16" f16 + // WGSL Extension "dual_source_blending" + dual_source_blending // A Chromium-specific extension for disabling uniformity analysis. chromium_disable_uniformity_analysis // A Chromium-specific extension for push constants @@ -84,6 +86,8 @@ // A Chromium-specific extension that relaxes memory layout requirements for uniform storage. chromium_internal_relaxed_uniform_layout // A Chromium-specific extension that enables dual source blending. + // TODO(chromium:341973423): Remove `chromium_internal_dual_source_blending` when it is totally + // replaced with `dual_source_blending`. chromium_internal_dual_source_blending // A Chromium-specific extension that enables pixel local storage. // Cannot be used with chromium_experimental_framebuffer_fetch @@ -581,6 +585,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..e856dd4 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); } @@ -1004,7 +1014,7 @@ ast_attrs.Push(b.Location(u32(*location))); } if (auto blend_src = ir_attrs.blend_src) { - Enable(wgsl::Extension::kChromiumInternalDualSourceBlending); + Enable(wgsl::Extension::kDualSourceBlending); ast_attrs.Push(b.BlendSrc(u32(*blend_src))); } if (auto builtin = ir_attrs.builtin) {
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; }