Variables
constexpr const char *	async_commit_queue_scope = "async_commit_queue_scope"
	Annotations for invoking and synchronizing asynchronous operations.

constexpr const char *	async_wait_queue_scope = "async_wait_queue_scope"

constexpr const char *	async_wait_inflight_count = "async_wait_inflight_count"

constexpr const char *	async_scope = "async_scope"
	Mark that the attached statement runs asynchronously.

constexpr const char *	buffer_bound = "buffer_bound"
	Mark stores/loads with their bounds.

constexpr const char *	double_buffer_scope = "double_buffer_scope"
	Marks production of double buffer data.

constexpr const char *	double_buffer_write = "double_buffer_write"
	Marks region used by double buffer write.

constexpr const char *	fragment_shape = "fragment_shape"
	Mark that the shape of TensorCore fragment.

constexpr const char *	fragment_layout = "fragment_layout"
	Mark that the layout of TensorCore fragment.

constexpr const char *	pragma_loop_partition_hint = "pragma_loop_partition_hint"
	Mark that the loop should be partitioned.

constexpr const char *	reduce_scope = "reduce_scope"
	Mark of reduce scope.

constexpr const char *	virtual_thread = "virtual_thread"
	Mark launching of a virtual thread.

constexpr const char *	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 *	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 *	meta_schedule_thread_extent_low_inclusive
	The allowed range of thread extent in thread bindings.

constexpr const char *	meta_schedule_thread_extent_high_inclusive
	The allowed range of thread extent in thread bindings.

constexpr const char *	meta_schedule_random_compute_producer
	Mark the block whose producer needs to be applied by rule Random-Compute-Location.

constexpr const char *	meta_schedule_parallel = "meta_schedule.parallel"
	Mark auto-parallel setting on the block.

constexpr const char *	meta_schedule_vectorize = "meta_schedule.vectorize"
	Mark auto-vectorize setting on the block.

constexpr const char *	meta_schedule_unroll_explicit = "meta_schedule.unroll_explicit"
	Mark auto-unroll setting on the block.

constexpr const char *	meta_schedule_unroll_implicit = "meta_schedule.unroll_implicit"
	Mark auto-unroll setting on the block.

constexpr const char *	meta_schedule_auto_tensorize = "meta_schedule.auto_tensorize"
	Mark that a block should be further rewritten using tensorization.

constexpr const char *	meta_schedule_layout_rewrite_preproc = "meta_schedule.layout_rewrite_preproc"
	Mark that a block is a preprocessor block for layout rewrite.

constexpr const char *	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 *	meta_schedule_inline_rule = "meta_schedule.inline_rule"
	Mark that a block is disallowed in auto inline.

constexpr const char *	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 *	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 *	software_pipeline_stage = "software_pipeline_stage"
	Mark the stage of a statement in the software pipeline.

constexpr const char *	software_pipeline_order = "software_pipeline_order"
	Mark the order of a statement in the software pipeline.

constexpr const char *	software_pipeline_async_stages = "software_pipeline_async_stages"
	List stages in the software pipeline that should run asynchronously.

constexpr const char *	layout_free_buffers = "layout_free_buffers"
	Mark the buffers which is const access and can be transformed layout.

constexpr const char *	manifest_shared_memory_local_stage
	Mark the local stage for the shared memory access should be added.

constexpr const char *	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 *	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 *	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 *	irregular_loop_mark = "irregular_loop_mark"
	,ark a ForNode represent an irregular loop of non-structural control flow edges.

constexpr const char *	auto_copy = "auto_copy"
	Mark auto copy for memhammer.

constexpr const char *	local_stage = "local_stage"
	Mark local stage constraint on data copy.

constexpr const char *	vector_bytes = "vector_bytes"
	Mark vectorization length constraint on block.

constexpr const char *	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 *	layout_transforms = "layout_transforms"
	Marks the layout transforms to be used for a tensor.

constexpr const char *	hand_threaded = "hand_threaded"
	Mark that the kernel is hand threaded and doesn't need syncs inserted.

Variable Documentation

◆ async_commit_queue_scope

constexpr const char* tvm::s_tir::attr::async_commit_queue_scope = "async_commit_queue_scope"

constexpr

Annotations for invoking and synchronizing asynchronous operations.

◆ async_scope

constexpr const char* tvm::s_tir::attr::async_scope = "async_scope"

constexpr

Mark that the attached statement runs asynchronously.

◆ async_wait_inflight_count

constexpr const char* tvm::s_tir::attr::async_wait_inflight_count = "async_wait_inflight_count"

constexpr

◆ async_wait_queue_scope

constexpr const char* tvm::s_tir::attr::async_wait_queue_scope = "async_wait_queue_scope"

constexpr

◆ auto_copy

constexpr const char* tvm::s_tir::attr::auto_copy = "auto_copy"

constexpr

Mark auto copy for memhammer.

◆ buffer_bound

constexpr const char* tvm::s_tir::attr::buffer_bound = "buffer_bound"

constexpr

Mark stores/loads with their bounds.

◆ buffer_dim_align

constexpr const char* tvm::s_tir::attr::buffer_dim_align = "buffer_dim_align"

constexpr

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.

◆ double_buffer_scope

constexpr const char* tvm::s_tir::attr::double_buffer_scope = "double_buffer_scope"

constexpr

Marks production of double buffer data.

◆ double_buffer_write

constexpr const char* tvm::s_tir::attr::double_buffer_write = "double_buffer_write"

constexpr

Marks region used by double buffer write.

◆ explicit_read_region

constexpr const char* tvm::s_tir::attr::explicit_read_region = "explicit_read_region"

constexpr

Mark that a block has an explicitly specified read region. This is used to override the default read region inference in TIR.

◆ explicit_write_region

constexpr const char* tvm::s_tir::attr::explicit_write_region = "explicit_write_region"

constexpr

Mark that a block has an explicitly specified write region. This is used to override the default write region inference in TIR.

◆ fragment_layout

constexpr const char* tvm::s_tir::attr::fragment_layout = "fragment_layout"

constexpr

Mark that the layout of TensorCore fragment.

◆ fragment_shape

constexpr const char* tvm::s_tir::attr::fragment_shape = "fragment_shape"

constexpr

Mark that the shape of TensorCore fragment.

◆ hand_threaded

constexpr const char* tvm::s_tir::attr::hand_threaded = "hand_threaded"

constexpr

Mark that the kernel is hand threaded and doesn't need syncs inserted.

◆ irregular_loop_mark

constexpr const char* tvm::s_tir::attr::irregular_loop_mark = "irregular_loop_mark"

constexpr

,ark a ForNode represent an irregular loop of non-structural control flow edges.

◆ layout_free_buffers

constexpr const char* tvm::s_tir::attr::layout_free_buffers = "layout_free_buffers"

constexpr

Mark the buffers which is const access and can be transformed layout.

◆ layout_transforms

constexpr const char* tvm::s_tir::attr::layout_transforms = "layout_transforms"

constexpr

Marks the layout transforms to be used for a tensor.

Only applies to a tensor-like input, as it should be made part of the PrimFunc attributes for TIR.

◆ local_stage

constexpr const char* tvm::s_tir::attr::local_stage = "local_stage"

constexpr

Mark local stage constraint on data copy.

◆ manifest_shared_memory_local_stage

constexpr const char* tvm::s_tir::attr::manifest_shared_memory_local_stage

constexpr

Initial value:

=

"tirx.manifest_shared_memory_local_stage"

Mark the local stage for the shared memory access should be added.

◆ meta_schedule_auto_tensorize

constexpr const char* tvm::s_tir::attr::meta_schedule_auto_tensorize = "meta_schedule.auto_tensorize"

constexpr

Mark that a block should be further rewritten using tensorization.

◆ meta_schedule_auto_tensorize_init

constexpr const char* tvm::s_tir::attr::meta_schedule_auto_tensorize_init = "meta_schedule.auto_tensorize_init"

constexpr

Mark that the init statement of a block should be further rewritten using tensorization.

◆ meta_schedule_cooperative_fetch

constexpr const char* tvm::s_tir::attr::meta_schedule_cooperative_fetch = "meta_schedule.cooperative_fetch"

constexpr

Mark that the loop should be further skip and bound to environment threads to enable cooperative fetching.

◆ meta_schedule_inline_rule

constexpr const char* tvm::s_tir::attr::meta_schedule_inline_rule = "meta_schedule.inline_rule"

constexpr

Mark that a block is disallowed in auto inline.

◆ meta_schedule_layout_rewrite_preproc

constexpr const char* tvm::s_tir::attr::meta_schedule_layout_rewrite_preproc = "meta_schedule.layout_rewrite_preproc"

constexpr

Mark that a block is a preprocessor block for layout rewrite.

◆ meta_schedule_parallel

constexpr const char* tvm::s_tir::attr::meta_schedule_parallel = "meta_schedule.parallel"

constexpr

Mark auto-parallel setting on the block.

◆ meta_schedule_random_compute_producer

constexpr const char* tvm::s_tir::attr::meta_schedule_random_compute_producer

constexpr

Initial value:

=

"meta_schedule.random_compute_producer"

Mark the block whose producer needs to be applied by rule Random-Compute-Location.

◆ meta_schedule_thread_extent_high_inclusive

constexpr const char* tvm::s_tir::attr::meta_schedule_thread_extent_high_inclusive

constexpr

Initial value:

=

"meta_schedule.thread_extent_high_inclusive"

The allowed range of thread extent in thread bindings.

◆ meta_schedule_thread_extent_low_inclusive

constexpr const char* tvm::s_tir::attr::meta_schedule_thread_extent_low_inclusive

constexpr

Initial value:

=

"meta_schedule.thread_extent_low_inclusive"

The allowed range of thread extent in thread bindings.

◆ meta_schedule_tiling_structure

constexpr const char* tvm::s_tir::attr::meta_schedule_tiling_structure = "meta_schedule.tiling_structure"

constexpr

Mark the tiling structure of blocks that are applied by rule Multi-Level-Tiling.

◆ meta_schedule_unroll_explicit

constexpr const char* tvm::s_tir::attr::meta_schedule_unroll_explicit = "meta_schedule.unroll_explicit"

constexpr

Mark auto-unroll setting on the block.

◆ meta_schedule_unroll_implicit

constexpr const char* tvm::s_tir::attr::meta_schedule_unroll_implicit = "meta_schedule.unroll_implicit"

constexpr

Mark auto-unroll setting on the block.

◆ meta_schedule_vectorize

constexpr const char* tvm::s_tir::attr::meta_schedule_vectorize = "meta_schedule.vectorize"

constexpr

Mark auto-vectorize setting on the block.

◆ pragma_loop_partition_hint

constexpr const char* tvm::s_tir::attr::pragma_loop_partition_hint = "pragma_loop_partition_hint"

constexpr

Mark that the loop should be partitioned.

◆ reduce_scope

constexpr const char* tvm::s_tir::attr::reduce_scope = "reduce_scope"

constexpr

Mark of reduce scope.

◆ require_block_var_bound_predicate

constexpr const char* tvm::s_tir::attr::require_block_var_bound_predicate = "require_bound_predicate"

constexpr

Mark that the block need to add predicate for block var bounds during lowering.

◆ script_parsing_detect_access

constexpr const char* tvm::s_tir::attr::script_parsing_detect_access = "tirx.script_parsing_detect_access"

constexpr

Mark whether the script-completer need to fill in missing access region during script parsing.

Note: The result should be a integer mask with range [0, 4). if (mask & 1) the read region should be detected, if (mask & 2) the write region should be detected.

◆ software_pipeline_async_stages

constexpr const char* tvm::s_tir::attr::software_pipeline_async_stages = "software_pipeline_async_stages"

constexpr

List stages in the software pipeline that should run asynchronously.

Note: All statements in the provided stages are assumed to have asynchronous semantics (e.g. CUDA async global to shared memory copy).

◆ software_pipeline_order

constexpr const char* tvm::s_tir::attr::software_pipeline_order = "software_pipeline_order"

constexpr

Mark the order of a statement in the software pipeline.

◆ software_pipeline_stage

constexpr const char* tvm::s_tir::attr::software_pipeline_stage = "software_pipeline_stage"

constexpr

Mark the stage of a statement in the software pipeline.

◆ vector_bytes

constexpr const char* tvm::s_tir::attr::vector_bytes = "vector_bytes"

constexpr

Mark vectorization length constraint on block.

◆ virtual_thread

constexpr const char* tvm::s_tir::attr::virtual_thread = "virtual_thread"

constexpr

Mark launching of a virtual thread.

◆ warp_execution

constexpr const char* tvm::s_tir::attr::warp_execution = "warp_execution"

constexpr

Mark that a block is executed by a warp. This implies the extend of threadIdx.x is warp size.

Variables

Variable Documentation

◆ async_commit_queue_scope

◆ async_scope

◆ async_wait_inflight_count

◆ async_wait_queue_scope

◆ auto_copy

◆ buffer_bound

◆ buffer_dim_align

◆ double_buffer_scope

◆ double_buffer_write

◆ explicit_read_region

◆ explicit_write_region

◆ fragment_layout

◆ fragment_shape

◆ hand_threaded

◆ irregular_loop_mark

◆ layout_free_buffers

◆ layout_transforms

◆ local_stage

◆ manifest_shared_memory_local_stage

◆ meta_schedule_auto_tensorize

◆ meta_schedule_auto_tensorize_init

◆ meta_schedule_cooperative_fetch

◆ meta_schedule_inline_rule

◆ meta_schedule_layout_rewrite_preproc

◆ meta_schedule_parallel

◆ meta_schedule_random_compute_producer

◆ meta_schedule_thread_extent_high_inclusive

◆ meta_schedule_thread_extent_low_inclusive

◆ meta_schedule_tiling_structure

◆ meta_schedule_unroll_explicit

◆ meta_schedule_unroll_implicit

◆ meta_schedule_vectorize

◆ pragma_loop_partition_hint

◆ reduce_scope

◆ require_block_var_bound_predicate

◆ script_parsing_detect_access

◆ software_pipeline_async_stages

◆ software_pipeline_order

◆ software_pipeline_stage

◆ vector_bytes

◆ virtual_thread

◆ warp_execution