Skip to content

Commit

Permalink
update Convention
Browse files Browse the repository at this point in the history
  • Loading branch information
@tornadomeet
tornadomeet committed Oct 30, 2015
1 parent 1ceeb10 commit cee2587
Show file tree
Hide file tree
Showing 2 changed files with 45 additions and 43 deletions.
32 changes: 16 additions & 16 deletions src/operator/cudnn_deconvolution-inl.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -47,16 +47,16 @@ class CuDNNDeconvolutionOp : public Operator {
CHECK_EQ(in_data.size(), expected);
CHECK_EQ(out_data.size(), 1);
Stream<gpu> *s = ctx.get_stream<gpu>();
Tensor<gpu, 4> data = in_data[kData].get<gpu, 4, real_t>(s);
Tensor<gpu, 4> wmat = in_data[kWeight].get<gpu, 4, real_t>(s);
Tensor<gpu, 4> out = out_data[kOut].get<gpu, 4, real_t>(s);
Tensor<gpu, 4> data = in_data[deconv::kData].get<gpu, 4, real_t>(s);
Tensor<gpu, 4> wmat = in_data[deconv::kWeight].get<gpu, 4, real_t>(s);
Tensor<gpu, 4> out = out_data[deconv::kOut].get<gpu, 4, real_t>(s);
CHECK_EQ(data.CheckContiguous(), true);
CHECK_EQ(wmat.CheckContiguous(), true);
CHECK_EQ(out.CheckContiguous(), true);
if (!init_cudnn_) {
Init(s, in_data, out_data);
}
Tensor<gpu, 1> workspace = ctx.requested[kTempSpace].get_space<gpu>(
Tensor<gpu, 1> workspace = ctx.requested[deconv::kTempSpace].get_space<gpu>(
mshadow::Shape1(forward_workspace_), s);
CHECK_EQ(cudnnConvolutionBackwardData_v3(s->dnn_handle_,
&alpha,
Expand All @@ -73,7 +73,7 @@ class CuDNNDeconvolutionOp : public Operator {
out.dptr_), CUDNN_STATUS_SUCCESS);
if (!param_.no_bias) {
beta = 1.0f;
Tensor<gpu, 1> bias = in_data[kBias].get<gpu, 1, real_t>(s);
Tensor<gpu, 1> bias = in_data[deconv::kBias].get<gpu, 1, real_t>(s);
CHECK_EQ(cudnnAddTensor(s->dnn_handle_,
CUDNN_ADD_SAME_C,
&alpha,
Expand All @@ -100,17 +100,17 @@ class CuDNNDeconvolutionOp : public Operator {
CHECK_EQ(out_grad.size(), 1);
CHECK(in_data.size() == expected && in_grad.size() == expected);
// TODO(bing): think about how to support add to
CHECK_EQ(req[kWeight], kWriteTo);
CHECK_EQ(req[deconv::kWeight], kWriteTo);
Stream<gpu> *s = ctx.get_stream<gpu>();
Tensor<gpu, 4> grad = out_grad[kOut].get<gpu, 4, real_t>(s);
Tensor<gpu, 4> wmat = in_data[kWeight].get<gpu, 4, real_t>(s);
Tensor<gpu, 4> gwmat = in_grad[kWeight].get<gpu, 4, real_t>(s);
Tensor<gpu, 4> data = in_data[kData].get<gpu, 4, real_t>(s);
Tensor<gpu, 4> gdata = in_grad[kData].get<gpu, 4, real_t>(s);
Tensor<gpu, 1> workspace = ctx.requested[kTempSpace].get_space<gpu>(
Tensor<gpu, 4> grad = out_grad[deconv::kOut].get<gpu, 4, real_t>(s);
Tensor<gpu, 4> wmat = in_data[deconv::kWeight].get<gpu, 4, real_t>(s);
Tensor<gpu, 4> gwmat = in_grad[deconv::kWeight].get<gpu, 4, real_t>(s);
Tensor<gpu, 4> data = in_data[deconv::kData].get<gpu, 4, real_t>(s);
Tensor<gpu, 4> gdata = in_grad[deconv::kData].get<gpu, 4, real_t>(s);
Tensor<gpu, 1> workspace = ctx.requested[deconv::kTempSpace].get_space<gpu>(
mshadow::Shape1(backward_workspace_), s);
if (!param_.no_bias) {
Tensor<gpu, 1> gbias = in_grad[kBias].get<gpu, 1, real_t>(s);
Tensor<gpu, 1> gbias = in_grad[deconv::kBias].get<gpu, 1, real_t>(s);
CHECK_EQ(cudnnConvolutionBackwardBias(s->dnn_handle_,
&alpha,
out_desc_,
Expand Down Expand Up @@ -160,8 +160,8 @@ class CuDNNDeconvolutionOp : public Operator {
size_t workspace_byte = static_cast<size_t>(param_.workspace * sizeof(real_t));
size_t back_size = 0;
size_t back_size_w = 0;
Tensor<gpu, 4> data = in_data[kData].get<gpu, 4, real_t>(s);
Tensor<gpu, 4> out = out_data[kOut].get<gpu, 4, real_t>(s);
Tensor<gpu, 4> data = in_data[deconv::kData].get<gpu, 4, real_t>(s);
Tensor<gpu, 4> out = out_data[deconv::kOut].get<gpu, 4, real_t>(s);
CHECK_EQ(cudnnCreateTensorDescriptor(&in_desc_), CUDNN_STATUS_SUCCESS);
CHECK_EQ(cudnnCreateTensorDescriptor(&out_desc_), CUDNN_STATUS_SUCCESS);
CHECK_EQ(cudnnCreateTensorDescriptor(&bias_desc_), CUDNN_STATUS_SUCCESS);
Expand Down Expand Up @@ -196,7 +196,7 @@ class CuDNNDeconvolutionOp : public Operator {
out.shape_[2],
out.shape_[3]), CUDNN_STATUS_SUCCESS);
if (!param_.no_bias) {
Tensor<gpu, 1> bias = in_data[kBias].get<gpu, 1, real_t>(s);
Tensor<gpu, 1> bias = in_data[deconv::kBias].get<gpu, 1, real_t>(s);
CHECK_EQ(cudnnSetTensor4dDescriptor(bias_desc_,
CUDNN_TENSOR_NCHW,
dtype_,
Expand Down
56 changes: 29 additions & 27 deletions src/operator/deconvolution-inl.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -21,9 +21,11 @@
namespace mxnet {
namespace op {

enum DeconvolutionOpInputs {kData, kWeight, kBias};
enum DeconvolutionOpOutputs {kOut};
enum DeconvolutionOpResource {kTempSpace};
namespace deconv {
enum DeconvolutionOpInputs {kData, kWeight, kBias};
enum DeconvolutionOpOutputs {kOut};
enum DeconvolutionOpResource {kTempSpace};
}

struct DeconvolutionParam : public dmlc::Parameter<DeconvolutionParam> {
TShape kernel;
Expand Down Expand Up @@ -68,24 +70,24 @@ class DeconvolutionOp : public Operator {
const std::vector<TBlob> &aux_args) {
using namespace mshadow;
using namespace mshadow::expr;
CHECK_EQ(req[kOut], kWriteTo);
CHECK_EQ(req[deconv::kOut], kWriteTo);
size_t expected = param_.no_bias ? 2 : 3;
CHECK_EQ(in_data.size(), expected);
CHECK_EQ(out_data.size(), 1);
Stream<xpu> *s = ctx.get_stream<xpu>();
Tensor<xpu, 4> data = in_data[kData].get<xpu, 4, real_t>(s);
Tensor<xpu, 4> out = out_data[kOut].get<xpu, 4, real_t>(s);
Tensor<xpu, 4> data = in_data[deconv::kData].get<xpu, 4, real_t>(s);
Tensor<xpu, 4> out = out_data[deconv::kOut].get<xpu, 4, real_t>(s);
Shape<3> wmat_shape =
Shape3(param_.num_group,
data.shape_[1] / param_.num_group,
param_.num_filter / param_.num_group * param_.kernel[0] * param_.kernel[1]);
Tensor<xpu, 3> wmat = in_data[kWeight].get_with_shape<xpu, 3, real_t>(wmat_shape, s);
Tensor<xpu, 3> wmat = in_data[deconv::kWeight].get_with_shape<xpu, 3, real_t>(wmat_shape, s);
#if defined(__CUDACC__)
CHECK_EQ(s->blas_handle_ownership_, Stream<xpu>::OwnHandle)
<< "Must init CuBLAS handle in stream";
#endif
const index_t nbatch = data.size(0);
Tensor<xpu, 1> workspace = ctx.requested[kTempSpace].get_space<xpu>(
Tensor<xpu, 1> workspace = ctx.requested[deconv::kTempSpace].get_space<xpu>(
Shape1(this->InitTemp(out.shape_, data.shape_)), s);
for (index_t i = 0; i < nbatch; i += nstep_) {
const index_t step = std::min(nstep_, nbatch - i);
Expand Down Expand Up @@ -137,7 +139,7 @@ class DeconvolutionOp : public Operator {
}
if (!param_.no_bias) {
// add bias, broadcast bias to dim 1: channel
Tensor<xpu, 1> bias = in_data[kBias].get<xpu, 1, real_t>(s);
Tensor<xpu, 1> bias = in_data[deconv::kBias].get<xpu, 1, real_t>(s);
out += broadcast<1>(bias, out.shape_);
}
}
Expand All @@ -156,24 +158,24 @@ class DeconvolutionOp : public Operator {
size_t expected = param_.no_bias == 0 ? 3 : 2;
CHECK(in_data.size() == expected && in_grad.size() == expected);
CHECK_EQ(req.size(), expected);
CHECK_EQ(in_data[kWeight].CheckContiguous(), true);
CHECK_EQ(in_data[deconv::kWeight].CheckContiguous(), true);
// get data
Stream<xpu> *s = ctx.get_stream<xpu>();
Tensor<xpu, 4> data = in_data[kData].get<xpu, 4, real_t>(s);
Tensor<xpu, 4> grad = out_grad[kOut].get<xpu, 4, real_t>(s);
Tensor<xpu, 4> gdata = in_grad[kData].get<xpu, 4, real_t>(s);
Tensor<xpu, 4> data = in_data[deconv::kData].get<xpu, 4, real_t>(s);
Tensor<xpu, 4> grad = out_grad[deconv::kOut].get<xpu, 4, real_t>(s);
Tensor<xpu, 4> gdata = in_grad[deconv::kData].get<xpu, 4, real_t>(s);
Shape<3> wmat_shape =
Shape3(param_.num_group,
data.shape_[1] / param_.num_group,
param_.num_filter / param_.num_group * param_.kernel[0] * param_.kernel[1]);
Tensor<xpu, 3> wmat = in_data[kWeight].get_with_shape<xpu, 3, real_t>(wmat_shape, s);
Tensor<xpu, 3> gwmat = in_grad[kWeight].get_with_shape<xpu, 3, real_t>(wmat_shape, s);
Tensor<xpu, 3> wmat = in_data[deconv::kWeight].get_with_shape<xpu, 3, real_t>(wmat_shape, s);
Tensor<xpu, 3> gwmat = in_grad[deconv::kWeight].get_with_shape<xpu, 3, real_t>(wmat_shape, s);
#if defined(__CUDACC__)
CHECK_EQ(s->blas_handle_ownership_, Stream<xpu>::OwnHandle)
<< "Must init CuBLAS handle in stream";
#endif
const index_t nbatch = data.size(0);
Tensor<xpu, 1> workspace = ctx.requested[kTempSpace].get_space<xpu>(
Tensor<xpu, 1> workspace = ctx.requested[deconv::kTempSpace].get_space<xpu>(
Shape1(this->InitTemp(grad.shape_, data.shape_)), s);
for (index_t i = 0; i < nbatch; i += nstep_) {
const index_t step = std::min(nstep_, nbatch - i);
Expand Down Expand Up @@ -203,12 +205,12 @@ class DeconvolutionOp : public Operator {
Tensor<xpu, 2> tmpc = temp_col.Slice(gstride * gid, gstride * (gid + 1));
if (i == 0) {
Tensor<xpu, 2> tmp_gwmat = gwmat[gid];
Assign(tmp_gwmat, req[kWeight], dot(temp_dst[gid], tmpc.T()));
Assign(tmp_gwmat, req[deconv::kWeight], dot(temp_dst[gid], tmpc.T()));
} else {
gwmat[gid] += dot(temp_dst[gid], tmpc.T());
}
}
if (req[kData] == kWriteTo || req[kData] == kWriteInplace) {
if (req[deconv::kData] == kWriteTo || req[deconv::kData] == kWriteInplace) {
for (uint32_t gid = 0; gid < param_.num_group; ++gid) {
Tensor<xpu, 2> tmpc = temp_col.Slice(gstride * gid, gstride * (gid + 1));
temp_dst[gid] = dot(wmat[gid], tmpc);
Expand All @@ -221,8 +223,8 @@ class DeconvolutionOp : public Operator {
}
}
if (!param_.no_bias) {
Tensor<xpu, 1> gbias = in_grad[kBias].get<xpu, 1, real_t>(s);
Assign(gbias, req[kBias], sumall_except_dim<1>(grad));
Tensor<xpu, 1> gbias = in_grad[deconv::kBias].get<xpu, 1, real_t>(s);
Assign(gbias, req[deconv::kBias], sumall_except_dim<1>(grad));
}
}

Expand Down Expand Up @@ -289,15 +291,15 @@ class DeconvolutionProp : public OperatorProperty {
} else {
CHECK_EQ(in_shape->size(), 2) << "Input:[data, weight]";
}
const TShape &dshape = (*in_shape)[kData];
const TShape &dshape = (*in_shape)[deconv::kData];
if (dshape.ndim() == 0) return false;
CHECK_EQ(dshape.ndim(), 4) \
<< "Input data should be 4D in batch-num_filter-y-x";
SHAPE_ASSIGN_CHECK(*in_shape,
kWeight,
deconv::kWeight,
Shape4(dshape[1], param_.num_filter, param_.kernel[0], param_.kernel[1]));
if (!param_.no_bias) {
SHAPE_ASSIGN_CHECK(*in_shape, kBias, Shape1(param_.num_filter));
SHAPE_ASSIGN_CHECK(*in_shape, deconv::kBias, Shape1(param_.num_filter));
}
out_shape->clear();
out_shape->push_back(dshape);
Expand All @@ -313,9 +315,9 @@ class DeconvolutionProp : public OperatorProperty {
<< "incorrect stride size: " << param_.stride;
CHECK(ksize_x <= dshape[3] && ksize_y <= dshape[2])
<< "kernel size exceed input";
(*out_shape)[kOut][1] = param_.num_filter;
(*out_shape)[kOut][2] = param_.stride[0] * (dshape[2] - 1) + ksize_y - 2 * param_.pad[0];
(*out_shape)[kOut][3] = param_.stride[1] * (dshape[3] - 1) + ksize_x - 2 * param_.pad[1];
(*out_shape)[deconv::kOut][1] = param_.num_filter;
(*out_shape)[deconv::kOut][2] = param_.stride[0] * (dshape[2] - 1) + ksize_y - 2 * param_.pad[0];
(*out_shape)[deconv::kOut][3] = param_.stride[1] * (dshape[3] - 1) + ksize_x - 2 * param_.pad[1];
return true;
}

Expand All @@ -333,7 +335,7 @@ class DeconvolutionProp : public OperatorProperty {
const std::vector<int> &out_grad,
const std::vector<int> &in_data,
const std::vector<int> &out_data) const override {
return {out_grad[kOut], in_data[kData], in_data[kWeight]};
return {out_grad[deconv::kOut], in_data[deconv::kData], in_data[deconv::kWeight]};
}

std::vector<ResourceRequest> ForwardResource(
Expand Down

0 comments on commit cee2587

Please sign in to comment.