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 cuda/injective.h |
22 | * \brief CUDA schedule for injective operations |
23 | */ |
24 | #ifndef TVM_TOPI_CUDA_INJECTIVE_H_ |
25 | #define TVM_TOPI_CUDA_INJECTIVE_H_ |
26 | |
27 | #include <tvm/target/generic_func.h> |
28 | #include <tvm/te/operation.h> |
29 | #include <tvm/te/schedule_pass.h> |
30 | #include <tvm/topi/detail/fuse.h> |
31 | #include <tvm/topi/tags.h> |
32 | |
33 | namespace tvm { |
34 | namespace topi { |
35 | |
36 | using namespace tvm::te; |
37 | |
38 | namespace cuda { |
39 | |
40 | /*! |
41 | * \brief Updates an existing schedule for the given injective ops. |
42 | * |
43 | * \param sch The schedule to update. |
44 | * \param out The tensor representing the injective op. |
45 | * |
46 | * \return The updated schedule. |
47 | */ |
48 | inline Schedule schedule_injective_from_existing(Schedule sch, const Tensor& out) { |
49 | auto fused = detail::Fuse(sch[out], sch[out]->op.as<ComputeOpNode>()->axis); |
50 | auto target = Target::Current(false); |
51 | int num_thread = target->GetAttr<Integer>("max_num_threads" ).value().IntValue(); |
52 | IterVar bx, tx; |
53 | sch[out].split(fused, num_thread, &bx, &tx); |
54 | sch[out].bind(bx, thread_axis(Range(), "blockIdx.x" )); |
55 | sch[out].bind(tx, thread_axis(Range(), "threadIdx.x" )); |
56 | return sch; |
57 | } |
58 | |
59 | /*! |
60 | * \brief Create a CUDA schedule for the given output tensors. |
61 | * |
62 | * \param target The target to generate a schedule for. |
63 | * \param outs The output tensors. |
64 | * |
65 | * \return A schedule for the given ops. |
66 | */ |
67 | inline Schedule schedule_injective(const Target& target, const Array<Tensor>& outs) { |
68 | Array<Operation> out_ops; |
69 | for (auto t : outs) { |
70 | out_ops.push_back(t->op); |
71 | } |
72 | auto s = create_schedule(out_ops); |
73 | tvm::te::AutoInlineInjective(s); |
74 | for (auto out : outs) { |
75 | schedule_injective_from_existing(s, out); |
76 | } |
77 | return s; |
78 | } |
79 | |
80 | } // namespace cuda |
81 | } // namespace topi |
82 | } // namespace tvm |
83 | #endif // TVM_TOPI_CUDA_INJECTIVE_H_ |
84 | |