tvm
Loading...
Searching...
No Matches
Variables
tvm::s_tir::attr Namespace Reference

Variables

constexpr const char * kIsScheduled = "tirx.is_scheduled"
 Mark the function as scheduled, so the default schedule will pass will skip it.
 
constexpr const char * kLayoutFreeTensors = "layout_free_tensors"
 Mark the tensors which is const access and can be transformed layout.
 
constexpr const char * kLayoutTransforms = "layout_transforms"
 Marks the layout transforms to be used for a tensor.
 
constexpr const char * kEstimatedFlops = "estimated_flops"
 
constexpr const char * kThreadBinding = "thread_binding"
 Hardware thread axis bound to a parallel loop during S-TIR scheduling.
 
constexpr const char * kDoubleBufferScope = "double_buffer_scope"
 SBlock annotation selecting a write-tensor index for double buffering in InjectSoftwarePipeline.
 
constexpr const char * kFragmentShape = "fragment_shape"
 String-valued allocation Call attribute containing the TensorCore fragment shape.
 
constexpr const char * kFragmentLayout = "fragment_layout"
 String-valued allocation Call attribute containing the TensorCore fragment layout.
 
constexpr const char * kLoopPartitionHint = "loop_partition_hint"
 For annotation requesting partitioning when its PrimExpr value is provably true.
 
constexpr const char * kMetaScheduleTilingStructure = "meta_schedule.tiling_structure"
 Mark the tiling structure of blocks that are applied by rule Multi-Level-Tiling.
 
constexpr const char * 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 * kMetaScheduleThreadExtentLowInclusive
 The allowed range of thread extent in thread bindings.
 
constexpr const char * kMetaScheduleThreadExtentHighInclusive
 The allowed range of thread extent in thread bindings.
 
constexpr const char * kMetaScheduleRandomComputeProducer = "meta_schedule.random_compute_producer"
 Mark the block whose producer needs to be applied by rule Random-Compute-Location.
 
constexpr const char * kMetaScheduleParallel = "meta_schedule.parallel"
 Mark auto-parallel setting on the block.
 
constexpr const char * kMetaScheduleVectorize = "meta_schedule.vectorize"
 Mark auto-vectorize setting on the block.
 
constexpr const char * kMetaScheduleUnrollExplicit = "meta_schedule.unroll_explicit"
 Mark auto-unroll setting on the block.
 
constexpr const char * kMetaScheduleUnrollImplicit = "meta_schedule.unroll_implicit"
 Mark auto-unroll setting on the block.
 
constexpr const char * kMetaScheduleAutoTensorize = "meta_schedule.auto_tensorize"
 Mark that a block should be further rewritten using tensorization.
 
constexpr const char * kMetaScheduleLayoutRewritePreproc = "meta_schedule.layout_rewrite_preproc"
 Mark that a block is a preprocessor block for layout rewrite.
 
constexpr const char * kMetaScheduleAutoTensorizeInit = "meta_schedule.auto_tensorize_init"
 Mark that the init statement of a block should be further rewritten using tensorization.
 
constexpr const char * kMetaScheduleInlineRule = "meta_schedule.inline_rule"
 Mark that a block is disallowed in auto inline.
 
constexpr const char * kTensorAllocatedAddr = "s_tir.tensor_allocated_addr"
 TensorAllocatedAddresses for block-owned allocations and match tensors.
 
constexpr const char * kScriptParsingDetectAccess = "tirx.script_parsing_detect_access"
 Mark whether the script-completer need to fill in missing access region during script parsing.
 
constexpr const char * kRequireBlockVarBoundPredicate = "require_bound_predicate"
 Mark that the block need to add predicate for block var bounds during lowering.
 
constexpr const char * kSoftwarePipelineStage = "software_pipeline_stage"
 Mark the stage of a statement in the software pipeline.
 
constexpr const char * kSoftwarePipelineOrder = "software_pipeline_order"
 Mark the order of a statement in the software pipeline.
 
constexpr const char * kSoftwarePipelineAsyncStages = "software_pipeline_async_stages"
 List stages in the software pipeline that should run asynchronously.
 
constexpr const char * kManifestSharedMemoryLocalStage = "tirx.manifest_shared_memory_local_stage"
 Mark the local stage for the shared memory access should be added.
 
constexpr const char * 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 * 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 * 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 * kIrregularLoopMark = "irregular_loop_mark"
 ,ark a ForNode represent an irregular loop of non-structural control flow edges.
 
constexpr const char * kAutoCopy = "auto_copy"
 Mark auto copy for memhammer.
 
constexpr const char * kLocalStage = "local_stage"
 Mark local stage constraint on data copy.
 
constexpr const char * kVectorBytes = "vector_bytes"
 Mark vectorization length constraint on block.
 
constexpr const char * 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 * kPermutedLayout = "permuted_layout"
 
constexpr const char * kScheduleRule = "schedule_rule"
 
constexpr const char * kMetaScheduleWriteCacheLevel = "s_tir.meta_schedule.write_cache_level"
 

Variable Documentation

◆ kAutoCopy

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

Mark auto copy for memhammer.

◆ kDoubleBufferScope

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

SBlock annotation selecting a write-tensor index for double buffering in InjectSoftwarePipeline.

◆ kEstimatedFlops

constexpr const char* tvm::s_tir::attr::kEstimatedFlops = "estimated_flops"
constexpr

◆ kExplicitReadRegion

constexpr const char* tvm::s_tir::attr::kExplicitReadRegion = "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.

◆ kExplicitWriteRegion

constexpr const char* tvm::s_tir::attr::kExplicitWriteRegion = "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.

◆ kFragmentLayout

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

String-valued allocation Call attribute containing the TensorCore fragment layout.

◆ kFragmentShape

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

String-valued allocation Call attribute containing the TensorCore fragment shape.

◆ kIrregularLoopMark

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

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

◆ kIsScheduled

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

Mark the function as scheduled, so the default schedule will pass will skip it.

Type: IntImm

◆ kLayoutFreeTensors

constexpr const char* tvm::s_tir::attr::kLayoutFreeTensors = "layout_free_tensors"
constexpr

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

◆ kLayoutTransforms

constexpr const char* tvm::s_tir::attr::kLayoutTransforms = "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 Function attributes for TIR.

◆ kLocalStage

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

Mark local stage constraint on data copy.

◆ kLoopPartitionHint

constexpr const char* tvm::s_tir::attr::kLoopPartitionHint = "loop_partition_hint"
constexpr

For annotation requesting partitioning when its PrimExpr value is provably true.

◆ kManifestSharedMemoryLocalStage

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

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

◆ kMetaScheduleAutoTensorize

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

Mark that a block should be further rewritten using tensorization.

◆ kMetaScheduleAutoTensorizeInit

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

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

◆ kMetaScheduleCooperativeFetch

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

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

◆ kMetaScheduleInlineRule

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

Mark that a block is disallowed in auto inline.

◆ kMetaScheduleLayoutRewritePreproc

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

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

◆ kMetaScheduleParallel

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

Mark auto-parallel setting on the block.

◆ kMetaScheduleRandomComputeProducer

constexpr const char* tvm::s_tir::attr::kMetaScheduleRandomComputeProducer = "meta_schedule.random_compute_producer"
constexpr

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

◆ kMetaScheduleThreadExtentHighInclusive

constexpr const char* tvm::s_tir::attr::kMetaScheduleThreadExtentHighInclusive
constexpr
Initial value:
=
"meta_schedule.thread_extent_high_inclusive"

The allowed range of thread extent in thread bindings.

◆ kMetaScheduleThreadExtentLowInclusive

constexpr const char* tvm::s_tir::attr::kMetaScheduleThreadExtentLowInclusive
constexpr
Initial value:
=
"meta_schedule.thread_extent_low_inclusive"

The allowed range of thread extent in thread bindings.

◆ kMetaScheduleTilingStructure

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

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

◆ kMetaScheduleUnrollExplicit

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

Mark auto-unroll setting on the block.

◆ kMetaScheduleUnrollImplicit

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

Mark auto-unroll setting on the block.

◆ kMetaScheduleVectorize

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

Mark auto-vectorize setting on the block.

◆ kMetaScheduleWriteCacheLevel

constexpr const char* tvm::s_tir::attr::kMetaScheduleWriteCacheLevel = "s_tir.meta_schedule.write_cache_level"
constexpr

◆ kPermutedLayout

constexpr const char* tvm::s_tir::attr::kPermutedLayout = "permuted_layout"
constexpr

◆ kRequireBlockVarBoundPredicate

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

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

◆ kScheduleRule

constexpr const char* tvm::s_tir::attr::kScheduleRule = "schedule_rule"
constexpr

◆ kScriptParsingDetectAccess

constexpr const char* tvm::s_tir::attr::kScriptParsingDetectAccess = "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.

◆ kSoftwarePipelineAsyncStages

constexpr const char* tvm::s_tir::attr::kSoftwarePipelineAsyncStages = "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).

◆ kSoftwarePipelineOrder

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

Mark the order of a statement in the software pipeline.

◆ kSoftwarePipelineStage

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

Mark the stage of a statement in the software pipeline.

◆ kTensorAllocatedAddr

constexpr const char* tvm::s_tir::attr::kTensorAllocatedAddr = "s_tir.tensor_allocated_addr"
constexpr

TensorAllocatedAddresses for block-owned allocations and match tensors.

◆ kTensorDimAlign

constexpr const char* tvm::s_tir::attr::kTensorDimAlign = "tensor_dim_align"
constexpr

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.

◆ kThreadBinding

constexpr const char* tvm::s_tir::attr::kThreadBinding = "thread_binding"
constexpr

Hardware thread axis bound to a parallel loop during S-TIR scheduling.

◆ kVectorBytes

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

Mark vectorization length constraint on block.

◆ kWarpExecution

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

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