conv_cudnn_helper.h 25.4 KB
Newer Older
Q
qingqing01 已提交
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16
/* Copyright (c) 2019 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

17
#include <algorithm>
18
#include <array>
19
#include <memory>
20
#include <string>
Q
qingqing01 已提交
21
#include <vector>
22

23
#include "paddle/fluid/framework/conv_search_cache.h"
Q
qingqing01 已提交
24 25 26 27 28 29
#include "paddle/fluid/framework/operator_kernel_configs.h"
#include "paddle/fluid/operators/conv_cudnn_op_cache.h"
#include "paddle/fluid/platform/cudnn_desc.h"
namespace paddle {
namespace operators {

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 55 56 57 58 59 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
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>
static void RemovePaddingSlice(const framework::ExecutionContext& context,
                               const Tensor* input, Tensor* out,
                               const std::vector<int>& starts,
                               const std::vector<int>& axes) {
  auto& place =
      *context.template device_context<DeviceContext>().eigen_device();
  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);
}

87 88 89 90 91 92 93 94
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;
}

95 96 97 98 99 100 101 102 103 104
inline int MaxBwdFilterAlgos(cudnnHandle_t cudnn_handle) {
  int max_algos = 0;
#if CUDNN_VERSION_MIN(7, 0, 1)
  PADDLE_ENFORCE_CUDA_SUCCESS(
      platform::dynload::cudnnGetConvolutionBackwardFilterAlgorithmMaxCount(
          cudnn_handle, &max_algos));
#endif
  return max_algos;
}

105 106 107 108 109 110 111 112 113 114 115 116 117 118 119 120 121 122
template <typename PerfType, typename AlgoType>
void ChooseAlgoByWorkspace(PerfType* perf_results, size_t perf_num,
                           size_t workspace_byte, AlgoType* algo) {
  for (size_t i = 0; i < perf_num; ++i) {
    auto result = perf_results[i];
    if (result.status == CUDNN_STATUS_SUCCESS &&
        result.memory < workspace_byte) {
      *algo = result.algo;
      VLOG(3) << "    algo: " << result.algo << ", time: " << result.time
              << " ms, wksp = " << result.memory
              << ", status = " << result.status;
      return;
    }
  }
  VLOG(3) << "Can not find alog that requires memory < "
          << static_cast<double>(workspace_byte) / (1 << 20) << " MB";
}

123 124 125 126 127 128 129 130 131 132 133 134 135 136 137 138 139 140 141 142 143 144 145 146 147 148 149 150 151 152 153 154 155 156 157 158 159 160 161 162 163 164 165 166 167
template <typename PerfType, typename AlgoType>
void ChooseAlgo(const std::vector<PerfType>& perf_results,
                size_t workspace_byte, AlgoType* algo) {
  VLOG(3) << "=========BwdFilterAlgo Perf result=========";
  for (const auto& result : perf_results) {
    auto math_type_str = "False";
    if (result.mathType == CUDNN_TENSOR_OP_MATH) {
      math_type_str = "True";
    }
    VLOG(3) << "    algo: " << result.algo << ", TensorCore: " << math_type_str
            << ", time: " << result.time << " ms"
            << ", wksp = " << result.memory << ", status = " << result.status;
  }

  for (size_t i = 0; i != perf_results.size(); ++i) {
    const auto& result = perf_results[i];
    if (result.status == CUDNN_STATUS_SUCCESS &&
        (result.memory <= workspace_byte)) {
      if ((result.mathType == CUDNN_TENSOR_OP_MATH) &&
          (i != perf_results.size() - 1)) {
        const auto& next_result = perf_results[i + 1];
        if (next_result.status == CUDNN_STATUS_SUCCESS &&
            next_result.algo == result.algo &&
            next_result.memory == result.memory &&
            next_result.mathType != CUDNN_TENSOR_OP_MATH &&
            next_result.time < 1.01 * result.time) {
          // Skip over this result- it's not really a Tensor Core algo.
          // Because it is only 1% performance difference.
          // Prefer to choose the next equivalent non-Tensor Core algo.
          continue;
        }
      }
      *algo = result.algo;
      auto math_type_str = "0";
      if (result.mathType == CUDNN_TENSOR_OP_MATH) {
        math_type_str = "1";
      }
      VLOG(3) << "    choose algo: " << result.algo << ", TC: " << math_type_str
              << ", time: " << result.time << " ms"
              << ", wksp = " << result.memory << ", status = " << result.status;
      return;
    }
  }
}

168
using framework::ConvSearchCache;
Q
qingqing01 已提交
169 170 171 172 173 174 175

struct ConvArgs {
  cudnnHandle_t handle;
  platform::TensorDescriptor idesc, odesc;
  platform::FilterDescriptor wdesc;
  platform::ConvolutionDescriptor cdesc;
  const framework::Tensor *x, *w, *o;
176
  cudnnDataType_t cudnn_dtype;
Q
qingqing01 已提交
177 178 179 180 181 182 183 184 185 186

  // 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,
187 188 189
           const std::vector<int> p, const std::vector<int> d,
           cudnnDataType_t dtype)
      : x(x), w(w), o(o), s(s), p(p), d(d), cudnn_dtype(dtype) {}
Q
qingqing01 已提交
190 191 192 193 194 195 196 197 198 199 200 201
};

template <typename perf_t>
struct SearchAlgorithm {};

template <>
struct SearchAlgorithm<cudnnConvolutionFwdAlgoPerf_t> {
  using perf_t = cudnnConvolutionFwdAlgoPerf_t;
  using algo_t = cudnnConvolutionFwdAlgo_t;

  template <typename T>
  static algo_t Find(const ConvArgs& args, bool exhaustive_search,
202
                     bool deterministic,
Q
qingqing01 已提交
203 204
                     const framework::ExecutionContext& ctx) {
    auto dtype = platform::CudnnDataType<T>::type;
205
    bool has_got_workspace_size = true;
Q
qingqing01 已提交
206 207
    bool exhaustive = (exhaustive_search) & (dtype != CUDNN_DATA_HALF);
    size_t workspace_size_limit = FLAGS_conv_workspace_size_limit * 1024 * 1024;
208
    size_t workspace_size = 0;
Q
qingqing01 已提交
209
    algo_t algo;
210 211 212 213

#if CUDA_VERSION >= 9000 && CUDNN_VERSION_MIN(7, 0, 1)
    auto& dev_ctx = ctx.template device_context<platform::CUDADeviceContext>();
    if (dev_ctx.GetComputeCapability() >= 70 && dtype == CUDNN_DATA_HALF) {
214 215 216
      PADDLE_ENFORCE_CUDA_SUCCESS(
          platform::dynload::cudnnSetConvolutionMathType(args.cdesc.desc(),
                                                         CUDNN_TENSOR_OP_MATH));
217 218
      VLOG(5) << "use cudnn_tensor_op_math";
    } else {
219 220 221
      PADDLE_ENFORCE_CUDA_SUCCESS(
          platform::dynload::cudnnSetConvolutionMathType(args.cdesc.desc(),
                                                         CUDNN_DEFAULT_MATH));
222 223 224 225
      VLOG(5) << "NOT use cudnn_tensor_op_math";
    }
#endif

226
    if (!exhaustive && !deterministic) {
227 228 229 230
#if CUDNN_VERSION >= 7001
      int perf_count;
      int best_algo_idx = 0;
      std::unique_ptr<perf_t[]> perf_results(new perf_t[kNUM_CUDNN_FWD_ALGS]);
231 232 233 234 235
      PADDLE_ENFORCE_CUDA_SUCCESS(
          platform::dynload::cudnnGetConvolutionForwardAlgorithm_v7(
              args.handle, args.idesc.desc(), args.wdesc.desc(),
              args.cdesc.desc(), args.odesc.desc(), kNUM_CUDNN_FWD_ALGS,
              &perf_count, perf_results.get()));
236 237 238 239
      algo = (perf_results.get())[best_algo_idx].algo;
      workspace_size = GetWorkspaceSize(args, algo);

      if (workspace_size > workspace_size_limit) {
240
#if CUDNN_VERSION >= 8000
241 242 243 244
        // cudnnGetConvolutionForwardAlgorithm is removed in CUDNN-8
        ChooseAlgoByWorkspace<perf_t, algo_t>(perf_results.get(),
                                              kNUM_CUDNN_FWD_ALGS,
                                              workspace_size_limit, &algo);
245 246 247 248 249 250 251 252 253 254 255 256
#else
        VLOG(1) << "Fallback to non-v7 method to find conv algorithm becasue "
                   "the workspace size request("
                << workspace_size << ") exceeds the limit("
                << workspace_size_limit << ")";
        PADDLE_ENFORCE_CUDA_SUCCESS(
            platform::dynload::cudnnGetConvolutionForwardAlgorithm(
                args.handle, args.idesc.desc(), args.wdesc.desc(),
                args.cdesc.desc(), args.odesc.desc(),
                CUDNN_CONVOLUTION_FWD_SPECIFY_WORKSPACE_LIMIT,
                workspace_size_limit, &algo));
#endif
257 258
      }
#else
259 260 261 262 263 264
      PADDLE_ENFORCE_CUDA_SUCCESS(
          platform::dynload::cudnnGetConvolutionForwardAlgorithm(
              args.handle, args.idesc.desc(), args.wdesc.desc(),
              args.cdesc.desc(), args.odesc.desc(),
              CUDNN_CONVOLUTION_FWD_SPECIFY_WORKSPACE_LIMIT,
              workspace_size_limit, &algo));
265
#endif
Q
qingqing01 已提交
266
      VLOG(3) << "choose algo " << algo;
267 268
    } else if (deterministic) {
      algo = static_cast<cudnnConvolutionFwdAlgo_t>(1);
Q
qingqing01 已提交
269 270 271 272 273
    } else {
      auto& dev_ctx =
          ctx.template device_context<platform::CUDADeviceContext>();
      auto workspace_handle = dev_ctx.cudnn_workspace_handle();

274 275
      auto& temp = ctx.cuda_device_context();
      AlgorithmsCache<algo_t>& algo_cache =
276
          *(framework::ConvSearchCache::Instance().GetForward());
277

Q
qingqing01 已提交
278 279 280
      auto x_dims = framework::vectorize(args.x->dims());
      auto w_dims = framework::vectorize(args.w->dims());

281 282 283
      VLOG(10) << "cudnnConvolutionFwdAlgoPerf_t:"
               << ", x_dims:" << x_dims << ", w_dims:" << w_dims << ", args.s"
               << args.s << ", args.p" << args.p << ", args.d" << args.d;
284

Q
qingqing01 已提交
285
      algo = algo_cache.GetAlgorithm(
286 287
          x_dims, w_dims, args.s, args.p, args.d, 0,
          static_cast<int64_t>(args.cudnn_dtype), [&]() {
Q
qingqing01 已提交
288 289 290 291
            int returned_algo_count;
            std::array<perf_t, kNUM_CUDNN_FWD_ALGS> perf_stat;

            auto cudnn_find_func = [&](void* cudnn_workspace_ptr) {
292
              PADDLE_ENFORCE_CUDA_SUCCESS(
Q
qingqing01 已提交
293 294 295 296 297 298 299 300 301 302 303 304 305 306 307 308 309 310 311 312 313 314 315 316 317
                  platform::dynload::cudnnFindConvolutionForwardAlgorithmEx(
                      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, &returned_algo_count,
                      perf_stat.data(), cudnn_workspace_ptr,
                      workspace_size_limit));
            };
            workspace_handle.RunFuncSync(cudnn_find_func, workspace_size_limit);

            VLOG(3) << "FwdAlgo Perf result: (algo: stat, time, memory)";
            for (int i = 0; i < returned_algo_count; ++i) {
              const auto& stat = perf_stat[i];
              VLOG(3) << stat.algo << ": " << stat.status << " " << stat.time
                      << " " << stat.memory;
            }
            return perf_stat[0].algo;
          });
    }
    VLOG(3) << "choose algo " << algo;
    return algo;
  }

  static size_t GetWorkspaceSize(const ConvArgs& args, algo_t algo) {
    size_t workspace_size = 0;
318 319 320 321
    PADDLE_ENFORCE_CUDA_SUCCESS(
        platform::dynload::cudnnGetConvolutionForwardWorkspaceSize(
            args.handle, args.idesc.desc(), args.wdesc.desc(),
            args.cdesc.desc(), args.odesc.desc(), algo, &workspace_size));
Q
qingqing01 已提交
322 323 324 325 326 327 328 329 330 331 332
    return workspace_size;
  }
};

template <>
struct SearchAlgorithm<cudnnConvolutionBwdDataAlgoPerf_t> {
  using perf_t = cudnnConvolutionBwdDataAlgoPerf_t;
  using algo_t = cudnnConvolutionBwdDataAlgo_t;

  template <typename T>
  static algo_t Find(const ConvArgs& args, bool exhaustive_search,
333
                     bool deterministic,
Q
qingqing01 已提交
334 335 336 337
                     const framework::ExecutionContext& ctx) {
    auto dtype = platform::CudnnDataType<T>::type;
    bool exhaustive = (exhaustive_search) & (dtype != CUDNN_DATA_HALF);
    size_t workspace_size_limit = FLAGS_conv_workspace_size_limit * 1024 * 1024;
338 339
    size_t workspace_size = 0;
    bool has_got_workspace_size = true;
Q
qingqing01 已提交
340
    algo_t algo;
341 342 343
#if CUDA_VERSION >= 9000 && CUDNN_VERSION_MIN(7, 0, 1)
    auto& dev_ctx = ctx.template device_context<platform::CUDADeviceContext>();
    if (dev_ctx.GetComputeCapability() >= 70 && dtype == CUDNN_DATA_HALF) {
344 345 346
      PADDLE_ENFORCE_CUDA_SUCCESS(
          platform::dynload::cudnnSetConvolutionMathType(args.cdesc.desc(),
                                                         CUDNN_TENSOR_OP_MATH));
347 348
      VLOG(5) << "use cudnn_tensor_op_math";
    } else {
349 350 351
      PADDLE_ENFORCE_CUDA_SUCCESS(
          platform::dynload::cudnnSetConvolutionMathType(args.cdesc.desc(),
                                                         CUDNN_DEFAULT_MATH));
352 353 354 355
      VLOG(5) << "NOT use cudnn_tensor_op_math";
    }
#endif

Q
qingqing01 已提交
356
    if (!exhaustive && !deterministic) {
357 358 359 360 361
#if CUDNN_VERSION >= 7001
      int perf_count;
      int best_algo_idx = 0;
      std::unique_ptr<perf_t[]> perf_results(
          new perf_t[kNUM_CUDNN_BWD_DATA_ALGS]);
362
      PADDLE_ENFORCE_CUDA_SUCCESS(
363 364 365 366 367 368 369 370 371 372 373 374 375 376 377 378 379 380 381 382 383 384
          platform::dynload::cudnnGetConvolutionBackwardDataAlgorithm_v7(
              args.handle, args.wdesc.desc(), args.odesc.desc(),
              args.cdesc.desc(), args.idesc.desc(), kNUM_CUDNN_BWD_DATA_ALGS,
              &perf_count, perf_results.get()));
      algo = (perf_results.get())[best_algo_idx].algo;

#if CUDNN_VERSION < 7500
      int stride_dim = args.x->dims().size() - 2;
      bool blacklist = std::any_of(args.s.begin(), args.s.begin() + stride_dim,
                                   [=](int n) { return n != 1; });
      if (blacklist && (static_cast<cudnnConvolutionBwdDataAlgo_t>(
                            perf_results[best_algo_idx].algo) ==
                            CUDNN_CONVOLUTION_BWD_DATA_ALGO_FFT_TILING ||
                        static_cast<cudnnConvolutionBwdDataAlgo_t>(
                            perf_results[best_algo_idx].algo) ==
                            CUDNN_CONVOLUTION_BWD_DATA_ALGO_FFT)) {
        algo = CUDNN_CONVOLUTION_BWD_DATA_ALGO_1;
      }
#endif
      workspace_size = GetWorkspaceSize(args, algo);
      if (workspace_size > workspace_size_limit) {
        has_got_workspace_size = false;
385
#if CUDNN_VERSION >= 8000
386 387 388 389
        // cudnnGetConvolutionBackwardDataAlgorithm is removed in CUDNN-8
        ChooseAlgoByWorkspace<perf_t, algo_t>(perf_results.get(),
                                              kNUM_CUDNN_BWD_DATA_ALGS,
                                              workspace_size_limit, &algo);
390 391 392 393 394 395 396 397 398 399 400 401
#else
        VLOG(1) << "Fallback to non-v7 method to find conv algorithm becasue "
                   "the workspace size request("
                << workspace_size << ") exceeds the limit("
                << workspace_size_limit << ")";
        PADDLE_ENFORCE_CUDA_SUCCESS(
            platform::dynload::cudnnGetConvolutionBackwardDataAlgorithm(
                args.handle, args.wdesc.desc(), args.odesc.desc(),
                args.cdesc.desc(), args.idesc.desc(),
                CUDNN_CONVOLUTION_BWD_DATA_SPECIFY_WORKSPACE_LIMIT,
                workspace_size_limit, &algo));
#endif
402 403
      }
#else
404 405 406 407 408 409
      PADDLE_ENFORCE_CUDA_SUCCESS(
          platform::dynload::cudnnGetConvolutionBackwardDataAlgorithm(
              args.handle, args.wdesc.desc(), args.odesc.desc(),
              args.cdesc.desc(), args.idesc.desc(),
              CUDNN_CONVOLUTION_BWD_DATA_SPECIFY_WORKSPACE_LIMIT,
              workspace_size_limit, &algo));
410
#endif
Q
qingqing01 已提交
411 412 413 414 415 416 417
    } else if (deterministic) {
      return CUDNN_CONVOLUTION_BWD_DATA_ALGO_1;
    } else {
      auto& dev_ctx =
          ctx.template device_context<platform::CUDADeviceContext>();
      auto workspace_handle = dev_ctx.cudnn_workspace_handle();

418
      AlgorithmsCache<algo_t>& algo_cache =
419
          *(framework::ConvSearchCache::Instance().GetBackwardData());
420

Q
qingqing01 已提交
421 422 423
      auto x_dims = framework::vectorize(args.x->dims());
      auto w_dims = framework::vectorize(args.w->dims());

424 425 426
      VLOG(10) << "cudnnConvolutionFwdAlgoPerf_t"
               << ", x_dims:" << x_dims << ", w_dims:" << w_dims << ", args.s"
               << args.s << ", args.p" << args.p << ", args.d" << args.d;
427

Q
qingqing01 已提交
428
      algo = algo_cache.GetAlgorithm(
429 430
          x_dims, w_dims, args.s, args.p, args.d, 0,
          static_cast<int64_t>(args.cudnn_dtype), [&]() {
Q
qingqing01 已提交
431
            int returned_algo_count;
432
            std::array<perf_t, kNUM_CUDNN_BWD_DATA_ALGS> perf_stat;
Q
qingqing01 已提交
433 434

            auto cudnn_find_func = [&](void* cudnn_workspace_ptr) {
435
              PADDLE_ENFORCE_CUDA_SUCCESS(
Q
qingqing01 已提交
436 437 438 439 440 441 442 443 444 445 446 447 448 449 450 451 452 453 454 455 456 457 458 459 460 461 462 463
                  platform::dynload::
                      cudnnFindConvolutionBackwardDataAlgorithmEx(
                          args.handle, args.wdesc.desc(), args.w->data<T>(),
                          args.odesc.desc(), args.o->data<T>(),
                          args.cdesc.desc(), args.idesc.desc(),
                          const_cast<T*>(args.x->data<T>()),
                          kNUM_CUDNN_BWD_DATA_ALGS, &returned_algo_count,
                          perf_stat.data(), cudnn_workspace_ptr,
                          workspace_size_limit));
            };
            workspace_handle.RunFuncSync(cudnn_find_func, workspace_size_limit);

            VLOG(3) << "BwdDataAlgo Perf result: (algo: stat, time, memory)";
            for (int i = 0; i < returned_algo_count; ++i) {
              const auto& stat = perf_stat[i];
              VLOG(3) << stat.algo << ": " << stat.status << " " << stat.time
                      << " " << stat.memory;
            }

            return perf_stat[0].algo;
          });
    }
    VLOG(3) << "choose algo " << algo;
    return algo;
  }

  static size_t GetWorkspaceSize(const ConvArgs& args, algo_t algo) {
    size_t workspace_size = 0;
464
    PADDLE_ENFORCE_CUDA_SUCCESS(
Q
qingqing01 已提交
465
        platform::dynload::cudnnGetConvolutionBackwardDataWorkspaceSize(
466 467
            args.handle, args.wdesc.desc(), args.odesc.desc(),
            args.cdesc.desc(), args.idesc.desc(), algo, &workspace_size));
Q
qingqing01 已提交
468 469 470 471 472 473 474 475 476 477 478
    return workspace_size;
  }
};

template <>
struct SearchAlgorithm<cudnnConvolutionBwdFilterAlgoPerf_t> {
  using perf_t = cudnnConvolutionBwdFilterAlgoPerf_t;
  using algo_t = cudnnConvolutionBwdFilterAlgo_t;

  template <typename T>
  static algo_t Find(const ConvArgs& args, bool exhaustive_search,
479
                     bool deterministic,
Q
qingqing01 已提交
480 481 482
                     const framework::ExecutionContext& ctx) {
    auto dtype = platform::CudnnDataType<T>::type;
    size_t workspace_size_limit = FLAGS_conv_workspace_size_limit * 1024 * 1024;
483 484 485 486 487 488
    size_t workspace_size = 0;
    bool has_got_workspace_size = true;

#if CUDA_VERSION >= 9000 && CUDNN_VERSION_MIN(7, 0, 1)
    auto& dev_ctx = ctx.template device_context<platform::CUDADeviceContext>();
    if (dev_ctx.GetComputeCapability() >= 70 && dtype == CUDNN_DATA_HALF) {
489 490 491
      PADDLE_ENFORCE_CUDA_SUCCESS(
          platform::dynload::cudnnSetConvolutionMathType(args.cdesc.desc(),
                                                         CUDNN_TENSOR_OP_MATH));
492 493
      VLOG(5) << "use cudnn_tensor_op_math";
    } else {
494 495 496
      PADDLE_ENFORCE_CUDA_SUCCESS(
          platform::dynload::cudnnSetConvolutionMathType(args.cdesc.desc(),
                                                         CUDNN_DEFAULT_MATH));
497 498 499
      VLOG(5) << "NOT use cudnn_tensor_op_math";
    }
#endif
Q
qingqing01 已提交
500 501

    algo_t algo;
502
    if (!exhaustive_search && !deterministic) {
503 504 505 506 507 508
#if CUDNN_VERSION >= 7001
      using perf_t = cudnnConvolutionBwdFilterAlgoPerf_t;
      int perf_count;
      int best_algo_idx = 0;
      std::unique_ptr<perf_t[]> perf_results(
          new perf_t[kNUM_CUDNN_BWD_FILTER_ALGS]);
509
      PADDLE_ENFORCE_CUDA_SUCCESS(
510 511 512 513 514 515 516
          platform::dynload::cudnnGetConvolutionBackwardFilterAlgorithm_v7(
              args.handle, args.idesc.desc(), args.odesc.desc(),
              args.cdesc.desc(), args.wdesc.desc(), kNUM_CUDNN_BWD_FILTER_ALGS,
              &perf_count, perf_results.get()));
      algo = (perf_results.get())[best_algo_idx].algo;
      workspace_size = GetWorkspaceSize(args, algo);
      if (workspace_size > workspace_size_limit) {
517
        workspace_size = workspace_size_limit;
518 519 520 521 522 523 524 525 526 527 528 529 530 531 532 533 534
#if CUDNN_VERSION >= 8000
        // cudnnGetConvolutionBackwardFilterAlgorithm is removed in CUDNN-8
        ChooseAlgoByWorkspace<perf_t, algo_t>(perf_results.get(),
                                              kNUM_CUDNN_BWD_FILTER_ALGS,
                                              workspace_size_limit, &algo);
#else
        VLOG(1) << "Fallback to non-v7 method to find conv algorithm becasue "
                   "the workspace size request("
                << workspace_size << ") exceeds the limit("
                << workspace_size_limit << ")";
        PADDLE_ENFORCE_CUDA_SUCCESS(
            platform::dynload::cudnnGetConvolutionBackwardFilterAlgorithm(
                args.handle, args.idesc.desc(), args.odesc.desc(),
                args.cdesc.desc(), args.wdesc.desc(),
                CUDNN_CONVOLUTION_BWD_FILTER_SPECIFY_WORKSPACE_LIMIT,
                workspace_size_limit, &algo));
#endif
535 536
      }
#else
537
      PADDLE_ENFORCE_CUDA_SUCCESS(
Q
qingqing01 已提交
538 539 540 541 542
          platform::dynload::cudnnGetConvolutionBackwardFilterAlgorithm(
              args.handle, args.idesc.desc(), args.odesc.desc(),
              args.cdesc.desc(), args.wdesc.desc(),
              CUDNN_CONVOLUTION_BWD_FILTER_SPECIFY_WORKSPACE_LIMIT,
              workspace_size_limit, &algo));
543
#endif
Q
qingqing01 已提交
544 545 546 547 548 549
    } else if (deterministic) {
      return CUDNN_CONVOLUTION_BWD_FILTER_ALGO_1;
    } else {
      auto& dev_ctx =
          ctx.template device_context<platform::CUDADeviceContext>();
      auto workspace_handle = dev_ctx.cudnn_workspace_handle();
550
      AlgorithmsCache<algo_t>& algo_cache =
551
          *(framework::ConvSearchCache::Instance().GetBackwardFilter());
Q
qingqing01 已提交
552 553 554 555

      auto x_dims = framework::vectorize(args.x->dims());
      auto w_dims = framework::vectorize(args.w->dims());

556 557 558
      VLOG(10) << "cudnnConvolutionFwdAlgoPerf_t:"
               << ", x_dims:" << x_dims << ", w_dims:" << w_dims << ", args.s"
               << args.s << ", args.p" << args.p << ", args.d" << args.d;
559 560 561 562 563
      if (dtype != CUDNN_DATA_HALF) {
        algo = algo_cache.GetAlgorithm(
            x_dims, w_dims, args.s, args.p, args.d, 0,
            static_cast<int64_t>(args.cudnn_dtype), [&]() {
              int returned_algo_count;
564
              std::array<perf_t, kNUM_CUDNN_BWD_FILTER_ALGS> perf_stat;
565 566 567 568 569 570 571 572 573 574 575 576 577 578 579 580 581 582 583 584 585 586 587 588 589 590 591 592 593 594 595 596
              auto cudnn_find_func = [&](void* cudnn_workspace_ptr) {
                PADDLE_ENFORCE_CUDA_SUCCESS(
                    platform::dynload::
                        cudnnFindConvolutionBackwardFilterAlgorithmEx(
                            args.handle, args.idesc.desc(), args.x->data<T>(),
                            args.odesc.desc(), args.o->data<T>(),
                            args.cdesc.desc(), args.wdesc.desc(),
                            const_cast<T*>(args.w->data<T>()),
                            kNUM_CUDNN_BWD_FILTER_ALGS, &returned_algo_count,
                            perf_stat.data(), cudnn_workspace_ptr,
                            workspace_size_limit));
              };
              workspace_handle.RunFuncSync(cudnn_find_func,
                                           workspace_size_limit);

              VLOG(3)
                  << "BwdFilterAlgo Perf result: (algo: stat, time, memory)";
              for (int i = 0; i < returned_algo_count; ++i) {
                const auto& stat = perf_stat[i];
                VLOG(3) << stat.algo << ": " << stat.status << " " << stat.time
                        << " " << stat.memory;
              }
              return perf_stat[0].algo;
            });
      } else {
        auto max_algos = MaxBwdFilterAlgos(args.handle);
        algo = algo_cache.GetAlgorithm(
            x_dims, w_dims, args.s, args.p, args.d, 0,
            static_cast<int64_t>(args.cudnn_dtype), [&]() {
              algo_t chosen_algo;
              std::vector<perf_t> perf_results(max_algos);
              int actual_algos = 0;
597
              PADDLE_ENFORCE_CUDA_SUCCESS(
Q
qingqing01 已提交
598
                  platform::dynload::
599 600
                      cudnnFindConvolutionBackwardFilterAlgorithm(
                          args.handle, args.idesc.desc(), args.odesc.desc(),
Q
qingqing01 已提交
601
                          args.cdesc.desc(), args.wdesc.desc(),
602 603 604 605 606 607 608 609
                          perf_results.size(), &actual_algos,
                          perf_results.data()));
              perf_results.resize(actual_algos);
              ChooseAlgo<perf_t, algo_t>(perf_results, workspace_size_limit,
                                         &chosen_algo);
              return chosen_algo;
            });
      }
Q
qingqing01 已提交
610 611 612 613 614 615 616
    }
    VLOG(3) << "choose algo " << algo;
    return algo;
  }

  static size_t GetWorkspaceSize(const ConvArgs& args, algo_t algo) {
    size_t workspace_size = 0;
617
    PADDLE_ENFORCE_CUDA_SUCCESS(
Q
qingqing01 已提交
618 619 620 621 622 623 624 625 626
        platform::dynload::cudnnGetConvolutionBackwardFilterWorkspaceSize(
            args.handle, args.idesc.desc(), args.odesc.desc(),
            args.cdesc.desc(), args.wdesc.desc(), algo, &workspace_size));
    return workspace_size;
  }
};

}  // namespace operators
}  // namespace paddle