wehub-resource-sync
634 lines
25 KiB
Plaintext
634 lines
25 KiB
Plaintext
/* 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. */
|
|
|
|
#include "paddle/phi/kernels/conv_transpose_kernel.h"
|
|
|
|
#include <algorithm>
|
|
|
|
#include "paddle/common/ddim.h"
|
|
#include "paddle/phi/backends/context_pool.h"
|
|
#include "paddle/phi/backends/dynload/cudnn.h"
|
|
#include "paddle/phi/core/kernel_registry.h"
|
|
#include "paddle/phi/kernels/cpu/conv_util.h"
|
|
#include "paddle/phi/kernels/funcs/padding.h"
|
|
#include "paddle/phi/kernels/funcs/slice.h"
|
|
#include "paddle/phi/kernels/gpudnn/conv_gpudnn.h"
|
|
#include "paddle/phi/kernels/transpose_kernel.h"
|
|
|
|
#ifdef PADDLE_WITH_HIP
|
|
#include "paddle/phi/backends/gpu/rocm/miopen_helper.h"
|
|
#include "paddle/phi/kernels/gpudnn/conv_miopen_helper.h"
|
|
#else
|
|
#include "paddle/phi/backends/gpu/cuda/cudnn_helper.h"
|
|
#include "paddle/phi/backends/gpu/cuda/cudnn_workspace_helper.h"
|
|
#include "paddle/phi/kernels/gpudnn/conv_cudnn_v7.h"
|
|
#endif
|
|
#include "paddle/phi/kernels/full_kernel.h"
|
|
|
|
#ifdef PADDLE_WITH_CUDNN_FRONTEND
|
|
// clang-format off
|
|
#include "paddle/phi/backends/dynload/cudnn_frontend.h"
|
|
#include "paddle/phi/kernels/autotune/cache.h"
|
|
#include "paddle/phi/kernels/gpudnn/conv_cudnn_frontend.h"
|
|
// clang-format on
|
|
#endif
|
|
|
|
#include "paddle/common/flags.h"
|
|
|
|
COMMON_DECLARE_bool(use_accuracy_compatible_kernel);
|
|
|
|
namespace phi {
|
|
|
|
template <typename T, typename Context>
|
|
void ConvTransposeCudnnKernelImplV7(const DenseTensor* transformed_x,
|
|
const DenseTensor* filter,
|
|
const Context& dev_ctx,
|
|
const std::vector<int>& strides,
|
|
const std::vector<int>& padding_common,
|
|
const std::vector<int>& dilations_,
|
|
DataLayout data_layout,
|
|
DataLayout layout,
|
|
bool exhaustive_search,
|
|
bool deterministic,
|
|
int groups,
|
|
DenseTensor* transformed_out) {
|
|
int iwo_groups = 1;
|
|
int c_groups = groups;
|
|
groups = 1;
|
|
size_t workspace_size = 0;
|
|
|
|
const T* x_data = transformed_x->data<T>();
|
|
const T* filter_data = filter->data<T>();
|
|
T* transformed_out_data = transformed_out->data<T>();
|
|
#ifdef PADDLE_WITH_HIP
|
|
miopenConvBwdDataAlgorithm_t algo{};
|
|
#else
|
|
cudnnConvolutionBwdDataAlgo_t algo{};
|
|
#endif
|
|
// ------------------- cudnn conv algorithm ---------------------
|
|
auto handle = dev_ctx.cudnn_handle();
|
|
auto layout_tensor = phi::backends::gpu::GetCudnnTensorFormat(layout);
|
|
auto dtype = phi::backends::gpu::CudnnDataType<T>::type;
|
|
// ------------------- cudnn descriptors ---------------------
|
|
ConvArgs args{handle,
|
|
transformed_out,
|
|
filter,
|
|
transformed_x,
|
|
strides,
|
|
padding_common,
|
|
dilations_,
|
|
dtype,
|
|
groups,
|
|
data_layout};
|
|
args.idesc.set(*transformed_out, iwo_groups);
|
|
args.wdesc.set(*filter, layout_tensor, iwo_groups);
|
|
args.odesc.set(*transformed_x, iwo_groups);
|
|
args.cdesc.set(dtype,
|
|
padding_common,
|
|
strides,
|
|
dilations_,
|
|
phi::AllowTF32Cudnn(),
|
|
c_groups);
|
|
|
|
#ifdef PADDLE_WITH_HIP
|
|
SearchResult<miopenConvBwdDataAlgorithm_t> bwd_result;
|
|
using search = SearchAlgorithm<miopenConvBwdDataAlgorithm_t>;
|
|
workspace_size = std::max(workspace_size, search::GetWorkspaceSize(args));
|
|
bwd_result.algo = search::Find<T>(
|
|
args, exhaustive_search, deterministic, workspace_size, dev_ctx);
|
|
#else
|
|
SearchResult<cudnnConvolutionBwdDataAlgo_t> bwd_result;
|
|
using search = SearchAlgorithm<ConvKind::kBackwardData>;
|
|
bwd_result =
|
|
search::Find<T>(dev_ctx, args, exhaustive_search, deterministic, false);
|
|
workspace_size =
|
|
std::max(workspace_size, search::GetWorkspaceSize(args, bwd_result.algo));
|
|
#endif
|
|
|
|
// ------------------- cudnn conv transpose forward ---------------------
|
|
int64_t x_offset_64 =
|
|
transformed_x->numel() / transformed_x->dims()[0] / groups;
|
|
int64_t out_offset_64 =
|
|
transformed_out->numel() / transformed_out->dims()[0] / groups;
|
|
int64_t filter_offset_64 = filter->numel() / groups;
|
|
PADDLE_ENFORCE_LE_INT_MAX(x_offset_64, "x_offset");
|
|
PADDLE_ENFORCE_LE_INT_MAX(out_offset_64, "out_offset");
|
|
PADDLE_ENFORCE_LE_INT_MAX(filter_offset_64, "filter_offset");
|
|
int x_offset = static_cast<int>(x_offset_64);
|
|
int out_offset = static_cast<int>(out_offset_64);
|
|
int filter_offset = static_cast<int>(filter_offset_64);
|
|
ScalingParamType<T> alpha = 1.0f;
|
|
ScalingParamType<T> beta = 0.0f;
|
|
auto workspace_handle = dev_ctx.cudnn_workspace_handle();
|
|
#ifdef PADDLE_WITH_HIP
|
|
for (int g = 0; g < groups; g++) {
|
|
auto cudnn_func = [&](void* cudnn_workspace) {
|
|
PADDLE_ENFORCE_GPU_SUCCESS(dynload::miopenConvolutionBackwardData(
|
|
handle,
|
|
&alpha,
|
|
args.odesc.desc(),
|
|
x_data + x_offset * g,
|
|
args.wdesc.desc(),
|
|
filter_data + filter_offset * g,
|
|
args.cdesc.desc(),
|
|
bwd_result.algo,
|
|
&beta,
|
|
args.idesc.desc(),
|
|
transformed_out_data + out_offset * g,
|
|
cudnn_workspace,
|
|
workspace_size));
|
|
};
|
|
workspace_handle.RunFunc(cudnn_func, workspace_size);
|
|
}
|
|
#else
|
|
ConvRunner<T, ConvKind::kBackwardData>::Apply(dev_ctx,
|
|
args,
|
|
bwd_result,
|
|
x_data,
|
|
filter_data,
|
|
transformed_out_data,
|
|
groups,
|
|
out_offset,
|
|
filter_offset,
|
|
x_offset,
|
|
workspace_size,
|
|
&workspace_handle,
|
|
false);
|
|
#endif
|
|
}
|
|
#ifdef PADDLE_WITH_CUDNN_FRONTEND
|
|
template <typename T, typename Context>
|
|
void ConvTransposeCudnnKernelImplV8(const DenseTensor* transformed_x,
|
|
const DenseTensor* filter,
|
|
const Context& dev_ctx,
|
|
const std::vector<int>& strides,
|
|
const std::vector<int>& padding_common,
|
|
const std::vector<int>& dilations_,
|
|
DataLayout data_layout,
|
|
DataLayout layout,
|
|
bool exhaustive_search,
|
|
bool deterministic,
|
|
int groups,
|
|
DenseTensor* transformed_out) {
|
|
auto& plan_cache = phi::autotune::AutoTuneCache::Instance().GetConvV8(
|
|
phi::autotune::AlgorithmType::kConvBackwardDataV8);
|
|
|
|
T* input_data = const_cast<T*>(transformed_x->data<T>());
|
|
T* filter_data = const_cast<T*>(filter->data<T>());
|
|
T* output_data = transformed_out->data<T>();
|
|
cudnnHandle_t handle = const_cast<cudnnHandle_t>(dev_ctx.cudnn_handle());
|
|
auto workspace_handle = dev_ctx.cudnn_workspace_handle();
|
|
|
|
auto layout_format = phi::backends::gpu::GetCudnnTensorFormat(layout);
|
|
auto dtype = phi::backends::gpu::CudnnDataType<T>::type;
|
|
|
|
float alpha = 1.0f;
|
|
float beta = 0.0f;
|
|
|
|
using helper = CudnnFrontendConvHelper;
|
|
auto op_graph = helper::BuildConvOperationGraph<
|
|
CUDNN_BACKEND_OPERATION_CONVOLUTION_BACKWARD_DATA_DESCRIPTOR>(
|
|
transformed_out,
|
|
transformed_x,
|
|
filter,
|
|
layout_format,
|
|
strides,
|
|
padding_common,
|
|
dilations_,
|
|
dtype,
|
|
handle,
|
|
alpha,
|
|
beta);
|
|
if (plan_cache.FindPlan(op_graph, handle)) {
|
|
const cudnn_frontend::ExecutionPlan* cached_plan = nullptr;
|
|
int64_t workspace_size = 0;
|
|
plan_cache.GetPlanAndWorkspaceSize(
|
|
op_graph, &cached_plan, &workspace_size, handle);
|
|
helper::ExecutePlan(handle,
|
|
&workspace_handle,
|
|
output_data,
|
|
input_data,
|
|
filter_data,
|
|
cached_plan->get_raw_desc(),
|
|
workspace_size);
|
|
return;
|
|
}
|
|
|
|
auto plans = helper::FindExecutionPlans(&op_graph,
|
|
exhaustive_search,
|
|
deterministic,
|
|
output_data,
|
|
input_data,
|
|
filter_data,
|
|
handle,
|
|
&workspace_handle,
|
|
transformed_x->dtype());
|
|
|
|
helper::ExecutePlansAndCache(handle,
|
|
&workspace_handle,
|
|
output_data,
|
|
input_data,
|
|
filter_data,
|
|
&plans,
|
|
exhaustive_search,
|
|
op_graph,
|
|
&plan_cache);
|
|
}
|
|
#endif
|
|
|
|
template <typename T, typename Context>
|
|
void ConvTransposeRawGPUDNNKernel(const Context& dev_ctx,
|
|
const DenseTensor& x,
|
|
const DenseTensor& filter,
|
|
const std::vector<int>& strides,
|
|
const std::vector<int>& paddings,
|
|
const std::string& padding_algorithm,
|
|
int groups,
|
|
const std::vector<int>& dilations,
|
|
const std::string& data_format,
|
|
DenseTensor* out) {
|
|
if (x.numel() == 0 || filter.numel() == 0) {
|
|
Full<T, Context>(dev_ctx, out->dims(), 0, out);
|
|
return;
|
|
}
|
|
|
|
bool has_exhaustive_search = dev_ctx.HasDnnAttr("exhaustive_search");
|
|
bool exhaustive_search_attr =
|
|
has_exhaustive_search
|
|
? PADDLE_GET_CONST(bool, dev_ctx.GetDnnAttr("exhaustive_search"))
|
|
: false;
|
|
bool exhaustive_search =
|
|
FLAGS_cudnn_exhaustive_search || exhaustive_search_attr;
|
|
bool deterministic = FLAGS_cudnn_deterministic;
|
|
PADDLE_ENFORCE_EQ(exhaustive_search && deterministic,
|
|
false,
|
|
common::errors::InvalidArgument(
|
|
"Can't set exhaustive_search True and "
|
|
"FLAGS_cudnn_deterministic True at same time."));
|
|
|
|
std::vector<int> paddings_ = paddings;
|
|
std::vector<int> dilations_ = dilations;
|
|
const DataLayout data_layout =
|
|
(data_format != "NHWC" ? DataLayout::NCHW : DataLayout::NHWC);
|
|
std::vector<int64_t> x_vec = vectorize<int64_t>(x.dims());
|
|
std::vector<int64_t> out_vec = vectorize<int64_t>(out->dims());
|
|
// if channel_last, transpose to channel_first
|
|
DenseTensor x_transpose;
|
|
if (data_layout == DataLayout::NHWC) {
|
|
if (strides.size() == 2U) {
|
|
std::vector<int> axis = {0, 3, 1, 2};
|
|
for (size_t i = 0; i < axis.size(); ++i) {
|
|
x_vec[i] = x.dims()[axis[i]];
|
|
out_vec[i] = out->dims()[axis[i]];
|
|
}
|
|
x_transpose = Transpose<T, Context>(dev_ctx, x, axis);
|
|
} else if (strides.size() == 3U) {
|
|
std::vector<int> axis = {0, 4, 1, 2, 3};
|
|
for (size_t i = 0; i < axis.size(); ++i) {
|
|
x_vec[i] = x.dims()[axis[i]];
|
|
out_vec[i] = out->dims()[axis[i]];
|
|
}
|
|
x_transpose = Transpose<T, Context>(dev_ctx, x, axis);
|
|
}
|
|
} else {
|
|
x_transpose = x;
|
|
}
|
|
|
|
// update padding and dilation
|
|
auto x_dims = x_transpose.dims();
|
|
auto filter_dims = filter.dims();
|
|
DDim x_data_dims;
|
|
x_data_dims = slice_ddim(x_dims, 2, x_dims.size());
|
|
DDim filter_data_dims = slice_ddim(filter_dims, 2, filter_dims.size());
|
|
std::vector<int> ksize = vectorize<int>(filter_data_dims);
|
|
UpdatePaddingAndDilation(
|
|
&paddings_, &dilations_, padding_algorithm, x_data_dims, strides, ksize);
|
|
|
|
int data_dim = strides.size(); // 2d or 3d
|
|
bool is_sys_pad = funcs::IsSymmetricPadding(paddings_, data_dim);
|
|
|
|
std::vector<int> x_pad(x_dims.size() * 2, 0);
|
|
DenseTensor transformed_x;
|
|
std::vector<int> padding_common(data_dim, 0);
|
|
if (!is_sys_pad) {
|
|
std::vector<int> padding_diff(data_dim);
|
|
std::vector<int64_t> new_x_shape_vec(data_dim + 2);
|
|
new_x_shape_vec[0] = x_dims[0];
|
|
new_x_shape_vec[1] = x_dims[1];
|
|
|
|
for (size_t i = 0; i < data_dim; ++i) {
|
|
padding_diff[i] = std::abs(paddings_[2 * i] - paddings_[2 * i + 1]);
|
|
padding_common[i] = std::min(paddings_[2 * i], paddings_[2 * i + 1]);
|
|
new_x_shape_vec[i + 2] = x_dims[i + 2] + padding_diff[i];
|
|
x_pad[2 * i + 4] = paddings_[2 * i] - padding_common[i];
|
|
x_pad[2 * i + 4 + 1] = paddings_[2 * i + 1] - padding_common[i];
|
|
}
|
|
DDim new_x_shape(make_ddim(new_x_shape_vec));
|
|
transformed_x.Resize(new_x_shape);
|
|
dev_ctx.template Alloc<T>(&transformed_x);
|
|
|
|
const int rank = x_dims.size();
|
|
T pad_value(0.0);
|
|
switch (rank) {
|
|
case 4: {
|
|
funcs::PadFunction<Context, T, 4>(
|
|
dev_ctx, x_pad, x_transpose, pad_value, &transformed_x);
|
|
} break;
|
|
case 5: {
|
|
funcs::PadFunction<Context, T, 5>(
|
|
dev_ctx, x_pad, x_transpose, pad_value, &transformed_x);
|
|
} break;
|
|
default:
|
|
PADDLE_THROW(errors::InvalidArgument(
|
|
"Op(ConvTranspose) only supports 4-D or 5-D x DenseTensor."));
|
|
}
|
|
} else {
|
|
transformed_x = x_transpose;
|
|
if (paddings_.size() == data_dim) {
|
|
for (size_t i = 0; i < data_dim; ++i) {
|
|
padding_common[i] = paddings_[i];
|
|
}
|
|
} else {
|
|
for (size_t i = 0; i < data_dim; ++i) {
|
|
padding_common[i] = paddings_[2 * i];
|
|
}
|
|
}
|
|
}
|
|
|
|
std::vector<int64_t> starts(data_dim, 0);
|
|
std::vector<int64_t> ends(data_dim, 0);
|
|
std::vector<int64_t> axes(data_dim, 0);
|
|
for (size_t i = 0; i < data_dim; ++i) {
|
|
starts[i] = x_pad[2 * i + 4] * (strides[i] + 1);
|
|
ends[i] = starts[i] + out_vec[i + 2];
|
|
axes[i] = i + 2;
|
|
}
|
|
|
|
x_vec = vectorize<int64_t>(transformed_x.dims());
|
|
|
|
std::vector<int64_t> transformed_out_vec = out_vec;
|
|
for (size_t i = 0; i < data_dim; ++i) {
|
|
transformed_out_vec[i + 2] =
|
|
out_vec[i + 2] + (x_pad[2 * i + 4] + x_pad[2 * i + 5]) * strides[i] -
|
|
2 * padding_common[i] + paddings_[2 * i] + paddings_[2 * i + 1];
|
|
}
|
|
|
|
DenseTensor transformed_out;
|
|
if (!is_sys_pad) {
|
|
transformed_out.Resize(transformed_out_vec);
|
|
dev_ctx.template Alloc<T>(&transformed_out);
|
|
} else {
|
|
dev_ctx.template Alloc<T>(out);
|
|
transformed_out.ShareDataWith(*out);
|
|
transformed_out.Resize(transformed_out_vec);
|
|
}
|
|
|
|
DataLayout layout;
|
|
if (strides.size() == 2U) {
|
|
layout = DataLayout::NCHW;
|
|
} else {
|
|
layout = DataLayout::NCDHW;
|
|
}
|
|
|
|
#ifdef PADDLE_WITH_CUDNN_FRONTEND
|
|
if (dynload::IsCudnnFrontendEnabled())
|
|
ConvTransposeCudnnKernelImplV8<T>(&transformed_x,
|
|
&filter,
|
|
dev_ctx,
|
|
strides,
|
|
padding_common,
|
|
dilations_,
|
|
data_layout,
|
|
layout,
|
|
exhaustive_search,
|
|
deterministic,
|
|
groups,
|
|
&transformed_out);
|
|
else
|
|
ConvTransposeCudnnKernelImplV7<T>(&transformed_x,
|
|
&filter,
|
|
dev_ctx,
|
|
strides,
|
|
padding_common,
|
|
dilations_,
|
|
data_layout,
|
|
layout,
|
|
exhaustive_search,
|
|
deterministic,
|
|
groups,
|
|
&transformed_out);
|
|
#else
|
|
ConvTransposeCudnnKernelImplV7<T>(&transformed_x,
|
|
&filter,
|
|
dev_ctx,
|
|
strides,
|
|
padding_common,
|
|
dilations_,
|
|
data_layout,
|
|
layout,
|
|
exhaustive_search,
|
|
deterministic,
|
|
groups,
|
|
&transformed_out);
|
|
#endif
|
|
|
|
if (!is_sys_pad && strides.size() == 2U) {
|
|
funcs::Slice<Context, T, 4>(
|
|
dev_ctx, &transformed_out, out, starts, ends, axes);
|
|
} else if (!is_sys_pad && strides.size() == 3U) {
|
|
funcs::Slice<Context, T, 5>(
|
|
dev_ctx, &transformed_out, out, starts, ends, axes);
|
|
}
|
|
|
|
if (data_layout == DataLayout::NHWC) {
|
|
DenseTensor out_transpose;
|
|
DenseTensor out_nchw;
|
|
out_nchw.ShareDataWith(*out);
|
|
out_nchw.Resize(out_vec);
|
|
|
|
if (strides.size() == 2U) {
|
|
out_transpose = Transpose<T, Context>(dev_ctx, out_nchw, {0, 2, 3, 1});
|
|
} else if (strides.size() == 3U) {
|
|
out_transpose = Transpose<T, Context>(dev_ctx, out_nchw, {0, 2, 3, 4, 1});
|
|
}
|
|
*out = out_transpose;
|
|
}
|
|
}
|
|
|
|
#ifdef PADDLE_WITH_CUDNN_FRONTEND
|
|
template <typename T, typename Context>
|
|
void ConvTransposeRawGPUDNNKernelV8(const Context& dev_ctx,
|
|
const DenseTensor& x,
|
|
const DenseTensor& filter,
|
|
const std::vector<int>& strides,
|
|
const std::vector<int>& paddings,
|
|
const std::string& padding_algorithm,
|
|
int groups,
|
|
const std::vector<int>& dilations,
|
|
const std::string& data_format,
|
|
DenseTensor* out) {
|
|
DenseTensor dummy_input;
|
|
dummy_input.Resize(out->dims());
|
|
|
|
dev_ctx.template Alloc<T>(out);
|
|
dummy_input.ShareDataWith(*out);
|
|
|
|
DenseTensor* filter_grad_ptr = nullptr;
|
|
|
|
ConvCudnnGradKernel<T, Context>(dev_ctx,
|
|
dummy_input,
|
|
filter,
|
|
x,
|
|
strides,
|
|
paddings,
|
|
padding_algorithm,
|
|
dilations,
|
|
groups,
|
|
data_format,
|
|
out,
|
|
filter_grad_ptr);
|
|
}
|
|
#endif
|
|
|
|
template <typename T, typename Context>
|
|
void Conv2dTransposeGPUDNNKernel(const Context& dev_ctx,
|
|
const DenseTensor& x,
|
|
const DenseTensor& filter,
|
|
const std::vector<int>& strides,
|
|
const std::vector<int>& paddings,
|
|
const std::vector<int>& output_padding,
|
|
const IntArray& output_size,
|
|
const std::string& padding_algorithm,
|
|
int groups,
|
|
const std::vector<int>& dilations,
|
|
const std::string& data_format,
|
|
DenseTensor* out) {
|
|
#ifdef PADDLE_WITH_CUDNN_FRONTEND
|
|
if (dynload::IsCudnnFrontendEnabled() && FLAGS_use_accuracy_compatible_kernel)
|
|
ConvTransposeRawGPUDNNKernelV8<T, Context>(dev_ctx,
|
|
x,
|
|
filter,
|
|
strides,
|
|
paddings,
|
|
padding_algorithm,
|
|
groups,
|
|
dilations,
|
|
data_format,
|
|
out);
|
|
else
|
|
ConvTransposeRawGPUDNNKernel<T, Context>(dev_ctx,
|
|
x,
|
|
filter,
|
|
strides,
|
|
paddings,
|
|
padding_algorithm,
|
|
groups,
|
|
dilations,
|
|
data_format,
|
|
out);
|
|
|
|
#else
|
|
ConvTransposeRawGPUDNNKernel<T, Context>(dev_ctx,
|
|
x,
|
|
filter,
|
|
strides,
|
|
paddings,
|
|
padding_algorithm,
|
|
groups,
|
|
dilations,
|
|
data_format,
|
|
out);
|
|
#endif
|
|
}
|
|
|
|
template <typename T, typename Context>
|
|
void Conv3dTransposeGPUDNNKernel(const Context& dev_ctx,
|
|
const DenseTensor& x,
|
|
const DenseTensor& filter,
|
|
const std::vector<int>& strides,
|
|
const std::vector<int>& paddings,
|
|
const std::vector<int>& output_padding,
|
|
const std::vector<int>& output_size,
|
|
const std::string& padding_algorithm,
|
|
int groups,
|
|
const std::vector<int>& dilations,
|
|
const std::string& data_format,
|
|
DenseTensor* out) {
|
|
ConvTransposeRawGPUDNNKernel<T, Context>(dev_ctx,
|
|
x,
|
|
filter,
|
|
strides,
|
|
paddings,
|
|
padding_algorithm,
|
|
groups,
|
|
dilations,
|
|
data_format,
|
|
out);
|
|
}
|
|
|
|
} // namespace phi
|
|
|
|
using float16 = phi::float16;
|
|
|
|
#ifdef PADDLE_WITH_HIP
|
|
// MIOPEN do not support double
|
|
PD_REGISTER_KERNEL(conv2d_transpose,
|
|
GPUDNN,
|
|
ALL_LAYOUT,
|
|
phi::Conv2dTransposeGPUDNNKernel,
|
|
float,
|
|
float16) {}
|
|
PD_REGISTER_KERNEL(conv3d_transpose,
|
|
GPUDNN,
|
|
ALL_LAYOUT,
|
|
phi::Conv3dTransposeGPUDNNKernel,
|
|
float,
|
|
float16) {}
|
|
#else
|
|
#if CUDNN_VERSION_MIN(8, 1, 0)
|
|
PD_REGISTER_KERNEL(conv2d_transpose,
|
|
GPUDNN,
|
|
ALL_LAYOUT,
|
|
phi::Conv2dTransposeGPUDNNKernel,
|
|
float,
|
|
double,
|
|
float16,
|
|
phi::bfloat16) {}
|
|
PD_REGISTER_KERNEL(conv3d_transpose,
|
|
GPUDNN,
|
|
ALL_LAYOUT,
|
|
phi::Conv3dTransposeGPUDNNKernel,
|
|
float,
|
|
double,
|
|
float16,
|
|
phi::bfloat16) {}
|
|
#else
|
|
PD_REGISTER_KERNEL(conv2d_transpose,
|
|
GPUDNN,
|
|
ALL_LAYOUT,
|
|
phi::Conv2dTransposeGPUDNNKernel,
|
|
float,
|
|
double,
|
|
float16) {}
|
|
PD_REGISTER_KERNEL(conv3d_transpose,
|
|
GPUDNN,
|
|
ALL_LAYOUT,
|
|
phi::Conv3dTransposeGPUDNNKernel,
|
|
float,
|
|
double,
|
|
float16) {}
|
|
#endif
|
|
|
|
#endif
|