24 #ifndef TVM_TIRX_STMT_H_
25 #define TVM_TIRX_STMT_H_
27 #include <tvm/ffi/reflection/registry.h>
34 #include <type_traits>
53 namespace refl = tvm::ffi::reflection;
64 class Stmt :
public ffi::ObjectRef {
84 namespace refl = tvm::ffi::reflection;
85 refl::ObjectDef<BindNode>()
87 .def_ro(
"var", &
BindNode::var, refl::AttachFieldFlag::SEqHashDefRecursive())
127 namespace refl = tvm::ffi::reflection;
128 refl::ObjectDef<AttrStmtNode>()
169 namespace refl = tvm::ffi::reflection;
170 refl::ObjectDef<AssertStmtNode>()
213 namespace refl = tvm::ffi::reflection;
214 refl::ObjectDef<BufferStoreNode>()
230 ffi::Optional<PrimExpr> predicate = std::nullopt,
244 namespace refl = tvm::ffi::reflection;
272 namespace refl = tvm::ffi::reflection;
273 refl::ObjectDef<AllocBufferNode>()
286 ffi::Map<ffi::String, ffi::Any> annotations = ffi::Map<ffi::String, ffi::Any>(),
294 for (
const PrimExpr& extent : (*this)->buffer->shape) {
295 if (
const auto* int_size = extent.as<
IntImmNode>()) {
296 result *= int_size->value;
325 namespace refl = tvm::ffi::reflection;
343 namespace refl = tvm::ffi::reflection;
374 size_t size()
const {
return operator->()->size(); }
399 template <
typename... Args>
401 ffi::Array<Stmt> seq;
403 ffi::details::for_each(
Flattener(&seq), std::forward<Args>(seq_args)...);
407 }
else if (seq.size() == 1) {
414 if constexpr (
sizeof...(seq_args) == 1) {
416 SeqStmt original = opt.value();
417 bool all_same = [&]() {
418 if (original->seq.
size() != seq.size()) {
421 for (
size_t i = 0; i < seq.size(); i++) {
422 if (!original->seq[i].same_as(seq[i])) {
439 explicit Flattener(ffi::Array<Stmt>* seq) : seq_(seq) {}
441 template <
typename T>
443 if constexpr (std::is_same_v<T, SeqStmt>) {
446 if constexpr (!std::is_base_of_v<T, SeqStmt>) {
449 if constexpr (std::is_base_of_v<Stmt, T>) {
450 if (
const SeqStmtNode* ptr = t.template as<SeqStmtNode>()) {
451 return ffi::GetRef<SeqStmt>(ptr);
459 template <
typename T>
461 if constexpr (std::is_base_of_v<ObjectRef, T>) {
463 if (!stmt_or_seq.defined()) {
468 if constexpr (std::is_same_v<T, SeqStmt>) {
470 (*this)(0, stmt_or_seq->seq);
474 if constexpr (std::is_base_of_v<T, SeqStmt>) {
477 if (
auto* op = stmt_or_seq.template as<SeqStmtNode>()) {
483 if constexpr (std::is_base_of_v<T, Evaluate>) {
488 if (
auto* op = stmt_or_seq.template as<EvaluateNode>()) {
489 if (
auto* as_int = op->value.template as<IntImmNode>(); as_int && as_int->value == 0) {
495 if constexpr (std::is_base_of_v<Stmt, T>) {
497 seq_->push_back(stmt_or_seq);
500 for (
auto v : stmt_or_seq) {
507 ffi::Array<Stmt>* seq_;
527 namespace refl = tvm::ffi::reflection;
528 refl::ObjectDef<IfThenElseNode>()
543 ffi::Optional<Stmt> else_case = std::nullopt,
Span span =
Span());
619 namespace refl = tvm::ffi::reflection;
620 refl::ObjectDef<ForNode>()
621 .def_ro(
"loop_var", &
ForNode::loop_var, refl::AttachFieldFlag::SEqHashDefRecursive())
644 ffi::Optional<IterVar> thread_binding = std::nullopt,
645 ffi::Map<ffi::String, ffi::Any> annotations = {},
646 ffi::Optional<PrimExpr> step = std::nullopt,
Span span =
Span());
670 namespace refl = tvm::ffi::reflection;
671 refl::ObjectDef<WhileNode>()
696 namespace refl = tvm::ffi::reflection;
697 refl::ObjectDef<BreakNode>();
721 namespace refl = tvm::ffi::reflection;
722 refl::ObjectDef<ContinueNode>();
751 namespace refl = tvm::ffi::reflection;
752 refl::ObjectDef<BufferRegionNode>()
807 namespace refl = tvm::ffi::reflection;
808 refl::ObjectDef<MatchBufferRegionNode>()
879 namespace refl = tvm::ffi::reflection;
880 refl::ObjectDef<SBlockNode>()
901 ffi::Array<IterVar> iter_vars, ffi::Array<BufferRegion> reads,
902 ffi::Array<BufferRegion> writes, ffi::String name_hint,
Stmt body,
903 ffi::Optional<Stmt> init = std::nullopt,
904 ffi::Array<Buffer> alloc_buffers = ffi::Array<Buffer>(),
905 ffi::Array<MatchBufferRegion> match_buffers = ffi::Array<MatchBufferRegion>(),
906 ffi::Map<ffi::String, ffi::Any> annotations = ffi::Map<ffi::String, ffi::Any>(),
910 ffi::Array<Buffer> alloc_buffers = ffi::Array<Buffer>(),
933 namespace refl = tvm::ffi::reflection;
934 refl::ObjectDef<SBlockRealizeNode>()
969 namespace refl = tvm::ffi::reflection;
1157 "tirx.manifest_shared_memory_local_stage";
1170 "meta_schedule.thread_extent_low_inclusive";
1174 "meta_schedule.thread_extent_high_inclusive";
1178 "meta_schedule.random_compute_producer";
1275 return attr_key.compare(0, 7,
"pragma_") == 0;
1298 return "vectorized";
1302 return "thread_binding";
1304 TVM_FFI_THROW(InternalError) <<
"Unknown ForKind" << t;
1305 TVM_FFI_UNREACHABLE();
Managed reference to ExprNode.
Definition: base_expr.h:311
Constant integer literals in the program.
Definition: expr.h:361
Base class for other IR constructs that can be converted to PrimExpr. This is useful for the FFI to c...
Definition: base_expr.h:391
Managed reference to PrimExprConvertibleNode.
Definition: base_expr.h:402
Typed reference/view over any Expr whose ExprNode::ty is PrimType.
Definition: base_expr.h:354
Definition: base_expr.h:113
Definition: source_map.h:111
Allocate a buffer and declare it in scope.
Definition: stmt.h:259
ffi::Map< ffi::String, ffi::Any > annotations
Additional annotations about the allocation.
Definition: stmt.h:269
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("tirx.AllocBuffer", AllocBufferNode, StmtNode)
static void RegisterReflection()
Definition: stmt.h:271
Buffer buffer
The buffer being allocated and declared.
Definition: stmt.h:262
Managed reference to AllocBufferNode.
Definition: stmt.h:282
AllocBuffer(Buffer buffer, ffi::Map< ffi::String, ffi::Any > annotations=ffi::Map< ffi::String, ffi::Any >(), Span span=Span())
std::optional< int64_t > ConstantAllocationSize() const
If the buffer's shape is constant, return the total number of elements.
Definition: stmt.h:292
TVM_DEFINE_OBJECT_REF_COW_METHOD(AllocBufferNode)
TVM_FFI_DEFINE_OBJECT_REF_METHODS_NULLABLE(AllocBuffer, Stmt, AllocBufferNode)
Assert condition, if an error occurs, return the error message.
Definition: stmt.h:159
PrimExpr condition
Condition to be checked.
Definition: stmt.h:162
StringImm error_kind
The error kind, e.g. "RuntimeError", "TypeError", "ValueError".
Definition: stmt.h:164
static void RegisterReflection()
Definition: stmt.h:168
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("tirx.AssertStmt", AssertStmtNode, StmtNode)
ffi::Array< StringImm > message_parts
Error message fragments, concatenated at runtime when assertion fails.
Definition: stmt.h:166
Managed reference to AssertStmtNode.
Definition: stmt.h:182
TVM_FFI_DEFINE_OBJECT_REF_METHODS_NULLABLE(AssertStmt, Stmt, AssertStmtNode)
AssertStmt(PrimExpr condition, StringImm error_kind, ffi::Array< StringImm > message_parts, Span span=Span())
TVM_DEFINE_OBJECT_REF_COW_METHOD(AssertStmtNode)
Define certain auxiliary attribute for the body to be a symbolic value. This provide auxiliary inform...
Definition: stmt.h:115
static void RegisterReflection()
Definition: stmt.h:126
PrimExpr value
The attribute value, value is well defined at current scope.
Definition: stmt.h:122
ffi::Any node
this is attribute about certain node
Definition: stmt.h:118
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("tirx.AttrStmt", AttrStmtNode, StmtNode)
ffi::String attr_key
the type key of the attribute
Definition: stmt.h:120
Stmt body
The body statement to be executed.
Definition: stmt.h:124
Managed reference to AttrStmtNode.
Definition: stmt.h:141
TVM_FFI_DEFINE_OBJECT_REF_METHODS_NULLABLE(AttrStmt, Stmt, AttrStmtNode)
AttrStmt(ffi::Any node, ffi::String attr_key, PrimExpr value, Stmt body, Span span=Span())
TVM_DEFINE_OBJECT_REF_COW_METHOD(AttrStmtNode)
Bind a variable to a value in the enclosing scope.
Definition: stmt.h:76
Expr value
The value to bind to the variable.
Definition: stmt.h:81
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("tirx.Bind", BindNode, StmtNode)
static void RegisterReflection()
Definition: stmt.h:83
Var var
The variable being bound.
Definition: stmt.h:79
Managed reference to BindNode.
Definition: stmt.h:97
Bind(Var var, Expr value, Span span=Span())
TVM_DEFINE_OBJECT_REF_COW_METHOD(BindNode)
TVM_FFI_DEFINE_OBJECT_REF_METHODS_NULLABLE(Bind, Stmt, BindNode)
A Break in control flow.
Definition: stmt.h:693
static void RegisterReflection()
Definition: stmt.h:695
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("tirx.Break", BreakNode, StmtNode)
Managed reference to BreakNode.
Definition: stmt.h:707
TVM_FFI_DEFINE_OBJECT_REF_METHODS_NULLABLE(Break, Stmt, BreakNode)
TVM_DEFINE_OBJECT_REF_COW_METHOD(BreakNode)
Representing the region of multi-dimensional buffer access.
Definition: stmt.h:743
ffi::Array< Range > region
The region array of the buffer region.
Definition: stmt.h:748
PrimExpr ToPrimExpr() const final
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("tirx.BufferRegion", BufferRegionNode, PrimExprConvertibleNode)
static void RegisterReflection()
Definition: stmt.h:750
static constexpr TVMFFISEqHashKind _type_s_eq_hash_kind
Definition: stmt.h:759
Buffer buffer
The buffer of the buffer region.
Definition: stmt.h:746
Managed reference to BufferRegionNode.
Definition: stmt.h:767
static BufferRegion FullRegion(Buffer buffer)
Create a BufferRegion which is full region of the given buffer.
TVM_FFI_DEFINE_OBJECT_REF_METHODS_NULLABLE(BufferRegion, PrimExprConvertible, BufferRegionNode)
BufferRegion(Buffer buffer, ffi::Array< Range > region)
static BufferRegion FromPoint(Buffer buffer, ffi::Array< PrimExpr > indices)
Create a BufferRegion which is a single point of the given buffer.
TVM_DEFINE_OBJECT_REF_COW_METHOD(BufferRegionNode)
Store value to the high dimension buffer.
Definition: stmt.h:201
Buffer buffer
The buffer variable.
Definition: stmt.h:204
ffi::Array< PrimExpr > indices
The indices location to be stored.
Definition: stmt.h:208
static void RegisterReflection()
Definition: stmt.h:212
PrimExpr value
The value to be stored.
Definition: stmt.h:206
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("tirx.BufferStore", BufferStoreNode, StmtNode)
ffi::Optional< PrimExpr > predicate
The predicate mask for storing values.
Definition: stmt.h:210
Managed reference to BufferStoreNode.
Definition: stmt.h:227
BufferStore(Buffer buffer, PrimExpr value, ffi::Array< PrimExpr > indices, ffi::Optional< PrimExpr > predicate=std::nullopt, Span span=Span())
TVM_DEFINE_OBJECT_REF_COW_METHOD(BufferStoreNode)
TVM_FFI_DEFINE_OBJECT_REF_METHODS_NULLABLE(BufferStore, Stmt, BufferStoreNode)
Buffer is a symbolic n-darray structure. It is a composition of primitive symbolic types,...
Definition: buffer.h:187
A Continue in control flow.
Definition: stmt.h:718
static void RegisterReflection()
Definition: stmt.h:720
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("tirx.Continue", ContinueNode, StmtNode)
Managed reference to ContinueNode.
Definition: stmt.h:732
TVM_DEFINE_OBJECT_REF_COW_METHOD(ContinueNode)
TVM_FFI_DEFINE_OBJECT_REF_METHODS_NULLABLE(Continue, Stmt, ContinueNode)
Declare a buffer that can be used in the body.
Definition: stmt.h:238
Buffer buffer
The buffer being declared.
Definition: stmt.h:241
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("tirx.DeclBuffer", DeclBufferNode, StmtNode)
static void RegisterReflection()
Definition: stmt.h:243
Managed reference to DeclBufferNode.
Definition: stmt.h:251
TVM_FFI_DEFINE_OBJECT_REF_METHODS_NULLABLE(DeclBuffer, Stmt, DeclBufferNode)
DeclBuffer(Buffer buffer, Span span=Span())
TVM_DEFINE_OBJECT_REF_COW_METHOD(DeclBufferNode)
Evaluates an expression. This is mostly used for putting a Call node into Stmt.
Definition: stmt.h:337
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("tirx.Evaluate", EvaluateNode, StmtNode)
Expr value
The expression to be evaluated.
Definition: stmt.h:340
static void RegisterReflection()
Definition: stmt.h:342
Managed reference to EvaluateNode.
Definition: stmt.h:353
Evaluate(Expr value, Span span=Span())
Evaluate(int value, Span span=Span())
Definition: stmt.h:357
TVM_DEFINE_OBJECT_REF_COW_METHOD(EvaluateNode)
TVM_FFI_DEFINE_OBJECT_REF_METHODS_NULLABLE(Evaluate, Stmt, EvaluateNode)
A for loop, with possible type annotations.
Definition: stmt.h:587
PrimExpr min
The minimum value of iteration.
Definition: stmt.h:592
ffi::Optional< IterVar > thread_binding
Only valid when kind == ForKind::kThreadBinding The context thread that this loop variable bounds to.
Definition: stmt.h:603
ffi::Optional< PrimExpr > step
The loop step. It is one if not specified.
Definition: stmt.h:616
bool HasTrivialStep() const
Check it is a loop without nontrivial loop step.
Stmt body
The body of the for loop.
Definition: stmt.h:598
PrimExpr extent
The extent of the iteration.
Definition: stmt.h:594
ForKind kind
The kind of the for loop.
Definition: stmt.h:596
ffi::Map< ffi::String, ffi::Any > annotations
Additional annotations about the loop.
Definition: stmt.h:612
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("tirx.For", ForNode, StmtNode)
PrimVar loop_var
The loop variable.
Definition: stmt.h:590
static void RegisterReflection()
Definition: stmt.h:618
Managed reference to ForNode.
Definition: stmt.h:641
For(PrimVar loop_var, PrimExpr min, PrimExpr extent, ForKind kind, Stmt body, ffi::Optional< IterVar > thread_binding=std::nullopt, ffi::Map< ffi::String, ffi::Any > annotations={}, ffi::Optional< PrimExpr > step=std::nullopt, Span span=Span())
TVM_FFI_DEFINE_OBJECT_REF_METHODS_NULLABLE(For, Stmt, ForNode)
TVM_DEFINE_OBJECT_REF_COW_METHOD(ForNode)
IfThenElse statement.
Definition: stmt.h:517
static void RegisterReflection()
Definition: stmt.h:526
ffi::Optional< Stmt > else_case
The branch to be executed when condition is false, can be null.
Definition: stmt.h:524
PrimExpr condition
The condition.
Definition: stmt.h:520
Stmt then_case
The branch to be executed when condition is true.
Definition: stmt.h:522
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("tirx.IfThenElse", IfThenElseNode, StmtNode)
Managed reference to IfThenElseNode.
Definition: stmt.h:540
TVM_DEFINE_OBJECT_REF_COW_METHOD(IfThenElseNode)
TVM_FFI_DEFINE_OBJECT_REF_METHODS_NULLABLE(IfThenElse, Stmt, IfThenElseNode)
IfThenElse(PrimExpr condition, Stmt then_case, ffi::Optional< Stmt > else_case=std::nullopt, Span span=Span())
Match introduces a constraint that the source buffer region can be remapped to the data layout specif...
Definition: stmt.h:799
Buffer buffer
The target buffer.
Definition: stmt.h:802
BufferRegion source
The source buffer region.
Definition: stmt.h:804
static void RegisterReflection()
Definition: stmt.h:806
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("tirx.MatchBufferRegion", MatchBufferRegionNode, ffi::Object)
Managed reference to MatchBufferRegionNode.
Definition: stmt.h:821
TVM_DEFINE_OBJECT_REF_COW_METHOD(MatchBufferRegionNode)
MatchBufferRegion(Buffer buffer, BufferRegion source)
TVM_FFI_DEFINE_OBJECT_REF_METHODS_NULLABLE(MatchBufferRegion, ffi::ObjectRef, MatchBufferRegionNode)
Checked scalar view over a VarNode.
Definition: var.h:127
A block is a basic schedule unit in TIR.
Definition: stmt.h:851
ffi::Array< Buffer > alloc_buffers
The buffer allocated in the block.
Definition: stmt.h:862
Stmt body
The body of the block.
Definition: stmt.h:876
ffi::Array< BufferRegion > reads
The read buffer regions of the block.
Definition: stmt.h:856
static void RegisterReflection()
Definition: stmt.h:878
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("tirx.SBlock", SBlockNode, StmtNode)
ffi::Array< MatchBufferRegion > match_buffers
The match buffer regions.
Definition: stmt.h:864
ffi::String name_hint
The name_hint of the block.
Definition: stmt.h:860
ffi::Optional< Stmt > init
The init statement is executed during the first iteration of reduction loops in a reduction block....
Definition: stmt.h:874
ffi::Array< IterVar > iter_vars
The variables of the block.
Definition: stmt.h:854
ffi::Array< BufferRegion > writes
The write buffer regions of the block.
Definition: stmt.h:858
ffi::Map< ffi::String, ffi::Any > annotations
The annotation of the block.
Definition: stmt.h:866
A block realization node represents execution of the block at the binding values.
Definition: stmt.h:920
static void RegisterReflection()
Definition: stmt.h:932
PrimExpr predicate
The predicate of the block realization, the block will only be executed when the predicate is true.
Definition: stmt.h:928
ffi::Array< PrimExpr > iter_values
The corresponding values of the iter vars.
Definition: stmt.h:923
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("tirx.SBlockRealize", SBlockRealizeNode, StmtNode)
SBlock block
The block to be realized.
Definition: stmt.h:930
Managed reference to BlockRealizeNode.
Definition: stmt.h:946
TVM_DEFINE_OBJECT_REF_COW_METHOD(SBlockRealizeNode)
SBlockRealize(ffi::Array< PrimExpr > iter_values, PrimExpr predicate, SBlock block, Span span=Span())
TVM_FFI_DEFINE_OBJECT_REF_METHODS_NULLABLE(SBlockRealize, Stmt, SBlockRealizeNode)
Managed reference to SBlockNode.
Definition: stmt.h:898
SBlock(ffi::String name_hint, Stmt body, ffi::Array< Buffer > alloc_buffers=ffi::Array< Buffer >(), Span span=Span())
TVM_FFI_DEFINE_OBJECT_REF_METHODS_NULLABLE(SBlock, Stmt, SBlockNode)
TVM_DEFINE_OBJECT_REF_COW_METHOD(SBlockNode)
SBlock(ffi::Array< IterVar > iter_vars, ffi::Array< BufferRegion > reads, ffi::Array< BufferRegion > writes, ffi::String name_hint, Stmt body, ffi::Optional< Stmt > init=std::nullopt, ffi::Array< Buffer > alloc_buffers=ffi::Array< Buffer >(), ffi::Array< MatchBufferRegion > match_buffers=ffi::Array< MatchBufferRegion >(), ffi::Map< ffi::String, ffi::Any > annotations=ffi::Map< ffi::String, ffi::Any >(), Span span=Span())
Standalone statement that declares a scope-id binding (e.g. cta_id, warp_id, lane_id)....
Definition: stmt.h:963
ScopeIdDef def
The scope-id definition (Vars + extents + binding).
Definition: stmt.h:966
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("tirx.ScopeIdDefStmt", ScopeIdDefStmtNode, StmtNode)
static void RegisterReflection()
Definition: stmt.h:968
Managed reference to ScopeIdDefStmtNode.
Definition: stmt.h:976
ScopeIdDefStmt(ScopeIdDef def, Span span=Span())
TVM_FFI_DEFINE_OBJECT_REF_METHODS_NULLABLE(ScopeIdDefStmt, Stmt, ScopeIdDefStmtNode)
TVM_DEFINE_OBJECT_REF_COW_METHOD(ScopeIdDefStmtNode)
Definition: exec_scope.h:137
The container of seq statement. Represent a sequence of statements.
Definition: stmt.h:312
size_t size() const
Definition: stmt.h:318
Stmt operator[](size_t index) const
Get the index-th element in the sequence.
Definition: stmt.h:322
ffi::Array< Stmt > seq
internal sequence content.
Definition: stmt.h:315
static void RegisterReflection()
Definition: stmt.h:324
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("tirx.SeqStmt", SeqStmtNode, StmtNode)
Helper class to flatten sequence of arguments into Array.
Definition: stmt.h:437
static ffi::Optional< SeqStmt > AsSeqStmt(const T &t)
Definition: stmt.h:442
void operator()(size_t i, const T &stmt_or_seq) const
Definition: stmt.h:460
Flattener(ffi::Array< Stmt > *seq)
Definition: stmt.h:439
Sequence statement.
Definition: stmt.h:364
TVM_DEFINE_OBJECT_REF_COW_METHOD(SeqStmtNode)
Stmt operator[](size_t index) const
Get the index-th element in the sequence.
Definition: stmt.h:378
static Stmt Flatten(Args &&... seq_args)
Construct a sequence statement by flattening all the arrays and sequences in the arguments recursivel...
Definition: stmt.h:400
size_t size() const
Definition: stmt.h:374
SeqStmt(ffi::Array< Stmt > seq, Span span=Span())
Construct SeqStmt.
TVM_FFI_DEFINE_OBJECT_REF_METHODS_NULLABLE(SeqStmt, Stmt, SeqStmtNode)
Base node of all statements.
Definition: stmt.h:41
static constexpr const uint32_t _type_child_slots
Definition: stmt.h:59
static constexpr TVMFFISEqHashKind _type_s_eq_hash_kind
Definition: stmt.h:57
Span span
Span that points to the original source code. Reserved debug information.
Definition: stmt.h:47
StmtNode(Span span)
Definition: stmt.h:50
TVM_FFI_DECLARE_OBJECT_INFO("tirx.Stmt", StmtNode, ffi::Object)
static void RegisterReflection()
Definition: stmt.h:52
Container of all statements.
Definition: stmt.h:64
TVM_FFI_DEFINE_OBJECT_REF_METHODS_NULLABLE(Stmt, ffi::ObjectRef, StmtNode)
Managed reference to StringImmNode.
Definition: expr.h:69
a named variable in TIR
Definition: var.h:68
A While loop.
Definition: stmt.h:662
static void RegisterReflection()
Definition: stmt.h:669
PrimExpr condition
The termination condition.
Definition: stmt.h:665
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("tirx.While", WhileNode, StmtNode)
Stmt body
The body of the while loop.
Definition: stmt.h:667
Managed reference to WhileNode.
Definition: stmt.h:682
TVM_DEFINE_OBJECT_REF_COW_METHOD(WhileNode)
While(PrimExpr condition, Stmt body, Span span=Span())
TVM_FFI_DEFINE_OBJECT_REF_METHODS_NULLABLE(While, Stmt, WhileNode)
void Evaluate(Expr value)
Evaluate the input expression.
PrimVar var(std::string name_hint, PrimType t=PrimType::Int(32))
Construct a new Var expression.
constexpr const char * hand_threaded
Mark that the kernel is hand threaded and doesn't need syncs inserted.
Definition: stmt.h:1124
constexpr const char * script_parsing_detect_access
Mark whether the script-completer need to fill in missing access region during script parsing.
Definition: stmt.h:1133
constexpr const char * software_pipeline_stage
Mark the stage of a statement in the software pipeline.
Definition: stmt.h:1141
constexpr const char * irregular_loop_mark
,ark a ForNode represent an irregular loop of non-structural control flow edges.
Definition: stmt.h:1254
constexpr const char * fragment_layout
Mark that the layout of TensorCore fragment.
Definition: stmt.h:1119
constexpr const char * manifest_shared_memory_local_stage
Mark the local stage for the shared memory access should be added.
Definition: stmt.h:1156
constexpr const char * meta_schedule_vectorize
Mark auto-vectorize setting on the block.
Definition: stmt.h:1184
constexpr const char * explicit_write_region
Mark that a block has an explicitly specified write region. This is used to override the default writ...
Definition: stmt.h:1250
constexpr const int meta_schedule_cache_type_read
Definition: stmt.h:1219
constexpr const char * channel_write_scope
channel write scope
Definition: stmt.h:1077
constexpr const char * auto_copy
Mark auto copy for memhammer.
Definition: stmt.h:1225
constexpr const char * thread_extent
Mark launching extent of thread, used by device API.
Definition: stmt.h:1020
bool IsPragmaKey(const std::string &attr_key)
Check if attr_key is a pragma key extension.
Definition: stmt.h:1274
constexpr const char * double_buffer_scope
Marks production of double buffer data.
Definition: stmt.h:1041
constexpr const char * software_pipeline_async_stages
List stages in the software pipeline that should run asynchronously.
Definition: stmt.h:1150
constexpr const char * meta_schedule_cache_type
Mark a block as generated by cache_read or cache_write block. 0 means cache_read; 1 means cache_write...
Definition: stmt.h:1216
constexpr const char * meta_schedule_unroll_implicit
Mark auto-unroll setting on the block.
Definition: stmt.h:1190
constexpr const char * meta_schedule_thread_extent_low_inclusive
The allowed range of thread extent in thread bindings.
Definition: stmt.h:1169
constexpr const char * warp_execution
Mark that a block is executed by a warp. This implies the extend of threadIdx.x is warp size.
Definition: stmt.h:1237
constexpr const char * explicit_read_region
Mark that a block has an explicitly specified read region. This is used to override the default read ...
Definition: stmt.h:1245
constexpr const char * software_pipeline_order
Mark the order of a statement in the software pipeline.
Definition: stmt.h:1144
constexpr const char * pragma_unroll_explicit
Pragma: unroll explicit.
Definition: stmt.h:1016
constexpr const char * storage_alignment
Mark storage alignment requirement of buffers.
Definition: stmt.h:1018
constexpr const char * buffer_bound
Mark stores/loads with their bounds.
Definition: stmt.h:987
constexpr const char * meta_schedule_cooperative_fetch
Mark that the loop should be further skip and bound to environment threads to enable cooperative fetc...
Definition: stmt.h:1166
constexpr const char * scan_init_scope
Mark of scan init scope.
Definition: stmt.h:1049
constexpr const char * layout_free_buffers
Mark the buffers which is const access and can be transformed layout.
Definition: stmt.h:1153
constexpr const char * buffer_data_alignment
Mark buffer initial addr alignment in bytes.
Definition: stmt.h:1058
constexpr const char * meta_schedule_inline_rule
Mark that a block is disallowed in auto inline.
Definition: stmt.h:1240
constexpr const char * kPersistentKernel
Mark the kernel as persistent.
Definition: stmt.h:1259
constexpr const char * device_id
The allocation device for global malloc in host.
Definition: stmt.h:994
constexpr const char * meta_schedule_layout_rewrite_preproc
Mark that a block is a preprocessor block for layout rewrite.
Definition: stmt.h:1196
constexpr const char * scan_update_scope
Mark of scan update scope.
Definition: stmt.h:1047
constexpr const char * pragma_auto_unroll_max_step
Pragma: auto-unroll, max_step.
Definition: stmt.h:1006
constexpr const char * async_commit_queue_scope
Annotations for invoking and synchronizing asynchronous operations.
Definition: stmt.h:1107
constexpr const char * tensorized_nki_instruction
Definition: stmt.h:1251
constexpr const char * layout_transforms
Marks the layout transforms to be used for a tensor.
Definition: stmt.h:1029
constexpr const char * pragma_import_llvm
Import llvm source or file into the final code gen module.
Definition: stmt.h:1010
constexpr const char * meta_schedule_auto_tensorize_init
Mark that the init statement of a block should be further rewritten using tensorization.
Definition: stmt.h:1200
constexpr const char * pragma_import_c
Import C source or file into the final code gen module.
Definition: stmt.h:1008
constexpr const char * meta_schedule_tensor_core_enabled
Mark that tensor core is enabled in the PrimExpr.
Definition: stmt.h:1208
constexpr const char * compute_scope
Mark the scope as when computation start to happen. This can hint some code generator to create a new...
Definition: stmt.h:992
constexpr const char * async_scope
Mark that the attached statement runs asynchronously.
Definition: stmt.h:1088
constexpr const char * pragma_tensor_core
Try to modify the AST to support Tensor Core.
Definition: stmt.h:1014
constexpr const char * buffer_allocated_addr
Mark buffer allocated addr in bytes.
Definition: stmt.h:1060
constexpr const char * buffer_bind_scope
Bind the buffer specification to the region of the op When this scope occurs, the stmt....
Definition: stmt.h:1070
constexpr const char * meta_schedule_unroll_explicit
Mark auto-unroll setting on the block.
Definition: stmt.h:1187
constexpr const char * async_wait_queue_scope
Definition: stmt.h:1108
constexpr const char * meta_schedule_auto_tensorize
Mark that a block should be further rewritten using tensorization.
Definition: stmt.h:1193
constexpr const char * channel_read_scope
channel read scope
Definition: stmt.h:1073
constexpr const char * meta_schedule_parallel
Mark auto-parallel setting on the block.
Definition: stmt.h:1181
constexpr const char * device_type
The device type.
Definition: stmt.h:998
constexpr const char * axis_separators
Marks the physical axis separators.
Definition: stmt.h:1037
constexpr const char * async_wait_inflight_count
Definition: stmt.h:1109
constexpr const char * channel_read_advance
Advance step of channel after end of scope.
Definition: stmt.h:1075
constexpr const char * kVolatile
Annotation key on AllocBuffer marking the allocation as volatile.
Definition: stmt.h:1022
constexpr const char * device_scope
Mark that it is in the device scope.
Definition: stmt.h:996
constexpr const char * extern_scope
Mark the scope as generated by extern primitive. Such scope can contain arbitrary ir program and we n...
Definition: stmt.h:1004
constexpr const char * double_buffer_write
Marks region used by double buffer write.
Definition: stmt.h:1045
constexpr const char * pipeline_exec_scope
pipeline execution scope, implies the scope can be pipelined.
Definition: stmt.h:1083
constexpr const char * local_stage
Mark local stage constraint on data copy.
Definition: stmt.h:1228
constexpr const char * fragment_shape
Mark that the shape of TensorCore fragment.
Definition: stmt.h:1114
constexpr const char * pipeline_stage_scope
pipeline stage scope, implies always execution
Definition: stmt.h:1081
constexpr const int meta_schedule_cache_type_write
Definition: stmt.h:1222
constexpr const char * kDeviceEntry
Mark the device-region entry within a PrimFunc body. The AttrStmt so-keyed has a body that is the dev...
Definition: stmt.h:1267
constexpr const char * channel_write_advance
Advance step of channel after end of scope.
Definition: stmt.h:1079
constexpr const char * pragma_scope_prefix
Mark region is guarded by the pragma extension.
Definition: stmt.h:1012
constexpr const char * buffer_dim_align
Mark alignment of buffer dimension stmt.node is Tensor stmt.value is tvm_tuple(dim,...
Definition: stmt.h:1056
constexpr const char * pragma_loop_partition_hint
Mark that the loop should be partitioned.
Definition: stmt.h:1138
constexpr const char * meta_schedule_tiling_structure
Mark the tiling structure of blocks that are applied by rule Multi-Level-Tiling.
Definition: stmt.h:1160
constexpr const char * meta_schedule_random_compute_producer
Mark the block whose producer needs to be applied by rule Random-Compute-Location.
Definition: stmt.h:1177
constexpr const char * meta_schedule_thread_extent_high_inclusive
The allowed range of thread extent in thread bindings.
Definition: stmt.h:1173
constexpr const char * vector_bytes
Mark vectorization length constraint on block.
Definition: stmt.h:1231
constexpr const char * require_block_var_bound_predicate
Mark that the block need to add predicate for block var bounds during lowering.
Definition: stmt.h:1205
@ kUnrolled
The execution is unrolled.
Definition: var.h:210
@ kVectorized
The loop is vectorized.
Definition: var.h:214
PrimExpr TypeAnnotation(PrimType dtype, Span span=Span())
Create a type annotation expression.
std::ostream & operator<<(std::ostream &os, CallEffectKind side_effect)
Definition: op_attr_types.h:140
ForKind
The kind of the loop.
Definition: stmt.h:556
@ kThreadBinding
The loop variable is bound to a thread in an environment. In the final stage of lowering,...
@ kParallel
Parallel execution on CPU.
@ kSerial
default semantics – serial execution.
const char * ForKind2String(ForKind t)
Definition: stmt.h:1291
An object that builds and maintains block scope and StmtSref mapping for Dependence analysis.
Definition: analyzer.h:40