conv_miopen_helper.h 8.2 KB
Newer Older
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25
/* Copyright (c) 2020 PaddlePaddle Authors. All Rights Reserved.

Licensed under the Apache License, Version 2.0 (the "License");
you may not use this file except in compliance with the License.
You may obtain a copy of the License at

    http://www.apache.org/licenses/LICENSE-2.0

Unless required by applicable law or agreed to in writing, software
distributed under the License is distributed on an "AS IS" BASIS,
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and
limitations under the License. */

#pragma once

#include <algorithm>
#include <array>
#include <memory>
#include <string>
#include <vector>

#include "paddle/fluid/framework/conv_search_cache.h"
#include "paddle/fluid/framework/operator_kernel_configs.h"
#include "paddle/fluid/operators/conv_cudnn_op_cache.h"
26
#include "paddle/fluid/platform/device/gpu/gpu_dnn.h"
H
hong 已提交
27
#include "paddle/phi/backends/gpu/gpu_context.h"
28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54

namespace paddle {
namespace operators {

using Tensor = framework::Tensor;
using DataLayout = platform::DataLayout;
template <typename T>
using ScalingParamType = typename platform::CudnnDataType<T>::ScalingParamType;
using framework::AlgorithmsCache;
static inline void GetNCDHW(const framework::DDim& dims,
                            const DataLayout& layout, int* N, int* C, int* D,
                            int* H, int* W) {
  *N = dims[0];
  *C = layout == DataLayout::kNCHW ? dims[1] : dims[dims.size() - 1];
  int i = layout == DataLayout::kNCHW ? 0 : 1;
  if (dims.size() == 5) {
    *D = dims[2 - i];
    *H = dims[3 - i];
    *W = dims[4 - i];
  } else {
    *D = 1;
    *H = dims[2 - i];
    *W = dims[3 - i];
  }
}

template <typename DeviceContext, typename T, size_t D>
H
hong 已提交
55
static void RemovePaddingSlice(const phi::GPUContext& context,
56 57 58
                               const Tensor* input, Tensor* out,
                               const std::vector<int>& starts,
                               const std::vector<int>& axes) {
H
hong 已提交
59
  auto& place = *context.eigen_device();
60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90 91 92 93 94 95 96 97 98 99 100 101 102 103 104 105 106 107 108 109 110 111 112 113 114 115 116 117 118 119 120 121 122 123 124 125 126 127 128 129
  auto in_dims = input->dims();
  auto new_out_dims = out->dims();
  auto offsets = Eigen::array<int, D>();
  auto extents = Eigen::array<int, D>();
  for (size_t i = 0; i < D; ++i) {
    offsets[i] = 0;
    extents[i] = new_out_dims[i];
  }

  int start;
  for (size_t i = 0; i < axes.size(); ++i) {
    start = starts[i];
    if (start < 0) {
      start = (start + in_dims[axes[i]]);
    }
    start = std::max(start, 0);
    offsets[axes[i]] = start;
  }
  auto in_t =
      framework::EigenTensor<T, D, Eigen::RowMajor, Eigen::DenseIndex>::From(
          *input);

  auto out_t =
      framework::EigenTensor<T, D, Eigen::RowMajor, Eigen::DenseIndex>::From(
          *out, new_out_dims);
  out_t.device(place) = in_t.slice(offsets, extents);
}

template <typename T>
std::ostream& operator<<(std::ostream& out, const std::vector<T>& v) {
  out << "[";
  for (auto const& tmp : v) out << tmp << ",";
  out << "]";
  return out;
}

using framework::ConvSearchCache;

struct ConvArgs {
  miopenHandle_t handle;
  platform::TensorDescriptor idesc, odesc;
  platform::FilterDescriptor wdesc;
  platform::ConvolutionDescriptor cdesc;
  const framework::Tensor *x, *w, *o;
  miopenDataType_t cudnn_dtype;

  // strides
  std::vector<int> s;
  // paddings
  std::vector<int> p;
  // dilations
  std::vector<int> d;

  ConvArgs(const framework::Tensor* x, const framework::Tensor* w,
           const framework::Tensor* o, const std::vector<int> s,
           const std::vector<int> p, const std::vector<int> d,
           miopenDataType_t dtype)
      : x(x), w(w), o(o), s(s), p(p), d(d), cudnn_dtype(dtype) {}
};

template <typename algo_t>
struct SearchAlgorithm {};

template <>
struct SearchAlgorithm<miopenConvFwdAlgorithm_t> {
  using perf_t = miopenConvAlgoPerf_t;
  using algo_t = miopenConvFwdAlgorithm_t;

  template <typename T>
  static algo_t Find(const ConvArgs& args, bool exhaustive_search,
130
                     bool deterministic, size_t workspace_size,
H
hong 已提交
131
                     const phi::GPUContext& ctx) {
132 133
    algo_t algo;

H
hong 已提交
134
    auto workspace_handle = ctx.cudnn_workspace_handle();
135

136 137 138
    int find_count;
    miopenConvAlgoPerf_t find_result;
    auto cudnn_find_func = [&](void* cudnn_workspace_ptr) {
139
      PADDLE_ENFORCE_GPU_SUCCESS(
140 141 142 143 144 145 146 147
          platform::dynload::miopenFindConvolutionForwardAlgorithm(
              args.handle, args.idesc.desc(), args.x->data<T>(),
              args.wdesc.desc(), args.w->data<T>(), args.cdesc.desc(),
              args.odesc.desc(), const_cast<T*>(args.o->data<T>()),
              kNUM_CUDNN_FWD_ALGS, &find_count, &find_result,
              cudnn_workspace_ptr, workspace_size, false));
    };

R
ronnywang 已提交
148 149
    workspace_handle.RunFuncSync(cudnn_find_func, workspace_size);
    algo = find_result.fwd_algo;
150 151 152 153
    VLOG(3) << "choose algo " << algo;
    return algo;
  }

154
  static size_t GetWorkspaceSize(const ConvArgs& args) {
155
    size_t workspace_size = 0;
156
    PADDLE_ENFORCE_GPU_SUCCESS(
157 158 159 160 161 162 163 164 165 166 167 168 169 170
        platform::dynload::miopenConvolutionForwardGetWorkSpaceSize(
            args.handle, args.wdesc.desc(), args.idesc.desc(),
            args.cdesc.desc(), args.odesc.desc(), &workspace_size));
    return workspace_size;
  }
};

template <>
struct SearchAlgorithm<miopenConvBwdDataAlgorithm_t> {
  using perf_t = miopenConvAlgoPerf_t;
  using algo_t = miopenConvBwdDataAlgorithm_t;

  template <typename T>
  static algo_t Find(const ConvArgs& args, bool exhaustive_search,
171
                     bool deterministic, size_t workspace_size,
H
hong 已提交
172
                     const phi::GPUContext& ctx) {
173 174
    algo_t algo;

H
hong 已提交
175
    auto workspace_handle = ctx.cudnn_workspace_handle();
176

177 178 179
    int find_count;
    miopenConvAlgoPerf_t find_result;
    auto cudnn_find_func = [&](void* cudnn_workspace_ptr) {
180
      PADDLE_ENFORCE_GPU_SUCCESS(
181 182 183 184 185 186 187 188
          platform::dynload::miopenFindConvolutionBackwardDataAlgorithm(
              args.handle, args.odesc.desc(), args.o->data<T>(),
              args.wdesc.desc(), args.w->data<T>(), args.cdesc.desc(),
              args.idesc.desc(), const_cast<T*>(args.x->data<T>()),
              kNUM_CUDNN_BWD_DATA_ALGS, &find_count, &find_result,
              cudnn_workspace_ptr, workspace_size, false));
    };

R
ronnywang 已提交
189 190
    workspace_handle.RunFuncSync(cudnn_find_func, workspace_size);
    algo = find_result.bwd_data_algo;
191 192 193 194
    VLOG(3) << "choose algo " << algo;
    return algo;
  }

195
  static size_t GetWorkspaceSize(const ConvArgs& args) {
196
    size_t workspace_size = 0;
197
    PADDLE_ENFORCE_GPU_SUCCESS(
198 199 200 201 202 203 204 205 206 207 208 209 210 211
        platform::dynload::miopenConvolutionBackwardDataGetWorkSpaceSize(
            args.handle, args.odesc.desc(), args.wdesc.desc(),
            args.cdesc.desc(), args.idesc.desc(), &workspace_size));
    return workspace_size;
  }
};

template <>
struct SearchAlgorithm<miopenConvBwdWeightsAlgorithm_t> {
  using perf_t = miopenConvAlgoPerf_t;
  using algo_t = miopenConvBwdWeightsAlgorithm_t;

  template <typename T>
  static algo_t Find(const ConvArgs& args, bool exhaustive_search,
212
                     bool deterministic, size_t workspace_size,
H
hong 已提交
213
                     const phi::GPUContext& ctx) {
214 215
    algo_t algo;

H
hong 已提交
216
    auto workspace_handle = ctx.cudnn_workspace_handle();
217 218 219 220

    int find_count;
    miopenConvAlgoPerf_t find_result;
    auto cudnn_find_func = [&](void* cudnn_workspace_ptr) {
221
      PADDLE_ENFORCE_GPU_SUCCESS(
222 223 224 225 226 227 228 229
          platform::dynload::miopenFindConvolutionBackwardWeightsAlgorithm(
              args.handle, args.odesc.desc(), args.o->data<T>(),
              args.idesc.desc(), args.x->data<T>(), args.cdesc.desc(),
              args.wdesc.desc(), const_cast<T*>(args.w->data<T>()),
              kNUM_CUDNN_BWD_FILTER_ALGS, &find_count, &find_result,
              cudnn_workspace_ptr, workspace_size, false));
    };

R
ronnywang 已提交
230 231
    workspace_handle.RunFuncSync(cudnn_find_func, workspace_size);
    algo = find_result.bwd_weights_algo;
232 233 234 235
    VLOG(3) << "choose algo " << algo;
    return algo;
  }

236
  static size_t GetWorkspaceSize(const ConvArgs& args) {
237
    size_t workspace_size = 0;
238
    PADDLE_ENFORCE_GPU_SUCCESS(
239 240 241 242 243 244 245 246 247
        platform::dynload::miopenConvolutionBackwardWeightsGetWorkSpaceSize(
            args.handle, args.odesc.desc(), args.idesc.desc(),
            args.cdesc.desc(), args.wdesc.desc(), &workspace_size));
    return workspace_size;
  }
};

}  // namespace operators
}  // namespace paddle