stmt.h File Reference

S-TIR (Schedulable TIR) statements and attributes. More...

#include <tvm/ffi/container/tuple.h>
#include <tvm/s_tir/iter_var.h>
#include <tvm/tirx/stmt.h>

Include dependency graph for stmt.h:

This graph shows which files directly or indirectly include this file:

Go to the source code of this file.

Classes
class	tvm::s_tir::MatchTensorRegionNode
	Match introduces a constraint that the source tensor region can be remapped to the data layout specified by the tensor field. The constraint can be checked in later part of lowering (or optionally during runtime). More...
class	tvm::s_tir::MatchTensorRegion
	Managed reference to MatchTensorRegionNode. More...
class	tvm::s_tir::SBlockNode
	A block is a basic schedule unit in TIR. More...
class	tvm::s_tir::SBlock
	Managed reference to SBlockNode. More...
class	tvm::s_tir::SBlockRealizeNode
	A block realization node represents execution of the block at the binding values. More...
class	tvm::s_tir::SBlockRealize
	Managed reference to BlockRealizeNode. More...

Namespaces
namespace	tvm
	An object that builds and maintains block scope and StmtSref mapping for Dependence analysis.
namespace	tvm::s_tir
namespace	tvm::s_tir::attr

Typedefs
using	tvm::s_tir::TensorAllocatedAddresses = ffi::Array< ffi::Tuple< Var, ffi::Array< PrimExpr > > >
	Placement entries owned by a block, traversable with their tensor references.

Functions
const Op &	tvm::s_tir::async_copy_scope ()
	Region marking copies eligible for asynchronous lowering.
const Op &	tvm::s_tir::async_commit ()
	Commit the preceding asynchronous copies to a queue.
const Op &	tvm::s_tir::async_wait ()
	Wait until at most the given number of committed groups remain in flight.
const Op &	tvm::s_tir::manual_sync ()
	Region whose author manages cross-thread synchronization explicitly.
ffi::Optional< ffi::String >	tvm::s_tir::GetThreadBinding (const ForNode *loop)
	Interpret and validate S-TIR thread placement on a shared loop.
ffi::Optional< ffi::String >	tvm::s_tir::GetThreadBinding (const For &loop)

Variables
constexpr const char *	tvm::s_tir::attr::kThreadBinding = "thread_binding"
	Hardware thread axis bound to a parallel loop during S-TIR scheduling.
constexpr const char *	tvm::s_tir::attr::kDoubleBufferScope = "double_buffer_scope"
	SBlock annotation selecting a write-tensor index for double buffering in InjectSoftwarePipeline.
constexpr const char *	tvm::s_tir::attr::kFragmentShape = "fragment_shape"
	String-valued allocation Call attribute containing the TensorCore fragment shape.
constexpr const char *	tvm::s_tir::attr::kFragmentLayout = "fragment_layout"
	String-valued allocation Call attribute containing the TensorCore fragment layout.
constexpr const char *	tvm::s_tir::attr::kLoopPartitionHint = "loop_partition_hint"
	For annotation requesting partitioning when its PrimExpr value is provably true.
constexpr const char *	tvm::s_tir::attr::kMetaScheduleTilingStructure = "meta_schedule.tiling_structure"
	Mark the tiling structure of blocks that are applied by rule Multi-Level-Tiling.
constexpr const char *	tvm::s_tir::attr::kMetaScheduleCooperativeFetch = "meta_schedule.cooperative_fetch"
	Mark that the loop should be further skip and bound to environment threads to enable cooperative fetching.
constexpr const char *	tvm::s_tir::attr::kMetaScheduleThreadExtentLowInclusive
	The allowed range of thread extent in thread bindings.
constexpr const char *	tvm::s_tir::attr::kMetaScheduleThreadExtentHighInclusive
	The allowed range of thread extent in thread bindings.
constexpr const char *	tvm::s_tir::attr::kMetaScheduleRandomComputeProducer = "meta_schedule.random_compute_producer"
	Mark the block whose producer needs to be applied by rule Random-Compute-Location.
constexpr const char *	tvm::s_tir::attr::kMetaScheduleParallel = "meta_schedule.parallel"
	Mark auto-parallel setting on the block.
constexpr const char *	tvm::s_tir::attr::kMetaScheduleVectorize = "meta_schedule.vectorize"
	Mark auto-vectorize setting on the block.
constexpr const char *	tvm::s_tir::attr::kMetaScheduleUnrollExplicit = "meta_schedule.unroll_explicit"
	Mark auto-unroll setting on the block.
constexpr const char *	tvm::s_tir::attr::kMetaScheduleUnrollImplicit = "meta_schedule.unroll_implicit"
	Mark auto-unroll setting on the block.
constexpr const char *	tvm::s_tir::attr::kMetaScheduleAutoTensorize = "meta_schedule.auto_tensorize"
	Mark that a block should be further rewritten using tensorization.
constexpr const char *	tvm::s_tir::attr::kMetaScheduleLayoutRewritePreproc = "meta_schedule.layout_rewrite_preproc"
	Mark that a block is a preprocessor block for layout rewrite.
constexpr const char *	tvm::s_tir::attr::kMetaScheduleAutoTensorizeInit = "meta_schedule.auto_tensorize_init"
	Mark that the init statement of a block should be further rewritten using tensorization.
constexpr const char *	tvm::s_tir::attr::kMetaScheduleInlineRule = "meta_schedule.inline_rule"
	Mark that a block is disallowed in auto inline.
constexpr const char *	tvm::s_tir::attr::kTensorAllocatedAddr = "s_tir.tensor_allocated_addr"
	TensorAllocatedAddresses for block-owned allocations and match tensors.
constexpr const char *	tvm::s_tir::attr::kScriptParsingDetectAccess = "tirx.script_parsing_detect_access"
	Mark whether the script-completer need to fill in missing access region during script parsing.
constexpr const char *	tvm::s_tir::attr::kRequireBlockVarBoundPredicate = "require_bound_predicate"
	Mark that the block need to add predicate for block var bounds during lowering.
constexpr const char *	tvm::s_tir::attr::kSoftwarePipelineStage = "software_pipeline_stage"
	Mark the stage of a statement in the software pipeline.
constexpr const char *	tvm::s_tir::attr::kSoftwarePipelineOrder = "software_pipeline_order"
	Mark the order of a statement in the software pipeline.
constexpr const char *	tvm::s_tir::attr::kSoftwarePipelineAsyncStages = "software_pipeline_async_stages"
	List stages in the software pipeline that should run asynchronously.
constexpr const char *	tvm::s_tir::attr::kManifestSharedMemoryLocalStage = "tirx.manifest_shared_memory_local_stage"
	Mark the local stage for the shared memory access should be added.
constexpr const char *	tvm::s_tir::attr::kTensorDimAlign = "tensor_dim_align"
	Mark alignment of tensor dimension The annotation value is an array of explicit tuples (tensor_index, axis, factor, offset). This requires the stride of an axis to be k * factor + offset.
constexpr const char *	tvm::s_tir::attr::kExplicitReadRegion = "explicit_read_region"
	Mark that a block has an explicitly specified read region. This is used to override the default read region inference in TIR.
constexpr const char *	tvm::s_tir::attr::kExplicitWriteRegion = "explicit_write_region"
	Mark that a block has an explicitly specified write region. This is used to override the default write region inference in TIR.
constexpr const char *	tvm::s_tir::attr::kIrregularLoopMark = "irregular_loop_mark"
	,ark a ForNode represent an irregular loop of non-structural control flow edges.
constexpr const char *	tvm::s_tir::attr::kAutoCopy = "auto_copy"
	Mark auto copy for memhammer.
constexpr const char *	tvm::s_tir::attr::kLocalStage = "local_stage"
	Mark local stage constraint on data copy.
constexpr const char *	tvm::s_tir::attr::kVectorBytes = "vector_bytes"
	Mark vectorization length constraint on block.
constexpr const char *	tvm::s_tir::attr::kWarpExecution = "warp_execution"
	Mark that a block is executed by a warp. This implies the extend of threadIdx.x is warp size.
constexpr const char *	tvm::s_tir::attr::kPermutedLayout = "permuted_layout"
constexpr const char *	tvm::s_tir::attr::kScheduleRule = "schedule_rule"
constexpr const char *	tvm::s_tir::attr::kMetaScheduleWriteCacheLevel = "s_tir.meta_schedule.write_cache_level"

Detailed Description

S-TIR (Schedulable TIR) statements and attributes.

This file contains attribute keys that are specific to the schedulable TIR (S-TIR) layer, including meta_schedule annotations and schedule primitive / SBlock annotations.