tvm
transform.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  */
19 
24 #ifndef TVM_S_TIR_TRANSFORM_H_
25 #define TVM_S_TIR_TRANSFORM_H_
26 
27 #include <tvm/ir/transform.h>
28 #include <tvm/target/target.h>
29 #include <tvm/tir/transform.h>
30 
31 #include <string>
32 #include <vector>
33 
34 namespace tvm {
35 namespace s_tir {
36 namespace transform {
37 
41 
47 
54 
59 TVM_DLL Pass LowerInitBlock();
60 
68 
76 
82 
121 TVM_DLL Pass CompactBufferAllocation(bool is_strict = true);
122 
128 
134 
140 
146 
158 
196 
201 TVM_DLL Pass LowerAutoCopy();
202 
208 
211 
217 TVM_DLL Pass LoopPartition();
218 
225 
232 
233 } // namespace transform
234 } // namespace s_tir
235 } // namespace tvm
236 
237 #endif // TVM_S_TIR_TRANSFORM_H_
Definition: transform.h:400
tvm::transform::PassContext PassContext
Definition: transform.h:40
tvm::transform::Pass Pass
Definition: transform.h:38
Pass CanonicalizeLoop()
Canonicalize loop to start from zero .
Pass LowerAutoCopy()
Automatically do memory optimizations for auto copy blocks.
Pass CompactBufferAllocation(bool is_strict=true)
Compact the buffer access region by removing the buffer regions that are not accessed,...
Pass LowerInitBlock()
Lower block init stmt into IfThenElse stmts.
Pass InjectVirtualThread()
Inject virtual thread loops.
Pass PlanAndUpdateBufferAllocationLocation()
Locate the buffer allocation to the exact position (usually is the lca of buffer access)....
Pass LowerCrossThreadReduction()
Lower cross-thread reduction from thread bindings to intrinsic function calls.
Pass InjectSoftwarePipeline()
This pass transforms annotated loops into pipelined ones where producers and consumers are overlapped...
Pass InjectDoubleBuffer()
Inject double buffer statements.
Pass TransformMmaBufferLayout()
Transform Mma scope (m16n8k8.matrixA/B/C) to local scope with layout transformation.
Pass LowerMatchBuffer()
Remove match buffers inside the block. Also, it will validate the binding.
Pass LoopPartition()
partition loops in the stmt.
Pass AnnotateIrregularLoop()
Annotate irregular loop mark.
Pass LowerOpaqueBlock()
Remove the block to ensure that the TIR can not be scheduled again.
Pass InjectPermutedLayout()
Inject permuted layout for shared memory.
Pass ConvertBlocksToOpaque()
Substitute all the block vars with the PrimExprs they are bound to, indicated by the corresponding it...
Pass ManifestSharedMemoryLocalStage()
Add the explicit local stage for the shared memory access on GPU.
Pass LiftThreadBinding()
Lift the same thread bindings to their LCA loops.
Pass UnifyThreadBinding()
Unify all the thread bindings for "blockIdx.x/y/z", "threadIdx.x/y/z", and "vthread....
Pass CreatePrimFuncPass(std::function< PrimFunc(PrimFunc, IRModule, PassContext)> pass_func, int opt_level, ffi::String name, tvm::ffi::Array< ffi::String > required, bool traceable=false)
Performance counters for profiling via the PAPI library.
Definition: analyzer.h:37
Compilation target object.
TIR specific transformation passes.