tvm
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 namespace tvm {
31 namespace s_tir {
32 namespace attr {
33 
37 constexpr const char* async_commit_queue_scope = "async_commit_queue_scope";
38 constexpr const char* async_wait_queue_scope = "async_wait_queue_scope";
39 constexpr const char* async_wait_inflight_count = "async_wait_inflight_count";
40 
44 constexpr const char* async_scope = "async_scope";
45 
49 constexpr const char* double_buffer_scope = "double_buffer_scope";
50 
54 constexpr const char* double_buffer_write = "double_buffer_write";
55 
59 constexpr const char* fragment_shape = "fragment_shape";
60 
64 constexpr const char* fragment_layout = "fragment_layout";
65 
69 constexpr const char* pragma_loop_partition_hint = "pragma_loop_partition_hint";
70 
72 constexpr const char* reduce_scope = "reduce_scope";
73 
75 constexpr const char* virtual_thread = "virtual_thread";
76 
77 // -----------------------------------------------------------------------
78 // meta_schedule annotations
79 // -----------------------------------------------------------------------
80 
82 constexpr const char* meta_schedule_tiling_structure = "meta_schedule.tiling_structure";
83 
88 constexpr const char* meta_schedule_cooperative_fetch = "meta_schedule.cooperative_fetch";
89 
92  "meta_schedule.thread_extent_low_inclusive";
93 
96  "meta_schedule.thread_extent_high_inclusive";
97 
100  "meta_schedule.random_compute_producer";
101 
103 constexpr const char* meta_schedule_parallel = "meta_schedule.parallel";
104 
106 constexpr const char* meta_schedule_vectorize = "meta_schedule.vectorize";
107 
109 constexpr const char* meta_schedule_unroll_explicit = "meta_schedule.unroll_explicit";
110 
112 constexpr const char* meta_schedule_unroll_implicit = "meta_schedule.unroll_implicit";
113 
115 constexpr const char* meta_schedule_auto_tensorize = "meta_schedule.auto_tensorize";
116 
118 constexpr const char* meta_schedule_layout_rewrite_preproc = "meta_schedule.layout_rewrite_preproc";
119 
123 constexpr const char* meta_schedule_auto_tensorize_init = "meta_schedule.auto_tensorize_init";
124 
126 constexpr const char* meta_schedule_tensor_core_enabled = "meta_schedule.tensor_core_enabled";
127 
134 constexpr const char* meta_schedule_cache_type = "meta_schedule.cache_type";
135 
137 constexpr const int meta_schedule_cache_type_read = 0;
138 
140 constexpr const int meta_schedule_cache_type_write = 1;
141 
143 constexpr const char* meta_schedule_inline_rule = "meta_schedule.inline_rule";
144 
145 // -----------------------------------------------------------------------
146 // Schedule primitive / SBlock annotations
147 // -----------------------------------------------------------------------
148 
156 constexpr const char* script_parsing_detect_access = "tirx.script_parsing_detect_access";
157 
161 constexpr const char* require_block_var_bound_predicate = "require_bound_predicate";
162 
164 constexpr const char* software_pipeline_stage = "software_pipeline_stage";
165 
167 constexpr const char* software_pipeline_order = "software_pipeline_order";
168 
173 constexpr const char* software_pipeline_async_stages = "software_pipeline_async_stages";
174 
176 constexpr const char* layout_free_buffers = "layout_free_buffers";
177 
179 constexpr const char* manifest_shared_memory_local_stage =
180  "tirx.manifest_shared_memory_local_stage";
181 
188 constexpr const char* buffer_dim_align = "buffer_dim_align";
189 
193 constexpr const char* explicit_read_region = "explicit_read_region";
194 
198 constexpr const char* explicit_write_region = "explicit_write_region";
199 
201 constexpr const char* irregular_loop_mark = "irregular_loop_mark";
202 
204 constexpr const char* auto_copy = "auto_copy";
205 
207 constexpr const char* local_stage = "local_stage";
208 
210 constexpr const char* vector_bytes = "vector_bytes";
211 
216 constexpr const char* warp_execution = "warp_execution";
217 
224 constexpr const char* layout_transforms = "layout_transforms";
225 
233 constexpr const char* axis_separators = "axis_separators";
234 
238 constexpr const char* hand_threaded = "hand_threaded";
239 
240 } // namespace attr
241 } // namespace s_tir
242 } // namespace tvm
243 #endif // TVM_S_TIR_STMT_H_
constexpr const char * meta_schedule_auto_tensorize
Mark that a block should be further rewritten using tensorization.
Definition: stmt.h:115
constexpr const char * meta_schedule_tensor_core_enabled
Mark that tensor core is enabled in the PrimExpr.
Definition: stmt.h:126
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:123
constexpr const char * async_scope
Mark that the attached statement runs asynchronously.
Definition: stmt.h:44
constexpr const char * meta_schedule_unroll_explicit
Mark auto-unroll setting on the block.
Definition: stmt.h:109
constexpr const char * vector_bytes
Mark vectorization length constraint on block.
Definition: stmt.h:210
constexpr const char * reduce_scope
Mark of reduce scope.
Definition: stmt.h:72
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:88
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:156
constexpr const int meta_schedule_cache_type_read
Definition: stmt.h:137
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:99
constexpr const char * axis_separators
Marks the physical axis separators.
Definition: stmt.h:233
constexpr const char * software_pipeline_async_stages
List stages in the software pipeline that should run asynchronously.
Definition: stmt.h:173
constexpr const char * layout_free_buffers
Mark the buffers which is const access and can be transformed layout.
Definition: stmt.h:176
constexpr const char * local_stage
Mark local stage constraint on data copy.
Definition: stmt.h:207
constexpr const char * pragma_loop_partition_hint
Mark that the loop should be partitioned.
Definition: stmt.h:69
constexpr const char * meta_schedule_parallel
Mark auto-parallel setting on the block.
Definition: stmt.h:103
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:161
constexpr const char * virtual_thread
Mark launching of a virtual thread.
Definition: stmt.h:75
constexpr const char * buffer_dim_align
Mark alignment of buffer dimension stmt.node is Tensor stmt.value is tvm_tuple(dim,...
Definition: stmt.h:188
constexpr const char * async_wait_inflight_count
Definition: stmt.h:39
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:198
constexpr const char * async_commit_queue_scope
Annotations for invoking and synchronizing asynchronous operations.
Definition: stmt.h:37
constexpr const int meta_schedule_cache_type_write
Definition: stmt.h:140
constexpr const char * meta_schedule_cache_type
Mark a block as generated by cache_read or cache_write block. 0 means cache_read; 1 means cache_write...
Definition: stmt.h:134
constexpr const char * manifest_shared_memory_local_stage
Mark the local stage for the shared memory access should be added.
Definition: stmt.h:179
constexpr const char * layout_transforms
Marks the layout transforms to be used for a tensor.
Definition: stmt.h:224
constexpr const char * meta_schedule_layout_rewrite_preproc
Mark that a block is a preprocessor block for layout rewrite.
Definition: stmt.h:118
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:216
constexpr const char * software_pipeline_stage
Mark the stage of a statement in the software pipeline.
Definition: stmt.h:164
constexpr const char * auto_copy
Mark auto copy for memhammer.
Definition: stmt.h:204
constexpr const char * meta_schedule_vectorize
Mark auto-vectorize setting on the block.
Definition: stmt.h:106
constexpr const char * meta_schedule_thread_extent_high_inclusive
The allowed range of thread extent in thread bindings.
Definition: stmt.h:95
constexpr const char * fragment_shape
Mark that the shape of TensorCore fragment.
Definition: stmt.h:59
constexpr const char * hand_threaded
Mark that the kernel is hand threaded and doesn't need syncs inserted.
Definition: stmt.h:238
constexpr const char * double_buffer_write
Marks region used by double buffer write.
Definition: stmt.h:54
constexpr const char * software_pipeline_order
Mark the order of a statement in the software pipeline.
Definition: stmt.h:167
constexpr const char * meta_schedule_thread_extent_low_inclusive
The allowed range of thread extent in thread bindings.
Definition: stmt.h:91
constexpr const char * irregular_loop_mark
,ark a ForNode represent an irregular loop of non-structural control flow edges.
Definition: stmt.h:201
constexpr const char * meta_schedule_unroll_implicit
Mark auto-unroll setting on the block.
Definition: stmt.h:112
constexpr const char * meta_schedule_tiling_structure
Mark the tiling structure of blocks that are applied by rule Multi-Level-Tiling.
Definition: stmt.h:82
constexpr const char * fragment_layout
Mark that the layout of TensorCore fragment.
Definition: stmt.h:64
constexpr const char * double_buffer_scope
Marks production of double buffer data.
Definition: stmt.h:49
constexpr const char * async_wait_queue_scope
Definition: stmt.h:38
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:193
constexpr const char * meta_schedule_inline_rule
Mark that a block is disallowed in auto inline.
Definition: stmt.h:143
An object that builds and maintains block scope and StmtSref mapping for Dependence analysis.
Definition: analyzer.h:37