Add Conv3DLayer, DeConv3DLayer #3348

chengduoZH · 2017-08-09T04:40:07Z

fix #3359
Add Conv3DLayer, DeConv3DLayer.

hedaoyuan

本地环境装一下pre-commit，提交代码前需要format.
https://github.com/PaddlePaddle/Paddle/blob/e4c97e48e58edd7d449b77abd5b8caf6b658577c/doc/howto/dev/contribute_to_paddle_cn.md#使用-pre-commit-钩子

hedaoyuan · 2017-08-14T12:45:56Z

paddle/math/Matrix.h

@@ -1039,6 +1039,42 @@ class Matrix : public BaseMatrix {
    LOG(FATAL) << "Not implemented";
  }

+  virtual void vol2Col(real* data,


在test_matrixCompare.cpp里面增加一下vol2Col，col2Vol的CPU和GPU实现check吧。

hedaoyuan · 2017-08-14T12:56:10Z

paddle/gserver/layers/Conv3DLayer.cpp

+      conf.filter_size_y() * conf.filter_size());
+      weights_[index]->getW()->reshape(
+              weights_[index]->getW()->getWidth(),
+              weights_[index]->getW()->getHeight());


34-36行的reshape是做什么用的？

在ConqBaseLayer中创建的weights_的shape在Conv3DLayer中不适用，所以在Conv3DLayer中需要把weights_的width和height颠倒过来，不过这个在最新的提交中做了做了调整，即把创建weights_的操作写在派生类中

hedaoyuan · 2017-08-14T13:04:43Z

paddle/gserver/layers/Conv3DLayer.cpp

+              MatrixPtr wMatSub = wMat->subMatrix(g * M, M);
+              MatrixPtr in = colBuf_->subMatrix(g * K, K);
+              MatrixPtr out = outMatSub->subMatrix(g * M, M);
+              out->mul(*wMatSub, *in, 1.0, 0.0);


0.0需要修改成1.0，否则inputLayers_.size() > 1时就错了。

hedaoyuan · 2017-08-14T13:09:30Z

paddle/gserver/layers/ConvBaseLayer.cpp

+  isDeconv_ = (config_.type() == "exconv" ||
+          config_.type() == "cudnn_conv" ||
+          config_.type() == "conv3d" ||
+          config_.type() == "deconv3d"   )


27行不能有。

这个得有，不然会出错，在DeConv3D中filter数据排布方式是： (output * depth * height * weigth) x channel，所以
height = filterPixels_[i] * filterChannels_[i]; width = (!isDeconv_) ? numFilters_ : channels_[i];中width要等与numFilters_

hedaoyuan · 2017-08-14T13:13:46Z

paddle/gserver/layers/DeConv3DLayer.cpp

+          wGradData += K * M;
+          inData += M * N;
+        }
+        weights_[i]->getParameterPtr()->incUpdate(callback);


171行放到batchSize循环外面。

qingqing01

Still need:

parser in python/paddle/trainer/config_parser.py
Python interface

qingqing01 · 2017-08-09T11:15:13Z

paddle/cuda/src/hl_cuda_matrix.cu

+
+
+__global__ void keMatrixVol2Col(
+        int num_kernels, real*data, real* data_col,


realdata -> const real data

qingqing01 · 2017-08-18T08:50:59Z

paddle/cuda/src/hl_cuda_matrix.cu

+  const int threads = 512;
+  const int blocks = DIVUP(num_kernels, threads);
+
+  keMatrixVol2Col<<<blocks, threads>>>(num_kernels,


Use the STREAM_DEFAULT:

keMatrixVol2Col <<< blocks, threads, 0, STREAM_DEFAULT>>>

The STREAM_DEFAULT is defined as follows. It's a thread local stream.

./cuda/include/hl_base.h:215:extern __thread cudaStream_t default_stream; ./cuda/include/hl_base.h:216:#define STREAM_DEFAULT default_stream

qingqing01 · 2017-08-18T08:51:29Z

paddle/cuda/src/hl_cuda_matrix.cu

+  const int threads = 512;
+  const int blocks = DIVUP(num_kernels, threads);
+
+  keMatrixCol2Vol<<<blocks, threads>>>(num_kernels,


same as above comments.

qingqing01 · 2017-08-18T08:58:10Z

paddle/gserver/layers/Conv3DLayer.cpp

+    if (nullptr != weights_[index]->getWGrad())
+      weights_[index]->getWGrad()->reshape(
+          weights_[index]->getWGrad()->getWidth(),
+          weights_[index]->getWGrad()->getHeight());


It's better to modify the weight shape in the base layer. Or move the creation in the base class and create weight in the inherited class.

qingqing01 · 2017-08-18T10:18:44Z

paddle/gserver/layers/Conv3DLayer.cpp

+          weights_[index]->getWGrad()->getHeight());
+    ++index;
+  }
+  CHECK(inputLayers_.size() == parameters_.size());


This line is not correct. Remove this check. If inputLayers_.size() == 1 and this layer has bias, the parameters_.size() == 2, so it is not correct.

So the 58 line of ConqbaseLayer.cpp should be removed?

qingqing01 · 2017-08-18T10:33:28Z

paddle/gserver/layers/DeConv3DLayer.cpp

+  } else {
+    outMat->addBias(*(biases_->getW()), 1.0f);
+  }
+}


In fact, the same problem as the implementing 2D-Deconv before. Many similar code with 3D-Conv.

@hedaoyuan How do you think of this problem?

qingqing01 · 2017-08-18T10:36:17Z

paddle/gserver/layers/DeConv3DLayer.cpp

+REGISTER_LAYER(deconv3d, DeConv3DLayer);
+
+#define DECONV_OUTPUT_SIZE(IN_SIZE, STRID, PAD, KSIZE) \
+  (((IN_SIZE)-1) * (STRID)-2 * (PAD) + (KSIZE))


Why not the use imageSize() in ./paddle/math/MathUtils.h?

Done. Sorry, I'm not familiar with ./paddle/math/MathUtils.h

qingqing01 · 2017-08-18T10:37:08Z

paddle/gserver/tests/test_LayerGrad.cpp

+}
+
+int deConvOutputSize(int inSize, int kSize, int pad, int stride) {
+  return (inSize - 1) * stride - 2 * pad + kSize;


Why not the use imageSize() in ./paddle/math/MathUtils.h?

It has been modified

qingqing01 · 2017-08-18T10:38:52Z

paddle/math/tests/test_matrixCompare.cpp

@@ -1203,4 +1203,120 @@ TEST(Matrix, warpCTC) {
  }
 }

+int outputSizeCol2Vol(


same as above comments. ./paddle/math/MathUtils.h has contains the calculation.

qingqing01 · 2017-08-18T10:40:47Z

paddle/math/tests/test_matrixCompare.cpp

+  for (auto depth : {9, 16, 64, 128}) {
+    for (auto height : {9, 11, 73, 128, 256}) {
+      for (auto width : {
+               9, 32, 100, 512,


For speeding testing, please use the smaller number for depth, height, width. But please test larger number off-line.

chengduoZH requested review from qingqing01, hedaoyuan, luotao1 and lcy-seso August 9, 2017 04:40

chengduoZH added the developing label Aug 9, 2017

chengduoZH mentioned this pull request Aug 9, 2017

Support 3-Dimensional Conv3DLayer, DeConv3DLayer #3359

Closed

chengduoZH changed the title ~~add Conv3DLayer, DeConv3DLayer, Test_Conv3D~~ Add Conv3DLayer, DeConv3DLayer Aug 10, 2017

chengduoZH force-pushed the develop branch 2 times, most recently from b9f7008 to defb286 Compare August 13, 2017 01:25

Add vol2col and col2vol cuda kernel

cfb86c4

chengduoZH force-pushed the develop branch from defb286 to cfb86c4 Compare August 13, 2017 01:30

chengduoZH added 6 commits August 13, 2017 09:38

Modify ConvConfig, Add depth dimension

8cc0eb9

Add depth dimension information to ConvBaseLayer

5d7f6dd

Add Conv3DLayer

11975b4

Add DeConv3DLayer

23cf0c6

Add col2vol and vol2col CPU funtion

52ceeed

Add depth dimension information to Argument

9b3d6ac

hedaoyuan requested changes Aug 14, 2017

View reviewed changes

chengduoZH added 5 commits August 18, 2017 14:10

add unit test DeConv3D, Conv3D, col2vol, vol2col

424b325

fix DeConv3D, Conv3D

c792ef7

fix Matrix

43f6cdc

fix col2vol vol2col kernel

0a7516d

modified bias shape of ConvLayer

38cc5da

qingqing01 requested changes Aug 18, 2017

View reviewed changes

chengduoZH force-pushed the develop branch 4 times, most recently from f629486 to bbbedf9 Compare August 21, 2017 06:56

fix above comments

d5768eb

chengduoZH force-pushed the develop branch from bbbedf9 to d5768eb Compare August 21, 2017 08:07

chengduoZH force-pushed the develop branch 3 times, most recently from fd3da80 to b608d39 Compare August 28, 2017 03:35

fix above comments

2710584

chengduoZH force-pushed the develop branch from b608d39 to 2710584 Compare August 28, 2017 07:58

fix previous comments(c++)

6053f7e

chengduoZH force-pushed the develop branch 8 times, most recently from c5246f1 to fe63eb1 Compare August 29, 2017 14:45

fix conflict

4dcc533

chengduoZH force-pushed the develop branch 2 times, most recently from 56d669f to 4dcc533 Compare August 30, 2017 01:13

Update networks.py

34f4f76

hedaoyuan previously approved these changes Aug 30, 2017

View reviewed changes

qingqing01 previously approved these changes Aug 30, 2017

View reviewed changes

fix data_layer for 3D data

2ae37a4

chengduoZH dismissed stale reviews from qingqing01 and hedaoyuan via 2ae37a4 August 30, 2017 16:28

hedaoyuan previously approved these changes Aug 31, 2017

View reviewed changes

fix layers_test.py

2e97045

chengduoZH dismissed hedaoyuan’s stale review via 2e97045 August 31, 2017 06:37

chengduoZH added 2 commits August 31, 2017 16:33

Add test_conv3d_layer.protostr,test_deconv3d_layer.protostr

a4e1e12

fix conflict

455d2bd

qingqing01 approved these changes Aug 31, 2017

View reviewed changes

chengduoZH merged commit 1e6c992 into PaddlePaddle:develop Aug 31, 2017

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Add Conv3DLayer, DeConv3DLayer #3348

Add Conv3DLayer, DeConv3DLayer #3348

chengduoZH commented Aug 9, 2017 •

edited

Loading

hedaoyuan left a comment

hedaoyuan Aug 14, 2017

chengduoZH Aug 19, 2017

hedaoyuan Aug 14, 2017

chengduoZH Aug 19, 2017

hedaoyuan Aug 14, 2017

chengduoZH Aug 19, 2017

hedaoyuan Aug 14, 2017

chengduoZH Aug 18, 2017 •

edited

Loading

hedaoyuan Aug 14, 2017

chengduoZH Aug 19, 2017

qingqing01 left a comment

qingqing01 Aug 9, 2017

chengduoZH Aug 19, 2017

qingqing01 Aug 18, 2017

chengduoZH Aug 19, 2017

qingqing01 Aug 18, 2017

chengduoZH Aug 19, 2017

qingqing01 Aug 18, 2017

chengduoZH Aug 19, 2017

qingqing01 Aug 18, 2017

chengduoZH Aug 18, 2017

qingqing01 Aug 18, 2017

qingqing01 Aug 22, 2017

qingqing01 Aug 18, 2017

chengduoZH Aug 19, 2017

qingqing01 Aug 18, 2017

chengduoZH Aug 19, 2017

qingqing01 Aug 18, 2017

chengduoZH Aug 19, 2017

qingqing01 Aug 18, 2017

chengduoZH Aug 19, 2017



		__global__ void keMatrixVol2Col(
		int num_kernels, realdata, real data_col,

Add Conv3DLayer, DeConv3DLayer #3348

Add Conv3DLayer, DeConv3DLayer #3348

Conversation

chengduoZH commented Aug 9, 2017 • edited Loading

hedaoyuan left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

chengduoZH Aug 18, 2017 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

qingqing01 left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

chengduoZH commented Aug 9, 2017 •

edited

Loading

chengduoZH Aug 18, 2017 •

edited

Loading