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