From 32c9ecebc9e719c37cf778bfe7b79af5967dc036 Mon Sep 17 00:00:00 2001 From: Zhang Date: Mon, 4 Jun 2018 17:45:51 -0700 Subject: [PATCH 1/6] add sample ratio --- src/operator/contrib/roi_align-inl.h | 3 +++ src/operator/contrib/roi_align.cc | 4 ++-- src/operator/contrib/roi_align.cu | 4 ++-- 3 files changed, 7 insertions(+), 4 deletions(-) diff --git a/src/operator/contrib/roi_align-inl.h b/src/operator/contrib/roi_align-inl.h index 5ac420cc3d4a..263f72a6abc0 100644 --- a/src/operator/contrib/roi_align-inl.h +++ b/src/operator/contrib/roi_align-inl.h @@ -47,6 +47,7 @@ enum ROIAlignOpOutputs {kOut}; struct ROIAlignParam : public dmlc::Parameter { TShape pooled_size; float spatial_scale; + int sample_ratio; DMLC_DECLARE_PARAMETER(ROIAlignParam) { DMLC_DECLARE_FIELD(pooled_size) .set_expect_ndim(2).enforce_nonzero() @@ -54,6 +55,8 @@ struct ROIAlignParam : public dmlc::Parameter { DMLC_DECLARE_FIELD(spatial_scale).set_range(0.0, 1.0) .describe("Ratio of input feature map height (or w) to raw image height (or w). " "Equals the reciprocal of total stride in convolutional layers"); + DMLC_DECLARE_FIELD(sample_ratio).set_default(-1) + .describe("Optional sampling ratio of ROI align, using adaptive size by default."); } }; diff --git a/src/operator/contrib/roi_align.cc b/src/operator/contrib/roi_align.cc index c2cb929966a4..58395312ca9f 100644 --- a/src/operator/contrib/roi_align.cc +++ b/src/operator/contrib/roi_align.cc @@ -440,7 +440,7 @@ void ROIAlignForwardCompute(const nnvm::NodeAttrs& attrs, DType *top_data = out_data[roialign::kOut].dptr(); ROIAlignForward(count, bottom_data, param.spatial_scale, channels, - height, width, pooled_height, pooled_width, -1, bottom_rois, + height, width, pooled_height, pooled_width, param.sample_ratio, bottom_rois, rois_cols, top_data); }) } @@ -490,7 +490,7 @@ void ROIAlignBackwardCompute(const nnvm::NodeAttrs& attrs, } ROIAlignBackward(count, top_diff, num_rois, param.spatial_scale, channels, height, width, pooled_height, pooled_width, - -1, grad_in, bottom_rois, rois_cols); + param.sample_ratio, grad_in, bottom_rois, rois_cols); } if (kWriteTo == req[roialign::kBox]) { Fill(s, outputs[1], kWriteTo, static_cast(0)); diff --git a/src/operator/contrib/roi_align.cu b/src/operator/contrib/roi_align.cu index 21066ea15fab..7c600308e7bd 100644 --- a/src/operator/contrib/roi_align.cu +++ b/src/operator/contrib/roi_align.cu @@ -399,7 +399,7 @@ void ROIAlignForwardCompute(const nnvm::NodeAttrs& attrs, width, pooled_height, pooled_width, - -1, + param.sample_ratio, bottom_rois, top_data); }) @@ -467,7 +467,7 @@ void ROIAlignBackwardCompute(const nnvm::NodeAttrs& attrs, width, pooled_height, pooled_width, - -1, + param.sample_ratio, grad_in, bottom_rois); }) From 5d1b5d23cf55ef6c232c7a1d2627e218ae970191 Mon Sep 17 00:00:00 2001 From: Zhang Date: Mon, 4 Jun 2018 17:55:39 -0700 Subject: [PATCH 2/6] pylint --- src/operator/contrib/roi_align.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/operator/contrib/roi_align.cc b/src/operator/contrib/roi_align.cc index 58395312ca9f..22611273cf50 100644 --- a/src/operator/contrib/roi_align.cc +++ b/src/operator/contrib/roi_align.cc @@ -440,8 +440,8 @@ void ROIAlignForwardCompute(const nnvm::NodeAttrs& attrs, DType *top_data = out_data[roialign::kOut].dptr(); ROIAlignForward(count, bottom_data, param.spatial_scale, channels, - height, width, pooled_height, pooled_width, param.sample_ratio, bottom_rois, - rois_cols, top_data); + height, width, pooled_height, pooled_width, param.sample_ratio, + bottom_rois, rois_cols, top_data); }) } From f976e24a1e4abf03b89ae3053167b8382b22720c Mon Sep 17 00:00:00 2001 From: Zhang Date: Mon, 4 Jun 2018 22:39:16 -0700 Subject: [PATCH 3/6] increase size limit for bilinearup --- src/operator/contrib/bilinear_resize-inl.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/operator/contrib/bilinear_resize-inl.h b/src/operator/contrib/bilinear_resize-inl.h index b73ead9eba5e..c096f0149751 100644 --- a/src/operator/contrib/bilinear_resize-inl.h +++ b/src/operator/contrib/bilinear_resize-inl.h @@ -51,9 +51,9 @@ struct BilinearSampleParam : public dmlc::Parameter { int height; int width; DMLC_DECLARE_PARAMETER(BilinearSampleParam) { - DMLC_DECLARE_FIELD(height).set_range(1, 1000) + DMLC_DECLARE_FIELD(height).set_range(1, 10000) .describe("output height (required)"); - DMLC_DECLARE_FIELD(width).set_range(1, 1000) + DMLC_DECLARE_FIELD(width).set_range(1, 10000) .describe("output width (required)"); } }; From 5981b8b2dd5fa7689647a56922b9362dc5f37d09 Mon Sep 17 00:00:00 2001 From: Zhang Date: Tue, 5 Jun 2018 10:52:16 -0700 Subject: [PATCH 4/6] add test case --- tests/python/unittest/test_operator.py | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/tests/python/unittest/test_operator.py b/tests/python/unittest/test_operator.py index 1eb23cc92281..d7f3c7b18376 100644 --- a/tests/python/unittest/test_operator.py +++ b/tests/python/unittest/test_operator.py @@ -6162,7 +6162,7 @@ def roialign_forward_backward(data, rois, pooled_size, spatial_scale, sampling_r out[r, c, ph, pw] = val * 1.0 / count return out, [dx, drois] - def test_roi_align_value(): + def test_roi_align_value(sampling_ratio=0): ctx=default_context() dtype = np.float32 @@ -6173,7 +6173,6 @@ def test_roi_align_value(): pooled_size = (3, 4) spatial_scale = H * 1.0 / dlen - sampling_ratio = 0 data = mx.nd.array(np.arange(N*C*W*H).reshape((N,C,H,W)), ctx=ctx, dtype = dtype) # data = mx.nd.random.uniform(0, 1, (N, C, H, W), dtype = dtype) center_xy = mx.nd.random.uniform(0, dlen, (R, 2), ctx=ctx, dtype = dtype) @@ -6196,11 +6195,11 @@ def test_roi_align_value(): assert np.allclose(rois.grad.asnumpy(), drois) # modified from test_roipooling() - def test_roi_align_autograd(): + def test_roi_align_autograd(sampling_ratio=-1): ctx=default_context() data = mx.symbol.Variable(name='data') rois = mx.symbol.Variable(name='rois') - test = mx.symbol.contrib.ROIAlign(data=data, rois=rois, pooled_size=(4, 4), spatial_scale=1) + test = mx.symbol.contrib.ROIAlign(data=data, rois=rois, pooled_size=(4, 4), spatial_scale=1, sample_ratio=sampling_ratio) x1 = np.random.rand(4, 1, 12, 12).astype('float64') x2 = np.array([[0, 1.1, 1.1, 6.2, 6.2], [2, 6.1, 2.1, 8.2, 11.2], @@ -6214,7 +6213,9 @@ def test_roi_align_autograd(): numeric_eps=1e-4, rtol=1e-1, atol=1e-4, ctx=ctx) test_roi_align_value() + test_roi_align_value(2) test_roi_align_autograd() + test_roi_align_autograd(2) if __name__ == '__main__': From 34c3655c79ff37c27df379441af671010d61cd36 Mon Sep 17 00:00:00 2001 From: Zhang Date: Thu, 7 Jun 2018 13:13:53 -0700 Subject: [PATCH 5/6] fix typo --- tests/python/unittest/test_operator.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tests/python/unittest/test_operator.py b/tests/python/unittest/test_operator.py index d7f3c7b18376..377c83f7f66b 100644 --- a/tests/python/unittest/test_operator.py +++ b/tests/python/unittest/test_operator.py @@ -6185,7 +6185,7 @@ def test_roi_align_value(sampling_ratio=0): rois.attach_grad() with mx.autograd.record(): output = mx.nd.contrib.ROIAlign(data, rois, pooled_size=pooled_size, - spatial_scale=spatial_scale) + spatial_scale=spatial_scale, sample_ratio=sampling_ratio) dy = mx.nd.random.uniform(-1, 1, (R, C) + pooled_size, ctx=ctx, dtype = dtype) output.backward(dy) real_output, [dx, drois] = roialign_forward_backward(data.asnumpy(), rois.asnumpy(), pooled_size, spatial_scale, sampling_ratio, dy.asnumpy()) @@ -6195,7 +6195,7 @@ def test_roi_align_value(sampling_ratio=0): assert np.allclose(rois.grad.asnumpy(), drois) # modified from test_roipooling() - def test_roi_align_autograd(sampling_ratio=-1): + def test_roi_align_autograd(sampling_ratio=0): ctx=default_context() data = mx.symbol.Variable(name='data') rois = mx.symbol.Variable(name='rois') From 36c2a16cf2114d21dd46c07cc73ded2aeeaf4150 Mon Sep 17 00:00:00 2001 From: Zhang Date: Thu, 7 Jun 2018 17:18:51 -0700 Subject: [PATCH 6/6] rm comments and cpu back --- src/operator/contrib/roi_align.cu | 17 ----------------- tests/python/unittest/test_operator.py | 9 +++++---- 2 files changed, 5 insertions(+), 21 deletions(-) diff --git a/src/operator/contrib/roi_align.cu b/src/operator/contrib/roi_align.cu index 7c600308e7bd..d3db70b73b1a 100644 --- a/src/operator/contrib/roi_align.cu +++ b/src/operator/contrib/roi_align.cu @@ -231,13 +231,6 @@ __device__ void bilinear_interpolate_gradient( T lx = x - *x_low; T hy = 1. - ly, hx = 1. - lx; - // reference in forward - // T v1 = bottom_data[*y_low * width + *x_low]; - // T v2 = bottom_data[*y_low * width + *x_high]; - // T v3 = bottom_data[*y_high * width + *x_low]; - // T v4 = bottom_data[*y_high * width + *x_high]; - // T val = (w1 * v1 + *w2 * v2 + *w3 * v3 + *w4 * v4); - *w1 = hy * hx, *w2 = hy * lx, *w3 = ly * hx, *w4 = ly * lx; return; @@ -341,16 +334,6 @@ __global__ void RoIAlignBackwardKernel( offset_bottom_diff + y_high * width + x_low, static_cast(g3)); atomicAdd( offset_bottom_diff + y_high * width + x_high, static_cast(g4)); - /* - gpu_atomic_add( - static_cast(g1), offset_bottom_diff + y_low * width + x_low); - gpu_atomic_add( - static_cast(g2), offset_bottom_diff + y_low * width + x_high); - gpu_atomic_add( - static_cast(g3), offset_bottom_diff + y_high * width + x_low); - gpu_atomic_add( - static_cast(g4), offset_bottom_diff + y_high * width + x_high); - */ } // if } // ix } // iy diff --git a/tests/python/unittest/test_operator.py b/tests/python/unittest/test_operator.py index a34db73b5d22..b28c10253fcf 100644 --- a/tests/python/unittest/test_operator.py +++ b/tests/python/unittest/test_operator.py @@ -6233,7 +6233,8 @@ def test_roi_align_value(sampling_ratio=0): spatial_scale=spatial_scale, sample_ratio=sampling_ratio) dy = mx.nd.random.uniform(-1, 1, (R, C) + pooled_size, ctx=ctx, dtype = dtype) output.backward(dy) - real_output, [dx, drois] = roialign_forward_backward(data.asnumpy(), rois.asnumpy(), pooled_size, spatial_scale, sampling_ratio, dy.asnumpy()) + real_output, [dx, drois] = roialign_forward_backward(data.asnumpy(), rois.asnumpy(), pooled_size, + spatial_scale, sampling_ratio, dy.asnumpy()) assert np.allclose(output.asnumpy(), real_output) # It seems that the precision between Cfloat and Pyfloat is different. assert np.allclose(data.grad.asnumpy(), dx, atol = 1e-5), np.abs(data.grad.asnumpy() - dx).max() @@ -6241,10 +6242,11 @@ def test_roi_align_value(sampling_ratio=0): # modified from test_roipooling() def test_roi_align_autograd(sampling_ratio=0): - ctx=default_context() + ctx = default_context() data = mx.symbol.Variable(name='data') rois = mx.symbol.Variable(name='rois') - test = mx.symbol.contrib.ROIAlign(data=data, rois=rois, pooled_size=(4, 4), spatial_scale=1, sample_ratio=sampling_ratio) + test = mx.symbol.contrib.ROIAlign(data=data, rois=rois, pooled_size=(4, 4), spatial_scale=1, + sample_ratio=sampling_ratio) x1 = np.random.rand(4, 1, 12, 12).astype('float64') x2 = np.array([[0, 1.1, 1.1, 6.2, 6.2], [2, 6.1, 2.1, 8.2, 11.2], @@ -6260,7 +6262,6 @@ def test_roi_align_autograd(sampling_ratio=0): test_roi_align_value() test_roi_align_value(2) test_roi_align_autograd() - test_roi_align_autograd(2) if __name__ == '__main__':