From 7deef68ab2a0fa633589777dcb7e982b32000ee8 Mon Sep 17 00:00:00 2001
From: zhangting2020 <709968123@qq.com>
Date: Tue, 22 Sep 2020 05:58:25 +0000
Subject: [PATCH] use exhaustive_search for float16

---
 paddle/fluid/operators/conv_cudnn_helper.h | 32 ++++++++++++++++++++--
 python/paddle/fluid/layers/nn.py           |  2 +-
 2 files changed, 31 insertions(+), 3 deletions(-)

diff --git a/paddle/fluid/operators/conv_cudnn_helper.h b/paddle/fluid/operators/conv_cudnn_helper.h
index c2d91b31c74..fd892094a2f 100644
--- a/paddle/fluid/operators/conv_cudnn_helper.h
+++ b/paddle/fluid/operators/conv_cudnn_helper.h
@@ -443,7 +443,8 @@ struct SearchAlgorithm<cudnnConvolutionBwdFilterAlgoPerf_t> {
                      bool deterministic,
                      const framework::ExecutionContext& ctx) {
     auto dtype = platform::CudnnDataType<T>::type;
-    bool exhaustive = (exhaustive_search) & (dtype != CUDNN_DATA_HALF);
+    // bool exhaustive = (exhaustive_search) & (dtype != CUDNN_DATA_HALF);
+    bool exhaustive = exhaustive_search;
     size_t workspace_size_limit = FLAGS_conv_workspace_size_limit * 1024 * 1024;
     size_t workspace_size = 0;
     bool has_got_workspace_size = true;
@@ -539,7 +540,7 @@ struct SearchAlgorithm<cudnnConvolutionBwdFilterAlgoPerf_t> {
       VLOG(10) << "cudnnConvolutionFwdAlgoPerf_t:"
                << ", x_dims:" << x_dims << ", w_dims:" << w_dims << ", args.s"
                << args.s << ", args.p" << args.p << ", args.d" << args.d;
-
+      /*
       algo = algo_cache.GetAlgorithm(
           x_dims, w_dims, args.s, args.p, args.d, 0,
           static_cast<int64_t>(args.cudnn_dtype), [&]() {
@@ -567,7 +568,34 @@ struct SearchAlgorithm<cudnnConvolutionBwdFilterAlgoPerf_t> {
             }
             return perf_stat[0].algo;
           });
+      */
+      algo = algo_cache.GetAlgorithm(
+          x_dims, w_dims, args.s, args.p, args.d, 0,
+          static_cast<int64_t>(args.cudnn_dtype), [&]() {
+            algo_t sel_algo;
+            auto max_bwd_filt_algos = MaxBackwardFilterAlgos(args.handle);
+            std::vector<cudnnConvolutionBwdFilterAlgoPerf_t> bwd_filt_results(
+                max_bwd_filt_algos);
+            int actual_bwd_filter_algos = 0;
+            PADDLE_ENFORCE_CUDA_SUCCESS(
+                platform::dynload::cudnnFindConvolutionBackwardFilterAlgorithm(
+                    args.handle, args.idesc.desc(), args.odesc.desc(),
+                    args.cdesc.desc(), args.wdesc.desc(),
+                    bwd_filt_results.size(), &actual_bwd_filter_algos,
+                    bwd_filt_results.data()));
+            bwd_filt_results.resize(actual_bwd_filter_algos);
+            AlgoFinalSelect<cudnnConvolutionBwdFilterAlgoPerf_t,
+                            cudnnConvolutionBwdFilterAlgo_t>(
+                bwd_filt_results, "backprop-to-filter", -1,
+                workspace_size_limit, &sel_algo, deterministic);
+            workspace_size = GetWorkspaceSize(args, sel_algo);
+            if (workspace_size > workspace_size_limit) {
+              workspace_size = workspace_size_limit;
+            }
+            return sel_algo;
+          });
     }
+
     VLOG(3) << "choose algo " << algo;
     return algo;
   }
diff --git a/python/paddle/fluid/layers/nn.py b/python/paddle/fluid/layers/nn.py
index bc9f182d95e..75d11648d20 100755
--- a/python/paddle/fluid/layers/nn.py
+++ b/python/paddle/fluid/layers/nn.py
@@ -1582,7 +1582,7 @@ def conv2d(input,
             'use_mkldnn': False,
             'fuse_relu_before_depthwise_conv': False,
             "padding_algorithm": padding_algorithm,
-            "data_format": data_format,
+            "data_format": data_format
         })
 
     if data_format == 'NCHW':
-- 
GitLab