|
|
|
@ -20,11 +20,12 @@ namespace operators {
|
|
|
|
|
namespace math {
|
|
|
|
|
template <typename T>
|
|
|
|
|
__global__ void KernelUnpool2dMax(const int nthreads, const T* input_data,
|
|
|
|
|
const int* indices_data,
|
|
|
|
|
const int input_height, const int input_width,
|
|
|
|
|
const int channels, T* output_data,
|
|
|
|
|
const int output_height,
|
|
|
|
|
const int output_width) {
|
|
|
|
|
const int* indices_data,
|
|
|
|
|
const int input_height,
|
|
|
|
|
const int input_width,
|
|
|
|
|
const int channels, T* output_data,
|
|
|
|
|
const int output_height,
|
|
|
|
|
const int output_width) {
|
|
|
|
|
int in_n_stride = input_height * input_width * channels;
|
|
|
|
|
int in_c_stride = input_height * input_width;
|
|
|
|
|
int out_n_stride = output_height * output_width * channels;
|
|
|
|
@ -42,12 +43,11 @@ __global__ void KernelUnpool2dMax(const int nthreads, const T* input_data,
|
|
|
|
|
}
|
|
|
|
|
}
|
|
|
|
|
template <typename T>
|
|
|
|
|
__global__ void KernelUnpool2dMaxGrad(const int nthreads, const T* input_data,
|
|
|
|
|
const int* indices_data,
|
|
|
|
|
const int input_height, const int input_width,
|
|
|
|
|
const int channels, const T* output_data,
|
|
|
|
|
const T* output_grad, const int output_height,
|
|
|
|
|
const int output_width, T* input_grad) {
|
|
|
|
|
__global__ void KernelUnpool2dMaxGrad(
|
|
|
|
|
const int nthreads, const T* input_data, const int* indices_data,
|
|
|
|
|
const int input_height, const int input_width, const int channels,
|
|
|
|
|
const T* output_data, const T* output_grad, const int output_height,
|
|
|
|
|
const int output_width, T* input_grad) {
|
|
|
|
|
int in_n_stride = input_height * input_width * channels;
|
|
|
|
|
int in_c_stride = input_height * input_width;
|
|
|
|
|
int out_n_stride = output_height * output_width * channels;
|
|
|
|
@ -71,8 +71,8 @@ template <typename T>
|
|
|
|
|
class Unpool2dMaxFunctor<platform::GPUPlace, T> {
|
|
|
|
|
public:
|
|
|
|
|
void operator()(const platform::DeviceContext& context,
|
|
|
|
|
const framework::Tensor& input, const framework::Tensor& indices,
|
|
|
|
|
framework::Tensor* output) {
|
|
|
|
|
const framework::Tensor& input,
|
|
|
|
|
const framework::Tensor& indices, framework::Tensor* output) {
|
|
|
|
|
const int batch_size = input.dims()[0];
|
|
|
|
|
const int input_height = input.dims()[2];
|
|
|
|
|
const int input_width = input.dims()[3];
|
|
|
|
@ -88,8 +88,8 @@ class Unpool2dMaxFunctor<platform::GPUPlace, T> {
|
|
|
|
|
T><<<grid, threads, 0,
|
|
|
|
|
reinterpret_cast<const platform::CUDADeviceContext&>(context)
|
|
|
|
|
.stream()>>>(input.numel(), input_data, indices_data,
|
|
|
|
|
input_height, input_width, output_channels,
|
|
|
|
|
output_data, output_height, output_width);
|
|
|
|
|
input_height, input_width, output_channels,
|
|
|
|
|
output_data, output_height, output_width);
|
|
|
|
|
}
|
|
|
|
|
};
|
|
|
|
|
/*
|
|
|
|
@ -121,9 +121,9 @@ class Unpool2dMaxGradFunctor<platform::GPUPlace, T> {
|
|
|
|
|
T><<<grid, threads, 0,
|
|
|
|
|
reinterpret_cast<const platform::CUDADeviceContext&>(context)
|
|
|
|
|
.stream()>>>(input.numel(), input_data, indices_data,
|
|
|
|
|
input_height, input_width, output_channels, output_data,
|
|
|
|
|
output_grad_data, output_height, output_width,
|
|
|
|
|
input_grad_data);
|
|
|
|
|
input_height, input_width, output_channels,
|
|
|
|
|
output_data, output_grad_data, output_height,
|
|
|
|
|
output_width, input_grad_data);
|
|
|
|
|
}
|
|
|
|
|
};
|
|
|
|
|
template class Unpool2dMaxGradFunctor<platform::GPUPlace, float>;
|
|
|
|
|