提交 c93e044a 编写于 作者: D dengkaipeng 提交者: dengkaipeng

add inclusive/exclusive mode in PoolOp avg pool type

上级 0a80f06e
...@@ -29,8 +29,8 @@ class Pool2dFunctor<platform::CPUDeviceContext, PoolProcess, T> { ...@@ -29,8 +29,8 @@ class Pool2dFunctor<platform::CPUDeviceContext, PoolProcess, T> {
public: public:
void operator()(const platform::CPUDeviceContext& context, void operator()(const platform::CPUDeviceContext& context,
const framework::Tensor& input, const std::vector<int>& ksize, const framework::Tensor& input, const std::vector<int>& ksize,
const std::vector<int>& strides, const std::vector<int>& strides, const std::vector<int>& paddings,
const std::vector<int>& paddings, PoolProcess pool_process, PoolProcess pool_process, bool exclusive,
framework::Tensor* output) { framework::Tensor* output) {
const int batch_size = input.dims()[0]; const int batch_size = input.dims()[0];
const int input_height = input.dims()[2]; const int input_height = input.dims()[2];
...@@ -68,7 +68,8 @@ class Pool2dFunctor<platform::CPUDeviceContext, PoolProcess, T> { ...@@ -68,7 +68,8 @@ class Pool2dFunctor<platform::CPUDeviceContext, PoolProcess, T> {
pool_process.compute(input_data[h * input_width + w], &ele); pool_process.compute(input_data[h * input_width + w], &ele);
} }
} }
int pool_size = (hend - hstart) * (wend - wstart); int pool_size = exclusive ? (hend - hstart) * (wend - wstart)
: ksize_height * ksize_width;
pool_process.finalize(static_cast<T>(pool_size), &ele); pool_process.finalize(static_cast<T>(pool_size), &ele);
output_data[ph * output_width + pw] = ele; output_data[ph * output_width + pw] = ele;
} }
...@@ -93,7 +94,7 @@ class Pool2dGradFunctor<platform::CPUDeviceContext, PoolProcess, T> { ...@@ -93,7 +94,7 @@ class Pool2dGradFunctor<platform::CPUDeviceContext, PoolProcess, T> {
const framework::Tensor& output, const framework::Tensor& output_grad, const framework::Tensor& output, const framework::Tensor& output_grad,
const std::vector<int>& ksize, const std::vector<int>& strides, const std::vector<int>& ksize, const std::vector<int>& strides,
const std::vector<int>& paddings, PoolProcess pool_grad_process, const std::vector<int>& paddings, PoolProcess pool_grad_process,
framework::Tensor* input_grad) { bool exclusive, framework::Tensor* input_grad) {
const int batch_size = input.dims()[0]; const int batch_size = input.dims()[0];
const int input_height = input.dims()[2]; const int input_height = input.dims()[2];
const int input_width = input.dims()[3]; const int input_width = input.dims()[3];
...@@ -124,7 +125,8 @@ class Pool2dGradFunctor<platform::CPUDeviceContext, PoolProcess, T> { ...@@ -124,7 +125,8 @@ class Pool2dGradFunctor<platform::CPUDeviceContext, PoolProcess, T> {
int wstart = pw * stride_width - padding_width; int wstart = pw * stride_width - padding_width;
int wend = std::min(wstart + ksize_width, input_width); int wend = std::min(wstart + ksize_width, input_width);
wstart = std::max(wstart, 0); wstart = std::max(wstart, 0);
int pool_size = (hend - hstart) * (wend - wstart); int pool_size = exclusive ? (hend - hstart) * (wend - wstart)
: ksize_height * ksize_width;
float scale = 1.0 / pool_size; float scale = 1.0 / pool_size;
for (int h = hstart; h < hend; ++h) { for (int h = hstart; h < hend; ++h) {
for (int w = wstart; w < wend; ++w) { for (int w = wstart; w < wend; ++w) {
...@@ -247,9 +249,9 @@ class Pool3dFunctor<platform::CPUDeviceContext, PoolProcess, T> { ...@@ -247,9 +249,9 @@ class Pool3dFunctor<platform::CPUDeviceContext, PoolProcess, T> {
public: public:
void operator()(const platform::CPUDeviceContext& context, void operator()(const platform::CPUDeviceContext& context,
const framework::Tensor& input, const std::vector<int>& ksize, const framework::Tensor& input, const std::vector<int>& ksize,
const std::vector<int>& strides, const std::vector<int>& strides, const std::vector<int>& paddings,
const std::vector<int>& paddings, PoolProcess pool_process, PoolProcess pool_process,
framework::Tensor* output) { bool exclusive, framework::Tensor* output) {
const int batch_size = input.dims()[0]; const int batch_size = input.dims()[0];
const int input_depth = input.dims()[2]; const int input_depth = input.dims()[2];
const int input_height = input.dims()[3]; const int input_height = input.dims()[3];
...@@ -299,8 +301,9 @@ class Pool3dFunctor<platform::CPUDeviceContext, PoolProcess, T> { ...@@ -299,8 +301,9 @@ class Pool3dFunctor<platform::CPUDeviceContext, PoolProcess, T> {
} }
} }
} }
int pool_size = int pool_size = exclusive ?
(dend - dstart) * (hend - hstart) * (wend - wstart); (dend - dstart) * (hend - hstart) * (wend - wstart)
: ksize_depth * ksize_height * ksize_width;
pool_process.finalize(static_cast<T>(pool_size), &ele); pool_process.finalize(static_cast<T>(pool_size), &ele);
output_data[output_idx] = ele; output_data[output_idx] = ele;
} }
...@@ -326,7 +329,7 @@ class Pool3dGradFunctor<platform::CPUDeviceContext, PoolProcess, T> { ...@@ -326,7 +329,7 @@ class Pool3dGradFunctor<platform::CPUDeviceContext, PoolProcess, T> {
const framework::Tensor& output, const framework::Tensor& output_grad, const framework::Tensor& output, const framework::Tensor& output_grad,
const std::vector<int>& ksize, const std::vector<int>& strides, const std::vector<int>& ksize, const std::vector<int>& strides,
const std::vector<int>& paddings, PoolProcess pool_grad_process, const std::vector<int>& paddings, PoolProcess pool_grad_process,
framework::Tensor* input_grad) { bool exclusive, framework::Tensor* input_grad) {
const int batch_size = input.dims()[0]; const int batch_size = input.dims()[0];
const int input_depth = input.dims()[2]; const int input_depth = input.dims()[2];
const int input_height = input.dims()[3]; const int input_height = input.dims()[3];
...@@ -368,8 +371,9 @@ class Pool3dGradFunctor<platform::CPUDeviceContext, PoolProcess, T> { ...@@ -368,8 +371,9 @@ class Pool3dGradFunctor<platform::CPUDeviceContext, PoolProcess, T> {
int wend = std::min(wstart + ksize_width, input_width); int wend = std::min(wstart + ksize_width, input_width);
wstart = std::max(wstart, 0); wstart = std::max(wstart, 0);
int pool_size = int pool_size = exclusive ?
(dend - dstart) * (hend - hstart) * (wend - wstart); (dend - dstart) * (hend - hstart) * (wend - wstart)
: ksize_depth * ksize_height * ksize_width;
float scale = 1.0 / pool_size; float scale = 1.0 / pool_size;
for (int d = dstart; d < dend; ++d) { for (int d = dstart; d < dend; ++d) {
for (int h = hstart; h < hend; ++h) { for (int h = hstart; h < hend; ++h) {
......
...@@ -29,7 +29,7 @@ __global__ void KernelPool2D(const int nthreads, const T* input_data, ...@@ -29,7 +29,7 @@ __global__ void KernelPool2D(const int nthreads, const T* input_data,
const int ksize_width, const int stride_height, const int ksize_width, const int stride_height,
const int stride_width, const int padding_height, const int stride_width, const int padding_height,
const int padding_width, PoolProcess pool_process, const int padding_width, PoolProcess pool_process,
T* output_data) { bool exclusive, T* output_data) {
for (int index = blockIdx.x * blockDim.x + threadIdx.x; index < nthreads; for (int index = blockIdx.x * blockDim.x + threadIdx.x; index < nthreads;
index += blockDim.x * gridDim.x) { index += blockDim.x * gridDim.x) {
int pw = index % output_width; int pw = index % output_width;
...@@ -52,7 +52,8 @@ __global__ void KernelPool2D(const int nthreads, const T* input_data, ...@@ -52,7 +52,8 @@ __global__ void KernelPool2D(const int nthreads, const T* input_data,
pool_process.compute(input_data[h * input_width + w], &ele); pool_process.compute(input_data[h * input_width + w], &ele);
} }
} }
int pool_size = (hend - hstart) * (wend - wstart); int pool_size = exclusive ? (hend - hstart) * (wend - wstart)
: ksize_height * ksize_width;
pool_process.finalize(static_cast<T>(pool_size), &ele); pool_process.finalize(static_cast<T>(pool_size), &ele);
output_data[index] = ele; output_data[index] = ele;
} }
...@@ -65,7 +66,7 @@ __global__ void KernelPool2DGrad( ...@@ -65,7 +66,7 @@ __global__ void KernelPool2DGrad(
const int input_width, const int output_height, const int output_width, const int input_width, const int output_height, const int output_width,
const int ksize_height, const int ksize_width, const int stride_height, const int ksize_height, const int ksize_width, const int stride_height,
const int stride_width, const int padding_height, const int padding_width, const int stride_width, const int padding_height, const int padding_width,
PoolProcess pool_process, T* input_grad) { PoolProcess pool_process, bool exclusive, T* input_grad) {
for (int index = blockIdx.x * blockDim.x + threadIdx.x; index < nthreads; for (int index = blockIdx.x * blockDim.x + threadIdx.x; index < nthreads;
index += blockDim.x * gridDim.x) { index += blockDim.x * gridDim.x) {
int offsetW = index % input_width + padding_width; int offsetW = index % input_width + padding_width;
...@@ -95,7 +96,8 @@ __global__ void KernelPool2DGrad( ...@@ -95,7 +96,8 @@ __global__ void KernelPool2DGrad(
int wend = min(wstart + ksize_width, input_width); int wend = min(wstart + ksize_width, input_width);
hstart = max(hstart, 0); hstart = max(hstart, 0);
wstart = max(wstart, 0); wstart = max(wstart, 0);
int pool_size = (hend - hstart) * (wend - wstart); int pool_size = exclusive ? (hend - hstart) * (wend - wstart)
: ksize_height * ksize_width;
int output_sub_idx = ph * output_width + pw; int output_sub_idx = ph * output_width + pw;
pool_process.compute(input, output_data[output_sub_idx], pool_process.compute(input, output_data[output_sub_idx],
output_grad[output_sub_idx], output_grad[output_sub_idx],
...@@ -163,7 +165,7 @@ class Pool2dFunctor<platform::CUDADeviceContext, PoolProcess, T> { ...@@ -163,7 +165,7 @@ class Pool2dFunctor<platform::CUDADeviceContext, PoolProcess, T> {
const framework::Tensor& input, const std::vector<int>& ksize, const framework::Tensor& input, const std::vector<int>& ksize,
const std::vector<int>& strides, const std::vector<int>& strides,
const std::vector<int>& paddings, PoolProcess pool_process, const std::vector<int>& paddings, PoolProcess pool_process,
framework::Tensor* output) { bool exclusive, framework::Tensor* output) {
const int batch_size = input.dims()[0]; const int batch_size = input.dims()[0];
const int input_channels = input.dims()[1]; const int input_channels = input.dims()[1];
const int input_height = input.dims()[2]; const int input_height = input.dims()[2];
...@@ -189,7 +191,8 @@ class Pool2dFunctor<platform::CUDADeviceContext, PoolProcess, T> { ...@@ -189,7 +191,8 @@ class Pool2dFunctor<platform::CUDADeviceContext, PoolProcess, T> {
KernelPool2D<PoolProcess, T><<<grid, threads, 0, context.stream()>>>( KernelPool2D<PoolProcess, T><<<grid, threads, 0, context.stream()>>>(
nthreads, input_data, input_channels, input_height, input_width, nthreads, input_data, input_channels, input_height, input_width,
output_height, output_width, ksize_height, ksize_width, stride_height, output_height, output_width, ksize_height, ksize_width, stride_height,
stride_width, padding_height, padding_width, pool_process, output_data); stride_width, padding_height, padding_width, pool_process, exclusive,
output_data);
} }
}; };
...@@ -208,7 +211,7 @@ class Pool2dGradFunctor<platform::CUDADeviceContext, PoolProcess, T> { ...@@ -208,7 +211,7 @@ class Pool2dGradFunctor<platform::CUDADeviceContext, PoolProcess, T> {
const std::vector<int>& ksize, const std::vector<int>& ksize,
const std::vector<int>& strides, const std::vector<int>& strides,
const std::vector<int>& paddings, PoolProcess pool_process, const std::vector<int>& paddings, PoolProcess pool_process,
framework::Tensor* input_grad) { bool exclusive, framework::Tensor* input_grad) {
const int batch_size = input.dims()[0]; const int batch_size = input.dims()[0];
const int input_channels = input.dims()[1]; const int input_channels = input.dims()[1];
const int input_height = input.dims()[2]; const int input_height = input.dims()[2];
...@@ -236,7 +239,7 @@ class Pool2dGradFunctor<platform::CUDADeviceContext, PoolProcess, T> { ...@@ -236,7 +239,7 @@ class Pool2dGradFunctor<platform::CUDADeviceContext, PoolProcess, T> {
nthreads, input_data, output_data, output_grad_data, input_channels, nthreads, input_data, output_data, output_grad_data, input_channels,
input_height, input_width, output_height, output_width, ksize_height, input_height, input_width, output_height, output_width, ksize_height,
ksize_width, stride_height, stride_width, padding_height, padding_width, ksize_width, stride_height, stride_width, padding_height, padding_width,
pool_process, input_grad_data); pool_process, exclusive, input_grad_data);
} }
}; };
...@@ -313,16 +316,14 @@ template class Pool2dGradFunctor<platform::CUDADeviceContext, ...@@ -313,16 +316,14 @@ template class Pool2dGradFunctor<platform::CUDADeviceContext,
double>; double>;
template <typename PoolProcess, typename T> template <typename PoolProcess, typename T>
__global__ void KernelPool3D(const int nthreads, const T* input_data, __global__ void KernelPool3D(
const int channels, const int input_depth, const int nthreads, const T* input_data, const int channels,
const int input_height, const int input_width, const int input_depth, const int input_height, const int input_width,
const int output_depth, const int output_height, const int output_depth, const int output_height, const int output_width,
const int output_width, const int ksize_depth, const int ksize_depth, const int ksize_height, const int ksize_width,
const int ksize_height, const int ksize_width, const int stride_depth, const int stride_height, const int stride_width,
const int stride_depth, const int stride_height, const int padding_depth, const int padding_height, const int padding_width,
const int stride_width, const int padding_depth, PoolProcess pool_process, bool exclusive, T* output_data) {
const int padding_height, const int padding_width,
PoolProcess pool_process, T* output_data) {
for (int index = blockIdx.x * blockDim.x + threadIdx.x; index < nthreads; for (int index = blockIdx.x * blockDim.x + threadIdx.x; index < nthreads;
index += blockDim.x * gridDim.x) { index += blockDim.x * gridDim.x) {
int pw = index % output_width; int pw = index % output_width;
...@@ -351,7 +352,9 @@ __global__ void KernelPool3D(const int nthreads, const T* input_data, ...@@ -351,7 +352,9 @@ __global__ void KernelPool3D(const int nthreads, const T* input_data,
} }
} }
} }
int pool_size = (dend - dstart) * (hend - hstart) * (wend - wstart); int pool_size = exclusive ?
(dend - dstart) * (hend - hstart) * (wend - wstart)
: ksize_depth * ksize_height * ksize_width;
pool_process.finalize(static_cast<T>(pool_size), &ele); pool_process.finalize(static_cast<T>(pool_size), &ele);
output_data[index] = ele; output_data[index] = ele;
} }
...@@ -366,7 +369,7 @@ __global__ void KernelPool3DGrad( ...@@ -366,7 +369,7 @@ __global__ void KernelPool3DGrad(
const int ksize_height, const int ksize_width, const int stride_depth, const int ksize_height, const int ksize_width, const int stride_depth,
const int stride_height, const int stride_width, const int padding_depth, const int stride_height, const int stride_width, const int padding_depth,
const int padding_height, const int padding_width, PoolProcess pool_process, const int padding_height, const int padding_width, PoolProcess pool_process,
T* input_grad) { bool exclusive, T* input_grad) {
for (int index = blockIdx.x * blockDim.x + threadIdx.x; index < nthreads; for (int index = blockIdx.x * blockDim.x + threadIdx.x; index < nthreads;
index += blockDim.x * gridDim.x) { index += blockDim.x * gridDim.x) {
int offsetW = index % input_width + padding_width; int offsetW = index % input_width + padding_width;
...@@ -409,7 +412,9 @@ __global__ void KernelPool3DGrad( ...@@ -409,7 +412,9 @@ __global__ void KernelPool3DGrad(
dstart = max(dstart, 0); dstart = max(dstart, 0);
hstart = max(hstart, 0); hstart = max(hstart, 0);
wstart = max(wstart, 0); wstart = max(wstart, 0);
int pool_size = (dend - dstart) * (hend - hstart) * (wend - wstart); int pool_size = exclusive ?
(dend - dstart) * (hend - hstart) * (wend - wstart)
: ksize_depth * ksize_height * ksize_width;
int output_sub_idx = (pd * output_height + ph) * output_width + pw; int output_sub_idx = (pd * output_height + ph) * output_width + pw;
pool_process.compute(input, output_data[output_sub_idx], pool_process.compute(input, output_data[output_sub_idx],
output_grad[output_sub_idx], output_grad[output_sub_idx],
...@@ -484,7 +489,7 @@ class Pool3dFunctor<platform::CUDADeviceContext, PoolProcess, T> { ...@@ -484,7 +489,7 @@ class Pool3dFunctor<platform::CUDADeviceContext, PoolProcess, T> {
const framework::Tensor& input, const std::vector<int>& ksize, const framework::Tensor& input, const std::vector<int>& ksize,
const std::vector<int>& strides, const std::vector<int>& strides,
const std::vector<int>& paddings, PoolProcess pool_process, const std::vector<int>& paddings, PoolProcess pool_process,
framework::Tensor* output) { bool exclusive, framework::Tensor* output) {
const int batch_size = input.dims()[0]; const int batch_size = input.dims()[0];
const int input_channels = input.dims()[1]; const int input_channels = input.dims()[1];
const int input_depth = input.dims()[2]; const int input_depth = input.dims()[2];
...@@ -518,7 +523,7 @@ class Pool3dFunctor<platform::CUDADeviceContext, PoolProcess, T> { ...@@ -518,7 +523,7 @@ class Pool3dFunctor<platform::CUDADeviceContext, PoolProcess, T> {
input_width, output_depth, output_height, output_width, ksize_depth, input_width, output_depth, output_height, output_width, ksize_depth,
ksize_height, ksize_width, stride_depth, stride_height, stride_width, ksize_height, ksize_width, stride_depth, stride_height, stride_width,
padding_depth, padding_height, padding_width, pool_process, padding_depth, padding_height, padding_width, pool_process,
output_data); exclusive, output_data);
} }
}; };
...@@ -537,7 +542,7 @@ class Pool3dGradFunctor<platform::CUDADeviceContext, PoolProcess, T> { ...@@ -537,7 +542,7 @@ class Pool3dGradFunctor<platform::CUDADeviceContext, PoolProcess, T> {
const std::vector<int>& ksize, const std::vector<int>& ksize,
const std::vector<int>& strides, const std::vector<int>& strides,
const std::vector<int>& paddings, PoolProcess pool_process, const std::vector<int>& paddings, PoolProcess pool_process,
framework::Tensor* input_grad) { bool exclusive, framework::Tensor* input_grad) {
const int batch_size = input.dims()[0]; const int batch_size = input.dims()[0];
const int input_channels = input.dims()[1]; const int input_channels = input.dims()[1];
const int input_depth = input.dims()[2]; const int input_depth = input.dims()[2];
...@@ -573,7 +578,7 @@ class Pool3dGradFunctor<platform::CUDADeviceContext, PoolProcess, T> { ...@@ -573,7 +578,7 @@ class Pool3dGradFunctor<platform::CUDADeviceContext, PoolProcess, T> {
input_depth, input_height, input_width, output_depth, output_height, input_depth, input_height, input_width, output_depth, output_height,
output_width, ksize_depth, ksize_height, ksize_width, stride_depth, output_width, ksize_depth, ksize_height, ksize_width, stride_depth,
stride_height, stride_width, padding_depth, padding_height, stride_height, stride_width, padding_depth, padding_height,
padding_width, pool_process, input_grad_data); padding_width, pool_process, exclusive, input_grad_data);
} }
}; };
......
...@@ -89,7 +89,7 @@ class Pool2dFunctor { ...@@ -89,7 +89,7 @@ class Pool2dFunctor {
const std::vector<int>& ksize, const std::vector<int>& ksize,
const std::vector<int>& strides, const std::vector<int>& strides,
const std::vector<int>& paddings, PoolProcess pool_compute, const std::vector<int>& paddings, PoolProcess pool_compute,
framework::Tensor* output); bool exclusive, framework::Tensor* output);
}; };
template <typename DeviceContext, typename PoolProcess, typename T> template <typename DeviceContext, typename PoolProcess, typename T>
...@@ -101,7 +101,7 @@ class Pool2dGradFunctor { ...@@ -101,7 +101,7 @@ class Pool2dGradFunctor {
const std::vector<int>& ksize, const std::vector<int>& ksize,
const std::vector<int>& strides, const std::vector<int>& strides,
const std::vector<int>& paddings, PoolProcess pool_compute, const std::vector<int>& paddings, PoolProcess pool_compute,
framework::Tensor* input_grad); bool exclusive, framework::Tensor* input_grad);
}; };
template <typename DeviceContext, class T> template <typename DeviceContext, class T>
...@@ -123,7 +123,7 @@ class Pool3dFunctor { ...@@ -123,7 +123,7 @@ class Pool3dFunctor {
const std::vector<int>& ksize, const std::vector<int>& ksize,
const std::vector<int>& strides, const std::vector<int>& strides,
const std::vector<int>& paddings, PoolProcess pool_compute, const std::vector<int>& paddings, PoolProcess pool_compute,
framework::Tensor* output); bool exclusive, framework::Tensor* output);
}; };
template <typename DeviceContext, typename PoolProcess, typename T> template <typename DeviceContext, typename PoolProcess, typename T>
...@@ -135,7 +135,7 @@ class Pool3dGradFunctor { ...@@ -135,7 +135,7 @@ class Pool3dGradFunctor {
const std::vector<int>& ksize, const std::vector<int>& ksize,
const std::vector<int>& strides, const std::vector<int>& strides,
const std::vector<int>& paddings, PoolProcess pool_compute, const std::vector<int>& paddings, PoolProcess pool_compute,
framework::Tensor* input_grad); bool exclusive, framework::Tensor* input_grad);
}; };
template <typename DeviceContext, class T> template <typename DeviceContext, class T>
......
...@@ -41,6 +41,7 @@ class PoolCUDNNOpKernel : public framework::OpKernel<T> { ...@@ -41,6 +41,7 @@ class PoolCUDNNOpKernel : public framework::OpKernel<T> {
T *output_data = output->mutable_data<T>(ctx.GetPlace()); T *output_data = output->mutable_data<T>(ctx.GetPlace());
std::string pooling_type = ctx.Attr<std::string>("pooling_type"); std::string pooling_type = ctx.Attr<std::string>("pooling_type");
bool exclusive = ctx.Attr<bool>("exclusive");
std::vector<int> ksize = ctx.Attr<std::vector<int>>("ksize"); std::vector<int> ksize = ctx.Attr<std::vector<int>>("ksize");
std::vector<int> strides = ctx.Attr<std::vector<int>>("strides"); std::vector<int> strides = ctx.Attr<std::vector<int>>("strides");
std::vector<int> paddings = ctx.Attr<std::vector<int>>("paddings"); std::vector<int> paddings = ctx.Attr<std::vector<int>>("paddings");
...@@ -72,7 +73,7 @@ class PoolCUDNNOpKernel : public framework::OpKernel<T> { ...@@ -72,7 +73,7 @@ class PoolCUDNNOpKernel : public framework::OpKernel<T> {
if (pooling_type == "max") { if (pooling_type == "max") {
pooling_mode = PoolingMode::kMaximum; pooling_mode = PoolingMode::kMaximum;
} else { } else {
pooling_mode = PoolingMode::kAverage; pooling_mode = exclusive ? PoolingMode::kAverageExclusive : PoolingMode::kAverageInclusive;
} }
cudnnPoolingDescriptor_t cudnn_pool_desc = cudnnPoolingDescriptor_t cudnn_pool_desc =
...@@ -101,6 +102,7 @@ class PoolCUDNNGradOpKernel : public framework::OpKernel<T> { ...@@ -101,6 +102,7 @@ class PoolCUDNNGradOpKernel : public framework::OpKernel<T> {
Tensor *input_grad = ctx.Output<Tensor>(framework::GradVarName("X")); Tensor *input_grad = ctx.Output<Tensor>(framework::GradVarName("X"));
std::string pooling_type = ctx.Attr<std::string>("pooling_type"); std::string pooling_type = ctx.Attr<std::string>("pooling_type");
bool exclusive = ctx.Attr<bool>("exclusive");
std::vector<int> ksize = ctx.Attr<std::vector<int>>("ksize"); std::vector<int> ksize = ctx.Attr<std::vector<int>>("ksize");
std::vector<int> strides = ctx.Attr<std::vector<int>>("strides"); std::vector<int> strides = ctx.Attr<std::vector<int>>("strides");
std::vector<int> paddings = ctx.Attr<std::vector<int>>("paddings"); std::vector<int> paddings = ctx.Attr<std::vector<int>>("paddings");
...@@ -141,7 +143,7 @@ class PoolCUDNNGradOpKernel : public framework::OpKernel<T> { ...@@ -141,7 +143,7 @@ class PoolCUDNNGradOpKernel : public framework::OpKernel<T> {
pooling_mode = PoolingMode::kMaximum; pooling_mode = PoolingMode::kMaximum;
} }
} else { } else {
pooling_mode = PoolingMode::kAverage; pooling_mode = exclusive ? PoolingMode::kAverageExclusive : PoolingMode::kAverageInclusive;
} }
cudnnPoolingDescriptor_t cudnn_pool_desc = cudnnPoolingDescriptor_t cudnn_pool_desc =
......
...@@ -180,6 +180,12 @@ void Pool2dOpMaker::Make() { ...@@ -180,6 +180,12 @@ void Pool2dOpMaker::Make() {
"operator." "operator."
"If global_pooling = true, paddings and ksize will be ignored.") "If global_pooling = true, paddings and ksize will be ignored.")
.SetDefault({0, 0}); .SetDefault({0, 0});
AddAttr<bool>(
"exclusive",
"(bool, default True) When true, will exclude the zero-padding in the "
"averaging calculating, otherwise, include the zero-padding. Note, it "
"is only used when pooling_type is avg. The defalut is True.")
.SetDefault(true);
AddAttr<bool>( AddAttr<bool>(
"use_cudnn", "use_cudnn",
"(bool, default false) Only used in cudnn kernel, need install cudnn") "(bool, default false) Only used in cudnn kernel, need install cudnn")
...@@ -283,6 +289,12 @@ void Pool3dOpMaker::Make() { ...@@ -283,6 +289,12 @@ void Pool3dOpMaker::Make() {
"If global_pooling = true, ksize and paddings will be ignored.") "If global_pooling = true, ksize and paddings will be ignored.")
.SetDefault({0, 0, 0}); // TODO(Chengduo): Add checker. (Currently, .SetDefault({0, 0, 0}); // TODO(Chengduo): Add checker. (Currently,
// TypedAttrChecker don't support vector type.) // TypedAttrChecker don't support vector type.)
AddAttr<bool>(
"exclusive",
"(bool, default True) When true, will exclude the zero-padding in the "
"averaging calculating, otherwise, include the zero-padding. Note, it "
"is only used when pooling_type is avg. The defalut is True.")
.SetDefault(true);
AddAttr<bool>( AddAttr<bool>(
"use_cudnn", "use_cudnn",
......
...@@ -69,6 +69,7 @@ class PoolKernel : public framework::OpKernel<T> { ...@@ -69,6 +69,7 @@ class PoolKernel : public framework::OpKernel<T> {
std::vector<int> ksize = context.Attr<std::vector<int>>("ksize"); std::vector<int> ksize = context.Attr<std::vector<int>>("ksize");
std::vector<int> strides = context.Attr<std::vector<int>>("strides"); std::vector<int> strides = context.Attr<std::vector<int>>("strides");
std::vector<int> paddings = context.Attr<std::vector<int>>("paddings"); std::vector<int> paddings = context.Attr<std::vector<int>>("paddings");
bool exclusive = context.Attr<bool>("exclusive");
if (context.Attr<bool>("global_pooling")) { if (context.Attr<bool>("global_pooling")) {
for (size_t i = 0; i < ksize.size(); ++i) { for (size_t i = 0; i < ksize.size(); ++i) {
paddings[i] = 0; paddings[i] = 0;
...@@ -84,7 +85,7 @@ class PoolKernel : public framework::OpKernel<T> { ...@@ -84,7 +85,7 @@ class PoolKernel : public framework::OpKernel<T> {
pool2d_forward; pool2d_forward;
paddle::operators::math::MaxPool<T> pool_process; paddle::operators::math::MaxPool<T> pool_process;
pool2d_forward(dev_ctx, *in_x, ksize, strides, paddings, pool_process, pool2d_forward(dev_ctx, *in_x, ksize, strides, paddings, pool_process,
out); true, out);
} else if (pooling_type == "avg") { } else if (pooling_type == "avg") {
paddle::operators::math::Pool2dFunctor< paddle::operators::math::Pool2dFunctor<
...@@ -92,7 +93,7 @@ class PoolKernel : public framework::OpKernel<T> { ...@@ -92,7 +93,7 @@ class PoolKernel : public framework::OpKernel<T> {
pool2d_forward; pool2d_forward;
paddle::operators::math::AvgPool<T> pool_process; paddle::operators::math::AvgPool<T> pool_process;
pool2d_forward(dev_ctx, *in_x, ksize, strides, paddings, pool_process, pool2d_forward(dev_ctx, *in_x, ksize, strides, paddings, pool_process,
out); exclusive, out);
} }
} break; } break;
case 3: { case 3: {
...@@ -102,14 +103,14 @@ class PoolKernel : public framework::OpKernel<T> { ...@@ -102,14 +103,14 @@ class PoolKernel : public framework::OpKernel<T> {
pool3d_forward; pool3d_forward;
paddle::operators::math::MaxPool<T> pool_process; paddle::operators::math::MaxPool<T> pool_process;
pool3d_forward(dev_ctx, *in_x, ksize, strides, paddings, pool_process, pool3d_forward(dev_ctx, *in_x, ksize, strides, paddings, pool_process,
out); true, out);
} else if (pooling_type == "avg") { } else if (pooling_type == "avg") {
paddle::operators::math::Pool3dFunctor< paddle::operators::math::Pool3dFunctor<
DeviceContext, paddle::operators::math::AvgPool<T>, T> DeviceContext, paddle::operators::math::AvgPool<T>, T>
pool3d_forward; pool3d_forward;
paddle::operators::math::AvgPool<T> pool_process; paddle::operators::math::AvgPool<T> pool_process;
pool3d_forward(dev_ctx, *in_x, ksize, strides, paddings, pool_process, pool3d_forward(dev_ctx, *in_x, ksize, strides, paddings, pool_process,
out); exclusive, out);
} }
} break; } break;
default: { PADDLE_THROW("Pool op only supports 2D and 3D input."); } default: { PADDLE_THROW("Pool op only supports 2D and 3D input."); }
...@@ -131,6 +132,7 @@ class PoolGradKernel : public framework::OpKernel<T> { ...@@ -131,6 +132,7 @@ class PoolGradKernel : public framework::OpKernel<T> {
std::vector<int> ksize = context.Attr<std::vector<int>>("ksize"); std::vector<int> ksize = context.Attr<std::vector<int>>("ksize");
std::vector<int> strides = context.Attr<std::vector<int>>("strides"); std::vector<int> strides = context.Attr<std::vector<int>>("strides");
std::vector<int> paddings = context.Attr<std::vector<int>>("paddings"); std::vector<int> paddings = context.Attr<std::vector<int>>("paddings");
bool exclusive = context.Attr<bool>("exclusive");
if (context.Attr<bool>("global_pooling")) { if (context.Attr<bool>("global_pooling")) {
for (size_t i = 0; i < ksize.size(); ++i) { for (size_t i = 0; i < ksize.size(); ++i) {
...@@ -157,7 +159,7 @@ class PoolGradKernel : public framework::OpKernel<T> { ...@@ -157,7 +159,7 @@ class PoolGradKernel : public framework::OpKernel<T> {
pool2d_backward; pool2d_backward;
paddle::operators::math::AvgPoolGrad<T> pool_process; paddle::operators::math::AvgPoolGrad<T> pool_process;
pool2d_backward(dev_ctx, *in_x, *out, *out_grad, ksize, strides, pool2d_backward(dev_ctx, *in_x, *out, *out_grad, ksize, strides,
paddings, pool_process, in_x_grad); paddings, pool_process, exclusive, in_x_grad);
} }
} break; } break;
case 3: { case 3: {
...@@ -172,7 +174,7 @@ class PoolGradKernel : public framework::OpKernel<T> { ...@@ -172,7 +174,7 @@ class PoolGradKernel : public framework::OpKernel<T> {
pool3d_backward; pool3d_backward;
paddle::operators::math::AvgPoolGrad<T> pool_process; paddle::operators::math::AvgPoolGrad<T> pool_process;
pool3d_backward(dev_ctx, *in_x, *out, *out_grad, ksize, strides, pool3d_backward(dev_ctx, *in_x, *out, *out_grad, ksize, strides,
paddings, pool_process, in_x_grad); paddings, pool_process, exclusive, in_x_grad);
} }
} break; } break;
default: { PADDLE_THROW("Pool op only supports 2D and 3D input."); } default: { PADDLE_THROW("Pool op only supports 2D and 3D input."); }
......
...@@ -56,12 +56,14 @@ class SppKernel : public framework::OpKernel<T> { ...@@ -56,12 +56,14 @@ class SppKernel : public framework::OpKernel<T> {
math::Pool2dFunctor<DeviceContext, math::MaxPool<T>, T> pool_forward; math::Pool2dFunctor<DeviceContext, math::MaxPool<T>, T> pool_forward;
math::MaxPool<T> max_process; math::MaxPool<T> max_process;
pool_forward(context.template device_context<DeviceContext>(), *in_x, pool_forward(context.template device_context<DeviceContext>(), *in_x,
kernel_size, strides, paddings, max_process, &out_level); kernel_size, strides, paddings, max_process, true,
&out_level);
} else if (pooling_type == "avg") { } else if (pooling_type == "avg") {
math::Pool2dFunctor<DeviceContext, math::AvgPool<T>, T> pool_forward; math::Pool2dFunctor<DeviceContext, math::AvgPool<T>, T> pool_forward;
math::AvgPool<T> avg_process; math::AvgPool<T> avg_process;
pool_forward(context.template device_context<DeviceContext>(), *in_x, pool_forward(context.template device_context<DeviceContext>(), *in_x,
kernel_size, strides, paddings, avg_process, &out_level); kernel_size, strides, paddings, avg_process, true,
&out_level);
} }
// flatten pooling output shape // flatten pooling output shape
int output_flatten_w = in_x->dims()[1] * bins * bins; int output_flatten_w = in_x->dims()[1] * bins * bins;
...@@ -154,7 +156,7 @@ class SppGradKernel : public framework::OpKernel<T> { ...@@ -154,7 +156,7 @@ class SppGradKernel : public framework::OpKernel<T> {
math::AvgPoolGrad<T> avg_process; math::AvgPoolGrad<T> avg_process;
pool_backward(context.template device_context<DeviceContext>(), *in_x, pool_backward(context.template device_context<DeviceContext>(), *in_x,
*&out_level, *&outgrad_level, kernel_size, strides, *&out_level, *&outgrad_level, kernel_size, strides,
paddings, avg_process, in_x_grad); paddings, avg_process, true, in_x_grad);
} }
} }
} }
......
...@@ -76,8 +76,9 @@ enum class DataLayout { // Not use ...@@ -76,8 +76,9 @@ enum class DataLayout { // Not use
enum class PoolingMode { enum class PoolingMode {
kMaximum, kMaximum,
kAverage,
kMaximumDeterministic, kMaximumDeterministic,
kAverageExclusive,
kAverageInclusive,
}; };
#if CUDNN_VERSION < 6000 #if CUDNN_VERSION < 6000
...@@ -91,8 +92,10 @@ inline cudnnPoolingMode_t GetPoolingMode(const PoolingMode& mode) { ...@@ -91,8 +92,10 @@ inline cudnnPoolingMode_t GetPoolingMode(const PoolingMode& mode) {
switch (mode) { switch (mode) {
case PoolingMode::kMaximumDeterministic: case PoolingMode::kMaximumDeterministic:
return CUDNN_POOLING_MAX; return CUDNN_POOLING_MAX;
case PoolingMode::kAverage: case PoolingMode::kAverageExclusive:
return CUDNN_POOLING_AVERAGE_COUNT_EXCLUDE_PADDING; return CUDNN_POOLING_AVERAGE_COUNT_EXCLUDE_PADDING;
case PoolingMode::kAverageInclusive:
return CUDNN_POOLING_AVERAGE_COUNT_INCLUDE_PADDING;
case PoolingMode::kMaximum: case PoolingMode::kMaximum:
return CUDNN_POOLING_MAX; return CUDNN_POOLING_MAX;
default: default:
...@@ -105,8 +108,10 @@ inline cudnnPoolingMode_t GetPoolingMode(const PoolingMode& mode) { ...@@ -105,8 +108,10 @@ inline cudnnPoolingMode_t GetPoolingMode(const PoolingMode& mode) {
switch (mode) { switch (mode) {
case PoolingMode::kMaximumDeterministic: case PoolingMode::kMaximumDeterministic:
return CUDNN_POOLING_MAX_DETERMINISTIC; return CUDNN_POOLING_MAX_DETERMINISTIC;
case PoolingMode::kAverage: case PoolingMode::kAverageExclusive:
return CUDNN_POOLING_AVERAGE_COUNT_EXCLUDE_PADDING; return CUDNN_POOLING_AVERAGE_COUNT_EXCLUDE_PADDING;
case PoolingMode::kAverageInclusive:
return CUDNN_POOLING_AVERAGE_COUNT_INCLUDE_PADDING;
case PoolingMode::kMaximum: case PoolingMode::kMaximum:
return CUDNN_POOLING_MAX; return CUDNN_POOLING_MAX;
default: default:
......
...@@ -2067,6 +2067,7 @@ def pool2d(input, ...@@ -2067,6 +2067,7 @@ def pool2d(input,
global_pooling=False, global_pooling=False,
use_cudnn=True, use_cudnn=True,
ceil_mode=False, ceil_mode=False,
exclusive=True,
name=None): name=None):
""" """
${comment} ${comment}
...@@ -2081,9 +2082,11 @@ def pool2d(input, ...@@ -2081,9 +2082,11 @@ def pool2d(input,
pool_type: ${pooling_type_comment} pool_type: ${pooling_type_comment}
pool_stride (int): stride of the pooling layer. pool_stride (int): stride of the pooling layer.
pool_padding (int): padding size. pool_padding (int): padding size.
global_pooling: ${global_pooling_comment} global_pooling (bool): ${global_pooling_comment}
use_cudnn: ${use_cudnn_comment} use_cudnn (bool): ${use_cudnn_comment}
ceil_mode: ${ceil_mode_comment} ceil_mode (bool): ${ceil_mode_comment}
exclusive (bool): Whether to exclude padding points in average pooling
mode, default is true
name (str|None): A name for this layer(optional). If set None, the name (str|None): A name for this layer(optional). If set None, the
layer will be named automatically. layer will be named automatically.
...@@ -2143,7 +2146,8 @@ def pool2d(input, ...@@ -2143,7 +2146,8 @@ def pool2d(input,
"paddings": pool_padding, "paddings": pool_padding,
"use_cudnn": use_cudnn, "use_cudnn": use_cudnn,
"ceil_mode": ceil_mode, "ceil_mode": ceil_mode,
"use_mkldnn": False "use_mkldnn": False,
"exclusive": exclusive,
}) })
return pool_out return pool_out
...@@ -2157,6 +2161,7 @@ def pool3d(input, ...@@ -2157,6 +2161,7 @@ def pool3d(input,
global_pooling=False, global_pooling=False,
use_cudnn=True, use_cudnn=True,
ceil_mode=False, ceil_mode=False,
exclusive=True,
name=None): name=None):
""" """
This function adds the operator for pooling in 3-dimensions, using the This function adds the operator for pooling in 3-dimensions, using the
...@@ -2171,6 +2176,8 @@ def pool3d(input, ...@@ -2171,6 +2176,8 @@ def pool3d(input,
global_pooling (bool): ${global_pooling_comment} global_pooling (bool): ${global_pooling_comment}
use_cudnn (bool): ${use_cudnn_comment} use_cudnn (bool): ${use_cudnn_comment}
ceil_mode (bool): ${ceil_mode_comment} ceil_mode (bool): ${ceil_mode_comment}
exclusive (bool): Whether to exclude padding points in average pooling
mode, default is true
name (str): A name for this layer(optional). If set None, the layer name (str): A name for this layer(optional). If set None, the layer
will be named automatically. will be named automatically.
...@@ -2211,7 +2218,8 @@ def pool3d(input, ...@@ -2211,7 +2218,8 @@ def pool3d(input,
"paddings": pool_padding, "paddings": pool_padding,
"use_cudnn": use_cudnn, "use_cudnn": use_cudnn,
"ceil_mode": ceil_mode, "ceil_mode": ceil_mode,
"use_mkldnn": False "use_mkldnn": False,
"exclusive": exclusive,
}) })
return pool_out return pool_out
......
...@@ -26,7 +26,8 @@ def max_pool2D_forward_naive(x, ...@@ -26,7 +26,8 @@ def max_pool2D_forward_naive(x,
strides, strides,
paddings, paddings,
global_pool=0, global_pool=0,
ceil_mode=False): ceil_mode=False,
exclusive=True):
N, C, H, W = x.shape N, C, H, W = x.shape
if global_pool == 1: if global_pool == 1:
ksize = [H, W] ksize = [H, W]
...@@ -54,7 +55,8 @@ def avg_pool2D_forward_naive(x, ...@@ -54,7 +55,8 @@ def avg_pool2D_forward_naive(x,
strides, strides,
paddings, paddings,
global_pool=0, global_pool=0,
ceil_mode=False): ceil_mode=False,
exclusive=True):
N, C, H, W = x.shape N, C, H, W = x.shape
if global_pool == 1: if global_pool == 1:
ksize = [H, W] ksize = [H, W]
...@@ -73,8 +75,9 @@ def avg_pool2D_forward_naive(x, ...@@ -73,8 +75,9 @@ def avg_pool2D_forward_naive(x,
c_end = np.min((j * strides[1] + ksize[1] - paddings[1], W)) c_end = np.min((j * strides[1] + ksize[1] - paddings[1], W))
x_masked = x[:, :, r_start:r_end, c_start:c_end] x_masked = x[:, :, r_start:r_end, c_start:c_end]
out[:, :, i, j] = np.sum(x_masked, axis=(2, 3)) / ( field_size = ((r_end - r_start) * (c_end - c_start)) if exclusive \
(r_end - r_start) * (c_end - c_start)) else (ksize[0] * ksize[1])
out[:, :, i, j] = np.sum(x_masked, axis=(2, 3)) / field_size
return out return out
...@@ -89,12 +92,13 @@ class TestPool2d_Op(OpTest): ...@@ -89,12 +92,13 @@ class TestPool2d_Op(OpTest):
self.init_kernel_type() self.init_kernel_type()
self.init_pool_type() self.init_pool_type()
self.init_ceil_mode() self.init_ceil_mode()
self.init_exclusive()
if self.global_pool: if self.global_pool:
self.paddings = [0 for _ in range(len(self.paddings))] self.paddings = [0 for _ in range(len(self.paddings))]
input = np.random.random(self.shape).astype(self.dtype) input = np.random.random(self.shape).astype(self.dtype)
output = self.pool2D_forward_naive(input, self.ksize, self.strides, output = self.pool2D_forward_naive(input, self.ksize, self.strides,
self.paddings, self.global_pool, self.paddings, self.global_pool,
self.ceil_mode).astype(self.dtype) self.ceil_mode, self.exclusive).astype(self.dtype)
self.inputs = {'X': OpTest.np_dtype_to_fluid_dtype(input)} self.inputs = {'X': OpTest.np_dtype_to_fluid_dtype(input)}
self.attrs = { self.attrs = {
...@@ -106,7 +110,8 @@ class TestPool2d_Op(OpTest): ...@@ -106,7 +110,8 @@ class TestPool2d_Op(OpTest):
'use_cudnn': self.use_cudnn, 'use_cudnn': self.use_cudnn,
'use_mkldnn': self.use_mkldnn, 'use_mkldnn': self.use_mkldnn,
'ceil_mode': self.ceil_mode, 'ceil_mode': self.ceil_mode,
'data_format': 'AnyLayout' # TODO(dzhwinter) : should be fix latter 'data_format': 'AnyLayout', # TODO(dzhwinter) : should be fix latter
'exclusive': self.exclusive
} }
self.outputs = {'Out': output} self.outputs = {'Out': output}
...@@ -150,6 +155,9 @@ class TestPool2d_Op(OpTest): ...@@ -150,6 +155,9 @@ class TestPool2d_Op(OpTest):
def init_ceil_mode(self): def init_ceil_mode(self):
self.ceil_mode = False self.ceil_mode = False
def init_exclusive(self):
self.exclusive = True
class TestCase1(TestPool2d_Op): class TestCase1(TestPool2d_Op):
def init_test_case(self): def init_test_case(self):
...@@ -321,6 +329,14 @@ class TestCeilModeCase4(TestCase2): ...@@ -321,6 +329,14 @@ class TestCeilModeCase4(TestCase2):
def init_ceil_mode(self): def init_ceil_mode(self):
self.ceil_mode = True self.ceil_mode = True
class TestAvgInclude(TestCase2):
def init_exclusive(self):
self.exclusive = False
class TestCUDNNAvgInclude(TestCUDNNCase3):
def init_exclusive(self):
self.exclusive = False
if __name__ == '__main__': if __name__ == '__main__':
unittest.main() unittest.main()
...@@ -26,7 +26,8 @@ def max_pool3D_forward_naive(x, ...@@ -26,7 +26,8 @@ def max_pool3D_forward_naive(x,
strides, strides,
paddings, paddings,
global_pool=0, global_pool=0,
ceil_mode=False): ceil_mode=False,
exclusive=True):
N, C, D, H, W = x.shape N, C, D, H, W = x.shape
if global_pool == 1: if global_pool == 1:
ksize = [D, H, W] ksize = [D, H, W]
...@@ -60,7 +61,8 @@ def avg_pool3D_forward_naive(x, ...@@ -60,7 +61,8 @@ def avg_pool3D_forward_naive(x,
strides, strides,
paddings, paddings,
global_pool=0, global_pool=0,
ceil_mode=False): ceil_mode=False,
exclusive=True):
N, C, D, H, W = x.shape N, C, D, H, W = x.shape
if global_pool == 1: if global_pool == 1:
ksize = [D, H, W] ksize = [D, H, W]
...@@ -85,8 +87,9 @@ def avg_pool3D_forward_naive(x, ...@@ -85,8 +87,9 @@ def avg_pool3D_forward_naive(x,
w_end = np.min((j * strides[1] + ksize[1] - paddings[1], W)) w_end = np.min((j * strides[1] + ksize[1] - paddings[1], W))
x_masked = x[:, :, d_start:d_end, h_start:h_end, w_start:w_end] x_masked = x[:, :, d_start:d_end, h_start:h_end, w_start:w_end]
out[:, :, k, i, j] = np.sum(x_masked, axis=(2, 3, 4)) / ( field_size = (d_end - d_start) * (h_end - h_start) * (w_end - w_start) \
(d_end - d_start) * (h_end - h_start) * (w_end - w_start)) if exclusive else ksize[0] * ksize[1] * ksize[2]
out[:, :, k, i, j] = np.sum(x_masked, axis=(2, 3, 4)) / field_size
return out return out
...@@ -100,13 +103,14 @@ class TestPool3d_Op(OpTest): ...@@ -100,13 +103,14 @@ class TestPool3d_Op(OpTest):
self.init_kernel_type() self.init_kernel_type()
self.init_pool_type() self.init_pool_type()
self.init_ceil_mode() self.init_ceil_mode()
self.init_exclusive()
if self.global_pool: if self.global_pool:
self.paddings = [0 for _ in range(len(self.paddings))] self.paddings = [0 for _ in range(len(self.paddings))]
input = np.random.random(self.shape).astype(self.dtype) input = np.random.random(self.shape).astype(self.dtype)
output = self.pool3D_forward_naive(input, self.ksize, self.strides, output = self.pool3D_forward_naive(input, self.ksize, self.strides,
self.paddings, self.global_pool, self.paddings, self.global_pool,
self.ceil_mode).astype(self.dtype) self.ceil_mode, self.exclusive).astype(self.dtype)
self.inputs = {'X': OpTest.np_dtype_to_fluid_dtype(input)} self.inputs = {'X': OpTest.np_dtype_to_fluid_dtype(input)}
self.attrs = { self.attrs = {
...@@ -117,7 +121,8 @@ class TestPool3d_Op(OpTest): ...@@ -117,7 +121,8 @@ class TestPool3d_Op(OpTest):
'global_pooling': self.global_pool, 'global_pooling': self.global_pool,
'use_cudnn': self.use_cudnn, 'use_cudnn': self.use_cudnn,
'ceil_mode': self.ceil_mode, 'ceil_mode': self.ceil_mode,
'data_format': 'AnyLayout' # TODO(dzhwinter) : should be fix latter 'data_format': 'AnyLayout', # TODO(dzhwinter) : should be fix latter
'exclusive': self.exclusive
} }
self.outputs = {'Out': output} self.outputs = {'Out': output}
...@@ -161,6 +166,9 @@ class TestPool3d_Op(OpTest): ...@@ -161,6 +166,9 @@ class TestPool3d_Op(OpTest):
def init_ceil_mode(self): def init_ceil_mode(self):
self.ceil_mode = False self.ceil_mode = False
def init_exclusive(self):
self.exclusive = True
class TestCase1(TestPool3d_Op): class TestCase1(TestPool3d_Op):
def init_test_case(self): def init_test_case(self):
...@@ -332,6 +340,14 @@ class TestCeilModeCase4(TestCase2): ...@@ -332,6 +340,14 @@ class TestCeilModeCase4(TestCase2):
def init_ceil_mode(self): def init_ceil_mode(self):
self.ceil_mode = True self.ceil_mode = True
class TestAvgInclude(TestCase2):
def init_exclusive(self):
self.exclusive = False
class TestCUDNNAvgInclude(TestCUDNNCase3):
def init_exclusive(self):
self.exclusive = False
if __name__ == '__main__': if __name__ == '__main__':
unittest.main() unittest.main()
Markdown is supported
0% .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑!
想要评论请 注册