tvm
Namespaces | Functions
transform.h File Reference

S-TIR specific transformation passes. More...

#include <tvm/ir/transform.h>
#include <tvm/target/target.h>
#include <tvm/tir/transform.h>
#include <string>
#include <vector>
Include dependency graph for transform.h:

Go to the source code of this file.

Namespaces

 tvm
 Performance counters for profiling via the PAPI library.
 
 tvm::s_tir
 
 tvm::s_tir::transform
 

Functions

Pass tvm::s_tir::transform::CanonicalizeLoop ()
 Canonicalize loop to start from zero . More...
 
Pass tvm::s_tir::transform::LowerCrossThreadReduction ()
 Lower cross-thread reduction from thread bindings to intrinsic function calls. More...
 
Pass tvm::s_tir::transform::LowerInitBlock ()
 Lower block init stmt into IfThenElse stmts. More...
 
Pass tvm::s_tir::transform::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. More...
 
Pass tvm::s_tir::transform::ConvertBlocksToOpaque ()
 Substitute all the block vars with the PrimExprs they are bound to, indicated by the corresponding iter_values in BlockRealize, for opaque blocks by removing all . the iter_values in BlockRealize and iter_vars in Block. More...
 
Pass tvm::s_tir::transform::LiftThreadBinding ()
 Lift the same thread bindings to their LCA loops. More...
 
Pass tvm::s_tir::transform::CompactBufferAllocation (bool is_strict=true)
 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. More...
 
Pass tvm::s_tir::transform::LowerMatchBuffer ()
 Remove match buffers inside the block. Also, it will validate the binding. More...
 
Pass tvm::s_tir::transform::InjectPermutedLayout ()
 Inject permuted layout for shared memory. More...
 
Pass tvm::s_tir::transform::TransformMmaBufferLayout ()
 Transform Mma scope (m16n8k8.matrixA/B/C) to local scope with layout transformation. More...
 
Pass tvm::s_tir::transform::LowerOpaqueBlock ()
 Remove the block to ensure that the TIR can not be scheduled again. More...
 
Pass tvm::s_tir::transform::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. More...
 
Pass tvm::s_tir::transform::InjectSoftwarePipeline ()
 This pass transforms annotated loops into pipelined ones where producers and consumers are overlapped with the information provided in loop annotations, which enables optimization techniques like prefetching and pipeline parallelism. More...
 
Pass tvm::s_tir::transform::LowerAutoCopy ()
 Automatically do memory optimizations for auto copy blocks. More...
 
Pass tvm::s_tir::transform::ManifestSharedMemoryLocalStage ()
 Add the explicit local stage for the shared memory access on GPU. More...
 
Pass tvm::s_tir::transform::AnnotateIrregularLoop ()
 Annotate irregular loop mark. More...
 
Pass tvm::s_tir::transform::LoopPartition ()
 partition loops in the stmt. More...
 
Pass tvm::s_tir::transform::InjectVirtualThread ()
 Inject virtual thread loops. More...
 
Pass tvm::s_tir::transform::InjectDoubleBuffer ()
 Inject double buffer statements. More...
 

Detailed Description

S-TIR specific transformation passes.