diff --git a/include/tvm/ir/expr.h b/include/tvm/ir/expr.h index efde52385177..6c81f06561d2 100644 --- a/include/tvm/ir/expr.h +++ b/include/tvm/ir/expr.h @@ -866,28 +866,6 @@ struct PackedFuncValueConverter { } }; -/* \brief Backwards compatibility wrapper for IntImm arguments - * - * In previous versions of TVM, IntImm was the default FFI type for - * integer arguments, instead of runtime::Int. For backwards - * compatibility where the callee has been updated to expected a - * runtime::Int, the caller has not been updated to provide a - * runtime::Int (e.g. relay script parsing), and the auto-unboxing of - * runtime::Int does not apply (e.g. making an `Array`), - * allow the IntImm to be generated. - */ -template <> -struct PackedFuncValueConverter { - template - static runtime::Int From(const PODSubclass& val) { - if (val.template IsObjectRef()) { - return runtime::Int(val.template AsObjectRef()->value); - } else { - return val.template AsObjectRef(); - } - } -}; - } // namespace runtime } // namespace tvm diff --git a/include/tvm/meta_schedule/schedule/cuda/thread_bind.h b/include/tvm/meta_schedule/schedule/cuda/thread_bind.h index 125d6dc11fc8..c1d65f9e738b 100644 --- a/include/tvm/meta_schedule/schedule/cuda/thread_bind.h +++ b/include/tvm/meta_schedule/schedule/cuda/thread_bind.h @@ -36,7 +36,7 @@ namespace meta_schedule { * \return A sampler that returns a random thread extent. */ std::function MakeFactorSampler(tir::Schedule sch, - Array thread_extents); + Array thread_extents); /*! * \brief Bind blockIdx.x and threadIdx.x to the given loop diff --git a/include/tvm/meta_schedule/schedule_rule.h b/include/tvm/meta_schedule/schedule_rule.h index 90aec05187eb..6c22f5e9e12f 100644 --- a/include/tvm/meta_schedule/schedule_rule.h +++ b/include/tvm/meta_schedule/schedule_rule.h @@ -154,11 +154,11 @@ class ScheduleRule : public runtime::ObjectRef { * ignored by default. This function should return True for a block that should be tiled. * \return The schedule rule created */ - TVM_DLL static ScheduleRule MultiLevelTiling(String structure, // - Optional> tile_binds, // - Optional max_innermost_factor, // - Optional> vector_load_lens, // - Optional> reuse_read, // + TVM_DLL static ScheduleRule MultiLevelTiling(String structure, // + Optional> tile_binds, // + Optional max_innermost_factor, // + Optional> vector_load_lens, // + Optional> reuse_read, // Optional> reuse_write, Optional filter_fn = NullOpt); @@ -181,7 +181,7 @@ class ScheduleRule : public runtime::ObjectRef { */ TVM_DLL static ScheduleRule MultiLevelTilingWithIntrin( String intrin_name, String structure, Optional> tile_binds, - Optional max_innermost_factor, Optional> vector_load_lens, + Optional max_innermost_factor, Optional> vector_load_lens, Optional> reuse_read, Optional> reuse_write); /*! @@ -206,8 +206,8 @@ class ScheduleRule : public runtime::ObjectRef { */ TVM_DLL static ScheduleRule MultiLevelTilingTensorCore( Array> intrin_groups, String structure, - Optional> tile_binds, Optional max_innermost_factor, - Optional> vector_load_lens, Optional> reuse_read, + Optional> tile_binds, Optional max_innermost_factor, + Optional> vector_load_lens, Optional> reuse_read, Optional> reuse_write, bool use_software_pipeline); /*! @@ -222,8 +222,9 @@ class ScheduleRule : public runtime::ObjectRef { * \return The schedule rule created */ TVM_DLL static ScheduleRule MultiLevelTilingWideVector( - String structure, Integer vector_length_in_bits, Optional max_innermost_factor, - Optional> reuse_read, Optional> reuse_write); + String structure, runtime::Int vector_length_in_bits, + Optional max_innermost_factor, Optional> reuse_read, + Optional> reuse_write); /*! * \brief Create a rule: add-rfactor to some blocks if needed @@ -234,7 +235,7 @@ class ScheduleRule : public runtime::ObjectRef { * \return The schedule rule created */ TVM_DLL static ScheduleRule AddRFactor(int max_jobs_per_core, // - Optional max_innermost_factor); + Optional max_innermost_factor); /*! * \brief Create a schedule rule which applies cross-thread reduction to some reduction blocks * correspondingly when needed @@ -272,7 +273,7 @@ class ScheduleRule : public runtime::ObjectRef { * when this schedule rule is created. * \return The schedule rule created */ - TVM_DLL static ScheduleRule AutoBind(int max_threadblocks, Array thread_extents, + TVM_DLL static ScheduleRule AutoBind(int max_threadblocks, Array thread_extents, int max_threads_per_block = -1); /*! * \brief Create a schedule rule with customized methods on the python-side. diff --git a/include/tvm/relay/attrs/transform.h b/include/tvm/relay/attrs/transform.h index 91020fc7443b..555ed58f8ecf 100644 --- a/include/tvm/relay/attrs/transform.h +++ b/include/tvm/relay/attrs/transform.h @@ -530,6 +530,9 @@ struct MatrixSetDiagAttrs : public tvm::AttrsNode { } }; // struct MatrixSetDiagAttrs +template +using Identity = T; + /*! \brief Attributes used in cumsum and cumprod operator */ struct ScanopAttrs : public tvm::AttrsNode { Integer axis; @@ -542,7 +545,7 @@ struct ScanopAttrs : public tvm::AttrsNode { // Default is 0 which is "false" TVM_ATTR_FIELD(exclusive) .describe("The first element is not included") - .set_default(Bool(false)); + .set_default(Identity(false)); } }; // struct ScanopAttrs diff --git a/src/meta_schedule/schedule/cuda/thread_bind.cc b/src/meta_schedule/schedule/cuda/thread_bind.cc index 110cae96cb53..682f2d9fed0e 100644 --- a/src/meta_schedule/schedule/cuda/thread_bind.cc +++ b/src/meta_schedule/schedule/cuda/thread_bind.cc @@ -31,14 +31,14 @@ namespace meta_schedule { using namespace tvm::tir; -std::function MakeFactorSampler(Schedule sch, Array thread_extents) { +std::function MakeFactorSampler(Schedule sch, Array thread_extents) { return [sch = std::move(sch), thread_extents = std::move(thread_extents)](int64_t max_extent) -> ExprRV { Array extents; extents.reserve(thread_extents.size()); - for (const Integer extent : thread_extents) { + for (auto extent : thread_extents) { if (extent->value <= max_extent) { - extents.push_back(runtime::Int(extent->value)); + extents.push_back(extent); } } int n = extents.size(); @@ -64,7 +64,7 @@ Array BindSpatialLoop(Schedule sch, LoopRV loop, int64_t max_threadblock } if (extent <= max_threadblocks * max_threads_per_block) { if (!get_factor) { - get_factor = MakeFactorSampler(sch, {32, 64, 128, 256, 512, 1024}); + get_factor = MakeFactorSampler(sch, Array{32, 64, 128, 256, 512, 1024}); } ExprRV factor = get_factor(std::min(extent, max_threads_per_block)); Array splits = sch->Split(loop, {NullOpt, factor}); diff --git a/src/meta_schedule/schedule_rule/add_rfactor.cc b/src/meta_schedule/schedule_rule/add_rfactor.cc index 2fc1352677cb..1820d588eb73 100644 --- a/src/meta_schedule/schedule_rule/add_rfactor.cc +++ b/src/meta_schedule/schedule_rule/add_rfactor.cc @@ -68,10 +68,10 @@ class AddRFactorNode : public ScheduleRuleNode { }; ScheduleRule ScheduleRule::AddRFactor(int max_jobs_per_core, - Optional max_innermost_factor) { + Optional max_innermost_factor) { ObjectPtr n = make_object(); n->max_jobs_per_core = max_jobs_per_core; - n->max_innermost_factor = max_innermost_factor.value_or(Integer(-1))->value; + n->max_innermost_factor = max_innermost_factor.value_or(runtime::Int(-1))->value; n->max_parallel_extent_ = -1; n->max_parallel_basic_ = -1; return ScheduleRule(n); diff --git a/src/meta_schedule/schedule_rule/auto_bind.cc b/src/meta_schedule/schedule_rule/auto_bind.cc index fa47d1edb860..19f14eed519a 100644 --- a/src/meta_schedule/schedule_rule/auto_bind.cc +++ b/src/meta_schedule/schedule_rule/auto_bind.cc @@ -31,11 +31,11 @@ class AutoBindNode : public ScheduleRuleNode { // Inherited from ScheduleRuleNode void InitializeWithTuneContext(const TuneContext& context) final { CHECK(context->target.defined()) << "ValueError: target is not defined"; - Optional max_threads_per_block = - context->target.value()->GetAttr("max_threads_per_block"); + Optional max_threads_per_block = + context->target.value()->GetAttr("max_threads_per_block"); CHECK(max_threads_per_block.defined()) << "ValueError: missing attribute `max_threads_per_block` in the target"; - this->max_threads_per_block_ = max_threads_per_block.value().IntValue(); + this->max_threads_per_block_ = max_threads_per_block.value(); } // Inherited from ScheduleRuleNode @@ -53,7 +53,7 @@ class AutoBindNode : public ScheduleRuleNode { /*! \brief The max number of threadblocks in the cuda device */ int64_t max_threadblocks_ = -1; /*! \brief thread_extents Candidates of thread axis extent. */ - Array thread_extents_; + Array thread_extents_; void VisitAttrs(tvm::AttrVisitor* v) { // `max_threads_per_block_` is not visited @@ -72,7 +72,7 @@ Array AutoBindNode::Apply(const tir::Schedule& sch, const tir::Bl return {sch}; } -ScheduleRule ScheduleRule::AutoBind(int max_threadblocks, Array thread_extents, +ScheduleRule ScheduleRule::AutoBind(int max_threadblocks, Array thread_extents, int max_threads_per_block) { ObjectPtr n = make_object(); n->max_threadblocks_ = max_threadblocks; diff --git a/src/meta_schedule/schedule_rule/multi_level_tiling.cc b/src/meta_schedule/schedule_rule/multi_level_tiling.cc index 2979e4229bdd..07524543f7bd 100644 --- a/src/meta_schedule/schedule_rule/multi_level_tiling.cc +++ b/src/meta_schedule/schedule_rule/multi_level_tiling.cc @@ -392,8 +392,8 @@ void MultiLevelTilingNode::AnnotateCooperativeFetching(Schedule* sch, // Constructor ScheduleRule ScheduleRule::MultiLevelTiling(String structure, Optional> tile_binds, - Optional max_innermost_factor, - Optional> vector_load_lens, + Optional max_innermost_factor, + Optional> vector_load_lens, Optional> reuse_read, Optional> reuse_write, Optional filter_fn) { diff --git a/src/meta_schedule/schedule_rule/multi_level_tiling.h b/src/meta_schedule/schedule_rule/multi_level_tiling.h index 23d6599a2538..ee3355bde2b8 100644 --- a/src/meta_schedule/schedule_rule/multi_level_tiling.h +++ b/src/meta_schedule/schedule_rule/multi_level_tiling.h @@ -239,16 +239,16 @@ class MultiLevelTilingNode : public ScheduleRuleNode { template ObjectPtr MultiLevelTilingInitCommon(String structure, Optional> tile_binds, - Optional max_innermost_factor, - Optional> vector_load_lens, + Optional max_innermost_factor, + Optional> vector_load_lens, Optional> reuse_read, Optional> reuse_write) { ObjectPtr n = make_object(); n->structure = structure; n->tile_binds = tile_binds.value_or({}); - n->max_innermost_factor = max_innermost_factor.value_or(Integer(-1))->value; + n->max_innermost_factor = max_innermost_factor.value_or(runtime::Int(-1))->value; n->vector_load_lens = vector_load_lens.defined() - ? support::AsVector(vector_load_lens.value()) + ? support::AsVector(vector_load_lens.value()) : std::vector(); n->reuse_read_ = reuse_read.defined() ? ReuseConfig(reuse_read.value()) : ReuseConfig(); n->reuse_write_ = reuse_write.defined() ? ReuseConfig(reuse_write.value()) : ReuseConfig(); diff --git a/src/meta_schedule/schedule_rule/multi_level_tiling_tensor_core.cc b/src/meta_schedule/schedule_rule/multi_level_tiling_tensor_core.cc index e038ab908dd8..89a774f7af66 100644 --- a/src/meta_schedule/schedule_rule/multi_level_tiling_tensor_core.cc +++ b/src/meta_schedule/schedule_rule/multi_level_tiling_tensor_core.cc @@ -890,7 +890,7 @@ inline std::vector MultiLevelTilingTensorCoreNode::TransformForTensorizat ScheduleRule ScheduleRule::MultiLevelTilingTensorCore( Array> intrin_groups, String structure, Optional> tile_binds, - Optional max_innermost_factor, Optional> vector_load_lens, + Optional max_innermost_factor, Optional> vector_load_lens, Optional> reuse_read, Optional> reuse_write, bool use_software_pipeline) { if (tile_binds.defined()) { diff --git a/src/meta_schedule/schedule_rule/multi_level_tiling_wide_vector.cc b/src/meta_schedule/schedule_rule/multi_level_tiling_wide_vector.cc index e68b64ea2d3a..0689bd320509 100644 --- a/src/meta_schedule/schedule_rule/multi_level_tiling_wide_vector.cc +++ b/src/meta_schedule/schedule_rule/multi_level_tiling_wide_vector.cc @@ -113,8 +113,9 @@ std::pair, Array> MultiLevelTilingWideVectorNode } ScheduleRule ScheduleRule::MultiLevelTilingWideVector( - String structure, Integer vector_length_in_bits, Optional max_innermost_factor, - Optional> reuse_read, Optional> reuse_write) { + String structure, runtime::Int vector_length_in_bits, + Optional max_innermost_factor, Optional> reuse_read, + Optional> reuse_write) { auto node = MultiLevelTilingInitCommon( structure, NullOpt, max_innermost_factor, NullOpt, reuse_read, reuse_write); node->vector_length_in_bits = vector_length_in_bits->value; diff --git a/src/meta_schedule/schedule_rule/multi_level_tiling_with_intrin.cc b/src/meta_schedule/schedule_rule/multi_level_tiling_with_intrin.cc index 428a1206a4ca..5d95f480c4f8 100644 --- a/src/meta_schedule/schedule_rule/multi_level_tiling_with_intrin.cc +++ b/src/meta_schedule/schedule_rule/multi_level_tiling_with_intrin.cc @@ -93,7 +93,7 @@ class MultiLevelTilingWithIntrinNode : public MultiLevelTilingNode { ScheduleRule ScheduleRule::MultiLevelTilingWithIntrin( String intrin_name, String structure, Optional> tile_binds, - Optional max_innermost_factor, Optional> vector_load_lens, + Optional max_innermost_factor, Optional> vector_load_lens, Optional> reuse_read, Optional> reuse_write) { ICHECK(tir::TensorIntrin::Get(intrin_name).defined()) << "Provided tensor intrinsic " << intrin_name << " is not registered."; diff --git a/src/meta_schedule/schedule_rule/schedule_rule.cc b/src/meta_schedule/schedule_rule/schedule_rule.cc index 83f5d073cb32..c7afbc3fbd4f 100644 --- a/src/meta_schedule/schedule_rule/schedule_rule.cc +++ b/src/meta_schedule/schedule_rule/schedule_rule.cc @@ -65,16 +65,16 @@ Array ScheduleRule::DefaultLLVM() { /*disallow_op=*/Array{"tir.exp"}), ScheduleRule::AddRFactor( /*max_jobs_per_core=*/16, - /*max_innermost_factor=*/Integer(64)), + /*max_innermost_factor=*/runtime::Int(64)), ScheduleRule::MultiLevelTiling( /*structure=*/"SSRSRS", /*tile_binds=*/NullOpt, - /*max_innermost_factor=*/Integer(64), + /*max_innermost_factor=*/runtime::Int(64), /*vector_load_lens=*/NullOpt, /*reuse_read=*/NullOpt, /*reuse_write=*/ Map{{"req", String("may")}, - {"levels", Array{1, 2}}, + {"levels", Array{1, 2}}, {"scope", String("global")}}), ScheduleRule::ParallelizeVectorizeUnroll( /*max_jobs_per_core=*/16, @@ -101,27 +101,27 @@ Array ScheduleRule::DefaultX86(const String& type) { /*disallow_op=*/Array{"tir.exp"}), ScheduleRule::AddRFactor( /*max_jobs_per_core=*/16, - /*max_innermost_factor=*/Integer(64)), + /*max_innermost_factor=*/runtime::Int(64)), ScheduleRule::MultiLevelTilingWithIntrin( /*intrin_name=*/intrins[type], /*structure=*/"SSRSRS", /*tile_binds=*/NullOpt, - /*max_innermost_factor=*/Integer(64), + /*max_innermost_factor=*/runtime::Int(64), /*vector_load_lens=*/NullOpt, /*reuse_read=*/NullOpt, /*reuse_write=*/ Map{{"req", String("may")}, - {"levels", Array{1, 2}}, + {"levels", Array{1, 2}}, {"scope", String("global")}}), ScheduleRule::MultiLevelTiling( /*structure=*/"SSRSRS", /*tile_binds=*/NullOpt, - /*max_innermost_factor=*/Integer(64), + /*max_innermost_factor=*/runtime::Int(64), /*vector_load_lens=*/NullOpt, /*reuse_read=*/NullOpt, /*reuse_write=*/ Map{{"req", String("may")}, - {"levels", Array{1, 2}}, + {"levels", Array{1, 2}}, {"scope", String("global")}}), ScheduleRule::ParallelizeVectorizeUnroll( /*max_jobs_per_core=*/16, @@ -138,15 +138,15 @@ Array ScheduleRule::DefaultCUDA() { ScheduleRule::MultiLevelTiling( /*structure=*/"SSSRRSRS", /*tile_binds=*/Array{"blockIdx.x", "vthread.x", "threadIdx.x"}, - /*max_innermost_factor=*/Integer(64), - /*vector_load_lens=*/Array{1, 2, 3, 4, 8, 16}, + /*max_innermost_factor=*/runtime::Int(64), + /*vector_load_lens=*/Array{1, 2, 3, 4, 8, 16}, /*reuse_read=*/ Map{{"req", String("must")}, - {"levels", Array{4}}, // + {"levels", Array{4}}, // {"scope", String("shared")}}, /*reuse_write=*/ Map{{"req", String("must")}, - {"levels", Array{3}}, // + {"levels", Array{3}}, // {"scope", String("local")}}), ScheduleRule::InlineConstantScalars(), ScheduleRule::AutoInline( @@ -166,7 +166,7 @@ Array ScheduleRule::DefaultCUDA() { /*unroll_explicit=*/true), ScheduleRule::AutoBind( /*max_threadblocks=*/256, - /*thread_extents*/ Array{32, 64, 128, 256, 512, 1024}), + /*thread_extents*/ Array{32, 64, 128, 256, 512, 1024}), }; } @@ -241,30 +241,30 @@ Array ScheduleRule::DefaultCUDATensorCore() { /*intrin_groups=*/wmma_intrin_groups, /*structure=*/"SSSRRSRS", /*tile_binds=*/Array{"blockIdx.y", "blockIdx.x", "threadIdx.y"}, - /*max_innermost_factor=*/Integer(4), - /*vector_load_lens=*/Array{1, 2, 3, 4, 8, 16}, + /*max_innermost_factor=*/runtime::Int(4), + /*vector_load_lens=*/Array{1, 2, 3, 4, 8, 16}, /*reuse_read=*/ Map{{"req", String("must")}, - {"levels", Array{4}}, // + {"levels", Array{4}}, // {"scope", String("shared.dyn")}}, /*reuse_write=*/ Map{{"req", String("must")}, - {"levels", Array{2}}, // + {"levels", Array{2}}, // {"scope", String("shared.dyn")}}, /*use_software_pipeline=*/false), // ScheduleRule::MultiLevelTilingTensorCore( /*intrin_groups=*/mma_intrin_groups, /*structure=*/"SSSRRSRS", /*tile_binds=*/Array{"blockIdx.y", "blockIdx.x", "threadIdx.y"}, - /*max_innermost_factor=*/Integer(4), - /*vector_load_lens=*/Array{1, 2, 3, 4, 8, 16}, + /*max_innermost_factor=*/runtime::Int(4), + /*vector_load_lens=*/Array{1, 2, 3, 4, 8, 16}, /*reuse_read=*/ Map{{"req", String("must")}, - {"levels", Array{4}}, // + {"levels", Array{4}}, // {"scope", String("shared.dyn")}}, /*reuse_write=*/ Map{{"req", String("no")}, - {"levels", Array{2}}, // + {"levels", Array{2}}, // {"scope", String("shared.dyn")}}, /*use_software_pipeline=*/true) // }; @@ -288,11 +288,11 @@ Array ScheduleRule::DefaultHexagon() { ScheduleRule::MultiLevelTilingWideVector( /*structure=*/"SRSRS", /*vector_length_in_bits=*/1024, - /*max_innermost_factor=*/Integer(128), + /*max_innermost_factor=*/runtime::Int(128), /*reuse_read=*/NullOpt, /*reuse_write=*/ Map{{"req", String("may")}, - {"levels", Array{1, 2}}, + {"levels", Array{1, 2}}, {"scope", String("global")}}), ScheduleRule::ParallelizeVectorizeUnroll( /*max_jobs_per_core=*/16, @@ -317,12 +317,12 @@ Array ScheduleRule::DefaultMicro() { ScheduleRule::MultiLevelTiling( /*structure=*/"SSRSRS", /*tile_binds=*/NullOpt, - /*max_innermost_factor=*/Integer(64), + /*max_innermost_factor=*/runtime::Int(64), /*vector_load_lens=*/NullOpt, /*reuse_read=*/NullOpt, /*reuse_write=*/ Map{{"req", String("may")}, - {"levels", Array{1, 2}}, + {"levels", Array{1, 2}}, {"scope", String("global")}}), }; } @@ -333,12 +333,12 @@ Array GetARMNeonSpecificRules() { /*intrin_name=*/String("dot_4x4_i8i8s32_neon"), /*structure=*/"SSRSRS", /*tile_binds=*/NullOpt, - /*max_innermost_factor=*/Integer(32), + /*max_innermost_factor=*/runtime::Int(32), /*vector_load_lens=*/NullOpt, /*reuse_read=*/NullOpt, /*reuse_write=*/ Map{{"req", String("may")}, - {"levels", Array{1, 2}}, + {"levels", Array{1, 2}}, {"scope", String("global")}}), }; } @@ -349,34 +349,34 @@ Array GetARMDotprodSpecificRules() { /*intrin_name=*/String("dot_4x4_i8i8s32_sdot"), /*structure=*/"SSRSRS", /*tile_binds=*/NullOpt, - /*max_innermost_factor=*/Integer(32), + /*max_innermost_factor=*/runtime::Int(32), /*vector_load_lens=*/NullOpt, /*reuse_read=*/NullOpt, /*reuse_write=*/ Map{{"req", String("may")}, - {"levels", Array{1, 2}}, + {"levels", Array{1, 2}}, {"scope", String("global")}}), ScheduleRule::MultiLevelTilingWithIntrin( /*intrin_name=*/String("dot_4x4_u8u8u32_udot"), /*structure=*/"SSRSRS", /*tile_binds=*/NullOpt, - /*max_innermost_factor=*/Integer(32), + /*max_innermost_factor=*/runtime::Int(32), /*vector_load_lens=*/NullOpt, /*reuse_read=*/NullOpt, /*reuse_write=*/ Map{{"req", String("may")}, - {"levels", Array{1, 2}}, + {"levels", Array{1, 2}}, {"scope", String("global")}}), ScheduleRule::MultiLevelTilingWithIntrin( /*intrin_name=*/String("dot_4x4_u8u8i32_hdot"), /*structure=*/"SSRSRS", /*tile_binds=*/NullOpt, - /*max_innermost_factor=*/Integer(32), + /*max_innermost_factor=*/runtime::Int(32), /*vector_load_lens=*/NullOpt, /*reuse_read=*/NullOpt, /*reuse_write=*/ Map{{"req", String("may")}, - {"levels", Array{1, 2}}, + {"levels", Array{1, 2}}, {"scope", String("global")}}), }; } @@ -394,18 +394,18 @@ Array ScheduleRule::DefaultARM(const String& type) { /*disallow_op=*/Array{"tir.exp"}), ScheduleRule::AddRFactor( /*max_jobs_per_core=*/8, - /*max_innermost_factor=*/Integer(32)), + /*max_innermost_factor=*/runtime::Int(32)), "neon" == type ? GetARMNeonSpecificRules() : Array{}, "dotprod" == type ? GetARMDotprodSpecificRules() : Array{}, ScheduleRule::MultiLevelTiling( /*structure=*/"SSRSRS", /*tile_binds=*/NullOpt, - /*max_innermost_factor=*/Integer(32), + /*max_innermost_factor=*/runtime::Int(32), /*vector_load_lens=*/NullOpt, /*reuse_read=*/NullOpt, /*reuse_write=*/ Map{{"req", String("may")}, - {"levels", Array{1, 2}}, + {"levels", Array{1, 2}}, {"scope", String("global")}}), ScheduleRule::ParallelizeVectorizeUnroll( /*max_jobs_per_core=*/8, diff --git a/src/meta_schedule/trace_apply.cc b/src/meta_schedule/trace_apply.cc index e60fdf5b9d2b..706b6e0f1c6d 100644 --- a/src/meta_schedule/trace_apply.cc +++ b/src/meta_schedule/trace_apply.cc @@ -242,13 +242,13 @@ void ScheduleUsingAnchorTrace(Schedule sch, const Trace& anchor_trace, const tvm } else if (target->kind->name == "llvm" || target->kind->name == "hexagon") { sch->Parallel(sch->Fuse(sch->GetLoops(last_block))); } else if (IsGPUTarget(target->kind->name)) { - auto max_threads_per_block = target->GetAttr("max_threads_per_block"); + auto max_threads_per_block = target->GetAttr("max_threads_per_block"); ICHECK(max_threads_per_block.defined()) << "ValueError: missing attribute `max_threads_per_block` in the target"; auto auto_bind_rule = ScheduleRule::AutoBind(/*max_threadblocks=*/256, - /*thread_extents*/ Array{32, 64, 128, 256, 512, 1024}, + /*thread_extents*/ Array{32, 64, 128, 256, 512, 1024}, max_threads_per_block.value()->value); auto_bind_rule->Apply(sch, last_block); } diff --git a/src/relay/parser/parser.cc b/src/relay/parser/parser.cc index 233455bf89ba..fe7caa1671d4 100644 --- a/src/relay/parser/parser.cc +++ b/src/relay/parser/parser.cc @@ -21,6 +21,7 @@ * \file parser.cc * \brief A parser for TVM IR. */ + #include #include #include @@ -1348,8 +1349,11 @@ class Parser { auto next = Peek(); switch (next->token_type) { case TokenType::kFloat: + return runtime::Float(Downcast(Match(next->token_type)->data)->value); case TokenType::kInteger: + return runtime::Int(Downcast(Match(next->token_type)->data)->value); case TokenType::kBoolean: + return runtime::Bool(Downcast(Match(next->token_type)->data)->value); case TokenType::kStringLiteral: return Match(next->token_type)->data; case TokenType::kMetaReference: