tvm
Loading...
Searching...
No Matches
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/prim/expr.h>
28#include <tvm/ir/transform.h>
29#include <tvm/target/target.h>
30#include <tvm/tirx/transform.h>
31
32#include <string>
33#include <vector>
34
35namespace tvm {
36namespace s_tir {
37using namespace tvm::prim;
38
47
48namespace transform {
49
53
59
66
72
80
88
94
134
140
146
152
158
170
208
214
220
223
230
237
244
253TVM_DLL Pass HoistIfThenElse(tvm::ffi::String variant = "");
254
266
272
278
284
291
297
303
309TVM_DLL Pass ThreadSync(tvm::ffi::String storage_scope);
310
316
322
328
334
340
346
353
359
365
371
372} // namespace transform
373} // namespace s_tir
374} // namespace tvm
375
376#endif // TVM_S_TIR_TRANSFORM_H_
RAII wrapper function to enter and exit a context object similar to python's with syntax.
Definition with_context.h:59
Managed reference to PrimFuncNode.
Definition function.h:105
PassContext that is used to configure the pass behavior.
Definition transform.h:151
Definition transform.h:417
TIR expressions.
Definition builtin.h:25
Pass HoistIfThenElse(tvm::ffi::String variant="")
Hoist loop-invariant IfThenElse nodes to outside the eligible loops.
Pass DecorateDeviceScope()
Decorate all the function's body as device function.
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 LowerVtcmAlloc()
Lower VTCM allocations.
Pass LowerInitBlock()
Lower block init stmt into IfThenElse stmts.
Pass UseAssumeToReduceBranches()
Eliminate branches by leveraging buffer assumptions (T.assume).
Pass InjectVirtualThread()
Inject virtual thread loops.
Pass HoistExpression()
Hoist loop-invariant expressions to outside the eligible 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 RemoveStoreUndef()
Remove stores of tirx::builtin::undef.
Pass InjectDoubleBuffer()
Inject double buffer statements.
Pass TransformMmaBufferLayout()
Transform Mma scope (m16n8k8.matrixA/B/C) to local scope with layout transformation.
Pass InjectPTXAsyncCopy()
Rewrite global to shared memory copy on CUDA with asynchronous copy.
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 InstrumentProfileIntrinsics()
Insert intrinsic calls to instrument function and loop level profiling.
Pass LowerOpaqueBlock()
Remove the block to ensure that the TIR can not be scheduled again.
Pass LowerAsyncDMA()
Lower Async TIR primitives to DMA copy and wait builtins.
Pass InjectPTXLDG32(bool enable_inject=true)
Rewrite global to local memory copy on CUDA with ldg32 instruction.
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 LowerThreadAllreduce()
Lower cross thread allreduce.
Pass RenormalizeSplitPattern()
Renormalize the split pattern from floordiv(floormod()) to floormod(floordiv()).
Pass ManifestSharedMemoryLocalStage()
Add the explicit local stage for the shared memory access on GPU.
Pass InstrumentBoundCheckers()
Instruments bound checkers.
Pass LiftThreadBinding()
Lift the same thread bindings to their LCA loops.
Pass DefaultGPUSchedule()
Set default thread bindings for GPU PrimFuncs.
Pass RemoveWeightLayoutRewriteBlock(bool skip_tensor_rewrite=false)
Remove weight layout rewrite block before benchmark.
Pass RewriteUnsafeSelect()
Detect and rewrite unsafe select that contains memory access.
Pass UnifyThreadBinding()
Unify all the thread bindings for "blockIdx.x/y/z", "threadIdx.x/y/z", and "vthread....
Pass MergeSharedMemoryAllocations()
Merge multiple TIR-level shared memory allocations into one.
Pass ThreadSync(tvm::ffi::String storage_scope)
Insert sync between parallel read/write of shared buffers.
Pass InferFragment()
Infer the TensorCore fragment information using tensor intrinsics.
tirx::PrimFunc RenewDefs(const tirx::PrimFunc &func)
Renew the definition nodes for a TIR, including Var, Buffer and IterVar. This pass works as a simple ...
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)
An object that builds and maintains block scope and StmtSref mapping for Dependence analysis.
Definition analyzer.h:40
Compilation target object.
TIR specific transformation passes.