|
tvm
|
S-TIR (Schedulable TIR) statement attribute declarations. More...
Go to the source code of this file.
Namespaces | |
| tvm | |
| An object that builds and maintains block scope and StmtSref mapping for Dependence analysis. | |
| tvm::s_tir | |
| 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. More... | |
| 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. More... | |
| constexpr const char * | tvm::s_tir::attr::double_buffer_scope = "double_buffer_scope" |
| Marks production of double buffer data. More... | |
| constexpr const char * | tvm::s_tir::attr::double_buffer_write = "double_buffer_write" |
| Marks region used by double buffer write. More... | |
| constexpr const char * | tvm::s_tir::attr::fragment_shape = "fragment_shape" |
| Mark that the shape of TensorCore fragment. More... | |
| constexpr const char * | tvm::s_tir::attr::fragment_layout = "fragment_layout" |
| Mark that the layout of TensorCore fragment. More... | |
| constexpr const char * | tvm::s_tir::attr::pragma_loop_partition_hint = "pragma_loop_partition_hint" |
| Mark that the loop should be partitioned. More... | |
| constexpr const char * | tvm::s_tir::attr::reduce_scope = "reduce_scope" |
| Mark of reduce scope. More... | |
| constexpr const char * | tvm::s_tir::attr::virtual_thread = "virtual_thread" |
| Mark launching of a virtual thread. More... | |
| 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. More... | |
| 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. More... | |
| constexpr const char * | tvm::s_tir::attr::meta_schedule_thread_extent_low_inclusive |
| The allowed range of thread extent in thread bindings. More... | |
| constexpr const char * | tvm::s_tir::attr::meta_schedule_thread_extent_high_inclusive |
| The allowed range of thread extent in thread bindings. More... | |
| 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. More... | |
| constexpr const char * | tvm::s_tir::attr::meta_schedule_parallel = "meta_schedule.parallel" |
| Mark auto-parallel setting on the block. More... | |
| constexpr const char * | tvm::s_tir::attr::meta_schedule_vectorize = "meta_schedule.vectorize" |
| Mark auto-vectorize setting on the block. More... | |
| constexpr const char * | tvm::s_tir::attr::meta_schedule_unroll_explicit = "meta_schedule.unroll_explicit" |
| Mark auto-unroll setting on the block. More... | |
| constexpr const char * | tvm::s_tir::attr::meta_schedule_unroll_implicit = "meta_schedule.unroll_implicit" |
| Mark auto-unroll setting on the block. More... | |
| 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. More... | |
| 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. More... | |
| 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. More... | |
| constexpr const char * | tvm::s_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::s_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::s_tir::attr::meta_schedule_cache_type_read = 0 |
| constexpr const int | tvm::s_tir::attr::meta_schedule_cache_type_write = 1 |
| constexpr const char * | tvm::s_tir::attr::meta_schedule_inline_rule = "meta_schedule.inline_rule" |
| Mark that a block is disallowed in auto inline. More... | |
| 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. More... | |
| 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. More... | |
| constexpr const char * | tvm::s_tir::attr::software_pipeline_stage = "software_pipeline_stage" |
| Mark the stage of a statement in the software pipeline. More... | |
| constexpr const char * | tvm::s_tir::attr::software_pipeline_order = "software_pipeline_order" |
| Mark the order of a statement in the software pipeline. More... | |
| 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. More... | |
| 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. More... | |
| constexpr const char * | tvm::s_tir::attr::manifest_shared_memory_local_stage |
| Mark the local stage for the shared memory access should be added. More... | |
| 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. More... | |
| 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. More... | |
| 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. More... | |
| 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. More... | |
| constexpr const char * | tvm::s_tir::attr::auto_copy = "auto_copy" |
| Mark auto copy for memhammer. More... | |
| constexpr const char * | tvm::s_tir::attr::local_stage = "local_stage" |
| Mark local stage constraint on data copy. More... | |
| constexpr const char * | tvm::s_tir::attr::vector_bytes = "vector_bytes" |
| Mark vectorization length constraint on block. More... | |
| 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. More... | |
| constexpr const char * | tvm::s_tir::attr::layout_transforms = "layout_transforms" |
| Marks the layout transforms to be used for a tensor. More... | |
| constexpr const char * | tvm::s_tir::attr::axis_separators = "axis_separators" |
| Marks the physical axis separators. More... | |
| constexpr const char * | tvm::s_tir::attr::hand_threaded = "hand_threaded" |
| Mark that the kernel is hand threaded and doesn't need syncs inserted. More... | |
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.
1.9.1