stmt.h File Reference

S-TIR (Schedulable TIR) statement attribute declarations. More...

#include <tvm/ir/prim/expr.h>

Include dependency graph for stmt.h:

Go to the source code of this file.

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

Variables
constexpr const char *	tvm::s_tir::attr::async_commit_queue_scope = "async_commit_queue_scope"
	Annotations for invoking and synchronizing asynchronous operations.
constexpr const char *	tvm::s_tir::attr::async_wait_queue_scope = "async_wait_queue_scope"
constexpr const char *	tvm::s_tir::attr::async_wait_inflight_count = "async_wait_inflight_count"
constexpr const char *	tvm::s_tir::attr::async_scope = "async_scope"
	Mark that the attached statement runs asynchronously.
constexpr const char *	tvm::s_tir::attr::buffer_bound = "buffer_bound"
	Mark stores/loads with their bounds.
constexpr const char *	tvm::s_tir::attr::double_buffer_scope = "double_buffer_scope"
	Marks production of double buffer data.
constexpr const char *	tvm::s_tir::attr::double_buffer_write = "double_buffer_write"
	Marks region used by double buffer write.
constexpr const char *	tvm::s_tir::attr::fragment_shape = "fragment_shape"
	Mark that the shape of TensorCore fragment.
constexpr const char *	tvm::s_tir::attr::fragment_layout = "fragment_layout"
	Mark that the layout of TensorCore fragment.
constexpr const char *	tvm::s_tir::attr::pragma_loop_partition_hint = "pragma_loop_partition_hint"
	Mark that the loop should be partitioned.
constexpr const char *	tvm::s_tir::attr::reduce_scope = "reduce_scope"
	Mark of reduce scope.
constexpr const char *	tvm::s_tir::attr::virtual_thread = "virtual_thread"
	Mark launching of a virtual thread.
constexpr const char *	tvm::s_tir::attr::meta_schedule_tiling_structure = "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::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.
constexpr const char *	tvm::s_tir::attr::meta_schedule_thread_extent_low_inclusive
	The allowed range of thread extent in thread bindings.
constexpr const char *	tvm::s_tir::attr::meta_schedule_thread_extent_high_inclusive
	The allowed range of thread extent in thread bindings.
constexpr const char *	tvm::s_tir::attr::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::meta_schedule_parallel = "meta_schedule.parallel"
	Mark auto-parallel setting on the block.
constexpr const char *	tvm::s_tir::attr::meta_schedule_vectorize = "meta_schedule.vectorize"
	Mark auto-vectorize setting on the block.
constexpr const char *	tvm::s_tir::attr::meta_schedule_unroll_explicit = "meta_schedule.unroll_explicit"
	Mark auto-unroll setting on the block.
constexpr const char *	tvm::s_tir::attr::meta_schedule_unroll_implicit = "meta_schedule.unroll_implicit"
	Mark auto-unroll setting on the block.
constexpr const char *	tvm::s_tir::attr::meta_schedule_auto_tensorize = "meta_schedule.auto_tensorize"
	Mark that a block should be further rewritten using tensorization.
constexpr const char *	tvm::s_tir::attr::meta_schedule_layout_rewrite_preproc = "meta_schedule.layout_rewrite_preproc"
	Mark that a block is a preprocessor block for layout rewrite.
constexpr const char *	tvm::s_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.
constexpr const char *	tvm::s_tir::attr::meta_schedule_inline_rule = "meta_schedule.inline_rule"
	Mark that a block is disallowed in auto inline.
constexpr const char *	tvm::s_tir::attr::script_parsing_detect_access = "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::require_block_var_bound_predicate = "require_bound_predicate"
	Mark that the block need to add predicate for block var bounds during lowering.
constexpr const char *	tvm::s_tir::attr::software_pipeline_stage = "software_pipeline_stage"
	Mark the stage of a statement in the software pipeline.
constexpr const char *	tvm::s_tir::attr::software_pipeline_order = "software_pipeline_order"
	Mark the order of a statement in the software pipeline.
constexpr const char *	tvm::s_tir::attr::software_pipeline_async_stages = "software_pipeline_async_stages"
	List stages in the software pipeline that should run asynchronously.
constexpr const char *	tvm::s_tir::attr::layout_free_buffers = "layout_free_buffers"
	Mark the buffers which is const access and can be transformed layout.
constexpr const char *	tvm::s_tir::attr::manifest_shared_memory_local_stage
	Mark the local stage for the shared memory access should be added.
constexpr const char *	tvm::s_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.
constexpr const char *	tvm::s_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.
constexpr const char *	tvm::s_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.
constexpr const char *	tvm::s_tir::attr::irregular_loop_mark = "irregular_loop_mark"
	,ark a ForNode represent an irregular loop of non-structural control flow edges.
constexpr const char *	tvm::s_tir::attr::auto_copy = "auto_copy"
	Mark auto copy for memhammer.
constexpr const char *	tvm::s_tir::attr::local_stage = "local_stage"
	Mark local stage constraint on data copy.
constexpr const char *	tvm::s_tir::attr::vector_bytes = "vector_bytes"
	Mark vectorization length constraint on block.
constexpr const char *	tvm::s_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.
constexpr const char *	tvm::s_tir::attr::layout_transforms = "layout_transforms"
	Marks the layout transforms to be used for a tensor.
constexpr const char *	tvm::s_tir::attr::hand_threaded = "hand_threaded"
	Mark that the kernel is hand threaded and doesn't need syncs inserted.

Detailed Description

S-TIR (Schedulable TIR) statement attribute declarations.

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