wehub-resource-sync
694 lines
32 KiB
Plaintext
694 lines
32 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/pad3d_grad_kernel.h"
|
|
|
|
#include "paddle/phi/backends/gpu/gpu_context.h"
|
|
#include "paddle/phi/backends/gpu/gpu_primitives.h"
|
|
#include "paddle/phi/core/kernel_registry.h"
|
|
#include "paddle/phi/kernels/funcs/math_function.h"
|
|
|
|
namespace phi {
|
|
|
|
template <typename T, typename IndexType>
|
|
__global__ void Pad3DGradConstNCDHW(const IndexType in_size,
|
|
T* d_in_data,
|
|
const IndexType num,
|
|
const IndexType channels,
|
|
const IndexType in_depth,
|
|
const IndexType in_height,
|
|
const IndexType in_width,
|
|
const IndexType out_depth,
|
|
const IndexType out_height,
|
|
const IndexType out_width,
|
|
const IndexType pad_front,
|
|
const IndexType pad_top,
|
|
const IndexType pad_left,
|
|
const T* d_out_data) {
|
|
CUDA_KERNEL_LOOP_TYPE(in_index, in_size, IndexType) {
|
|
const IndexType in_w = in_index % in_width;
|
|
|
|
IndexType nc = in_index / in_width;
|
|
const IndexType in_h = nc % in_height;
|
|
|
|
nc /= in_height;
|
|
const IndexType in_d = nc % in_depth;
|
|
|
|
nc /= in_depth;
|
|
|
|
const IndexType out_d = in_d + pad_front;
|
|
const IndexType out_h = in_h + pad_top;
|
|
const IndexType out_w = in_w + pad_left;
|
|
bool out_of_bound = out_d < 0 || out_h < 0 || out_w < 0;
|
|
|
|
d_in_data[in_index] =
|
|
out_of_bound ? static_cast<T>(0)
|
|
: d_out_data[nc * out_depth * out_height * out_width +
|
|
out_d * out_height * out_width +
|
|
out_h * out_width + out_w];
|
|
}
|
|
}
|
|
|
|
template <typename T, typename IndexType>
|
|
__global__ void Pad3DGradConstNDHWC(const IndexType in_size,
|
|
T* d_in_data,
|
|
const IndexType num,
|
|
const IndexType channels,
|
|
const IndexType in_depth,
|
|
const IndexType in_height,
|
|
const IndexType in_width,
|
|
const IndexType out_depth,
|
|
const IndexType out_height,
|
|
const IndexType out_width,
|
|
const IndexType pad_front,
|
|
const IndexType pad_top,
|
|
const IndexType pad_left,
|
|
const T* d_out_data) {
|
|
CUDA_KERNEL_LOOP_TYPE(in_index, in_size, IndexType) {
|
|
const IndexType c = in_index % channels;
|
|
IndexType n = in_index / channels;
|
|
|
|
const IndexType in_w = n % in_width;
|
|
n /= in_width;
|
|
|
|
const IndexType in_h = n % in_height;
|
|
n /= in_height;
|
|
|
|
const IndexType in_d = n % in_depth;
|
|
n /= in_depth;
|
|
|
|
const IndexType out_d = in_d + pad_front;
|
|
const IndexType out_h = in_h + pad_top;
|
|
const IndexType out_w = in_w + pad_left;
|
|
bool out_of_bound = out_d < 0 || out_h < 0 || out_w < 0;
|
|
d_in_data[in_index] =
|
|
out_of_bound
|
|
? static_cast<T>(0)
|
|
: d_out_data[n * out_depth * out_height * out_width * channels +
|
|
out_d * out_height * out_width * channels +
|
|
out_h * out_width * channels + out_w * channels + c];
|
|
}
|
|
}
|
|
|
|
template <typename T, typename IndexType>
|
|
__global__ void Pad3DGradReflectNCDHW(const IndexType out_size,
|
|
T* d_in_data,
|
|
const IndexType num,
|
|
const IndexType channels,
|
|
const IndexType in_depth,
|
|
const IndexType in_height,
|
|
const IndexType in_width,
|
|
const IndexType out_depth,
|
|
const IndexType out_height,
|
|
const IndexType out_width,
|
|
const IndexType pad_front,
|
|
const IndexType pad_top,
|
|
const IndexType pad_left,
|
|
const T* d_out_data) {
|
|
CUDA_KERNEL_LOOP_TYPE(out_index, out_size, IndexType) {
|
|
IndexType nc = out_index / out_width;
|
|
const IndexType out_w = out_index % out_width;
|
|
const IndexType out_h = nc % out_height;
|
|
nc /= out_height;
|
|
const IndexType out_d = nc % out_depth;
|
|
nc /= out_depth;
|
|
|
|
IndexType in_d = out_d - pad_front;
|
|
IndexType in_h = out_h - pad_top;
|
|
IndexType in_w = out_w - pad_left;
|
|
|
|
in_d = max(in_d, -in_d);
|
|
in_h = max(in_h, -in_h);
|
|
in_w = max(in_w, -in_w);
|
|
|
|
in_d = min(in_d, 2 * in_depth - in_d - 2);
|
|
in_h = min(in_h, 2 * in_height - in_h - 2);
|
|
in_w = min(in_w, 2 * in_width - in_w - 2);
|
|
|
|
CudaAtomicAdd(
|
|
&d_in_data[nc * in_depth * in_height * in_width +
|
|
in_d * in_height * in_width + in_h * in_width + in_w],
|
|
d_out_data[out_index]);
|
|
}
|
|
}
|
|
|
|
template <typename T, typename IndexType>
|
|
__global__ void Pad3DGradReflectNDHWC(const IndexType out_size,
|
|
T* d_in_data,
|
|
const IndexType num,
|
|
const IndexType channels,
|
|
const IndexType in_depth,
|
|
const IndexType in_height,
|
|
const IndexType in_width,
|
|
const IndexType out_depth,
|
|
const IndexType out_height,
|
|
const IndexType out_width,
|
|
const IndexType pad_front,
|
|
const IndexType pad_top,
|
|
const IndexType pad_left,
|
|
const T* d_out_data) {
|
|
CUDA_KERNEL_LOOP_TYPE(out_index, out_size, IndexType) {
|
|
const IndexType c = out_index % channels;
|
|
IndexType n = out_index / channels;
|
|
const IndexType out_w = n % out_width;
|
|
n /= out_width;
|
|
const IndexType out_h = n % out_height;
|
|
n /= out_height;
|
|
const IndexType out_d = n % out_depth;
|
|
n /= out_depth;
|
|
|
|
IndexType in_d = out_d - pad_front;
|
|
IndexType in_h = out_h - pad_top;
|
|
IndexType in_w = out_w - pad_left;
|
|
|
|
in_d = max(in_d, -in_d);
|
|
in_h = max(in_h, -in_h);
|
|
in_w = max(in_w, -in_w);
|
|
|
|
in_d = min(in_d, in_depth * 2 - in_d - 2);
|
|
in_h = min(in_h, in_height * 2 - in_h - 2);
|
|
in_w = min(in_w, in_width * 2 - in_w - 2);
|
|
CudaAtomicAdd(&d_in_data[n * in_depth * in_height * in_width * channels +
|
|
in_d * in_height * in_width * channels +
|
|
in_h * in_width * channels + in_w * channels + c],
|
|
d_out_data[out_index]);
|
|
}
|
|
}
|
|
|
|
template <typename T, typename IndexType>
|
|
__global__ void Pad3DGradReplicateNCDHW(const IndexType out_size,
|
|
T* d_in_data,
|
|
const IndexType num,
|
|
const IndexType channels,
|
|
const IndexType in_depth,
|
|
const IndexType in_height,
|
|
const IndexType in_width,
|
|
const IndexType out_depth,
|
|
const IndexType out_height,
|
|
const IndexType out_width,
|
|
const IndexType pad_front,
|
|
const IndexType pad_top,
|
|
const IndexType pad_left,
|
|
const T* d_out_data) {
|
|
CUDA_KERNEL_LOOP_TYPE(out_index, out_size, IndexType) {
|
|
IndexType nc = out_index / out_width;
|
|
const IndexType out_w = out_index % out_width;
|
|
const IndexType out_h = nc % out_height;
|
|
nc /= out_height;
|
|
const IndexType out_d = nc % out_depth;
|
|
nc /= out_depth;
|
|
|
|
const IndexType in_d =
|
|
min(in_depth - 1, max(out_d - pad_front, static_cast<IndexType>(0)));
|
|
const IndexType in_h =
|
|
min(in_height - 1, max(out_h - pad_top, static_cast<IndexType>(0)));
|
|
const IndexType in_w =
|
|
min(in_width - 1, max(out_w - pad_left, static_cast<IndexType>(0)));
|
|
|
|
CudaAtomicAdd(
|
|
&d_in_data[nc * in_depth * in_height * in_width +
|
|
in_d * in_height * in_width + in_h * in_width + in_w],
|
|
d_out_data[out_index]);
|
|
}
|
|
}
|
|
|
|
template <typename T, typename IndexType>
|
|
__global__ void Pad3DGradReplicateNDHWC(const IndexType out_size,
|
|
T* d_in_data,
|
|
const IndexType num,
|
|
const IndexType channels,
|
|
const IndexType in_depth,
|
|
const IndexType in_height,
|
|
const IndexType in_width,
|
|
const IndexType out_depth,
|
|
const IndexType out_height,
|
|
const IndexType out_width,
|
|
const IndexType pad_front,
|
|
const IndexType pad_top,
|
|
const IndexType pad_left,
|
|
const T* d_out_data) {
|
|
CUDA_KERNEL_LOOP_TYPE(out_index, out_size, IndexType) {
|
|
const IndexType c = out_index % channels;
|
|
IndexType n = out_index / channels;
|
|
const IndexType out_w = n % out_width;
|
|
n /= out_width;
|
|
const IndexType out_h = n % out_height;
|
|
n /= out_height;
|
|
const IndexType out_d = n % out_depth;
|
|
n /= out_depth;
|
|
|
|
const IndexType in_d =
|
|
min(in_depth - 1, max(out_d - pad_front, static_cast<IndexType>(0)));
|
|
const IndexType in_h =
|
|
min(in_height - 1, max(out_h - pad_top, static_cast<IndexType>(0)));
|
|
const IndexType in_w =
|
|
min(in_width - 1, max(out_w - pad_left, static_cast<IndexType>(0)));
|
|
|
|
CudaAtomicAdd(&d_in_data[n * in_depth * in_height * in_width * channels +
|
|
in_d * in_height * in_width * channels +
|
|
in_h * in_width * channels + in_w * channels + c],
|
|
d_out_data[out_index]);
|
|
}
|
|
}
|
|
|
|
template <typename T, typename IndexType>
|
|
__global__ void Pad3DGradCircularNCDHW(const IndexType out_size,
|
|
T* d_in_data,
|
|
const IndexType num,
|
|
const IndexType channels,
|
|
const IndexType in_depth,
|
|
const IndexType in_height,
|
|
const IndexType in_width,
|
|
const IndexType out_depth,
|
|
const IndexType out_height,
|
|
const IndexType out_width,
|
|
const IndexType pad_front,
|
|
const IndexType pad_top,
|
|
const IndexType pad_left,
|
|
const T* d_out_data) {
|
|
CUDA_KERNEL_LOOP_TYPE(out_index, out_size, IndexType) {
|
|
IndexType nc = out_index / out_width;
|
|
const IndexType out_w = out_index % out_width;
|
|
const IndexType out_h = nc % out_height;
|
|
nc /= out_height;
|
|
const IndexType out_d = nc % out_depth;
|
|
nc /= out_depth;
|
|
|
|
IndexType in_d = ((out_d - pad_front) % in_depth + in_depth) % in_depth;
|
|
IndexType in_h = ((out_h - pad_top) % in_height + in_height) % in_height;
|
|
IndexType in_w = ((out_w - pad_left) % in_width + in_width) % in_width;
|
|
|
|
CudaAtomicAdd(
|
|
&d_in_data[nc * in_depth * in_height * in_width +
|
|
in_d * in_height * in_width + in_h * in_width + in_w],
|
|
d_out_data[out_index]);
|
|
}
|
|
}
|
|
|
|
template <typename T, typename IndexType>
|
|
__global__ void Pad3DGradCircularNDHWC(const IndexType out_size,
|
|
T* d_in_data,
|
|
const IndexType num,
|
|
const IndexType channels,
|
|
const IndexType in_depth,
|
|
const IndexType in_height,
|
|
const IndexType in_width,
|
|
const IndexType out_depth,
|
|
const IndexType out_height,
|
|
const IndexType out_width,
|
|
const IndexType pad_front,
|
|
const IndexType pad_top,
|
|
const IndexType pad_left,
|
|
const T* d_out_data) {
|
|
CUDA_KERNEL_LOOP_TYPE(out_index, out_size, IndexType) {
|
|
const IndexType c = out_index % channels;
|
|
IndexType n = out_index / channels;
|
|
const IndexType out_w = n % out_width;
|
|
n /= out_width;
|
|
const IndexType out_h = n % out_height;
|
|
n /= out_height;
|
|
const IndexType out_d = n % out_depth;
|
|
n /= out_depth;
|
|
|
|
IndexType in_d = ((out_d - pad_front) % in_depth + in_depth) % in_depth;
|
|
IndexType in_h = ((out_h - pad_top) % in_height + in_height) % in_height;
|
|
IndexType in_w = ((out_w - pad_left) % in_width + in_width) % in_width;
|
|
|
|
CudaAtomicAdd(&d_in_data[n * in_depth * in_height * in_width * channels +
|
|
in_d * in_height * in_width * channels +
|
|
in_h * in_width * channels + in_w * channels + c],
|
|
d_out_data[out_index]);
|
|
}
|
|
}
|
|
|
|
template <typename T, typename Context>
|
|
void Pad3dGradKernel(const Context& dev_ctx,
|
|
const DenseTensor& x,
|
|
const DenseTensor& out_grad,
|
|
const IntArray& paddings,
|
|
const std::string& mode,
|
|
double pad_value,
|
|
const std::string& data_format,
|
|
DenseTensor* x_grad) {
|
|
std::vector<int64_t> pads = paddings.GetData();
|
|
auto* d_out = &out_grad;
|
|
auto* d_in = x_grad;
|
|
auto d_in_dims = d_in->dims();
|
|
auto d_out_dims = d_out->dims();
|
|
const T* d_out_data = d_out->data<T>();
|
|
T* d_in_data = dev_ctx.template Alloc<T>(d_in);
|
|
if (x.numel() == 0) return;
|
|
|
|
funcs::SetConstant<Context, T>()(dev_ctx, d_in, static_cast<T>(0));
|
|
|
|
const int64_t pad_left = pads[0];
|
|
const int64_t pad_top = pads[2];
|
|
const int64_t pad_front = pads[4];
|
|
|
|
const int64_t num = d_in_dims[0];
|
|
|
|
auto stream = dev_ctx.stream();
|
|
int block = PADDLE_CUDA_NUM_THREADS;
|
|
const size_t out_size = d_out->numel();
|
|
const size_t in_size = d_in->numel();
|
|
uint32_t grid = (out_size + block - 1) / block;
|
|
|
|
bool use_int32_index = true;
|
|
if (out_size > std::numeric_limits<int32_t>::max()) {
|
|
use_int32_index = false;
|
|
} else {
|
|
for (int i = 0; i < d_out_dims.size(); ++i) {
|
|
if (d_out_dims[i] > std::numeric_limits<int32_t>::max()) {
|
|
use_int32_index = false;
|
|
break;
|
|
}
|
|
}
|
|
}
|
|
if (use_int32_index) {
|
|
if (data_format == "NCDHW") {
|
|
const int channels = d_in_dims[1];
|
|
const int in_depth = d_in_dims[2];
|
|
const int in_height = d_in_dims[3];
|
|
const int in_width = d_in_dims[4];
|
|
const int out_depth = d_out_dims[2];
|
|
const int out_height = d_out_dims[3];
|
|
const int out_width = d_out_dims[4];
|
|
|
|
if (mode == "reflect") {
|
|
Pad3DGradReflectNCDHW<T, int32_t>
|
|
<<<grid, block, 0, stream>>>(out_size,
|
|
d_in_data,
|
|
num,
|
|
channels,
|
|
in_depth,
|
|
in_height,
|
|
in_width,
|
|
out_depth,
|
|
out_height,
|
|
out_width,
|
|
pad_front,
|
|
pad_top,
|
|
pad_left,
|
|
d_out_data);
|
|
} else if (mode == "replicate") {
|
|
Pad3DGradReplicateNCDHW<T, int32_t>
|
|
<<<grid, block, 0, stream>>>(out_size,
|
|
d_in_data,
|
|
num,
|
|
channels,
|
|
in_depth,
|
|
in_height,
|
|
in_width,
|
|
out_depth,
|
|
out_height,
|
|
out_width,
|
|
pad_front,
|
|
pad_top,
|
|
pad_left,
|
|
d_out_data);
|
|
} else if (mode == "circular") {
|
|
Pad3DGradCircularNCDHW<T, int32_t>
|
|
<<<grid, block, 0, stream>>>(out_size,
|
|
d_in_data,
|
|
num,
|
|
channels,
|
|
in_depth,
|
|
in_height,
|
|
in_width,
|
|
out_depth,
|
|
out_height,
|
|
out_width,
|
|
pad_front,
|
|
pad_top,
|
|
pad_left,
|
|
d_out_data);
|
|
} else {
|
|
grid = (in_size + block - 1) / block;
|
|
Pad3DGradConstNCDHW<T, int32_t><<<grid, block, 0, stream>>>(in_size,
|
|
d_in_data,
|
|
num,
|
|
channels,
|
|
in_depth,
|
|
in_height,
|
|
in_width,
|
|
out_depth,
|
|
out_height,
|
|
out_width,
|
|
pad_front,
|
|
pad_top,
|
|
pad_left,
|
|
d_out_data);
|
|
}
|
|
} else {
|
|
const int channels = d_in_dims[4];
|
|
const int in_depth = d_in_dims[1];
|
|
const int in_height = d_in_dims[2];
|
|
const int in_width = d_in_dims[3];
|
|
const int out_depth = d_out_dims[1];
|
|
const int out_height = d_out_dims[2];
|
|
const int out_width = d_out_dims[3];
|
|
if (mode == "reflect") {
|
|
Pad3DGradReflectNDHWC<T, int32_t>
|
|
<<<grid, block, 0, stream>>>(out_size,
|
|
d_in_data,
|
|
num,
|
|
channels,
|
|
in_depth,
|
|
in_height,
|
|
in_width,
|
|
out_depth,
|
|
out_height,
|
|
out_width,
|
|
pad_front,
|
|
pad_top,
|
|
pad_left,
|
|
d_out_data);
|
|
} else if (mode == "replicate") {
|
|
Pad3DGradReplicateNDHWC<T, int32_t>
|
|
<<<grid, block, 0, stream>>>(out_size,
|
|
d_in_data,
|
|
num,
|
|
channels,
|
|
in_depth,
|
|
in_height,
|
|
in_width,
|
|
out_depth,
|
|
out_height,
|
|
out_width,
|
|
pad_front,
|
|
pad_top,
|
|
pad_left,
|
|
d_out_data);
|
|
} else if (mode == "circular") {
|
|
Pad3DGradCircularNDHWC<T, int32_t>
|
|
<<<grid, block, 0, stream>>>(out_size,
|
|
d_in_data,
|
|
num,
|
|
channels,
|
|
in_depth,
|
|
in_height,
|
|
in_width,
|
|
out_depth,
|
|
out_height,
|
|
out_width,
|
|
pad_front,
|
|
pad_top,
|
|
pad_left,
|
|
d_out_data);
|
|
} else {
|
|
grid = (in_size + block - 1) / block;
|
|
Pad3DGradConstNDHWC<T, int32_t><<<grid, block, 0, stream>>>(in_size,
|
|
d_in_data,
|
|
num,
|
|
channels,
|
|
in_depth,
|
|
in_height,
|
|
in_width,
|
|
out_depth,
|
|
out_height,
|
|
out_width,
|
|
pad_front,
|
|
pad_top,
|
|
pad_left,
|
|
d_out_data);
|
|
}
|
|
}
|
|
|
|
} else {
|
|
if (data_format == "NCDHW") {
|
|
const int64_t channels = d_in_dims[1];
|
|
const int64_t in_depth = d_in_dims[2];
|
|
const int64_t in_height = d_in_dims[3];
|
|
const int64_t in_width = d_in_dims[4];
|
|
const int64_t out_depth = d_out_dims[2];
|
|
const int64_t out_height = d_out_dims[3];
|
|
const int64_t out_width = d_out_dims[4];
|
|
|
|
if (mode == "reflect") {
|
|
Pad3DGradReflectNCDHW<T, int64_t>
|
|
<<<grid, block, 0, stream>>>(out_size,
|
|
d_in_data,
|
|
num,
|
|
channels,
|
|
in_depth,
|
|
in_height,
|
|
in_width,
|
|
out_depth,
|
|
out_height,
|
|
out_width,
|
|
pad_front,
|
|
pad_top,
|
|
pad_left,
|
|
d_out_data);
|
|
} else if (mode == "replicate") {
|
|
Pad3DGradReplicateNCDHW<T, int64_t>
|
|
<<<grid, block, 0, stream>>>(out_size,
|
|
d_in_data,
|
|
num,
|
|
channels,
|
|
in_depth,
|
|
in_height,
|
|
in_width,
|
|
out_depth,
|
|
out_height,
|
|
out_width,
|
|
pad_front,
|
|
pad_top,
|
|
pad_left,
|
|
d_out_data);
|
|
} else if (mode == "circular") {
|
|
Pad3DGradCircularNCDHW<T, int64_t>
|
|
<<<grid, block, 0, stream>>>(out_size,
|
|
d_in_data,
|
|
num,
|
|
channels,
|
|
in_depth,
|
|
in_height,
|
|
in_width,
|
|
out_depth,
|
|
out_height,
|
|
out_width,
|
|
pad_front,
|
|
pad_top,
|
|
pad_left,
|
|
d_out_data);
|
|
} else {
|
|
grid = (in_size + block - 1) / block;
|
|
Pad3DGradConstNCDHW<T, int64_t><<<grid, block, 0, stream>>>(in_size,
|
|
d_in_data,
|
|
num,
|
|
channels,
|
|
in_depth,
|
|
in_height,
|
|
in_width,
|
|
out_depth,
|
|
out_height,
|
|
out_width,
|
|
pad_front,
|
|
pad_top,
|
|
pad_left,
|
|
d_out_data);
|
|
}
|
|
} else {
|
|
const int64_t channels = d_in_dims[4];
|
|
const int64_t in_depth = d_in_dims[1];
|
|
const int64_t in_height = d_in_dims[2];
|
|
const int64_t in_width = d_in_dims[3];
|
|
const int64_t out_depth = d_out_dims[1];
|
|
const int64_t out_height = d_out_dims[2];
|
|
const int64_t out_width = d_out_dims[3];
|
|
if (mode == "reflect") {
|
|
Pad3DGradReflectNDHWC<T, int64_t>
|
|
<<<grid, block, 0, stream>>>(out_size,
|
|
d_in_data,
|
|
num,
|
|
channels,
|
|
in_depth,
|
|
in_height,
|
|
in_width,
|
|
out_depth,
|
|
out_height,
|
|
out_width,
|
|
pad_front,
|
|
pad_top,
|
|
pad_left,
|
|
d_out_data);
|
|
} else if (mode == "replicate") {
|
|
Pad3DGradReplicateNDHWC<T, int64_t>
|
|
<<<grid, block, 0, stream>>>(out_size,
|
|
d_in_data,
|
|
num,
|
|
channels,
|
|
in_depth,
|
|
in_height,
|
|
in_width,
|
|
out_depth,
|
|
out_height,
|
|
out_width,
|
|
pad_front,
|
|
pad_top,
|
|
pad_left,
|
|
d_out_data);
|
|
} else if (mode == "circular") {
|
|
Pad3DGradCircularNDHWC<T, int64_t>
|
|
<<<grid, block, 0, stream>>>(out_size,
|
|
d_in_data,
|
|
num,
|
|
channels,
|
|
in_depth,
|
|
in_height,
|
|
in_width,
|
|
out_depth,
|
|
out_height,
|
|
out_width,
|
|
pad_front,
|
|
pad_top,
|
|
pad_left,
|
|
d_out_data);
|
|
} else {
|
|
grid = (in_size + block - 1) / block;
|
|
Pad3DGradConstNDHWC<T, int64_t><<<grid, block, 0, stream>>>(in_size,
|
|
d_in_data,
|
|
num,
|
|
channels,
|
|
in_depth,
|
|
in_height,
|
|
in_width,
|
|
out_depth,
|
|
out_height,
|
|
out_width,
|
|
pad_front,
|
|
pad_top,
|
|
pad_left,
|
|
d_out_data);
|
|
}
|
|
}
|
|
}
|
|
}
|
|
|
|
} // namespace phi
|
|
|
|
PD_REGISTER_KERNEL(pad3d_grad,
|
|
GPU,
|
|
ALL_LAYOUT,
|
|
phi::Pad3dGradKernel,
|
|
float,
|
|
double,
|
|
int,
|
|
int64_t,
|
|
phi::float16,
|
|
phi::bfloat16,
|
|
phi::complex64,
|
|
phi::complex128) {}
|