27#ifndef TVM_S_TIR_STMT_H_
28#define TVM_S_TIR_STMT_H_
30#include <tvm/ffi/container/tuple.h>
62 namespace refl = tvm::ffi::reflection;
63 refl::ObjectDef<MatchTensorRegionNode>()
65 refl::AttachFieldFlag::SEqHashDefPattern())
80 : ffi::ObjectRef(tag) {
81 data_ = std::move(data);
144 namespace refl = tvm::ffi::reflection;
145 refl::ObjectDef<SBlockNode>()
151 refl::AttachFieldFlag::SEqHashDefSimple())
167 ffi::Array<s_tir::IterVar> iter_vars, ffi::Array<TensorRegion> reads,
168 ffi::Array<TensorRegion> writes, ffi::String name_hint,
SeqStmt body,
169 ffi::Optional<SeqStmt> init = std::nullopt,
170 ffi::Array<tirx::TensorVar> alloc_tensors = ffi::Array<tirx::TensorVar>(),
171 ffi::Array<MatchTensorRegion> match_tensors = ffi::Array<MatchTensorRegion>(),
172 ffi::Map<ffi::String, ffi::Any> annotations = ffi::Map<ffi::String, ffi::Any>(),
176 ffi::Array<tirx::TensorVar> alloc_tensors = ffi::Array<tirx::TensorVar>(),
179 explicit SBlock(ffi::ObjectPtr<SBlockNode> node) :
Stmt(std::move(node)) {}
206 namespace refl = tvm::ffi::reflection;
207 refl::ObjectDef<SBlockRealizeNode>()
279 "meta_schedule.thread_extent_low_inclusive";
283 "meta_schedule.thread_extent_high_inclusive";
398 <<
"thread_binding requires a parallel loop";
399 TVM_FFI_CHECK(tag->as<ffi::String>().has_value(), TypeError)
400 <<
"thread_binding annotation must be a string";
401 TVM_FFI_CHECK(loop->
HasTrivialStep(), ValueError) <<
"Thread binding loops require a unit step";
402 return tag->cast<ffi::String>();
A for loop, with possible type annotations.
Definition stmt.h:457
ForKind kind
The kind of the for loop.
Definition stmt.h:474
ffi::Map< ffi::String, ffi::Any > annotations
Additional annotations about the loop.
Definition stmt.h:483
bool HasTrivialStep() const
Check it is a loop without nontrivial loop step.
Managed reference to ForNode.
Definition stmt.h:511
Non-nullable source-location metadata.
Definition location.h:91
Reference-counted handle to a canonical named operator.
Definition op.h:147
Typed reference/view over any Expr whose ExprNode::ty is PrimType.
Definition base_expr.h:447
Sequence statement.
Definition stmt.h:92
Base node of all statements.
Definition stmt.h:37
Container of all statements.
Definition stmt.h:62
Managed reference to TensorRegionNode.
Definition expr.h:1079
Reference to the shared immutable unknown location.
Definition location.h:105
Match introduces a constraint that the source tensor region can be remapped to the data layout specif...
Definition stmt.h:49
MatchTensorRegionNode(ffi::UnsafeInit tag)
Definition stmt.h:51
MatchTensorRegionNode(tirx::TensorVar tensor, TensorRegion source)
Definition stmt.h:53
tirx::TensorVar tensor
The target tensor.
Definition stmt.h:57
static constexpr TVMFFISEqHashKind _type_s_eq_hash_kind
Definition stmt.h:69
static void RegisterReflection()
Definition stmt.h:61
TensorRegion source
The source tensor region.
Definition stmt.h:59
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("s_tir.MatchTensorRegion", MatchTensorRegionNode, ffi::Object)
Managed reference to MatchTensorRegionNode.
Definition stmt.h:77
MatchTensorRegion(tirx::TensorVar tensor, TensorRegion source)
MatchTensorRegion(ffi::UnsafeInit tag, ffi::ObjectPtr< MatchTensorRegionNode > data)
Definition stmt.h:79
TVM_FFI_DEFINE_OBJECT_REF_METHODS_NOTNULLABLE(MatchTensorRegion, ffi::ObjectRef, MatchTensorRegionNode)
TVM_DEFINE_OBJECT_REF_COW_METHOD(MatchTensorRegionNode)
A block is a basic schedule unit in TIR.
Definition stmt.h:112
ffi::Array< TensorRegion > writes
The write tensor regions of the block.
Definition stmt.h:123
ffi::Array< TensorRegion > reads
The read tensor regions of the block.
Definition stmt.h:121
ffi::Array< tirx::TensorVar > alloc_tensors
The tensor allocated in the block.
Definition stmt.h:127
ffi::Array< s_tir::IterVar > iter_vars
The variables of the block.
Definition stmt.h:119
SBlockNode(ffi::UnsafeInit tag)
Definition stmt.h:114
ffi::String name_hint
The name_hint of the block.
Definition stmt.h:125
ffi::Array< MatchTensorRegion > match_tensors
The match tensor regions.
Definition stmt.h:129
static void RegisterReflection()
Definition stmt.h:143
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("s_tir.SBlock", SBlockNode, StmtNode)
SeqStmt body
The body of the block.
Definition stmt.h:141
ffi::Optional< SeqStmt > init
The init statement is executed during the first iteration of reduction loops in a reduction block....
Definition stmt.h:139
ffi::Map< ffi::String, ffi::Any > annotations
The annotation of the block.
Definition stmt.h:131
SBlockNode(SeqStmt body)
Definition stmt.h:116
A block realization node represents execution of the block at the binding values.
Definition stmt.h:188
PrimExpr predicate
The predicate of the block realization, the block will only be executed when the predicate is true.
Definition stmt.h:201
TVM_FFI_DECLARE_OBJECT_INFO_FINAL("s_tir.SBlockRealize", SBlockRealizeNode, StmtNode)
SBlock block
The block to be realized.
Definition stmt.h:203
SBlockRealizeNode(ffi::UnsafeInit tag)
Definition stmt.h:190
SBlockRealizeNode(PrimExpr predicate, SBlock block)
Definition stmt.h:192
static void RegisterReflection()
Definition stmt.h:205
ffi::Array< PrimExpr > iter_values
The corresponding values of the iter vars.
Definition stmt.h:196
Managed reference to BlockRealizeNode.
Definition stmt.h:219
SBlockRealize(ffi::ObjectPtr< SBlockRealizeNode > node)
Definition stmt.h:224
SBlockRealize(ffi::Array< PrimExpr > iter_values, PrimExpr predicate, SBlock block, Location loc=UnknownLoc())
TVM_DEFINE_OBJECT_REF_COW_METHOD(SBlockRealizeNode)
TVM_FFI_DEFINE_OBJECT_REF_METHODS_NOTNULLABLE(SBlockRealize, Stmt, SBlockRealizeNode)
Managed reference to SBlockNode.
Definition stmt.h:164
SBlock(ffi::ObjectPtr< SBlockNode > node)
Definition stmt.h:179
TVM_DEFINE_OBJECT_REF_COW_METHOD(SBlockNode)
SBlock(ffi::Array< s_tir::IterVar > iter_vars, ffi::Array< TensorRegion > reads, ffi::Array< TensorRegion > writes, ffi::String name_hint, SeqStmt body, ffi::Optional< SeqStmt > init=std::nullopt, ffi::Array< tirx::TensorVar > alloc_tensors=ffi::Array< tirx::TensorVar >(), ffi::Array< MatchTensorRegion > match_tensors=ffi::Array< MatchTensorRegion >(), ffi::Map< ffi::String, ffi::Any > annotations=ffi::Map< ffi::String, ffi::Any >(), Location loc=UnknownLoc())
TVM_FFI_DEFINE_OBJECT_REF_METHODS_NOTNULLABLE(SBlock, Stmt, SBlockNode)
SBlock(ffi::String name_hint, SeqStmt body, ffi::Array< tirx::TensorVar > alloc_tensors=ffi::Array< tirx::TensorVar >(), Location loc=UnknownLoc())
Checked zero-state view over an ordinary VarNode with TensorType.
Definition expr.h:50
Axis metadata for schedulable tensor computations.
constexpr const char * kExplicitReadRegion
Mark that a block has an explicitly specified read region. This is used to override the default read ...
Definition stmt.h:361
constexpr const char * kMetaScheduleTilingStructure
Mark the tiling structure of blocks that are applied by rule Multi-Level-Tiling.
Definition stmt.h:269
constexpr const char * kThreadBinding
Hardware thread axis bound to a parallel loop during S-TIR scheduling.
Definition stmt.h:241
constexpr const char * kMetaScheduleParallel
Mark auto-parallel setting on the block.
Definition stmt.h:289
constexpr const char * kAutoCopy
Mark auto copy for memhammer.
Definition stmt.h:372
constexpr const char * kIrregularLoopMark
,ark a ForNode represent an irregular loop of non-structural control flow edges.
Definition stmt.h:369
constexpr const char * kTensorAllocatedAddr
TensorAllocatedAddresses for block-owned allocations and match tensors.
Definition stmt.h:319
constexpr const char * kSoftwarePipelineAsyncStages
List stages in the software pipeline that should run asynchronously.
Definition stmt.h:345
constexpr const char * kMetaScheduleUnrollImplicit
Mark auto-unroll setting on the block.
Definition stmt.h:298
constexpr const char * kMetaScheduleAutoTensorize
Mark that a block should be further rewritten using tensorization.
Definition stmt.h:301
constexpr const char * kMetaScheduleWriteCacheLevel
Definition stmt.h:390
constexpr const char * kMetaScheduleVectorize
Mark auto-vectorize setting on the block.
Definition stmt.h:292
constexpr const char * kScriptParsingDetectAccess
Mark whether the script-completer need to fill in missing access region during script parsing.
Definition stmt.h:328
constexpr const char * kMetaScheduleInlineRule
Mark that a block is disallowed in auto inline.
Definition stmt.h:312
constexpr const char * kLocalStage
Mark local stage constraint on data copy.
Definition stmt.h:375
constexpr const char * kWarpExecution
Mark that a block is executed by a warp. This implies the extend of threadIdx.x is warp size.
Definition stmt.h:384
constexpr const char * kMetaScheduleCooperativeFetch
Mark that the loop should be further skip and bound to environment threads to enable cooperative fetc...
Definition stmt.h:275
constexpr const char * kVectorBytes
Mark vectorization length constraint on block.
Definition stmt.h:378
constexpr const char * kLoopPartitionHint
For annotation requesting partitioning when its PrimExpr value is provably true.
Definition stmt.h:262
constexpr const char * kMetaScheduleRandomComputeProducer
Mark the block whose producer needs to be applied by rule Random-Compute-Location.
Definition stmt.h:286
constexpr const char * kTensorDimAlign
Mark alignment of tensor dimension The annotation value is an array of explicit tuples (tensor_index,...
Definition stmt.h:356
constexpr const char * kMetaScheduleAutoTensorizeInit
Mark that the init statement of a block should be further rewritten using tensorization.
Definition stmt.h:309
constexpr const char * kMetaScheduleThreadExtentHighInclusive
The allowed range of thread extent in thread bindings.
Definition stmt.h:282
constexpr const char * kMetaScheduleUnrollExplicit
Mark auto-unroll setting on the block.
Definition stmt.h:295
constexpr const char * kMetaScheduleThreadExtentLowInclusive
The allowed range of thread extent in thread bindings.
Definition stmt.h:278
constexpr const char * kManifestSharedMemoryLocalStage
Mark the local stage for the shared memory access should be added.
Definition stmt.h:348
constexpr const char * kFragmentLayout
String-valued allocation Call attribute containing the TensorCore fragment layout.
Definition stmt.h:257
constexpr const char * kFragmentShape
String-valued allocation Call attribute containing the TensorCore fragment shape.
Definition stmt.h:252
constexpr const char * kSoftwarePipelineOrder
Mark the order of a statement in the software pipeline.
Definition stmt.h:339
constexpr const char * kMetaScheduleLayoutRewritePreproc
Mark that a block is a preprocessor block for layout rewrite.
Definition stmt.h:304
constexpr const char * kRequireBlockVarBoundPredicate
Mark that the block need to add predicate for block var bounds during lowering.
Definition stmt.h:333
constexpr const char * kExplicitWriteRegion
Mark that a block has an explicitly specified write region. This is used to override the default writ...
Definition stmt.h:366
constexpr const char * kDoubleBufferScope
SBlock annotation selecting a write-tensor index for double buffering in InjectSoftwarePipeline.
Definition stmt.h:247
constexpr const char * kPermutedLayout
Definition stmt.h:386
constexpr const char * kSoftwarePipelineStage
Mark the stage of a statement in the software pipeline.
Definition stmt.h:336
constexpr const char * kScheduleRule
Definition stmt.h:388
const Op & async_commit()
Commit the preceding asynchronous copies to a queue.
const Op & async_copy_scope()
Region marking copies eligible for asynchronous lowering.
ffi::Optional< ffi::String > GetThreadBinding(const ForNode *loop)
Interpret and validate S-TIR thread placement on a shared loop.
Definition stmt.h:395
const Op & manual_sync()
Region whose author manages cross-thread synchronization explicitly.
const Op & async_wait()
Wait until at most the given number of committed groups remain in flight.
ffi::Array< ffi::Tuple< Var, ffi::Array< PrimExpr > > > TensorAllocatedAddresses
Placement entries owned by a block, traversable with their tensor references.
Definition stmt.h:38
An object that builds and maintains block scope and StmtSref mapping for Dependence analysis.
Definition attr.h:23
@ kParallel
Parallel execution.
TIRx statements and shared statement entry header.