Skip to content
This repository has been archived by the owner on Nov 17, 2023. It is now read-only.

[MXNET-80] Fix average pooling kernel size assignment error #10000

Merged
merged 4 commits into from
Apr 9, 2018
Merged

[MXNET-80] Fix average pooling kernel size assignment error #10000

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
d6c0da8
fix average pooling kernel size assignment error
CoinCheung Mar 6, 2018
6d88c35
no order change and test kernel=
CoinCheung Mar 23, 2018
f882365
change order
CoinCheung Mar 23, 2018
04b60df
Merge branch 'master' into Coin-branch
CoinCheung Apr 9, 2018
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
49 changes: 31 additions & 18 deletions src/operator/nn/pooling-inl.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -50,22 +50,22 @@ struct PoolingParam : public dmlc::Parameter<PoolingParam> {
bool global_pool;
bool cudnn_off;
DMLC_DECLARE_PARAMETER(PoolingParam) {
DMLC_DECLARE_FIELD(global_pool).set_default(false)
.describe("Ignore kernel size, do global pooling based on current input feature map. ");

DMLC_DECLARE_FIELD(cudnn_off).set_default(false)
.describe("Turn off cudnn pooling and use MXNet pooling operator. ");

DMLC_DECLARE_FIELD(kernel)
DMLC_DECLARE_FIELD(kernel).set_default(TShape()) // add default value here
.enforce_nonzero()
.describe("Pooling kernel size: (y, x) or (d, y, x)");

DMLC_DECLARE_FIELD(pool_type)
DMLC_DECLARE_FIELD(pool_type).set_default(pool_enum::kMaxPooling) // add default pooling method
.add_enum("max", pool_enum::kMaxPooling)
.add_enum("avg", pool_enum::kAvgPooling)
.add_enum("sum", pool_enum::kSumPooling)
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I realized that we need to change the order of the DMLC_DECLARE_FIELD here. In the original version, the parameters that do not have default values will be set first and then goes the params with default values. So the order will be kernel, pool_type, global_pool, cudnn_off, ... After we add default values to kernel and pool_type, the order becomes global_pool, cudnn_off, kernel, pool_type, ... Thus, the way to solve the problem is to change the order:

    DMLC_DECLARE_FIELD(kernel).set_default(TShape())  // add default value here
    .enforce_nonzero()
    .describe("Pooling kernel size: (y, x) or (d, y, x)");

    DMLC_DECLARE_FIELD(pool_type).set_default(pool_enum::kMaxPooling)  // add default pooling method
    .add_enum("max", pool_enum::kMaxPooling)
    .add_enum("avg", pool_enum::kAvgPooling)
    .add_enum("sum", pool_enum::kSumPooling)
    .describe("Pooling type to be applied.");

    DMLC_DECLARE_FIELD(global_pool).set_default(false)
    .describe("Ignore kernel size, do global pooling based on current input feature map. ");

    DMLC_DECLARE_FIELD(cudnn_off).set_default(false)
    .describe("Turn off cudnn pooling and use MXNet pooling operator. ");

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

But I see in the original version that the order is: global_pool(false), cudnn_off(false), kernel(no default value), pool_type(no default value), ....

In the original version, "kernel" and "pool_type" do not have default value, but they go after "global_pool" and "cudnn_off" which are with default values.

https://github.com/apache/incubator-mxnet/blob/94f68fc8fd21611b7f5c148cb0e5d134efe58f87/src/operator/nn/pooling-inl.h#L59

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Sorry, maybe I have misunderstood what you said, I will have a try.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I tried a few times, and it failed at this position:
https://github.com/apache/incubator-mxnet/blob/94f68fc8fd21611b7f5c148cb0e5d134efe58f87/src/operator/nn/pooling.cc#L55
But I do not understand why it requires stride and kernel have the same length.

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

There's no need to check this if global_pool is turned on.

.describe("Pooling type to be applied.");

DMLC_DECLARE_FIELD(global_pool).set_default(false)
.describe("Ignore kernel size, do global pooling based on current input feature map. ");

DMLC_DECLARE_FIELD(cudnn_off).set_default(false)
.describe("Turn off cudnn pooling and use MXNet pooling operator. ");

DMLC_DECLARE_FIELD(pooling_convention).set_default(pool_enum::kValid)
.add_enum("full", pool_enum::kFull)
.add_enum("valid", pool_enum::kValid)
Expand Down Expand Up @@ -132,19 +132,23 @@ class PoolingOp {
using namespace mshadow;
Stream<xpu> *s = ctx.get_stream<xpu>();
const TShape& ishape = in_data.shape_;
TShape kernel = param_.kernel;
TShape padding = param_.pad;
TShape stride = param_.stride;
if (param_.global_pool) {
for (index_t i = 0; i < padding.ndim(); i++) {
kernel = TShape(ishape.data() + 2,
ishape.data() + ishape.ndim());
padding = TShape(ishape.ndim() - 2);
for (index_t i = 0; i < ishape.ndim() - 2; i++) {
padding[i] = 0;
}
stride = TShape(ishape.ndim() - 2);
}

pool(s, in_data.dptr<DType>(), in_data.shape_, out_data.shape_,
param_.global_pool?
TShape(ishape.data()+ishape.ndim()-param_.kernel.ndim(), ishape.data()+ishape.ndim())
: param_.kernel,
kernel,
padding,
param_.global_pool? TShape(param_.kernel.ndim()) : param_.stride,
stride,
param_.pool_type, req, out_data.dptr<DType>());
}

Expand All @@ -154,20 +158,24 @@ class PoolingOp {
using namespace mshadow;
Stream<xpu> *s = ctx.get_stream<xpu>();
const TShape& ishape = in_data.shape_;
TShape kernel = param_.kernel;
TShape padding = param_.pad;
TShape stride = param_.stride;
if (param_.global_pool) {
for (index_t i = 0; i < padding.ndim(); i++) {
kernel = TShape(ishape.data() + 2,
ishape.data() + ishape.ndim());
padding = TShape(ishape.ndim() - 2);
for (index_t i = 0; i < ishape.ndim() - 2; i++) {
padding[i] = 0;
}
stride = TShape(ishape.ndim() - 2);
}

unpool(s, out_grad.dptr<DType>(), in_data.dptr<DType>(), out_data.dptr<DType>(),
in_grad.shape_, out_grad.shape_,
param_.global_pool?
TShape(ishape.data()+ishape.ndim()-param_.kernel.ndim(), ishape.data()+ishape.ndim())
: param_.kernel,
kernel,
padding,
param_.global_pool? TShape(param_.kernel.ndim()) : param_.stride,
stride,
param_.pool_type, req, in_grad.dptr<DType>());
}

Expand All @@ -178,6 +186,11 @@ class PoolingOp {
template<typename xpu, typename DType>
PoolingOp<xpu, DType> &GetPoolingOp(const PoolingParam &param) {
static thread_local PoolingOp<xpu, DType> op;
// check if filter size assigned correctly
if (param.global_pool == false) {
CHECK_GT(param.kernel.ndim(), 0U)
<< "You need to set the kernel size if global pooling is not used";
}
op.Init(param);
return op;
}
Expand Down
155 changes: 76 additions & 79 deletions src/operator/nn/pooling.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -46,15 +46,14 @@ static void PoolingParamParser(nnvm::NodeAttrs *attrs) {
if (param.stride.ndim() == 0) param.stride = Shape2(1, 1);
if (param.pad.ndim() == 0) param.pad = Shape2(0, 0);
} else {
CHECK_EQ(param.kernel.ndim(), 3U) << param.kernel.ndim()
<< "D pooling not supported";
// ignore kernel size only if global_pool not assigned false
if (param.global_pool == false) {
CHECK_EQ(param.kernel.ndim(), 3U) << param.kernel.ndim()
<< "D pooling not supported";
}
if (param.stride.ndim() == 0) param.stride = Shape3(1, 1, 1);
if (param.pad.ndim() == 0) param.pad = Shape3(0, 0, 0);
}
CHECK_EQ(param.stride.ndim(), param.kernel.ndim())
<< "stride and kernel should have the same length";
CHECK_EQ(param.pad.ndim(), param.kernel.ndim())
<< "pad and kernel should have the same length";
attrs->parsed = std::move(param);
}

Expand Down Expand Up @@ -98,28 +97,37 @@ static bool PoolingShape(const nnvm::NodeAttrs &attrs,
<< "Pooling: Input data should be 3D in (batch, channel, x)"
<< " Or 4D in (batch, channel, y, x) "
<< " Or 5D in (batch, channel, d, y, x)";
CHECK_LE(dshape.ndim(), 5U)
<< "Pooling: Input data should be 3D in (batch, channel, x)"
<< " Or 4D in (batch, channel, y, x) "
<< " Or 5D in (batch, channel, d, y, x)";
TShape oshape = dshape;
if (dshape.ndim() == 0) return false;
if (param.kernel.ndim() == 1) {
if (param.global_pool) {
for (size_t i{2}; i < dshape.ndim(); i++)
oshape[i] = 1;
out_shape->clear();
out_shape->push_back(oshape); // save output shape
#if MXNET_USE_MKLDNN == 1
if (MKLDNNRequireWorkspace(param) && SupportMKLDNNPooling(param))
out_shape->push_back(oshape); // for workspace
#endif
} else if (param.kernel.ndim() == 1) {
CHECK_EQ(dshape.ndim(), 3U)
<< "Pooling: Input data should be 3D in (batch, channel, x)";
if (param.global_pool) {
oshape[2] = 1;
CHECK(param.kernel[0] <= dshape[2] + 2 * param.pad[0])
<< "kernel size (" << param.kernel[0] << ") exceeds input ("
<< dshape[2] << " padded to " << (dshape[2] + 2 * param.pad[0])
<< ")";
if (param.pooling_convention == pool_enum::kValid) {
oshape[2] = 1 +
(dshape[2] + 2 * param.pad[0] - param.kernel[0]) /
param.stride[0];
} else {
CHECK(param.kernel[0] <= dshape[2] + 2 * param.pad[0])
<< "kernel size (" << param.kernel[0] << ") exceeds input ("
<< dshape[2] << " padded to " << (dshape[2] + 2 * param.pad[0])
<< ")";
if (param.pooling_convention == pool_enum::kValid) {
oshape[2] = 1 +
(dshape[2] + 2 * param.pad[0] - param.kernel[0]) /
param.stride[0];
} else {
oshape[2] = 1 + static_cast<int>(ceil(
static_cast<float>(dshape[2] + 2 * param.pad[0] -
param.kernel[0]) /
param.stride[0]));
}
oshape[2] = 1 + static_cast<int>(ceil(
static_cast<float>(dshape[2] + 2 * param.pad[0] -
param.kernel[0]) /
param.stride[0]));
}
out_shape->clear();
out_shape->push_back(oshape); // save output shape
Expand All @@ -130,35 +138,30 @@ static bool PoolingShape(const nnvm::NodeAttrs &attrs,
} else if (param.kernel.ndim() == 2) {
CHECK_EQ(dshape.ndim(), 4U)
<< "Pooling: Input data should be 4D in (batch, channel, y, x)";
if (param.global_pool) {
oshape[2] = 1;
oshape[3] = 1;
CHECK(param.kernel[0] <= dshape[2] + 2 * param.pad[0])
<< "kernel size (" << param.kernel[0] << ") exceeds input ("
<< dshape[2] << " padded to " << (dshape[2] + 2 * param.pad[0])
<< ")";
CHECK(param.kernel[1] <= dshape[3] + 2 * param.pad[1])
<< "kernel size (" << param.kernel[1] << ") exceeds input ("
<< dshape[3] << " padded to " << (dshape[3] + 2 * param.pad[1])
<< ")";
if (param.pooling_convention == pool_enum::kValid) {
oshape[2] = 1 +
(dshape[2] + 2 * param.pad[0] - param.kernel[0]) /
param.stride[0];
oshape[3] = 1 +
(dshape[3] + 2 * param.pad[1] - param.kernel[1]) /
param.stride[1];
} else {
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The if (param.global_pool) is removed and the else should also be removed.

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Ignore this comment. I've misinterpreted the code.

CHECK(param.kernel[0] <= dshape[2] + 2 * param.pad[0])
<< "kernel size (" << param.kernel[0] << ") exceeds input ("
<< dshape[2] << " padded to " << (dshape[2] + 2 * param.pad[0])
<< ")";
CHECK(param.kernel[1] <= dshape[3] + 2 * param.pad[1])
<< "kernel size (" << param.kernel[1] << ") exceeds input ("
<< dshape[3] << " padded to " << (dshape[3] + 2 * param.pad[1])
<< ")";
if (param.pooling_convention == pool_enum::kValid) {
oshape[2] = 1 +
(dshape[2] + 2 * param.pad[0] - param.kernel[0]) /
param.stride[0];
oshape[3] = 1 +
(dshape[3] + 2 * param.pad[1] - param.kernel[1]) /
param.stride[1];
} else {
oshape[2] = 1 + static_cast<int>(ceil(
static_cast<float>(dshape[2] + 2 * param.pad[0] -
param.kernel[0]) /
param.stride[0]));
oshape[3] = 1 + static_cast<int>(ceil(
static_cast<float>(dshape[3] + 2 * param.pad[1] -
param.kernel[1]) /
param.stride[1]));
}
oshape[2] = 1 + static_cast<int>(ceil(
static_cast<float>(dshape[2] + 2 * param.pad[0] -
param.kernel[0]) /
param.stride[0]));
oshape[3] = 1 + static_cast<int>(ceil(
static_cast<float>(dshape[3] + 2 * param.pad[1] -
param.kernel[1]) /
param.stride[1]));
}
out_shape->clear();
out_shape->push_back(oshape); // save output shape
Expand All @@ -175,35 +178,29 @@ static bool PoolingShape(const nnvm::NodeAttrs &attrs,
<< "kernel size exceeds input";
CHECK_LE(param.kernel[2], dshape[4] + 2 * param.pad[2])
<< "kernel size exceeds input";
if (param.global_pool) {
oshape[2] = 1;
oshape[3] = 1;
oshape[4] = 1;
if (param.pooling_convention == pool_enum::kValid) {
oshape[2] = 1 +
(dshape[2] + 2 * param.pad[0] - param.kernel[0]) /
param.stride[0];
oshape[3] = 1 +
(dshape[3] + 2 * param.pad[1] - param.kernel[1]) /
param.stride[1];
oshape[4] = 1 +
(dshape[4] + 2 * param.pad[2] - param.kernel[2]) /
param.stride[2];
} else {
if (param.pooling_convention == pool_enum::kValid) {
oshape[2] = 1 +
(dshape[2] + 2 * param.pad[0] - param.kernel[0]) /
param.stride[0];
oshape[3] = 1 +
(dshape[3] + 2 * param.pad[1] - param.kernel[1]) /
param.stride[1];
oshape[4] = 1 +
(dshape[4] + 2 * param.pad[2] - param.kernel[2]) /
param.stride[2];
} else {
oshape[2] = 1 + static_cast<int>(ceil(
static_cast<float>(dshape[2] + 2 * param.pad[0] -
param.kernel[0]) /
param.stride[0]));
oshape[3] = 1 + static_cast<int>(ceil(
static_cast<float>(dshape[3] + 2 * param.pad[1] -
param.kernel[1]) /
param.stride[1]));
oshape[4] = 1 + static_cast<int>(ceil(
static_cast<float>(dshape[4] + 2 * param.pad[2] -
param.kernel[2]) /
param.stride[2]));
}
oshape[2] = 1 + static_cast<int>(ceil(
static_cast<float>(dshape[2] + 2 * param.pad[0] -
param.kernel[0]) /
param.stride[0]));
oshape[3] = 1 + static_cast<int>(ceil(
static_cast<float>(dshape[3] + 2 * param.pad[1] -
param.kernel[1]) /
param.stride[1]));
oshape[4] = 1 + static_cast<int>(ceil(
static_cast<float>(dshape[4] + 2 * param.pad[2] -
param.kernel[2]) /
param.stride[2]));
}

out_shape->clear();
Expand Down
Loading