SINGA-371 Implement functional operations in c++ for autograd - fixed some bugs.
Project: http://git-wip-us.apache.org/repos/asf/incubator-singa/repo Commit: http://git-wip-us.apache.org/repos/asf/incubator-singa/commit/4a45ee6f Tree: http://git-wip-us.apache.org/repos/asf/incubator-singa/tree/4a45ee6f Diff: http://git-wip-us.apache.org/repos/asf/incubator-singa/diff/4a45ee6f Branch: refs/heads/master Commit: 4a45ee6f080bb9eacdeb1294047a78a3dbd4635a Parents: 5340b65 Author: xuewanqi <[email protected]> Authored: Wed Jul 4 06:46:09 2018 +0000 Committer: xuewanqi <[email protected]> Committed: Wed Jul 4 07:31:58 2018 +0000 ---------------------------------------------------------------------- python/singa/autograd.py | 28 ++++++++++++++++------------ src/api/model_operation.i | 18 +++++++++--------- src/model/layer/convolution.cc | 4 ++-- src/model/layer/convolution.h | 21 +++++++++++---------- src/model/operation/convolution.cc | 28 ++++++++++++++-------------- src/model/operation/convolution.h | 18 +++++++++--------- test/python/test_operation.py | 4 ++-- 7 files changed, 63 insertions(+), 58 deletions(-) ---------------------------------------------------------------------- http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/4a45ee6f/python/singa/autograd.py ---------------------------------------------------------------------- diff --git a/python/singa/autograd.py b/python/singa/autograd.py index 2a10608..b05f701 100755 --- a/python/singa/autograd.py +++ b/python/singa/autograd.py @@ -463,7 +463,10 @@ class _Conv2D(Operation): #assert 0 == 0, 'invalid padding' if training: - self.inputs = (x, W, b) + if self.handle.bias_term_: + self.inputs = (x, W, b) + else: + self.inputs = (x, W) if self.handle.device_id == -1: return singa.CpuConvForward(x, W, b, self.handle) @@ -717,32 +720,33 @@ class Conv2D(NewLayer): else: # to keep consistency when to do forward. self.b = Tensor(data=CTensor( - [1]), requires_grad=False, stores_grad=False) - self.b.set_value(0.0) + []), requires_grad=False, stores_grad=False) def __call__(self, x): - assert x.shape[1] == self.in_channels,'in_channels dismatched' - assert (x.shape[2]+2*self.padding[0]-self.kernel_size[0])%self.stride[0] == 0, 'invalid padding or strides.' - assert (x.shape[3]+2*self.padding[1]-self.kernel_size[1])%self.stride[1] == 0, 'invalid padding or stride.' + assert x.shape[1] == self.in_channels, 'in_channels dismatched' + assert (x.shape[2] + 2 * self.padding[0] - self.kernel_size[0] + ) % self.stride[0] == 0, 'invalid padding or strides.' + assert (x.shape[3] + 2 * self.padding[1] - self.kernel_size[1] + ) % self.stride[1] == 0, 'invalid padding or stride.' self.device_check(x, self.W, self.b) if x.device.id() == -1: if not hasattr(self, 'handle'): self.handle = singa.ConvHandle(x.data, self.kernel_size, self.stride, - self.padding, self.in_channels, self.out_channels, self.bias) + self.padding, self.in_channels, self.out_channels, self.bias) elif x.shape[0] != self.handle.batchsize: self.handle = singa.ConvHandle(x.data, self.kernel_size, self.stride, - self.padding, self.in_channels, self.out_channels, self.bias) + self.padding, self.in_channels, self.out_channels, self.bias) else: if not hasattr(self, 'handle'): self.handle = singa.CudnnConvHandle(x.data, self.kernel_size, self.stride, - self.padding, self.in_channels, self.out_channels, self.bias, - self.inner_params['workspace_MB_limit'] * 1024 * 1024, self.inner_params['cudnn_prefer']) + self.padding, self.in_channels, self.out_channels, self.bias, + self.inner_params['workspace_MB_limit'] * 1024 * 1024, self.inner_params['cudnn_prefer']) elif x.shape[0] != self.handle.batchsize: self.handle = singa.CudnnConvHandle(x.data, self.kernel_size, self.stride, - self.padding, self.in_channels, self.out_channels, self.bias, - self.inner_params['workspace_MB_limit'] * 1024 * 1024, self.inner_params['cudnn_prefer']) + self.padding, self.in_channels, self.out_channels, self.bias, + self.inner_params['workspace_MB_limit'] * 1024 * 1024, self.inner_params['cudnn_prefer']) self.handle.device_id = x.device.id() y = conv2d(x, self.W, self.b, self.handle) http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/4a45ee6f/src/api/model_operation.i ---------------------------------------------------------------------- diff --git a/src/api/model_operation.i b/src/api/model_operation.i index 58e5270..26f5c69 100755 --- a/src/api/model_operation.i +++ b/src/api/model_operation.i @@ -10,10 +10,10 @@ struct ConvHandle{ size_t batchsize; const bool bias_term_; - ConvHandle(const Tensor &input, const std::vector<size_t> kernel_size, - const std::vector<size_t> stride, const std::vector<size_t> padding, - const size_t in_channels, const size_t out_channels, - const bool bias_term_); + ConvHandle(const Tensor &input, const std::vector<size_t>& kernel_size, + const std::vector<size_t>& stride, const std::vector<size_t>& padding, + const size_t in_channels, const size_t out_channels, + const bool bias); }; struct CudnnConvHandle{ @@ -21,11 +21,11 @@ struct CudnnConvHandle{ size_t batchsize; const bool bias_term_; - CudnnConvHandle(const Tensor &input, const std::vector<size_t> kernel_size, - const std::vector<size_t> stride, const std::vector<size_t> padding, - const size_t in_channels, const size_t out_channels, - const bool bias_term_, const size_t workspace_byte_limit_=1024*1024*1024, - const std::string prefer_="fastest"); + CudnnConvHandle(const Tensor &input, const std::vector<size_t>& kernel_size, + const std::vector<size_t>& stride, const std::vector<size_t>& padding, + const size_t in_channels, const size_t out_channels, + const bool bias, const size_t workspace_byte_limit_ = 1024 * 1024 * 1024, + const std::string& prefer_ = "fastest"); }; Tensor GpuConvForward(const Tensor &x, const Tensor &W, const Tensor &b, const CudnnConvHandle &cch); http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/4a45ee6f/src/model/layer/convolution.cc ---------------------------------------------------------------------- diff --git a/src/model/layer/convolution.cc b/src/model/layer/convolution.cc old mode 100644 new mode 100755 index 3fc7afb..cc77433 --- a/src/model/layer/convolution.cc +++ b/src/model/layer/convolution.cc @@ -194,7 +194,7 @@ void Convolution::ToDevice(std::shared_ptr<Device> device) { bias_.ToDevice(device); } -void Convolution::Im2col(const float *data_im, const int channels, +void Im2col(const float *data_im, const int channels, const int height, const int width, const int kernel_h, const int kernel_w, const int pad_h, const int pad_w, @@ -221,7 +221,7 @@ void Convolution::Im2col(const float *data_im, const int channels, } } -void Convolution::Col2im(const float *data_col, const int channels, +void Col2im(const float *data_col, const int channels, const int height, const int width, const int kernel_h, const int kernel_w, const int pad_h, const int pad_w, http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/4a45ee6f/src/model/layer/convolution.h ---------------------------------------------------------------------- diff --git a/src/model/layer/convolution.h b/src/model/layer/convolution.h old mode 100644 new mode 100755 index 89b5319..d11cdeb --- a/src/model/layer/convolution.h +++ b/src/model/layer/convolution.h @@ -46,16 +46,6 @@ class Convolution : public Layer { void ToDevice(std::shared_ptr<Device> device) override; - void Im2col(const float* data_im, const int channels, const int height, - const int width, const int kernel_h, const int kernel_w, - const int pad_h, const int pad_w, const int stride_h, - const int stride_w, float* data_col); - - void Col2im(const float* data_col, const int channels, const int height, - const int width, const int kernel_h, const int kernel_w, - const int pad_h, const int pad_w, const int stride_h, - const int stride_w, float* data_im); - const std::vector<Tensor> param_values() override { if (bias_term_) return std::vector<Tensor>{weight_, bias_}; @@ -97,5 +87,16 @@ class Convolution : public Layer { bool bias_term_; vector<size_t> out_sample_shape_; }; + +void Im2col(const float* data_im, const int channels, const int height, + const int width, const int kernel_h, const int kernel_w, + const int pad_h, const int pad_w, const int stride_h, + const int stride_w, float* data_col); + +void Col2im(const float* data_col, const int channels, const int height, + const int width, const int kernel_h, const int kernel_w, + const int pad_h, const int pad_w, const int stride_h, + const int stride_w, float* data_im); + } // namespace singa #endif // SRC_MODEL_LAYER_CONVOLUTION_H_ http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/4a45ee6f/src/model/operation/convolution.cc ---------------------------------------------------------------------- diff --git a/src/model/operation/convolution.cc b/src/model/operation/convolution.cc index d64fbc1..9a702fa 100755 --- a/src/model/operation/convolution.cc +++ b/src/model/operation/convolution.cc @@ -1,10 +1,10 @@ #include "./convolution.h" -// #include "../layer/convolution.h" -#include<iostream> +#include "../layer/convolution.h" + namespace singa { -ConvHandle::ConvHandle(const Tensor &input, const std::vector<size_t> kernel_size, +ConvHandle::ConvHandle(const Tensor &input, const std::vector<size_t>& kernel_size, const std::vector<size_t>& stride, const std::vector<size_t>& padding, const size_t in_channels, const size_t out_channels, const bool bias) { @@ -37,7 +37,7 @@ ConvHandle::ConvHandle(const Tensor &input, const std::vector<size_t> kernel_siz imagesize = input.Size() / batchsize; } -// Convolution C; + Tensor CpuConvForward(const Tensor &x, Tensor &W, Tensor &b, const ConvHandle &ch) { CHECK_EQ(x.device()->lang(), kCpp); @@ -67,7 +67,7 @@ Tensor CpuConvForward(const Tensor &x, Tensor &W, Tensor &b, const ConvHandle & float *data_col = new float[ch.col_height_ * ch.col_width_]; auto in_data = x.data<float>(); for (size_t num = 0; num < ch.batchsize; num++) { - C.Im2col(in_data + num * ch.imagesize, ch.channels_, ch.height_, ch.width_, ch.kernel_h_, + Im2col(in_data + num * ch.imagesize, ch.channels_, ch.height_, ch.width_, ch.kernel_h_, ch.kernel_w_, ch.pad_h_, ch.pad_w_, ch.stride_h_, ch.stride_w_, data_col); col_data.CopyDataFromHostPtr(data_col, ch.col_height_ * ch.col_width_); @@ -105,7 +105,7 @@ Tensor CpuConvBackwardx(const Tensor &dy, Tensor &W, const Tensor &x, const Conv CopyDataToFrom(&grad_b, dy, grad_b.Size(), 0, num * grad_b.Size()); Tensor dcol_b = Mult(W.T(), grad_b); auto dcol_data = dcol_b.data<float>(); - C.Col2im(dcol_data, ch.channels_, ch.height_, ch.width_, ch.kernel_h_, ch.kernel_w_, ch.pad_h_, + Col2im(dcol_data, ch.channels_, ch.height_, ch.width_, ch.kernel_h_, ch.kernel_w_, ch.pad_h_, ch.pad_w_, ch.stride_h_, ch.stride_w_, dx_b); dx.CopyDataFromHostPtr(dx_b, ch.imagesize, num * ch.imagesize); } @@ -134,7 +134,7 @@ Tensor CpuConvBackwardW(const Tensor &dy, const Tensor &x, const Tensor &W, cons float *data_col = new float[ch.col_height_ * ch.col_width_]; auto in_data = dy.data<float>(); for (size_t num = 0; num < ch.batchsize; num++) { - C.Im2col(in_data + num * ch.imagesize, ch.channels_, ch.height_, ch.width_, ch.kernel_h_, + Im2col(in_data + num * ch.imagesize, ch.channels_, ch.height_, ch.width_, ch.kernel_h_, ch.kernel_w_, ch.pad_h_, ch.pad_w_, ch.stride_h_, ch.stride_w_, data_col); col_data.CopyDataFromHostPtr(data_col, ch.col_height_ * ch.col_width_); Tensor grad_b(Shape{ch.num_filters_, ch.conv_height_ * ch.conv_width_}); @@ -171,9 +171,9 @@ Tensor CpuConvBackwardb(const Tensor &dy, const Tensor &b, const ConvHandle &ch) #ifdef USE_CUDNN CudnnConvHandle::CudnnConvHandle(const Tensor &input, const std::vector<size_t>& kernel_size, const std::vector<size_t>& stride, const std::vector<size_t>& padding, - const size_t in_channels, const size_t out_channels, const bool bias_term_, + const size_t in_channels, const size_t out_channels, const bool bias, const size_t workspace_byte_limit_, const std::string& prefer_) - : ConvHandle(input, kernel_size, stride, padding, in_channels, out_channels, bias_term_) { + : ConvHandle(input, kernel_size, stride, padding, in_channels, out_channels, bias) { DataType dtype = input.data_type(); auto dev = input.device(); @@ -295,7 +295,7 @@ Tensor GpuConvForward(const Tensor &x, const Tensor &W, const Tensor &b, const C Shape shape{cch.batchsize, cch.num_filters_, cch.conv_height_, cch.conv_width_}; Tensor output(shape, dev, dtype); - output.device()->Exec([output, x, W, cch](Context * ctx) { + output.device()->Exec([&output, &x, &W, &cch](Context * ctx) { Block *inblock = x.block(), *outblock = output.block(), *wblock = W.block(); float alpha = 1.f, beta = 0.f; @@ -308,7 +308,7 @@ Tensor GpuConvForward(const Tensor &x, const Tensor &W, const Tensor &b, const C }, {x.block(), W.block()}, {output.block()}, cch.workspace_.block()); if (cch.bias_term_) { - output.device()->Exec([output, b, cch](Context * ctx) { + output.device()->Exec([&output, &b, &cch](Context * ctx) { float beta = 1.f, alpha = 1.0f; Block *outblock = output.block(), *bblock = b.block(); cudnnAddTensor(ctx->cudnn_handle, &alpha, cch.bias_desc_, @@ -326,7 +326,7 @@ Tensor GpuConvBackwardx(const Tensor &dy, const Tensor &W, const Tensor &x, cons Tensor dx; dx.ResetLike(x); - dy.device()->Exec([dx, dy, W, cch](Context * ctx) { + dy.device()->Exec([&dx, &dy, &W, &cch](Context * ctx) { Block *wblock = W.block(), *dyblock = dy.block(), *dxblock = dx.block(); float alpha = 1.f, beta = 0.f; @@ -347,7 +347,7 @@ Tensor GpuConvBackwardW(const Tensor &dy, const Tensor &x, const Tensor &W, cons Tensor dW; dW.ResetLike(W); - dy.device()->Exec([dW, dy, x, W, cch](Context * ctx) { + dy.device()->Exec([&dW, &dy, &x, &cch](Context * ctx) { Block *inblock = x.block(), *dyblock = dy.block(), *dwblock = dW.block(); float alpha = 1.f, beta = 0.f; @@ -369,7 +369,7 @@ Tensor GpuConvBackwardb(const Tensor &dy, const Tensor &b, const CudnnConvHandle Tensor db; db.ResetLike(b); - dy.device()->Exec([db, dy, b, cch](Context * ctx) { + dy.device()->Exec([&db, &dy, &cch](Context * ctx) { Block *dyblock = dy.block(), *dbblock = db.block(); float alpha = 1.f, beta = 0.f; cudnnConvolutionBackwardBias(ctx->cudnn_handle, &alpha, cch.y_desc_, http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/4a45ee6f/src/model/operation/convolution.h ---------------------------------------------------------------------- diff --git a/src/model/operation/convolution.h b/src/model/operation/convolution.h index a114b47..93f7775 100755 --- a/src/model/operation/convolution.h +++ b/src/model/operation/convolution.h @@ -3,12 +3,12 @@ #include <string> #include <vector> +#include "singa/core/tensor.h" #include "singa/utils/logging.h" #ifdef USE_CUDNN #include <cudnn.h> -// #include "../layer/cudnn_convolution.h" -// #include "../layer/cudnn_utils.h" +#include "../layer/cudnn_utils.h" #endif // USE_CUDNN @@ -21,7 +21,7 @@ class ConvHandle { const std::vector<size_t>& stride, const std::vector<size_t>& padding, const size_t in_channels, const size_t out_channels, const bool bias); - protected: + size_t kernel_w_; size_t pad_w_; size_t stride_w_; @@ -66,12 +66,12 @@ class CudnnConvHandle: public ConvHandle { const std::string& prefer_ = "fastest"); ~CudnnConvHandle(); // TODO(wangwei) add the destructor - protected: - cudnnTensorDescriptor_t x_desc_ ; - cudnnTensorDescriptor_t y_desc_ ; - cudnnTensorDescriptor_t bias_desc_ ; - cudnnFilterDescriptor_t filter_desc_ ; - cudnnConvolutionDescriptor_t conv_desc_ ; + + cudnnTensorDescriptor_t x_desc_ = nullptr; + cudnnTensorDescriptor_t y_desc_ = nullptr; + cudnnTensorDescriptor_t bias_desc_ = nullptr; + cudnnFilterDescriptor_t filter_desc_ = nullptr; + cudnnConvolutionDescriptor_t conv_desc_ = nullptr; cudnnConvolutionFwdAlgo_t fp_alg_; cudnnConvolutionBwdFilterAlgo_t bp_filter_alg_; cudnnConvolutionBwdDataAlgo_t bp_data_alg_; http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/4a45ee6f/test/python/test_operation.py ---------------------------------------------------------------------- diff --git a/test/python/test_operation.py b/test/python/test_operation.py old mode 100644 new mode 100755 index 1bbc70c..315a992 --- a/test/python/test_operation.py +++ b/test/python/test_operation.py @@ -48,7 +48,7 @@ class TestPythonOperation(unittest.TestCase): # forward without bias y_without_bias = conv_without_bias_0(gpu_input_tensor) - self.check_shape(y.shape, (2, 1, 2, 2)) + self.check_shape(y_without_bias.shape, (2, 1, 2, 2)) def test_conv2d_cpu(self): # (in_channels, out_channels, kernel_size) @@ -68,7 +68,7 @@ class TestPythonOperation(unittest.TestCase): # forward without bias y_without_bias = conv_without_bias_1(cpu_input_tensor) - self.check_shape(y.shape, (2, 1, 2, 2)) + self.check_shape(y_without_bias.shape, (2, 1, 2, 2)) if __name__ == '__main__': unittest.main()
