tvm
Loading...
Searching...
No Matches
stmt.h
Go to the documentation of this file.
1/*
2 * Licensed to the Apache Software Foundation (ASF) under one
3 * or more contributor license agreements. See the NOTICE file
4 * distributed with this work for additional information
5 * regarding copyright ownership. The ASF licenses this file
6 * to you under the Apache License, Version 2.0 (the
7 * "License"); you may not use this file except in compliance
8 * with the License. You may obtain a copy of the License at
9 *
10 * http://www.apache.org/licenses/LICENSE-2.0
11 *
12 * Unless required by applicable law or agreed to in writing,
13 * software distributed under the License is distributed on an
14 * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
15 * KIND, either express or implied. See the License for the
16 * specific language governing permissions and limitations
17 * under the License.
18 */
27#ifndef TVM_S_TIR_STMT_H_
28#define TVM_S_TIR_STMT_H_
29
30#include <tvm/ir/prim/expr.h>
31
32namespace tvm {
33namespace s_tir {
34using namespace tvm::prim;
35namespace attr {
36
40constexpr const char* async_commit_queue_scope = "async_commit_queue_scope";
41constexpr const char* async_wait_queue_scope = "async_wait_queue_scope";
42constexpr const char* async_wait_inflight_count = "async_wait_inflight_count";
43
47constexpr const char* async_scope = "async_scope";
48
50constexpr const char* buffer_bound = "buffer_bound";
51
55constexpr const char* double_buffer_scope = "double_buffer_scope";
56
60constexpr const char* double_buffer_write = "double_buffer_write";
61
65constexpr const char* fragment_shape = "fragment_shape";
66
70constexpr const char* fragment_layout = "fragment_layout";
71
75constexpr const char* pragma_loop_partition_hint = "pragma_loop_partition_hint";
76
78constexpr const char* reduce_scope = "reduce_scope";
79
81constexpr const char* virtual_thread = "virtual_thread";
82
83// -----------------------------------------------------------------------
84// meta_schedule annotations
85// -----------------------------------------------------------------------
86
88constexpr const char* meta_schedule_tiling_structure = "meta_schedule.tiling_structure";
89
94constexpr const char* meta_schedule_cooperative_fetch = "meta_schedule.cooperative_fetch";
95
98 "meta_schedule.thread_extent_low_inclusive";
99
102 "meta_schedule.thread_extent_high_inclusive";
103
106 "meta_schedule.random_compute_producer";
107
109constexpr const char* meta_schedule_parallel = "meta_schedule.parallel";
110
112constexpr const char* meta_schedule_vectorize = "meta_schedule.vectorize";
113
115constexpr const char* meta_schedule_unroll_explicit = "meta_schedule.unroll_explicit";
116
118constexpr const char* meta_schedule_unroll_implicit = "meta_schedule.unroll_implicit";
119
121constexpr const char* meta_schedule_auto_tensorize = "meta_schedule.auto_tensorize";
122
124constexpr const char* meta_schedule_layout_rewrite_preproc = "meta_schedule.layout_rewrite_preproc";
125
129constexpr const char* meta_schedule_auto_tensorize_init = "meta_schedule.auto_tensorize_init";
130
132constexpr const char* meta_schedule_inline_rule = "meta_schedule.inline_rule";
133
134// -----------------------------------------------------------------------
135// Schedule primitive / SBlock annotations
136// -----------------------------------------------------------------------
137
145constexpr const char* script_parsing_detect_access = "tirx.script_parsing_detect_access";
146
150constexpr const char* require_block_var_bound_predicate = "require_bound_predicate";
151
153constexpr const char* software_pipeline_stage = "software_pipeline_stage";
154
156constexpr const char* software_pipeline_order = "software_pipeline_order";
157
162constexpr const char* software_pipeline_async_stages = "software_pipeline_async_stages";
163
165constexpr const char* layout_free_buffers = "layout_free_buffers";
166
169 "tirx.manifest_shared_memory_local_stage";
170
177constexpr const char* buffer_dim_align = "buffer_dim_align";
178
182constexpr const char* explicit_read_region = "explicit_read_region";
183
187constexpr const char* explicit_write_region = "explicit_write_region";
188
190constexpr const char* irregular_loop_mark = "irregular_loop_mark";
191
193constexpr const char* auto_copy = "auto_copy";
194
196constexpr const char* local_stage = "local_stage";
197
199constexpr const char* vector_bytes = "vector_bytes";
200
205constexpr const char* warp_execution = "warp_execution";
206
213constexpr const char* layout_transforms = "layout_transforms";
214
218constexpr const char* hand_threaded = "hand_threaded";
219
220} // namespace attr
221} // namespace s_tir
222} // namespace tvm
223#endif // TVM_S_TIR_STMT_H_
TIR expressions.
Definition builtin.h:25
constexpr const char * meta_schedule_auto_tensorize
Mark that a block should be further rewritten using tensorization.
Definition stmt.h:121
constexpr const char * meta_schedule_auto_tensorize_init
Mark that the init statement of a block should be further rewritten using tensorization.
Definition stmt.h:129
constexpr const char * async_scope
Mark that the attached statement runs asynchronously.
Definition stmt.h:47
constexpr const char * meta_schedule_unroll_explicit
Mark auto-unroll setting on the block.
Definition stmt.h:115
constexpr const char * vector_bytes
Mark vectorization length constraint on block.
Definition stmt.h:199
constexpr const char * reduce_scope
Mark of reduce scope.
Definition stmt.h:78
constexpr const char * meta_schedule_cooperative_fetch
Mark that the loop should be further skip and bound to environment threads to enable cooperative fetc...
Definition stmt.h:94
constexpr const char * script_parsing_detect_access
Mark whether the script-completer need to fill in missing access region during script parsing.
Definition stmt.h:145
constexpr const char * meta_schedule_random_compute_producer
Mark the block whose producer needs to be applied by rule Random-Compute-Location.
Definition stmt.h:105
constexpr const char * software_pipeline_async_stages
List stages in the software pipeline that should run asynchronously.
Definition stmt.h:162
constexpr const char * layout_free_buffers
Mark the buffers which is const access and can be transformed layout.
Definition stmt.h:165
constexpr const char * local_stage
Mark local stage constraint on data copy.
Definition stmt.h:196
constexpr const char * pragma_loop_partition_hint
Mark that the loop should be partitioned.
Definition stmt.h:75
constexpr const char * meta_schedule_parallel
Mark auto-parallel setting on the block.
Definition stmt.h:109
constexpr const char * require_block_var_bound_predicate
Mark that the block need to add predicate for block var bounds during lowering.
Definition stmt.h:150
constexpr const char * virtual_thread
Mark launching of a virtual thread.
Definition stmt.h:81
constexpr const char * buffer_dim_align
Mark alignment of buffer dimension stmt.node is Tensor stmt.value is tvm_tuple(dim,...
Definition stmt.h:177
constexpr const char * async_wait_inflight_count
Definition stmt.h:42
constexpr const char * explicit_write_region
Mark that a block has an explicitly specified write region. This is used to override the default writ...
Definition stmt.h:187
constexpr const char * async_commit_queue_scope
Annotations for invoking and synchronizing asynchronous operations.
Definition stmt.h:40
constexpr const char * manifest_shared_memory_local_stage
Mark the local stage for the shared memory access should be added.
Definition stmt.h:168
constexpr const char * layout_transforms
Marks the layout transforms to be used for a tensor.
Definition stmt.h:213
constexpr const char * meta_schedule_layout_rewrite_preproc
Mark that a block is a preprocessor block for layout rewrite.
Definition stmt.h:124
constexpr const char * warp_execution
Mark that a block is executed by a warp. This implies the extend of threadIdx.x is warp size.
Definition stmt.h:205
constexpr const char * software_pipeline_stage
Mark the stage of a statement in the software pipeline.
Definition stmt.h:153
constexpr const char * auto_copy
Mark auto copy for memhammer.
Definition stmt.h:193
constexpr const char * meta_schedule_vectorize
Mark auto-vectorize setting on the block.
Definition stmt.h:112
constexpr const char * meta_schedule_thread_extent_high_inclusive
The allowed range of thread extent in thread bindings.
Definition stmt.h:101
constexpr const char * fragment_shape
Mark that the shape of TensorCore fragment.
Definition stmt.h:65
constexpr const char * hand_threaded
Mark that the kernel is hand threaded and doesn't need syncs inserted.
Definition stmt.h:218
constexpr const char * double_buffer_write
Marks region used by double buffer write.
Definition stmt.h:60
constexpr const char * software_pipeline_order
Mark the order of a statement in the software pipeline.
Definition stmt.h:156
constexpr const char * meta_schedule_thread_extent_low_inclusive
The allowed range of thread extent in thread bindings.
Definition stmt.h:97
constexpr const char * irregular_loop_mark
,ark a ForNode represent an irregular loop of non-structural control flow edges.
Definition stmt.h:190
constexpr const char * meta_schedule_unroll_implicit
Mark auto-unroll setting on the block.
Definition stmt.h:118
constexpr const char * meta_schedule_tiling_structure
Mark the tiling structure of blocks that are applied by rule Multi-Level-Tiling.
Definition stmt.h:88
constexpr const char * fragment_layout
Mark that the layout of TensorCore fragment.
Definition stmt.h:70
constexpr const char * buffer_bound
Mark stores/loads with their bounds.
Definition stmt.h:50
constexpr const char * double_buffer_scope
Marks production of double buffer data.
Definition stmt.h:55
constexpr const char * async_wait_queue_scope
Definition stmt.h:41
constexpr const char * explicit_read_region
Mark that a block has an explicitly specified read region. This is used to override the default read ...
Definition stmt.h:182
constexpr const char * meta_schedule_inline_rule
Mark that a block is disallowed in auto inline.
Definition stmt.h:132
An object that builds and maintains block scope and StmtSref mapping for Dependence analysis.
Definition analyzer.h:40