|
tvm
|
Functions | |
| Pass | VerifySSA () |
| Pass variant of VerifySSA. More... | |
| Pass | VerifyMemory () |
| Pass variant of VerifyMemory. More... | |
| Pass | VerifyGPUCode (ffi::Map< ffi::String, PrimExpr > constraints) |
| Pass variant of VerifyGPUCode. More... | |
| Pass | VerifyVTCMLimit (ffi::Optional< Target > target=std::nullopt) |
| Pass to checks if the size of the allocated vtcm memory satisfies the limit. More... | |
| Pass | OOBChecker () |
| Statically check TIR code for out of bounds array access. More... | |
| 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) |
| Pass | VectorizeLoop (bool enable_vectorize=true) |
| Lower vectorization loops. More... | |
| 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. More... | |
| Pass | UnrollLoop () |
| unroll the constant loop marked by unroll. This pass also automatically attach pragma unroll tag to loops which meets the standard. More... | |
| Pass | RemoveNoOp () |
| Remove No Op from the Stmt. More... | |
| Pass | RewriteUnsafeSelect () |
| Detect and rewrite unsafe select that contains memory access. More... | |
| Pass | Simplify () |
| Run arithmetic simplifications on the statements and expressions. More... | |
| Pass | ConvertSSA () |
| Convert an IRModule to be SSA form. More... | |
| Pass | InstrumentBoundCheckers () |
| Instruments bound checkers. More... | |
| Pass | MakePackedAPI () |
| Transform the high-level PrimFunc to a low-level version that can be used as an API function. More... | |
| Pass | MakeUnpackedAPI () |
| Transform the high-level PrimFunc to a C signature that can be used to call the operator directly. More... | |
| Pass | RemapThreadAxis (ffi::Map< ffi::String, IterVar > axis_map) |
| Remap the thread axis. More... | |
| Pass | LowerCustomDatatypes () |
| Lower custom datatypes. More... | |
| Pass | DecorateDeviceScope () |
| Decorate all the function's body as device function. More... | |
| Pass | AnnotateDeviceRegions () |
| Annotate locations that should be run on the device. More... | |
| Pass | SplitHostDevice () |
| Split the function into a host function and device functions. More... | |
| Pass | LowerDeviceKernelLaunch () |
| Lower cross-device function calls. More... | |
| Pass | SkipAssert () |
| skip assert stmt. More... | |
| Pass | ThreadSync (ffi::String storage_scope) |
| Insert sync between parallel read/write of shared buffers. More... | |
| Pass | LowerThreadAllreduce () |
| Lower cross thread alleduce. More... | |
| Pass | InferFragment () |
| Infer the TensorCore fragment infomation using tensor intrinsics. More... | |
| Pass | LowerTVMBuiltin () |
| Lower builtin intrinsics. More... | |
| Pass | LowerIntrin () |
| Lower the target specific function intrinsics in each of the function. More... | |
| Pass | LowerWarpMemory () |
| Lower warp memory access to low-level device related function calls. More... | |
| Pass | LowerDeviceStorageAccessInfo () |
| Lower attached storage access information on device. More... | |
| Pass | CombineContextCall () |
| Combine context calls in the host function. More... | |
| Pass | NarrowDataType (int target_bits) |
| Narrow down PrimExpr datatype in stmt to target_bits. More... | |
| Pass | ForceNarrowIndexToInt32 () |
| Force to narrow down indexing expressions and integer buffers to int32 dtype. More... | |
| Pass | BF16ComputeLegalize () |
| Legalize bf16 compute Ops. Add a cast to fp32 before Ops, then add a cast back to bf16. More... | |
| Pass | FP8ComputeLegalize (ffi::String promote_dtype="float16") |
| Legalize fp8 compute Ops. Add a cast to fp16/fp32 before Ops, then add a cast back to fp8. More... | |
| Pass | BF16StorageLegalize () |
| Legalize bf16 storage types to u16. More... | |
| Pass | FP8StorageLegalize () |
| Legalize fp8 storage types to u8. More... | |
| Pass | InlinePrivateFunctions () |
| Inline calls to private functions. More... | |
| Pass | PointerValueTypeRewrite () |
| Rewrite the pointer content type of arguments, as well as Alloc internal to the function to use the most frequently accessed type for load/store to avoid pointer casting in backend when possible. More... | |
| Pass | HoistIfThenElse () |
| Hoist loop-invariant IfThenElse nodes to outside the elligible loops. More... | |
| Pass | HoistExpression () |
| Hoist loop-invariant expressions nodes to outside the elligible loops. More... | |
| Pass | FlattenBuffer () |
| Flatten the multi-dimensional BufferLoad and BufferStore to single dimensional BufferLoad/BufferStore for the TIR not contains opaque block. More... | |
| Pass | LowerVtcmAlloc () |
| Pass | LowerAsyncDMA () |
| Lower Async TIR primitives to DMA copy and wait builtins. More... | |
| 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 duplicated sub-expressions. More... | |
| Pass | MergeSharedMemoryAllocations () |
| 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. More... | |
| Pass | UnifiedStaticMemoryPlanner () |
| This is the unified static memory planner pass that will plan for memory intra- and inter- PrimFuncs together. The pass requires all the function to be PrimFuncs including the main. More... | |
| Pass | BindParams (const ffi::Array< runtime::Tensor > &constants) |
| Pass | ExtractPrimFuncConstants () |
| Pass to collect tir non-scalar constants into module's 'Constants' attribute. More... | |
| Pass | RenormalizeSplitPattern () |
| Renormalize the split pattern from floordiv(floormod()) to floormod(floordiv()) More... | |
| Pass | BindTarget (Target target) |
| Annotate a PrimFunc with a given target. More... | |
| Pass | AnnotateEntryFunc () |
| Set a PrimFunc as the entry point if it is only function in IRModule. More... | |
| Pass | Filter (ffi::TypedFunction< bool(PrimFunc)> fcond) |
| Filter PrimFuncs with a given condition. More... | |
| Pass | InjectPTXAsyncCopy () |
| Pass to rewrite global to shared memory copy on CUDA with asyncronous copy. More... | |
| Pass | InjectPTXLDG32 (bool enable_ptx_ldg32=true) |
| Pass to rewrite global to local memory copy on CUDA with ldg32 instruction. More... | |
| Pass | RemoveWeightLayoutRewriteBlock (bool skip_tensor_rewrite=false) |
| Remove the weight layout rewrite block. More... | |
| Pass | InstrumentProfileIntrinsics () |
| Insert intrinsic calls to instrument function and loop level profiling. More... | |
| Pass | DefaultGPUSchedule () |
| The pass sets default thread bindings for PrimFuncs, including symbolic shape functions, allowing their build and execution on GPU devices. It examines all the blocks within the PrimFunc and conducts loop fusion, splitting, and reordering operations based on the loop extent and target information, such as the maximum thread block number and maximum thread per block. More... | |
| Pass | UseAssumeToReduceBranches () |
| This pass analyzes primfunc & eliminates branch introdued due to layout specific padding. It leverages from the buffer assumptions and use the information to eliminate the branch. More... | |
| Pass tvm::tir::transform::AnnotateDeviceRegions | ( | ) |
Annotate locations that should be run on the device.
Insert AttrStmt nodes specifying a target on which regions within the PrimFunc should be executed. Only modifies functions that have a tvm::attr::kTarget attribute, and where that target defines a host.
| Pass tvm::tir::transform::AnnotateEntryFunc | ( | ) |
| Pass tvm::tir::transform::BF16ComputeLegalize | ( | ) |
Legalize bf16 compute Ops. Add a cast to fp32 before Ops, then add a cast back to bf16.
| Pass tvm::tir::transform::BF16StorageLegalize | ( | ) |
Legalize bf16 storage types to u16.
| Pass tvm::tir::transform::BindParams | ( | const ffi::Array< runtime::Tensor > & | constants | ) |
Annotate a PrimFunc with a given target.
| Pass tvm::tir::transform::CombineContextCall | ( | ) |
Combine context calls in the host function.
| Pass tvm::tir::transform::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 duplicated sub-expressions.
| enable_cse_tir | Whether common subexpression elimination is enabled. |
| identify_equiv_terms | Whether equivalent terms should be identified. |
| Pass tvm::tir::transform::ConvertForLoopsToSerial | ( | ) |
| Pass tvm::tir::transform::ConvertSSA | ( | ) |
Convert an IRModule to be SSA form.
This pass handles cases where the same tir::Var appears in multiple functions within the same module. For example, after extracting a fragment from one function into another, where the same tir::Var may be defined both as within the body of the original function, and as a parameter within the hoisted function.
| Pass tvm::tir::transform::CreatePrimFuncPass | ( | std::function< PrimFunc(PrimFunc, IRModule, PassContext)> | pass_func, |
| int | opt_level, | ||
| ffi::String | name, | ||
| tvm::ffi::Array< ffi::String > | required, | ||
| bool | traceable = false |
||
| ) |
| Pass tvm::tir::transform::DecorateDeviceScope | ( | ) |
Decorate all the function's body as device function.
| Pass tvm::tir::transform::DefaultGPUSchedule | ( | ) |
The pass sets default thread bindings for PrimFuncs, including symbolic shape functions, allowing their build and execution on GPU devices. It examines all the blocks within the PrimFunc and conducts loop fusion, splitting, and reordering operations based on the loop extent and target information, such as the maximum thread block number and maximum thread per block.
| Pass tvm::tir::transform::ExtractPrimFuncConstants | ( | ) |
Pass to collect tir non-scalar constants into module's 'Constants' attribute.
Filter PrimFuncs with a given condition.
| Pass tvm::tir::transform::FlattenBuffer | ( | ) |
Flatten the multi-dimensional BufferLoad and BufferStore to single dimensional BufferLoad/BufferStore for the TIR not contains opaque block.
| Pass tvm::tir::transform::ForceNarrowIndexToInt32 | ( | ) |
Force to narrow down indexing expressions and integer buffers to int32 dtype.
| Pass tvm::tir::transform::FP8ComputeLegalize | ( | ffi::String | promote_dtype = "float16" | ) |
Legalize fp8 compute Ops. Add a cast to fp16/fp32 before Ops, then add a cast back to fp8.
| promote_dtype | The data type used for type promotion, defaults to float16 |
| Pass tvm::tir::transform::FP8StorageLegalize | ( | ) |
Legalize fp8 storage types to u8.
| Pass tvm::tir::transform::HoistExpression | ( | ) |
Hoist loop-invariant expressions nodes to outside the elligible loops.
Can hoist conditionals used in IfThenElse statements and expressions, bindings of variables in Let statements and expressions, or boolean expressions, configurable to enable/disable each hoistable type.
| Pass tvm::tir::transform::HoistIfThenElse | ( | ) |
Hoist loop-invariant IfThenElse nodes to outside the elligible loops.
| Pass tvm::tir::transform::InferFragment | ( | ) |
Infer the TensorCore fragment infomation using tensor intrinsics.
| Pass tvm::tir::transform::InjectPTXAsyncCopy | ( | ) |
Pass to rewrite global to shared memory copy on CUDA with asyncronous copy.
| Pass tvm::tir::transform::InjectPTXLDG32 | ( | bool | enable_ptx_ldg32 = true | ) |
Pass to rewrite global to local memory copy on CUDA with ldg32 instruction.
| Pass tvm::tir::transform::InlinePrivateFunctions | ( | ) |
Inline calls to private functions.
| Pass tvm::tir::transform::InstrumentBoundCheckers | ( | ) |
Instruments bound checkers.
| Pass tvm::tir::transform::InstrumentProfileIntrinsics | ( | ) |
Insert intrinsic calls to instrument function and loop level profiling.
| Pass tvm::tir::transform::LowerAsyncDMA | ( | ) |
Lower Async TIR primitives to DMA copy and wait builtins.
| Pass tvm::tir::transform::LowerCustomDatatypes | ( | ) |
Lower custom datatypes.
See tvm::datatypes::Registry for more information on adding custom datatypes.
| Pass tvm::tir::transform::LowerDeviceKernelLaunch | ( | ) |
Lower cross-device function calls.
Prior to this pass, host to device calls are represented as subroutine calls, with environment parameters (e.g. env_thread) specified internally. The device function is an internal function, without a tvm::attr::kGlobalSymbol attribute.
After this pass, host to device calls are represented as tvm_call_packed built-in. The device function is an externally-exposed function, with a non-empty tvm::attr::kGlobalSymbol attribute.
| Pass tvm::tir::transform::LowerDeviceStorageAccessInfo | ( | ) |
Lower attached storage access information on device.
| Pass tvm::tir::transform::LowerIntrin | ( | ) |
Lower the target specific function intrinsics in each of the function.
| Pass tvm::tir::transform::LowerThreadAllreduce | ( | ) |
Lower cross thread alleduce.
| Pass tvm::tir::transform::LowerTVMBuiltin | ( | ) |
Lower builtin intrinsics.
| Pass tvm::tir::transform::LowerVtcmAlloc | ( | ) |
| Pass tvm::tir::transform::LowerWarpMemory | ( | ) |
Lower warp memory access to low-level device related function calls.
| Pass tvm::tir::transform::MakePackedAPI | ( | ) |
Transform the high-level PrimFunc to a low-level version that can be used as an API function.
The main task of this function is to create code to :
let num_packed_args = len(api_args);
if num_packed_args is zero: f()
if num_packed_args is not zero: f(void *, TVMFFIAny* packed_args, int num_packed_args, api_arg_k, api_arg_k+1, ... api_arg_n, TVMFFIAny* out_ret_val)
where n == len(api_args), k == num_packed_args
| Pass tvm::tir::transform::MakeUnpackedAPI | ( | ) |
| Pass tvm::tir::transform::MergeSharedMemoryAllocations | ( | ) |
A pass to merge multiple TIR-level shared memory allocations into one
| Pass tvm::tir::transform::NarrowDataType | ( | int | target_bits | ) |
Narrow down PrimExpr datatype in stmt to target_bits.
| target_bits | The target bits |
| Pass tvm::tir::transform::OOBChecker | ( | ) |
Statically check TIR code for out of bounds array access.
This analysis is conservative: it will only raise errors if it can prove that out of bounds access occurs. Cases that are uncertain do not raise errors.
| Pass tvm::tir::transform::PointerValueTypeRewrite | ( | ) |
Rewrite the pointer content type of arguments, as well as Alloc internal to the function to use the most frequently accessed type for load/store to avoid pointer casting in backend when possible.
Remap the thread axis.
This can be used to get equivalent program which uses threadIdx.y in place of threadIdx.x by passing {"threadIdx.x": thread_axis("threadIdx.y")}
| Pass tvm::tir::transform::RemoveWeightLayoutRewriteBlock | ( | bool | skip_tensor_rewrite = false | ) |
Remove the weight layout rewrite block.
| skip_tensor_rewrite | If True, exact rewrite of Tensor, according to the given index map, will be skipped. Only the shape of the Tensor is transformed correctly, and the content of the destination array will be filled with random values. |
When this pass is called many times during MetaSchedule tuning, the raw data of Tensor, before and after rewrite, does not matter. Since Tensor layout rewrite, using IndexMap's MapTensor, is currently slow, skipping the exact rewrite is sometimes necessary.
| Pass tvm::tir::transform::RenormalizeSplitPattern | ( | ) |
Renormalize the split pattern from floordiv(floormod()) to floormod(floordiv())
| Pass tvm::tir::transform::RewriteUnsafeSelect | ( | ) |
Detect and rewrite unsafe select that contains memory access.
| Pass tvm::tir::transform::Simplify | ( | ) |
Run arithmetic simplifications on the statements and expressions.
| Pass tvm::tir::transform::SkipAssert | ( | ) |
skip assert stmt.
| Pass tvm::tir::transform::SplitHostDevice | ( | ) |
Split the function into a host function and device functions.
The resulting host-side function will keep the same tvm::attr::kTarget attribute (e.g. T.target("cuda", host=T.target("llvm"))). This ensures that MakePackedAPI knows which device type should be used for the input buffers.
The resulting device-side function will have the host stripped from its target attribute (e.g. T.target("cuda")).
| Pass tvm::tir::transform::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 tvm::tir::transform::ThreadSync | ( | ffi::String | storage_scope | ) |
Insert sync between parallel read/write of shared buffers.
| storage_scope | The storage scope considered. |
| Pass tvm::tir::transform::UnifiedStaticMemoryPlanner | ( | ) |
This is the unified static memory planner pass that will plan for memory intra- and inter- PrimFuncs together. The pass requires all the function to be PrimFuncs including the main.
| Pass tvm::tir::transform::UnrollLoop | ( | ) |
unroll the constant loop marked by unroll. This pass also automatically attach pragma unroll tag to loops which meets the standard.
| Pass tvm::tir::transform::UseAssumeToReduceBranches | ( | ) |
This pass analyzes primfunc & eliminates branch introdued due to layout specific padding. It leverages from the buffer assumptions and use the information to eliminate the branch.
| Pass tvm::tir::transform::VectorizeLoop | ( | bool | enable_vectorize = true | ) |
Lower vectorization loops.
| enable_vectorize | Whether vectorization is enabled. |
Pass variant of VerifyGPUCode.
| constraints | The dict to specify constraints to check. |
| Pass tvm::tir::transform::VerifyMemory | ( | ) |
| Pass tvm::tir::transform::VerifySSA | ( | ) |
Pass to checks if the size of the allocated vtcm memory satisfies the limit.
| target | The target whose VTCM limit should be used for any functions not already annotated with tvm::attr::kTarget. |
1.9.1