You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@singa.apache.org by wa...@apache.org on 2018/07/05 03:10:01 UTC
[06/18] incubator-singa git commit: SINGA-371 Implement functional
operations in c++ for autograd
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 <xu...@u.nus.edu>
Authored: Tue Jun 19 11:01:45 2018 +0000
Committer: xuewanqi <xu...@u.nus.edu>
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);