/*
 * Licensed to the Apache Software Foundation (ASF) under one
 * or more contributor license agreements.  See the NOTICE file
 * distributed with this work for additional information
 * regarding copyright ownership.  The ASF licenses this file
 * to you under the Apache License, Version 2.0 (the
 * "License"); you may not use this file except in compliance
 * with the License.  You may obtain a copy of the License at
 *
 *   http://www.apache.org/licenses/LICENSE-2.0
 *
 * Unless required by applicable law or agreed to in writing,
 * software distributed under the License is distributed on an
 * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
 * KIND, either express or implied.  See the License for the
 * specific language governing permissions and limitations
 * under the License.
 */

#ifndef TVM_S_TIR_META_SCHEDULE_SCHEDULE_RULE_H_
#define TVM_S_TIR_META_SCHEDULE_SCHEDULE_RULE_H_

#include <tvm/ffi/container/array.h>
#include <tvm/ffi/container/map.h>
#include <tvm/ffi/function.h>
#include <tvm/ffi/optional.h>
#include <tvm/ffi/reflection/registry.h>
#include <tvm/ffi/string.h>
#include <tvm/ir/expr.h>
#include <tvm/s_tir/schedule/schedule.h>

namespace tvm {
namespace s_tir {
namespace meta_schedule {

class TuneContext;
class ScheduleRule;

/*! \brief Rules to modify a block in a schedule. */
class ScheduleRuleNode : public ffi::Object {
 public:
  /*! \brief Virtual destructor. */
  virtual ~ScheduleRuleNode() = default;

  static void RegisterReflection() {
    namespace refl = tvm::ffi::reflection;
    refl::ObjectDef<ScheduleRuleNode>();
  }

  /*!
   * \brief Initialize the design space generator with tuning context.
   * \param context The tuning context for initialization.
   * \note This method is supposed to be called only once before every other method.
   */
  virtual void InitializeWithTuneContext(const TuneContext& context) = 0;

  /*!
   * \brief Apply a schedule rule to the specific block in the given schedule.
   * \param sch The schedule to be modified.
   * \param block The specific block to apply the schedule rule.
   * \return The list of schedules generated by applying the schedule rule.
   */
  virtual ffi::Array<s_tir::Schedule> Apply(const s_tir::Schedule& sch,
                                            const s_tir::SBlockRV& block) = 0;

  /*!
   * \brief Deep clone the schedule rule.
   * \return The cloned schedule rule.
   */
  virtual ScheduleRule Clone() const = 0;

  static constexpr const bool _type_mutable = true;
  TVM_FFI_DECLARE_OBJECT_INFO("s_tir.meta_schedule.ScheduleRule", ScheduleRuleNode, ffi::Object);
};

/*!
 * \brief Managed reference to ScheduleRuleNode
 * \sa ScheduleRuleNode
 */
class ScheduleRule : public ffi::ObjectRef {
 public:
  /*!
   * \brief The function type of `InitializeWithTuneContext` method.
   * \param context The tuning context for initialization.
   */
  using FInitializeWithTuneContext = ffi::TypedFunction<void(const TuneContext&)>;
  /*!
   * \brief The function type of `Apply` method.
   * \param sch The schedule to be modified.
   * \param block The specific block to apply the schedule rule.
   * \return The list of schedules generated by applying the schedule rule.
   */
  using FApply = ffi::TypedFunction<ffi::Array<s_tir::Schedule>(const s_tir::Schedule&,
                                                                const s_tir::SBlockRV&)>;
  /*!
   * \brief The function type of `Clone` method.
   * \return The cloned schedule rule.
   */
  using FClone = ffi::TypedFunction<ScheduleRule()>;
  /*!
   * \brief Create a rule that applies customized rules registered using block attribute
   * `schedule_rule`. The rule will be dispatched according to target keys.
   * \return The created schedule rule.
   */
  TVM_DLL static ScheduleRule ApplyCustomRule();
  /*! \brief Check if the rule is `ApplyCustomRule` */
  TVM_DLL static bool IsApplyCustomRule(const ScheduleRule& rule);
  /*!
   * \brief Create an auto-inline rule that inlines spatial blocks if it satisfies some conditions
   * \param into_producer If allows to inline a block into its producer
   * \param into_consumer If allows to inline a block into its consumer
   * \param inline_const_tensor Always inline constant tensors
   * \param disallow_if_then_else Always disallow if-then-else-like constructs
   * \param require_ordered Always require the read-to-write mapping to be ordered
   * \param require_injective Always require the read-to-write mapping to be injective
   * \param disallow_op The operators that are disallowed in auto inline
   * \return The schedule rule created
   */
  TVM_DLL static ScheduleRule AutoInline(bool into_producer,          //
                                         bool into_consumer,          //
                                         bool inline_const_tensor,    //
                                         bool disallow_if_then_else,  //
                                         bool require_injective,      //
                                         bool require_ordered,        //
                                         ffi::Optional<ffi::Array<ffi::String>> disallow_op);

  /*!
   * \brief Inline blocks that produce a constant scalar. Such blocks get in the way of
   * ReverseComputeInline during AutoInline, since they are also counted as a producer block
   * unless they are inlined first. So it is recommended to run InlineConstantScalars before
   * AutoInline.
   * \return The schedule rule created
   */
  TVM_DLL static ScheduleRule InlineConstantScalars();

  /*!
   * \brief Create a mega rule: multi-level tiling with data reuse
   * \param structure The tiling structure. Recommended:
   * - 'SSRSRS' on CPU
   * - 'SSSRRSRS' on GPU
   * \param tile_binds For each level of tiles, which thread axis it is bound to. Recommended:
   * - std::nullopt on CPU
   * - [blockIdx.x, vthread.x, threadIdx.x] on GPU
   * \param max_innermost_factor The maximum size of the innermost factor. std::nullopt means no
   * limit \param vector_load_lens The length of vector lane in vectorized cooperative fetching.
   * std::nullopt means disable vectorization
   * \param reuse_read Data reuse configuration for reading. std::nullopt means no reuse.
   * \param reuse_write Data reuse configuration for writing. std::nullopt means no reuse.
   * \param filter_fn A function that can be passed to overwrite the default condition for applying
   * MultiLevelTiling to a block. Its signature must be (Schedule, SBlockRV) -> bool.
   * This is useful if there is a need to apply MultiLevelTiling to an operation / block which is
   * 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(
      ffi::String structure,                                      //
      ffi::Optional<ffi::Array<ffi::String>> tile_binds,          //
      ffi::Optional<int64_t> max_innermost_factor,                //
      ffi::Optional<ffi::Array<int64_t>> vector_load_lens,        //
      ffi::Optional<ffi::Map<ffi::String, ffi::Any>> reuse_read,  //
      ffi::Optional<ffi::Map<ffi::String, ffi::Any>> reuse_write,
      ffi::Optional<ffi::Function> filter_fn = std::nullopt);

  /*!
   * \brief Extension of MultiLevelTiling for auto-tensorization with a single intrinsic.
   * \param intrin_name The name of a tensor intrinsic, must be registered via
   * TensorIntrin.register(...) beforehand
   * \param structure The tiling structure. Recommended:
   * - 'SSRSRS' on CPU
   * - 'SSSRRSRS' on GPU
   * \param tile_binds For each level of tiles, which thread axis it is bound to. Recommended:
   * - std::nullopt on CPU
   * - [blockIdx.x, vthread.x, threadIdx.x] on GPU
   * \param max_innermost_factor The maximum size of the innermost factor. std::nullopt means no
   * limit \param vector_load_lens The length of vector lane in vectorized cooperative fetching.
   * std::nullopt means disable vectorization
   * \param reuse_read Data reuse configuration for reading. std::nullopt means no reuse.
   * \param reuse_write Data reuse configuration for writing. std::nullopt means no reuse.
   * \return The schedule rule created
   */
  TVM_DLL static ScheduleRule MultiLevelTilingWithIntrin(
      ffi::String intrin_name, ffi::String structure,
      ffi::Optional<ffi::Array<ffi::String>> tile_binds,
      ffi::Optional<int64_t> max_innermost_factor,
      ffi::Optional<ffi::Array<int64_t>> vector_load_lens,
      ffi::Optional<ffi::Map<ffi::String, ffi::Any>> reuse_read,
      ffi::Optional<ffi::Map<ffi::String, ffi::Any>> reuse_write);

  /*!
   * \brief Extension of MultiLevelTiling for auto-tensorization with multiple groups of candidate
   * tensor core intrinsics
   * \param intrin_groups A list of groups of tensor core intrinsics. The map should contains key
   * "init", "load_a", "load_b", "compute", "store", which represent the tensor intrin for
   * initialization, loading operand A, loading operand B, tensor core computation, storing the
   * result. The value of the map should be names of tensor intrinsics, must be registered via
   * TensorIntrin.register(...) beforehand
   * \param structure The tiling structure. Recommended:
   * - 'SSSRRSRS' on GPU
   * \param tile_binds For each level of tiles, which thread axis it is bound to. Recommended:
   * - [blockIdx.y, blockIdx.x, threadIdx.y] on GPU
   * \param max_innermost_factor The maximum size of the innermost factor. std::nullopt means no
   * limit \param vector_load_lens The length of vector lane in vectorized cooperative fetching.
   * std::nullopt means disable vectorization
   * \param reuse_read Data reuse configuration for reading. std::nullopt means no reuse.
   * \param reuse_write Data reuse configuration for writing. std::nullopt means no reuse.
   * \param use_software_pipeline Whether use the software pipeline.
   * \return The schedule rule created
   */
  TVM_DLL static ScheduleRule MultiLevelTilingTensorCore(
      ffi::Array<ffi::Map<ffi::String, ffi::String>> intrin_groups, ffi::String structure,
      ffi::Optional<ffi::Array<ffi::String>> tile_binds,
      ffi::Optional<int64_t> max_innermost_factor,
      ffi::Optional<ffi::Array<int64_t>> vector_load_lens,
      ffi::Optional<ffi::Map<ffi::String, ffi::Any>> reuse_read,
      ffi::Optional<ffi::Map<ffi::String, ffi::Any>> reuse_write, bool use_software_pipeline);

  /*!
   * \brief Extension of MultiLevelTiling for backends with wide vectors.
   * The loop over the innermost spatial axis of the output buffer is always vectorized with the
   * maximum vector length.
   * \param structure The tiling structure. 'SSRSRS' is recommended.
   * \param vector_length_in_bits The length of a vector register in bits.
   * \param max_innermost_factor The maximum size of the innermost factor. std::nullopt means no
   * limit \param reuse_read Data reuse configuration for reading. std::nullopt means no reuse.
   * \param reuse_write Data reuse configuration for writing. std::nullopt means no reuse.
   * \return The schedule rule created
   */
  TVM_DLL static ScheduleRule MultiLevelTilingWideVector(
      ffi::String structure, int64_t vector_length_in_bits,
      ffi::Optional<int64_t> max_innermost_factor,
      ffi::Optional<ffi::Map<ffi::String, ffi::Any>> reuse_read,
      ffi::Optional<ffi::Map<ffi::String, ffi::Any>> reuse_write);

  /*!
   * \brief Create a rule: add-rfactor to some blocks if needed
   * \param max_jobs_per_core The maximum number of jobs to be launched per CPU core. It sets the
   * uplimit of CPU parallelism, i.e. `num_cores * max_jobs_per_core`. Use -1 to disable
   * parallelism.
   * \param max_innermost_factor The maximum size of the innermost factor. std::nullopt means no
   * limit \return The schedule rule created
   */
  TVM_DLL static ScheduleRule AddRFactor(int max_jobs_per_core,  //
                                         ffi::Optional<int64_t> max_innermost_factor);
  /*!
   * \brief Create a schedule rule which applies cross-thread reduction to some reduction blocks
   * correspondingly when needed
   * \param thread_extents Candidates of thread axis extent (values are required to be positive).
   * \return The schedule rule created
   */
  TVM_DLL static ScheduleRule CrossThreadReduction(ffi::Array<int64_t> thread_extents);
  /*!
   * \brief A rule that randomly select a compute-at location for a free block
   * \return The schedule rule created
   */
  TVM_DLL static ScheduleRule RandomComputeLocation();
  /*!
   * \brief Mark parallelize, vectorize and unroll to the root block. The mark will be applied to
   * each block in a follow-up post processor
   * \param max_jobs_per_core The maximum number of jobs to be launched per CPU core. It sets the
   * upper limit of CPU parallelism, i.e. `num_cores * max_jobs_per_core`. Use -1 to disable
   * parallelism.
   * \param max_vectorize_extent The maximum extent to be vectorized.
   * It sets the upper limit of the hardware target vectorization. Use -1 to disable vectorization.
   * \param unroll_max_steps The options of the maximum number of unroll steps to be done.
   * Use an empty array to disable unroll.
   * \param unroll_explicit Whether to explicitly unroll the loop, or just add an "unroll" pragma.
   * \return The schedule rule created
   */
  TVM_DLL static ScheduleRule ParallelizeVectorizeUnroll(int max_jobs_per_core,                 //
                                                         int max_vectorize_extent,              //
                                                         ffi::Array<int64_t> unroll_max_steps,  //
                                                         bool unroll_explicit);
  /*!
   * \brief Auto bind loops around the block to BlockIdx and ThreadIdx
   * \param max_threadblocks The maximum number of threadblock on GPU
   * \param thread_extents Candidates of thread axis extent.
   * \param max_threads_per_block The maximum number of threads per block, if it is known
   * when this schedule rule is created.
   * \return The schedule rule created
   */
  TVM_DLL static ScheduleRule AutoBind(int max_threadblocks, ffi::Array<int64_t> thread_extents,
                                       int max_threads_per_block = -1);
  /*!
   * \brief Create a schedule rule with customized methods on the python-side.
   * \param f_initialize_with_tune_context The packed function of `InitializeWithTuneContext`.
   * \param f_apply The packed function of `Apply`.
   * \param f_clone The packed function of `Clone`.
   * \return The schedule rule created.
   */
  TVM_DLL static ScheduleRule PyScheduleRule(
      FInitializeWithTuneContext f_initialize_with_tune_context,  //
      FApply f_apply,                                             //
      FClone f_clone);

  /*! \brief Create default schedule rules for LLVM */
  TVM_DLL static ffi::Array<ScheduleRule, void> DefaultLLVM();
  /*! \brief Create default schedule rules for x86 (AVX512 and VNNI) */
  TVM_DLL static ffi::Array<ScheduleRule, void> DefaultX86(const ffi::String& type);
  /*! \brief Create default schedule rules for CUDA */
  TVM_DLL static ffi::Array<ScheduleRule, void> DefaultCUDA();
  /*! \brief Create default postprocessors for CUDA with TensorCore */
  TVM_DLL static ffi::Array<ScheduleRule, void> DefaultCUDATensorCore();
  /*! \brief Create default schedule rules for Hexagon */
  TVM_DLL static ffi::Array<ScheduleRule, void> DefaultHexagon();
  /*! \brief Create default schedule rules for ARM CPU (NEON and DOTPROD) */
  TVM_DLL static ffi::Array<ScheduleRule, void> DefaultARM(const ffi::String& type);
  /*! \brief Create default schedule rules for RISCV CPU (RVV) */
  TVM_DLL static ffi::Array<ScheduleRule, void> DefaultRISCV(int vlen);

  TVM_FFI_DEFINE_OBJECT_REF_METHODS_NULLABLE(ScheduleRule, ffi::ObjectRef, ScheduleRuleNode);
};

/*! \brief The schedule rule with customized methods on the python-side. */
class PyScheduleRuleNode : public ScheduleRuleNode {
 public:
  using FInitializeWithTuneContext = ScheduleRule::FInitializeWithTuneContext;
  using FApply = ScheduleRule::FApply;
  using FClone = ScheduleRule::FClone;

  /*! \brief The packed function to the `InitializeWithTuneContext` function. */
  FInitializeWithTuneContext f_initialize_with_tune_context;
  /*! \brief The packed function to the `Apply` function. */
  FApply f_apply;
  /*! \brief The packed function to the `Clone` function. */
  FClone f_clone;

  static void RegisterReflection() {
    // `f_initialize_with_tune_context` is not registered
    // `f_apply` is not registered
    // `f_clone` is not registered
    namespace refl = tvm::ffi::reflection;
    refl::ObjectDef<PyScheduleRuleNode>();
  }

  void InitializeWithTuneContext(const TuneContext& context) final;
  ffi::Array<s_tir::Schedule> Apply(const s_tir::Schedule& sch, const s_tir::SBlockRV& block) final;
  ScheduleRule Clone() const final;
  TVM_FFI_DECLARE_OBJECT_INFO_FINAL("s_tir.meta_schedule.PyScheduleRule", PyScheduleRuleNode,
                                    ScheduleRuleNode);
};

}  // namespace meta_schedule
}  // namespace s_tir
}  // namespace tvm

#endif  // TVM_S_TIR_META_SCHEDULE_SCHEDULE_RULE_H_
