SINGA-371 Implement functional operations in c++ for autograd - fix some bugs - rewrite the function SetupConv(), the reason is that LayerConf cannot be created in python,thus unable to send this arguement to InitConv() - these functions have been tested and are workable
Project: http://git-wip-us.apache.org/repos/asf/incubator-singa/repo Commit: http://git-wip-us.apache.org/repos/asf/incubator-singa/commit/2cac057d Tree: http://git-wip-us.apache.org/repos/asf/incubator-singa/tree/2cac057d Diff: http://git-wip-us.apache.org/repos/asf/incubator-singa/diff/2cac057d Branch: refs/heads/master Commit: 2cac057dffa77bbdeffdaa8aabdf37a03e82dd00 Parents: d48dea0 Author: xuewanqi <[email protected]> Authored: Tue Jun 19 11:01:45 2018 +0000 Committer: xuewanqi <[email protected]> Committed: Wed Jun 20 14:47:37 2018 +0000 ---------------------------------------------------------------------- src/api/model_operation.i | 11 ++-- src/model/convolution_functions.cc | 96 +++++---------------------------- src/model/convolution_functions.h | 10 +++- 3 files changed, 30 insertions(+), 87 deletions(-) ---------------------------------------------------------------------- http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/2cac057d/src/api/model_operation.i ---------------------------------------------------------------------- diff --git a/src/api/model_operation.i b/src/api/model_operation.i index 79707eb..77ef6bb 100644 --- a/src/api/model_operation.i +++ b/src/api/model_operation.i @@ -2,8 +2,6 @@ %{ #include "../src/model/convolution_functions.h" -using singa::Tensor; -using singa::CudnnConvHandle; %} namespace singa{ @@ -11,7 +9,13 @@ struct ConvHandle{}; struct CudnnConvHandle{}; -ConvHandle SetupConv(const size_t in_channels, const LayerConf &conf); +ConvHandle SetupConv( + const size_t kernel_h_, const size_t kernel_w_, + const size_t pad_h_, const size_t pad_w_, + const size_t stride_h_,const size_t stride_w_, + const size_t channels_, const size_t num_filters_, + const bool bias_term_ = true, const size_t workspace_byte_limit_ =1024*1024, + const std::string prefer_="fastest"); CudnnConvHandle InitCudnn(const Tensor &input, const ConvHandle ch); @@ -23,4 +27,5 @@ Tensor CudnnConvBackwardW(const Tensor &dy, const Tensor &x, const Tensor &W, co Tensor CudnnConvBackwardb(const Tensor &dy, const Tensor &b, const CudnnConvHandle cch); Tensor CudnnConvBackwardx(const Tensor &dy, const Tensor &W, const Tensor &x, const CudnnConvHandle cch); + } http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/2cac057d/src/model/convolution_functions.cc ---------------------------------------------------------------------- diff --git a/src/model/convolution_functions.cc b/src/model/convolution_functions.cc index 0fc8e65..7ff399a 100644 --- a/src/model/convolution_functions.cc +++ b/src/model/convolution_functions.cc @@ -4,87 +4,18 @@ //#include "./layer/cudnn_utils.h" //#include "singa/utils/logging.h" #include "./convolution_functions.h" - +#include<iostream> namespace singa{ // Done in conv2d.__init__() -ConvHandle SetupConv(const size_t in_channels, const LayerConf &conf){ - - size_t kernel_w_, pad_w_, stride_w_; - size_t kernel_h_, pad_h_, stride_h_; - - size_t channels_, num_filters_; - - bool bias_term_; - - size_t workspace_byte_limit_; - string prefer_; - - ConvolutionConf conv_conf = conf.convolution_conf(); - - workspace_byte_limit_ = conv_conf.workspace_byte_limit() << 20; - prefer_ = ToLowerCase(conv_conf.prefer()); - CHECK(prefer_ == "fastest" || prefer_ == "limited_workspace" || - prefer_ == "no_workspace" || prefer_ == "autotune") - << "CudnnConvolution only supports four algorithm preferences: fastest, " - "limited_workspace, no_workspace and autotune"; - - - // kernel_size, pad, and stride are repeated fields. - if (conv_conf.kernel_size_size() > 0) { - if (conv_conf.kernel_size_size() == 1) { - kernel_w_ = kernel_h_ = conv_conf.kernel_size(0); - } else { - kernel_w_ = conv_conf.kernel_size(0); - kernel_h_ = conv_conf.kernel_size(1); - } - } else { - kernel_w_ = conv_conf.kernel_w(); - kernel_h_ = conv_conf.kernel_h(); - } - CHECK_GT(kernel_w_, 0u); - CHECK_GT(kernel_h_, 0u); - - if (conv_conf.pad_size() > 0) { - if (conv_conf.pad_size() == 1) { - pad_w_ = pad_h_ = conv_conf.pad(0); - } else { - pad_w_ = conv_conf.pad(0); - pad_h_ = conv_conf.pad(1); - } - } else { - pad_w_ = conv_conf.pad_w(); - pad_h_ = conv_conf.pad_h(); - } - CHECK_GE(pad_w_, 0u); - CHECK_GE(pad_h_, 0u); - - const int kStrideDefault = 1; - if (conv_conf.stride_size() > 0) { - if (conv_conf.stride_size() == 1) { - stride_w_ = stride_h_ = conv_conf.stride(0); - } else { - stride_w_ = conv_conf.stride(0); - stride_h_ = conv_conf.stride(1); - } - } else { - stride_w_ = kStrideDefault; - stride_h_ = kStrideDefault; - if (conv_conf.has_stride_w()) { - stride_w_ = conv_conf.stride_w(); - } - if (conv_conf.has_stride_h()) { - stride_h_ = conv_conf.stride_h(); - } - } - CHECK_GT(stride_w_, 0u); - CHECK_GE(stride_h_, 0u); // 0 for 1D conv - - channels_ = in_channels; - num_filters_ = conv_conf.num_output(); - bias_term_ = conv_conf.bias_term(); - - return ConvHandle{ +ConvHandle SetupConv( + const size_t kernel_h_, const size_t kernel_w_, + const size_t pad_h_, const size_t pad_w_, + const size_t stride_h_,const size_t stride_w_, + const size_t channels_, const size_t num_filters_, + const bool bias_term_ , const size_t workspace_byte_limit_, + const std::string prefer_){ + return ConvHandle{ kernel_w_, pad_w_, stride_w_, @@ -103,7 +34,6 @@ ConvHandle SetupConv(const size_t in_channels, const LayerConf &conf){ }; - // Done in conv2d.__call__(): // if self.cudnnconvhandle is None: // self.cudnnconvhandle= InitCudnn(...) @@ -126,11 +56,11 @@ CudnnConvHandle InitCudnn(const Tensor &input, const ConvHandle ch){ size_t width_; size_t conv_height_; size_t conv_width_; - + DataType dtype = input.data_type(); auto dev = input.device(); Context *ctx = dev->context(0); - + size_t batchsize, channels_; batchsize = input.shape(0); channels_ = input.shape(1); @@ -143,7 +73,7 @@ CudnnConvHandle InitCudnn(const Tensor &input, const ConvHandle ch){ if (ch.stride_h_ > 0) conv_height_ = (height_ + 2 * ch.pad_h_ - ch.kernel_h_) / ch.stride_h_ + 1; conv_width_ = (width_ + 2 * ch.pad_w_ - ch.kernel_w_) / ch.stride_w_ + 1; - + CUDNN_CHECK(cudnnCreateTensorDescriptor(&x_desc_)); CUDNN_CHECK(cudnnCreateTensorDescriptor(&y_desc_)); if (ch.bias_term_) @@ -197,7 +127,7 @@ CudnnConvHandle InitCudnn(const Tensor &input, const ConvHandle ch){ CUDNN_CHECK(cudnnGetConvolutionBackwardDataAlgorithm( ctx->cudnn_handle, filter_desc_, y_desc_, conv_desc_, x_desc_, bwd_data_pref, ch.workspace_byte_limit_, &bp_data_alg_)); - } else if (ch.prefer_ == "autotune") { + } else if (ch.prefer_ == "autotune") { const int topk = 1; int num_fp_alg, num_bp_filt_alg, num_bp_data_alg; cudnnConvolutionFwdAlgoPerf_t fp_alg_perf[topk]; http://git-wip-us.apache.org/repos/asf/incubator-singa/blob/2cac057d/src/model/convolution_functions.h ---------------------------------------------------------------------- diff --git a/src/model/convolution_functions.h b/src/model/convolution_functions.h index eba0e50..9462805 100644 --- a/src/model/convolution_functions.h +++ b/src/model/convolution_functions.h @@ -43,7 +43,15 @@ struct CudnnConvHandle{ size_t batchsize; }; -ConvHandle SetupConv(const size_t in_channels, const LayerConf &conf); +ConvHandle SetupConv( + const size_t kernel_h_, const size_t kernel_w_, + const size_t pad_h_, const size_t pad_w_, + const size_t stride_h_,const size_t stride_w_, + const size_t channels_, const size_t num_filters_, + const bool bias_term_ = true ,const size_t workspace_byte_limit_=1024*1024, + const std::string prefer_="fastest"); + +void testInitCudnn(const Tensor &input, const ConvHandle ch); CudnnConvHandle InitCudnn(const Tensor &input, const ConvHandle ch);
