From ad2bbd69d9334ff9dc98c3bb9f08b4e262f67f21 Mon Sep 17 00:00:00 2001 From: eric-haibin-lin Date: Tue, 8 May 2018 06:38:44 +0000 Subject: [PATCH 1/7] inplace version of activation(relu) --- src/operator/nn/activation-inl.h | 148 +++++++++---------- src/operator/nn/activation.cc | 47 ++++-- src/operator/nn/activation.cu | 36 +++-- src/operator/nn/cudnn/cudnn_activation-inl.h | 3 + tests/python/gpu/test_operator_gpu.py | 19 ++- tests/python/unittest/test_operator.py | 4 + 6 files changed, 141 insertions(+), 116 deletions(-) diff --git a/src/operator/nn/activation-inl.h b/src/operator/nn/activation-inl.h index 32a7a5ad6176..403ac48a18d4 100644 --- a/src/operator/nn/activation-inl.h +++ b/src/operator/nn/activation-inl.h @@ -83,7 +83,7 @@ struct hash { namespace mxnet { namespace op { -template +template void ActivationForward(const OpContext &ctx, const TBlob &in_data, const OpReqType &req, const TBlob &out_data) { using namespace mshadow; @@ -91,16 +91,16 @@ void ActivationForward(const OpContext &ctx, const TBlob &in_data, Stream *s = ctx.get_stream(); const size_t sz = in_data.shape_.Size(); if (sz) { - MXNET_ASSIGN_REQ_SWITCH(req, Req, { - mxnet_op::Kernel, xpu>::Launch( - s, sz, - out_data.dptr(), - in_data.dptr()); + MSHADOW_REAL_TYPE_SWITCH(in_data.type_flag_, DType, { + MXNET_ASSIGN_REQ_SWITCH(req, Req, { + mxnet_op::Kernel, xpu>::Launch( + s, sz, out_data.dptr(), in_data.dptr()); + }); }); } } -template +template void ActivationBackward(const OpContext &ctx, const TBlob &out_grad, const TBlob &out_data, const OpReqType &req, const TBlob &in_grad) { @@ -109,13 +109,12 @@ void ActivationBackward(const OpContext &ctx, const TBlob &out_grad, Stream *s = ctx.get_stream(); const size_t sz = out_data.shape_.Size(); if (sz) { - MXNET_ASSIGN_REQ_SWITCH(req, Req, { - mxnet_op::Kernel, Req>, xpu>::Launch( - s, sz, - in_grad.dptr(), - out_grad.dptr(), - out_data.dptr()); + MSHADOW_REAL_TYPE_SWITCH(out_grad.type_flag_, DType, { + MXNET_ASSIGN_REQ_SWITCH(req, Req, { + mxnet_op::Kernel, Req>, xpu>::Launch( + s, sz, in_grad.dptr(), out_grad.dptr(), out_data.dptr()); + }); }); } } @@ -123,72 +122,68 @@ void ActivationBackward(const OpContext &ctx, const TBlob &out_grad, template void ActivationComputeImpl(const ActivationParam ¶m, const OpContext &ctx, const TBlob &input, OpReqType req, const TBlob &output) { - MSHADOW_REAL_TYPE_SWITCH(input.type_flag_, DType, { - switch (param.act_type) { - case activation::kReLU: - ActivationForward( - ctx, input, req, output); - break; - case activation::kSigmoid: - ActivationForward( - ctx, input, req, output); - break; - case activation::kTanh: - ActivationForward( - ctx, input, req, output); - break; - case activation::kSoftReLU: - ActivationForward( - ctx, input, req, output); - break; - case activation::kSoftSign: - ActivationForward( - ctx, input, req, output); - break; - default: - LOG(FATAL) << "unknown activation type"; - } - }); + switch (param.act_type) { + case activation::kReLU: + ActivationForward( + ctx, input, req, output); + break; + case activation::kSigmoid: + ActivationForward( + ctx, input, req, output); + break; + case activation::kTanh: + ActivationForward( + ctx, input, req, output); + break; + case activation::kSoftReLU: + ActivationForward( + ctx, input, req, output); + break; + case activation::kSoftSign: + ActivationForward( + ctx, input, req, output); + break; + default: + LOG(FATAL) << "unknown activation type"; + } } template void ActivationGradComputeImpl(const ActivationParam ¶m, const OpContext &ctx, const TBlob &out_grad, const TBlob &out_data, OpReqType req, const TBlob &output) { - MSHADOW_REAL_TYPE_SWITCH(out_grad.type_flag_, DType, { - switch (param.act_type) { - case activation::kReLU: - ActivationBackward( - ctx, out_grad, out_data, req, output); - break; - case activation::kSigmoid: - ActivationBackward( - ctx, out_grad, out_data, req, output); - break; - case activation::kTanh: - ActivationBackward( - ctx, out_grad, out_data, req, output); - break; - case activation::kSoftReLU: - ActivationBackward( - ctx, out_grad, out_data, req, output); - break; - case activation::kSoftSign: - ActivationBackward( - ctx, out_grad, out_data, req, output); - break; - default: - LOG(FATAL) << "unknown activation type"; - } - }); + switch (param.act_type) { + case activation::kReLU: + ActivationBackward( + ctx, out_grad, out_data, req, output); + break; + case activation::kSigmoid: + ActivationBackward( + ctx, out_grad, out_data, req, output); + break; + case activation::kTanh: + ActivationBackward( + ctx, out_grad, out_data, req, output); + break; + case activation::kSoftReLU: + ActivationBackward( + ctx, out_grad, out_data, req, output); + break; + case activation::kSoftSign: + ActivationBackward( + ctx, out_grad, out_data, req, output); + break; + default: + LOG(FATAL) << "unknown activation type"; + } } template void ActivationCompute(const nnvm::NodeAttrs& attrs, - const OpContext& ctx, - const std::vector& inputs, - const std::vector& req, - const std::vector& outputs) { + const OpContext& ctx, + const std::vector& inputs, + const std::vector& req, + const std::vector& outputs) { CHECK_EQ(inputs.size(), 1U); CHECK_EQ(outputs.size(), 1U); const ActivationParam& param = nnvm::get(attrs.parsed); @@ -197,18 +192,19 @@ void ActivationCompute(const nnvm::NodeAttrs& attrs, template void ActivationGradCompute(const nnvm::NodeAttrs& attrs, - const OpContext& ctx, - const std::vector& inputs, - const std::vector& req, - const std::vector& outputs) { + const OpContext& ctx, + const std::vector& inputs, + const std::vector& req, + const std::vector& outputs) { + const ActivationParam& param = nnvm::get(attrs.parsed); + bool relu = param.act_type == activation::kReLU; #if (MXNET_USE_CUDNN == 1 || MXNET_USE_MKLDNN == 1) - CHECK_EQ(inputs.size(), 3U); + CHECK_EQ(inputs.size(), relu ? 2U : 3U); #else CHECK_EQ(inputs.size(), 2U); #endif CHECK_EQ(outputs.size(), 1U); CHECK_EQ(req.size(), 1U); - const ActivationParam& param = nnvm::get(attrs.parsed); ActivationGradComputeImpl(param, ctx, inputs[0], inputs[1], req[0], outputs[0]); } diff --git a/src/operator/nn/activation.cc b/src/operator/nn/activation.cc index 382efeb14478..d98548ebc0ac 100644 --- a/src/operator/nn/activation.cc +++ b/src/operator/nn/activation.cc @@ -45,7 +45,12 @@ struct ActivationGrad { std::vector heads(ograds.begin(), ograds.end()); heads.emplace_back(nnvm::NodeEntry{n, activation::kOut, 0}); #if (MXNET_USE_CUDNN == 1 || MXNET_USE_MKLDNN == 1) - heads.push_back(n->inputs[activation::kData]); + const NodeAttrs& attrs = n->attrs; + // for ReLU, no need to pass input data. This enables inplace optimization during the + // forward pass. + if (dmlc::get(attrs.parsed).act_type != activation::kReLU) { + heads.push_back(n->inputs[activation::kData]); + } #endif return MakeGradNode(op_name, n, heads, n->attrs.dict); } @@ -74,13 +79,14 @@ void ActivationGradComputeExCPU(const nnvm::NodeAttrs& attrs, const std::vector& inputs, const std::vector& req, const std::vector& outputs) { - CHECK_EQ(inputs.size(), 3U); const ActivationParam& param = nnvm::get(attrs.parsed); + bool relu = param.act_type == activation::kReLU; + CHECK_EQ(inputs.size(), relu ? 2U : 3U); if (SupportMKLDNN(inputs[0])) { MKLDNN_OPCHECK_INIT(true, outputs.size(), inputs, outputs); - MKLDNNActivationBackward(attrs, ctx, inputs[0], inputs[2], req[0], + MKLDNNActivationBackward(attrs, ctx, inputs[0], relu ? input[1] : inputs[2], req[0], outputs[0]); - MKLDNN_OPCHECK_RUN(ActivationGradCompute, attrs, ctx, inputs, req, outputs); + MKLDNN_OPCHECK_RUN(ActivationGradCompute, attrs, ctx, inputs, req, outputs); return; } ActivationGradComputeImpl(param, ctx, inputs[0].data(), inputs[1].data(), @@ -112,23 +118,28 @@ inline static bool BackwardActStorageType(const nnvm::NodeAttrs& attrs, DispatchMode* dispatch_mode, std::vector *in_attrs, std::vector *out_attrs) { + const ActivationParam& param = nnvm::get(attrs.parsed); + bool ret = false; #if (MXNET_USE_CUDNN == 1 || MXNET_USE_MKLDNN == 1) - CHECK_EQ(in_attrs->size(), 3U); + if (param.act_type != activation::kReLU) { + CHECK_EQ(in_attrs->size(), 3U); + ret = ElemwiseStorageType<3, 1, false, false, false>(attrs, dev_mask, + dispatch_mode, + in_attrs, out_attrs); + } else { + CHECK_EQ(in_attrs->size(), 2U); + ret = ElemwiseStorageType<2, 1, false, false, false>(attrs, dev_mask, + dispatch_mode, + in_attrs, out_attrs); + } #else CHECK_EQ(in_attrs->size(), 2U); + ret = ElemwiseStorageType<2, 1, false, false, false>(attrs, dev_mask, + dispatch_mode, + in_attrs, out_attrs); #endif CHECK_EQ(out_attrs->size(), 1U); -#if (MXNET_USE_CUDNN == 1 || MXNET_USE_MKLDNN == 1) - bool ret = ElemwiseStorageType<3, 1, false, false, false>(attrs, dev_mask, - dispatch_mode, - in_attrs, out_attrs); -#else - bool ret = ElemwiseStorageType<2, 1, false, false, false>(attrs, dev_mask, - dispatch_mode, - in_attrs, out_attrs); -#endif #if MXNET_USE_MKLDNN == 1 - const ActivationParam& param = nnvm::get(attrs.parsed); if (dev_mask == mshadow::cpu::kDevMask && SupportMKLDNNAct(param)) { *dispatch_mode = DispatchMode::kFComputeEx; } @@ -162,7 +173,11 @@ The following activation functions are supported: .add_arguments(ActivationParam::__FIELDS__()); NNVM_REGISTER_OP(_backward_Activation) -.set_num_inputs(3) +.set_num_inputs([](const nnvm::NodeAttrs& attrs) { + int act_type = dmlc::get(attrs.parsed).act_type; + if (act_type == activation::kReLU) return 2; + return 3; + }) .set_num_outputs(1) .set_attr("TIsBackward", true) .set_attr("FInferStorageType", BackwardActStorageType) diff --git a/src/operator/nn/activation.cu b/src/operator/nn/activation.cu index dc435b2acc17..7aa23d0af3b3 100644 --- a/src/operator/nn/activation.cu +++ b/src/operator/nn/activation.cu @@ -55,12 +55,13 @@ void ActivationCompute(const nnvm::NodeAttrs& attrs, CHECK_EQ(outputs.size(), 1U); const ActivationParam& param = nnvm::get(attrs.parsed); - // SoftReLU not supported by CUDNN yet + // SoftReLU and kSoftSign are both not supported by CUDNN yet if (param.act_type == activation::kSoftReLU) { - MSHADOW_REAL_TYPE_SWITCH(inputs[0].type_flag_, DType, { - ActivationForward(ctx, - inputs[0], req[0], outputs[0]); - }); + ActivationForward(ctx, + inputs[0], req[0], outputs[0]); + } else if (param.act_type == activation::kSoftSign) { + ActivationForward(ctx, + inputs[0], req[0], outputs[0]); } else { MSHADOW_REAL_TYPE_SWITCH(inputs[0].type_flag_, DType, { get_cudnn_op(param).Forward(ctx, inputs[0], req[0], outputs[0]); @@ -70,24 +71,27 @@ void ActivationCompute(const nnvm::NodeAttrs& attrs, template<> void ActivationGradCompute(const nnvm::NodeAttrs& attrs, - const OpContext& ctx, - const std::vector& inputs, - const std::vector& req, - const std::vector& outputs) { - CHECK_EQ(inputs.size(), 3U); + const OpContext& ctx, + const std::vector& inputs, + const std::vector& req, + const std::vector& outputs) { + const ActivationParam& param = nnvm::get(attrs.parsed); + bool relu = param.act_type == activation::kReLU; + CHECK_EQ(inputs.size(), relu ? 2U : 3U); CHECK_EQ(outputs.size(), 1U); CHECK_EQ(req.size(), 1U); - const ActivationParam& param = nnvm::get(attrs.parsed); // SoftReLU not supported by CUDNN yet if (param.act_type == activation::kSoftReLU) { - MSHADOW_REAL_TYPE_SWITCH(inputs[0].type_flag_, DType, { - ActivationBackward( - ctx, inputs[0], inputs[1], req[0], outputs[0]); - }); + ActivationBackward( + ctx, inputs[0], inputs[1], req[0], outputs[0]); + } else if (param.act_type == activation::kSoftSign) { + ActivationBackward( + ctx, inputs[0], inputs[1], req[0], outputs[0]); } else { MSHADOW_REAL_TYPE_SWITCH(inputs[0].type_flag_, DType, { - get_cudnn_op(param).Backward(ctx, inputs[0], inputs[2], inputs[1], req[0], outputs[0]); + get_cudnn_op(param).Backward(ctx, inputs[0], relu ? inputs[1] : inputs[2], + inputs[1], req[0], outputs[0]); }); } } diff --git a/src/operator/nn/cudnn/cudnn_activation-inl.h b/src/operator/nn/cudnn/cudnn_activation-inl.h index a89e7bfaf080..2c1f442808c1 100644 --- a/src/operator/nn/cudnn/cudnn_activation-inl.h +++ b/src/operator/nn/cudnn/cudnn_activation-inl.h @@ -130,6 +130,9 @@ class CuDNNActivationOp { #endif } + // backward computation for cudnn activation operator. Note that for relu + // it's okay to pass "out_data" as "in_data", since it doesn't make any + // difference in terms of computing the gradient of relu. void Backward(const OpContext &ctx, const TBlob &out_grad, const TBlob &in_data, const TBlob &out_data, const OpReqType &req, const TBlob &in_grad) { diff --git a/tests/python/gpu/test_operator_gpu.py b/tests/python/gpu/test_operator_gpu.py index 08c749e597eb..83dfc4281107 100644 --- a/tests/python/gpu/test_operator_gpu.py +++ b/tests/python/gpu/test_operator_gpu.py @@ -1133,14 +1133,17 @@ def test_fullyconnected_with_type(): @with_seed() def test_activation_with_type(): - sym = mx.sym.Activation(name='act', act_type='sigmoid') - ctx_list = [{'ctx': mx.gpu(0), 'act_data': (2, 2, 10, 10), 'type_dict': {'act_data': np.float64}}, - {'ctx': mx.gpu(0), 'act_data': (2, 2, 10, 10), 'type_dict': {'act_data': np.float32}}, - {'ctx': mx.gpu(0), 'act_data': (2, 2, 10, 10), 'type_dict': {'act_data': np.float16}}, - {'ctx': mx.cpu(0), 'act_data': (2, 2, 10, 10), 'type_dict': {'act_data': np.float64}}, - {'ctx': mx.cpu(0), 'act_data': (2, 2, 10, 10), 'type_dict': {'act_data': np.float32}}, - {'ctx': mx.cpu(0), 'act_data': (2, 2, 10, 10), 'type_dict': {'act_data': np.float16}}] - check_consistency(sym, ctx_list) + act_types = ['relu', 'sigmoid', 'tanh', 'softrelu', 'softsign'] + shape = (2, 2, 10, 10) + for act_type in act_types: + sym = mx.sym.Activation(name='act', act_type=act_type) + ctx_list = [{'ctx': mx.gpu(0), 'act_data': shape, 'type_dict': {'act_data': np.float64}}, + {'ctx': mx.gpu(0), 'act_data': shape, 'type_dict': {'act_data': np.float32}}, + {'ctx': mx.gpu(0), 'act_data': shape, 'type_dict': {'act_data': np.float16}}, + {'ctx': mx.cpu(0), 'act_data': shape, 'type_dict': {'act_data': np.float64}}, + {'ctx': mx.cpu(0), 'act_data': shape, 'type_dict': {'act_data': np.float32}}, + {'ctx': mx.cpu(0), 'act_data': shape, 'type_dict': {'act_data': np.float16}}] + check_consistency(sym, ctx_list) @with_seed() diff --git a/tests/python/unittest/test_operator.py b/tests/python/unittest/test_operator.py index 7ee67dd20660..67e2e663db1a 100644 --- a/tests/python/unittest/test_operator.py +++ b/tests/python/unittest/test_operator.py @@ -5816,6 +5816,10 @@ def get_output_names_callback(name, arr): name='pooling') check_name(us_sym, ['pooling_output']) +@with_seed() +def test_activation(): + pass + if __name__ == '__main__': import nose nose.runmodule() From c17b5ca471920ffccb47e02654a6b7e5b92a8785 Mon Sep 17 00:00:00 2001 From: eric-haibin-lin Date: Tue, 8 May 2018 06:43:50 +0000 Subject: [PATCH 2/7] inplace relu --- src/operator/tensor/elemwise_unary_op_basic.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/operator/tensor/elemwise_unary_op_basic.cc b/src/operator/tensor/elemwise_unary_op_basic.cc index fda47fc5f95e..93519d3ee8f7 100644 --- a/src/operator/tensor/elemwise_unary_op_basic.cc +++ b/src/operator/tensor/elemwise_unary_op_basic.cc @@ -86,7 +86,7 @@ The storage type of ``relu`` output depends upon the input storage type: .set_attr("FInferStorageType", ElemwiseStorageType<1, 1, false, true, false>) .set_attr("FCompute", UnaryOp::Compute) .set_attr("FComputeEx", UnaryOp::ComputeEx) -.set_attr("FGradient", ElemwiseGradUseIn{"_backward_relu"}); +.set_attr("FGradient", ElemwiseGradUseOut{"_backward_relu"}); MXNET_OPERATOR_REGISTER_BINARY_WITH_SPARSE_CPU(_backward_relu, unary_bwd); From 3acbe79028872ea6475be81907ca133546c0e435 Mon Sep 17 00:00:00 2001 From: eric-haibin-lin Date: Tue, 8 May 2018 06:46:20 +0000 Subject: [PATCH 3/7] add comments --- src/operator/nn/mkldnn/mkldnn_act.cc | 2 ++ tests/python/unittest/test_operator.py | 4 ---- 2 files changed, 2 insertions(+), 4 deletions(-) diff --git a/src/operator/nn/mkldnn/mkldnn_act.cc b/src/operator/nn/mkldnn/mkldnn_act.cc index 9be5bfbc150d..a057527d473f 100644 --- a/src/operator/nn/mkldnn/mkldnn_act.cc +++ b/src/operator/nn/mkldnn/mkldnn_act.cc @@ -165,6 +165,8 @@ void MKLDNNActivationForward(const nnvm::NodeAttrs& attrs, const OpContext &ctx, stream->Submit(); } +// For backward relu activation, it's okay to pass "out_data" as "in_data" to this +// function, since the computation only involes non-zeros. void MKLDNNActivationBackward(const nnvm::NodeAttrs& attrs, const OpContext &ctx, const NDArray &out_grad, const NDArray &in_data, const OpReqType &req, const NDArray &in_grad) { diff --git a/tests/python/unittest/test_operator.py b/tests/python/unittest/test_operator.py index 67e2e663db1a..7ee67dd20660 100644 --- a/tests/python/unittest/test_operator.py +++ b/tests/python/unittest/test_operator.py @@ -5816,10 +5816,6 @@ def get_output_names_callback(name, arr): name='pooling') check_name(us_sym, ['pooling_output']) -@with_seed() -def test_activation(): - pass - if __name__ == '__main__': import nose nose.runmodule() From abd1ce2374de27c6c1cf38f56010012a85b50954 Mon Sep 17 00:00:00 2001 From: eric-haibin-lin Date: Tue, 8 May 2018 06:51:50 +0000 Subject: [PATCH 4/7] add commnet --- src/operator/nn/activation.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/operator/nn/activation.cu b/src/operator/nn/activation.cu index 7aa23d0af3b3..74d4e4e32b0d 100644 --- a/src/operator/nn/activation.cu +++ b/src/operator/nn/activation.cu @@ -81,7 +81,7 @@ void ActivationGradCompute(const nnvm::NodeAttrs& attrs, CHECK_EQ(outputs.size(), 1U); CHECK_EQ(req.size(), 1U); - // SoftReLU not supported by CUDNN yet + // both SoftReLU and SoftSign not supported by CUDNN yet if (param.act_type == activation::kSoftReLU) { ActivationBackward( ctx, inputs[0], inputs[1], req[0], outputs[0]); From d9283f3ed85165395cefbfaa1cd4be026d4012f7 Mon Sep 17 00:00:00 2001 From: eric-haibin-lin Date: Tue, 8 May 2018 06:58:12 +0000 Subject: [PATCH 5/7] comments --- src/operator/nn/activation.cc | 5 ++++- src/operator/nn/activation.cu | 1 + 2 files changed, 5 insertions(+), 1 deletion(-) diff --git a/src/operator/nn/activation.cc b/src/operator/nn/activation.cc index d98548ebc0ac..4a0eeeb95a0f 100644 --- a/src/operator/nn/activation.cc +++ b/src/operator/nn/activation.cc @@ -84,9 +84,10 @@ void ActivationGradComputeExCPU(const nnvm::NodeAttrs& attrs, CHECK_EQ(inputs.size(), relu ? 2U : 3U); if (SupportMKLDNN(inputs[0])) { MKLDNN_OPCHECK_INIT(true, outputs.size(), inputs, outputs); + // XXX: for y = relu(x), y is passed as "in_data" to Backward() MKLDNNActivationBackward(attrs, ctx, inputs[0], relu ? input[1] : inputs[2], req[0], outputs[0]); - MKLDNN_OPCHECK_RUN(ActivationGradCompute, attrs, ctx, inputs, req, outputs); + MKLDNN_OPCHECK_RUN(ActivationGradCompute, attrs, ctx, inputs, req, outputs); return; } ActivationGradComputeImpl(param, ctx, inputs[0].data(), inputs[1].data(), @@ -127,6 +128,7 @@ inline static bool BackwardActStorageType(const nnvm::NodeAttrs& attrs, dispatch_mode, in_attrs, out_attrs); } else { + // for ReLU activation, the backward pass only needs ograd and output CHECK_EQ(in_attrs->size(), 2U); ret = ElemwiseStorageType<2, 1, false, false, false>(attrs, dev_mask, dispatch_mode, @@ -175,6 +177,7 @@ The following activation functions are supported: NNVM_REGISTER_OP(_backward_Activation) .set_num_inputs([](const nnvm::NodeAttrs& attrs) { int act_type = dmlc::get(attrs.parsed).act_type; + // for ReLU activation, the backward pass only needs ograd and output if (act_type == activation::kReLU) return 2; return 3; }) diff --git a/src/operator/nn/activation.cu b/src/operator/nn/activation.cu index 74d4e4e32b0d..68b4053efdda 100644 --- a/src/operator/nn/activation.cu +++ b/src/operator/nn/activation.cu @@ -90,6 +90,7 @@ void ActivationGradCompute(const nnvm::NodeAttrs& attrs, ctx, inputs[0], inputs[1], req[0], outputs[0]); } else { MSHADOW_REAL_TYPE_SWITCH(inputs[0].type_flag_, DType, { + // XXX: for y = relu(x), y is passed as "in_data" to Backward() get_cudnn_op(param).Backward(ctx, inputs[0], relu ? inputs[1] : inputs[2], inputs[1], req[0], outputs[0]); }); From 02b677dcbe3ec693017563fd453f30e7cd8686e2 Mon Sep 17 00:00:00 2001 From: eric-haibin-lin Date: Tue, 8 May 2018 16:26:08 +0000 Subject: [PATCH 6/7] fix compilation error --- src/operator/nn/activation-inl.h | 2 +- src/operator/nn/activation.cc | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/src/operator/nn/activation-inl.h b/src/operator/nn/activation-inl.h index 403ac48a18d4..a9f6dbeda894 100644 --- a/src/operator/nn/activation-inl.h +++ b/src/operator/nn/activation-inl.h @@ -197,8 +197,8 @@ void ActivationGradCompute(const nnvm::NodeAttrs& attrs, const std::vector& req, const std::vector& outputs) { const ActivationParam& param = nnvm::get(attrs.parsed); - bool relu = param.act_type == activation::kReLU; #if (MXNET_USE_CUDNN == 1 || MXNET_USE_MKLDNN == 1) + bool relu = param.act_type == activation::kReLU; CHECK_EQ(inputs.size(), relu ? 2U : 3U); #else CHECK_EQ(inputs.size(), 2U); diff --git a/src/operator/nn/activation.cc b/src/operator/nn/activation.cc index 4a0eeeb95a0f..595b8912ccc7 100644 --- a/src/operator/nn/activation.cc +++ b/src/operator/nn/activation.cc @@ -85,7 +85,7 @@ void ActivationGradComputeExCPU(const nnvm::NodeAttrs& attrs, if (SupportMKLDNN(inputs[0])) { MKLDNN_OPCHECK_INIT(true, outputs.size(), inputs, outputs); // XXX: for y = relu(x), y is passed as "in_data" to Backward() - MKLDNNActivationBackward(attrs, ctx, inputs[0], relu ? input[1] : inputs[2], req[0], + MKLDNNActivationBackward(attrs, ctx, inputs[0], relu ? inputs[1] : inputs[2], req[0], outputs[0]); MKLDNN_OPCHECK_RUN(ActivationGradCompute, attrs, ctx, inputs, req, outputs); return; @@ -119,9 +119,9 @@ inline static bool BackwardActStorageType(const nnvm::NodeAttrs& attrs, DispatchMode* dispatch_mode, std::vector *in_attrs, std::vector *out_attrs) { - const ActivationParam& param = nnvm::get(attrs.parsed); bool ret = false; #if (MXNET_USE_CUDNN == 1 || MXNET_USE_MKLDNN == 1) + const ActivationParam& param = nnvm::get(attrs.parsed); if (param.act_type != activation::kReLU) { CHECK_EQ(in_attrs->size(), 3U); ret = ElemwiseStorageType<3, 1, false, false, false>(attrs, dev_mask, From 66b66f967dabab10e210b83a176881060781ed56 Mon Sep 17 00:00:00 2001 From: eric-haibin-lin Date: Tue, 8 May 2018 19:40:15 +0000 Subject: [PATCH 7/7] add check_numerical_grad test --- tests/python/unittest/test_operator.py | 42 ++++++++++++++++++++++++++ 1 file changed, 42 insertions(+) diff --git a/tests/python/unittest/test_operator.py b/tests/python/unittest/test_operator.py index 7ee67dd20660..d6486e3781ca 100644 --- a/tests/python/unittest/test_operator.py +++ b/tests/python/unittest/test_operator.py @@ -5816,6 +5816,48 @@ def get_output_names_callback(name, arr): name='pooling') check_name(us_sym, ['pooling_output']) +@with_seed() +def test_activation(): + shape=(9, 10) + dtype_l = [np.float64, np.float32, np.float16] + rtol_l = [1e-7, 1e-6, 1e-2] + atol_l = [1e-7, 1e-6, 1e-2] + rtol_fd = 1e-5 + atol_fd = 1e-6 + num_eps = 1e-6 + unary_ops = { + 'relu': [lambda x: mx.sym.Activation(x, act_type='relu'), + lambda x: np.maximum(x, 0.), + lambda x: 1. * (x > 0.), + -5.0, 5.0], + 'sigmoid': [lambda x: mx.sym.Activation(x, act_type='sigmoid'), + lambda x: 1. / (np.exp(-x) + 1.), + lambda x: 1. / (np.exp(-x) + 1.) / (np.exp(x) + 1.), + -3.0, 3.0], + 'tanh': [lambda x: mx.sym.Activation(x, act_type='tanh'), + lambda x: np.tanh(x), + lambda x: 1. - np.tanh(x) ** 2, + -4.0, 4.0], + 'softrelu': [lambda x: mx.sym.Activation(x, act_type='softrelu'), + lambda x: np.log(1. + np.exp(x)), + lambda x: 1. - 1 / (1 + np.exp(x)), + -3.0, 3.0], + } + # Loop over operators + for name, op in unary_ops.items(): + # Loop over dtype's + for ind in range(len(dtype_l)): + dtype = dtype_l[ind] + rtol = rtol_l[ind] + atol = atol_l[ind] + compare_forw_backw_unary_op( + name, op[0], op[1], op[2], shape, op[3], op[4], rtol, atol, + dtype) + # Finite difference testing + finite_diff_unary_op( + name, op[0], shape, op[3], op[4], rtol_fd, atol_fd, num_eps) + + if __name__ == '__main__': import nose nose.runmodule()