tvm
Loading...
Searching...
No Matches
stmt.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 */
27#ifndef TVM_S_TIR_STMT_H_
28#define TVM_S_TIR_STMT_H_
29
30#include <tvm/ffi/container/tuple.h>
31#include <tvm/s_tir/iter_var.h>
32#include <tvm/tirx/stmt.h>
33
34namespace tvm {
35namespace s_tir {
36
38using TensorAllocatedAddresses = ffi::Array<ffi::Tuple<Var, ffi::Array<PrimExpr>>>;
39
49class MatchTensorRegionNode : public ffi::Object {
50 public:
51 explicit MatchTensorRegionNode(ffi::UnsafeInit tag) : tensor(tag), source(tag) {}
52
55
60
61 static void RegisterReflection() {
62 namespace refl = tvm::ffi::reflection;
63 refl::ObjectDef<MatchTensorRegionNode>()
64 .def_ro("tensor", &MatchTensorRegionNode::tensor,
65 refl::AttachFieldFlag::SEqHashDefPattern())
66 .def_ro("source", &MatchTensorRegionNode::source);
67 }
68
69 static constexpr TVMFFISEqHashKind _type_s_eq_hash_kind = kTVMFFISEqHashKindTreeNode;
70 TVM_FFI_DECLARE_OBJECT_INFO_FINAL("s_tir.MatchTensorRegion", MatchTensorRegionNode, ffi::Object);
71};
72
77class MatchTensorRegion : public ffi::ObjectRef {
78 public:
79 explicit MatchTensorRegion(ffi::UnsafeInit tag, ffi::ObjectPtr<MatchTensorRegionNode> data)
80 : ffi::ObjectRef(tag) {
81 data_ = std::move(data);
82 }
83
84 TVM_DLL explicit MatchTensorRegion(tirx::TensorVar tensor, TensorRegion source);
85
89};
90
112class SBlockNode : public StmtNode {
113 public:
114 explicit SBlockNode(ffi::UnsafeInit tag) : body(tag) {}
115
116 explicit SBlockNode(SeqStmt body) : body(std::move(body)) {}
117
119 ffi::Array<s_tir::IterVar> iter_vars;
121 ffi::Array<TensorRegion> reads;
123 ffi::Array<TensorRegion> writes;
125 ffi::String name_hint;
127 ffi::Array<tirx::TensorVar> alloc_tensors;
129 ffi::Array<MatchTensorRegion> match_tensors;
131 ffi::Map<ffi::String, ffi::Any> annotations;
139 ffi::Optional<SeqStmt> init;
142
143 static void RegisterReflection() {
144 namespace refl = tvm::ffi::reflection;
145 refl::ObjectDef<SBlockNode>()
146 .def_ro("iter_vars", &SBlockNode::iter_vars)
147 .def_ro("reads", &SBlockNode::reads)
148 .def_ro("writes", &SBlockNode::writes)
149 .def_ro("name_hint", &SBlockNode::name_hint, refl::AttachFieldFlag::SEqHashIgnore())
150 .def_ro("alloc_tensors", &SBlockNode::alloc_tensors,
151 refl::AttachFieldFlag::SEqHashDefSimple())
152 .def_ro("match_tensors", &SBlockNode::match_tensors)
153 .def_ro("annotations", &SBlockNode::annotations)
154 .def_ro("init", &SBlockNode::init)
155 .def_ro("body", &SBlockNode::body);
156 }
158};
159
164class SBlock : public Stmt {
165 public:
166 TVM_DLL explicit SBlock(
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>(),
173 Location loc = UnknownLoc());
174
175 TVM_DLL explicit SBlock(ffi::String name_hint, SeqStmt body,
176 ffi::Array<tirx::TensorVar> alloc_tensors = ffi::Array<tirx::TensorVar>(),
177 Location loc = UnknownLoc());
178
179 explicit SBlock(ffi::ObjectPtr<SBlockNode> node) : Stmt(std::move(node)) {}
180
183};
184
189 public:
190 explicit SBlockRealizeNode(ffi::UnsafeInit tag) : predicate(tag), block(tag) {}
191
194
196 ffi::Array<PrimExpr> iter_values;
204
205 static void RegisterReflection() {
206 namespace refl = tvm::ffi::reflection;
207 refl::ObjectDef<SBlockRealizeNode>()
208 .def_ro("iter_values", &SBlockRealizeNode::iter_values)
209 .def_ro("predicate", &SBlockRealizeNode::predicate)
210 .def_ro("block", &SBlockRealizeNode::block);
211 }
213};
214
219class SBlockRealize : public Stmt {
220 public:
221 TVM_DLL explicit SBlockRealize(ffi::Array<PrimExpr> iter_values, PrimExpr predicate, SBlock block,
222 Location loc = UnknownLoc());
223
224 explicit SBlockRealize(ffi::ObjectPtr<SBlockRealizeNode> node) : Stmt(std::move(node)) {}
225
228};
229
231TVM_DLL const Op& async_copy_scope();
233TVM_DLL const Op& async_commit();
235TVM_DLL const Op& async_wait();
237TVM_DLL const Op& manual_sync();
238
239namespace attr {
241constexpr const char* kThreadBinding = "thread_binding";
242
247constexpr const char* kDoubleBufferScope = "double_buffer_scope";
248
252constexpr const char* kFragmentShape = "fragment_shape";
253
257constexpr const char* kFragmentLayout = "fragment_layout";
258
262constexpr const char* kLoopPartitionHint = "loop_partition_hint";
263
264// -----------------------------------------------------------------------
265// meta_schedule annotations
266// -----------------------------------------------------------------------
267
269constexpr const char* kMetaScheduleTilingStructure = "meta_schedule.tiling_structure";
270
275constexpr const char* kMetaScheduleCooperativeFetch = "meta_schedule.cooperative_fetch";
276
279 "meta_schedule.thread_extent_low_inclusive";
280
283 "meta_schedule.thread_extent_high_inclusive";
284
286constexpr const char* kMetaScheduleRandomComputeProducer = "meta_schedule.random_compute_producer";
287
289constexpr const char* kMetaScheduleParallel = "meta_schedule.parallel";
290
292constexpr const char* kMetaScheduleVectorize = "meta_schedule.vectorize";
293
295constexpr const char* kMetaScheduleUnrollExplicit = "meta_schedule.unroll_explicit";
296
298constexpr const char* kMetaScheduleUnrollImplicit = "meta_schedule.unroll_implicit";
299
301constexpr const char* kMetaScheduleAutoTensorize = "meta_schedule.auto_tensorize";
302
304constexpr const char* kMetaScheduleLayoutRewritePreproc = "meta_schedule.layout_rewrite_preproc";
305
309constexpr const char* kMetaScheduleAutoTensorizeInit = "meta_schedule.auto_tensorize_init";
310
312constexpr const char* kMetaScheduleInlineRule = "meta_schedule.inline_rule";
313
314// -----------------------------------------------------------------------
315// Schedule primitive / SBlock annotations
316// -----------------------------------------------------------------------
317
319constexpr const char* kTensorAllocatedAddr = "s_tir.tensor_allocated_addr";
320
328constexpr const char* kScriptParsingDetectAccess = "tirx.script_parsing_detect_access";
329
333constexpr const char* kRequireBlockVarBoundPredicate = "require_bound_predicate";
334
336constexpr const char* kSoftwarePipelineStage = "software_pipeline_stage";
337
339constexpr const char* kSoftwarePipelineOrder = "software_pipeline_order";
340
345constexpr const char* kSoftwarePipelineAsyncStages = "software_pipeline_async_stages";
346
348constexpr const char* kManifestSharedMemoryLocalStage = "tirx.manifest_shared_memory_local_stage";
349
356constexpr const char* kTensorDimAlign = "tensor_dim_align";
357
361constexpr const char* kExplicitReadRegion = "explicit_read_region";
362
366constexpr const char* kExplicitWriteRegion = "explicit_write_region";
367
369constexpr const char* kIrregularLoopMark = "irregular_loop_mark";
370
372constexpr const char* kAutoCopy = "auto_copy";
373
375constexpr const char* kLocalStage = "local_stage";
376
378constexpr const char* kVectorBytes = "vector_bytes";
379
384constexpr const char* kWarpExecution = "warp_execution";
385
386constexpr const char* kPermutedLayout = "permuted_layout";
387
388constexpr const char* kScheduleRule = "schedule_rule";
389
390constexpr const char* kMetaScheduleWriteCacheLevel = "s_tir.meta_schedule.write_cache_level";
391
392} // namespace attr
393
395inline ffi::Optional<ffi::String> GetThreadBinding(const ForNode* loop) {
396 if (auto tag = loop->annotations.Get(tvm::s_tir::attr::kThreadBinding)) {
397 TVM_FFI_CHECK(loop->kind == ForKind::kParallel, ValueError)
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>();
403 }
404 return std::nullopt;
405}
406inline ffi::Optional<ffi::String> GetThreadBinding(const For& loop) {
407 return GetThreadBinding(loop.get());
408}
409
410} // namespace s_tir
411} // namespace tvm
412#endif // TVM_S_TIR_STMT_H_
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.