diff --git a/csrc/device_lower/id_model_options.h b/csrc/device_lower/id_model_options.h index 180ae91ddb0..ae1ceed5be3 100644 --- a/csrc/device_lower/id_model_options.h +++ b/csrc/device_lower/id_model_options.h @@ -7,146 +7,37 @@ // clang-format on #pragma once -#include - #include +#include "id_model/utils.h" +#include "options.h" + namespace nvfuser { class IdModelOptions { public: IdModelOptions() - : build_id_model_(!isOptionDisabled(DisableOption::IdModel)), - consumer_index_( - isIdModelOptionEnabled(IdModelEnableOption::ConsumerIndex)), - producer_index_( - isIdModelOptionEnabled(IdModelEnableOption::ProducerIndex)), - inline_predicate_( - isIdModelOptionEnabled(IdModelEnableOption::InlinePredicate)), - unswitch_predicate_( - isIdModelOptionEnabled(IdModelEnableOption::UnswitchPredicate)), - loop_(isIdModelOptionEnabled(IdModelEnableOption::Loop)) { - ensureConsistency(); - } - - bool buildIdModel() const { - return build_id_model_; - } - - void setBuildIdModel(bool b) { - build_id_model_ = b; - ensureConsistency(); - } - - bool buildTensorIndexer() const { - return build_tensor_indexer_; - } - - void setBuildTensorIndexer(bool b) { - build_tensor_indexer_ = b; - ensureConsistency(); - } - - bool consumerIndex() const { - return consumer_index_; - } - - void setConsumerIndex(bool b) { - consumer_index_ = b; - ensureConsistency(); - } - - bool producerIndex() const { - return producer_index_; - } + : tensor_indexer_enabled_(isOptionEnabled(EnableOption::IdModel)) {} - void setProducerIndex(bool b) { - producer_index_ = b; - ensureConsistency(); + void setTensorIndexer(bool b) { + tensor_indexer_enabled_ = b; } - void setIndex(bool b) { - setConsumerIndex(b); - setProducerIndex(b); - } - - bool inlinePredicate() const { - return inline_predicate_; - } - - void setInlinePredicate(bool b) { - inline_predicate_ = b; - ensureConsistency(); - } - - bool unswitchPredicate() const { - return unswitch_predicate_; - } - - void setUnswitchPredicate(bool b) { - unswitch_predicate_ = b; - ensureConsistency(); - } - - void setPredicate(bool b) { - setInlinePredicate(b); - setUnswitchPredicate(b); - } - - bool loop() const { - return loop_; - } - - void setLoop(bool b) { - loop_ = b; - ensureConsistency(); + bool isTensorIndexerEnabled() const { + return tensor_indexer_enabled_; } std::string toString() const { auto bool2str = [](bool b) { return b ? "true" : "false"; }; std::stringstream ss; - ss << "build_id_model=" << bool2str(build_id_model_) - << ", build_tensor_indexer=" << bool2str(build_tensor_indexer_) - << ", consumer_index=" << bool2str(consumer_index_) - << ", producer_index=" << bool2str(producer_index_) - << ", inline_predicate=" << bool2str(inline_predicate_) - << ", unswitch_predicate=" << bool2str(unswitch_predicate_) - << ", loop=" << bool2str(loop_); + ss << "enable_tensor_indexer=" << bool2str(tensor_indexer_enabled_); return ss.str(); } private: - void ensureConsistency() { - if (!build_id_model_) { - build_tensor_indexer_ = false; - consumer_index_ = false; - producer_index_ = false; - inline_predicate_ = false; - unswitch_predicate_ = false; - loop_ = false; - } else { - // TensorIndexer is required if these options are enabled - build_tensor_indexer_ = build_tensor_indexer_ || consumer_index_ || - producer_index_ || inline_predicate_ || unswitch_predicate_ || loop_; - } - } - - private: - // Build IdModel - bool build_id_model_ = true; - // Build TensorIndexer - bool build_tensor_indexer_ = false; - // Globally enables consumer indexing. - bool consumer_index_ = false; - // Globally enables producer indexing. - bool producer_index_ = false; - // Globally enables inline predicate - bool inline_predicate_ = false; - // Globally enables unswitch predicate - bool unswitch_predicate_ = false; - // Generate loops using IdModel - bool loop_ = false; + // Enable TensorIndexer + bool tensor_indexer_enabled_ = false; }; } // namespace nvfuser diff --git a/csrc/device_lower/lower2device.cpp b/csrc/device_lower/lower2device.cpp index fe244d2d7da..942154b1827 100644 --- a/csrc/device_lower/lower2device.cpp +++ b/csrc/device_lower/lower2device.cpp @@ -298,24 +298,13 @@ IdModelOptions getIdModelOptions(Fusion* fusion) { for (auto expr : fusion->exprs()) { if (auto ldst = dynamic_cast(expr)) { - if (ldst->opType() == LoadStoreOpType::CpAsyncBulkTensorTile || - ldst->opType() == LoadStoreOpType::CpAsyncBulk) { - options.setBuildTensorIndexer(true); - if (ldst->opType() == LoadStoreOpType::CpAsyncBulk) { - options.setInlinePredicate(true); - } + if (ldst->opType() == LoadStoreOpType::CpAsyncBulk) { + options.setTensorIndexer(true); continue; } - } else if (expr->isA()) { - options.setBuildTensorIndexer(true); - continue; } else if ( expr->isOneOf()) { - options.setProducerIndex(true); - options.setConsumerIndex(true); - options.setInlinePredicate(true); - options.setUnswitchPredicate(true); - options.setLoop(true); + options.setTensorIndexer(true); continue; } else if (auto reshape = dynamic_cast(expr)) { // The legacy indexer has an issue when an expand broadcast is @@ -369,31 +358,16 @@ IdModelOptions getIdModelOptions(Fusion* fusion) { return consumer_expanded_root_ids.count(input); }); })) { - options.setProducerIndex(true); - options.setConsumerIndex(true); - options.setInlinePredicate(true); - options.setUnswitchPredicate(true); + options.setTensorIndexer(true); } } } - // If a tensor does not have a nice root->logical/allocation->loop - // linear transformation history, use TensorIndexer - for (auto tv : fusion->allTvs()) { - if (tv->getMemoryType() == MemoryType::Tensor || - !ir_utils::hasRootToLoopLinearTransformations(tv)) { - options.setBuildTensorIndexer(true); - } - } - // If not supported, disable use of TensorIndexer by default. It is // still used if explicitly opted-in (see, for example, // Index::getConsumerIndex) if (!TensorIndexer::isSupported(fusion)) { - // Do not disable building of TensorIndexer as it may be still used - options.setIndex(false); - options.setPredicate(false); - options.setLoop(false); + options.setTensorIndexer(false); } return options; @@ -461,14 +435,12 @@ void GpuLower::analysis(Fusion* fusion) { // New IterDomains may be created, so it is expected that generated // code may use diffrent variable names - if (idModelOptions().buildIdModel()) { - info().set(std::make_unique( - fusion_, - /*build_graphs=*/true, - /*allow_self_mapping=*/false, - /*validate=*/false)); - info().idModel().validateAndPropagatePType(); - } + info().set(std::make_unique( + fusion_, + /*build_graphs=*/true, + /*allow_self_mapping=*/false, + /*validate=*/false)); + info().idModel().validateAndPropagatePType(); // Build what's refered to as the compute at map. This map contains the // mappings of all iteration domains across the fusion. There are three types @@ -578,16 +550,15 @@ void GpuLower::analysis(Fusion* fusion) { info().caMap().allocateIndexVariables(); dumpExprsIfEnabled(fusion_->exprs(), "allocateIndexVariables"); - if (idModelOptions().loop()) { + if (idModelOptions().isTensorIndexerEnabled()) { // Depends on CircularBufferInfo and compute_at_map_->allocateIndexVariables info().idModel().allocateLoopIndexVariables(); } - if (idModelOptions().buildTensorIndexer()) { - tensor_indexer_ = std::make_unique(info().idModel()); - non_divisible_predicate_info_ = - std::make_unique(fusion_); - } + tensor_indexer_ = std::make_unique(info().idModel()); + + non_divisible_predicate_info_ = + std::make_unique(fusion_); // Detects all exprssions that don't need predicates. Depends on // nonDivisibleSplitInfo. @@ -663,7 +634,7 @@ bool GpuLower::resolveComputeWith(Fusion* fusion) { Val* GpuLower::getLoopIndexVariable( IterDomain* id, CircularBufferLoopStage stage) const { - if (idModelOptions().loop()) { + if (idModelOptions().isTensorIndexerEnabled()) { return info().idModel().getLoopIndexVariable(id, stage); } else { return info().caMap().getIndexVariable(id, stage); diff --git a/csrc/device_lower/utils.cpp b/csrc/device_lower/utils.cpp index b86e39f8b15..8da85897f11 100644 --- a/csrc/device_lower/utils.cpp +++ b/csrc/device_lower/utils.cpp @@ -2126,7 +2126,7 @@ IterDomain* getConcreteLoopID(IterDomain* id) { // Currently, the concrete loop ID uses the IdModel loop // promotion only when opted in. if ((GpuLower::hasCurrent() && - GpuLower::current()->idModelOptions().loop()) || + GpuLower::current()->idModelOptions().isTensorIndexerEnabled()) || (!GpuLower::hasCurrent() && FusionInfoGuard::current()->hasIdModel() && FusionInfoGuard::current()->idModel().hasIdGraph(IdMappingMode::LOOP))) { // If enabled, the concret ID should be basically just the diff --git a/csrc/id_model/id_model.cpp b/csrc/id_model/id_model.cpp index 47c788c7543..56b454dbb71 100644 --- a/csrc/id_model/id_model.cpp +++ b/csrc/id_model/id_model.cpp @@ -1384,7 +1384,7 @@ void IdModel::allocateLoopIndexVariables() { // If enabled, allocate own indices. Otherwise, use the one // generated for ComputeAtMap for compatibility with the legacy // indexing - if (GpuLower::current()->idModelOptions().loop()) { + if (GpuLower::current()->idModelOptions().isTensorIndexerEnabled()) { loop_index = IrBuilder::create(DataType::Index); } else { const auto& ca_map = FusionInfoGuard::current()->caMap(); diff --git a/csrc/id_model/utils.h b/csrc/id_model/utils.h index 9750bbd592d..4752ff2362a 100644 --- a/csrc/id_model/utils.h +++ b/csrc/id_model/utils.h @@ -16,74 +16,6 @@ namespace nvfuser { -// Options to enable the IdModel-based tensor indexer selectively -enum class IdModelEnableOption { - ConsumerIndex, - ProducerIndex, - InlinePredicate, - UnswitchPredicate, - // Uses the loop promotion to generate loops. Indexing and - // predication need to be enabled as well. - Loop, -}; - -inline std::unordered_set getIdModelEnabledOptions() { - std::unordered_set opts; - - if (hasEnableOptionArgument(EnableOption::IdModel, "consumer_index") || - hasEnableOptionArgument(EnableOption::IdModel, "index") || - hasEnableOptionArgument(EnableOption::IdModel, "all")) { - opts.insert(IdModelEnableOption::ConsumerIndex); - } - - if (hasEnableOptionArgument(EnableOption::IdModel, "producer_index") || - hasEnableOptionArgument(EnableOption::IdModel, "index") || - hasEnableOptionArgument(EnableOption::IdModel, "all")) { - opts.insert(IdModelEnableOption::ProducerIndex); - } - - if (hasEnableOptionArgument(EnableOption::IdModel, "inline_predicate") || - hasEnableOptionArgument(EnableOption::IdModel, "predicate") || - hasEnableOptionArgument(EnableOption::IdModel, "all")) { - opts.insert(IdModelEnableOption::InlinePredicate); - } - - if (hasEnableOptionArgument(EnableOption::IdModel, "unswitch_predicate") || - hasEnableOptionArgument(EnableOption::IdModel, "predicate") || - hasEnableOptionArgument(EnableOption::IdModel, "all")) { - opts.insert(IdModelEnableOption::UnswitchPredicate); - } - - if (hasEnableOptionArgument(EnableOption::IdModel, "loop") || - hasEnableOptionArgument(EnableOption::IdModel, "all")) { - opts.insert(IdModelEnableOption::Loop); - } - - // Loop requires ConsumerIndex, ProducerIndex, InlinePredicate and - // UnswitchPredicate - if (opts.find(IdModelEnableOption::Loop) != opts.end()) { - NVF_ERROR( - opts.find(IdModelEnableOption::ConsumerIndex) != opts.end(), - "ConsumerIndex required for Loop"); - NVF_ERROR( - opts.find(IdModelEnableOption::ProducerIndex) != opts.end(), - "ProducerIndex required for Loop"); - NVF_ERROR( - opts.find(IdModelEnableOption::InlinePredicate) != opts.end(), - "InlinePredicate required for Loop"); - NVF_ERROR( - opts.find(IdModelEnableOption::UnswitchPredicate) != opts.end(), - "UnswitchPredicate required for Loop"); - } - - return opts; -} - -inline bool isIdModelOptionEnabled(IdModelEnableOption option) { - const auto opts = getIdModelEnabledOptions(); - return opts.find(option) != opts.end(); -} - // Get the promotion domain of a given loop domain. inline IterDomain* getLoopPromotion( IterDomain* loop_id, diff --git a/csrc/index_compute.cpp b/csrc/index_compute.cpp index 30f503c5462..63264b21252 100644 --- a/csrc/index_compute.cpp +++ b/csrc/index_compute.cpp @@ -1600,7 +1600,7 @@ Val* Index::getLinearLogicalIndex( const std::vector& loops) { if (!ir_utils::hasRootToLoopLinearTransformations(consumer_tv) || ir_utils::isCpAsyncBulkLoad(consumer_tv->definition()) || - GpuLower::current()->idModelOptions().consumerIndex() || + GpuLower::current()->idModelOptions().isTensorIndexerEnabled() || GpuLower::current()->tmemInfo().hasTMemTensor()) { const TensorIndexer& indexer = GpuLower::current()->tensorIndexer(); auto per_dim_indices = indexer.getIndexFor( @@ -1629,7 +1629,7 @@ std::vector Index::getConsumerPerDimLogicalIndex( TensorView* consumer_tv, const std::vector& loops) { if (!ir_utils::hasRootToLoopLinearTransformations(consumer_tv) || - GpuLower::current()->idModelOptions().consumerIndex() || + GpuLower::current()->idModelOptions().isTensorIndexerEnabled() || GpuLower::current()->tmemInfo().hasTMemTensor()) { const TensorIndexer& indexer = GpuLower::current()->tensorIndexer(); return indexer.getIndexFor( @@ -1652,7 +1652,7 @@ std::vector Index::getProducerPerDimLogicalIndex( const std::vector& loops, const std::unordered_map& override_index) { if (!ir_utils::hasRootToLoopLinearTransformations(producer_tv) || - GpuLower::current()->idModelOptions().producerIndex() || + GpuLower::current()->idModelOptions().isTensorIndexerEnabled() || GpuLower::current()->tmemInfo().hasTMemTensor()) { const TensorIndexer& indexer = GpuLower::current()->tensorIndexer(); return indexer.getIndexFor( @@ -2166,7 +2166,7 @@ bool shouldUseTensorIndexer( } // If opted in, TensorIndexer is used as long as it's supported - if (GpuLower::current()->idModelOptions().producerIndex() && + if (GpuLower::current()->idModelOptions().isTensorIndexerEnabled() && is_tensor_indexer_supported(/*assert=*/false)) { return true; } @@ -2284,7 +2284,7 @@ kir::TensorIndex* Index::getConsumerIndex( Val* index = nullptr; if (!ir_utils::hasRootToLoopLinearTransformations(consumer) || ir_utils::isCpAsyncBulkLoad(consumer->definition()) || - GpuLower::current()->idModelOptions().consumerIndex() || + GpuLower::current()->idModelOptions().isTensorIndexerEnabled() || GpuLower::current()->tmemInfo().hasTMemTensor()) { index = GpuLower::current()->tensorIndexer().getLinearIndex( consumer, consumer->definition(), loops, override_index, ld_st_matrix); diff --git a/csrc/options.cpp b/csrc/options.cpp index 81564414f4c..0cc602bb5c8 100644 --- a/csrc/options.cpp +++ b/csrc/options.cpp @@ -217,7 +217,6 @@ const std::unordered_map& getDisableOptions() { {"greedy_scheduler", DisableOption::GreedyScheduler}, {"grouped_grid_welford_outer_opt", DisableOption::GroupedGridWelfordOuterOpt}, - {"id_model", DisableOption::IdModel}, {"index_hoist", DisableOption::IndexHoist}, {"magic_zero", DisableOption::MagicZero}, {"matmul_expr_eval", DisableOption::MatmulExprEval}, diff --git a/csrc/options.h b/csrc/options.h index 6f722c1227c..3f21c3d9392 100644 --- a/csrc/options.h +++ b/csrc/options.h @@ -147,7 +147,6 @@ enum class DisableOption { GreedyScheduler, //! Disable the greedy scheduler GroupedGridWelfordOuterOpt, //! Disable use of outer-optimized //! grouped grid welford kernel - IdModel, //! Disable IdModel IndexHoist, //! Disable index hoisting MagicZero, //! Disable nvfuser_zero MatmulExprEval, //! Disable ATen evaluation for the entire fusion containing diff --git a/csrc/predicate_compute.cpp b/csrc/predicate_compute.cpp index 3ae6e8cca9b..532f2ff0d6e 100644 --- a/csrc/predicate_compute.cpp +++ b/csrc/predicate_compute.cpp @@ -878,7 +878,7 @@ Val* PredicateCompute::getInlinePredicate( std::vector pred_info_vec; if (!ir_utils::hasRootToLoopLinearTransformations(out_tv) || - GpuLower::current()->idModelOptions().inlinePredicate()) { + GpuLower::current()->idModelOptions().isTensorIndexerEnabled()) { pred_info_vec = gpu_lower->tensorIndexer().getPredicates(out_tv, expr, loops); } else { @@ -981,7 +981,7 @@ void UnswitchPredicate::predicateOn(Expr* tv_expr) { std::vector ref_pred_info; if (!ir_utils::hasRootToLoopLinearTransformations(out_tv) || - GpuLower::current()->idModelOptions().unswitchPredicate()) { + GpuLower::current()->idModelOptions().isTensorIndexerEnabled()) { ref_pred_info = gpu_lower->tensorIndexer().getPredicates( out_tv, tv_expr, for_loops_, unrolled_loop_); } else { diff --git a/python/nvfuser_direct/__init__.py b/python/nvfuser_direct/__init__.py index 317a780225f..34780436eac 100644 --- a/python/nvfuser_direct/__init__.py +++ b/python/nvfuser_direct/__init__.py @@ -366,10 +366,15 @@ def execute( # A copy of fusion is created after construction FusionExecutorCache # Delete the _fusion and reference the fusion inside FusionExecutorCache del self._fusion + + # Add "id_model" as a default enable option + default_enable_options = ["id_model"] + merged_enable_options = default_enable_options + _enable_options + return self.fec.execute( inputs, device=self._get_device_index(device), - _enable_options=_enable_options, + _enable_options=merged_enable_options, _disable_options=_disable_options, ) diff --git a/tests/cpp/test_abstract_tensor.cpp b/tests/cpp/test_abstract_tensor.cpp index 9e6bdd1902e..b91845f7453 100644 --- a/tests/cpp/test_abstract_tensor.cpp +++ b/tests/cpp/test_abstract_tensor.cpp @@ -21,7 +21,7 @@ class AbstractTensorTest : public NVFuserTest { void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); fusion_ptr_ = std::make_unique(); fusion_guard_ptr_ = std::make_unique(fusion_ptr_.get()); auto size = IrBuilder::create(16, DataType::Index); diff --git a/tests/cpp/test_allocation_domain.cpp b/tests/cpp/test_allocation_domain.cpp index 1d09a439ae0..b4a2ae52787 100644 --- a/tests/cpp/test_allocation_domain.cpp +++ b/tests/cpp/test_allocation_domain.cpp @@ -26,7 +26,7 @@ class AllocationDomainTest : public NVFuserTest { protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_allocation_order_inference.cpp b/tests/cpp/test_allocation_order_inference.cpp index b396ec5a37a..f229bbbceb9 100644 --- a/tests/cpp/test_allocation_order_inference.cpp +++ b/tests/cpp/test_allocation_order_inference.cpp @@ -27,7 +27,7 @@ class AllocationOrderInferenceTest : public NVFuserTest { protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_argsort.cpp b/tests/cpp/test_argsort.cpp index 5648eb9d262..1f24f96dbd7 100644 --- a/tests/cpp/test_argsort.cpp +++ b/tests/cpp/test_argsort.cpp @@ -26,7 +26,7 @@ class ArgsortTest : public NVFuserTest { protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_bfs.cpp b/tests/cpp/test_bfs.cpp index 92459457c6f..e4deed959c2 100644 --- a/tests/cpp/test_bfs.cpp +++ b/tests/cpp/test_bfs.cpp @@ -865,7 +865,7 @@ TEST_F(FindAllExprsTest, Rotation) { std::vector shape({16, 100}); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto tv0 = makeConcreteTensor(shape); fusion.addInput(tv0); diff --git a/tests/cpp/test_circular_buffering.cpp b/tests/cpp/test_circular_buffering.cpp index 16f29940d1e..5c332a8666b 100644 --- a/tests/cpp/test_circular_buffering.cpp +++ b/tests/cpp/test_circular_buffering.cpp @@ -21,7 +21,7 @@ TEST_F(NVFuserTest, BarSyncWarpSpecializedPointwise) { std::unique_ptr fusion = std::make_unique(); FusionGuard fg(fusion.get()); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); int64_t number_of_stages = 4; int64_t prefetch_distance = 1; @@ -98,7 +98,7 @@ TEST_F(NVFuserTest, RegisterSharingCircularBufferingPointwiseCustom) { std::unique_ptr fusion = std::make_unique(); FusionGuard fg(fusion.get()); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); int64_t number_of_stages = 4; int64_t prefetch_distance = 1; @@ -182,7 +182,7 @@ TEST_F(NVFuserTest, RegisterSharingCircularBufferingPointwiseNested) { std::unique_ptr fusion = std::make_unique(); FusionGuard fg(fusion.get()); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); int64_t number_of_stages = 4; int64_t prefetch_distance = 1; @@ -266,7 +266,7 @@ class CircularBufferingTest : public NVFuserFixtureParamTest { number_of_stages = std::get<0>(GetParam()); prefetch_distance = std::get<1>(GetParam()); NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; @@ -1100,7 +1100,7 @@ class TmaCircularBufferingTest // NOTE: Multiple of 16 required for inner dimension NVF_ERROR(tensor_inner_dim % 16 == 0); NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } bool testEnablesWarpSpecialization() { @@ -1219,7 +1219,7 @@ TEST_F(NVFuserTest, ElectSyncCompatibility) { std::unique_ptr fusion = std::make_unique(); FusionGuard fg(fusion.get()); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); TensorView* input = makeContigTensor(3); fusion->addInput(input); @@ -2413,7 +2413,7 @@ TEST_F(NVFuserTest, TmaRegisterSharingDynamicShapesExpectFail) { std::unique_ptr fusion = std::make_unique(); FusionGuard fg(fusion.get()); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto tv0 = makeContigTensor(2); fusion->addInput(tv0); @@ -2477,7 +2477,7 @@ TEST_P(TmaRegisterSharing, CtaShapeShmoo) { std::unique_ptr fusion = std::make_unique(); FusionGuard fg(fusion.get()); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto tv0 = makeContigTensor(2); fusion->addInput(tv0); diff --git a/tests/cpp/test_circular_buffering_ping_pong.cpp b/tests/cpp/test_circular_buffering_ping_pong.cpp index d0725f2d742..0085963e598 100644 --- a/tests/cpp/test_circular_buffering_ping_pong.cpp +++ b/tests/cpp/test_circular_buffering_ping_pong.cpp @@ -328,7 +328,7 @@ TEST_P(SiblingPingPongCircularBuffering, TwoTmaLoads) { auto [use_id_model, stage_slice_position] = GetParam(); if (use_id_model) { - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } std::unique_ptr fusion = std::make_unique(); diff --git a/tests/cpp/test_cluster.cpp b/tests/cpp/test_cluster.cpp index 0ab4f40e091..bfa341b74e9 100644 --- a/tests/cpp/test_cluster.cpp +++ b/tests/cpp/test_cluster.cpp @@ -28,7 +28,7 @@ class ClusterReductionTest : public NVFuserTest, protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); NVFUSER_TEST_CUDA_ARCH_GUARD(9, 0); } }; diff --git a/tests/cpp/test_combined_inner_outer_reduction.cpp b/tests/cpp/test_combined_inner_outer_reduction.cpp index 31849df0b8a..48a0312ad84 100644 --- a/tests/cpp/test_combined_inner_outer_reduction.cpp +++ b/tests/cpp/test_combined_inner_outer_reduction.cpp @@ -43,7 +43,7 @@ class CombinedSchedulerTest protected: void SetUp() override { NVFuserFixtureParamTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_compute_with.cpp b/tests/cpp/test_compute_with.cpp index 2841ada1f13..1ac8bb86f25 100644 --- a/tests/cpp/test_compute_with.cpp +++ b/tests/cpp/test_compute_with.cpp @@ -58,7 +58,7 @@ using namespace at::indexing; class ComputeWithTest : public NVFuserTest { protected: void SetUp() override { - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); NVFuserTest::SetUp(); } }; diff --git a/tests/cpp/test_contiguity_id_model.cpp b/tests/cpp/test_contiguity_id_model.cpp index 03989b9d822..aba28781e8e 100644 --- a/tests/cpp/test_contiguity_id_model.cpp +++ b/tests/cpp/test_contiguity_id_model.cpp @@ -31,7 +31,7 @@ class ContigIDGroupsTest : public NVFuserTest { protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } // Helper to construct a simple backward path for testing diff --git a/tests/cpp/test_gather.cpp b/tests/cpp/test_gather.cpp index 69d1cf9749d..1292c24f6e6 100644 --- a/tests/cpp/test_gather.cpp +++ b/tests/cpp/test_gather.cpp @@ -29,7 +29,7 @@ class GatherTest : public NVFuserTest { void SetUp() override { // To make the tests using std::rand deterministic std::srand(0); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_gpu1.cpp b/tests/cpp/test_gpu1.cpp index f0b78653d7b..3ec4bd2f7cb 100644 --- a/tests/cpp/test_gpu1.cpp +++ b/tests/cpp/test_gpu1.cpp @@ -61,7 +61,7 @@ class Gpu1Test : public NVFuserTest { protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; @@ -2805,7 +2805,7 @@ class Fp4CastTest : public NVFuserTest, public: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); std::tie(dtype_highp, vectorization_factor) = GetParam(); NVFUSER_TEST_CUDA_ARCH_GUARD(10, 0); } @@ -3010,7 +3010,7 @@ class AdvancedDtypeTest : public NVFuserFixtureParamTest { bool use_dynamic_shape; void SetUp() override { NVFuserFixtureParamTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); use_dynamic_shape = GetParam(); } }; @@ -3088,7 +3088,7 @@ class Float4E2m1ManualScheduleTestAllArch bool dynamic_shape; void SetUp() override { NVFuserFixtureParamTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); std::tie(vectorize_factor, dynamic_shape) = GetParam(); } }; diff --git a/tests/cpp/test_gpu2.cpp b/tests/cpp/test_gpu2.cpp index d21b0595847..644744529f5 100644 --- a/tests/cpp/test_gpu2.cpp +++ b/tests/cpp/test_gpu2.cpp @@ -62,7 +62,7 @@ class Gpu2Test : public NVFuserTest { protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_gpu3.cpp b/tests/cpp/test_gpu3.cpp index 682af1b3d67..02b8f65812c 100644 --- a/tests/cpp/test_gpu3.cpp +++ b/tests/cpp/test_gpu3.cpp @@ -63,7 +63,7 @@ class Gpu3Test : public NVFuserTest { protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; @@ -3398,7 +3398,7 @@ TEST_F(Gpu3Test, FusionIssueRepro1844_CUDA) { } TEST_F(Gpu3Test, FusionInsertMagicZero1_CUDA) { - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); Fusion fusion; FusionGuard fg(&fusion); diff --git a/tests/cpp/test_greedy.cpp b/tests/cpp/test_greedy.cpp index 4035c3629af..fc4804b4aab 100644 --- a/tests/cpp/test_greedy.cpp +++ b/tests/cpp/test_greedy.cpp @@ -27,7 +27,7 @@ class GreedySchedulerTest : public NVFuserTest { protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_index_put.cpp b/tests/cpp/test_index_put.cpp index 063304e8f90..0d529336b8a 100644 --- a/tests/cpp/test_index_put.cpp +++ b/tests/cpp/test_index_put.cpp @@ -42,7 +42,7 @@ std::vector generateSizeOneParams() { class IndexPut : public NVFuserFixtureParamTest { protected: void SetUp() override { - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); NVFuserTest::SetUp(); } }; diff --git a/tests/cpp/test_index_select.cpp b/tests/cpp/test_index_select.cpp index 2accf8217aa..860bc9195e9 100644 --- a/tests/cpp/test_index_select.cpp +++ b/tests/cpp/test_index_select.cpp @@ -74,7 +74,7 @@ void checkIndexSelectVectorization( class IndexSelectTest : public NVFuserTest { protected: void SetUp() override { - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); NVFuserTest::SetUp(); } }; diff --git a/tests/cpp/test_indexing.cpp b/tests/cpp/test_indexing.cpp index 918e0e7059e..407d73fac5c 100644 --- a/tests/cpp/test_indexing.cpp +++ b/tests/cpp/test_indexing.cpp @@ -426,7 +426,7 @@ class PredicateIndexValidator : public kir::IrVisitor { bool enable_contig_indexing, Args... args) { EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); // Disable simplifications to make the pattern matching of sameAs work DisableOptionsGuard disable_options_guard; @@ -864,8 +864,9 @@ TEST_F(IndexingTest, Reshape) { // to provide the extent of the group. However, since everything // should be deterministic, string match should also work. return std::string( - "( ( ( ( ( i98 * 20 ) + ( ( i99 * 10 ) + i100 ) ) / 25 ) * 25 ) " - "+ ( ( ( i98 * 20 ) + ( ( i99 * 10 ) + i100 ) ) % 25 ) )"); + "( ( ( ( ( i114 * 20 ) + ( ( i115 * 10 ) + i116 ) ) / 25 ) * 25 " + ") " + "+ ( ( ( i114 * 20 ) + ( ( i115 * 10 ) + i116 ) ) % 25 ) )"); } default: return std::string(); @@ -3095,7 +3096,7 @@ TEST_F(PredicateIndexingTest, DoubleBuffering1) { at::Tensor t0 = at::randn({1000}, options); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); KernelExecutor ke; ke.compile(&fusion, {t0}); @@ -3193,7 +3194,7 @@ TEST_F(PredicateIndexingTest, CircularBuffering1) { at::Tensor t0 = at::randn({1000}, options); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); KernelExecutor ke; ke.compile(&fusion, {t0}); @@ -3360,7 +3361,7 @@ TEST_F(PredicateIndexingTest, UnrolledCircularBuffering) { at::Tensor t0 = at::randn({1000}, options); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); KernelExecutor ke; ke.compile(&fusion, {t0}); @@ -3435,7 +3436,7 @@ TEST_F(PredicateIndexingTest, UnswitchedCircularBuffering1) { at::Tensor t0 = at::randn({99}, options); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); KernelExecutor ke; ke.compile(&fusion, {t0}); @@ -3520,7 +3521,7 @@ TEST_F(PredicateIndexingTest, UnswitchedCircularBuffering2) { at::Tensor t0 = at::randn({1000}, options); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); KernelExecutor ke; ke.compile(&fusion, {t0}); @@ -3622,7 +3623,7 @@ TEST_P(PredicateIndexingTest, UnswitchedCircularBuffering3) { at::Tensor t1 = at::randn({1000}, options); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); KernelExecutor ke; ke.compile(&fusion, {t0, t1}); @@ -3698,7 +3699,7 @@ TEST_F(PredicateIndexingTest, UnswitchedCircularBuffering4) { // Running this fusion with the legacy indexer would result in an // error if run with compute-sanitizer. EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto options = at::TensorOptions().dtype(at::kFloat).device(at::kCUDA, 0); auto t0 = at::randn({16}, options); @@ -3790,7 +3791,7 @@ TEST_F(PredicateIndexingTest, NonDivisibleSplit1) { PredicateIndexValidator::validate(&fusion, false); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto options = at::TensorOptions().dtype(at::kFloat).device(at::kCUDA, 0); at::Tensor t0 = at::randn({999}, options); @@ -3881,7 +3882,7 @@ TEST_F(PredicateIndexingTest, NonDivisibleSplitWithUnswitch) { PredicateIndexValidator::validate(&fusion, false); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto options = at::TensorOptions().dtype(at::kFloat).device(at::kCUDA, 0); at::Tensor t0 = at::randn({999}, options); @@ -3975,7 +3976,7 @@ TEST_F(PredicateIndexingTest, NonDivisibleSplitWithCircularBuffering) { PredicateIndexValidator::validate(&fusion, false); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto options = at::TensorOptions().dtype(at::kFloat).device(at::kCUDA, 0); at::Tensor t0 = at::randn({999}, options); @@ -4085,7 +4086,7 @@ TEST_F( PredicateIndexValidator::validate(&fusion, false); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto options = at::TensorOptions().dtype(at::kFloat).device(at::kCUDA, 0); at::Tensor t0 = at::randn({999}, options); @@ -4169,7 +4170,7 @@ TEST_P(PredicateIndexingTest, UnswitchPredicateIssueRepro681) { EnableOptionsGuard enable_options_guard; if (GetParam()) { - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } else { EnableOptionsGuard::getCurOptions().unset(EnableOption::IdModel); } @@ -4331,7 +4332,7 @@ TEST_F(PredicateIndexingTest, NonDivisibleSplitWithUnswitchAndBroadcast) { at::Tensor t1 = at::randn({5, 100}, options); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); KernelExecutor ke; ke.compile(&fusion, {t0, t1}); @@ -4346,7 +4347,7 @@ TEST_F(PredicateIndexingTest, NonDivisibleSplitWithNonLogicalToLoopDomains) { Fusion fusion; FusionGuard fg(&fusion); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); std::vector shape{5, 2}; @@ -4564,7 +4565,7 @@ TEST_F(PredicateIndexingTest, UnswitchConsolidationDifferentThreading) { at::Tensor t1 = at::randn({1000}, options); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); KernelExecutor ke; ke.compile(&fusion, {t0, t1}); @@ -4746,7 +4747,7 @@ TEST_F( PredicateIndexingTest, ParallelDimensionPredicateWithUnswitchAndSetLoopDomain) { // EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); Fusion fusion; FusionGuard fg(&fusion); @@ -5242,7 +5243,7 @@ TEST_F(ContigIndexingTest, ConcretizedBroadcastMerge) { IndexValidator::validate(&fusion, true); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto options = at::TensorOptions().dtype(at::kFloat).device(at::kCUDA, 0); auto t0 = at::randn({5, 6}, options); @@ -5355,7 +5356,7 @@ TEST_F(ContigIndexingTest, Transpose) { IndexValidator::validate(&fusion, true); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto options = at::TensorOptions().dtype(at::kFloat).device(at::kCUDA, 0); auto t0 = at::randn({100, 100}, options); @@ -5583,7 +5584,7 @@ TEST_F(ContigPredicateIndexingTest, NonDivisibleSplit1) { PredicateIndexValidator::validate(&fusion, true); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto options = at::TensorOptions().dtype(at::kFloat).device(at::kCUDA, 0); at::Tensor t0 = at::randn({10, 20}, options); @@ -5680,7 +5681,7 @@ TEST_F(IndexingTest, PerDimLogicalIndices) { }; EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); DisableOptionsGuard disable_options_guard; DisableOptionsGuard::getCurOptions().set(DisableOption::ExprSimplify); DisableOptionsGuard::getCurOptions().set(DisableOption::IndexHoist); @@ -5800,7 +5801,7 @@ TEST_F(IndexingTest, ResizeRotation) { const int64_t i0 = 32; EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto zero = fusion.zeroVal(); @@ -5873,7 +5874,7 @@ TEST_F(PredicateIndexingTest, VectorizedResizeRotation) { const int64_t i0 = 32; EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto zero = fusion.zeroVal(); @@ -5976,7 +5977,7 @@ TEST_F(IndexingTest, Issue3505Repro1) { const auto zero = fusion.zeroVal(); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto tv0 = makeContigConcreteTensor({i1, i2}); fusion.addInput(tv0); @@ -6019,7 +6020,7 @@ TEST_F(IndexingTest, Issue3505Repro2) { const auto zero = fusion.zeroVal(); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto tv0 = makeContigConcreteTensor({i0}); fusion.addInput(tv0); @@ -6058,7 +6059,7 @@ TEST_F(IndexingTest, Issue3505Repro2) { TEST_F(IndexingTest, AlmostExactIndexingUpdate) { EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); Fusion fusion; FusionGuard fg(&fusion); @@ -6131,7 +6132,7 @@ TEST_F(IndexingTest, BroadcastLogicalDomainIndexing) { TEST_F(IndexingTest, Rng) { EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto fusion_ptr = std::make_unique(); auto& fusion = *fusion_ptr; @@ -6164,7 +6165,7 @@ TEST_F(IndexingTest, Rng) { // loop may not be unrolled. TEST_F(IndexingTest, StaticIndexing) { EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto fusion_ptr = std::make_unique(); auto& fusion = *fusion_ptr; @@ -6252,7 +6253,7 @@ TEST_F(PredicateIndexingTest, NonTrivialSizeOneDomain) { PredicateIndexValidator::validate(&fusion, false); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto options = at::TensorOptions().dtype(at::kFloat).device(at::kCUDA, 0); at::Tensor t0 = at::randn({8}, options); @@ -6269,7 +6270,7 @@ TEST_F(PredicateIndexingTest, AdditionalNonDivisibleSplit) { Fusion fusion; FusionGuard fg(&fusion); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto tv0 = makeContigConcreteTensor({8}); fusion.addInput(tv0); @@ -6322,7 +6323,7 @@ TEST_F(PredicateIndexingTest, AdditionalNonDivisibleSplitAfterDivisibleSplit) { Fusion fusion; FusionGuard fg(&fusion); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto tv0 = makeContigConcreteTensor({8}); fusion.addInput(tv0); diff --git a/tests/cpp/test_indexing_advanced.cpp b/tests/cpp/test_indexing_advanced.cpp index 610844ffc96..346ce35d658 100644 --- a/tests/cpp/test_indexing_advanced.cpp +++ b/tests/cpp/test_indexing_advanced.cpp @@ -23,7 +23,7 @@ class AdvancedIndexingTest : public NVFuserFixtureParamTest { protected: void SetUp() override { if (GetParam()) { - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } else { EnableOptionsGuard::getCurOptions().unset(EnableOption::IdModel); } @@ -33,7 +33,7 @@ class AdvancedIndexingTest : public NVFuserFixtureParamTest { class AdvancedIndexingIdModelTest : public NVFuserTest { protected: void SetUp() override { - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_inlining.cpp b/tests/cpp/test_inlining.cpp index ced80f67493..b7127196d1f 100644 --- a/tests/cpp/test_inlining.cpp +++ b/tests/cpp/test_inlining.cpp @@ -25,7 +25,7 @@ class InliningTest : public NVFuserTest { protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_layout_op.cpp b/tests/cpp/test_layout_op.cpp index 1a2fa739f27..10dd1bee197 100644 --- a/tests/cpp/test_layout_op.cpp +++ b/tests/cpp/test_layout_op.cpp @@ -86,7 +86,7 @@ class LayoutOpTest : public NVFuserTest { protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_loop_domain_scheduling.cpp b/tests/cpp/test_loop_domain_scheduling.cpp index 2d3f120a1af..542b9634130 100644 --- a/tests/cpp/test_loop_domain_scheduling.cpp +++ b/tests/cpp/test_loop_domain_scheduling.cpp @@ -40,7 +40,7 @@ void checkGetAllStmts(Fusion* fusion) { class LoopDomainSchedulingTest : public NVFuserTest { protected: void SetUp() override { - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_matmul.cpp b/tests/cpp/test_matmul.cpp index 8364cadbdb8..83eaaeb4d89 100644 --- a/tests/cpp/test_matmul.cpp +++ b/tests/cpp/test_matmul.cpp @@ -663,7 +663,7 @@ TEST_P(MatmulTestWithLayout, AmpereMatmulRegCircularBuffer) { TEST_F(MatmulTest, MatmulMatmulAmpere) { NVFUSER_TEST_CUDA_ARCH_GUARD(8, 0); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); Fusion fusion; FusionGuard fg(&fusion); @@ -3657,7 +3657,7 @@ class HopperMatmulTest : public HopperBase { protected: void SetUp() override { HopperBase::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_matmul_scheduler.cpp b/tests/cpp/test_matmul_scheduler.cpp index fccb8c01fd7..544f2ed41c5 100644 --- a/tests/cpp/test_matmul_scheduler.cpp +++ b/tests/cpp/test_matmul_scheduler.cpp @@ -32,7 +32,7 @@ class MatmulSchedulerTest : public NVFuserTest { void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } private: @@ -2489,7 +2489,7 @@ class MatmulSchedulerPluginTest : public NVFuserTest { void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } private: @@ -2929,7 +2929,7 @@ class AllocationDomainTest void SetUp() override { NVFuserFixtureParamTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } std::pair getInputTVs( @@ -3383,7 +3383,7 @@ class HopperPlusMatmulSchedulerTest mparams.circular_buffer_options.circular_buffer_smem_read = true; mparams.circular_buffer_options.smem_circular_buffer_stage = 2; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } void TearDown() { diff --git a/tests/cpp/test_memory.cpp b/tests/cpp/test_memory.cpp index 2ca1880bd52..38c27689e7c 100644 --- a/tests/cpp/test_memory.cpp +++ b/tests/cpp/test_memory.cpp @@ -452,7 +452,7 @@ class TMASimpleLdstTest NVF_THROW("Invalid dimension"); } - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; @@ -2858,7 +2858,7 @@ class LdMatrixTest : public NVFuserFixtureParamTest { GTEST_SKIP() << "skipping tests on pre-Turing GPUs"; } NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; @@ -2912,7 +2912,7 @@ class StMatrixTest : public NVFuserFixtureParamTest { GTEST_SKIP() << "skipping tests on pre-Hopper GPUs"; } NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_mma.cpp b/tests/cpp/test_mma.cpp index ffc5ccac01b..8cfd88271e2 100644 --- a/tests/cpp/test_mma.cpp +++ b/tests/cpp/test_mma.cpp @@ -87,7 +87,7 @@ class MmaTest : public NVFuserFixtureParamTest { } NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; @@ -310,7 +310,7 @@ class HopperRS : public HopperBase, void SetUp() override { HopperBase::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); macro = std::get<0>(GetParam()); dtype = std::get<1>(GetParam()); @@ -423,7 +423,7 @@ class HopperRSStmatrix void SetUp() override { HopperBase::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); macro = std::get<0>(GetParam()); dtype = std::get<1>(GetParam()); layout = std::get<2>(GetParam()); @@ -669,7 +669,7 @@ class SSTest : public Base, void SetUp() override { Base::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); macro = std::get<0>(GetParam()); dtype = std::get<1>(GetParam()); diff --git a/tests/cpp/test_moe.cpp b/tests/cpp/test_moe.cpp index b1524e1575b..ee232a4aecf 100644 --- a/tests/cpp/test_moe.cpp +++ b/tests/cpp/test_moe.cpp @@ -43,7 +43,7 @@ class SgLangMoETest : public NVFuserFixtureParamTest { protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); std::tie( num_experts, num_tokens, topk, rounding_factor, manual_scheduling) = diff --git a/tests/cpp/test_move_pad.cpp b/tests/cpp/test_move_pad.cpp index 6390ad84ecd..1aae7f46acf 100644 --- a/tests/cpp/test_move_pad.cpp +++ b/tests/cpp/test_move_pad.cpp @@ -22,7 +22,7 @@ class MovePadTest : public NVFuserTest { protected: void SetUp() override { DisableOptionsGuard::getCurOptions().set(DisableOption::ResizeScheduler); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_move_repeat_forward.cpp b/tests/cpp/test_move_repeat_forward.cpp index 0806a3b05df..095f47edbd0 100644 --- a/tests/cpp/test_move_repeat_forward.cpp +++ b/tests/cpp/test_move_repeat_forward.cpp @@ -23,7 +23,7 @@ class MoveRepeatForwardTest : public NVFuserTest { protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_move_split_cat.cpp b/tests/cpp/test_move_split_cat.cpp index 52c69d35ddd..4ae3d4bcf0b 100644 --- a/tests/cpp/test_move_split_cat.cpp +++ b/tests/cpp/test_move_split_cat.cpp @@ -22,7 +22,7 @@ class MoveSplitCatTest : public NVFuserTest { protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_outer_reduction.cpp b/tests/cpp/test_outer_reduction.cpp index 5c63556bde2..dbcb6cdcc82 100644 --- a/tests/cpp/test_outer_reduction.cpp +++ b/tests/cpp/test_outer_reduction.cpp @@ -36,7 +36,7 @@ class OuterReductionTest : public NVFuserTest { protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_persistent_buffer.cpp b/tests/cpp/test_persistent_buffer.cpp index 5d463551e9b..420cd06c7c2 100644 --- a/tests/cpp/test_persistent_buffer.cpp +++ b/tests/cpp/test_persistent_buffer.cpp @@ -31,7 +31,7 @@ class PersistentBufferTest : public NVFuserTest { protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; @@ -2181,7 +2181,7 @@ class TmaPersistentTestP void SetUp() override { NVFUSER_TEST_CUDA_ARCH_GUARD(9, 0); NVFuserFixtureParamTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); EnableOptionsGuard::getCurOptions().set(EnableOption::TmaInnerPersistent); } }; diff --git a/tests/cpp/test_pointwise.cpp b/tests/cpp/test_pointwise.cpp index 259baf663bc..8386ae5cedd 100644 --- a/tests/cpp/test_pointwise.cpp +++ b/tests/cpp/test_pointwise.cpp @@ -25,7 +25,7 @@ namespace nvfuser { class PointwiseTest : public NVFuserTest { protected: void SetUp() override { - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; @@ -36,7 +36,7 @@ class PointwiseTestP : public NVFuserFixtureParamTest { protected: void SetUp() override { NVFuserFixtureParamTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_predicate_elimination.cpp b/tests/cpp/test_predicate_elimination.cpp index cb675fea4ec..9290f47bfe3 100644 --- a/tests/cpp/test_predicate_elimination.cpp +++ b/tests/cpp/test_predicate_elimination.cpp @@ -21,7 +21,7 @@ class PredicateEliminationTest : public NVFuserTest { protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_reduction.cpp b/tests/cpp/test_reduction.cpp index 511585b0add..3b91978597a 100644 --- a/tests/cpp/test_reduction.cpp +++ b/tests/cpp/test_reduction.cpp @@ -78,7 +78,7 @@ class ReductionTest : public NVFuserTest { protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_reduction_pointwise.cpp b/tests/cpp/test_reduction_pointwise.cpp index 3cac57a0ab4..0428ef7b07d 100644 --- a/tests/cpp/test_reduction_pointwise.cpp +++ b/tests/cpp/test_reduction_pointwise.cpp @@ -20,7 +20,7 @@ class PointwiseFusedReductionTest : public NVFuserTest { protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_remove_bcast_squeeze.cpp b/tests/cpp/test_remove_bcast_squeeze.cpp index 26bb4d537f6..2b2570778c4 100644 --- a/tests/cpp/test_remove_bcast_squeeze.cpp +++ b/tests/cpp/test_remove_bcast_squeeze.cpp @@ -26,7 +26,7 @@ class RemoveBcastSqueezeTest : public NVFuserTest { protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_replay.cpp b/tests/cpp/test_replay.cpp index b2ccda2d4e5..468a73904d6 100644 --- a/tests/cpp/test_replay.cpp +++ b/tests/cpp/test_replay.cpp @@ -28,7 +28,7 @@ class ReplayTest : public NVFuserTest { protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_reshape.cpp b/tests/cpp/test_reshape.cpp index 0d70471755e..e8ba6572ed9 100644 --- a/tests/cpp/test_reshape.cpp +++ b/tests/cpp/test_reshape.cpp @@ -55,7 +55,7 @@ class ReshapeTest : public NVFuserTest { protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; @@ -396,7 +396,7 @@ class ReshapeReduction : public NVFuserFixtureParamTest { protected: void SetUp() override { NVFuserFixtureParamTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_resize.cpp b/tests/cpp/test_resize.cpp index 9370eb2408e..260e55cde50 100644 --- a/tests/cpp/test_resize.cpp +++ b/tests/cpp/test_resize.cpp @@ -88,7 +88,7 @@ TEST_F(ResizeTest, Pad1) { auto t0 = at::randn(shape, options); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); KernelExecutor ke; ke.compile(&fusion, {t0}); @@ -119,7 +119,7 @@ TEST_F(ResizeTest, Pad2) { auto t0 = at::randn(shape, options); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); KernelExecutor ke; ke.compile(&fusion, {t0}); @@ -167,7 +167,7 @@ TEST_F(ResizeTest, Pad3) { auto t1 = at::randn(padded_shape, options); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); KernelExecutor ke; ke.compile(&fusion, {t0, t1}); @@ -196,7 +196,7 @@ TEST_F(ResizeTest, Pad4) { auto t0 = at::randn(shape, options); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); KernelExecutor ke; ke.compile(&fusion, {t0}); @@ -246,7 +246,7 @@ TEST_F(ResizeTest, Pad5) { auto t0 = at::randn(shape, options); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); KernelExecutor ke; ke.compile(&fusion, {t0}); @@ -292,7 +292,7 @@ TEST_F(ResizeTest, Pad6) { auto t1 = at::randn(padded_shape, options); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); KernelExecutor ke; ke.compile(&fusion, {t0, t1}); @@ -338,7 +338,7 @@ TEST_F(ResizeTest, Pad7) { auto t0 = at::randn(shape, options); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); KernelExecutor ke; ke.compile(&fusion, {t0}); @@ -411,7 +411,7 @@ TEST_F(ResizeTest, PadScheduler1) { auto t0 = at::randn(shape, options); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); FusionExecutorCache executor_cache(std::move(fusion)); auto cg_outputs = executor_cache.runFusionWithInputs({t0}); @@ -445,7 +445,7 @@ TEST_F(ResizeTest, PadScheduler2) { auto t1 = at::randn(padded_shape, options); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); FusionExecutorCache executor_cache(std::move(fusion_ptr)); auto cg_outputs = executor_cache.runFusionWithInputs({t0, t1}); @@ -521,7 +521,7 @@ TEST_F(ResizeTest, PadScheduler4) { std::vector pad_extents{1, 1}; EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); FusionExecutorCache executor_cache(std::move(fusion)); auto cg_outputs = executor_cache.runFusionWithInputs({t0, 1, 1}); @@ -556,7 +556,7 @@ TEST_F(ResizeTest, PadBroadcastInput) { auto t0 = at::randn(shape, options); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); FusionExecutorCache executor_cache(std::move(fusion)); auto cg_outputs = executor_cache.runFusionWithInputs({t0}); @@ -1396,7 +1396,7 @@ TEST_F(ResizeTest, PadReduceScheduler1) { } EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); FusionExecutorCache executor_cache(std::move(fusion_ptr)); auto cg_outputs = executor_cache.runFusionWithInputs(inputs); @@ -1686,7 +1686,7 @@ TEST_F(ResizeTest, PadWithValue) { auto t0 = at::randn(shape, options); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); KernelExecutor ke; ke.compile(&fusion, {t0}); @@ -1720,7 +1720,7 @@ TEST_F(ResizeTest, PadToEmptyTensor) { auto t0 = at::randn(shape, options); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); FusionExecutorCache executor_cache(std::move(fusion)); auto cg_outputs = executor_cache.runFusionWithInputs({t0}); @@ -1751,7 +1751,7 @@ TEST_F(ResizeTest, PadHalfWithDoubleValue) { auto t0 = at::ones(shape, options); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); KernelExecutor ke; ke.compile(&fusion, {t0}); @@ -2357,7 +2357,7 @@ TEST_F(ResizeTest, ResizePadToBroadcastStatic) { auto t1 = at::randn(t1_size, options); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); FusionExecutorCache executor_cache(std::move(fusion)); auto cg_outputs = executor_cache.runFusionWithInputs({t0, t1}); @@ -2424,7 +2424,7 @@ TEST_F(ResizeTest, ResizePadToBroadcastDynamic) { } EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); FusionExecutorCache executor_cache(std::move(fusion)); auto cg_outputs = executor_cache.runFusionWithInputs(inputs); @@ -2468,7 +2468,7 @@ TEST_F(ResizeTest, ResizePadToBroadcastIssue596) { auto t1 = at::randn({3}, options); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); KernelArgumentHolder args({t0, t1}); FusionKernelRuntime runtime(std::move(fusion), args); @@ -3021,7 +3021,7 @@ TEST_F(ResizeTest, ReshapeToPad) { fusion.addOutput(tv2); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); FusionExecutorCache executor_cache(std::move(fusion_ptr)); @@ -3193,7 +3193,7 @@ TEST_F(ResizeTest, PadExpandedEmpty) { auto t0 = at::randn({0}, options).as_strided({2, 0, 3}, {0, 0, 0}); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); FusionExecutorCache executor_cache(std::move(fusion_ptr)); auto cg_outputs = executor_cache.runFusionWithInputs({t0}); @@ -3220,7 +3220,7 @@ TEST_F(ResizeTest, PadOfBroadcast) { auto t0 = at::randn(shape0, options); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); KernelExecutor ke; ke.compile(&fusion, {t0}); @@ -3251,7 +3251,7 @@ TEST_F(ResizeTest, PadOfExpandedBroadcast) { auto t0 = at::randn(shape0, options); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); KernelExecutor ke; ke.compile(&fusion, {t0}); @@ -3593,7 +3593,7 @@ TEST_F(ResizeTest, SliceScheduledLikeProducer) { std::vector shape({100}); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); // concrete shapes to avoid dynamic Fusion auto tv0 = makeConcreteTensor(shape); @@ -3641,7 +3641,7 @@ TEST_F(ResizeTest, PadScheduledLikeConsumer) { std::vector shape({100}); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); // concrete shapes to avoid dynamic Fusion auto tv0 = makeConcreteTensor(shape); @@ -3689,7 +3689,7 @@ TEST_F(ResizeTest, SliceThenPadLeftHalf) { std::vector shape({100}); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); // concrete shapes to avoid dynamic Fusion auto tv0 = makeContigConcreteTensor(shape); @@ -3741,7 +3741,7 @@ TEST_F(ResizeTest, SliceThenPadRightHalf) { std::vector shape({100}); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); // concrete shapes to avoid dynamic Fusion auto tv0 = makeContigConcreteTensor(shape); @@ -3795,7 +3795,7 @@ TEST_F(ResizeTest, SliceThenConcat) { std::vector shape({100}); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); // concrete shapes to avoid dynamic Fusion auto tv0 = makeContigConcreteTensor(shape); @@ -3856,7 +3856,7 @@ TEST_F(ResizeTest, SliceSliceConcatConcat) { const int64_t rope_size = 32; EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto zero = fusion.zeroVal(); @@ -3955,7 +3955,7 @@ TEST_P(ResizeSchedulerTest, PropagateSliceToInputs) { std::vector shape({-1, 100}); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto tv0 = makeConcreteTensor(shape); fusion.addInput(tv0); @@ -4044,7 +4044,7 @@ TEST_P(ResizeSchedulerTest, PropagateSliceToInputsWithReshape1) { std::vector shape({16, 100}); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto tv0 = makeConcreteTensor(shape); fusion.addInput(tv0); @@ -4135,7 +4135,7 @@ TEST_P(ResizeSchedulerTest, PropagateSliceToInputsWithReshape2) { std::vector shape({16, 100}); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto tv0 = makeConcreteTensor(shape); fusion.addInput(tv0); @@ -4221,7 +4221,7 @@ TEST_P(ResizeSchedulerTest, PropagateMultipleSlicesToInputs1) { std::vector shape({-1, 100}); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto tv0 = makeConcreteTensor(shape); fusion.addInput(tv0); @@ -4333,7 +4333,7 @@ TEST_F(ResizeSchedulerTest, PropagateMultipleSlicesToInputs2) { std::vector shape({-1, 100}); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto tv0 = makeConcreteTensor(shape); fusion.addInput(tv0); @@ -4445,7 +4445,7 @@ TEST_F(ResizeSchedulerTest, PropagateMultipleSlicesToInputs3) { std::vector shape({-1, 100}); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto tv0 = makeConcreteTensor(shape); fusion.addInput(tv0); @@ -4744,7 +4744,7 @@ TEST_P(ResizeSchedulerTest, SliceRotateCat) { std::vector shape({-1, 100}); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto tv0 = makeConcreteTensor(shape); fusion.addInput(tv0); @@ -4876,7 +4876,7 @@ TEST_P(ResizeSchedulerTest, SliceRotateCatResidual) { std::vector shape({16, 100}); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto tv0 = makeConcreteTensor(shape); fusion.addInput(tv0); @@ -5006,7 +5006,7 @@ TEST_F(ResizeSchedulerTest, SliceRotateCatTwice) { std::vector shape({-1, 100}); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto tv0 = makeConcreteTensor(shape); fusion.addInput(tv0); @@ -5101,7 +5101,7 @@ TEST_P(ResizeSchedulerTest, PropagatePadToInputs) { std::vector shape({-1, 100}); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto tv0 = makeConcreteTensor(shape); fusion.addInput(tv0); @@ -5192,7 +5192,7 @@ TEST_P(ResizeSchedulerTest, PropagateCatToInputs) { std::vector shape({-1, 100}); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto tv0 = makeConcreteTensor(shape); fusion.addInput(tv0); @@ -6125,7 +6125,7 @@ TEST_F(ResizeTest, ReshapeAfterRef) { std::vector shape({2, 16, 100}); EnableOptionsGuard enable_options_guard; - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto tv0 = makeConcreteTensor(shape); fusion.addInput(tv0); diff --git a/tests/cpp/test_rng.cpp b/tests/cpp/test_rng.cpp index 65664ca359b..078c77203e8 100644 --- a/tests/cpp/test_rng.cpp +++ b/tests/cpp/test_rng.cpp @@ -68,7 +68,7 @@ at::Tensor generate_normal(int64_t size, at::ScalarType dtype) { class RNGTest : public NVFuserTest { void SetUp() override { - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_rope.cpp b/tests/cpp/test_rope.cpp index c2368b1694c..0a2289a565f 100644 --- a/tests/cpp/test_rope.cpp +++ b/tests/cpp/test_rope.cpp @@ -49,7 +49,7 @@ struct RopeConfig { class RopeTest : public NVFuserFixtureParamTest { void SetUp() override { NVFuserFixtureParamTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_scatter.cpp b/tests/cpp/test_scatter.cpp index cbfa7f8bb9e..57678c800f9 100644 --- a/tests/cpp/test_scatter.cpp +++ b/tests/cpp/test_scatter.cpp @@ -54,7 +54,7 @@ class ScatterTest : public NVFuserFixtureParamTest { protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); manual_scheduling = GetParam(); } @@ -487,7 +487,7 @@ class ScatterAccumulateTest protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); std::tie(m, n, dtype, accumulate_op) = GetParam(); } diff --git a/tests/cpp/test_select.cpp b/tests/cpp/test_select.cpp index 3d8e45b4c10..97286e2d0c8 100644 --- a/tests/cpp/test_select.cpp +++ b/tests/cpp/test_select.cpp @@ -18,7 +18,7 @@ namespace nvfuser { class SelectTest : public NVFuserTest { protected: void SetUp() override { - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); NVFuserTest::SetUp(); } }; diff --git a/tests/cpp/test_serial_gridreduce.cpp b/tests/cpp/test_serial_gridreduce.cpp index 31ee3c5da34..d65ef0125e4 100644 --- a/tests/cpp/test_serial_gridreduce.cpp +++ b/tests/cpp/test_serial_gridreduce.cpp @@ -36,7 +36,7 @@ class SerialGridReductionTest : public NVFuserTest { protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_stream.cpp b/tests/cpp/test_stream.cpp index 84f11a63142..aa69a439b46 100644 --- a/tests/cpp/test_stream.cpp +++ b/tests/cpp/test_stream.cpp @@ -30,7 +30,7 @@ namespace nvfuser { class StreamTest : public NVFuserTest { public: StreamTest() { - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_topk.cpp b/tests/cpp/test_topk.cpp index 912af23b6f8..f8812405231 100644 --- a/tests/cpp/test_topk.cpp +++ b/tests/cpp/test_topk.cpp @@ -480,7 +480,7 @@ class TopKTest : public NVFuserTest { protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_transpose.cpp b/tests/cpp/test_transpose.cpp index bd21e4a5f1e..cda882b1a04 100644 --- a/tests/cpp/test_transpose.cpp +++ b/tests/cpp/test_transpose.cpp @@ -56,7 +56,7 @@ class TransposeTest : public NVFuserTest { void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } private: @@ -1351,7 +1351,7 @@ TEST_F(TransposeTest, ReductionIterDomainOnInputsIssue1659) { TEST_F(TransposeTest, DanglingBroadcastIssue4957) { // The issue is not specific to TensorIndexer but just make sure the // fix works with it - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); auto fusion_ptr = std::make_unique(); Fusion& fusion = *fusion_ptr.get(); diff --git a/tests/cpp/test_unary.cpp b/tests/cpp/test_unary.cpp index 93dd2d38111..1fda2bb7923 100644 --- a/tests/cpp/test_unary.cpp +++ b/tests/cpp/test_unary.cpp @@ -20,7 +20,7 @@ namespace nvfuser { class UnaryTest : public NVFuserFixtureParamTest { void SetUp() override { NVFuserFixtureParamTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/test_vectorization.cpp b/tests/cpp/test_vectorization.cpp index b052380d993..be99db5b56b 100644 --- a/tests/cpp/test_vectorization.cpp +++ b/tests/cpp/test_vectorization.cpp @@ -37,7 +37,7 @@ void checkMappedVal( class VectorizationAnalysisTest : public NVFuserTest { void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; @@ -697,7 +697,7 @@ class VectorizationCastTest void SetUp() override { NVFuserTest::SetUp(); std::tie(dtype_from, dtype_to, vectorization_factor) = GetParam(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } protected: @@ -854,7 +854,7 @@ class Vect256Test : public NVFuserFixtureParamTest { NVFuserFixtureParamTest::SetUp(); NVFUSER_TEST_CUDA_ARCH_GUARD(10, 0); std::tie(dtype, cache_op) = GetParam(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } protected: diff --git a/tests/cpp/test_welford.cpp b/tests/cpp/test_welford.cpp index 6aa6b401792..1e9b7c32ab5 100644 --- a/tests/cpp/test_welford.cpp +++ b/tests/cpp/test_welford.cpp @@ -21,7 +21,7 @@ class WelfordTest : public NVFuserTest { protected: void SetUp() override { NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/cpp/utils.cpp b/tests/cpp/utils.cpp index b7d7093cfc6..f3393f7235b 100644 --- a/tests/cpp/utils.cpp +++ b/tests/cpp/utils.cpp @@ -56,7 +56,7 @@ void NVFuserTest::SetUp() { if (!deviceMajorMinorCheck(6)) { GTEST_SKIP() << "skipping tests on pre-PASCAL GPUs"; } - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } NVFuserTest::~NVFuserTest() { diff --git a/tests/cpp/utils.h b/tests/cpp/utils.h index 0bc9e4a6f77..bc7c4bd6282 100644 --- a/tests/cpp/utils.h +++ b/tests/cpp/utils.h @@ -475,7 +475,7 @@ class BlackwellBase : public NVFuserTest { "sm_100/sm_104, not sm_110+)"; } NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; @@ -487,7 +487,7 @@ class TmaBase : public NVFuserTest { GTEST_SKIP() << "skipping tests on pre-Hopper GPUs"; } NVFuserTest::SetUp(); - EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel, {"all"}); + EnableOptionsGuard::getCurOptions().set(EnableOption::IdModel); } }; diff --git a/tests/python/direct/test_python_direct.py b/tests/python/direct/test_python_direct.py index 57b1c594778..de7ac506ab7 100644 --- a/tests/python/direct/test_python_direct.py +++ b/tests/python/direct/test_python_direct.py @@ -229,22 +229,20 @@ def test_fusion_execution_cache(): i2 = i0 / 8; nvfuser_index_t i3; i3 = i0 % 8; - nvfuser_index_t i4; - i4 = ((nvfuser_index_t)threadIdx.x) + (128 * ((nvfuser_index_t)blockIdx.x)); - if ((i4 < 64)) { + if ((((nvfuser_index_t)threadIdx.x) < 64)) { Array T4; T4[0] = 0; T4[0] - = T1[((((T1.alloc_stride[0LL] * i1) + (T1.alloc_stride[1LL] * i2)) + (T1.alloc_stride[2LL] * i3)) + ((4 * T1.alloc_stride[0LL]) * ((nvfuser_index_t)blockIdx.x)))]; + = T1[(((T1.alloc_stride[0LL] * i1) + (T1.alloc_stride[1LL] * i2)) + (T1.alloc_stride[2LL] * i3))]; Array T3; T3[0] = 0; T3[0] - = T0[((((T0.alloc_stride[0LL] * i1) + (T0.alloc_stride[1LL] * i2)) + (T0.alloc_stride[2LL] * i3)) + ((4 * T0.alloc_stride[0LL]) * ((nvfuser_index_t)blockIdx.x)))]; + = T0[(((T0.alloc_stride[0LL] * i1) + (T0.alloc_stride[1LL] * i2)) + (T0.alloc_stride[2LL] * i3))]; Array T5; T5[0] = T3[0] + T4[0]; - T2[i4] + T2[((nvfuser_index_t)threadIdx.x)] = T5[0]; } }\n"""