target.cpp 2.53 KB
Newer Older
Paul's avatar
Paul committed
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
#include <migraphx/gpu/target.hpp>
#include <migraphx/gpu/lowering.hpp>
#include <migraphx/memory_coloring.hpp>
#include <migraphx/gpu/write_literals.hpp>
#include <migraphx/gpu/context.hpp>
#include <migraphx/gpu/eliminate_workspace.hpp>
#include <migraphx/eliminate_allocation.hpp>
#include <migraphx/gpu/fuse_ops.hpp>
#include <migraphx/check_context.hpp>
#include <migraphx/auto_contiguous.hpp>
#include <migraphx/dead_code_elimination.hpp>
#include <migraphx/simplify_reshapes.hpp>
#include <migraphx/simplify_algebra.hpp>
#include <migraphx/constant_propagate.hpp>
#include <migraphx/eliminate_contiguous.hpp>
#include <migraphx/common_subexpression_elimination.hpp>
#include <migraphx/fwd_conv_batchnorm_rewrite.hpp>
Shucai Xiao's avatar
Shucai Xiao committed
18
#include <migraphx/rewrite_rnn.hpp>
Paul's avatar
Paul committed
19
#include <migraphx/eliminate_concat.hpp>
20
#include <migraphx/eliminate_identity.hpp>
Paul's avatar
Paul committed
21
#include <migraphx/gpu/concat_gpu_opt.hpp>
Paul's avatar
Paul committed
22
#include <migraphx/gpu/schedule_model.hpp>
23
#include <migraphx/pad_rewrite.hpp>
Paul's avatar
Paul committed
24
#include <migraphx/schedule.hpp>
Paul's avatar
Paul committed
25

Paul's avatar
Paul committed
26
namespace migraphx {
Paul's avatar
Paul committed
27
inline namespace MIGRAPHX_INLINE_NS {
Paul's avatar
Paul committed
28
namespace gpu {
mei-ye's avatar
mei-ye committed
29

Paul's avatar
Paul committed
30
std::vector<pass> target::get_passes(migraphx::context& gctx) const
Paul's avatar
Paul committed
31
{
Paul's avatar
Paul committed
32
    auto& ctx = any_cast<context>(gctx);
Paul's avatar
Paul committed
33
34
35
    // clang-format off
    return
    {
Paul's avatar
Paul committed
36
        dead_code_elimination{},
37
        eliminate_identity{},
38
39
        pad_rewrite{},
        dead_code_elimination{},
Paul's avatar
Paul committed
40
        fwd_conv_batchnorm_rewrite{},
Paul's avatar
Paul committed
41
        dead_code_elimination{},
Shucai Xiao's avatar
Shucai Xiao committed
42
43
        rewrite_rnn{},
        dead_code_elimination{},
44
45
        //common_subexpression_elimination{},
        //dead_code_elimination{},
Paul's avatar
Paul committed
46
47
48
49
        simplify_algebra{},
        dead_code_elimination{},
        constant_propagate{},
        dead_code_elimination{},
Paul's avatar
Paul committed
50
        auto_contiguous{},
51
        //simplify_reshapes{},
Paul's avatar
Paul committed
52
        dead_code_elimination{},
Paul's avatar
Paul committed
53
        lowering{ctx},
wsttiger's avatar
wsttiger committed
54
55
        eliminate_concat{concat_gpu_optimization{}},
        dead_code_elimination{},
56
        eliminate_contiguous{},
Paul's avatar
Paul committed
57
        dead_code_elimination{},
Paul's avatar
Paul committed
58
59
        fuse_ops{&ctx},
        dead_code_elimination{},
Paul's avatar
Paul committed
60
        write_literals{&ctx},
Paul's avatar
Paul committed
61
        schedule{gpu::schedule_model{ctx.get_current_device().nstreams()}},
Paul's avatar
Paul committed
62
        memory_coloring{"hip::allocate"},
Paul's avatar
Paul committed
63
        dead_code_elimination{},
Paul's avatar
Paul committed
64
65
        eliminate_workspace{},
        eliminate_allocation{"hip::allocate"},
66
        check_context<context>{},
67
68
        dead_code_elimination{},
        eliminate_identity{}
Paul's avatar
Paul committed
69
70
    };
    // clang-format on
Paul's avatar
Paul committed
71
}
Paul's avatar
Paul committed
72

Paul's avatar
Paul committed
73
std::string target::name() const { return "miopen"; }
Paul's avatar
Paul committed
74

Paul's avatar
Paul committed
75
migraphx::context target::get_context() const { return context{}; }
Paul's avatar
Paul committed
76
} // namespace gpu
Paul's avatar
Paul committed
77
} // namespace MIGRAPHX_INLINE_NS
Paul's avatar
Paul committed
78
} // namespace migraphx