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
33namespace tvm {
34namespace topi {
35
36using namespace tvm::te;
37
38namespace 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 */
48inline 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 */
67inline 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