tvm
All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros
transform_step.h
Go to the documentation of this file.
1 /*
2  * Licensed to the Apache Software Foundation (ASF) under one
3  * or more contributor license agreements. See the NOTICE file
4  * distributed with this work for additional information
5  * regarding copyright ownership. The ASF licenses this file
6  * to you under the Apache License, Version 2.0 (the
7  * "License"); you may not use this file except in compliance
8  * with the License. You may obtain a copy of the License at
9  *
10  * http://www.apache.org/licenses/LICENSE-2.0
11  *
12  * Unless required by applicable law or agreed to in writing,
13  * software distributed under the License is distributed on an
14  * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
15  * KIND, either express or implied. See the License for the
16  * specific language governing permissions and limitations
17  * under the License.
18  */
19 
45 #ifndef TVM_AUTO_SCHEDULER_TRANSFORM_STEP_H_
46 #define TVM_AUTO_SCHEDULER_TRANSFORM_STEP_H_
47 
48 #include <dmlc/common.h>
49 #include <dmlc/json.h>
50 #include <tvm/node/node.h>
51 #include <tvm/te/schedule.h>
52 
53 #include <vector>
54 
55 namespace tvm {
56 namespace auto_scheduler {
57 
59 
65 void UpdateStageToAxesMap(const te::Stage& stage, StageToAxesMap* stage_to_axes);
66 
68 enum class IteratorKind : int {
70  kSpatial = 0,
72  kReduction = 1,
74  kMixed = 2,
76  kSpecial = 3
77 };
78 
80 enum class IteratorAnnotation : int {
82  kNone = 0,
84  kUnroll = 1,
86  kVectorize = 2,
88  kParallel = 3,
90  kVThread = 4,
92  kBlockX = 5,
94  kThreadX = 6,
96  kBlockY = 7,
98  kThreadY = 8,
100  kBlockZ = 9,
102  kThreadZ = 10,
104  kTensorize = 11
105 };
106 
107 extern const char* IteratorAnnotationString[];
108 
109 // forward declaration
110 class Iterator;
111 
116 class IteratorNode : public Object {
117  public:
127  std::vector<Iterator> orig_iters;
128 
130  v->Visit("name", &name);
131  v->Visit("range", &range);
132  v->Visit("iter_kind", &iter_kind);
133  v->Visit("annotation", &annotation);
134  }
135 
136  static constexpr const char* _type_key = "auto_scheduler.Iterator";
138 };
139 
144 class Iterator : public ObjectRef {
145  public:
154  Iterator(String name, Range range, IteratorKind iter_kind, IteratorAnnotation annotation,
155  const std::vector<Iterator>* orig_iters = nullptr);
156 
158 };
159 
164 class StepNode : public Object {
165  public:
167  int stage_id;
168 
173  virtual void WriteToRecord(dmlc::JSONWriter* writer) const = 0;
174 
175  static constexpr const char* _type_key = "auto_scheduler.Step";
177 };
178 
183 class Step : public ObjectRef {
184  public:
199  StepNode* CopyOnWrite();
200 
202 };
203 
204 // Forward declaration
205 class State;
206 class ComputeDAG;
207 
212 Step StepReadFromRecord(dmlc::JSONReader* reader);
213 
220 void StepApplyToState(const Step& step, State* state, const ComputeDAG& dag);
221 
230 void StepApplyToSchedule(const Step& step, Array<te::Stage>* stages, StageToAxesMap* stage_to_axes,
231  te::Schedule* schedule, const Array<Step>& transform_steps);
232 
242 String StepPrintAsPythonAPI(const Step& step, Array<te::Stage>* stages,
243  StageToAxesMap* stage_to_axes, te::Schedule* schedule,
244  const Array<Step>& transform_steps);
245 
246 /********** Steps working on single stage **********/
247 
252 class AnnotationStepNode : public StepNode {
253  public:
255  int iter_id;
258 
259  void WriteToRecord(dmlc::JSONWriter* writer) const final;
260 
266  Iterator ApplyToState(State* state) const;
267 
273  void ApplyToSchedule(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes) const;
274 
281  String PrintAsPythonAPI(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes) const;
282 
283  static constexpr const char* record_prefix_str = "AN";
284 
285  static constexpr const char* _type_key = "auto_scheduler.AnnotationStep";
287 };
288 
293 class AnnotationStep : public Step {
294  public:
301  AnnotationStep(int stage_id, int iter_id, IteratorAnnotation ann);
302 
308  explicit AnnotationStep(dmlc::JSONReader* reader);
309 
311 };
312 
314 class FuseStepNode : public StepNode {
315  public:
318 
319  void WriteToRecord(dmlc::JSONWriter* writer) const final;
320 
328  Iterator ApplyToState(State* state) const;
329 
336  tir::IterVar ApplyToSchedule(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes) const;
337 
344  String PrintAsPythonAPI(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes) const;
345 
346  static constexpr const char* record_prefix_str = "FU";
347 
348  static constexpr const char* _type_key = "auto_scheduler.FuseStep";
350 };
351 
356 class FuseStep : public Step {
357  public:
363  FuseStep(int stage_id, const Array<Integer>& fused_ids);
364 
370  explicit FuseStep(dmlc::JSONReader* reader);
371 
373 };
374 
376 class PragmaStepNode : public StepNode {
377  public:
379  int iter_id;
382 
383  void WriteToRecord(dmlc::JSONWriter* writer) const final;
384 
389  void ApplyToState(State* state) const;
390 
396  void ApplyToSchedule(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes) const;
397 
404  String PrintAsPythonAPI(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes) const;
405 
406  static constexpr const char* record_prefix_str = "PR";
407 
408  static constexpr const char* _type_key = "auto_scheduler.PragmaStep";
410 };
411 
416 class PragmaStep : public Step {
417  public:
424  PragmaStep(int stage_id, int iter_id, String pragma_type);
425 
431  explicit PragmaStep(dmlc::JSONReader* reader);
432 
434 };
435 
437 class ReorderStepNode : public StepNode {
438  public:
444 
445  void WriteToRecord(dmlc::JSONWriter* writer) const final;
446 
451  void ApplyToState(State* state) const;
452 
458  void ApplyToSchedule(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes) const;
459 
466  String PrintAsPythonAPI(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes) const;
467 
468  static constexpr const char* record_prefix_str = "RE";
469 
470  static constexpr const char* _type_key = "auto_scheduler.ReorderStep";
472 };
473 
478 class ReorderStep : public Step {
479  public:
485  ReorderStep(int stage_id, const Array<Integer>& after_ids);
486 
492  explicit ReorderStep(dmlc::JSONReader* reader);
493 
495 };
496 
501 class SplitStepNode : public StepNode {
502  public:
504  int iter_id;
514 
515  void WriteToRecord(dmlc::JSONWriter* writer) const final;
516 
524  Array<Iterator> ApplyToState(State* state) const;
525 
532  Array<tir::IterVar> ApplyToSchedule(Array<te::Stage>* stages,
533  StageToAxesMap* stage_to_axes) const;
534 
541  String PrintAsPythonAPI(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes) const;
542 
543  static constexpr const char* record_prefix_str = "SP";
544 
545  static constexpr const char* _type_key = "auto_scheduler.SplitStep";
547 };
548 
553 class SplitStep : public Step {
554  public:
563  SplitStep(int stage_id, int iter_id, Optional<PrimExpr> extent,
564  const Array<Optional<Integer>>& lengths, bool inner_to_outer);
565 
571  explicit SplitStep(dmlc::JSONReader* reader);
572 
574 };
575 
579  public:
581  int iter_id;
585  int n_split;
586 
587  void WriteToRecord(dmlc::JSONWriter* writer) const final;
588 
594  Array<Optional<Integer>> ExtractSplitLengths(const Array<Step>& transform_steps) const;
595 
601  Array<Iterator> ApplyToState(State* state) const;
602 
610  Array<tir::IterVar> ApplyToSchedule(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes,
611  const Array<Step>& transform_steps) const;
612 
620  String PrintAsPythonAPI(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes,
621  const Array<Step>& transform_steps) const;
622 
623  static constexpr const char* record_prefix_str = "FSP";
624 
625  static constexpr const char* _type_key = "auto_scheduler.FollowSplitStep";
627 };
628 
633 class FollowSplitStep : public Step {
634  public:
642  FollowSplitStep(int stage_id, int iter_id, int src_step_id, int n_split);
643 
649  explicit FollowSplitStep(dmlc::JSONReader* reader);
650 
652 };
653 
658  public:
660  int iter_id;
664  int level;
667 
668  void WriteToRecord(dmlc::JSONWriter* writer) const final;
669 
675  Optional<Integer> ExtractSplitLength(const Array<Step>& transform_steps) const;
676 
682  Array<Iterator> ApplyToState(State* state) const;
683 
691  Array<tir::IterVar> ApplyToSchedule(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes,
692  const Array<Step>& transform_steps) const;
693 
701  String PrintAsPythonAPI(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes,
702  const Array<Step>& transform_steps) const;
703 
704  static constexpr const char* record_prefix_str = "FFSP";
705 
706  static constexpr const char* _type_key = "auto_scheduler.FollowFusedSplitStep";
708 };
709 
714 class FollowFusedSplitStep : public Step {
715  public:
724  FollowFusedSplitStep(int stage_id, int iter_id, const Array<Integer>& src_step_ids, int level,
725  bool factor_or_nparts);
726 
732  explicit FollowFusedSplitStep(dmlc::JSONReader* reader);
733 
735 };
736 
739  public:
741  int iter_id;
743  int factor;
745  int offset;
746 
747  void WriteToRecord(dmlc::JSONWriter* writer) const final;
748 
753  void ApplyToState(State* state) const;
754 
760  void ApplyToSchedule(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes) const;
761 
768  String PrintAsPythonAPI(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes) const;
769 
770  static constexpr const char* record_prefix_str = "SA";
771 
772  static constexpr const char* _type_key = "auto_scheduler.StorageAlignStep";
774 };
775 
780 class StorageAlignStep : public Step {
781  public:
789  StorageAlignStep(int stage_id, int iter_id, int factor, int offset);
790 
796  explicit StorageAlignStep(dmlc::JSONReader* reader);
797 
799 };
800 
801 /********** Steps working on multiple stages **********/
802 
804 class ComputeAtStepNode : public StepNode {
805  public:
810 
811  void WriteToRecord(dmlc::JSONWriter* writer) const final;
812 
821  void ApplyToState(State* state) const;
822 
828  void ApplyToSchedule(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes) const;
829 
836  String PrintAsPythonAPI(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes) const;
837 
838  static constexpr const char* record_prefix_str = "CA";
839 
840  static constexpr const char* _type_key = "auto_scheduler.ComputeAtStep";
842 };
843 
848 class ComputeAtStep : public Step {
849  public:
856  ComputeAtStep(int stage_id, int target_stage_id, int target_iter_id);
857 
863  explicit ComputeAtStep(dmlc::JSONReader* reader);
864 
866 };
867 
870  public:
871  void WriteToRecord(dmlc::JSONWriter* writer) const final;
872 
877  void ApplyToState(State* state) const;
878 
885  void ApplyToSchedule(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes) const;
886 
893  String PrintAsPythonAPI(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes) const;
894 
895  static constexpr const char* record_prefix_str = "CI";
896 
897  static constexpr const char* _type_key = "auto_scheduler.ComputeInlineStep";
899 };
900 
905 class ComputeInlineStep : public Step {
906  public:
911  explicit ComputeInlineStep(int stage_id);
912 
918  explicit ComputeInlineStep(dmlc::JSONReader* reader);
919 
921 };
922 
925  public:
926  void WriteToRecord(dmlc::JSONWriter* writer) const final;
927 
936  void ApplyToState(State* state) const;
937 
944  void ApplyToSchedule(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes) const;
945 
952  String PrintAsPythonAPI(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes) const;
953 
954  static constexpr const char* record_prefix_str = "CR";
955 
956  static constexpr const char* _type_key = "auto_scheduler.ComputeRootStep";
958 };
959 
964 class ComputeRootStep : public Step {
965  public:
970  explicit ComputeRootStep(int stage_id);
971 
977  explicit ComputeRootStep(dmlc::JSONReader* reader);
978 
980 };
981 
982 /********** Steps adding new stages **********/
983 
989 class CacheReadStepNode : public StepNode {
990  public:
995 
996  void WriteToRecord(dmlc::JSONWriter* writer) const final;
997 
1004  int ApplyToState(State* state, const ComputeDAG& dag) const;
1005 
1013  te::Tensor ApplyToSchedule(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes,
1014  te::Schedule* schedule) const;
1015 
1023  String PrintAsPythonAPI(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes,
1024  te::Schedule* schedule) const;
1025 
1026  static constexpr const char* record_prefix_str = "CHR";
1027 
1028  static constexpr const char* _type_key = "auto_scheduler.CacheReadStep";
1030 };
1031 
1036 class CacheReadStep : public Step {
1037  public:
1044  CacheReadStep(int stage_id, String scope_name, const Array<Integer>& reader_stage_ids);
1045 
1051  explicit CacheReadStep(dmlc::JSONReader* reader);
1052 
1054 };
1055 
1063  public:
1066 
1067  void WriteToRecord(dmlc::JSONWriter* writer) const final;
1068 
1075  int ApplyToState(State* state, const ComputeDAG& dag) const;
1076 
1084  Array<te::Tensor> ApplyToSchedule(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes,
1085  te::Schedule* schedule) const;
1086 
1094  String PrintAsPythonAPI(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes,
1095  te::Schedule* schedule) const;
1096 
1097  static constexpr const char* record_prefix_str = "CHW";
1098 
1099  static constexpr const char* _type_key = "auto_scheduler.CacheWriteStep";
1101 };
1102 
1107 class CacheWriteStep : public Step {
1108  public:
1114  CacheWriteStep(int stage_id, String scope_name);
1115 
1121  explicit CacheWriteStep(dmlc::JSONReader* reader);
1122 
1124 };
1125 
1127 class RfactorStepNode : public StepNode {
1128  public:
1130  int iter_id;
1133 
1134  void WriteToRecord(dmlc::JSONWriter* writer) const final;
1135 
1142  int ApplyToState(State* state, const ComputeDAG& dag) const;
1143 
1151  Array<te::Tensor> ApplyToSchedule(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes,
1152  te::Schedule* schedule) const;
1153 
1161  String PrintAsPythonAPI(Array<te::Stage>* stages, StageToAxesMap* stage_to_axes,
1162  te::Schedule* schedule) const;
1163 
1164  static constexpr const char* record_prefix_str = "RF";
1165 
1166  static constexpr const char* _type_key = "auto_scheduler.RfactorStep";
1168 };
1169 
1174 class RfactorStep : public Step {
1175  public:
1182  RfactorStep(int stage_id, int iter_id, int factor_iter_id);
1183 
1189  explicit RfactorStep(dmlc::JSONReader* reader);
1190 
1192 };
1193 
1194 } // namespace auto_scheduler
1195 } // namespace tvm
1196 
1197 #endif // TVM_AUTO_SCHEDULER_TRANSFORM_STEP_H_
String-aware ObjectRef hash functor.
Definition: base.h:50
Managed reference to ComputeRootStepNode.
Definition: transform_step.h:964
Optional< PrimExpr > extent
The extent length of the axis to split.
Definition: transform_step.h:506
void StepApplyToSchedule(const Step &step, Array< te::Stage > *stages, StageToAxesMap *stage_to_axes, te::Schedule *schedule, const Array< Step > &transform_steps)
Apply a general step to tvm.schedule with runtime dynamic dispatching.
Managed reference to SplitStepNode.
Definition: transform_step.h:553
Managed reference to ComputeDAGNode.
Definition: compute_dag.h:219
Global schedule container For operations and all the operations they depend on. The schedule per Oper...
Definition: schedule.h:318
Definitions and helper macros for IR/AST nodes.
IteratorAnnotation annotation
The annotation type of this iterator.
Definition: transform_step.h:125
Pragma step that corresponds to te::Stage::pragma.
Definition: transform_step.h:376
int iter_id
The id of the iter to be split.
Definition: transform_step.h:581
int factor_iter_id
The position where the new iterator is placed.
Definition: transform_step.h:1132
void StepApplyToState(const Step &step, State *state, const ComputeDAG &dag)
Apply a general step to a State with runtime dynamic dispatching.
Array< Integer > after_ids
The iterator ids after reorder. This array should specify the order of all iterators.
Definition: transform_step.h:443
Fuse step that corresponds to te::Stage::fuse.
Definition: transform_step.h:314
Managed reference to AnnotationStepNode.
Definition: transform_step.h:293
Cache read step that corresponds to te::Schedule::cache_read.
Definition: transform_step.h:989
runtime implementation for LibTorch/TorchScript.
Definition: analyzer.h:36
Map< tvm::te::Stage, Array< tir::IterVar >, ObjectHash, ObjectEqual > StageToAxesMap
Definition: transform_step.h:58
This iterator has been bind to threadIdx.x.
int iter_id
The index of the iterator to be factored.
Definition: transform_step.h:1130
Managed reference to CacheReadStepNode.
Definition: transform_step.h:1036
Managed reference to ComputeAtStepNode.
Definition: transform_step.h:848
int iter_id
The id of the iter to split.
Definition: transform_step.h:660
Iteration Variable, represents an iteration over an integer interval.
Definition: var.h:308
Managed reference to StateNode.
Definition: loop_state.h:272
Cache write step that corresponds to te::Schedule::cache_write.
Definition: transform_step.h:1062
Stage, contains scheduling for a stage of computation.
Definition: schedule.h:58
String name
The name of this iterator.
Definition: transform_step.h:119
IteratorKind
The type of an iterator.
Definition: transform_step.h:68
This iterator has been bind to threadIdx.y.
Managed reference to ComputeInlineStepNode.
Definition: transform_step.h:905
Compute at step that corresponds to te::Stage::compute_at.
Definition: transform_step.h:804
This iterator has been bind to blockIdx.x.
base class of all object containers.
Definition: object.h:167
String StepPrintAsPythonAPI(const Step &step, Array< te::Stage > *stages, StageToAxesMap *stage_to_axes, te::Schedule *schedule, const Array< Step > &transform_steps)
Print a general step as equivalent python schedule API with runtime dynamic dispatching.
String scope_name
The scope name of the newly added read stage. (e.g., local, shared, global)
Definition: transform_step.h:992
IteratorKind iter_kind
The iterator type of this iterator.
Definition: transform_step.h:123
int iter_id
The index of the iterator to add pragma.
Definition: transform_step.h:379
bool factor_or_nparts
If this is true, use factor. Otherwise, use nparts.
Definition: transform_step.h:666
int iter_id
The iterator to be aligned.
Definition: transform_step.h:741
Array< Optional< Integer > > lengths
The split factors.
Definition: transform_step.h:508
Visitor class to get the attributes of an AST/IR node. The content is going to be called for each fie...
Definition: reflection.h:52
int level
Use the length in this split level.
Definition: transform_step.h:664
Similar to SplitStepNode, but uses split factors from another step (i.e. Follow another split step) ...
Definition: transform_step.h:578
String scope_name
The scope name of the newly added compute stage. (e.g. local, shared, global)
Definition: transform_step.h:1065
Range constainer.
Definition: expr.h:715
Similar to FollowSplitStep, but uses split factors from multiple steps.
Definition: transform_step.h:657
An iterator of a for-loop Similar to tvm::IterVar in include/tvm/tir/expr.h
Definition: transform_step.h:116
IteratorAnnotation
The type of an iterator&#39;s annotation.
Definition: transform_step.h:80
std::vector< Iterator > orig_iters
Definition: transform_step.h:127
Fused spatial and reduction iterator.
Array, container representing a contiguous sequence of ObjectRefs.
Definition: array.h:289
Reorder step that corresponds to te::Stage::reorder.
Definition: transform_step.h:437
void VisitAttrs(tvm::AttrVisitor *v)
Definition: transform_step.h:129
IteratorAnnotation annotation
The annotation type of this step.
Definition: transform_step.h:257
bool inner_to_outer
If true, the lengths denote the lengths of iterators from inner level to outer level.
Definition: transform_step.h:513
This iterator has been bind to threadIdx.y.
void UpdateStageToAxesMap(const te::Stage &stage, StageToAxesMap *stage_to_axes)
Update the current stage IterVar information to StageToAxesMap.
The base class of transformation steps. Each step has its corresponding tvm.te schedule primitives...
Definition: transform_step.h:164
Reference to string objects.
Definition: string.h:98
Managed reference to CacheWriteStepNode.
Definition: transform_step.h:1107
String pragma_type
The pragma string.
Definition: transform_step.h:381
Array< Integer > fused_ids
The ids of iterators to fuse.
Definition: transform_step.h:317
This iterator has been vectorized.
String-aware ObjectRef equal functor.
Definition: base.h:40
This iterator has been paralleld.
#define TVM_DEFINE_OBJECT_REF_METHODS(TypeName, ParentType, ObjectName)
Definition: object.h:713
Annotation step that corresponds to vectorize, parallel, unroll and thread binding. (i.e. te::Stage::vectorize, te::Stage::parallel, te::Stage::vectorize, te::Stage::bind)
Definition: transform_step.h:252
Compute root step that corresponds to te::Stage::compute_root.
Definition: transform_step.h:924
const char * IteratorAnnotationString[]
int iter_id
The id of the iter to split.
Definition: transform_step.h:504
This iterator has no annotation.
Array< Integer > reader_stage_ids
The indices of read stages.
Definition: transform_step.h:994
int offset
The offset in the alignment specification.
Definition: transform_step.h:745
Base class of all object reference.
Definition: object.h:511
int n_split
The number of split level.
Definition: transform_step.h:585
Tensor structure representing a possible input, or intermediate computation result.
Definition: tensor.h:102
int target_stage_id
The index of stage that this step will compute at to.
Definition: transform_step.h:807
#define TVM_DECLARE_FINAL_OBJECT_INFO(TypeName, ParentType)
helper macro to declare type information in a final class.
Definition: object.h:671
Managed reference to FollowFusedSplitStepNode.
Definition: transform_step.h:714
This iterator has been bind to vthread.
Reduction factor step that corresponds to te::Schedule::rfactor.
Definition: transform_step.h:1127
Managed reference to RfactorStepNode.
Definition: transform_step.h:1174
int target_iter_id
The index of iterator in target stage that this step will compute at to.
Definition: transform_step.h:809
Step StepReadFromRecord(dmlc::JSONReader *reader)
Read a step record from JSONReader and create the corresponding step.
int iter_id
The index of the iterator to add annotation.
Definition: transform_step.h:255
Map container of NodeRef->NodeRef in DSL graph. Map implements copy on write semantics, which means map is mutable but copy will happen when array is referenced in more than two places.
Definition: map.h:1271
Managed reference to IteratorNode.
Definition: transform_step.h:144
Split step that corresponds to te::Stage::split with additional support of multiple-level of factors...
Definition: transform_step.h:501
Optional container that to represent to a Nullable variant of T.
Definition: optional.h:51
int src_step_id
The index of the split step to be followed in the history.
Definition: transform_step.h:583
int factor
The factor in alignment specification.
Definition: transform_step.h:743
Managed reference to ReorderStepNode.
Definition: transform_step.h:478
Managed reference to FuseStepNode.
Definition: transform_step.h:356
Compute inline step that corresponds to te::Stage::compute_inline.
Definition: transform_step.h:869
Storage align step that corresponds to te::Stage::storage_align.
Definition: transform_step.h:738
This iterator has been mapped with a tensorize intrinsic.
Special iterator. (e.g. virtual root iterator)
This iterator has been unrolled.
Managed reference to StorageAlignStepNode.
Definition: transform_step.h:780
Managed reference to StepNode.
Definition: transform_step.h:183
Managed reference to FollowSplitStepNode.
Definition: transform_step.h:633
#define TVM_DECLARE_BASE_OBJECT_INFO(TypeName, ParentType)
helper macro to declare a base object type that can be inherited.
Definition: object.h:648
Range range
The range of this iterator.
Definition: transform_step.h:121
Define a schedule.
This iterator has been bind to blockIdx.y.
Array< Integer > src_step_ids
The indices of the split steps to be followed in the history.
Definition: transform_step.h:662
Managed reference to PragmaStepNode.
Definition: transform_step.h:416
int stage_id
The index of the stage.
Definition: transform_step.h:167
This iterator has been bind to blockIdx.y.