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_TIR_TRANSFORM_H_
25 #define TVM_TIR_TRANSFORM_H_
26 
27 #include <tvm/ir/transform.h>
28 #include <tvm/target/target.h>
29 #include <tvm/tir/expr.h>
30 #include <tvm/tir/function.h>
31 
32 #include <string>
33 #include <vector>
34 
35 namespace tvm {
36 namespace tir {
37 namespace transform {
38 
46 
47 /*
48  * \brief Create a function pass that optimizes PrimFuncs.
49  *
50  * \param pass_func The packed function that contains the optimization.
51  * \param opt_level The optimization level of the function pass.
52  * \param name The name of the function pass.
53  * \param required The list of the passes that the function pass is dependent on.
54  *
55  * \return The created function pass.
56  */
57 TVM_DLL Pass CreatePrimFuncPass(
58  const runtime::TypedPackedFunc<PrimFunc(PrimFunc, IRModule, PassContext)>& pass_func,
59  int opt_level, String name, tvm::Array<String> required);
60 
66 TVM_DLL Pass InjectPrefetch();
67 
68 // TODO(tvm-team): consolidate configs to the PassContext
78 TVM_DLL Pass StorageFlatten(int cache_line_size, bool create_bound_attribute = false);
79 
93 TVM_DLL Pass InjectCopyIntrin(String pragma_key, runtime::PackedFunc fintrin);
94 
100 TVM_DLL Pass CoProcSync();
101 
108 TVM_DLL Pass LiftAttrScope(String attr_key);
109 
115 TVM_DLL Pass LoopPartition();
116 
124 TVM_DLL Pass VectorizeLoop(bool enable_vectorize = true);
125 
131 TVM_DLL Pass InjectVirtualThread();
132 
138 TVM_DLL Pass InjectDoubleBuffer();
139 
148 TVM_DLL Pass StorageRewrite();
149 
156 TVM_DLL Pass UnrollLoop();
157 
163 TVM_DLL Pass RemoveNoOp();
164 
170 TVM_DLL Pass RewriteUnsafeSelect();
171 
177 TVM_DLL Pass Simplify();
178 
184 TVM_DLL Pass InstrumentBoundCheckers();
185 
215 TVM_DLL Pass MakePackedAPI(int num_unpacked_args);
216 
226 TVM_DLL Pass MakeUnpackedAPI();
227 
238 TVM_DLL Pass RemapThreadAxis(Map<String, IterVar> axis_map);
239 
247 TVM_DLL Pass LowerCustomDatatypes();
248 
254 TVM_DLL Pass DecorateDeviceScope();
255 
261 TVM_DLL Pass SplitHostDevice();
262 
268 TVM_DLL Pass SkipAssert();
269 
276 TVM_DLL Pass ThreadSync(String storage_scope);
277 
283 TVM_DLL Pass LowerThreadAllreduce();
284 
290 TVM_DLL Pass InferFragment();
291 
295 static constexpr const char* kDisableLowerTVMBuiltin = "disable_lower_builtin";
296 
301 TVM_DLL Pass LowerTVMBuiltin();
302 
308 TVM_DLL Pass LowerIntrin();
309 
314 TVM_DLL Pass LowerWarpMemory();
315 
324 
330 TVM_DLL Pass CombineContextCall();
331 
340 TVM_DLL Pass NarrowDataType(int target_bits);
341 
347 TVM_DLL Pass BF16Legalize();
348 
357 TVM_DLL Pass PointerValueTypeRewrite();
358 
365 TVM_DLL Pass HoistIfThenElse();
366 
378 TVM_DLL Pass HoistExpression();
379 
386 
391 TVM_DLL Pass LowerInitBlock();
392 
400 
407 TVM_DLL Pass ConvertBlocksToOpaque();
408 
446 TVM_DLL Pass CompactBufferAllocation();
447 
451 TVM_DLL Pass LegalizePackedCalls();
452 
457 TVM_DLL Pass LowerMatchBuffer();
458 
465 TVM_DLL Pass FlattenBuffer();
466 
467 /*
468  * \brief Flatten the multi-dimensional read/write
469  * to two dimensional texture Load/Store and realize
470  * texture buffer allocations.
471  *
472  * \return The Pass
473  */
474 TVM_DLL Pass TextureFlatten();
475 
476 /*
477  * \brief Lower VTCM allocations
478  *
479  * \return The Pass
480  */
481 TVM_DLL Pass LowerVtcmAlloc();
482 
490 TVM_DLL Pass CommonSubexprElimTIR(bool enable_cse_tir = true, bool identify_equiv_terms = false);
491 
502 TVM_DLL Pass UnifyThreadBinding();
503 
508 
516 TVM_DLL Pass ConvertForLoopsToSerial();
517 
525 
625 TVM_DLL Pass InjectSoftwarePipeline();
626 
627 TVM_DLL Pass BindParams(const Array<runtime::NDArray>& constants);
628 
635 
640 TVM_DLL Pass RenormalizeSplitPattern();
641 
646 TVM_DLL Pass BindTarget(Target target);
647 
652 TVM_DLL Pass AnnotateEntryFunc();
653 
658 TVM_DLL Pass Filter(runtime::TypedPackedFunc<bool(PrimFunc)> fcond);
659 
664 TVM_DLL Pass InjectPTXAsyncCopy();
665 
671 
672 } // namespace transform
673 } // namespace tir
674 } // namespace tvm
675 
676 #endif // TVM_TIR_TRANSFORM_H_
Pass LowerDeviceStorageAccessInfo()
Lower attached storage access information on device.
Pass LowerIntrin()
Lower the target specific function intrinsics in each of the function.
Pass InferFragment()
Infer the TensorCore fragment infomation using tensor intrinsics.
Pass StorageFlatten(int cache_line_size, bool create_bound_attribute=false)
Flatten the multi-dimensional read/write to single dimensional Load/Store.
Pass LowerMatchBuffer()
Remove match buffers inside the block. Also, it will validate the binding.
Pass SplitHostDevice()
Split the function into a host function and device functions.
Pass ConvertForLoopsToSerial()
This pass is post-scheduling pass to convert all Parallel For loops to Serial ones. This is run to attain lesser memory and/or executor/backend does not support parallel launch of For loops.
Pass MakeUnpackedAPI()
Transform the high-level PrimFunc to a C signature that can be used to call the operator directly...
runtime implementation for LibTorch/TorchScript.
Definition: analyzer.h:36
Pass LiftAttrScope(String attr_key)
Lift common attrs with attr_key to outer scope.
Pass InjectVirtualThread()
Inject virtual thread loops.
Pass CombineContextCall()
Combine context calls in the host function.
Pass InstrumentBoundCheckers()
Instruments bound checkers.
Pass LowerInitBlock()
Lower block init stmt into IfThenElse stmts.
Pass DecorateDeviceScope()
Decorate all the function&#39;s body as device function.
tvm::transform::Sequential Sequential
Definition: transform.h:49
Pass HoistIfThenElse()
Hoist loop-invariant IfThenElse nodes to outside the elligible loops.
Pass MakePackedAPI(int num_unpacked_args)
Transform the high-level PrimFunc to a low-level version that can be used as an API function...
Pass NarrowDataType(int target_bits)
Narrow down PrimExpr datatype in stmt to target_bits.
Pass LowerCrossThreadReduction()
Lower cross-thread reduction from thread bindings to intrinsic function calls.
Pass InjectCopyIntrin(String pragma_key, runtime::PackedFunc fintrin)
Inject copy intrinsics with optional pad.
TIR Function.
Pass LoopPartition()
partition loops in the stmt.
Pass BindTarget(Target target)
Annotate a PrimFunc with a given target.
Pass ConvertBlocksToOpaque()
Substitute all the block vars with the PrimExprs they are bound to, indicated by the corresponding it...
tvm::transform::PassNode PassNode
Definition: transform.h:44
Pass UnrollLoop()
unroll the constant loop marked by unroll. This pass also automatically attach pragma unroll tag to l...
Pass CoProcSync()
Detect and insert sync points to co-processor.
Pass AnnotateEntryFunc()
Set a PrimFunc as the entry point if it is only function in IRModule.
Pass InjectSoftwarePipeline()
This pass transforms annotated loops into pipelined ones where producers and consumers are overlapped...
Pass HoistExpression()
Hoist loop-invariant expressions nodes to outside the elligible loops.
Pass RemoveWeightLayoutRewriteBlock()
Remove the weight layout rewrite block.
Pass RewriteUnsafeSelect()
Detect and rewrite unsafe select that contains memory access.
Pass FlattenBuffer()
Flatten the multi-dimensional BufferLoad and BufferStore to single dimensional Load/Store. Also remove Block to ensure that the flattened TIR can not be scheduled again.
Pass LowerThreadAllreduce()
Lower cross thread alleduce.
Pass ExtractPrimFuncConstants()
Pass to collect tir non-scalar constants into module&#39;s &#39;Constants&#39; attribute.
Pass LowerTVMBuiltin()
Lower builtin intrinsics.
Pass CreatePrimFuncPass(const runtime::TypedPackedFunc< PrimFunc(PrimFunc, IRModule, PassContext)> &pass_func, int opt_level, String name, tvm::Array< String > required)
TIR expressions.
Array, container representing a contiguous sequence of ObjectRefs.
Definition: array.h:270
tvm::transform::PassInfoNode PassInfoNode
Definition: transform.h:46
Pass MergeDynamicSharedMemoryAllocations()
Pass VectorizeLoop(bool enable_vectorize=true)
Lower vectorization loops.
Pass Filter(runtime::TypedPackedFunc< bool(PrimFunc)> fcond)
Filter PrimFuncs with a given condition.
Pass StorageRewrite()
Rewrite storage allocation pattern. Moves the allocation to outer most possible scope. Trying to share space between allocations to make a static allocation plan when possible.
Pass BF16Legalize()
Legalize bf16 typed Ops. Add a cast to fp32 before Ops, then add a cast back to bf16.
tvm::transform::Pass Pass
Definition: transform.h:35
Pass InjectPrefetch()
Inject prefetch instructions into stmt.
Pass Simplify()
Run arithmetic simplifications on the statements and expressions.
Pass RemapThreadAxis(Map< String, IterVar > axis_map)
Remap the thread axis.
Pass UnifiedStaticMemoryPlanner()
This is the unified static memory planner pass that will plan for memory intra- and inter- PrimFuncs ...
tvm::transform::PassContext PassContext
Definition: transform.h:47
Compilation target object.
Pass LowerWarpMemory()
Lower warp memory access to low-level device related function calls.
Pass CompactBufferAllocation()
Compact the buffer access region by removing the buffer regions that are not accessed, i.e. narrowing the buffer shape and adjust the access region if necessary.
Pass PlanAndUpdateBufferAllocationLocation()
Locate the buffer allocation to the exact position (usually is the lca of buffer access). This pass will inject opaque block with alloc_buffers at the allocation site.
Pass PointerValueTypeRewrite()
Rewrite the pointer content type of arguments, as well as Alloc internal to the function to use the m...
Pass RenormalizeSplitPattern()
Renormalize the split pattern from floordiv(floormod()) to floormod(floordiv())
Pass RemoveNoOp()
Remove No Op from the Stmt.
Pass BindParams(const Array< runtime::NDArray > &constants)
Pass LowerCustomDatatypes()
Lower custom datatypes.
Pass InjectDoubleBuffer()
Inject double buffer statements.
Pass ThreadSync(String storage_scope)
Insert sync between parallel read/write of shared buffers.
Pass InjectPTXAsyncCopy()
Pass to rewrite global to shared memory copy on CUDA with asyncronous copy.
Pass CommonSubexprElimTIR(bool enable_cse_tir=true, bool identify_equiv_terms=false)
Implements a Common Subexpression Elimination (CSE) for TIR which introduces let-in bindings for dupl...
tvm::transform::PassContextNode PassContextNode
Definition: transform.h:48
Pass SkipAssert()
skip assert stmt.
tvm::transform::PassInfo PassInfo
Definition: transform.h:45
Pass UnifyThreadBinding()
Unify all the thread bindings for "blockIdx.x/y/z", "threadIdx.x/y/z", and "vthread.x/y/z". Before the unification, two vars that are bound to a thread axis (e.g., "threadIdx.x") use different IterVars and variables in their AttrStmts. After the unification, we use a consolidated IterVar and a variable for them.