stmt.h File Reference

TIR statements. More...

#include <tvm/tir/expr.h>
#include <string>
#include <type_traits>
#include <utility>
#include <vector>

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::tir::StmtNode
	Base node of all statements. More...
class	tvm::tir::Stmt
	Container of all statements. More...
class	tvm::tir::LetStmtNode
	Let binding, bind var to value, then run body. More...
class	tvm::tir::LetStmt
	Managed reference to LetStmtNode. More...
class	tvm::tir::AttrStmtNode
	Define certain auxiliary attribute for the body to be a symbolic value. This provide auxiliary information for IR passes that transforms body. More...
class	tvm::tir::AttrStmt
	Managed reference to AttrStmtNode. More...
class	tvm::tir::AssertStmtNode
	Assert condition, if an error occurs, return the error message. More...
class	tvm::tir::AssertStmt
	Managed reference to AssertStmtNode. More...
class	tvm::tir::BufferStoreNode
	Store value to the high dimension buffer. More...
class	tvm::tir::BufferStore
	Managed reference to BufferStoreNode. More...
class	tvm::tir::BufferRealizeNode
	Annotate the region where the buffer need to be read and write in the body. We only need to allocate the space for the corresponding region. More...
class	tvm::tir::BufferRealize
	Managed reference to BufferRealizeNode. More...
class	tvm::tir::ProducerStoreNode
	Store value into mult-dimensional array that will be read by the consumer of the producer. More...
class	tvm::tir::ProducerStore
	Managed reference to ProducerStoreNode. More...
class	tvm::tir::ProducerRealizeNode
	Annotate the bounds where the data produced by the producer need to be written and read in body. We will need to allocate space for the corresponding regions. More...
class	tvm::tir::ProducerRealize
	Managed reference to ProducerRealizeNode. More...
class	tvm::tir::AllocateNode
	Allocate a buffer that can be used in body. More...
class	tvm::tir::Allocate
	Managed reference to AllocateNode. More...
class	tvm::tir::AllocateConstNode
	Allocate a buffer that can be used in body. More...
class	tvm::tir::AllocateConst
	Managed reference to AllocateConstNode. More...
class	tvm::tir::DeclBufferNode
	Declare a buffer that can be used in the body. More...
class	tvm::tir::DeclBuffer
	Managed reference to DeclBufferNode. More...
class	tvm::tir::SeqStmtNode
	The container of seq statement. Represent a sequence of statements. More...
class	tvm::tir::EvaluateNode
	Evaluates an expression. This is mostly used for putting a Call node into Stmt. More...
class	tvm::tir::Evaluate
	Managed reference to EvaluateNode. More...
class	tvm::tir::SeqStmt
	Sequence statement. More...
class	tvm::tir::SeqStmt::Flattener
	Helper class to flatten sequence of arguments into Array. More...
class	tvm::tir::IfThenElseNode
	IfThenElse statement. More...
class	tvm::tir::IfThenElse
	Managed reference to IfThenElseNode. More...
class	tvm::tir::ForNode
	A for loop, with possible type annotations. More...
class	tvm::tir::For
	Managed reference to ForNode. More...
class	tvm::tir::WhileNode
	A While loop. More...
class	tvm::tir::While
	Managed reference to WhileNode. More...
class	tvm::tir::PrefetchNode
	A prefetch hint for a buffer. More...
class	tvm::tir::Prefetch
	Managed reference to PrefetchNode. More...
class	tvm::tir::BufferRegionNode
	Representing the region of multi-dimensional buffer access. More...
class	tvm::tir::BufferRegion
	Managed reference to BufferRegionNode. More...
class	tvm::tir::MatchBufferRegionNode
	Match introduces a constraint that the source buffer region can be remapped to the data layout specified by the buffer field. The constraint can be checked in later part of lowering (or optionally during runtime). More...
class	tvm::tir::MatchBufferRegion
	Managed reference to MatchBufferRegionNode. More...
class	tvm::tir::BlockNode
	A block is a basic schedule unit in TIR. More...
class	tvm::tir::Block
	Managed reference to BlockNode. More...
class	tvm::tir::BlockRealizeNode
	A block realization node represents execution of the block at the binding values. More...
class	tvm::tir::BlockRealize
	Managed reference to BlockRealizeNode. More...

Namespaces
	tvm
	runtime implementation for LibTorch/TorchScript.
	tvm::tir
	tvm::tir::attr
	PrimFunc specific attribute names.

Enumerations
enum class	tvm::tir::ForKind : int {   tvm::tir::kSerial = 0 , tvm::tir::kParallel = 1 , tvm::tir::kVectorized = 2 , tvm::tir::kUnrolled = 3 ,   tvm::tir::kThreadBinding = 4 }
	The kind of the loop. More...

Functions
bool	tvm::tir::attr::IsPragmaKey (const std::string &attr_key)
	Check if attr_key is a pragma key extension. More...
PrimExpr	tvm::tir::TypeAnnotation (DataType dtype, Span span=Span())
	Create a type annotation expression. More...
std::ostream &	tvm::tir::operator<< (std::ostream &os, ForKind kind)
const char *	tvm::tir::ForKind2String (ForKind t)

Variables
constexpr const char *	tvm::tir::attr::thread_extent = "thread_extent"
	Mark launching extent of thread, used by device API. More...
constexpr const char *	tvm::tir::attr::virtual_thread = "virtual_thread"
	Mark launching of a virtual thread. More...
constexpr const char *	tvm::tir::attr::coproc_scope = "coproc_scope"
	Mark region is processed by a co-processor. More...
constexpr const char *	tvm::tir::attr::coproc_uop_scope = "coproc_uop_scope"
	Mark region creates coprocessor micro ops, can be reused if corresponding variable is independent. More...
constexpr const char *	tvm::tir::attr::volatile_scope = "volatile_scope"
	Mark the scope as volatile access for certain handle. More...
constexpr const char *	tvm::tir::attr::extern_scope = "extern_scope"
	Mark the scope as generated by extern primitive. such scope can contain arbitrary ir program and we need to be careful when make certain assumptions about the structure of the program. More...
constexpr const char *	tvm::tir::attr::compute_scope = "compute_scope"
	Mark the scope as when computation start to happen This can hint some code generator to create a new function for compute. More...
constexpr const char *	tvm::tir::attr::storage_alignment = "storage_alignment"
	Mark storage alignment requirement of buffers. More...
constexpr const char *	tvm::tir::attr::realize_scope = "realize_scope"
	Mark storage scope of realization. More...
constexpr const char *	tvm::tir::attr::device_id = "device_id"
	The allocation device for global malloc in host. More...
constexpr const char *	tvm::tir::attr::device_type = "device_type"
	The device type. More...
constexpr const char *	tvm::tir::attr::loop_scope = "loop_scope"
	Mark of loop scope. More...
constexpr const char *	tvm::tir::attr::reduce_scope = "reduce_scope"
	Mark of reduce scope. More...
constexpr const char *	tvm::tir::attr::pragma_auto_unroll_max_step = "pragma_auto_unroll_max_step"
	Pragma: auto-unroll, max_step. More...
constexpr const char *	tvm::tir::attr::pragma_unroll_explicit = "pragma_unroll_explicit"
	Pragma: unroll explicit. More...
constexpr const char *	tvm::tir::attr::pragma_scope_prefix = "pragma_"
	Mark region is guarded by the pragma extension. More...
constexpr const char *	tvm::tir::attr::pragma_import_c = "pragma_import_c"
	Import C source or file into the final code gen module. More...
constexpr const char *	tvm::tir::attr::pragma_import_llvm = "pragma_import_llvm"
	Import llvm source or file into the final code gen module. More...
constexpr const char *	tvm::tir::attr::pragma_tensor_core = "pragma_tensor_core"
	Try to modify the AST to support Tensor Core. More...
constexpr const char *	tvm::tir::attr::prefetch_scope = "prefetch_scope"
	Mark of prefetch scope, value=offset, run prefetch of Tensor on the current loop scope. More...
constexpr const char *	tvm::tir::attr::layout_transforms = "layout_transforms"
	Marks the layout transforms to be used for a tensor. More...
constexpr const char *	tvm::tir::attr::axis_separators = "axis_separators"
	Marks the physical axis separators. More...
constexpr const char *	tvm::tir::attr::double_buffer_scope = "double_buffer_scope"
	Marks production of double buffer data. More...
constexpr const char *	tvm::tir::attr::double_buffer_write = "double_buffer_write"
	Marks region used by double buffer write. More...
constexpr const char *	tvm::tir::attr::rolling_buffer_scope = "rolling_buffer_scope"
	Mark realization for rolling buffer optimization. More...
constexpr const char *	tvm::tir::attr::scan_update_scope = "scan_update_scope"
	Mark of scan update scope. More...
constexpr const char *	tvm::tir::attr::scan_init_scope = "scan_init_scope"
	Mark of scan init scope. More...
constexpr const char *	tvm::tir::attr::buffer_dim_align = "buffer_dim_align"
	Mark alignment of buffer dimension stmt.node is Tensor stmt.value is tvm_tuple(dim, align, offset) This gives hint to require stride of dim to be k * align + offset. More...
constexpr const char *	tvm::tir::attr::buffer_bound = "buffer_bound"
	Mark stores/loads with theirs bounds. More...
constexpr const char *	tvm::tir::attr::buffer_bind_scope = "buffer_bind_scope"
	Bind the buffer specification to the region of the op When this scope occurs, the stmt.node is a Array<NodeRef> = [buffer, tensor] stmt.value is a tvm_tuple(min0, extent0, min1, extent1, ...). The scope represents that we need to bind the storage region of tensor to buffer. This will affect replacement of some variables inside the scope that corresponds to field of buffer to be the actual expressions of tensor during storage flattening phase. More...
constexpr const char *	tvm::tir::attr::channel_read_scope = "channel_read_scope"
	channel read scope More...
constexpr const char *	tvm::tir::attr::channel_read_advance = "channel_read_advance"
	Advance step of channel after end of scope. More...
constexpr const char *	tvm::tir::attr::channel_write_scope = "channel_write_scope"
	channel write scope More...
constexpr const char *	tvm::tir::attr::channel_write_advance = "channel_write_advance"
	Advance step of channel after end of scope. More...
constexpr const char *	tvm::tir::attr::pipeline_stage_scope = "pipeline_stage_scope"
	pipeline stage scope, implies always execution More...
constexpr const char *	tvm::tir::attr::pipeline_exec_scope = "pipeline_exec_scope"
	pipeline execution scope, implies the scope can be pipelined. More...
constexpr const char *	tvm::tir::attr::device_scope = "device_scope"
	Mark that it is in the device scope. More...
constexpr const char *	tvm::tir::attr::async_scope = "async_scope"
	Mark that the attached statement runs asynchronously. More...
constexpr const char *	tvm::tir::attr::async_commit_queue_scope = "async_commit_queue_scope"
	Annotations for invoking and synchronizing asynchronous operations. More...
constexpr const char *	tvm::tir::attr::async_wait_queue_scope = "async_wait_queue_scope"
constexpr const char *	tvm::tir::attr::async_wait_inflight_count = "async_wait_inflight_count"
constexpr const char *	tvm::tir::attr::fragment_shape = "fragment_shape"
	Mark that the shape of TensorCore fragment. More...
constexpr const char *	tvm::tir::attr::fragment_layout = "fragment_layout"
	Mark that the layout of TensorCore fragment. More...
constexpr const char *	tvm::tir::attr::hand_threaded = "hand_threaded"
	Mark that the kernel is hand threaded and doesn't need syncs inserted. More...
constexpr const char *	tvm::tir::attr::script_parsing_detect_access = "tir.script_parsing_detect_access"
	Mark whether the script-completer need to fill in missing access region during script parsing. More...
constexpr const char *	tvm::tir::attr::pragma_loop_partition_hint = "pragma_loop_partition_hint"
	Mark that the loop should be partitioned. More...
constexpr const char *	tvm::tir::attr::software_pipeline_stage = "software_pipeline_stage"
	Mark the stage of a statement in the software pipeline. More...
constexpr const char *	tvm::tir::attr::software_pipeline_order = "software_pipeline_order"
	Mark the order of a statement in the software pipeline. More...
constexpr const char *	tvm::tir::attr::software_pipeline_async_stages = "software_pipeline_async_stages"
	List stages in the software pipeline that should run asynchronously. More...
constexpr const char *	tvm::tir::attr::layout_free_buffers = "layout_free_buffers"
	Mark the buffers which is const access and can be transformed layout. More...
constexpr const char *	tvm::tir::attr::manifest_shared_memory_local_stage = "tir.manifest_shared_memory_local_stage"
	Mark the local stage for the shared memory access should be added. More...
constexpr const char *	tvm::tir::attr::meta_schedule_tiling_structure = "meta_schedule.tiling_structure"
	Mark the tiling structure of blocks that are applied by rule Multi-Level-Tiling. More...
constexpr const char *	tvm::tir::attr::meta_schedule_cooperative_fetch = "meta_schedule.cooperative_fetch"
	Mark that the loop should be further skip and bound to environment threads to enable cooperative fetching. More...
constexpr const char *	tvm::tir::attr::meta_schedule_thread_extent_low_inclusive
	The allowed range of thread extent in thread bindings. More...
constexpr const char *	tvm::tir::attr::meta_schedule_thread_extent_high_inclusive
	The allowed range of thread extent in thread bindings. More...
constexpr const char *	tvm::tir::attr::meta_schedule_random_compute_producer
	Mark the block whose producer needs to be applied by rule Random-Compute-Location. More...
constexpr const char *	tvm::tir::attr::meta_schedule_parallel = "meta_schedule.parallel"
	Mark auto-parallel setting on the block. More...
constexpr const char *	tvm::tir::attr::meta_schedule_vectorize = "meta_schedule.vectorize"
	Mark auto-vectorize setting on the block. More...
constexpr const char *	tvm::tir::attr::meta_schedule_unroll_explicit = "meta_schedule.unroll_explicit"
	Mark auto-unroll setting on the block. More...
constexpr const char *	tvm::tir::attr::meta_schedule_unroll_implicit = "meta_schedule.unroll_implicit"
	Mark auto-unroll setting on the block. More...
constexpr const char *	tvm::tir::attr::meta_schedule_auto_tensorize = "meta_schedule.auto_tensorize"
	Mark that a block should be further rewritten using tensorization. More...
constexpr const char *	tvm::tir::attr::meta_schedule_layout_rewrite_preproc = "meta_schedule.layout_rewrite_preproc"
	Mark that a block is a preprocessor block for layout rewrite. More...
constexpr const char *	tvm::tir::attr::meta_schedule_auto_tensorize_init = "meta_schedule.auto_tensorize_init"
	Mark that the init statement of a block should be further rewritten using tensorization. More...
constexpr const char *	tvm::tir::attr::require_block_var_bound_predicate = "require_bound_predicate"
	Mark that the block need to add predicate for block var bounds during lowering. More...
constexpr const char *	tvm::tir::attr::meta_schedule_tensor_core_enabled = "meta_schedule.tensor_core_enabled"
	Mark that tensor core is enabled in the PrimExpr. More...
constexpr const char *	tvm::tir::attr::meta_schedule_cache_type = "meta_schedule.cache_type"
	Mark a block as generated by cache_read or cache_write block. 0 means cache_read; 1 means cache_write. More...
constexpr const int	tvm::tir::attr::meta_schedule_cache_type_read = 0
constexpr const int	tvm::tir::attr::meta_schedule_cache_type_write = 1
constexpr const char *	tvm::tir::attr::auto_copy = "auto_copy"
	Mark auto copy for memhammer. More...
constexpr const char *	tvm::tir::attr::local_stage = "local_stage"
	Mark local stage constraint on data copy. More...
constexpr const char *	tvm::tir::attr::vector_bytes = "vector_bytes"
	Mark vectorization length constraint on block. More...
constexpr const char *	tvm::tir::attr::warp_execution = "warp_execution"
	Mark that a block is executed by a warp. This implies the extend of threadIdx.x is warp size. More...
constexpr const char *	tvm::tir::attr::meta_schedule_inline_rule = "meta_schedule.inline_rule"
	Mark that a block is disallowed in auto inline. More...
constexpr const char *	tvm::tir::attr::explicit_read_region = "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. More...
constexpr const char *	tvm::tir::attr::explicit_write_region = "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. More...

Detailed Description

TIR statements.