From 40688d223e86741c13faba76bd4986491cacf9bd Mon Sep 17 00:00:00 2001 From: chengduoZH Date: Wed, 18 Oct 2017 14:24:28 +0800 Subject: [PATCH] refine im2col (up_pad,down_pad) --- paddle/operators/math/im2col.cc | 43 ++++++++---- paddle/operators/math/im2col.cu | 43 ++++++++---- paddle/operators/math/im2col_test.cc | 90 ++++++++++++++++++++------ paddle/operators/sequence_project_op.h | 37 ++--------- 4 files changed, 135 insertions(+), 78 deletions(-) diff --git a/paddle/operators/math/im2col.cc b/paddle/operators/math/im2col.cc index 15b223479f9..729ba8665cf 100644 --- a/paddle/operators/math/im2col.cc +++ b/paddle/operators/math/im2col.cc @@ -140,11 +140,8 @@ class Im2ColFunctor= down_pad) { + row_begin = 0; + } else { + row_begin = down_pad - up_pad; + } + row_end = row_begin + ((input_height + up_pad + down_pad - filter_height) / + stride_height + + 1); + const T* im_data = im.data(); T* col_data = col.data(); @@ -204,12 +213,8 @@ class Col2ImFunctor { public: void operator()(const platform::DeviceContext& context, framework::Tensor& im, - const framework::Tensor& col, int stride, int pad, - int row_start, int row_end) { - int stride_height = stride; - int stride_width = 0; - int padding_height = pad; - int padding_width = 0; + const framework::Tensor& col, int stride_height, + int stride_width, int up_pad, int down_pad) { PADDLE_ENFORCE(im.dims().size() == 3); PADDLE_ENFORCE(col.dims().size() == 5); int input_channels = im.dims()[0]; @@ -220,10 +225,22 @@ class Col2ImFunctor= down_pad) { + row_begin = 0; + } else { + row_begin = down_pad - up_pad; + } + row_end = row_begin + ((input_height + up_pad + down_pad - filter_height) / + stride_height + + 1); + T* im_data = im.data(); const T* col_data = col.data(); - for (int col_row_idx = row_start; col_row_idx < row_end; ++col_row_idx) { + for (int col_row_idx = row_begin; col_row_idx < row_end; ++col_row_idx) { for (int col_col_idx = 0; col_col_idx < output_width; ++col_col_idx) { for (int channel = 0; channel < input_channels; ++channel) { for (int filter_row_idx = 0; filter_row_idx < filter_height; @@ -235,7 +252,7 @@ class Col2ImFunctor= down_pad) { + row_begin = 0; + } else { + row_begin = down_pad - up_pad; + } + row_end = row_begin + ((input_height + up_pad + down_pad - filter_height) / + stride_height + + 1); + int output_height = row_end - row_begin; // col.dims()[0]; int output_width = col.dims()[1]; @@ -295,7 +304,6 @@ __global__ void col2imOCF(T* im_data, const T* col_data, int input_channels, int row_end) { int swid = blockIdx.x; int shid = blockIdx.y; - // if (shid < row_begin || shid > row_end) return; for (int channelid = threadIdx.z; channelid < input_channels; channelid += blockDim.z) { for (int idy = threadIdx.y; idy < filter_height; idy += blockDim.y) { @@ -331,12 +339,8 @@ class Col2ImFunctor { public: void operator()(const platform::DeviceContext& context, framework::Tensor& im, - const framework::Tensor& col, int stride, int pad, - int row_begin, int row_end) { - int stride_height = stride; - int stride_width = 0; - int padding_height = pad; - int padding_width = 0; + const framework::Tensor& col, int stride_height, + int stride_width, int up_pad, int down_pad) { PADDLE_ENFORCE(im.dims().size() == 3); PADDLE_ENFORCE(col.dims().size() == 5); int input_channels = im.dims()[0]; @@ -344,6 +348,19 @@ class Col2ImFunctor= down_pad) { + row_begin = 0; + } else { + row_begin = down_pad - up_pad; + } + row_end = row_begin + ((input_height + up_pad + down_pad - filter_height) / + stride_height + + 1); + int output_height = row_end - row_begin; // col.dims()[0]; int output_width = col.dims()[1]; diff --git a/paddle/operators/math/im2col_test.cc b/paddle/operators/math/im2col_test.cc index 46de79af8fd..6406d43a9bc 100644 --- a/paddle/operators/math/im2col_test.cc +++ b/paddle/operators/math/im2col_test.cc @@ -35,6 +35,12 @@ void testIm2col() { * * output_ocf = [0, 1, 3, 4 * 1, 2, 4, 5] + * + * col2im_cfo = [0, 2, 2 + * 3, 4, 5] + * + * col2im_ocf = [0, 2, 2 + * 3, 4, 5] */ int input_height = 2; int input_width = 3; @@ -59,7 +65,7 @@ void testIm2col() { new paddle::platform::CUDADeviceContext(paddle::platform::GPUPlace()); #else PADDLE_THROW("no GPU support"); -#endif // PADDLE_ONLY_CPU +#endif // PADDLE_WITH_CUDA } if (paddle::platform::is_cpu_place(*place)) { input = input_tmp; @@ -71,6 +77,7 @@ void testIm2col() { output_ocf.mutable_data( {output_height, output_width, 1, filter_size, filter_size}, *place); + // Im2Col paddle::operators::math::Im2ColFunctor< paddle::operators::math::ColFormat::kCFO, Place, float> im2col; @@ -79,8 +86,12 @@ void testIm2col() { im2col_ocf; im2col(*context, input, output_cfo, stride, stride, padding, padding); - im2col_ocf(*context, input, output_ocf, stride, padding, 0, - output_height * output_width); + im2col_ocf(*context, input, output_ocf, /*stride_height*/ stride, + /*stride_width*/ stride, /*up_pad*/ padding, + /*down_pad*/ padding); + + float out_cfo_data[] = {0, 1, 1, 2, 3, 4, 4, 5}; + float out_ocf_data[] = {0, 1, 3, 4, 1, 2, 4, 5}; float* out_cfo_ptr; if (paddle::platform::is_cpu_place(*place)) { @@ -90,14 +101,9 @@ void testIm2col() { *context); out_cfo_ptr = output_tmp.data(); } - EXPECT_EQ(out_cfo_ptr[0], 0); - EXPECT_EQ(out_cfo_ptr[1], 1); - EXPECT_EQ(out_cfo_ptr[2], 1); - EXPECT_EQ(out_cfo_ptr[3], 2); - EXPECT_EQ(out_cfo_ptr[4], 3); - EXPECT_EQ(out_cfo_ptr[5], 4); - EXPECT_EQ(out_cfo_ptr[6], 4); - EXPECT_EQ(out_cfo_ptr[7], 5); + for (int i = 0; i < 6; ++i) { + EXPECT_EQ(out_cfo_ptr[i], out_cfo_data[i]); + } float* out_ocf_ptr; if (paddle::platform::is_cpu_place(*place)) { @@ -107,14 +113,60 @@ void testIm2col() { *context); out_ocf_ptr = output_tmp.data(); } - EXPECT_EQ(out_ocf_ptr[0], 0); - EXPECT_EQ(out_ocf_ptr[1], 1); - EXPECT_EQ(out_ocf_ptr[2], 3); - EXPECT_EQ(out_ocf_ptr[3], 4); - EXPECT_EQ(out_ocf_ptr[4], 1); - EXPECT_EQ(out_ocf_ptr[5], 2); - EXPECT_EQ(out_ocf_ptr[6], 4); - EXPECT_EQ(out_ocf_ptr[7], 5); + for (int i = 0; i < 6; ++i) { + EXPECT_EQ(out_ocf_ptr[i], out_ocf_data[i]); + } + + // Col2Im: kCFO + paddle::operators::math::Col2ImFunctor< + paddle::operators::math::ColFormat::kCFO, Place, float> + col2im; + paddle::operators::math::Col2ImFunctor< + paddle::operators::math::ColFormat::kOCF, Place, float> + col2im_ocf; + float col2im_data[] = {0, 2, 2, 3, 8, 5}; + + memset(input_ptr, 0, 6 * sizeof(float)); + if (paddle::platform::is_cpu_place(*place)) { + input = input_tmp; + } else { + input.CopyFrom(input_tmp, *place, *context); + } + + col2im(*context, input, output_cfo, stride, stride, padding, padding); + + float* in_ptr; + if (paddle::platform::is_cpu_place(*place)) { + in_ptr = input.data(); + } else { + input_tmp.CopyFrom(input, paddle::platform::CPUPlace(), *context); + in_ptr = input_tmp.data(); + } + for (int i = 0; i < 6; ++i) { + EXPECT_EQ(in_ptr[i], col2im_data[i]); + } + + // Col2Im: kOCF + memset(input_ptr, 0, 6 * sizeof(float)); + if (paddle::platform::is_cpu_place(*place)) { + input = input_tmp; + } else { + input.CopyFrom(input_tmp, *place, *context); + } + + col2im_ocf(*context, input, output_ocf, /*stride_height*/ stride, + /*stride_width*/ stride, /*up_pad*/ padding, + /*down_pad*/ padding); + + if (paddle::platform::is_cpu_place(*place)) { + in_ptr = input.data(); + } else { + input_tmp.CopyFrom(input, paddle::platform::CPUPlace(), *context); + in_ptr = input_tmp.data(); + } + for (int i = 0; i < 6; ++i) { + EXPECT_EQ(in_ptr[i], col2im_data[i]); + } } TEST(math, im2col) { diff --git a/paddle/operators/sequence_project_op.h b/paddle/operators/sequence_project_op.h index 6e911137a76..0a1b647070d 100644 --- a/paddle/operators/sequence_project_op.h +++ b/paddle/operators/sequence_project_op.h @@ -87,24 +87,9 @@ class SequenceProjectKernel : public framework::OpKernel { sequence_width}); // input_channels, input_height, input_width in_t.Resize(framework::make_ddim(input_shape)); for (int j = 0; j < context_length; ++j) { - int pad; - int row_start; - - if (up_pad != 0) { - pad = up_pad; - row_start = 0; - } else if (down_pad != 0) { - pad = down_pad; - row_start = down_pad; - } else { - pad = 0; - row_start = 0; - } - im2col_ocf(context.device_context(), in_t, out_t, - /*stride*/ context_stride, /*pad*/ pad, - /*row_start*/ row_start, - /*row_end*/ row_start + sequence_height); + /*stride_height*/ context_stride, /*stride_width*/ 0, up_pad, + down_pad); if (padding_trainable) { // add up trainable data out_t.Resize(framework::make_ddim( @@ -229,23 +214,9 @@ class SequenceProjectGradKernel : public framework::OpKernel { out_g_t.Resize(framework::make_ddim( {sequence_height, 1, 1, context_length, sequence_width})); - int pad; - int row_start; - - if (up_pad != 0) { - pad = up_pad; - row_start = 0; - } else if (down_pad != 0) { - pad = down_pad; - row_start = down_pad; - } else { - pad = 0; - row_start = 0; - } col2im_ocf(context.device_context(), in_g_t, out_g_t, - /*stride*/ context_stride, /*pad*/ pad, - /*row_start*/ row_start, - /*row_end*/ row_start + sequence_height); + /*stride_height*/ context_stride, /*stride_width*/ 0, up_pad, + down_pad); // out_g_t back to orign size } -- GitLab