|
|
|
@ -19,10 +19,10 @@ namespace paddle {
|
|
|
|
|
namespace operators {
|
|
|
|
|
namespace math {
|
|
|
|
|
|
|
|
|
|
template <typename T>
|
|
|
|
|
template <typename T, typename T2>
|
|
|
|
|
__global__ void KernelUnpool2dMax(const int nthreads,
|
|
|
|
|
const T* input_data,
|
|
|
|
|
const T* indices_data,
|
|
|
|
|
const T2 * indices_data,
|
|
|
|
|
const int input_height,
|
|
|
|
|
const int input_width,
|
|
|
|
|
const int channels,
|
|
|
|
@ -45,10 +45,10 @@ __global__ void KernelUnpool2dMax(const int nthreads,
|
|
|
|
|
output_data[out_offset + out_index] = input_data[i];
|
|
|
|
|
}
|
|
|
|
|
}
|
|
|
|
|
template <typename T>
|
|
|
|
|
template <typename T, typename T2>
|
|
|
|
|
__global__ void KernelUnpool2dMaxGrad(const int nthreads,
|
|
|
|
|
const T* input_data,
|
|
|
|
|
const T* indices_data,
|
|
|
|
|
const T2* indices_data,
|
|
|
|
|
const int input_height,
|
|
|
|
|
const int input_width,
|
|
|
|
|
const int channels,
|
|
|
|
@ -76,8 +76,8 @@ __global__ void KernelUnpool2dMaxGrad(const int nthreads,
|
|
|
|
|
/*
|
|
|
|
|
* All tensors are in NCHW format.
|
|
|
|
|
*/
|
|
|
|
|
template <typename T>
|
|
|
|
|
class Unpool2dMaxFunctor<platform::GPUPlace, T> {
|
|
|
|
|
template <typename T, typename T2>
|
|
|
|
|
class Unpool2dMaxFunctor<platform::GPUPlace, T, T2> {
|
|
|
|
|
public:
|
|
|
|
|
void operator()(const platform::DeviceContext& context,
|
|
|
|
|
const framework::Tensor& input,
|
|
|
|
@ -90,7 +90,7 @@ class Unpool2dMaxFunctor<platform::GPUPlace, T> {
|
|
|
|
|
const int output_height = output->dims()[2];
|
|
|
|
|
const int output_width = output->dims()[3];
|
|
|
|
|
const T* input_data = input.data<T>();
|
|
|
|
|
const T* indices_data = indices.data<T>();
|
|
|
|
|
const T2 * indices_data = indices.data<T2>();
|
|
|
|
|
T* output_data = output->mutable_data<T>(context.GetPlace());
|
|
|
|
|
int nthreads = batch_size * output_channels * input_height * input_width;
|
|
|
|
|
int blocks = (nthreads + 1024 - 1) / 1024;
|
|
|
|
@ -98,7 +98,7 @@ class Unpool2dMaxFunctor<platform::GPUPlace, T> {
|
|
|
|
|
dim3 grid(blocks, 1);
|
|
|
|
|
|
|
|
|
|
KernelUnpool2dMax<
|
|
|
|
|
T><<<grid, threads, 0,
|
|
|
|
|
T, T2><<<grid, threads, 0,
|
|
|
|
|
reinterpret_cast<const platform::CUDADeviceContext&>(context)
|
|
|
|
|
.stream()>>>(nthreads, input_data, indices_data,
|
|
|
|
|
input_height, input_width, output_channels,
|
|
|
|
@ -108,8 +108,8 @@ class Unpool2dMaxFunctor<platform::GPUPlace, T> {
|
|
|
|
|
/*
|
|
|
|
|
* All tensors are in NCHW format.
|
|
|
|
|
*/
|
|
|
|
|
template <typename T>
|
|
|
|
|
class Unpool2dMaxGradFunctor<platform::GPUPlace, T> {
|
|
|
|
|
template <typename T, typename T2>
|
|
|
|
|
class Unpool2dMaxGradFunctor<platform::GPUPlace, T, T2> {
|
|
|
|
|
public:
|
|
|
|
|
void operator()(const platform::DeviceContext& context,
|
|
|
|
|
const framework::Tensor& input,
|
|
|
|
@ -124,7 +124,7 @@ class Unpool2dMaxGradFunctor<platform::GPUPlace, T> {
|
|
|
|
|
const int output_height = output.dims()[2];
|
|
|
|
|
const int output_width = output.dims()[3];
|
|
|
|
|
const T* input_data = input.data<T>();
|
|
|
|
|
const T* indices_data = indices.data<T>();
|
|
|
|
|
const T2 * indices_data = indices.data<T2>();
|
|
|
|
|
const T* output_data = output.data<T>();
|
|
|
|
|
const T* output_grad_data = output_grad.data<T>();
|
|
|
|
|
T* input_grad_data = input_grad->mutable_data<T>(context.GetPlace());
|
|
|
|
@ -134,7 +134,7 @@ class Unpool2dMaxGradFunctor<platform::GPUPlace, T> {
|
|
|
|
|
dim3 grid(blocks, 1);
|
|
|
|
|
|
|
|
|
|
KernelUnpool2dMaxGrad<
|
|
|
|
|
T><<<grid, threads, 0,
|
|
|
|
|
T, T2><<<grid, threads, 0,
|
|
|
|
|
reinterpret_cast<const platform::CUDADeviceContext&>(context)
|
|
|
|
|
.stream()>>>(
|
|
|
|
|
nthreads, input_data, indices_data,
|
|
|
|
@ -145,11 +145,11 @@ class Unpool2dMaxGradFunctor<platform::GPUPlace, T> {
|
|
|
|
|
}
|
|
|
|
|
};
|
|
|
|
|
|
|
|
|
|
template class Unpool2dMaxGradFunctor<platform::GPUPlace, float>;
|
|
|
|
|
template class Unpool2dMaxGradFunctor<platform::GPUPlace, double>;
|
|
|
|
|
template class Unpool2dMaxGradFunctor<platform::GPUPlace, float, int>;
|
|
|
|
|
template class Unpool2dMaxGradFunctor<platform::GPUPlace, double, int>;
|
|
|
|
|
|
|
|
|
|
template class Unpool2dMaxFunctor<platform::GPUPlace, float>;
|
|
|
|
|
template class Unpool2dMaxFunctor<platform::GPUPlace, double>;
|
|
|
|
|
template class Unpool2dMaxFunctor<platform::GPUPlace, float, int>;
|
|
|
|
|
template class Unpool2dMaxFunctor<platform::GPUPlace, double, int>;
|
|
|
|
|
|
|
|
|
|
} // namespace math
|
|
|
|
|
} // namespace operators
|
|
|
|
|