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 | |
20 | /*! |
21 | * \file tir/transforms/convert_for_loops_serial.cc |
22 | * \brief Convert all for loops to serial for lesser memory consumption |
23 | */ |
24 | #include <tvm/arith/analyzer.h> |
25 | #include <tvm/runtime/device_api.h> |
26 | #include <tvm/tir/function.h> |
27 | #include <tvm/tir/stmt_functor.h> |
28 | #include <tvm/tir/transform.h> |
29 | |
30 | namespace tvm { |
31 | namespace tir { |
32 | |
33 | class ForLoopSerialConverter : public StmtExprMutator { |
34 | public: |
35 | ForLoopSerialConverter() = default; |
36 | Stmt operator()(const PrimFunc& func); |
37 | |
38 | private: |
39 | Stmt VisitStmt_(const ForNode* op) override; |
40 | }; |
41 | |
42 | Stmt ForLoopSerialConverter::VisitStmt_(const ForNode* op) { |
43 | if (op->kind == ForKind::kParallel) { |
44 | return For(op->loop_var, op->min, op->extent, ForKind::kSerial, op->body, op->thread_binding, |
45 | op->annotations, op->span); |
46 | } |
47 | return StmtExprMutator::VisitStmt_(op); |
48 | } |
49 | |
50 | Stmt ForLoopSerialConverter::operator()(const PrimFunc& func) { |
51 | return this->VisitStmt(func->body); |
52 | } |
53 | |
54 | PrimFunc ConvertForLoopsToSerial(PrimFunc func) { |
55 | PrimFuncNode* fptr = func.CopyOnWrite(); |
56 | fptr->body = ForLoopSerialConverter()(func); |
57 | return func; |
58 | } |
59 | |
60 | namespace transform { |
61 | |
62 | Pass ConvertForLoopsToSerial() { |
63 | auto pass_func = [=](PrimFunc f, IRModule m, PassContext ctx) { |
64 | return ConvertForLoopsToSerial(std::move(f)); |
65 | }; |
66 | return CreatePrimFuncPass(pass_func, 0, "tir.ConvertForLoopsToSerial" , {}); |
67 | } |
68 | |
69 | TVM_REGISTER_GLOBAL("tir.transform.ConvertForLoopsToSerial" ) |
70 | .set_body_typed(ConvertForLoopsToSerial); |
71 | |
72 | } // namespace transform |
73 | |
74 | } // namespace tir |
75 | } // namespace tvm |
76 | |