Skip to content

P
Paddle

PaddlePaddle / Paddle
大约 2 年 前同步成功

通知 2325
Star 20933
Fork 5424
P
Paddle
“dbe4e6721044548268218ebfcd188ca87fb840c3”上不存在“python/examples/pipeline/PaddleClas/MobileNetV2/benchmark.py”
未验证 提交 b937cdc5 编写于 作者: L limingshu 提交者： GitHub
浏览文件
操作

Autotune the workspace_size_limit in conv. (#40338) 

* Using the maximum workspace_size of all alogirhms to limit the workspace size in exhaustive search mode.

* Use the system cudaMalloc and cudaFree to allocate workspace during searching.

* Enable switch of two kind of workspace setting methods.
Co-authored-by: NLiu Yiqun <liuyiqun01@baidu.com>
上级 e1792a31
隐藏空白更改
内联 并排
Showing with 620 addition and 513 deletion
paddle/fluid/framework/conv_search_cache.h
浏览文件 @ b937cdc5
...@@ -16,7 +16,6 @@ limitations under the License. */ ...@@ -16,7 +16,6 @@ limitations under the License. */
#include "paddle/fluid/framework/op_registry.h" #include "paddle/fluid/framework/op_registry.h"
#include "paddle/fluid/framework/operator_kernel_configs.h" #include "paddle/fluid/framework/operator_kernel_configs.h"
#include "paddle/fluid/platform/device/gpu/gpu_dnn.h" #include "paddle/fluid/platform/device/gpu/gpu_dnn.h"
namespace paddle { namespace paddle {
......
paddle/fluid/operators/conv_base_helper.h 0 → 100644
浏览文件 @ b937cdc5
/* Copyright (c) 2022 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/operators/conv_cudnn_op_cache.h"
#include "paddle/phi/backends/gpu/gpu_context.h"
namespace paddle {
namespace operators {
using Tensor = framework::Tensor;
using DataLayout = platform::DataLayout;
using framework::AlgorithmsCache;
using framework::ConvSearchCache;
template <typename T>
using ScalingParamType = typename platform::CudnnDataType<T>::ScalingParamType;
// As the basic for SearchAlgorithm struct.
template <typename PerfT>
struct SearchAlgorithm {};
// As the container of searchAlgorithm::Find() result.
template <typename AlgoT>
struct SearchResult {
public:
AlgoT algo = static_cast<AlgoT>(0);
float time = -1.f;
size_t workspace_size = 0;
};
// As the container of conv relevant descriptors.
template <typename HandleT, typename DataT>
struct ConvArgsBase {
HandleT handle;
platform::TensorDescriptor idesc, odesc;
platform::FilterDescriptor wdesc;
platform::ConvolutionDescriptor cdesc;
const framework::Tensor *x, *w, *o;
DataT cudnn_dtype;
// strides
std::vector<int> s;
// paddings
std::vector<int> p;
// dilations
std::vector<int> d;
ConvArgsBase(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, DataT dtype)
: x(x), w(w), o(o), s(s), p(p), d(d), cudnn_dtype(dtype) {}
};
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 T>
static std::ostream& operator<<(std::ostream& out, const std::vector<T>& v) {
out << "[";
for (auto const& tmp : v) out << tmp << ",";
out << "]";
return out;
}
} // namespace operators
} // namespace paddle
paddle/fluid/operators/conv_cudnn_helper.h
浏览文件 @ b937cdc5
...@@ -14,44 +14,15 @@ limitations under the License. */ ...@@ -14,44 +14,15 @@ limitations under the License. */
#pragma once #pragma once
#include <algorithm> #include "paddle/fluid/operators/conv_base_helper.h"
#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"
#include "paddle/fluid/operators/eigen/eigen_function.h"
#include "paddle/fluid/platform/cuda_graph_with_memory_pool.h" #include "paddle/fluid/platform/cuda_graph_with_memory_pool.h"
#include "paddle/fluid/platform/device/gpu/gpu_dnn.h" #include "paddle/fluid/platform/device/gpu/gpu_info.h"
#include "paddle/phi/backends/gpu/gpu_context.h" #include "paddle/phi/kernels/funcs/eigen/eigen_function.h"
namespace paddle { namespace paddle {
namespace operators { namespace operators {
using Tensor = framework::Tensor; using ConvArgs = ConvArgsBase<cudnnHandle_t, cudnnDataType_t>;
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> template <typename DeviceContext, typename T, size_t D>
static void RemovePaddingSlice(const phi::GPUContext& context, static void RemovePaddingSlice(const phi::GPUContext& context,
...@@ -68,121 +39,103 @@ static void RemovePaddingSlice(const phi::GPUContext& context, ...@@ -68,121 +39,103 @@ static void RemovePaddingSlice(const phi::GPUContext& context,
extents[i] = new_out_dims[i]; extents[i] = new_out_dims[i];
} }
int start;
for (size_t i = 0; i < axes.size(); ++i) { for (size_t i = 0; i < axes.size(); ++i) {
start = starts[i]; int start = starts[i];
if (start < 0) { if (start < 0) {
start = (start + in_dims[axes[i]]); start = (start + in_dims[axes[i]]);
} }
start = std::max(start, 0); start = std::max(start, 0);
offsets[axes[i]] = start; offsets[axes[i]] = start;
} }
auto in_t = auto in_t =
framework::EigenTensor<T, D, Eigen::RowMajor, Eigen::DenseIndex>::From( framework::EigenTensor<T, D, Eigen::RowMajor, Eigen::DenseIndex>::From(
*input); *input);
auto out_t = auto out_t =
framework::EigenTensor<T, D, Eigen::RowMajor, Eigen::DenseIndex>::From( framework::EigenTensor<T, D, Eigen::RowMajor, Eigen::DenseIndex>::From(
*out, new_out_dims); *out, new_out_dims);
EigenSlice<std::decay_t<decltype(place)>, T, D>::Eval(place, out_t, in_t,
offsets, extents); phi::funcs::EigenSlice<std::decay_t<decltype(place)>, T, D>::Eval(
place, out_t, in_t, offsets, extents);
} }
template <typename T> static inline double ToMegaBytes(size_t bytes) {
std::ostream& operator<<(std::ostream& out, const std::vector<T>& v) { return static_cast<double>(bytes) / (1 << 20);
out << "[";
for (auto const& tmp : v) out << tmp << ",";
out << "]";
return out;
} }
inline int MaxBwdFilterAlgos(cudnnHandle_t cudnn_handle) { static inline bool UseFixedWorkspace() {
int max_algos = 0; return FLAGS_conv_workspace_size_limit >= 0;
#if CUDNN_VERSION_MIN(7, 0, 1)
PADDLE_ENFORCE_GPU_SUCCESS(
platform::dynload::cudnnGetConvolutionBackwardFilterAlgorithmMaxCount(
cudnn_handle, &max_algos));
#endif
return max_algos;
} }
template <typename PerfType, typename AlgoType> static size_t CaclWorkspaceLimitInBytes(const phi::GPUContext& ctx) {
void ChooseAlgoByWorkspace(PerfType* perf_results, size_t perf_num, if (!UseFixedWorkspace()) {
size_t workspace_byte, AlgoType* algo) { int device_id = platform::GetCurrentDeviceId();
for (size_t i = 0; i < perf_num; ++i) { int64_t allocated = memory::StatGetCurrentValue("Allocated", device_id);
auto result = perf_results[i]; int64_t reserved = memory::StatGetCurrentValue("Reserved", device_id);
if (result.status == CUDNN_STATUS_SUCCESS && int64_t availble = platform::GpuAvailableMemToAlloc();
result.memory < workspace_byte) { int64_t cur_workspace_size = ctx.cudnn_workspace_handle().WorkspaceSize();
*algo = result.algo; VLOG(3) << "[memory] allocated=" << ToMegaBytes(allocated)
VLOG(3) << " algo: " << result.algo << ", time: " << result.time << " MB, reserved=" << ToMegaBytes(reserved)
<< " ms, wksp = " << result.memory << " MB, available_to_alloc=" << ToMegaBytes(availble)
<< ", status = " << result.status; << " MB, current_workspace_size=" << ToMegaBytes(cur_workspace_size)
return; << " MB.";
} return std::max(std::max(availble, cur_workspace_size),
reserved - allocated);
} else {
return FLAGS_conv_workspace_size_limit * 1024 * 1024;
} }
VLOG(3) << "Can not find alog that requires memory < "
<< static_cast<double>(workspace_byte) / (1 << 20) << " MB";
} }
template <typename PerfType, typename AlgoType> template <typename PerfT>
void ChooseAlgo(const std::vector<PerfType>& perf_results, std::string GetPerfResultString(std::string prefix,
size_t workspace_byte, AlgoType* algo) { const std::vector<PerfT>& perf_results,
VLOG(3) << "=========BwdFilterAlgo Perf result========="; int actual_algo_count, size_t workspace_limit) {
for (const auto& result : perf_results) { std::ostringstream out;
auto math_type_str = "False"; out << prefix << " (workspace limit=" << ToMegaBytes(workspace_limit)
if (result.mathType == CUDNN_TENSOR_OP_MATH) { << " MB):\n";
math_type_str = "True"; for (int i = 0; i < actual_algo_count; ++i) {
} const auto& result = perf_results[i];
VLOG(3) << " algo: " << result.algo << ", TensorCore: " << math_type_str auto math_type_str = (result.mathType == CUDNN_TENSOR_OP_MATH) ? "T" : "F";
<< ", time: " << result.time << " ms" out << " algo=" << result.algo << ": tensor_core=" << math_type_str
<< ", wksp = " << result.memory << ", status = " << result.status; << ", time=" << result.time
<< " ms, memory=" << ToMegaBytes(result.memory)
<< " MB, status=" << result.status << "\n";
} }
return out.str();
}
for (size_t i = 0; i != perf_results.size(); ++i) { template <typename PerfT, typename AlgoT>
const auto& result = perf_results[i]; void ChooseAlgoByWorkspace(const std::vector<PerfT>& perf_results,
size_t workspace_limit,
SearchResult<AlgoT>* algo_result) {
for (size_t i = 0; i < perf_results.size(); ++i) {
auto result = perf_results[i];
if (result.status == CUDNN_STATUS_SUCCESS && if (result.status == CUDNN_STATUS_SUCCESS &&
(result.memory <= workspace_byte)) { result.memory < workspace_limit) {
if ((result.mathType == CUDNN_TENSOR_OP_MATH) && algo_result->algo = result.algo;
(i != perf_results.size() - 1)) { algo_result->time = result.time;
const auto& next_result = perf_results[i + 1]; algo_result->workspace_size = result.memory;
if (next_result.status == CUDNN_STATUS_SUCCESS && VLOG(3) << " algo=" << result.algo << ", time=" << result.time
next_result.algo == result.algo && << " ms, memory=" << ToMegaBytes(result.memory)
next_result.memory == result.memory && << " MB (limit=" << ToMegaBytes(workspace_limit)
next_result.mathType != CUDNN_TENSOR_OP_MATH && << " MB), status=" << result.status;
next_result.time < 1.01 * result.time) { return;
// 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;
break;
} }
} }
VLOG(3) << "Can not find an algorithm that requires memory < "
<< ToMegaBytes(workspace_limit) << " MB";
} }
using framework::ConvSearchCache;
static void SetConvMathType(const phi::GPUContext& ctx, cudnnDataType_t dtype, static void SetConvMathType(const phi::GPUContext& ctx, cudnnDataType_t dtype,
const platform::ConvolutionDescriptor& cdesc) { const platform::ConvolutionDescriptor& cdesc) {
#if CUDA_VERSION >= 9000 && CUDNN_VERSION_MIN(7, 0, 1) #if CUDA_VERSION >= 9000 && CUDNN_VERSION_MIN(7, 0, 1)
auto& dev_ctx = ctx; if (ctx.GetComputeCapability() >= 70 && dtype == CUDNN_DATA_HALF) {
if (dev_ctx.GetComputeCapability() >= 70 && dtype == CUDNN_DATA_HALF) {
PADDLE_ENFORCE_GPU_SUCCESS(platform::dynload::cudnnSetConvolutionMathType( PADDLE_ENFORCE_GPU_SUCCESS(platform::dynload::cudnnSetConvolutionMathType(
cdesc.desc(), CUDNN_TENSOR_OP_MATH)); cdesc.desc(), CUDNN_TENSOR_OP_MATH));
VLOG(5) << "use cudnn_tensor_op_math"; VLOG(5) << "use cudnn_tensor_op_math";
#if CUDA_VERSION >= 11000 #if CUDA_VERSION >= 11000
#if CUDNN_VERSION_MIN(8, 1, 0) #if CUDNN_VERSION_MIN(8, 1, 0)
} else if (dev_ctx.GetComputeCapability() >= 80 && } else if (ctx.GetComputeCapability() >= 80 && dtype == CUDNN_DATA_BFLOAT16) {
dtype == CUDNN_DATA_BFLOAT16) {
PADDLE_ENFORCE_GPU_SUCCESS(platform::dynload::cudnnSetConvolutionMathType( PADDLE_ENFORCE_GPU_SUCCESS(platform::dynload::cudnnSetConvolutionMathType(
cdesc.desc(), CUDNN_TENSOR_OP_MATH)); cdesc.desc(), CUDNN_TENSOR_OP_MATH));
#endif // CUDNN_VERSION_MIN(8, 1, 0) #endif // CUDNN_VERSION_MIN(8, 1, 0)
...@@ -198,76 +151,49 @@ static void SetConvMathType(const phi::GPUContext& ctx, cudnnDataType_t dtype, ...@@ -198,76 +151,49 @@ static void SetConvMathType(const phi::GPUContext& ctx, cudnnDataType_t dtype,
#endif #endif
} }
struct ConvArgs {
cudnnHandle_t handle;
platform::TensorDescriptor idesc, odesc;
platform::FilterDescriptor wdesc;
platform::ConvolutionDescriptor cdesc;
const framework::Tensor *x, *w, *o;
cudnnDataType_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,
cudnnDataType_t dtype)
: x(x), w(w), o(o), s(s), p(p), d(d), cudnn_dtype(dtype) {}
};
template <typename perf_t>
struct SearchAlgorithm {};
template <> template <>
struct SearchAlgorithm<cudnnConvolutionFwdAlgoPerf_t> { struct SearchAlgorithm<cudnnConvolutionFwdAlgoPerf_t> {
using perf_t = cudnnConvolutionFwdAlgoPerf_t; using PerfT = cudnnConvolutionFwdAlgoPerf_t;
using algo_t = cudnnConvolutionFwdAlgo_t; using AlgoT = cudnnConvolutionFwdAlgo_t;
template <typename T> template <typename T>
static algo_t Find(const ConvArgs& args, bool exhaustive_search, static SearchResult<AlgoT> Find(const ConvArgs& args, bool exhaustive_search,
bool deterministic, const phi::GPUContext& ctx) { bool deterministic,
const phi::GPUContext& ctx) {
SearchResult<AlgoT> result;
auto dtype = platform::CudnnDataType<T>::type; auto dtype = platform::CudnnDataType<T>::type;
bool has_got_workspace_size = true; size_t workspace_size_limit = CaclWorkspaceLimitInBytes(ctx);
size_t workspace_size_limit = FLAGS_conv_workspace_size_limit * 1024 * 1024;
size_t workspace_size = 0;
algo_t algo;
SetConvMathType(ctx, dtype, args.cdesc); SetConvMathType(ctx, dtype, args.cdesc);
if (!exhaustive_search && !deterministic) { if (!exhaustive_search && !deterministic) {
#if CUDNN_VERSION >= 7001 #if CUDNN_VERSION >= 7001
int perf_count; int actual_perf_count;
int best_algo_idx = 0; int best_algo_idx = 0;
std::unique_ptr<perf_t[]> perf_results(new perf_t[kNUM_CUDNN_FWD_ALGS]); std::vector<PerfT> perf_results(kNUM_CUDNN_FWD_ALGS);
PADDLE_ENFORCE_GPU_SUCCESS( PADDLE_ENFORCE_GPU_SUCCESS(
platform::dynload::cudnnGetConvolutionForwardAlgorithm_v7( platform::dynload::cudnnGetConvolutionForwardAlgorithm_v7(
args.handle, args.idesc.desc(), args.wdesc.desc(), args.handle, args.idesc.desc(), args.wdesc.desc(),
args.cdesc.desc(), args.odesc.desc(), kNUM_CUDNN_FWD_ALGS, args.cdesc.desc(), args.odesc.desc(), kNUM_CUDNN_FWD_ALGS,
&perf_count, perf_results.get())); &actual_perf_count, perf_results.data()));
algo = (perf_results.get())[best_algo_idx].algo; result.algo = perf_results[best_algo_idx].algo;
workspace_size = (perf_results.get())[best_algo_idx].memory; result.workspace_size = perf_results[best_algo_idx].memory;
if (workspace_size > workspace_size_limit) { if (result.workspace_size > workspace_size_limit) {
#if CUDNN_VERSION >= 8000 #if CUDNN_VERSION >= 8000
// cudnnGetConvolutionForwardAlgorithm is removed in CUDNN-8 // cudnnGetConvolutionForwardAlgorithm is removed in CUDNN-8
ChooseAlgoByWorkspace<perf_t, algo_t>(perf_results.get(), ChooseAlgoByWorkspace<PerfT, AlgoT>(perf_results, workspace_size_limit,
kNUM_CUDNN_FWD_ALGS, &result);
workspace_size_limit, &algo);
#else #else
VLOG(1) << "Fallback to non-v7 method to find conv algorithm becasue " VLOG(3) << "Fallback to non-v7 method to find conv algorithm "
"the workspace size request(" "becasue the workspace size request("
<< workspace_size << ") exceeds the limit(" << result.workspace_size << ") exceeds the limit("
<< workspace_size_limit << ")"; << workspace_size_limit << ")";
PADDLE_ENFORCE_GPU_SUCCESS( PADDLE_ENFORCE_GPU_SUCCESS(
platform::dynload::cudnnGetConvolutionForwardAlgorithm( platform::dynload::cudnnGetConvolutionForwardAlgorithm(
args.handle, args.idesc.desc(), args.wdesc.desc(), args.handle, args.idesc.desc(), args.wdesc.desc(),
args.cdesc.desc(), args.odesc.desc(), args.cdesc.desc(), args.odesc.desc(),
CUDNN_CONVOLUTION_FWD_SPECIFY_WORKSPACE_LIMIT, CUDNN_CONVOLUTION_FWD_SPECIFY_WORKSPACE_LIMIT,
workspace_size_limit, &algo)); workspace_size_limit, &(result.algo)));
#endif #endif
} }
#else #else
...@@ -276,30 +202,30 @@ struct SearchAlgorithm<cudnnConvolutionFwdAlgoPerf_t> { ...@@ -276,30 +202,30 @@ struct SearchAlgorithm<cudnnConvolutionFwdAlgoPerf_t> {
args.handle, args.idesc.desc(), args.wdesc.desc(), args.handle, args.idesc.desc(), args.wdesc.desc(),
args.cdesc.desc(), args.odesc.desc(), args.cdesc.desc(), args.odesc.desc(),
CUDNN_CONVOLUTION_FWD_SPECIFY_WORKSPACE_LIMIT, CUDNN_CONVOLUTION_FWD_SPECIFY_WORKSPACE_LIMIT,
workspace_size_limit, &algo)); workspace_size_limit, &(result.algo)));
#endif #endif
VLOG(3) << "choose algo " << algo;
} else if (deterministic) { } else if (deterministic) {
algo = static_cast<cudnnConvolutionFwdAlgo_t>(1); result.algo = static_cast<AlgoT>(1);
} else { } else {
auto& dev_ctx = ctx; auto workspace_handle = ctx.cudnn_workspace_handle();
auto workspace_handle = dev_ctx.cudnn_workspace_handle();
AlgorithmsCache<algo_t>& algo_cache =
*(framework::ConvSearchCache::Instance().GetForward());
auto x_dims = phi::vectorize(args.x->dims()); auto x_dims = phi::vectorize(args.x->dims());
auto w_dims = phi::vectorize(args.w->dims()); auto w_dims = phi::vectorize(args.w->dims());
VLOG(10) << "cudnnConvolutionFwdAlgoPerf_t:" VLOG(10) << "cudnnConvolutionFwdAlgoPerf_t:"
<< ", x_dims:" << x_dims << ", w_dims:" << w_dims << ", args.s" << ", x_dims:" << x_dims << ", w_dims:" << w_dims << ", args.s"
<< args.s << ", args.p" << args.p << ", args.d" << args.d; << args.s << ", args.p" << args.p << ", args.d" << args.d;
algo = algo_cache.GetAlgorithm( AlgorithmsCache<AlgoT>& algo_cache =
*(framework::ConvSearchCache::Instance().GetForward());
result.algo = algo_cache.GetAlgorithm(
x_dims, w_dims, args.s, args.p, args.d, 0, x_dims, w_dims, args.s, args.p, args.d, 0,
static_cast<int64_t>(args.cudnn_dtype), [&]() { static_cast<int64_t>(args.cudnn_dtype), [&]() {
int returned_algo_count; int returned_algo_count;
std::array<perf_t, kNUM_CUDNN_FWD_ALGS> perf_stat; std::vector<PerfT> perf_results(kNUM_CUDNN_FWD_ALGS);
size_t max_workspace_size =
FindMaxWorkspaceSize(args, workspace_size_limit);
VLOG(4) << "max_workspace_size=" << ToMegaBytes(max_workspace_size)
<< " MB";
auto cudnn_find_func = [&](void* cudnn_workspace_ptr) { auto cudnn_find_func = [&](void* cudnn_workspace_ptr) {
PADDLE_ENFORCE_GPU_SUCCESS( PADDLE_ENFORCE_GPU_SUCCESS(
...@@ -308,25 +234,28 @@ struct SearchAlgorithm<cudnnConvolutionFwdAlgoPerf_t> { ...@@ -308,25 +234,28 @@ struct SearchAlgorithm<cudnnConvolutionFwdAlgoPerf_t> {
args.wdesc.desc(), args.w->data<T>(), args.cdesc.desc(), args.wdesc.desc(), args.w->data<T>(), args.cdesc.desc(),
args.odesc.desc(), const_cast<T*>(args.o->data<T>()), args.odesc.desc(), const_cast<T*>(args.o->data<T>()),
kNUM_CUDNN_FWD_ALGS, &returned_algo_count, kNUM_CUDNN_FWD_ALGS, &returned_algo_count,
perf_stat.data(), cudnn_workspace_ptr, perf_results.data(), cudnn_workspace_ptr,
workspace_size_limit)); max_workspace_size));
}; };
workspace_handle.RunFuncSync(cudnn_find_func, workspace_size_limit); workspace_handle.RunFuncSync(cudnn_find_func, max_workspace_size,
UseFixedWorkspace());
VLOG(3) << "FwdAlgo Perf result: (algo: stat, time, memory)";
for (int i = 0; i < returned_algo_count; ++i) { VLOG(4) << GetPerfResultString<PerfT>(
const auto& stat = perf_stat[i]; "[Exhaustive Search] FwdAlgo Perf result", perf_results,
VLOG(3) << stat.algo << ": " << stat.status << " " << stat.time returned_algo_count, workspace_size_limit);
<< " " << stat.memory; result.time = perf_results[0].time;
} return perf_results[0].algo;
return perf_stat[0].algo;
}); });
} }
VLOG(3) << "choose algo " << algo; VLOG(3) << "[cuDNN Convoltion] exhaustive_search=" << exhaustive_search
return algo; << ", deterministic=" << deterministic
<< ", choose algo=" << result.algo << ", workspace="
<< ToMegaBytes(GetWorkspaceSize(args, result.algo)) << " MB";
return result;
} }
static size_t GetWorkspaceSize(const ConvArgs& args, algo_t algo) { static size_t GetWorkspaceSize(const ConvArgs& args,
cudnnConvolutionFwdAlgo_t algo) {
size_t workspace_size = 0; size_t workspace_size = 0;
PADDLE_ENFORCE_GPU_SUCCESS( PADDLE_ENFORCE_GPU_SUCCESS(
platform::dynload::cudnnGetConvolutionForwardWorkspaceSize( platform::dynload::cudnnGetConvolutionForwardWorkspaceSize(
...@@ -334,68 +263,84 @@ struct SearchAlgorithm<cudnnConvolutionFwdAlgoPerf_t> { ...@@ -334,68 +263,84 @@ struct SearchAlgorithm<cudnnConvolutionFwdAlgoPerf_t> {
args.cdesc.desc(), args.odesc.desc(), algo, &workspace_size)); args.cdesc.desc(), args.odesc.desc(), algo, &workspace_size));
return workspace_size; return workspace_size;
} }
private:
static size_t FindMaxWorkspaceSize(const ConvArgs& args,
size_t workspace_size_limit) {
if (!UseFixedWorkspace()) {
size_t max_workspace_size = 0;
for (size_t algo = 0; algo < kNUM_CUDNN_FWD_ALGS; ++algo) {
size_t workspace_size = 0;
auto status =
platform::dynload::cudnnGetConvolutionForwardWorkspaceSize(
args.handle, args.idesc.desc(), args.wdesc.desc(),
args.cdesc.desc(), args.odesc.desc(),
static_cast<cudnnConvolutionFwdAlgo_t>(algo), &workspace_size);
if (status == CUDNN_STATUS_SUCCESS) {
max_workspace_size = std::max(workspace_size, max_workspace_size);
}
}
return std::min(max_workspace_size, workspace_size_limit);
} else {
return workspace_size_limit;
}
}
}; };
template <> template <>
struct SearchAlgorithm<cudnnConvolutionBwdDataAlgoPerf_t> { struct SearchAlgorithm<cudnnConvolutionBwdDataAlgoPerf_t> {
using perf_t = cudnnConvolutionBwdDataAlgoPerf_t; using PerfT = cudnnConvolutionBwdDataAlgoPerf_t;
using algo_t = cudnnConvolutionBwdDataAlgo_t; using AlgoT = cudnnConvolutionBwdDataAlgo_t;
template <typename T> template <typename T>
static algo_t Find(const ConvArgs& args, bool exhaustive_search, static SearchResult<AlgoT> Find(const ConvArgs& args, bool exhaustive_search,
bool deterministic, const phi::GPUContext& ctx) { bool deterministic,
const phi::GPUContext& ctx) {
SearchResult<AlgoT> result;
auto dtype = platform::CudnnDataType<T>::type; auto dtype = platform::CudnnDataType<T>::type;
size_t workspace_size_limit = FLAGS_conv_workspace_size_limit * 1024 * 1024; size_t workspace_size_limit = CaclWorkspaceLimitInBytes(ctx);
size_t workspace_size = 0;
bool has_got_workspace_size = true;
algo_t algo;
SetConvMathType(ctx, dtype, args.cdesc); SetConvMathType(ctx, dtype, args.cdesc);
if (!exhaustive_search && !deterministic) { if (!exhaustive_search && !deterministic) {
#if CUDNN_VERSION >= 7001 #if CUDNN_VERSION >= 7001
int perf_count; int actual_perf_count;
int best_algo_idx = 0; int best_algo_idx = 0;
std::unique_ptr<perf_t[]> perf_results( std::vector<PerfT> perf_results(kNUM_CUDNN_BWD_DATA_ALGS);
new perf_t[kNUM_CUDNN_BWD_DATA_ALGS]);
PADDLE_ENFORCE_GPU_SUCCESS( PADDLE_ENFORCE_GPU_SUCCESS(
platform::dynload::cudnnGetConvolutionBackwardDataAlgorithm_v7( platform::dynload::cudnnGetConvolutionBackwardDataAlgorithm_v7(
args.handle, args.wdesc.desc(), args.odesc.desc(), args.handle, args.wdesc.desc(), args.odesc.desc(),
args.cdesc.desc(), args.idesc.desc(), kNUM_CUDNN_BWD_DATA_ALGS, args.cdesc.desc(), args.idesc.desc(), kNUM_CUDNN_BWD_DATA_ALGS,
&perf_count, perf_results.get())); &actual_perf_count, perf_results.data()));
algo = (perf_results.get())[best_algo_idx].algo; result.algo = perf_results[best_algo_idx].algo;
#if CUDNN_VERSION < 7500 #if CUDNN_VERSION < 7500
int stride_dim = args.x->dims().size() - 2; int stride_dim = args.x->dims().size() - 2;
bool blacklist = std::any_of(args.s.begin(), args.s.begin() + stride_dim, bool blacklist = std::any_of(args.s.begin(), args.s.begin() + stride_dim,
[=](int n) { return n != 1; }); [=](int n) { return n != 1; });
if (blacklist && (static_cast<cudnnConvolutionBwdDataAlgo_t>( if (blacklist && (perf_results[best_algo_idx].algo ==
perf_results[best_algo_idx].algo) ==
CUDNN_CONVOLUTION_BWD_DATA_ALGO_FFT_TILING || CUDNN_CONVOLUTION_BWD_DATA_ALGO_FFT_TILING ||
static_cast<cudnnConvolutionBwdDataAlgo_t>( perf_results[best_algo_idx].algo ==
perf_results[best_algo_idx].algo) ==
CUDNN_CONVOLUTION_BWD_DATA_ALGO_FFT)) { CUDNN_CONVOLUTION_BWD_DATA_ALGO_FFT)) {
algo = CUDNN_CONVOLUTION_BWD_DATA_ALGO_1; result.algo = CUDNN_CONVOLUTION_BWD_DATA_ALGO_1;
} }
#endif #endif
workspace_size = GetWorkspaceSize(args, algo); result.workspace_size = GetWorkspaceSize(args, result.algo);
if (workspace_size > workspace_size_limit) { if (result.workspace_size > workspace_size_limit) {
has_got_workspace_size = false;
#if CUDNN_VERSION >= 8000 #if CUDNN_VERSION >= 8000
// cudnnGetConvolutionBackwardDataAlgorithm is removed in CUDNN-8 // cudnnGetConvolutionBackwardDataAlgorithm is removed in CUDNN-8
ChooseAlgoByWorkspace<perf_t, algo_t>(perf_results.get(), ChooseAlgoByWorkspace<PerfT, AlgoT>(perf_results, workspace_size_limit,
kNUM_CUDNN_BWD_DATA_ALGS, &result);
workspace_size_limit, &algo);
#else #else
VLOG(1) << "Fallback to non-v7 method to find conv algorithm becasue " VLOG(1) << "Fallback to non-v7 method to find conv algorithm becasue "
"the workspace size request(" "the workspace size request("
<< workspace_size << ") exceeds the limit(" << result.workspace_size << ") exceeds the limit("
<< workspace_size_limit << ")"; << workspace_size_limit << ")";
PADDLE_ENFORCE_GPU_SUCCESS( PADDLE_ENFORCE_GPU_SUCCESS(
platform::dynload::cudnnGetConvolutionBackwardDataAlgorithm( platform::dynload::cudnnGetConvolutionBackwardDataAlgorithm(
args.handle, args.wdesc.desc(), args.odesc.desc(), args.handle, args.wdesc.desc(), args.odesc.desc(),
args.cdesc.desc(), args.idesc.desc(), args.cdesc.desc(), args.idesc.desc(),
CUDNN_CONVOLUTION_BWD_DATA_SPECIFY_WORKSPACE_LIMIT, CUDNN_CONVOLUTION_BWD_DATA_SPECIFY_WORKSPACE_LIMIT,
workspace_size_limit, &algo)); workspace_size_limit, &(result.algo)));
#endif #endif
} }
#else #else
...@@ -404,29 +349,29 @@ struct SearchAlgorithm<cudnnConvolutionBwdDataAlgoPerf_t> { ...@@ -404,29 +349,29 @@ struct SearchAlgorithm<cudnnConvolutionBwdDataAlgoPerf_t> {
args.handle, args.wdesc.desc(), args.odesc.desc(), args.handle, args.wdesc.desc(), args.odesc.desc(),
args.cdesc.desc(), args.idesc.desc(), args.cdesc.desc(), args.idesc.desc(),
CUDNN_CONVOLUTION_BWD_DATA_SPECIFY_WORKSPACE_LIMIT, CUDNN_CONVOLUTION_BWD_DATA_SPECIFY_WORKSPACE_LIMIT,
workspace_size_limit, &algo)); workspace_size_limit, &(result.algo)));
#endif #endif
} else if (deterministic) { } else if (deterministic) {
return CUDNN_CONVOLUTION_BWD_DATA_ALGO_1; result.algo = CUDNN_CONVOLUTION_BWD_DATA_ALGO_1;
} else { } else {
auto& dev_ctx = ctx; auto workspace_handle = ctx.cudnn_workspace_handle();
auto workspace_handle = dev_ctx.cudnn_workspace_handle();
AlgorithmsCache<algo_t>& algo_cache =
*(framework::ConvSearchCache::Instance().GetBackwardData());
auto x_dims = phi::vectorize(args.x->dims()); auto x_dims = phi::vectorize(args.x->dims());
auto w_dims = phi::vectorize(args.w->dims()); auto w_dims = phi::vectorize(args.w->dims());
VLOG(10) << "cudnnConvolutionFwdAlgoPerf_t" VLOG(10) << "cudnnConvolutionFwdAlgoPerf_t"
<< ", x_dims:" << x_dims << ", w_dims:" << w_dims << ", args.s" << ", x_dims:" << x_dims << ", w_dims:" << w_dims << ", args.s"
<< args.s << ", args.p" << args.p << ", args.d" << args.d; << args.s << ", args.p" << args.p << ", args.d" << args.d;
algo = algo_cache.GetAlgorithm( AlgorithmsCache<AlgoT>& algo_cache =
*(framework::ConvSearchCache::Instance().GetBackwardData());
result.algo = algo_cache.GetAlgorithm(
x_dims, w_dims, args.s, args.p, args.d, 0, x_dims, w_dims, args.s, args.p, args.d, 0,
static_cast<int64_t>(args.cudnn_dtype), [&]() { static_cast<int64_t>(args.cudnn_dtype), [&]() {
int returned_algo_count; int returned_algo_count;
std::array<perf_t, kNUM_CUDNN_BWD_DATA_ALGS> perf_stat; std::vector<PerfT> perf_results(kNUM_CUDNN_BWD_DATA_ALGS);
size_t max_workspace_size =
FindMaxWorkspaceSize(args, workspace_size_limit);
VLOG(3) << "max_workspace_size=" << ToMegaBytes(max_workspace_size)
<< " MB";
auto cudnn_find_func = [&](void* cudnn_workspace_ptr) { auto cudnn_find_func = [&](void* cudnn_workspace_ptr) {
PADDLE_ENFORCE_GPU_SUCCESS( PADDLE_ENFORCE_GPU_SUCCESS(
...@@ -437,26 +382,28 @@ struct SearchAlgorithm<cudnnConvolutionBwdDataAlgoPerf_t> { ...@@ -437,26 +382,28 @@ struct SearchAlgorithm<cudnnConvolutionBwdDataAlgoPerf_t> {
args.cdesc.desc(), args.idesc.desc(), args.cdesc.desc(), args.idesc.desc(),
const_cast<T*>(args.x->data<T>()), const_cast<T*>(args.x->data<T>()),
kNUM_CUDNN_BWD_DATA_ALGS, &returned_algo_count, kNUM_CUDNN_BWD_DATA_ALGS, &returned_algo_count,
perf_stat.data(), cudnn_workspace_ptr, perf_results.data(), cudnn_workspace_ptr,
workspace_size_limit)); max_workspace_size));
}; };
workspace_handle.RunFuncSync(cudnn_find_func, workspace_size_limit); workspace_handle.RunFuncSync(cudnn_find_func, max_workspace_size,
UseFixedWorkspace());
VLOG(3) << "BwdDataAlgo Perf result: (algo: stat, time, memory)";
for (int i = 0; i < returned_algo_count; ++i) { VLOG(3) << GetPerfResultString<PerfT>(
const auto& stat = perf_stat[i]; "[Exhaustive Search] BwdDataAlgo Perf result", perf_results,
VLOG(3) << stat.algo << ": " << stat.status << " " << stat.time returned_algo_count, workspace_size_limit);
<< " " << stat.memory; result.time = perf_results[0].time;
} return perf_results[0].algo;
return perf_stat[0].algo;
}); });
} }
VLOG(3) << "choose algo " << algo; VLOG(3) << "[cuDNN Convoltion] exhaustive_search=" << exhaustive_search
return algo; << ", deterministic=" << deterministic
<< ", choose algo=" << result.algo << ", workspace="
<< ToMegaBytes(GetWorkspaceSize(args, result.algo)) << " MB";
return result;
} }
static size_t GetWorkspaceSize(const ConvArgs& args, algo_t algo) { static size_t GetWorkspaceSize(const ConvArgs& args,
cudnnConvolutionBwdDataAlgo_t algo) {
size_t workspace_size = 0; size_t workspace_size = 0;
PADDLE_ENFORCE_GPU_SUCCESS( PADDLE_ENFORCE_GPU_SUCCESS(
platform::dynload::cudnnGetConvolutionBackwardDataWorkspaceSize( platform::dynload::cudnnGetConvolutionBackwardDataWorkspaceSize(
...@@ -464,57 +411,75 @@ struct SearchAlgorithm<cudnnConvolutionBwdDataAlgoPerf_t> { ...@@ -464,57 +411,75 @@ struct SearchAlgorithm<cudnnConvolutionBwdDataAlgoPerf_t> {
args.cdesc.desc(), args.idesc.desc(), algo, &workspace_size)); args.cdesc.desc(), args.idesc.desc(), algo, &workspace_size));
return workspace_size; return workspace_size;
} }
private:
static size_t FindMaxWorkspaceSize(const ConvArgs& args,
size_t workspace_size_limit) {
if (!UseFixedWorkspace()) {
size_t max_workspace_size = 0;
for (size_t algo = 0; algo < kNUM_CUDNN_BWD_DATA_ALGS; ++algo) {
size_t workspace_size = 0;
auto status =
platform::dynload::cudnnGetConvolutionBackwardDataWorkspaceSize(
args.handle, args.wdesc.desc(), args.odesc.desc(),
args.cdesc.desc(), args.idesc.desc(),
static_cast<cudnnConvolutionBwdDataAlgo_t>(algo),
&workspace_size);
if (status == CUDNN_STATUS_SUCCESS) {
max_workspace_size = std::max(workspace_size, max_workspace_size);
}
}
return std::min(max_workspace_size, workspace_size_limit);
} else {
return workspace_size_limit;
}
}
}; };
template <> template <>
struct SearchAlgorithm<cudnnConvolutionBwdFilterAlgoPerf_t> { struct SearchAlgorithm<cudnnConvolutionBwdFilterAlgoPerf_t> {
using perf_t = cudnnConvolutionBwdFilterAlgoPerf_t; using PerfT = cudnnConvolutionBwdFilterAlgoPerf_t;
using algo_t = cudnnConvolutionBwdFilterAlgo_t; using AlgoT = cudnnConvolutionBwdFilterAlgo_t;
template <typename T> template <typename T>
static algo_t Find(const ConvArgs& args, bool exhaustive_search, static SearchResult<AlgoT> Find(const ConvArgs& args, bool exhaustive_search,
bool deterministic, const phi::GPUContext& ctx) { bool deterministic,
const phi::GPUContext& ctx) {
platform::CUDAGraphCaptureModeGuard guard; platform::CUDAGraphCaptureModeGuard guard;
SearchResult<AlgoT> result;
auto dtype = platform::CudnnDataType<T>::type; auto dtype = platform::CudnnDataType<T>::type;
size_t workspace_size_limit = FLAGS_conv_workspace_size_limit * 1024 * 1024; size_t workspace_size_limit = CaclWorkspaceLimitInBytes(ctx);
size_t workspace_size = 0;
bool has_got_workspace_size = true;
SetConvMathType(ctx, dtype, args.cdesc); SetConvMathType(ctx, dtype, args.cdesc);
algo_t algo;
if (!exhaustive_search && !deterministic) { if (!exhaustive_search && !deterministic) {
#if CUDNN_VERSION >= 7001 #if CUDNN_VERSION >= 7001
using perf_t = cudnnConvolutionBwdFilterAlgoPerf_t; int actual_perf_count;
int perf_count;
int best_algo_idx = 0; int best_algo_idx = 0;
std::unique_ptr<perf_t[]> perf_results( std::vector<PerfT> perf_results(kNUM_CUDNN_BWD_FILTER_ALGS);
new perf_t[kNUM_CUDNN_BWD_FILTER_ALGS]);
PADDLE_ENFORCE_GPU_SUCCESS( PADDLE_ENFORCE_GPU_SUCCESS(
platform::dynload::cudnnGetConvolutionBackwardFilterAlgorithm_v7( platform::dynload::cudnnGetConvolutionBackwardFilterAlgorithm_v7(
args.handle, args.idesc.desc(), args.odesc.desc(), args.handle, args.idesc.desc(), args.odesc.desc(),
args.cdesc.desc(), args.wdesc.desc(), kNUM_CUDNN_BWD_FILTER_ALGS, args.cdesc.desc(), args.wdesc.desc(), kNUM_CUDNN_BWD_FILTER_ALGS,
&perf_count, perf_results.get())); &actual_perf_count, perf_results.data()));
algo = (perf_results.get())[best_algo_idx].algo; result.algo = perf_results[best_algo_idx].algo;
workspace_size = (perf_results.get())[best_algo_idx].memory; result.workspace_size = perf_results[best_algo_idx].memory;
if (workspace_size > workspace_size_limit) { if (result.workspace_size > workspace_size_limit) {
workspace_size = workspace_size_limit;
#if CUDNN_VERSION >= 8000 #if CUDNN_VERSION >= 8000
// cudnnGetConvolutionBackwardFilterAlgorithm is removed in CUDNN-8 // cudnnGetConvolutionBackwardFilterAlgorithm is removed in CUDNN-8
ChooseAlgoByWorkspace<perf_t, algo_t>(perf_results.get(), ChooseAlgoByWorkspace<PerfT, AlgoT>(perf_results, workspace_size_limit,
kNUM_CUDNN_BWD_FILTER_ALGS, &result);
workspace_size_limit, &algo);
#else #else
VLOG(1) << "Fallback to non-v7 method to find conv algorithm becasue " VLOG(1) << "Fallback to non-v7 method to find conv algorithm becasue "
"the workspace size request(" "the workspace size request("
<< workspace_size << ") exceeds the limit(" << result.workspace_size << ") exceeds the limit("
<< workspace_size_limit << ")"; << workspace_size_limit << ")";
PADDLE_ENFORCE_GPU_SUCCESS( PADDLE_ENFORCE_GPU_SUCCESS(
platform::dynload::cudnnGetConvolutionBackwardFilterAlgorithm( platform::dynload::cudnnGetConvolutionBackwardFilterAlgorithm(
args.handle, args.idesc.desc(), args.odesc.desc(), args.handle, args.idesc.desc(), args.odesc.desc(),
args.cdesc.desc(), args.wdesc.desc(), args.cdesc.desc(), args.wdesc.desc(),
CUDNN_CONVOLUTION_BWD_FILTER_SPECIFY_WORKSPACE_LIMIT, CUDNN_CONVOLUTION_BWD_FILTER_SPECIFY_WORKSPACE_LIMIT,
workspace_size_limit, &algo)); workspace_size_limit, &(result.algo)));
#endif #endif
} }
#else #else
...@@ -523,28 +488,32 @@ struct SearchAlgorithm<cudnnConvolutionBwdFilterAlgoPerf_t> { ...@@ -523,28 +488,32 @@ struct SearchAlgorithm<cudnnConvolutionBwdFilterAlgoPerf_t> {
args.handle, args.idesc.desc(), args.odesc.desc(), args.handle, args.idesc.desc(), args.odesc.desc(),
args.cdesc.desc(), args.wdesc.desc(), args.cdesc.desc(), args.wdesc.desc(),
CUDNN_CONVOLUTION_BWD_FILTER_SPECIFY_WORKSPACE_LIMIT, CUDNN_CONVOLUTION_BWD_FILTER_SPECIFY_WORKSPACE_LIMIT,
workspace_size_limit, &algo)); workspace_size_limit, &(result.algo)));
#endif #endif
} else if (deterministic) { } else if (deterministic) {
return CUDNN_CONVOLUTION_BWD_FILTER_ALGO_1; result.algo = CUDNN_CONVOLUTION_BWD_FILTER_ALGO_1;
} else { } else {
auto& dev_ctx = ctx; auto workspace_handle = ctx.cudnn_workspace_handle();
auto workspace_handle = dev_ctx.cudnn_workspace_handle();
AlgorithmsCache<algo_t>& algo_cache =
*(framework::ConvSearchCache::Instance().GetBackwardFilter());
auto x_dims = phi::vectorize(args.x->dims()); auto x_dims = phi::vectorize(args.x->dims());
auto w_dims = phi::vectorize(args.w->dims()); auto w_dims = phi::vectorize(args.w->dims());
VLOG(10) << "cudnnConvolutionFwdAlgoPerf_t:" VLOG(10) << "cudnnConvolutionFwdAlgoPerf_t:"
<< ", x_dims:" << x_dims << ", w_dims:" << w_dims << ", args.s" << ", x_dims:" << x_dims << ", w_dims:" << w_dims << ", args.s"
<< args.s << ", args.p" << args.p << ", args.d" << args.d; << args.s << ", args.p" << args.p << ", args.d" << args.d;
AlgorithmsCache<AlgoT>& algo_cache =
*(framework::ConvSearchCache::Instance().GetBackwardFilter());
if (dtype != CUDNN_DATA_HALF) { if (dtype != CUDNN_DATA_HALF) {
algo = algo_cache.GetAlgorithm( result.algo = algo_cache.GetAlgorithm(
x_dims, w_dims, args.s, args.p, args.d, 0, x_dims, w_dims, args.s, args.p, args.d, 0,
static_cast<int64_t>(args.cudnn_dtype), [&]() { static_cast<int64_t>(args.cudnn_dtype), [&]() {
int returned_algo_count; int returned_algo_count;
std::array<perf_t, kNUM_CUDNN_BWD_FILTER_ALGS> perf_stat; std::vector<PerfT> perf_results(kNUM_CUDNN_BWD_FILTER_ALGS);
size_t max_workspace_size =
FindMaxWorkspaceSize(args, workspace_size_limit);
VLOG(3) << "max_workspace_size="
<< ToMegaBytes(max_workspace_size) << " MB";
auto cudnn_find_func = [&](void* cudnn_workspace_ptr) { auto cudnn_find_func = [&](void* cudnn_workspace_ptr) {
PADDLE_ENFORCE_GPU_SUCCESS( PADDLE_ENFORCE_GPU_SUCCESS(
platform::dynload:: platform::dynload::
...@@ -554,29 +523,26 @@ struct SearchAlgorithm<cudnnConvolutionBwdFilterAlgoPerf_t> { ...@@ -554,29 +523,26 @@ struct SearchAlgorithm<cudnnConvolutionBwdFilterAlgoPerf_t> {
args.cdesc.desc(), args.wdesc.desc(), args.cdesc.desc(), args.wdesc.desc(),
const_cast<T*>(args.w->data<T>()), const_cast<T*>(args.w->data<T>()),
kNUM_CUDNN_BWD_FILTER_ALGS, &returned_algo_count, kNUM_CUDNN_BWD_FILTER_ALGS, &returned_algo_count,
perf_stat.data(), cudnn_workspace_ptr, perf_results.data(), cudnn_workspace_ptr,
workspace_size_limit)); max_workspace_size));
}; };
workspace_handle.RunFuncSync(cudnn_find_func, workspace_handle.RunFuncSync(cudnn_find_func, max_workspace_size,
workspace_size_limit); UseFixedWorkspace());
VLOG(3) VLOG(3) << GetPerfResultString<PerfT>(
<< "BwdFilterAlgo Perf result: (algo: stat, time, memory)"; "[Exhaustive Search] BwdFilterAlgo Perf result", perf_results,
for (int i = 0; i < returned_algo_count; ++i) { returned_algo_count, workspace_size_limit);
const auto& stat = perf_stat[i]; result.time = perf_results[0].time;
VLOG(3) << stat.algo << ": " << stat.status << " " << stat.time return perf_results[0].algo;
<< " " << stat.memory;
}
return perf_stat[0].algo;
}); });
} else { } else {
auto max_algos = MaxBwdFilterAlgos(args.handle); result.algo = algo_cache.GetAlgorithm(
algo = algo_cache.GetAlgorithm(
x_dims, w_dims, args.s, args.p, args.d, 0, x_dims, w_dims, args.s, args.p, args.d, 0,
static_cast<int64_t>(args.cudnn_dtype), [&]() { static_cast<int64_t>(args.cudnn_dtype), [&]() {
algo_t chosen_algo; SearchResult<AlgoT> algo_result;
std::vector<perf_t> perf_results(max_algos);
int actual_algos = 0; int actual_algos = 0;
std::vector<PerfT> perf_results(kNUM_CUDNN_BWD_FILTER_ALGS);
PADDLE_ENFORCE_GPU_SUCCESS( PADDLE_ENFORCE_GPU_SUCCESS(
platform::dynload:: platform::dynload::
cudnnFindConvolutionBackwardFilterAlgorithm( cudnnFindConvolutionBackwardFilterAlgorithm(
...@@ -585,17 +551,21 @@ struct SearchAlgorithm<cudnnConvolutionBwdFilterAlgoPerf_t> { ...@@ -585,17 +551,21 @@ struct SearchAlgorithm<cudnnConvolutionBwdFilterAlgoPerf_t> {
perf_results.size(), &actual_algos, perf_results.size(), &actual_algos,
perf_results.data())); perf_results.data()));
perf_results.resize(actual_algos); perf_results.resize(actual_algos);
ChooseAlgo<perf_t, algo_t>(perf_results, workspace_size_limit, ChooseAlgo(perf_results, workspace_size_limit, &algo_result);
&chosen_algo); result.time = algo_result.time;
return chosen_algo; return algo_result.algo;
}); });
} }
} }
VLOG(3) << "choose algo " << algo; VLOG(3) << "[cuDNN Convoltion] exhaustive_search=" << exhaustive_search
return algo; << ", deterministic=" << deterministic
<< ", choose algo=" << result.algo << ", workspace="
<< ToMegaBytes(GetWorkspaceSize(args, result.algo)) << " MB";
return result;
} }
static size_t GetWorkspaceSize(const ConvArgs& args, algo_t algo) { static size_t GetWorkspaceSize(const ConvArgs& args,
cudnnConvolutionBwdFilterAlgo_t algo) {
platform::CUDAGraphCaptureModeGuard guard; platform::CUDAGraphCaptureModeGuard guard;
size_t workspace_size = 0; size_t workspace_size = 0;
PADDLE_ENFORCE_GPU_SUCCESS( PADDLE_ENFORCE_GPU_SUCCESS(
...@@ -604,6 +574,69 @@ struct SearchAlgorithm<cudnnConvolutionBwdFilterAlgoPerf_t> { ...@@ -604,6 +574,69 @@ struct SearchAlgorithm<cudnnConvolutionBwdFilterAlgoPerf_t> {
args.cdesc.desc(), args.wdesc.desc(), algo, &workspace_size)); args.cdesc.desc(), args.wdesc.desc(), algo, &workspace_size));
return workspace_size; return workspace_size;
} }
private:
static size_t FindMaxWorkspaceSize(const ConvArgs& args,
size_t workspace_size_limit) {
if (!UseFixedWorkspace()) {
size_t max_workspace_size = 0;
for (size_t algo = 0; algo < kNUM_CUDNN_BWD_FILTER_ALGS; ++algo) {
size_t workspace_size = 0;
auto status =
platform::dynload::cudnnGetConvolutionBackwardFilterWorkspaceSize(
args.handle, args.idesc.desc(), args.odesc.desc(),
args.cdesc.desc(), args.wdesc.desc(),
static_cast<cudnnConvolutionBwdFilterAlgo_t>(algo),
&workspace_size);
if (status == CUDNN_STATUS_SUCCESS) {
max_workspace_size = std::max(workspace_size, max_workspace_size);
}
}
return std::min(max_workspace_size, workspace_size_limit);
} else {
return workspace_size_limit;
}
}
static void ChooseAlgo(const std::vector<PerfT>& perf_results,
size_t workspace_limit,
SearchResult<AlgoT>* algo_result) {
VLOG(3) << GetPerfResultString<PerfT>(
"[Exhaustive Search] BwdFilterAlgo Perf result", perf_results,
perf_results.size(), workspace_limit);
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_limit)) {
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 = result.algo;
algo_result->time = result.time;
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;
break;
}
}
}
}; };
} // namespace operators } // namespace operators
......
paddle/fluid/operators/conv_cudnn_op_cache.h
浏览文件 @ b937cdc5
...@@ -20,7 +20,7 @@ limitations under the License. */ ...@@ -20,7 +20,7 @@ limitations under the License. */
#include "paddle/fluid/framework/operator.h" #include "paddle/fluid/framework/operator.h"
#include "paddle/fluid/platform/device/gpu/gpu_dnn.h" #include "paddle/fluid/platform/device/gpu/gpu_dnn.h"
DECLARE_uint64(conv_workspace_size_limit); DECLARE_int64(conv_workspace_size_limit);
DECLARE_bool(cudnn_exhaustive_search); DECLARE_bool(cudnn_exhaustive_search);
DECLARE_int64(cudnn_exhaustive_search_times); DECLARE_int64(cudnn_exhaustive_search_times);
......
paddle/fluid/operators/conv_miopen_helper.h
浏览文件 @ b937cdc5
...@@ -14,42 +14,12 @@ limitations under the License. */ ...@@ -14,42 +14,12 @@ limitations under the License. */
#pragma once #pragma once
#include <algorithm> #include "paddle/fluid/operators/conv_base_helper.h"
#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"
#include "paddle/fluid/platform/device/gpu/gpu_dnn.h"
#include "paddle/phi/backends/gpu/gpu_context.h"
namespace paddle { namespace paddle {
namespace operators { namespace operators {
using Tensor = framework::Tensor; using ConvArgs = ConvArgsBase<miopenHandle_t, miopenDataType_t>;
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> template <typename DeviceContext, typename T, size_t D>
static void RemovePaddingSlice(const phi::GPUContext& context, static void RemovePaddingSlice(const phi::GPUContext& context,
...@@ -66,9 +36,8 @@ static void RemovePaddingSlice(const phi::GPUContext& context, ...@@ -66,9 +36,8 @@ static void RemovePaddingSlice(const phi::GPUContext& context,
extents[i] = new_out_dims[i]; extents[i] = new_out_dims[i];
} }
int start;
for (size_t i = 0; i < axes.size(); ++i) { for (size_t i = 0; i < axes.size(); ++i) {
start = starts[i]; int start = starts[i];
if (start < 0) { if (start < 0) {
start = (start + in_dims[axes[i]]); start = (start + in_dims[axes[i]]);
} }
...@@ -85,41 +54,6 @@ static void RemovePaddingSlice(const phi::GPUContext& context, ...@@ -85,41 +54,6 @@ static void RemovePaddingSlice(const phi::GPUContext& context,
out_t.device(place) = in_t.slice(offsets, extents); 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 <> template <>
struct SearchAlgorithm<miopenConvFwdAlgorithm_t> { struct SearchAlgorithm<miopenConvFwdAlgorithm_t> {
using perf_t = miopenConvAlgoPerf_t; using perf_t = miopenConvAlgoPerf_t;
......
paddle/fluid/operators/fused/fusion_conv_inception_op.cu
浏览文件 @ b937cdc5
...@@ -16,8 +16,6 @@ limitations under the License. */ ...@@ -16,8 +16,6 @@ limitations under the License. */
#include "paddle/fluid/operators/conv_cudnn_op_cache.h" #include "paddle/fluid/operators/conv_cudnn_op_cache.h"
#include "paddle/fluid/platform/device/gpu/gpu_dnn.h" #include "paddle/fluid/platform/device/gpu/gpu_dnn.h"
DECLARE_uint64(conv_workspace_size_limit);
namespace paddle { namespace paddle {
namespace operators { namespace operators {
......
paddle/fluid/platform/device/gpu/gpu_info.cc
浏览文件 @ b937cdc5
...@@ -188,6 +188,8 @@ class RecordedGpuMallocHelper { ...@@ -188,6 +188,8 @@ class RecordedGpuMallocHelper {
if (UNLIKELY(malloc_managed_memory)) { if (UNLIKELY(malloc_managed_memory)) {
result = cudaMallocManaged(ptr, size); result = cudaMallocManaged(ptr, size);
} else { } else {
VLOG(10) << "[cudaMalloc] size=" << static_cast<double>(size) / (1 << 20)
<< " MB";
result = cudaMalloc(ptr, size); result = cudaMalloc(ptr, size);
} }
#endif #endif
...@@ -226,6 +228,8 @@ class RecordedGpuMallocHelper { ...@@ -226,6 +228,8 @@ class RecordedGpuMallocHelper {
if (err != hipErrorDeinitialized) { if (err != hipErrorDeinitialized) {
#else #else
auto err = cudaFree(ptr); auto err = cudaFree(ptr);
VLOG(10) << "[cudaFree] size=" << static_cast<double>(size) / (1 << 20)
<< " MB";
if (err != cudaErrorCudartUnloading) { if (err != cudaErrorCudartUnloading) {
#endif #endif
PADDLE_ENFORCE_GPU_SUCCESS(err); PADDLE_ENFORCE_GPU_SUCCESS(err);
......
paddle/fluid/platform/device_context.cc
浏览文件 @ b937cdc5
...@@ -522,8 +522,8 @@ CUDADeviceContext::CUDADeviceContext(CUDAPlace place) : phi::GPUContext(place) { ...@@ -522,8 +522,8 @@ CUDADeviceContext::CUDADeviceContext(CUDAPlace place) : phi::GPUContext(place) {
cuda_stream_.reset(new stream::CUDAStream(phi::GPUContext::stream(), place)); cuda_stream_.reset(new stream::CUDAStream(phi::GPUContext::stream(), place));
auto& instance = memory::allocation::AllocatorFacade::Instance(); auto& instance = memory::allocation::AllocatorFacade::Instance();
instance.SetDefaultStream(place, phi::GPUContext::stream()); instance.SetDefaultStream(place, phi::GPUContext::stream());
workspace_.reset( workspace_.reset(new phi::DnnWorkspaceHandle(
new phi::DnnWorkspaceHandle(instance.GetAllocator(place).get())); instance.GetAllocator(place).get(), stream()));
} }
CUDADeviceContext::~CUDADeviceContext() = default; CUDADeviceContext::~CUDADeviceContext() = default;
...@@ -623,7 +623,8 @@ phi::DnnWorkspaceHandle CUDADeviceContext::cudnn_workspace_handle() const { ...@@ -623,7 +623,8 @@ phi::DnnWorkspaceHandle CUDADeviceContext::cudnn_workspace_handle() const {
return phi::DnnWorkspaceHandle( return phi::DnnWorkspaceHandle(
memory::allocation::AllocatorFacade::Instance() memory::allocation::AllocatorFacade::Instance()
.GetAllocator(GetPlace()) .GetAllocator(GetPlace())
.get()); .get(),
stream());
} }
return phi::GPUContext::cudnn_workspace_handle(); return phi::GPUContext::cudnn_workspace_handle();
} }
......
paddle/fluid/platform/flags.cc
浏览文件 @ b937cdc5
...@@ -161,10 +161,9 @@ PADDLE_DEFINE_EXPORTED_bool( ...@@ -161,10 +161,9 @@ PADDLE_DEFINE_EXPORTED_bool(
* increased. * increased.
* Users need to balance memory and speed. * Users need to balance memory and speed.
*/ */
PADDLE_DEFINE_EXPORTED_uint64( PADDLE_DEFINE_EXPORTED_int64(conv_workspace_size_limit,
conv_workspace_size_limit, paddle::platform::kDefaultConvWorkspaceSizeLimitMB,
paddle::platform::kDefaultConvWorkspaceSizeLimitMB, "cuDNN convolution workspace limit in MB unit.");
"cuDNN convolution workspace limit in MB unit.");
/** /**
* CUDNN related FLAG * CUDNN related FLAG
......
paddle/phi/backends/gpu/gpu_context.cc
浏览文件 @ b937cdc5
...@@ -12,6 +12,7 @@ distributed under the License is distributed on an "AS IS" BASIS, ...@@ -12,6 +12,7 @@ distributed under the License is distributed on an "AS IS" BASIS,
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and See the License for the specific language governing permissions and
limitations under the License. */ limitations under the License. */
#include "paddle/phi/backends/gpu/gpu_context.h" #include "paddle/phi/backends/gpu/gpu_context.h"
#include <algorithm> #include <algorithm>
#include <array> #include <array>
...@@ -155,6 +156,39 @@ static void StreamCallbackFunc(gpuStream_t stream, ...@@ -155,6 +156,39 @@ static void StreamCallbackFunc(gpuStream_t stream,
} // namespace internal } // namespace internal
void DnnWorkspaceHandle::RunFuncSync(
const std::function<void(void*)>& cudnn_func,
size_t required_workspace_bytes,
bool use_cached_allocation) {
bool need_realloc = required_workspace_bytes > WorkspaceSize();
if (need_realloc && !use_cached_allocation) {
void* workspace_ptr = nullptr;
size_t size = ((required_workspace_bytes + 255) >> 8) << 8;
std::lock_guard<std::mutex> guard(*mtx_);
#ifdef PADDLE_WITH_HIP
auto status = hipMalloc(&workspace_ptr, size);
#else
auto status = cudaMalloc(&workspace_ptr, size);
#endif
if (status == gpuSuccess) {
cudnn_func(workspace_ptr);
phi::backends::gpu::GpuStreamSync(stream_);
#ifdef PADDLE_WITH_HIP
PADDLE_ENFORCE_GPU_SUCCESS(hipFree(workspace_ptr));
#else
PADDLE_ENFORCE_GPU_SUCCESS(cudaFree(workspace_ptr));
#endif
return;
}
}
RunFunc(cudnn_func, required_workspace_bytes);
if (need_realloc) {
// Release the workspace allocated in this running.
ResetWorkspace();
}
}
void DnnWorkspaceHandle::ResetWorkspace() { allocation_ = nullptr; } void DnnWorkspaceHandle::ResetWorkspace() { allocation_ = nullptr; }
void DnnWorkspaceHandle::ReallocWorkspace(size_t required_workspace_bytes) { void DnnWorkspaceHandle::ReallocWorkspace(size_t required_workspace_bytes) {
...@@ -295,13 +329,13 @@ struct GPUContext::Impl { ...@@ -295,13 +329,13 @@ struct GPUContext::Impl {
void InitDnnWorkspace() { void InitDnnWorkspace() {
PD_CHECK(allocator_ != nullptr, PD_CHECK(allocator_ != nullptr,
"the device allocator for gpu context is nullptr."); "the device allocator for gpu context is nullptr.");
workspace_ = new DnnWorkspaceHandle(allocator_); workspace_ = new DnnWorkspaceHandle(allocator_, stream_);
} }
void DestoryInternalWorkspace() { void DestoryInternalWorkspace() {
if (owned_ && workspace_ != nullptr) { if (owned_ && workspace_ != nullptr) {
delete workspace_; delete workspace_;
stream_ = nullptr; workspace_ = nullptr;
} }
} }
...@@ -313,7 +347,7 @@ struct GPUContext::Impl { ...@@ -313,7 +347,7 @@ struct GPUContext::Impl {
DnnWorkspaceHandle GetDnnWorkspace() { DnnWorkspaceHandle GetDnnWorkspace() {
PD_CHECK(allocator_ != nullptr, PD_CHECK(allocator_ != nullptr,
"the device allocator for gpu context is nullptr."); "the device allocator for gpu context is nullptr.");
return DnnWorkspaceHandle(allocator_); return DnnWorkspaceHandle(allocator_, stream_);
} }
void InitStream() { void InitStream() {
......
paddle/phi/backends/gpu/gpu_context.h
浏览文件 @ b937cdc5
...@@ -21,6 +21,7 @@ limitations under the License. */ ...@@ -21,6 +21,7 @@ limitations under the License. */
#include "paddle/phi/backends/gpu/forwards.h" #include "paddle/phi/backends/gpu/forwards.h"
#include "paddle/phi/backends/gpu/gpu_decls.h" #include "paddle/phi/backends/gpu/gpu_decls.h"
#include "paddle/phi/backends/gpu/gpu_helper.h" #include "paddle/phi/backends/gpu/gpu_helper.h"
#include "paddle/phi/backends/gpu/gpu_info.h"
#include "paddle/phi/common/place.h" #include "paddle/phi/common/place.h"
#include "paddle/phi/core/device_context.h" #include "paddle/phi/core/device_context.h"
...@@ -28,8 +29,8 @@ namespace phi { ...@@ -28,8 +29,8 @@ namespace phi {
class DnnWorkspaceHandle { class DnnWorkspaceHandle {
public: public:
explicit inline DnnWorkspaceHandle(Allocator* allocator) inline DnnWorkspaceHandle(Allocator* allocator, gpuStream_t stream)
: allocator_(allocator) { : allocator_(allocator), stream_(stream) {
mtx_.reset(new std::mutex()); mtx_.reset(new std::mutex());
} }
...@@ -48,11 +49,9 @@ class DnnWorkspaceHandle { ...@@ -48,11 +49,9 @@ class DnnWorkspaceHandle {
* running the function. Currently this function is only used when cudnn * running the function. Currently this function is only used when cudnn
* exhaustive searching and callers have to guarantee that the input function * exhaustive searching and callers have to guarantee that the input function
* is host blocking */ * is host blocking */
inline void RunFuncSync(const std::function<void(void*)>& cudnn_func, void RunFuncSync(const std::function<void(void*)>& cudnn_func,
size_t required_workspace_bytes) { size_t required_workspace_bytes,
RunFunc(cudnn_func, required_workspace_bytes); bool use_cached_allocation = true);
ResetWorkspace();
}
inline size_t WorkspaceSize() { inline size_t WorkspaceSize() {
if (allocation_ == nullptr) { if (allocation_ == nullptr) {
...@@ -70,7 +69,8 @@ class DnnWorkspaceHandle { ...@@ -70,7 +69,8 @@ class DnnWorkspaceHandle {
private: private:
Allocator::AllocationPtr allocation_{nullptr}; Allocator::AllocationPtr allocation_{nullptr};
Allocator* allocator_{nullptr}; Allocator* allocator_{nullptr}; // Not owned
gpuStream_t stream_{nullptr}; // Not owned
std::unique_ptr<std::mutex> mtx_; std::unique_ptr<std::mutex> mtx_;
}; };
......
paddle/phi/kernels/autotune/CMakeLists.txt
浏览文件 @ b937cdc5
if (WITH_GPU) if (WITH_GPU)
nv_test(gpu_timer_test SRCS gpu_timer_test.cu DEPS gtest) nv_test(gpu_timer_test SRCS gpu_timer_test.cu DEPS gtest)
nv_test(auto_tune_test SRCS auto_tune_test.cu DEPS gtest) nv_test(auto_tune_test SRCS auto_tune_test.cu DEPS gtest)
elseif (WITH_ROCM) elseif (WITH_ROCM)
hip_test(gpu_timer_test SRCS gpu_timer_test.cu DEPS gtest) hip_test(gpu_timer_test SRCS gpu_timer_test.cu DEPS gtest)
hip_test(auto_tune_test SRCS auto_tune_test.cu DEPS gtest) hip_test(auto_tune_test SRCS auto_tune_test.cu DEPS gtest)
......
paddle/phi/kernels/gpudnn/conv_grad_grad_kernel.cu
浏览文件 @ b937cdc5
...@@ -289,21 +289,17 @@ void ConvCudnnGradGradKernel( ...@@ -289,21 +289,17 @@ void ConvCudnnGradGradKernel(
dtype}; dtype};
#ifdef PADDLE_WITH_HIP #ifdef PADDLE_WITH_HIP
miopenConvFwdAlgorithm_t fwd_algo1 = static_cast<miopenConvFwdAlgorithm_t>(0); paddle::operators::SearchResult<miopenConvFwdAlgorithm_t> fwd_result1;
miopenConvFwdAlgorithm_t fwd_algo2 = static_cast<miopenConvFwdAlgorithm_t>(0); paddle::operators::SearchResult<miopenConvFwdAlgorithm_t> fwd_result2;
miopenConvBwdDataAlgorithm_t data_algo = paddle::operators::SearchResult<miopenConvBwdDataAlgorithm_t> data_result;
static_cast<miopenConvBwdDataAlgorithm_t>(0); paddle::operators::SearchResult<miopenConvBwdWeightsAlgorithm_t>
miopenConvBwdWeightsAlgorithm_t filter_algo = filter_result;
static_cast<miopenConvBwdWeightsAlgorithm_t>(0);
#else #else
cudnnConvolutionFwdAlgo_t fwd_algo1 = paddle::operators::SearchResult<cudnnConvolutionFwdAlgo_t> fwd_result1;
static_cast<cudnnConvolutionFwdAlgo_t>(0); paddle::operators::SearchResult<cudnnConvolutionFwdAlgo_t> fwd_result2;
cudnnConvolutionFwdAlgo_t fwd_algo2 = paddle::operators::SearchResult<cudnnConvolutionBwdDataAlgo_t> data_result;
static_cast<cudnnConvolutionFwdAlgo_t>(0); paddle::operators::SearchResult<cudnnConvolutionBwdFilterAlgo_t>
cudnnConvolutionBwdDataAlgo_t data_algo = filter_result;
static_cast<cudnnConvolutionBwdDataAlgo_t>(0);
cudnnConvolutionBwdFilterAlgo_t filter_algo =
static_cast<cudnnConvolutionBwdFilterAlgo_t>(0);
#endif #endif
auto layout = paddle::platform::GetCudnnTensorFormat( auto layout = paddle::platform::GetCudnnTensorFormat(
...@@ -332,13 +328,13 @@ void ConvCudnnGradGradKernel( ...@@ -332,13 +328,13 @@ void ConvCudnnGradGradKernel(
using search1 = using search1 =
paddle::operators::SearchAlgorithm<miopenConvFwdAlgorithm_t>; paddle::operators::SearchAlgorithm<miopenConvFwdAlgorithm_t>;
workspace_size = search1::GetWorkspaceSize(args1); workspace_size = search1::GetWorkspaceSize(args1);
fwd_algo1 = search1::Find<T>( fwd_result1.algo = search1::Find<T>(
args1, exhaustive_search, false, workspace_size, ctx); args1, exhaustive_search, false, workspace_size, ctx);
#else #else
using search1 = using search1 =
paddle::operators::SearchAlgorithm<cudnnConvolutionFwdAlgoPerf_t>; paddle::operators::SearchAlgorithm<cudnnConvolutionFwdAlgoPerf_t>;
fwd_algo1 = search1::Find<T>(args1, exhaustive_search, false, ctx); fwd_result1 = search1::Find<T>(args1, exhaustive_search, false, ctx);
workspace_size = search1::GetWorkspaceSize(args1, fwd_algo1); workspace_size = search1::GetWorkspaceSize(args1, fwd_result1.algo);
#endif #endif
} }
...@@ -360,14 +356,14 @@ void ConvCudnnGradGradKernel( ...@@ -360,14 +356,14 @@ void ConvCudnnGradGradKernel(
paddle::operators::SearchAlgorithm<miopenConvFwdAlgorithm_t>; paddle::operators::SearchAlgorithm<miopenConvFwdAlgorithm_t>;
workspace_size = workspace_size =
std::max(workspace_size, search2::GetWorkspaceSize(args2)); std::max(workspace_size, search2::GetWorkspaceSize(args2));
fwd_algo2 = search2::Find<T>( fwd_result2.algo = search2::Find<T>(
args2, exhaustive_search, false, workspace_size, ctx); args2, exhaustive_search, false, workspace_size, ctx);
#else #else
using search2 = using search2 =
paddle::operators::SearchAlgorithm<cudnnConvolutionFwdAlgoPerf_t>; paddle::operators::SearchAlgorithm<cudnnConvolutionFwdAlgoPerf_t>;
fwd_algo2 = search2::Find<T>(args2, exhaustive_search, false, ctx); fwd_result2 = search2::Find<T>(args2, exhaustive_search, false, ctx);
workspace_size = workspace_size = std::max(
std::max(workspace_size, search2::GetWorkspaceSize(args2, fwd_algo2)); workspace_size, search2::GetWorkspaceSize(args2, fwd_result2.algo));
#endif #endif
} }
} }
...@@ -389,15 +385,15 @@ void ConvCudnnGradGradKernel( ...@@ -389,15 +385,15 @@ void ConvCudnnGradGradKernel(
using search3 = using search3 =
paddle::operators::SearchAlgorithm<miopenConvBwdWeightsAlgorithm_t>; paddle::operators::SearchAlgorithm<miopenConvBwdWeightsAlgorithm_t>;
workspace_size = std::max(workspace_size, search3::GetWorkspaceSize(args3)); workspace_size = std::max(workspace_size, search3::GetWorkspaceSize(args3));
filter_algo = search3::Find<T>( filter_result.algo = search3::Find<T>(
args3, exhaustive_search, deterministic, workspace_size, ctx); args3, exhaustive_search, deterministic, workspace_size, ctx);
#else #else
using search3 = using search3 =
paddle::operators::SearchAlgorithm<cudnnConvolutionBwdFilterAlgoPerf_t>; paddle::operators::SearchAlgorithm<cudnnConvolutionBwdFilterAlgoPerf_t>;
filter_algo = filter_result =
search3::Find<T>(args3, exhaustive_search, deterministic, ctx); search3::Find<T>(args3, exhaustive_search, deterministic, ctx);
workspace_size = workspace_size = std::max(
std::max(workspace_size, search3::GetWorkspaceSize(args3, filter_algo)); workspace_size, search3::GetWorkspaceSize(args3, filter_result.algo));
#endif #endif
} }
...@@ -419,14 +415,15 @@ void ConvCudnnGradGradKernel( ...@@ -419,14 +415,15 @@ void ConvCudnnGradGradKernel(
using search4 = using search4 =
paddle::operators::SearchAlgorithm<miopenConvBwdDataAlgorithm_t>; paddle::operators::SearchAlgorithm<miopenConvBwdDataAlgorithm_t>;
workspace_size = std::max(workspace_size, search4::GetWorkspaceSize(args4)); workspace_size = std::max(workspace_size, search4::GetWorkspaceSize(args4));
data_algo = search4::Find<T>( data_result.algo = search4::Find<T>(
args4, exhaustive_search, deterministic, workspace_size, ctx); args4, exhaustive_search, deterministic, workspace_size, ctx);
#else #else
using search4 = using search4 =
paddle::operators::SearchAlgorithm<cudnnConvolutionBwdDataAlgoPerf_t>; paddle::operators::SearchAlgorithm<cudnnConvolutionBwdDataAlgoPerf_t>;
data_algo = search4::Find<T>(args4, exhaustive_search, deterministic, ctx); data_result =
workspace_size = search4::Find<T>(args4, exhaustive_search, deterministic, ctx);
std::max(workspace_size, search4::GetWorkspaceSize(args4, data_algo)); workspace_size = std::max(
workspace_size, search4::GetWorkspaceSize(args4, data_result.algo));
#endif #endif
} }
...@@ -471,7 +468,7 @@ void ConvCudnnGradGradKernel( ...@@ -471,7 +468,7 @@ void ConvCudnnGradGradKernel(
args1.wdesc.desc(), args1.wdesc.desc(),
w, w,
args1.cdesc.desc(), args1.cdesc.desc(),
fwd_algo1, fwd_result1.algo,
&beta, &beta,
args1.odesc.desc(), args1.odesc.desc(),
transformed_ddy_channel, transformed_ddy_channel,
...@@ -492,7 +489,7 @@ void ConvCudnnGradGradKernel( ...@@ -492,7 +489,7 @@ void ConvCudnnGradGradKernel(
args1.wdesc.desc(), args1.wdesc.desc(),
w + i * group_offset_filter, w + i * group_offset_filter,
args1.cdesc.desc(), args1.cdesc.desc(),
fwd_algo1, fwd_result1.algo,
workspace_ptr, workspace_ptr,
workspace_size, workspace_size,
&beta, &beta,
...@@ -517,7 +514,7 @@ void ConvCudnnGradGradKernel( ...@@ -517,7 +514,7 @@ void ConvCudnnGradGradKernel(
args2.wdesc.desc(), args2.wdesc.desc(),
ddw, ddw,
args2.cdesc.desc(), args2.cdesc.desc(),
fwd_algo2, fwd_result2.algo,
&beta, &beta,
args2.odesc.desc(), args2.odesc.desc(),
transformed_ddy_channel, transformed_ddy_channel,
...@@ -538,7 +535,7 @@ void ConvCudnnGradGradKernel( ...@@ -538,7 +535,7 @@ void ConvCudnnGradGradKernel(
args2.wdesc.desc(), args2.wdesc.desc(),
ddw + i * group_offset_filter, ddw + i * group_offset_filter,
args2.cdesc.desc(), args2.cdesc.desc(),
fwd_algo2, fwd_result2.algo,
workspace_ptr, workspace_ptr,
workspace_size, workspace_size,
&alpha, &alpha,
...@@ -568,7 +565,7 @@ void ConvCudnnGradGradKernel( ...@@ -568,7 +565,7 @@ void ConvCudnnGradGradKernel(
args3.idesc.desc(), args3.idesc.desc(),
ddx, ddx,
args3.cdesc.desc(), args3.cdesc.desc(),
filter_algo, filter_result.algo,
&beta, &beta,
args3.wdesc.desc(), args3.wdesc.desc(),
dw, dw,
...@@ -589,7 +586,7 @@ void ConvCudnnGradGradKernel( ...@@ -589,7 +586,7 @@ void ConvCudnnGradGradKernel(
args3.odesc.desc(), args3.odesc.desc(),
transformed_dy_channel + i * group_offset_out, transformed_dy_channel + i * group_offset_out,
args3.cdesc.desc(), args3.cdesc.desc(),
filter_algo, filter_result.algo,
workspace_ptr, workspace_ptr,
workspace_size, workspace_size,
&beta, &beta,
...@@ -615,7 +612,7 @@ void ConvCudnnGradGradKernel( ...@@ -615,7 +612,7 @@ void ConvCudnnGradGradKernel(
args4.wdesc.desc(), args4.wdesc.desc(),
ddw, ddw,
args4.cdesc.desc(), args4.cdesc.desc(),
data_algo, data_result.algo,
&beta, &beta,
args4.idesc.desc(), args4.idesc.desc(),
transformed_dx, transformed_dx,
...@@ -636,7 +633,7 @@ void ConvCudnnGradGradKernel( ...@@ -636,7 +633,7 @@ void ConvCudnnGradGradKernel(
args4.odesc.desc(), args4.odesc.desc(),
transformed_dy_channel + i * group_offset_out, transformed_dy_channel + i * group_offset_out,
args4.cdesc.desc(), args4.cdesc.desc(),
data_algo, data_result.algo,
workspace_ptr, workspace_ptr,
workspace_size, workspace_size,
&beta, &beta,
......
paddle/phi/kernels/gpudnn/conv_grad_kernel.cu
浏览文件 @ b937cdc5
...@@ -322,17 +322,16 @@ void ConvCudnnGradKernel(const Context& ctx, ...@@ -322,17 +322,16 @@ void ConvCudnnGradKernel(const Context& ctx,
int group_offset_in = i_c / groups * i_h * i_w * i_d; int group_offset_in = i_c / groups * i_h * i_w * i_d;
int group_offset_out = o_c / groups * o_h * o_w * o_d; int group_offset_out = o_c / groups * o_h * o_w * o_d;
int group_offset_filter = transformed_filter_channel.numel() / groups; int group_offset_filter = transformed_filter_channel.numel() / groups;
// ------------------- cudnn backward algorithm --------------------- // ------------------- cudnn backward algorithm ---------------------
#ifdef PADDLE_WITH_HIP #ifdef PADDLE_WITH_HIP
miopenConvBwdDataAlgorithm_t data_algo = paddle::operators::SearchResult<miopenConvBwdDataAlgorithm_t> bwd_result;
static_cast<miopenConvBwdDataAlgorithm_t>(0); paddle::operators::SearchResult<miopenConvBwdWeightsAlgorithm_t>
miopenConvBwdWeightsAlgorithm_t filter_algo = filter_result;
static_cast<miopenConvBwdWeightsAlgorithm_t>(0);
#else #else
cudnnConvolutionBwdDataAlgo_t data_algo = paddle::operators::SearchResult<cudnnConvolutionBwdDataAlgo_t> bwd_result;
static_cast<cudnnConvolutionBwdDataAlgo_t>(0); paddle::operators::SearchResult<cudnnConvolutionBwdFilterAlgo_t>
cudnnConvolutionBwdFilterAlgo_t filter_algo = filter_result;
static_cast<cudnnConvolutionBwdFilterAlgo_t>(0);
#endif #endif
// input data workspace_size // input data workspace_size
size_t workspace_size_d = 0; size_t workspace_size_d = 0;
...@@ -368,14 +367,14 @@ void ConvCudnnGradKernel(const Context& ctx, ...@@ -368,14 +367,14 @@ void ConvCudnnGradKernel(const Context& ctx,
paddle::operators::SearchAlgorithm<miopenConvBwdDataAlgorithm_t>; paddle::operators::SearchAlgorithm<miopenConvBwdDataAlgorithm_t>;
workspace_size_d = workspace_size_d =
std::max(workspace_size_d, search1::GetWorkspaceSize(args1)); std::max(workspace_size_d, search1::GetWorkspaceSize(args1));
data_algo = search1::Find<T>( bwd_result.algo = search1::Find<T>(
args1, exhaustive_search, deterministic, workspace_size_d, ctx); args1, exhaustive_search, deterministic, workspace_size_d, ctx);
#else #else
using search1 = using search1 =
paddle::operators::SearchAlgorithm<cudnnConvolutionBwdDataAlgoPerf_t>; paddle::operators::SearchAlgorithm<cudnnConvolutionBwdDataAlgoPerf_t>;
data_algo = search1::Find<T>(args1, exhaustive_search, deterministic, ctx); bwd_result = search1::Find<T>(args1, exhaustive_search, deterministic, ctx);
workspace_size_d = workspace_size_d = std::max(
std::max(workspace_size_d, search1::GetWorkspaceSize(args1, data_algo)); workspace_size_d, search1::GetWorkspaceSize(args1, bwd_result.algo));
#endif #endif
} }
...@@ -397,15 +396,17 @@ void ConvCudnnGradKernel(const Context& ctx, ...@@ -397,15 +396,17 @@ void ConvCudnnGradKernel(const Context& ctx,
paddle::operators::SearchAlgorithm<miopenConvBwdWeightsAlgorithm_t>; paddle::operators::SearchAlgorithm<miopenConvBwdWeightsAlgorithm_t>;
workspace_size_w = workspace_size_w =
std::max(workspace_size_w, search2::GetWorkspaceSize(args2)); std::max(workspace_size_w, search2::GetWorkspaceSize(args2));
filter_algo = search2::Find<T>( filter_result.algo = search2::Find<T>(
args2, exhaustive_search, deterministic, workspace_size_w, ctx); args2, exhaustive_search, deterministic, workspace_size_w, ctx);
#else #else
using search2 = using search2 =
paddle::operators::SearchAlgorithm<cudnnConvolutionBwdFilterAlgoPerf_t>; paddle::operators::SearchAlgorithm<cudnnConvolutionBwdFilterAlgoPerf_t>;
filter_algo = filter_result =
search2::Find<T>(args2, exhaustive_search, deterministic, ctx); search2::Find<T>(args2, exhaustive_search, deterministic, ctx);
workspace_size_w = std::max(workspace_size_w, VLOG(3) << "filter algo: " << filter_result.algo << ", time "
search2::GetWorkspaceSize(args2, filter_algo)); << filter_result.time;
workspace_size_w = std::max(
workspace_size_w, search2::GetWorkspaceSize(args2, filter_result.algo));
#endif #endif
} }
...@@ -439,7 +440,7 @@ void ConvCudnnGradKernel(const Context& ctx, ...@@ -439,7 +440,7 @@ void ConvCudnnGradKernel(const Context& ctx,
args1.wdesc.desc(), args1.wdesc.desc(),
filter_data, filter_data,
args1.cdesc.desc(), args1.cdesc.desc(),
data_algo, bwd_result.algo,
&beta, &beta,
args1.idesc.desc(), args1.idesc.desc(),
temp_tensor_data, temp_tensor_data,
...@@ -471,7 +472,7 @@ void ConvCudnnGradKernel(const Context& ctx, ...@@ -471,7 +472,7 @@ void ConvCudnnGradKernel(const Context& ctx,
args1.wdesc.desc(), args1.wdesc.desc(),
filter_data, filter_data,
args1.cdesc.desc(), args1.cdesc.desc(),
data_algo, bwd_result.algo,
&beta, &beta,
args1.idesc.desc(), args1.idesc.desc(),
transformed_input_grad_data, transformed_input_grad_data,
...@@ -494,7 +495,7 @@ void ConvCudnnGradKernel(const Context& ctx, ...@@ -494,7 +495,7 @@ void ConvCudnnGradKernel(const Context& ctx,
args1.odesc.desc(), args1.odesc.desc(),
output_grad_data + i * group_offset_out, output_grad_data + i * group_offset_out,
args1.cdesc.desc(), args1.cdesc.desc(),
data_algo, bwd_result.algo,
cudnn_workspace_ptr, cudnn_workspace_ptr,
workspace_size_d, workspace_size_d,
&beta, &beta,
...@@ -554,7 +555,7 @@ void ConvCudnnGradKernel(const Context& ctx, ...@@ -554,7 +555,7 @@ void ConvCudnnGradKernel(const Context& ctx,
args2.idesc.desc(), args2.idesc.desc(),
input_data, input_data,
args2.cdesc.desc(), args2.cdesc.desc(),
filter_algo, filter_result.algo,
&beta, &beta,
args2.wdesc.desc(), args2.wdesc.desc(),
filter_grad_data, filter_grad_data,
...@@ -575,7 +576,7 @@ void ConvCudnnGradKernel(const Context& ctx, ...@@ -575,7 +576,7 @@ void ConvCudnnGradKernel(const Context& ctx,
args2.odesc.desc(), args2.odesc.desc(),
output_grad_data + i * group_offset_out, output_grad_data + i * group_offset_out,
args2.cdesc.desc(), args2.cdesc.desc(),
filter_algo, filter_result.algo,
cudnn_workspace_ptr, cudnn_workspace_ptr,
workspace_size_w, workspace_size_w,
&beta_filter, &beta_filter,
......
paddle/phi/kernels/gpudnn/conv_kernel.cu
浏览文件 @ b937cdc5
...@@ -18,7 +18,6 @@ ...@@ -18,7 +18,6 @@
#include "paddle/phi/backends/gpu/gpu_context.h" #include "paddle/phi/backends/gpu/gpu_context.h"
#include "paddle/phi/core/kernel_registry.h" #include "paddle/phi/core/kernel_registry.h"
#include "paddle/fluid/framework/eigen.h"
#ifdef PADDLE_WITH_HIP #ifdef PADDLE_WITH_HIP
#include "paddle/fluid/operators/conv_miopen_helper.h" #include "paddle/fluid/operators/conv_miopen_helper.h"
#else #else
...@@ -68,7 +67,6 @@ void ConvCudnnKernel(const Context& ctx, ...@@ -68,7 +67,6 @@ void ConvCudnnKernel(const Context& ctx,
"FLAGS_cudnn_deterministic True at same time.")); "FLAGS_cudnn_deterministic True at same time."));
const bool channel_last = (data_format == "NHWC" || data_format == "NDHWC"); const bool channel_last = (data_format == "NHWC" || data_format == "NDHWC");
auto dtype = paddle::platform::CudnnDataType<T>::type; auto dtype = paddle::platform::CudnnDataType<T>::type;
#ifdef PADDLE_WITH_HIP #ifdef PADDLE_WITH_HIP
...@@ -309,17 +307,17 @@ void ConvCudnnKernel(const Context& ctx, ...@@ -309,17 +307,17 @@ void ConvCudnnKernel(const Context& ctx,
size_t workspace_size = 0; // final workspace to allocate. size_t workspace_size = 0; // final workspace to allocate.
// ------------------- cudnn conv algorithm --------------------- // ------------------- cudnn conv algorithm ---------------------
#ifdef PADDLE_WITH_HIP #ifdef PADDLE_WITH_HIP
miopenConvFwdAlgorithm_t algo{}; paddle::operators::SearchResult<miopenConvFwdAlgorithm_t> fwd_result;
using search = paddle::operators::SearchAlgorithm<miopenConvFwdAlgorithm_t>; using search = paddle::operators::SearchAlgorithm<miopenConvFwdAlgorithm_t>;
workspace_size = search::GetWorkspaceSize(args); workspace_size = search::GetWorkspaceSize(args);
algo = search::Find<T>( fwd_result.algo = search::Find<T>(
args, exhaustive_search, deterministic, workspace_size, ctx); args, exhaustive_search, deterministic, workspace_size, ctx);
#else #else
cudnnConvolutionFwdAlgo_t algo{}; paddle::operators::SearchResult<cudnnConvolutionFwdAlgo_t> fwd_result;
using search = using search =
paddle::operators::SearchAlgorithm<cudnnConvolutionFwdAlgoPerf_t>; paddle::operators::SearchAlgorithm<cudnnConvolutionFwdAlgoPerf_t>;
algo = search::Find<T>(args, exhaustive_search, deterministic, ctx); fwd_result = search::Find<T>(args, exhaustive_search, deterministic, ctx);
workspace_size = search::GetWorkspaceSize(args, algo); workspace_size = search::GetWorkspaceSize(args, fwd_result.algo);
#endif #endif
#if defined(PADDLE_WITH_CUDA) && CUDNN_VERSION_MIN(7, 0, 1) #if defined(PADDLE_WITH_CUDA) && CUDNN_VERSION_MIN(7, 0, 1)
...@@ -328,7 +326,7 @@ void ConvCudnnKernel(const Context& ctx, ...@@ -328,7 +326,7 @@ void ConvCudnnKernel(const Context& ctx,
// in forward computation, so change the algorithm to CUDNN_CONVOLUTION_\ // in forward computation, so change the algorithm to CUDNN_CONVOLUTION_\
// FWD_ALGO_IMPLICIT_GEMM manually. // FWD_ALGO_IMPLICIT_GEMM manually.
if (groups > 1) { if (groups > 1) {
algo = static_cast<cudnnConvolutionFwdAlgo_t>(0); fwd_result.algo = static_cast<cudnnConvolutionFwdAlgo_t>(0);
} }
#endif #endif
...@@ -352,7 +350,7 @@ void ConvCudnnKernel(const Context& ctx, ...@@ -352,7 +350,7 @@ void ConvCudnnKernel(const Context& ctx,
args.wdesc.desc(), args.wdesc.desc(),
filter_data, filter_data,
args.cdesc.desc(), args.cdesc.desc(),
algo, fwd_result.algo,
&beta, &beta,
args.odesc.desc(), args.odesc.desc(),
output_data, output_data,
...@@ -373,7 +371,7 @@ void ConvCudnnKernel(const Context& ctx, ...@@ -373,7 +371,7 @@ void ConvCudnnKernel(const Context& ctx,
args.wdesc.desc(), args.wdesc.desc(),
filter_data + i * group_offset_filter, filter_data + i * group_offset_filter,
args.cdesc.desc(), args.cdesc.desc(),
algo, fwd_result.algo,
workspace_ptr, workspace_ptr,
workspace_size, workspace_size,
&beta, &beta,
......
paddle/phi/kernels/gpudnn/conv_transpose_grad_kernel.cu
浏览文件 @ b937cdc5
...@@ -188,11 +188,13 @@ void ConvTransposeGradRawGPUDNNKernel(const Context& ctx, ...@@ -188,11 +188,13 @@ void ConvTransposeGradRawGPUDNNKernel(const Context& ctx,
dtype}; dtype};
#ifdef PADDLE_WITH_HIP #ifdef PADDLE_WITH_HIP
miopenConvFwdAlgorithm_t data_algo{}; paddle::operators::SearchResult<miopenConvFwdAlgorithm_t> fwd_result;
miopenConvBwdWeightsAlgorithm_t filter_algo{}; paddle::operators::SearchResult<miopenConvBwdWeightsAlgorithm_t>
filter_result;
#else #else
cudnnConvolutionFwdAlgo_t data_algo{}; paddle::operators::SearchResult<cudnnConvolutionFwdAlgo_t> fwd_result;
cudnnConvolutionBwdFilterAlgo_t filter_algo{}; paddle::operators::SearchResult<cudnnConvolutionBwdFilterAlgo_t>
filter_result;
#endif #endif
auto layout_tensor = paddle::platform::GetCudnnTensorFormat(layout); auto layout_tensor = paddle::platform::GetCudnnTensorFormat(layout);
...@@ -218,14 +220,14 @@ void ConvTransposeGradRawGPUDNNKernel(const Context& ctx, ...@@ -218,14 +220,14 @@ void ConvTransposeGradRawGPUDNNKernel(const Context& ctx,
using search1 = using search1 =
paddle::operators::SearchAlgorithm<miopenConvFwdAlgorithm_t>; paddle::operators::SearchAlgorithm<miopenConvFwdAlgorithm_t>;
workspace_size = std::max(workspace_size, search1::GetWorkspaceSize(args1)); workspace_size = std::max(workspace_size, search1::GetWorkspaceSize(args1));
data_algo = fwd_result.algo =
search1::Find<T>(args1, false, deterministic, workspace_size, ctx); search1::Find<T>(args1, false, deterministic, workspace_size, ctx);
#else #else
using search1 = using search1 =
paddle::operators::SearchAlgorithm<cudnnConvolutionFwdAlgoPerf_t>; paddle::operators::SearchAlgorithm<cudnnConvolutionFwdAlgoPerf_t>;
data_algo = search1::Find<T>(args1, false, deterministic, ctx); fwd_result = search1::Find<T>(args1, false, deterministic, ctx);
workspace_size = workspace_size = std::max(
std::max(workspace_size, search1::GetWorkspaceSize(args1, data_algo)); workspace_size, search1::GetWorkspaceSize(args1, fwd_result.algo));
#endif #endif
} }
...@@ -245,14 +247,14 @@ void ConvTransposeGradRawGPUDNNKernel(const Context& ctx, ...@@ -245,14 +247,14 @@ void ConvTransposeGradRawGPUDNNKernel(const Context& ctx,
using search2 = using search2 =
paddle::operators::SearchAlgorithm<miopenConvBwdWeightsAlgorithm_t>; paddle::operators::SearchAlgorithm<miopenConvBwdWeightsAlgorithm_t>;
workspace_size = std::max(workspace_size, search2::GetWorkspaceSize(args2)); workspace_size = std::max(workspace_size, search2::GetWorkspaceSize(args2));
filter_algo = filter_result.algo =
search2::Find<T>(args2, false, deterministic, workspace_size, ctx); search2::Find<T>(args2, false, deterministic, workspace_size, ctx);
#else #else
using search2 = using search2 =
paddle::operators::SearchAlgorithm<cudnnConvolutionBwdFilterAlgoPerf_t>; paddle::operators::SearchAlgorithm<cudnnConvolutionBwdFilterAlgoPerf_t>;
filter_algo = search2::Find<T>(args2, false, deterministic, ctx); filter_result = search2::Find<T>(args2, false, deterministic, ctx);
workspace_size = workspace_size = std::max(
std::max(workspace_size, search2::GetWorkspaceSize(args2, filter_algo)); workspace_size, search2::GetWorkspaceSize(args2, filter_result.algo));
#endif #endif
} }
...@@ -278,7 +280,7 @@ void ConvTransposeGradRawGPUDNNKernel(const Context& ctx, ...@@ -278,7 +280,7 @@ void ConvTransposeGradRawGPUDNNKernel(const Context& ctx,
args1.wdesc.desc(), args1.wdesc.desc(),
filter_data + filter_offset * g, filter_data + filter_offset * g,
args1.cdesc.desc(), args1.cdesc.desc(),
data_algo, fwd_result.algo,
&beta, &beta,
args1.odesc.desc(), args1.odesc.desc(),
dx_data + x_offset * g, dx_data + x_offset * g,
...@@ -295,7 +297,7 @@ void ConvTransposeGradRawGPUDNNKernel(const Context& ctx, ...@@ -295,7 +297,7 @@ void ConvTransposeGradRawGPUDNNKernel(const Context& ctx,
args1.wdesc.desc(), args1.wdesc.desc(),
filter_data + filter_offset * g, filter_data + filter_offset * g,
args1.cdesc.desc(), args1.cdesc.desc(),
data_algo, fwd_result.algo,
cudnn_workspace, cudnn_workspace,
workspace_size, workspace_size,
&beta, &beta,
...@@ -338,7 +340,7 @@ void ConvTransposeGradRawGPUDNNKernel(const Context& ctx, ...@@ -338,7 +340,7 @@ void ConvTransposeGradRawGPUDNNKernel(const Context& ctx,
args2.idesc.desc(), args2.idesc.desc(),
dout_data + dout_offset * g, dout_data + dout_offset * g,
args2.cdesc.desc(), args2.cdesc.desc(),
filter_algo, filter_result.algo,
&beta, &beta,
args2.wdesc.desc(), args2.wdesc.desc(),
dfilter_data + filter_offset * g, dfilter_data + filter_offset * g,
...@@ -355,7 +357,7 @@ void ConvTransposeGradRawGPUDNNKernel(const Context& ctx, ...@@ -355,7 +357,7 @@ void ConvTransposeGradRawGPUDNNKernel(const Context& ctx,
args2.odesc.desc(), args2.odesc.desc(),
x_data + x_offset * g, x_data + x_offset * g,
args2.cdesc.desc(), args2.cdesc.desc(),
filter_algo, filter_result.algo,
cudnn_workspace, cudnn_workspace,
workspace_size, workspace_size,
&beta, &beta,
...@@ -653,22 +655,17 @@ void Conv2dTransposeDoubleGradGPUDNNKernel( ...@@ -653,22 +655,17 @@ void Conv2dTransposeDoubleGradGPUDNNKernel(
dilations_, dilations_,
dtype}; dtype};
#ifdef PADDLE_WITH_HIP #ifdef PADDLE_WITH_HIP
miopenConvBwdDataAlgorithm_t bwd_algo1 = paddle::operators::SearchResult<miopenConvBwdDataAlgorithm_t> bwd_result1;
static_cast<miopenConvBwdDataAlgorithm_t>(0); paddle::operators::SearchResult<miopenConvBwdDataAlgorithm_t> bwd_result2;
miopenConvBwdDataAlgorithm_t bwd_algo2 = paddle::operators::SearchResult<miopenConvBwdWeightsAlgorithm_t>
static_cast<miopenConvBwdDataAlgorithm_t>(0); filter_result;
miopenConvFwdAlgorithm_t data_algo = static_cast<miopenConvFwdAlgorithm_t>(0); paddle::operators::SearchResult<miopenConvFwdAlgorithm_t> fwd_result;
miopenConvBwdWeightsAlgorithm_t filter_algo =
static_cast<miopenConvBwdWeightsAlgorithm_t>(0);
#else #else
cudnnConvolutionBwdDataAlgo_t bwd_algo1 = paddle::operators::SearchResult<cudnnConvolutionBwdDataAlgo_t> bwd_result1;
static_cast<cudnnConvolutionBwdDataAlgo_t>(0); paddle::operators::SearchResult<cudnnConvolutionBwdDataAlgo_t> bwd_result2;
cudnnConvolutionBwdDataAlgo_t bwd_algo2 = paddle::operators::SearchResult<cudnnConvolutionBwdFilterAlgo_t>
static_cast<cudnnConvolutionBwdDataAlgo_t>(0); filter_result;
cudnnConvolutionFwdAlgo_t data_algo = paddle::operators::SearchResult<cudnnConvolutionFwdAlgo_t> fwd_result;
static_cast<cudnnConvolutionFwdAlgo_t>(0);
cudnnConvolutionBwdFilterAlgo_t filter_algo =
static_cast<cudnnConvolutionBwdFilterAlgo_t>(0);
#endif #endif
auto layout = paddle::platform::GetCudnnTensorFormat(GPUDNNDataLayout::kNCHW); auto layout = paddle::platform::GetCudnnTensorFormat(GPUDNNDataLayout::kNCHW);
...@@ -696,13 +693,13 @@ void Conv2dTransposeDoubleGradGPUDNNKernel( ...@@ -696,13 +693,13 @@ void Conv2dTransposeDoubleGradGPUDNNKernel(
using search1 = using search1 =
paddle::operators::SearchAlgorithm<miopenConvBwdDataAlgorithm_t>; paddle::operators::SearchAlgorithm<miopenConvBwdDataAlgorithm_t>;
workspace_size = search1::GetWorkspaceSize(args1); workspace_size = search1::GetWorkspaceSize(args1);
bwd_algo1 = bwd_result1.algo =
search1::Find<T>(args1, false, deterministic, workspace_size, ctx); search1::Find<T>(args1, false, deterministic, workspace_size, ctx);
#else #else
using search1 = using search1 =
paddle::operators::SearchAlgorithm<cudnnConvolutionBwdDataAlgoPerf_t>; paddle::operators::SearchAlgorithm<cudnnConvolutionBwdDataAlgoPerf_t>;
bwd_algo1 = search1::Find<T>(args1, false, deterministic, ctx); bwd_result1 = search1::Find<T>(args1, false, deterministic, ctx);
workspace_size = search1::GetWorkspaceSize(args1, bwd_algo1); workspace_size = search1::GetWorkspaceSize(args1, bwd_result1.algo);
#endif #endif
ddfilter_ = ddfilter.data<T>(); ddfilter_ = ddfilter.data<T>();
...@@ -720,14 +717,14 @@ void Conv2dTransposeDoubleGradGPUDNNKernel( ...@@ -720,14 +717,14 @@ void Conv2dTransposeDoubleGradGPUDNNKernel(
using search2 = using search2 =
paddle::operators::SearchAlgorithm<miopenConvBwdDataAlgorithm_t>; paddle::operators::SearchAlgorithm<miopenConvBwdDataAlgorithm_t>;
workspace_size = std::max(workspace_size, search2::GetWorkspaceSize(args2)); workspace_size = std::max(workspace_size, search2::GetWorkspaceSize(args2));
bwd_algo2 = bwd_result2.algo =
search2::Find<T>(args2, false, deterministic, workspace_size, ctx); search2::Find<T>(args2, false, deterministic, workspace_size, ctx);
#else #else
using search2 = using search2 =
paddle::operators::SearchAlgorithm<cudnnConvolutionBwdDataAlgoPerf_t>; paddle::operators::SearchAlgorithm<cudnnConvolutionBwdDataAlgoPerf_t>;
bwd_algo2 = search2::Find<T>(args2, false, deterministic, ctx); bwd_result2 = search2::Find<T>(args2, false, deterministic, ctx);
workspace_size = workspace_size = std::max(
std::max(workspace_size, search2::GetWorkspaceSize(args2, bwd_algo2)); workspace_size, search2::GetWorkspaceSize(args2, bwd_result2.algo));
#endif #endif
} }
...@@ -736,9 +733,7 @@ void Conv2dTransposeDoubleGradGPUDNNKernel( ...@@ -736,9 +733,7 @@ void Conv2dTransposeDoubleGradGPUDNNKernel(
args3.handle = handle; args3.handle = handle;
args3.idesc.set(transformed_dout, iwo_group); args3.idesc.set(transformed_dout, iwo_group);
args3.wdesc.set(*dfilter, layout, iwo_group); args3.wdesc.set(*dfilter, layout, iwo_group);
args3.odesc.set(transformed_ddx_channel, iwo_group); args3.odesc.set(transformed_ddx_channel, iwo_group);
args3.cdesc.set(dtype, args3.cdesc.set(dtype,
padding_common, padding_common,
strides, strides,
...@@ -749,14 +744,14 @@ void Conv2dTransposeDoubleGradGPUDNNKernel( ...@@ -749,14 +744,14 @@ void Conv2dTransposeDoubleGradGPUDNNKernel(
using search3 = using search3 =
paddle::operators::SearchAlgorithm<miopenConvBwdWeightsAlgorithm_t>; paddle::operators::SearchAlgorithm<miopenConvBwdWeightsAlgorithm_t>;
workspace_size = std::max(workspace_size, search3::GetWorkspaceSize(args3)); workspace_size = std::max(workspace_size, search3::GetWorkspaceSize(args3));
filter_algo = filter_result.algo =
search3::Find<T>(args3, false, deterministic, workspace_size, ctx); search3::Find<T>(args3, false, deterministic, workspace_size, ctx);
#else #else
using search3 = using search3 =
paddle::operators::SearchAlgorithm<cudnnConvolutionBwdFilterAlgoPerf_t>; paddle::operators::SearchAlgorithm<cudnnConvolutionBwdFilterAlgoPerf_t>;
filter_algo = search3::Find<T>(args3, false, deterministic, ctx); filter_result = search3::Find<T>(args3, false, deterministic, ctx);
workspace_size = workspace_size = std::max(
std::max(workspace_size, search3::GetWorkspaceSize(args3, filter_algo)); workspace_size, search3::GetWorkspaceSize(args3, filter_result.algo));
#endif #endif
} }
...@@ -777,14 +772,14 @@ void Conv2dTransposeDoubleGradGPUDNNKernel( ...@@ -777,14 +772,14 @@ void Conv2dTransposeDoubleGradGPUDNNKernel(
using search4 = using search4 =
paddle::operators::SearchAlgorithm<miopenConvFwdAlgorithm_t>; paddle::operators::SearchAlgorithm<miopenConvFwdAlgorithm_t>;
workspace_size = std::max(workspace_size, search4::GetWorkspaceSize(args4)); workspace_size = std::max(workspace_size, search4::GetWorkspaceSize(args4));
data_algo = fwd_result.algo =
search4::Find<T>(args4, false, deterministic, workspace_size, ctx); search4::Find<T>(args4, false, deterministic, workspace_size, ctx);
#else #else
using search4 = using search4 =
paddle::operators::SearchAlgorithm<cudnnConvolutionFwdAlgoPerf_t>; paddle::operators::SearchAlgorithm<cudnnConvolutionFwdAlgoPerf_t>;
data_algo = search4::Find<T>(args4, false, deterministic, ctx); fwd_result = search4::Find<T>(args4, false, deterministic, ctx);
workspace_size = workspace_size = std::max(
std::max(workspace_size, search4::GetWorkspaceSize(args4, data_algo)); workspace_size, search4::GetWorkspaceSize(args4, fwd_result.algo));
#endif #endif
} }
...@@ -831,7 +826,7 @@ void Conv2dTransposeDoubleGradGPUDNNKernel( ...@@ -831,7 +826,7 @@ void Conv2dTransposeDoubleGradGPUDNNKernel(
args1.wdesc.desc(), args1.wdesc.desc(),
filter_ + i * group_offset_filter, filter_ + i * group_offset_filter,
args1.cdesc.desc(), args1.cdesc.desc(),
bwd_algo1, bwd_result1.algo,
&beta, &beta,
args1.idesc.desc(), args1.idesc.desc(),
transformed_ddout_channel_ + i * group_offset_out, transformed_ddout_channel_ + i * group_offset_out,
...@@ -850,7 +845,7 @@ void Conv2dTransposeDoubleGradGPUDNNKernel( ...@@ -850,7 +845,7 @@ void Conv2dTransposeDoubleGradGPUDNNKernel(
args1.odesc.desc(), args1.odesc.desc(),
ddx_ + i * group_offset_in, ddx_ + i * group_offset_in,
args1.cdesc.desc(), args1.cdesc.desc(),
bwd_algo1, bwd_result1.algo,
workspace_ptr, workspace_ptr,
workspace_size, workspace_size,
&beta, &beta,
...@@ -877,7 +872,7 @@ void Conv2dTransposeDoubleGradGPUDNNKernel( ...@@ -877,7 +872,7 @@ void Conv2dTransposeDoubleGradGPUDNNKernel(
args2.wdesc.desc(), args2.wdesc.desc(),
ddfilter_ + i * group_offset_filter, ddfilter_ + i * group_offset_filter,
args2.cdesc.desc(), args2.cdesc.desc(),
bwd_algo2, bwd_result2.algo,
&beta, &beta,
args2.idesc.desc(), args2.idesc.desc(),
conv_x_ddfilter_data + i * group_offset_out, conv_x_ddfilter_data + i * group_offset_out,
...@@ -908,7 +903,7 @@ void Conv2dTransposeDoubleGradGPUDNNKernel( ...@@ -908,7 +903,7 @@ void Conv2dTransposeDoubleGradGPUDNNKernel(
args2.odesc.desc(), args2.odesc.desc(),
x_ + i * group_offset_in, x_ + i * group_offset_in,
args2.cdesc.desc(), args2.cdesc.desc(),
bwd_algo2, bwd_result2.algo,
workspace_ptr, workspace_ptr,
workspace_size, workspace_size,
&alpha, &alpha,
...@@ -964,7 +959,7 @@ void Conv2dTransposeDoubleGradGPUDNNKernel( ...@@ -964,7 +959,7 @@ void Conv2dTransposeDoubleGradGPUDNNKernel(
args3.idesc.desc(), args3.idesc.desc(),
transformed_dout_channel_ + i * group_offset_out, transformed_dout_channel_ + i * group_offset_out,
args3.cdesc.desc(), args3.cdesc.desc(),
filter_algo, filter_result.algo,
&beta, &beta,
args3.wdesc.desc(), args3.wdesc.desc(),
dfilter_ + i * group_offset_filter, dfilter_ + i * group_offset_filter,
...@@ -983,7 +978,7 @@ void Conv2dTransposeDoubleGradGPUDNNKernel( ...@@ -983,7 +978,7 @@ void Conv2dTransposeDoubleGradGPUDNNKernel(
args3.odesc.desc(), args3.odesc.desc(),
ddx_ + i * group_offset_in, ddx_ + i * group_offset_in,
args3.cdesc.desc(), args3.cdesc.desc(),
filter_algo, filter_result.algo,
workspace_ptr, workspace_ptr,
workspace_size, workspace_size,
&beta, &beta,
...@@ -1009,7 +1004,7 @@ void Conv2dTransposeDoubleGradGPUDNNKernel( ...@@ -1009,7 +1004,7 @@ void Conv2dTransposeDoubleGradGPUDNNKernel(
args4.wdesc.desc(), args4.wdesc.desc(),
ddfilter_ + i * group_offset_filter, ddfilter_ + i * group_offset_filter,
args4.cdesc.desc(), args4.cdesc.desc(),
data_algo, fwd_result.algo,
&beta, &beta,
args4.odesc.desc(), args4.odesc.desc(),
transformed_dx_ + i * group_offset_in, transformed_dx_ + i * group_offset_in,
...@@ -1028,7 +1023,7 @@ void Conv2dTransposeDoubleGradGPUDNNKernel( ...@@ -1028,7 +1023,7 @@ void Conv2dTransposeDoubleGradGPUDNNKernel(
args4.wdesc.desc(), args4.wdesc.desc(),
ddfilter_ + i * group_offset_filter, ddfilter_ + i * group_offset_filter,
args4.cdesc.desc(), args4.cdesc.desc(),
data_algo, fwd_result.algo,
workspace_ptr, workspace_ptr,
workspace_size, workspace_size,
&beta, &beta,
......
paddle/phi/kernels/gpudnn/conv_transpose_kernel.cu
浏览文件 @ b937cdc5
...@@ -217,16 +217,19 @@ void ConvTransposeRawGPUDNNKernel(const Context& ctx, ...@@ -217,16 +217,19 @@ void ConvTransposeRawGPUDNNKernel(const Context& ctx,
c_groups); c_groups);
#ifdef PADDLE_WITH_HIP #ifdef PADDLE_WITH_HIP
paddle::operators::SearchResult<miopenConvBwdDataAlgorithm_t> bwd_result;
using search = using search =
paddle::operators::SearchAlgorithm<miopenConvBwdDataAlgorithm_t>; paddle::operators::SearchAlgorithm<miopenConvBwdDataAlgorithm_t>;
workspace_size = std::max(workspace_size, search::GetWorkspaceSize(args)); workspace_size = std::max(workspace_size, search::GetWorkspaceSize(args));
algo = search::Find<T>(args, false, deterministic, workspace_size, ctx); bwd_result.algo =
search::Find<T>(args, false, deterministic, workspace_size, ctx);
#else #else
paddle::operators::SearchResult<cudnnConvolutionBwdDataAlgo_t> bwd_result;
using search = using search =
paddle::operators::SearchAlgorithm<cudnnConvolutionBwdDataAlgoPerf_t>; paddle::operators::SearchAlgorithm<cudnnConvolutionBwdDataAlgoPerf_t>;
algo = search::Find<T>(args, false, deterministic, ctx); bwd_result = search::Find<T>(args, false, deterministic, ctx);
workspace_size = workspace_size =
std::max(workspace_size, search::GetWorkspaceSize(args, algo)); std::max(workspace_size, search::GetWorkspaceSize(args, bwd_result.algo));
#endif #endif
// ------------------- cudnn conv transpose forward --------------------- // ------------------- cudnn conv transpose forward ---------------------
...@@ -247,7 +250,7 @@ void ConvTransposeRawGPUDNNKernel(const Context& ctx, ...@@ -247,7 +250,7 @@ void ConvTransposeRawGPUDNNKernel(const Context& ctx,
args.wdesc.desc(), args.wdesc.desc(),
filter_data + filter_offset * g, filter_data + filter_offset * g,
args.cdesc.desc(), args.cdesc.desc(),
algo, bwd_result.algo,
&beta, &beta,
args.idesc.desc(), args.idesc.desc(),
transformed_out_data + out_offset * g, transformed_out_data + out_offset * g,
...@@ -264,7 +267,7 @@ void ConvTransposeRawGPUDNNKernel(const Context& ctx, ...@@ -264,7 +267,7 @@ void ConvTransposeRawGPUDNNKernel(const Context& ctx,
args.odesc.desc(), args.odesc.desc(),
x_data + x_offset * g, x_data + x_offset * g,
args.cdesc.desc(), args.cdesc.desc(),
algo, bwd_result.algo,
cudnn_workspace, cudnn_workspace,
workspace_size, workspace_size,
&beta, &beta,
......
paddle/phi/kernels/impl/conv_cudnn_impl.h
浏览文件 @ b937cdc5
...@@ -36,7 +36,7 @@ ...@@ -36,7 +36,7 @@
#include "paddle/phi/kernels/funcs/batch_norm_utils.h" #include "paddle/phi/kernels/funcs/batch_norm_utils.h"
DECLARE_bool(cudnn_deterministic); DECLARE_bool(cudnn_deterministic);
DECLARE_uint64(conv_workspace_size_limit); DECLARE_int64(conv_workspace_size_limit);
DECLARE_bool(cudnn_exhaustive_search); DECLARE_bool(cudnn_exhaustive_search);
namespace phi { namespace phi {
......
python/paddle/fluid/tests/unittests/test_switch_autotune.py
浏览文件 @ b937cdc5
...@@ -43,6 +43,16 @@ def static_program(net, data): ...@@ -43,6 +43,16 @@ def static_program(net, data):
return loss return loss
def set_flags(enable_autotune):
if paddle.is_compiled_with_cuda():
if enable_autotune:
paddle.set_flags({'FLAGS_conv_workspace_size_limit': -1})
paddle.set_flags({'FLAGS_cudnn_exhaustive_search': 1})
else:
paddle.set_flags({'FLAGS_conv_workspace_size_limit': 512})
paddle.set_flags({'FLAGS_cudnn_exhaustive_search': 0})
class TestAutoTune(unittest.TestCase): class TestAutoTune(unittest.TestCase):
def test_autotune(self): def test_autotune(self):
paddle.fluid.core.disable_autotune() paddle.fluid.core.disable_autotune()
...@@ -61,6 +71,7 @@ class TestAutoTune(unittest.TestCase): ...@@ -61,6 +71,7 @@ class TestAutoTune(unittest.TestCase):
class TestDygraphAutoTuneStatus(TestAutoTune): class TestDygraphAutoTuneStatus(TestAutoTune):
def run_program(self, enable_autotune): def run_program(self, enable_autotune):
set_flags(enable_autotune)
if enable_autotune: if enable_autotune:
paddle.fluid.core.enable_autotune() paddle.fluid.core.enable_autotune()
else: else:
...@@ -107,6 +118,7 @@ class TestDygraphAutoTuneStatus(TestAutoTune): ...@@ -107,6 +118,7 @@ class TestDygraphAutoTuneStatus(TestAutoTune):
class TestStaticAutoTuneStatus(TestAutoTune): class TestStaticAutoTuneStatus(TestAutoTune):
def run_program(self, enable_autotune): def run_program(self, enable_autotune):
paddle.enable_static() paddle.enable_static()
set_flags(enable_autotune)
if enable_autotune: if enable_autotune:
paddle.fluid.core.enable_autotune() paddle.fluid.core.enable_autotune()
else: else:
......
Markdown is supported
0% 或 .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑！
想要评论请 注册 或