resolver: Validate pipline stage use for intrinsics Use the new [[stage()]] decorations in intrinsics.def to validate that intrinsics are only called from the correct pipeline stages. Fixed: tint:657 Change-Id: I9efda26369c45c6f816bdaa53408d3909db403a1 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/53084 Reviewed-by: Ryan Harrison <rharrison@chromium.org> Reviewed-by: David Neto <dneto@google.com> Commit-Queue: Ben Clayton <bclayton@google.com> Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/src/BUILD.gn b/src/BUILD.gn index c39e38b..9498f0a 100644 --- a/src/BUILD.gn +++ b/src/BUILD.gn
@@ -511,6 +511,7 @@ "sem/node.h", "sem/parameter_usage.cc", "sem/parameter_usage.h", + "sem/pipeline_stage_set.h", "sem/pointer_type.cc", "sem/pointer_type.h", "sem/reference_type.cc",
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 7c6ed45..580c171 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt
@@ -254,6 +254,7 @@ sem/member_accessor_expression.cc sem/parameter_usage.cc sem/parameter_usage.h + sem/pipeline_stage_set.h sem/node.cc sem/node.h sem/statement.cc @@ -580,6 +581,7 @@ resolver/function_validation_test.cc resolver/host_shareable_validation_test.cc resolver/intrinsic_test.cc + resolver/intrinsic_validation_test.cc resolver/is_host_shareable_test.cc resolver/is_storeable_test.cc resolver/ptr_ref_test.cc
diff --git a/src/intrinsic_table.cc b/src/intrinsic_table.cc index 2073e80..4da8cbb 100644 --- a/src/intrinsic_table.cc +++ b/src/intrinsic_table.cc
@@ -23,6 +23,7 @@ #include "src/sem/depth_texture_type.h" #include "src/sem/external_texture_type.h" #include "src/sem/multisampled_texture_type.h" +#include "src/sem/pipeline_stage_set.h" #include "src/sem/sampled_texture_type.h" #include "src/sem/storage_texture_type.h" #include "src/utils/scoped_assignment.h" @@ -288,6 +289,8 @@ using AccessControl = ast::AccessControl::Access; using StorageClass = ast::StorageClass; using ParameterUsage = sem::ParameterUsage; +using PipelineStageSet = sem::PipelineStageSet; +using PipelineStage = ast::PipelineStage; bool match_bool(const sem::Type* ty) { return ty->IsAnyOf<Any, sem::Bool>(); @@ -608,6 +611,66 @@ return state.builder.create<sem::ExternalTexture>(); } +/// ParameterInfo describes a parameter +struct ParameterInfo { + /// The parameter usage (parameter name in definition file) + ParameterUsage const usage; + + /// Pointer to a list of indices that are used to match the parameter type. + /// The matcher indices index on Matchers::type and / or Matchers::number. + /// These indices are consumed by the matchers themselves. + /// The first index is always a TypeMatcher. + MatcherIndex const* const matcher_indices; +}; + +/// OpenTypeInfo describes an open type +struct OpenTypeInfo { + /// Name of the open type (e.g. 'T') + const char* name; + /// Optional type matcher constraint. + /// Either an index in Matchers::type, or kNoMatcher + MatcherIndex const matcher_index; +}; + +/// OpenNumberInfo describes an open number +struct OpenNumberInfo { + /// Name of the open number (e.g. 'N') + const char* name; + /// Optional number matcher constraint. + /// Either an index in Matchers::number, or kNoMatcher + MatcherIndex const matcher_index; +}; + +/// OverloadInfo describes a single function overload +struct OverloadInfo { + /// Total number of parameters for the overload + uint8_t const num_parameters; + /// Total number of open types for the overload + uint8_t const num_open_types; + /// Total number of open numbers for the overload + uint8_t const num_open_numbers; + /// Pointer to the first open type + OpenTypeInfo const* const open_types; + /// Pointer to the first open number + OpenNumberInfo const* const open_numbers; + /// Pointer to the first parameter + ParameterInfo const* const parameters; + /// Pointer to a list of matcher indices that index on Matchers::type and + /// Matchers::number, used to build the return type. If the function has no + /// return type then this is null. + MatcherIndex const* const return_matcher_indices; + /// The pipeline stages that this overload can be used in. + PipelineStageSet supported_stages; +}; + +/// IntrinsicInfo describes an intrinsic function +struct IntrinsicInfo { + /// Number of overloads of the intrinsic function + uint8_t const num_overloads; + /// Pointer to the start of the overloads for the function + OverloadInfo const* const overloads; +}; + #include "intrinsic_table.inl" /// Impl is the private implementation of the IntrinsicTable interface. @@ -807,9 +870,9 @@ return_type = builder.create<sem::Void>(); } - return builder.create<sem::Intrinsic>(intrinsic_type, - const_cast<sem::Type*>(return_type), - std::move(parameters)); + return builder.create<sem::Intrinsic>( + intrinsic_type, const_cast<sem::Type*>(return_type), + std::move(parameters), overload.supported_stages); } MatchState Impl::Match(ClosedState& closed,
diff --git a/src/intrinsic_table.inl b/src/intrinsic_table.inl index 87be0fd..925b23b 100644 --- a/src/intrinsic_table.inl +++ b/src/intrinsic_table.inl
@@ -24,64 +24,6 @@ // clang-format off -/// ParameterInfo describes a parameter -struct ParameterInfo { - /// The parameter usage (parameter name in definition file) - ParameterUsage const usage; - - /// Pointer to a list of indices that are used to match the parameter type. - /// The matcher indices index on Matchers::type and / or Matchers::number. - /// These indices are consumed by the matchers themselves. - /// The first index is always a TypeMatcher. - MatcherIndex const* const matcher_indices; -}; - -/// OpenTypeInfo describes an open type -struct OpenTypeInfo { - /// Name of the open type (e.g. 'T') - const char* name; - /// Optional type matcher constraint. - /// Either an index in Matchers::type, or kNoMatcher - MatcherIndex const matcher_index; -}; - -/// OpenNumberInfo describes an open number -struct OpenNumberInfo { - /// Name of the open number (e.g. 'N') - const char* name; - /// Optional number matcher constraint. - /// Either an index in Matchers::number, or kNoMatcher - MatcherIndex const matcher_index; -}; - -/// OverloadInfo describes a single function overload -struct OverloadInfo { - /// Total number of parameters for the overload - uint8_t const num_parameters; - /// Total number of open types for the overload - uint8_t const num_open_types; - /// Total number of open numbers for the overload - uint8_t const num_open_numbers; - /// Pointer to the first open type - OpenTypeInfo const* const open_types; - /// Pointer to the first open number - OpenNumberInfo const* const open_numbers; - /// Pointer to the first parameter - ParameterInfo const* const parameters; - /// Pointer to a list of matcher indices that index on Matchers::type and - /// Matchers::number, used to build the return type. If the function has no - /// return type then this is null. - MatcherIndex const* const return_matcher_indices; -}; - -/// IntrinsicInfo describes an intrinsic function -struct IntrinsicInfo { - /// Number of overloads of the intrinsic function - uint8_t const num_overloads; - /// Pointer to the start of the overloads for the function - OverloadInfo const* const overloads; -}; - /// TypeMatcher for 'type bool' /// @see src/intrinsics.def:68:6 class Bool : public TypeMatcher { @@ -4394,6 +4336,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[434], /* return matcher indices */ &kMatcherIndices[40], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [1] */ @@ -4404,6 +4347,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[432], /* return matcher indices */ &kMatcherIndices[113], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [2] */ @@ -4414,6 +4358,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[379], /* return matcher indices */ &kMatcherIndices[113], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [3] */ @@ -4424,6 +4369,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[433], /* return matcher indices */ &kMatcherIndices[113], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [4] */ @@ -4434,6 +4380,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[429], /* return matcher indices */ &kMatcherIndices[113], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [5] */ @@ -4444,6 +4391,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[435], /* return matcher indices */ &kMatcherIndices[75], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [6] */ @@ -4454,6 +4402,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[425], /* return matcher indices */ &kMatcherIndices[75], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [7] */ @@ -4464,6 +4413,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[437], /* return matcher indices */ &kMatcherIndices[75], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [8] */ @@ -4474,6 +4424,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[423], /* return matcher indices */ &kMatcherIndices[75], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [9] */ @@ -4484,6 +4435,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[439], /* return matcher indices */ &kMatcherIndices[75], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [10] */ @@ -4494,6 +4446,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[421], /* return matcher indices */ &kMatcherIndices[75], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [11] */ @@ -4504,6 +4457,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[441], /* return matcher indices */ &kMatcherIndices[113], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [12] */ @@ -4514,6 +4468,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[442], /* return matcher indices */ &kMatcherIndices[113], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [13] */ @@ -4524,6 +4479,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[419], /* return matcher indices */ &kMatcherIndices[113], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [14] */ @@ -4534,6 +4490,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[444], /* return matcher indices */ &kMatcherIndices[113], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [15] */ @@ -4544,6 +4501,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[417], /* return matcher indices */ &kMatcherIndices[113], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [16] */ @@ -4554,6 +4512,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[446], /* return matcher indices */ &kMatcherIndices[75], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [17] */ @@ -4564,6 +4523,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[415], /* return matcher indices */ &kMatcherIndices[75], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [18] */ @@ -4574,6 +4534,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[448], /* return matcher indices */ &kMatcherIndices[75], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [19] */ @@ -4584,6 +4545,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[413], /* return matcher indices */ &kMatcherIndices[75], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [20] */ @@ -4594,6 +4556,7 @@ /* open numbers */ &kOpenNumbers[0], /* parameters */ &kParameters[450], /* return matcher indices */ &kMatcherIndices[40], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [21] */ @@ -4604,6 +4567,7 @@ /* open numbers */ &kOpenNumbers[0], /* parameters */ &kParameters[451], /* return matcher indices */ &kMatcherIndices[113], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [22] */ @@ -4614,6 +4578,7 @@ /* open numbers */ &kOpenNumbers[0], /* parameters */ &kParameters[452], /* return matcher indices */ &kMatcherIndices[113], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [23] */ @@ -4624,6 +4589,7 @@ /* open numbers */ &kOpenNumbers[0], /* parameters */ &kParameters[453], /* return matcher indices */ &kMatcherIndices[75], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [24] */ @@ -4634,6 +4600,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[454], /* return matcher indices */ &kMatcherIndices[113], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [25] */ @@ -4644,6 +4611,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[286], /* return matcher indices */ &kMatcherIndices[47], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [26] */ @@ -4654,6 +4622,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[289], /* return matcher indices */ &kMatcherIndices[47], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [27] */ @@ -4664,6 +4633,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[150], /* return matcher indices */ &kMatcherIndices[47], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [28] */ @@ -4674,6 +4644,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[292], /* return matcher indices */ &kMatcherIndices[47], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [29] */ @@ -4684,6 +4655,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[295], /* return matcher indices */ &kMatcherIndices[47], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [30] */ @@ -4694,6 +4666,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[304], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [31] */ @@ -4704,6 +4677,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[194], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [32] */ @@ -4714,6 +4688,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[393], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [33] */ @@ -4724,6 +4699,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[395], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [34] */ @@ -4734,6 +4710,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[313], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [35] */ @@ -4744,6 +4721,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[397], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [36] */ @@ -4754,6 +4732,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[405], /* return matcher indices */ &kMatcherIndices[39], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [37] */ @@ -4764,6 +4743,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[407], /* return matcher indices */ &kMatcherIndices[39], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [38] */ @@ -4774,6 +4754,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[322], /* return matcher indices */ &kMatcherIndices[39], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [39] */ @@ -4784,6 +4765,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[409], /* return matcher indices */ &kMatcherIndices[39], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [40] */ @@ -4794,6 +4776,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[411], /* return matcher indices */ &kMatcherIndices[43], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [41] */ @@ -4804,6 +4787,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[427], /* return matcher indices */ &kMatcherIndices[43], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [42] */ @@ -4814,6 +4798,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[331], /* return matcher indices */ &kMatcherIndices[43], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [43] */ @@ -4824,6 +4809,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[403], /* return matcher indices */ &kMatcherIndices[43], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [44] */ @@ -4834,6 +4820,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[369], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [45] */ @@ -4844,6 +4831,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[328], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [46] */ @@ -4854,6 +4842,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[319], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [47] */ @@ -4864,6 +4853,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[234], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [48] */ @@ -4874,6 +4864,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[214], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [49] */ @@ -4884,6 +4875,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[60], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [50] */ @@ -4894,6 +4886,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[307], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [51] */ @@ -4904,6 +4897,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[170], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [52] */ @@ -4914,6 +4908,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[277], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [53] */ @@ -4924,6 +4919,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[154], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [54] */ @@ -4934,6 +4930,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[274], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [55] */ @@ -4944,6 +4941,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[166], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [56] */ @@ -4954,6 +4952,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[190], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [57] */ @@ -4964,6 +4963,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[85], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [58] */ @@ -4974,6 +4974,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[316], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [59] */ @@ -4984,6 +4985,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[226], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [60] */ @@ -4994,6 +4996,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[325], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [61] */ @@ -5004,6 +5007,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[174], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [62] */ @@ -5014,6 +5018,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[100], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [63] */ @@ -5024,6 +5029,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[90], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [64] */ @@ -5034,6 +5040,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[43], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [65] */ @@ -5044,6 +5051,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[218], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [66] */ @@ -5054,6 +5062,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[135], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [67] */ @@ -5064,6 +5073,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[202], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [68] */ @@ -5074,6 +5084,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[95], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [69] */ @@ -5084,6 +5095,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[222], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [70] */ @@ -5094,6 +5106,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[125], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [71] */ @@ -5104,6 +5117,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[65], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [72] */ @@ -5114,6 +5128,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[37], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [73] */ @@ -5124,6 +5139,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[198], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [74] */ @@ -5134,6 +5150,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[75], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [75] */ @@ -5144,6 +5161,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[256], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [76] */ @@ -5154,6 +5172,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[340], /* return matcher indices */ nullptr, + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [77] */ @@ -5164,6 +5183,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[259], /* return matcher indices */ nullptr, + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [78] */ @@ -5174,6 +5194,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[178], /* return matcher indices */ nullptr, + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [79] */ @@ -5184,6 +5205,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[262], /* return matcher indices */ nullptr, + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [80] */ @@ -5194,6 +5216,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[238], /* return matcher indices */ nullptr, + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [81] */ @@ -5204,6 +5227,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[346], /* return matcher indices */ nullptr, + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [82] */ @@ -5214,6 +5238,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[162], /* return matcher indices */ nullptr, + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [83] */ @@ -5224,6 +5249,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[265], /* return matcher indices */ nullptr, + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [84] */ @@ -5234,6 +5260,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[268], /* return matcher indices */ nullptr, + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [85] */ @@ -5244,6 +5271,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[271], /* return matcher indices */ nullptr, + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [86] */ @@ -5254,6 +5282,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[158], /* return matcher indices */ nullptr, + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [87] */ @@ -5264,6 +5293,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[283], /* return matcher indices */ nullptr, + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [88] */ @@ -5274,6 +5304,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[431], /* return matcher indices */ &kMatcherIndices[40], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [89] */ @@ -5284,6 +5315,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[461], /* return matcher indices */ &kMatcherIndices[40], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [90] */ @@ -5294,6 +5326,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[462], /* return matcher indices */ &kMatcherIndices[40], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [91] */ @@ -5304,6 +5337,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[463], /* return matcher indices */ &kMatcherIndices[40], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [92] */ @@ -5314,6 +5348,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[464], /* return matcher indices */ &kMatcherIndices[40], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [93] */ @@ -5324,6 +5359,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[465], /* return matcher indices */ &kMatcherIndices[40], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [94] */ @@ -5334,6 +5370,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[466], /* return matcher indices */ &kMatcherIndices[40], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [95] */ @@ -5344,6 +5381,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[467], /* return matcher indices */ &kMatcherIndices[40], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [96] */ @@ -5354,6 +5392,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[468], /* return matcher indices */ &kMatcherIndices[40], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [97] */ @@ -5364,6 +5403,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[230], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [98] */ @@ -5374,6 +5414,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[105], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [99] */ @@ -5384,6 +5425,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[130], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [100] */ @@ -5394,6 +5436,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[49], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [101] */ @@ -5404,6 +5447,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[210], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [102] */ @@ -5414,6 +5458,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[80], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [103] */ @@ -5424,6 +5469,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[186], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [104] */ @@ -5434,6 +5480,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[120], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [105] */ @@ -5444,6 +5491,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[115], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [106] */ @@ -5454,6 +5502,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[7], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [107] */ @@ -5464,6 +5513,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[19], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [108] */ @@ -5474,6 +5524,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[0], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [109] */ @@ -5484,6 +5535,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[145], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [110] */ @@ -5494,6 +5546,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[31], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [111] */ @@ -5504,6 +5557,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[140], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [112] */ @@ -5514,6 +5568,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[13], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [113] */ @@ -5524,6 +5579,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[206], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [114] */ @@ -5534,6 +5590,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[70], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [115] */ @@ -5544,6 +5601,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[55], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [116] */ @@ -5554,6 +5612,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[25], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [117] */ @@ -5564,6 +5623,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[182], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [118] */ @@ -5574,6 +5634,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[110], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [119] */ @@ -5584,6 +5645,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[455], /* return matcher indices */ &kMatcherIndices[40], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [120] */ @@ -5594,6 +5656,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[456], /* return matcher indices */ &kMatcherIndices[40], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [121] */ @@ -5604,6 +5667,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[457], /* return matcher indices */ &kMatcherIndices[40], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [122] */ @@ -5614,6 +5678,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[458], /* return matcher indices */ &kMatcherIndices[40], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [123] */ @@ -5624,6 +5689,7 @@ /* open numbers */ &kOpenNumbers[0], /* parameters */ &kParameters[459], /* return matcher indices */ &kMatcherIndices[40], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [124] */ @@ -5634,6 +5700,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[391], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [125] */ @@ -5644,6 +5711,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[389], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [126] */ @@ -5654,6 +5722,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[537], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [127] */ @@ -5664,6 +5733,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[536], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [128] */ @@ -5674,6 +5744,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[535], /* return matcher indices */ &kMatcherIndices[3], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [129] */ @@ -5684,6 +5755,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[534], /* return matcher indices */ &kMatcherIndices[2], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [130] */ @@ -5694,6 +5766,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[539], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [131] */ @@ -5704,6 +5777,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[538], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [132] */ @@ -5714,6 +5788,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[399], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [133] */ @@ -5724,6 +5799,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[401], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [134] */ @@ -5734,6 +5810,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[385], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [135] */ @@ -5744,6 +5821,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[383], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [136] */ @@ -5754,6 +5832,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[301], /* return matcher indices */ &kMatcherIndices[3], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [137] */ @@ -5764,6 +5843,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[298], /* return matcher indices */ &kMatcherIndices[2], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [138] */ @@ -5774,6 +5854,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[532], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [139] */ @@ -5784,6 +5865,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[531], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [140] */ @@ -5794,6 +5876,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[530], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [141] */ @@ -5804,6 +5887,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[529], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [142] */ @@ -5814,6 +5898,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[528], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [143] */ @@ -5824,6 +5909,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[527], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [144] */ @@ -5834,6 +5920,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[526], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [145] */ @@ -5844,6 +5931,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[525], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [146] */ @@ -5854,6 +5942,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[524], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [147] */ @@ -5864,6 +5953,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[523], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [148] */ @@ -5874,6 +5964,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[522], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [149] */ @@ -5884,6 +5975,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[521], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kFragment), }, { /* [150] */ @@ -5894,6 +5986,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[520], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [151] */ @@ -5904,6 +5997,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[519], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [152] */ @@ -5914,6 +6008,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[518], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [153] */ @@ -5924,6 +6019,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[517], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [154] */ @@ -5934,6 +6030,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[253], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [155] */ @@ -5944,6 +6041,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[250], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [156] */ @@ -5954,6 +6052,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[516], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [157] */ @@ -5964,6 +6063,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[515], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [158] */ @@ -5974,6 +6074,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[280], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [159] */ @@ -5984,6 +6085,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[241], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [160] */ @@ -5994,6 +6096,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[514], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [161] */ @@ -6004,6 +6107,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[513], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [162] */ @@ -6014,6 +6118,7 @@ /* open numbers */ &kOpenNumbers[3], /* parameters */ &kParameters[377], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [163] */ @@ -6024,6 +6129,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[375], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [164] */ @@ -6034,6 +6140,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[512], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [165] */ @@ -6044,6 +6151,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[511], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [166] */ @@ -6054,6 +6162,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[510], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [167] */ @@ -6064,6 +6173,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[509], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [168] */ @@ -6074,6 +6184,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[508], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [169] */ @@ -6084,6 +6195,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[507], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [170] */ @@ -6094,6 +6206,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[506], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [171] */ @@ -6104,6 +6217,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[505], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [172] */ @@ -6114,6 +6228,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[504], /* return matcher indices */ &kMatcherIndices[1], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [173] */ @@ -6124,6 +6239,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[503], /* return matcher indices */ &kMatcherIndices[24], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [174] */ @@ -6134,6 +6250,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[502], /* return matcher indices */ &kMatcherIndices[1], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [175] */ @@ -6144,6 +6261,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[501], /* return matcher indices */ &kMatcherIndices[24], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [176] */ @@ -6154,6 +6272,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[500], /* return matcher indices */ &kMatcherIndices[1], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [177] */ @@ -6164,6 +6283,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[499], /* return matcher indices */ &kMatcherIndices[24], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [178] */ @@ -6174,6 +6294,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[498], /* return matcher indices */ &kMatcherIndices[1], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [179] */ @@ -6184,6 +6305,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[497], /* return matcher indices */ &kMatcherIndices[24], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [180] */ @@ -6194,6 +6316,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[373], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [181] */ @@ -6204,6 +6327,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[371], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [182] */ @@ -6214,6 +6338,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[496], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [183] */ @@ -6224,6 +6349,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[495], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [184] */ @@ -6234,6 +6360,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[494], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [185] */ @@ -6244,6 +6371,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[493], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [186] */ @@ -6254,6 +6382,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[492], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [187] */ @@ -6264,6 +6393,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[491], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [188] */ @@ -6274,6 +6404,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[541], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [189] */ @@ -6284,6 +6415,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[540], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [190] */ @@ -6294,6 +6426,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[355], /* return matcher indices */ &kMatcherIndices[3], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [191] */ @@ -6304,6 +6437,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[353], /* return matcher indices */ &kMatcherIndices[2], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [192] */ @@ -6314,6 +6448,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[334], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [193] */ @@ -6324,6 +6459,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[337], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [194] */ @@ -6334,6 +6470,7 @@ /* open numbers */ &kOpenNumbers[3], /* parameters */ &kParameters[351], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [195] */ @@ -6344,6 +6481,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[349], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [196] */ @@ -6354,6 +6492,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[543], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [197] */ @@ -6364,6 +6503,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[542], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [198] */ @@ -6374,6 +6514,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[474], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [199] */ @@ -6384,6 +6525,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[473], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [200] */ @@ -6394,6 +6536,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[545], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [201] */ @@ -6404,6 +6547,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[544], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [202] */ @@ -6414,6 +6558,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[449], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [203] */ @@ -6424,6 +6569,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[447], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [204] */ @@ -6434,6 +6580,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[470], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [205] */ @@ -6444,6 +6591,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[460], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [206] */ @@ -6454,6 +6602,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[472], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [207] */ @@ -6464,6 +6613,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[471], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [208] */ @@ -6474,6 +6624,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[359], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [209] */ @@ -6484,6 +6635,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[361], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [210] */ @@ -6494,6 +6646,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[363], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [211] */ @@ -6504,6 +6657,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[365], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [212] */ @@ -6514,6 +6668,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[550], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [213] */ @@ -6524,6 +6679,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[549], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [214] */ @@ -6534,6 +6690,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[482], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [215] */ @@ -6544,6 +6701,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[481], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [216] */ @@ -6554,6 +6712,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[247], /* return matcher indices */ &kMatcherIndices[3], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [217] */ @@ -6564,6 +6723,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[244], /* return matcher indices */ &kMatcherIndices[2], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [218] */ @@ -6574,6 +6734,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[480], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [219] */ @@ -6584,6 +6745,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[479], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [220] */ @@ -6594,6 +6756,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[478], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [221] */ @@ -6604,6 +6767,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[477], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [222] */ @@ -6614,6 +6778,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[476], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [223] */ @@ -6624,6 +6789,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[475], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [224] */ @@ -6634,6 +6800,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[310], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [225] */ @@ -6644,6 +6811,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[343], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [226] */ @@ -6654,6 +6822,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[484], /* return matcher indices */ &kMatcherIndices[3], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [227] */ @@ -6664,6 +6833,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[483], /* return matcher indices */ &kMatcherIndices[2], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [228] */ @@ -6674,6 +6844,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[367], /* return matcher indices */ &kMatcherIndices[3], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [229] */ @@ -6684,6 +6855,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[357], /* return matcher indices */ &kMatcherIndices[2], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [230] */ @@ -6694,6 +6866,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[552], /* return matcher indices */ &kMatcherIndices[3], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [231] */ @@ -6704,6 +6877,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[551], /* return matcher indices */ &kMatcherIndices[2], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [232] */ @@ -6714,6 +6888,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[488], /* return matcher indices */ &kMatcherIndices[28], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [233] */ @@ -6724,6 +6899,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[486], /* return matcher indices */ &kMatcherIndices[28], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [234] */ @@ -6734,6 +6910,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[487], /* return matcher indices */ &kMatcherIndices[28], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [235] */ @@ -6744,6 +6921,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[445], /* return matcher indices */ &kMatcherIndices[115], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [236] */ @@ -6754,6 +6932,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[443], /* return matcher indices */ &kMatcherIndices[115], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [237] */ @@ -6764,6 +6943,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[440], /* return matcher indices */ &kMatcherIndices[115], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [238] */ @@ -6774,6 +6954,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[438], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [239] */ @@ -6784,6 +6965,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[436], /* return matcher indices */ &kMatcherIndices[29], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [240] */ @@ -6794,6 +6976,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[553], /* return matcher indices */ nullptr, + /* supported_stages */ PipelineStageSet(PipelineStage::kCompute), }, { /* [241] */ @@ -6804,6 +6987,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[485], /* return matcher indices */ &kMatcherIndices[28], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [242] */ @@ -6814,6 +6998,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[489], /* return matcher indices */ &kMatcherIndices[28], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [243] */ @@ -6824,6 +7009,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[490], /* return matcher indices */ &kMatcherIndices[7], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [244] */ @@ -6834,6 +7020,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[469], /* return matcher indices */ &kMatcherIndices[40], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [245] */ @@ -6844,6 +7031,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[381], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [246] */ @@ -6854,6 +7042,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[533], /* return matcher indices */ &kMatcherIndices[9], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [247] */ @@ -6864,6 +7053,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[387], /* return matcher indices */ &kMatcherIndices[16], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [248] */ @@ -6874,6 +7064,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[546], /* return matcher indices */ &kMatcherIndices[28], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [249] */ @@ -6884,6 +7075,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[547], /* return matcher indices */ &kMatcherIndices[1], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [250] */ @@ -6894,6 +7086,7 @@ /* open numbers */ &kOpenNumbers[2], /* parameters */ &kParameters[548], /* return matcher indices */ &kMatcherIndices[1], + /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute), }, { /* [251] */ @@ -6904,6 +7097,7 @@ /* open numbers */ &kOpenNumbers[4], /* parameters */ &kParameters[553], /* return matcher indices */ nullptr, + /* supported_stages */ PipelineStageSet(PipelineStage::kCompute), }, };
diff --git a/src/intrinsic_table.inl.tmpl b/src/intrinsic_table.inl.tmpl index aa30ef7..fc3a0b7 100644 --- a/src/intrinsic_table.inl.tmpl +++ b/src/intrinsic_table.inl.tmpl
@@ -11,64 +11,6 @@ // clang-format off -/// ParameterInfo describes a parameter -struct ParameterInfo { - /// The parameter usage (parameter name in definition file) - ParameterUsage const usage; - - /// Pointer to a list of indices that are used to match the parameter type. - /// The matcher indices index on Matchers::type and / or Matchers::number. - /// These indices are consumed by the matchers themselves. - /// The first index is always a TypeMatcher. - MatcherIndex const* const matcher_indices; -}; - -/// OpenTypeInfo describes an open type -struct OpenTypeInfo { - /// Name of the open type (e.g. 'T') - const char* name; - /// Optional type matcher constraint. - /// Either an index in Matchers::type, or kNoMatcher - MatcherIndex const matcher_index; -}; - -/// OpenNumberInfo describes an open number -struct OpenNumberInfo { - /// Name of the open number (e.g. 'N') - const char* name; - /// Optional number matcher constraint. - /// Either an index in Matchers::number, or kNoMatcher - MatcherIndex const matcher_index; -}; - -/// OverloadInfo describes a single function overload -struct OverloadInfo { - /// Total number of parameters for the overload - uint8_t const num_parameters; - /// Total number of open types for the overload - uint8_t const num_open_types; - /// Total number of open numbers for the overload - uint8_t const num_open_numbers; - /// Pointer to the first open type - OpenTypeInfo const* const open_types; - /// Pointer to the first open number - OpenNumberInfo const* const open_numbers; - /// Pointer to the first parameter - ParameterInfo const* const parameters; - /// Pointer to a list of matcher indices that index on Matchers::type and - /// Matchers::number, used to build the return type. If the function has no - /// return type then this is null. - MatcherIndex const* const return_matcher_indices; -}; - -/// IntrinsicInfo describes an intrinsic function -struct IntrinsicInfo { - /// Number of overloads of the intrinsic function - uint8_t const num_overloads; - /// Pointer to the start of the overloads for the function - OverloadInfo const* const overloads; -}; - {{ with .Sem -}} {{ range .Types -}} {{ template "Type" . }} @@ -155,6 +97,10 @@ {{- if $o.ReturnMatcherIndicesOffset }} &kMatcherIndices[{{$o.ReturnMatcherIndicesOffset}}] {{- else }} nullptr {{- end }}, + /* supported_stages */ PipelineStageSet( +{{- range $i, $u := $o.CanBeUsedInStage.List -}} +{{- if $i -}}, {{end}}PipelineStage::k{{Title $u}} +{{- end }}), }, {{- end }} };
diff --git a/src/resolver/builtins_validation_test.cc b/src/resolver/builtins_validation_test.cc index fe2b59d..d526f1e 100644 --- a/src/resolver/builtins_validation_test.cc +++ b/src/resolver/builtins_validation_test.cc
@@ -12,6 +12,7 @@ // See the License for the specific language governing permissions and // limitations under the License. +#include "src/ast/call_statement.h" #include "src/resolver/resolver_test_helper.h" namespace tint { @@ -283,7 +284,8 @@ params.push_back(Expr(1.0f)); } auto* builtin = Call(name, params); - WrapInFunction(builtin); + Func("func", {}, ty.void_(), {create<ast::CallStatement>(builtin)}, + {create<ast::StageDecoration>(ast::PipelineStage::kFragment)}); EXPECT_TRUE(r()->Resolve()) << r()->error(); EXPECT_TRUE(TypeOf(builtin)->Is<sem::F32>()); @@ -298,7 +300,8 @@ params.push_back(vec2<f32>(1.0f, 1.0f)); } auto* builtin = Call(name, params); - WrapInFunction(builtin); + Func("func", {}, ty.void_(), {create<ast::CallStatement>(builtin)}, + {create<ast::StageDecoration>(ast::PipelineStage::kFragment)}); EXPECT_TRUE(r()->Resolve()) << r()->error(); EXPECT_TRUE(TypeOf(builtin)->is_float_vector()); @@ -313,7 +316,8 @@ params.push_back(vec3<f32>(1.0f, 1.0f, 1.0f)); } auto* builtin = Call(name, params); - WrapInFunction(builtin); + Func("func", {}, ty.void_(), {create<ast::CallStatement>(builtin)}, + {create<ast::StageDecoration>(ast::PipelineStage::kFragment)}); EXPECT_TRUE(r()->Resolve()) << r()->error(); EXPECT_TRUE(TypeOf(builtin)->is_float_vector()); @@ -328,7 +332,8 @@ params.push_back(vec4<f32>(1.0f, 1.0f, 1.0f, 1.0f)); } auto* builtin = Call(name, params); - WrapInFunction(builtin); + Func("func", {}, ty.void_(), {create<ast::CallStatement>(builtin)}, + {create<ast::StageDecoration>(ast::PipelineStage::kFragment)}); EXPECT_TRUE(r()->Resolve()) << r()->error(); EXPECT_TRUE(TypeOf(builtin)->is_float_vector());
diff --git a/src/resolver/intrinsic_test.cc b/src/resolver/intrinsic_test.cc index 8c27acf..ba8af09 100644 --- a/src/resolver/intrinsic_test.cc +++ b/src/resolver/intrinsic_test.cc
@@ -55,7 +55,8 @@ Global("ident", ty.f32(), ast::StorageClass::kInput); auto* expr = Call(name, "ident"); - WrapInFunction(expr); + Func("func", {}, ty.void_(), {create<ast::CallStatement>(expr)}, + {create<ast::StageDecoration>(ast::PipelineStage::kFragment)}); EXPECT_TRUE(r()->Resolve()) << r()->error(); @@ -68,7 +69,8 @@ Global("ident", ty.vec4<f32>(), ast::StorageClass::kInput); auto* expr = Call(name, "ident"); - WrapInFunction(expr); + Func("func", {}, ty.void_(), {create<ast::CallStatement>(expr)}, + {create<ast::StageDecoration>(ast::PipelineStage::kFragment)}); EXPECT_TRUE(r()->Resolve()) << r()->error(); @@ -1927,7 +1929,8 @@ param.buildSamplerVariable(this); auto* call = Call(param.function, param.args(this)); - WrapInFunction(call); + Func("func", {}, ty.void_(), {create<ast::CallStatement>(call)}, + {create<ast::StageDecoration>(ast::PipelineStage::kFragment)}); ASSERT_TRUE(r()->Resolve()) << r()->error();
diff --git a/src/resolver/intrinsic_validation_test.cc b/src/resolver/intrinsic_validation_test.cc new file mode 100644 index 0000000..31b5d3d --- /dev/null +++ b/src/resolver/intrinsic_validation_test.cc
@@ -0,0 +1,97 @@ +// Copyright 2021 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include "src/resolver/resolver.h" + +#include "gmock/gmock.h" +#include "src/ast/assignment_statement.h" +#include "src/ast/bitcast_expression.h" +#include "src/ast/break_statement.h" +#include "src/ast/call_statement.h" +#include "src/ast/continue_statement.h" +#include "src/ast/if_statement.h" +#include "src/ast/intrinsic_texture_helper_test.h" +#include "src/ast/loop_statement.h" +#include "src/ast/return_statement.h" +#include "src/ast/stage_decoration.h" +#include "src/ast/struct_block_decoration.h" +#include "src/ast/switch_statement.h" +#include "src/ast/unary_op_expression.h" +#include "src/ast/variable_decl_statement.h" +#include "src/resolver/resolver_test_helper.h" +#include "src/sem/call.h" +#include "src/sem/function.h" +#include "src/sem/member_accessor_expression.h" +#include "src/sem/sampled_texture_type.h" +#include "src/sem/statement.h" +#include "src/sem/variable.h" + +using ::testing::ElementsAre; +using ::testing::HasSubstr; + +namespace tint { +namespace resolver { +namespace { + +using IntrinsicType = sem::IntrinsicType; + +using ResolverIntrinsicValidationTest = ResolverTest; + +TEST_F(ResolverIntrinsicValidationTest, InvalidPipelineStageDirect) { + // [[stage(compute)]] fn func { return dpdx(1.0); } + + auto* dpdx = create<ast::CallExpression>(Source{{3, 4}}, Expr("dpdx"), + ast::ExpressionList{Expr(1.0f)}); + Func(Source{{1, 2}}, "func", ast::VariableList{}, ty.void_(), + {create<ast::CallStatement>(dpdx)}, + {Stage(ast::PipelineStage::kCompute)}); + + EXPECT_FALSE(r()->Resolve()); + EXPECT_EQ(r()->error(), + "3:4 error: built-in cannot be used by compute pipeline stage"); +} + +TEST_F(ResolverIntrinsicValidationTest, InvalidPipelineStageIndirect) { + // fn f0 { return dpdx(1.0); } + // fn f1 { f0(); } + // fn f2 { f1(); } + // [[stage(compute)]] fn main { return f2(); } + + auto* dpdx = create<ast::CallExpression>(Source{{3, 4}}, Expr("dpdx"), + ast::ExpressionList{Expr(1.0f)}); + Func(Source{{1, 2}}, "f0", ast::VariableList{}, ty.void_(), + {create<ast::CallStatement>(dpdx)}); + + Func(Source{{3, 4}}, "f1", ast::VariableList{}, ty.void_(), + {create<ast::CallStatement>(Call("f0"))}); + + Func(Source{{5, 6}}, "f2", ast::VariableList{}, ty.void_(), + {create<ast::CallStatement>(Call("f1"))}); + + Func(Source{{7, 8}}, "main", ast::VariableList{}, ty.void_(), + {create<ast::CallStatement>(Call("f2"))}, + {Stage(ast::PipelineStage::kCompute)}); + + EXPECT_FALSE(r()->Resolve()); + EXPECT_EQ(r()->error(), + R"(3:4 error: built-in cannot be used by compute pipeline stage +1:2 note: called by function 'f0' +3:4 note: called by function 'f1' +5:6 note: called by function 'f2' +7:8 note: called by entry point 'main')"); +} + +} // namespace +} // namespace resolver +} // namespace tint
diff --git a/src/resolver/resolver.cc b/src/resolver/resolver.cc index 2ac42f9..7b59942 100644 --- a/src/resolver/resolver.cc +++ b/src/resolver/resolver.cc
@@ -237,6 +237,10 @@ } } + if (!ValidatePipelineStages()) { + return false; + } + bool result = true; for (auto* node : builder_->ASTNodes().Objects()) { @@ -1129,6 +1133,10 @@ bool Resolver::Function(ast::Function* func) { auto* info = function_infos_.Create<FunctionInfo>(func); + if (func->IsEntryPoint()) { + entry_points_.emplace_back(info); + } + TINT_SCOPED_ASSIGNMENT(current_function_, info); variable_stack_.push_scope(); @@ -1707,6 +1715,10 @@ builder_->Sem().Add( call, builder_->create<sem::Call>(call, result, current_statement_)); SetType(call, result->ReturnType()); + + current_function_->intrinsic_calls.emplace_back( + IntrinsicCallInfo{call, result}); + return true; } @@ -2460,25 +2472,76 @@ expr_info_.emplace(expr, ExpressionInfo{type, type_name, current_statement_}); } +bool Resolver::ValidatePipelineStages() { + auto check_intrinsic_calls = [&](FunctionInfo* func, + FunctionInfo* entry_point) { + auto stage = entry_point->declaration->pipeline_stage(); + for (auto& call : func->intrinsic_calls) { + if (!call.intrinsic->SupportedStages().Contains(stage)) { + std::stringstream err; + err << "built-in cannot be used by " << stage << " pipeline stage"; + diagnostics_.add_error(err.str(), call.call->source()); + if (func != entry_point) { + TraverseCallChain(entry_point, func, [&](FunctionInfo* f) { + diagnostics_.add_note( + "called by function '" + + builder_->Symbols().NameFor(f->declaration->symbol()) + "'", + f->declaration->source()); + }); + diagnostics_.add_note("called by entry point '" + + builder_->Symbols().NameFor( + entry_point->declaration->symbol()) + + "'", + entry_point->declaration->source()); + } + return false; + } + } + return true; + }; + + for (auto* entry_point : entry_points_) { + if (!check_intrinsic_calls(entry_point, entry_point)) { + return false; + } + for (auto* func : entry_point->transitive_calls) { + if (!check_intrinsic_calls(func, entry_point)) { + return false; + } + } + } + return true; +} + +template <typename CALLBACK> +void Resolver::TraverseCallChain(FunctionInfo* from, + FunctionInfo* to, + CALLBACK&& callback) const { + for (auto* f : from->transitive_calls) { + if (f == to) { + callback(f); + return; + } + if (f->transitive_calls.contains(to)) { + TraverseCallChain(f, to, callback); + callback(f); + return; + } + } + TINT_ICE(diagnostics_) + << "TraverseCallChain() 'from' does not transitively call 'to'"; +} + void Resolver::CreateSemanticNodes() const { auto& sem = builder_->Sem(); // Collate all the 'ancestor_entry_points' - this is a map of function // symbol to all the entry points that transitively call the function. std::unordered_map<Symbol, std::vector<Symbol>> ancestor_entry_points; - for (auto* func : builder_->AST().Functions()) { - auto it = function_to_info_.find(func); - if (it == function_to_info_.end()) { - continue; // Resolver has likely errored. Process what we can. - } - - auto* info = it->second; - if (!func->IsEntryPoint()) { - continue; - } - for (auto* call : info->transitive_calls) { + for (auto* entry_point : entry_points_) { + for (auto* call : entry_point->transitive_calls) { auto& vec = ancestor_entry_points[call->declaration->symbol()]; - vec.emplace_back(func->symbol()); + vec.emplace_back(entry_point->declaration->symbol()); } }
diff --git a/src/resolver/resolver.h b/src/resolver/resolver.h index 252f8bc..2efdf7c 100644 --- a/src/resolver/resolver.h +++ b/src/resolver/resolver.h
@@ -51,6 +51,7 @@ } // namespace ast namespace sem { class Array; +class Intrinsic; class Statement; } // namespace sem @@ -100,6 +101,11 @@ sem::BindingPoint binding_point; }; + struct IntrinsicCallInfo { + const ast::CallExpression* call; + const sem::Intrinsic* intrinsic; + }; + /// Structure holding semantic information about a function. /// Used to build the sem::Function nodes at the end of resolving. struct FunctionInfo { @@ -115,9 +121,13 @@ sem::Type* return_type = nullptr; std::string return_type_name; std::array<sem::WorkgroupDimension, 3> workgroup_size; + std::vector<IntrinsicCallInfo> intrinsic_calls; // List of transitive calls this function makes UniqueVector<FunctionInfo*> transitive_calls; + + // List of entry point functions that transitively call this function + UniqueVector<FunctionInfo*> ancestor_entry_points; }; /// Structure holding semantic information about an expression. @@ -183,6 +193,8 @@ /// @returns true on success, false on error bool ResolveInternal(); + bool ValidatePipelineStages(); + /// Creates the nodes and adds them to the sem::Info mappings of the /// ProgramBuilder. void CreateSemanticNodes() const; @@ -359,12 +371,18 @@ /// @param node the AST node. void Mark(const ast::Node* node); + template <typename CALLBACK> + void TraverseCallChain(FunctionInfo* from, + FunctionInfo* to, + CALLBACK&& callback) const; + ProgramBuilder* const builder_; diag::List& diagnostics_; std::unique_ptr<IntrinsicTable> const intrinsic_table_; sem::BlockStatement* current_block_ = nullptr; ScopeStack<VariableInfo*> variable_stack_; std::unordered_map<Symbol, FunctionInfo*> symbol_to_function_; + std::vector<FunctionInfo*> entry_points_; std::unordered_map<const ast::Function*, FunctionInfo*> function_to_info_; std::unordered_map<const ast::Variable*, VariableInfo*> variable_to_info_; std::unordered_map<ast::CallExpression*, FunctionCallInfo> function_calls_;
diff --git a/src/sem/intrinsic.cc b/src/sem/intrinsic.cc index 6471f43..bdcae1d 100644 --- a/src/sem/intrinsic.cc +++ b/src/sem/intrinsic.cc
@@ -89,8 +89,11 @@ Intrinsic::Intrinsic(IntrinsicType type, sem::Type* return_type, - const ParameterList& parameters) - : Base(return_type, parameters), type_(type) {} + const ParameterList& parameters, + PipelineStageSet supported_stages) + : Base(return_type, parameters), + type_(type), + supported_stages_(supported_stages) {} Intrinsic::~Intrinsic() = default;
diff --git a/src/sem/intrinsic.h b/src/sem/intrinsic.h index 2268680..83a7f6d 100644 --- a/src/sem/intrinsic.h +++ b/src/sem/intrinsic.h
@@ -19,6 +19,7 @@ #include "src/sem/call_target.h" #include "src/sem/intrinsic_type.h" +#include "src/sem/pipeline_stage_set.h" namespace tint { namespace sem { @@ -75,9 +76,12 @@ /// @param type the intrinsic type /// @param return_type the return type for the intrinsic call /// @param parameters the parameters for the intrinsic overload + /// @param supported_stages the pipeline stages that this intrinsic can be + /// used in Intrinsic(IntrinsicType type, sem::Type* return_type, - const ParameterList& parameters); + const ParameterList& parameters, + PipelineStageSet supported_stages); /// Destructor ~Intrinsic() override; @@ -85,6 +89,9 @@ /// @return the type of the intrinsic IntrinsicType Type() const { return type_; } + /// @return the pipeline stages that this intrinsic can be used in + PipelineStageSet SupportedStages() const { return supported_stages_; } + /// @returns the name of the intrinsic function type. The spelling, including /// case, matches the name in the WGSL spec. const char* str() const; @@ -118,6 +125,7 @@ private: IntrinsicType const type_; + PipelineStageSet const supported_stages_; }; /// Emits the name of the intrinsic function type. The spelling, including case,
diff --git a/src/sem/pipeline_stage_set.h b/src/sem/pipeline_stage_set.h new file mode 100644 index 0000000..304a2e1 --- /dev/null +++ b/src/sem/pipeline_stage_set.h
@@ -0,0 +1,29 @@ +// Copyright 2021 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#ifndef SRC_SEM_PIPELINE_STAGE_SET_H_ +#define SRC_SEM_PIPELINE_STAGE_SET_H_ + +#include "src/ast/pipeline_stage.h" +#include "src/utils/enum_set.h" + +namespace tint { +namespace sem { + +using PipelineStageSet = utils::EnumSet<ast::PipelineStage>; + +} // namespace sem +} // namespace tint + +#endif // SRC_SEM_PIPELINE_STAGE_SET_H_
diff --git a/src/utils/unique_vector.h b/src/utils/unique_vector.h index 5f8ce34..e8359fc 100644 --- a/src/utils/unique_vector.h +++ b/src/utils/unique_vector.h
@@ -37,6 +37,10 @@ } } + /// @returns true if the vector contains `item` + /// @param item the item + bool contains(const T& item) const { return set.count(item); } + /// @returns the number of items in the vector size_t size() const { return vector.size(); } @@ -47,7 +51,7 @@ ConstIterator end() const { return vector.end(); } /// @returns a const reference to the internal vector - operator const std::vector<T>&() const { return vector; } + operator const std::vector<T> &() const { return vector; } private: std::vector<T> vector;
diff --git a/src/writer/hlsl/generator_impl_intrinsic_test.cc b/src/writer/hlsl/generator_impl_intrinsic_test.cc index 8660326..a6d79c6 100644 --- a/src/writer/hlsl/generator_impl_intrinsic_test.cc +++ b/src/writer/hlsl/generator_impl_intrinsic_test.cc
@@ -163,7 +163,8 @@ auto* call = GenerateCall(param.intrinsic, param.type, this); ASSERT_NE(nullptr, call) << "Unhandled intrinsic"; - WrapInFunction(call); + Func("func", {}, ty.void_(), {create<ast::CallStatement>(call)}, + {create<ast::StageDecoration>(ast::PipelineStage::kFragment)}); GeneratorImpl& gen = Build();
diff --git a/src/writer/msl/generator_impl_intrinsic_test.cc b/src/writer/msl/generator_impl_intrinsic_test.cc index 65d3ca5..7a635a3 100644 --- a/src/writer/msl/generator_impl_intrinsic_test.cc +++ b/src/writer/msl/generator_impl_intrinsic_test.cc
@@ -12,6 +12,7 @@ // See the License for the specific language governing permissions and // limitations under the License. +#include "src/ast/call_statement.h" #include "src/sem/call.h" #include "src/writer/msl/test_helper.h" @@ -176,7 +177,8 @@ auto* call = GenerateCall(param.intrinsic, param.type, this); ASSERT_NE(nullptr, call) << "Unhandled intrinsic"; - WrapInFunction(call); + Func("func", {}, ty.void_(), {create<ast::CallStatement>(call)}, + {create<ast::StageDecoration>(ast::PipelineStage::kFragment)}); GeneratorImpl& gen = Build();
diff --git a/src/writer/spirv/builder_intrinsic_test.cc b/src/writer/spirv/builder_intrinsic_test.cc index d030067..4e4238a 100644 --- a/src/writer/spirv/builder_intrinsic_test.cc +++ b/src/writer/spirv/builder_intrinsic_test.cc
@@ -414,7 +414,8 @@ auto* var = Global("v", ty.f32(), ast::StorageClass::kPrivate); auto* expr = Call(param.name, "v"); - WrapInFunction(expr); + Func("func", {}, ty.void_(), {create<ast::CallStatement>(expr)}, + {create<ast::StageDecoration>(ast::PipelineStage::kFragment)}); spirv::Builder& b = Build(); @@ -439,7 +440,8 @@ auto* var = Global("v", ty.vec3<f32>(), ast::StorageClass::kPrivate); auto* expr = Call(param.name, "v"); - WrapInFunction(expr); + Func("func", {}, ty.void_(), {create<ast::CallStatement>(expr)}, + {create<ast::StageDecoration>(ast::PipelineStage::kFragment)}); spirv::Builder& b = Build();
diff --git a/src/writer/spirv/builder_intrinsic_texture_test.cc b/src/writer/spirv/builder_intrinsic_texture_test.cc index bbb4600..78795a2 100644 --- a/src/writer/spirv/builder_intrinsic_texture_test.cc +++ b/src/writer/spirv/builder_intrinsic_texture_test.cc
@@ -3462,7 +3462,9 @@ auto* call = create<ast::CallExpression>(Expr(param.function), param.args(this)); - WrapInFunction(call); + + Func("func", {}, ty.void_(), {create<ast::CallStatement>(call)}, + {create<ast::StageDecoration>(ast::PipelineStage::kFragment)}); spirv::Builder& b = Build(); @@ -3515,7 +3517,8 @@ auto* call = create<ast::CallExpression>(Expr(param.function), param.args(this)); - WrapInFunction(call); + Func("func", {}, ty.void_(), {create<ast::CallStatement>(call)}, + {create<ast::StageDecoration>(ast::PipelineStage::kFragment)}); spirv::Builder& b = Build();
diff --git a/test/BUILD.gn b/test/BUILD.gn index ef868a6..ac376f8 100644 --- a/test/BUILD.gn +++ b/test/BUILD.gn
@@ -232,6 +232,7 @@ "../src/resolver/function_validation_test.cc", "../src/resolver/host_shareable_validation_test.cc", "../src/resolver/intrinsic_test.cc", + "../src/resolver/intrinsic_validation_test.cc", "../src/resolver/is_host_shareable_test.cc", "../src/resolver/is_storeable_test.cc", "../src/resolver/pipeline_overridable_constant_test.cc",
diff --git a/tools/src/cmd/intrinsic-gen/gen/intrinsic_table.go b/tools/src/cmd/intrinsic-gen/gen/intrinsic_table.go index 84fa3ee..82dfa0d 100644 --- a/tools/src/cmd/intrinsic-gen/gen/intrinsic_table.go +++ b/tools/src/cmd/intrinsic-gen/gen/intrinsic_table.go
@@ -96,6 +96,8 @@ // These indices are consumed by the matchers themselves. // The first index is always a TypeMatcher. ReturnMatcherIndicesOffset *int + // StageUses describes the stages an overload can be used in + CanBeUsedInStage sem.StageUses } // Function is used to create the C++ IntrinsicInfo structure @@ -193,6 +195,7 @@ OpenNumbersOffset: b.lut.openNumbers.Add(ob.openNumbers), ParametersOffset: b.lut.parameters.Add(ob.parameters), ReturnMatcherIndicesOffset: ob.returnTypeMatcherIndicesOffset, + CanBeUsedInStage: o.CanBeUsedInStage, }, nil }
diff --git a/tools/src/cmd/intrinsic-gen/sem/sem.go b/tools/src/cmd/intrinsic-gen/sem/sem.go index 36548ae..cdc91bf 100644 --- a/tools/src/cmd/intrinsic-gen/sem/sem.go +++ b/tools/src/cmd/intrinsic-gen/sem/sem.go
@@ -146,6 +146,21 @@ Compute bool } +// List returns the stage uses as a string list +func (u StageUses) List() []string { + out := []string{} + if u.Vertex { + out = append(out, "vertex") + } + if u.Fragment { + out = append(out, "fragment") + } + if u.Compute { + out = append(out, "compute") + } + return out +} + // Format implements the fmt.Formatter interface func (o Overload) Format(w fmt.State, verb rune) { fmt.Fprintf(w, "fn %v", o.Function.Name)