From 79e6e50b7ebf16adf73c7d14719fef6ee2833381 Mon Sep 17 00:00:00 2001 From: liuqi Date: Mon, 5 Mar 2018 11:21:23 +0800 Subject: [PATCH] Fix PReLU bug: not fuse PReLU anymore. --- mace/kernels/activation.h | 65 +++++++++++-------- mace/kernels/batch_norm.h | 19 ++---- mace/kernels/conv_2d.h | 21 ++---- mace/kernels/depthwise_conv2d.h | 21 ++---- mace/kernels/fully_connected.h | 19 ++---- mace/kernels/opencl/activation_opencl.cc | 5 +- mace/kernels/opencl/batch_norm_opencl.cc | 4 -- mace/kernels/opencl/cl/activation.cl | 11 +++- mace/kernels/opencl/cl/batch_norm.cl | 7 +- mace/kernels/opencl/cl/common.h | 6 +- mace/kernels/opencl/cl/conv_2d.cl | 11 ++-- mace/kernels/opencl/cl/conv_2d_1x1.cl | 11 ++-- mace/kernels/opencl/cl/conv_2d_3x3.cl | 13 ++-- mace/kernels/opencl/cl/depthwise_conv2d.cl | 22 +++---- mace/kernels/opencl/cl/fully_connected.cl | 7 +- mace/kernels/opencl/cl/winograd_transform.cl | 15 ++--- mace/kernels/opencl/conv_2d_opencl.cc | 11 +--- mace/kernels/opencl/conv_2d_opencl_1x1.cc | 5 -- mace/kernels/opencl/conv_2d_opencl_3x3.cc | 5 -- mace/kernels/opencl/conv_2d_opencl_general.cc | 5 -- mace/kernels/opencl/depthwise_conv_opencl.cc | 11 +--- mace/kernels/opencl/fully_connected_opencl.cc | 4 -- mace/kernels/opencl/winograd_transform.cc | 1 - mace/kernels/winograd_transform.h | 17 ++--- mace/ops/activation.h | 8 +-- mace/ops/activation_test.cc | 13 ++-- mace/ops/batch_norm.h | 2 +- mace/ops/conv_2d.h | 1 - mace/ops/conv_2d_test.cc | 4 +- mace/ops/depthwise_conv2d.h | 3 +- mace/ops/folded_batch_norm.h | 3 +- mace/ops/fully_connected.h | 3 +- mace/ops/fused_conv_2d.h | 3 +- mace/ops/winograd_inverse_transform.h | 3 +- 34 files changed, 155 insertions(+), 204 deletions(-) diff --git a/mace/kernels/activation.h b/mace/kernels/activation.h index 83acf4fb..b768eb28 100644 --- a/mace/kernels/activation.h +++ b/mace/kernels/activation.h @@ -46,8 +46,7 @@ void DoActivation(const T *input_ptr, T *output_ptr, const index_t size, const ActivationType type, - const float relux_max_limit, - const float prelu_alpha) { + const float relux_max_limit) { MACE_CHECK(DataTypeToEnum::value != DataType::DT_HALF); switch (type) { @@ -66,17 +65,6 @@ void DoActivation(const T *input_ptr, static_cast(relux_max_limit)); } break; - case PRELU: -#pragma omp parallel for - for (index_t i = 0; i < size; ++i) { - T in = input_ptr[i]; - if (in < 0) { - output_ptr[i] = in * prelu_alpha; - } else { - output_ptr[i] = in; - } - } - break; case TANH: #pragma omp parallel for for (index_t i = 0; i < size; ++i) { @@ -95,45 +83,70 @@ void DoActivation(const T *input_ptr, } } +template +void PReLUActivation(const T *input_ptr, + const index_t size, + const index_t input_chan, + const T *alpha_ptr, + T *output_ptr) { +#pragma omp parallel for + for (index_t i = 0; i < size; ++i) { + const index_t chan_idx = i % input_chan; + T in = input_ptr[i]; + if (in < 0) { + output_ptr[i] = in * alpha_ptr[chan_idx]; + } else { + output_ptr[i] = in; + } + } + +} + template class ActivationFunctor { public: - ActivationFunctor(ActivationType type, T relux_max_limit, T prelu_alpha) + ActivationFunctor(ActivationType type, T relux_max_limit) : activation_(type), - relux_max_limit_(relux_max_limit), - prelu_alpha_(prelu_alpha) {} + relux_max_limit_(relux_max_limit){} - void operator()(const Tensor *input, Tensor *output, StatsFuture *future) { + void operator()(const Tensor *input, + const Tensor *alpha, + Tensor *output, + StatsFuture *future) { const T *input_ptr = input->data(); T *output_ptr = output->mutable_data(); - DoActivation(input_ptr, output_ptr, output->size(), activation_, relux_max_limit_, - prelu_alpha_); + if (activation_ == PRELU) { + const T *alpha_ptr = alpha == nullptr ? nullptr : alpha->data(); + PReLUActivation(input_ptr, output->size(), input->dim(3), alpha_ptr, output_ptr); + } else { + DoActivation(input_ptr, output_ptr, output->size(), activation_, relux_max_limit_); + } } private: ActivationType activation_; T relux_max_limit_; - T prelu_alpha_; }; template <> void ActivationFunctor::operator()( - const Tensor *input, Tensor *output, StatsFuture *future); + const Tensor *input, const Tensor *alpha, Tensor *output, StatsFuture *future); template class ActivationFunctor { public: - ActivationFunctor(ActivationType type, T relux_max_limit, T prelu_alpha) + ActivationFunctor(ActivationType type, T relux_max_limit) : activation_(type), - relux_max_limit_(relux_max_limit), - prelu_alpha_(prelu_alpha) {} + relux_max_limit_(relux_max_limit){} - void operator()(const Tensor *input, Tensor *output, StatsFuture *future); + void operator()(const Tensor *input, + const Tensor *alpha, + Tensor *output, + StatsFuture *future); private: ActivationType activation_; T relux_max_limit_; - T prelu_alpha_; cl::Kernel kernel_; std::string tuning_key_prefix_; }; diff --git a/mace/kernels/batch_norm.h b/mace/kernels/batch_norm.h index bf5035de..0d489f40 100644 --- a/mace/kernels/batch_norm.h +++ b/mace/kernels/batch_norm.h @@ -21,27 +21,23 @@ namespace kernels { struct BatchNormFunctorBase { BatchNormFunctorBase(bool folded_constant, const ActivationType activation, - const float relux_max_limit, - const float prelu_alpha) + const float relux_max_limit) : folded_constant_(folded_constant), activation_(activation), - relux_max_limit_(relux_max_limit), - prelu_alpha_(prelu_alpha) {} + relux_max_limit_(relux_max_limit){} const bool folded_constant_; const ActivationType activation_; const float relux_max_limit_; - const float prelu_alpha_; }; template struct BatchNormFunctor : BatchNormFunctorBase { BatchNormFunctor(const bool folded_constant, const ActivationType activation, - const float relux_max_limit, - const float prelu_alpha) + const float relux_max_limit) : BatchNormFunctorBase( - folded_constant, activation, relux_max_limit, prelu_alpha) {} + folded_constant, activation, relux_max_limit) {} void operator()(const Tensor *input, const Tensor *scale, @@ -132,7 +128,7 @@ struct BatchNormFunctor : BatchNormFunctorBase { } } DoActivation(output_ptr, output_ptr, output->NumElements(), activation_, - relux_max_limit_, prelu_alpha_); + relux_max_limit_); } }; @@ -150,10 +146,9 @@ template struct BatchNormFunctor : BatchNormFunctorBase { BatchNormFunctor(const bool folded_constant, const ActivationType activation, - const float relux_max_limit, - const float prelu_alpha) + const float relux_max_limit) : BatchNormFunctorBase( - folded_constant, activation, relux_max_limit, prelu_alpha) {} + folded_constant, activation, relux_max_limit) {} void operator()(const Tensor *input, const Tensor *scale, const Tensor *offset, diff --git a/mace/kernels/conv_2d.h b/mace/kernels/conv_2d.h index 91e1471f..e6c22cd9 100644 --- a/mace/kernels/conv_2d.h +++ b/mace/kernels/conv_2d.h @@ -182,15 +182,13 @@ struct Conv2dFunctorBase { const std::vector &paddings, const int *dilations, const ActivationType activation, - const float relux_max_limit, - const float prelu_alpha) + const float relux_max_limit) : strides_(strides), padding_type_(padding_type), paddings_(paddings), dilations_(dilations), activation_(activation), - relux_max_limit_(relux_max_limit), - prelu_alpha_(prelu_alpha) {} + relux_max_limit_(relux_max_limit){} const int *strides_; // [stride_h, stride_w] const Padding padding_type_; @@ -198,7 +196,6 @@ struct Conv2dFunctorBase { const int *dilations_; // [dilation_h, dilation_w] const ActivationType activation_; const float relux_max_limit_; - const float prelu_alpha_; }; template @@ -208,15 +205,13 @@ struct Conv2dFunctor : Conv2dFunctorBase { const std::vector &paddings, const int *dilations, const ActivationType activation, - const float relux_max_limit, - const float prelu_alpha) + const float relux_max_limit) : Conv2dFunctorBase(strides, padding_type, paddings, dilations, activation, - relux_max_limit, - prelu_alpha) {} + relux_max_limit) {} void operator()(const Tensor *input, // NHWC const Tensor *filter, // HWOI @@ -622,7 +617,7 @@ struct Conv2dFunctor : Conv2dFunctorBase { } } DoActivation(output_data, output_data, output->NumElements(), activation_, - relux_max_limit_, prelu_alpha_); + relux_max_limit_); } }; @@ -640,15 +635,13 @@ struct Conv2dFunctor : Conv2dFunctorBase { const std::vector &paddings, const int *dilations, const ActivationType activation, - const float relux_max_limit, - const float prelu_alpha) + const float relux_max_limit) : Conv2dFunctorBase(strides, padding_type, paddings, dilations, activation, - relux_max_limit, - prelu_alpha) {} + relux_max_limit) {} void operator()(const Tensor *input, const Tensor *filter, diff --git a/mace/kernels/depthwise_conv2d.h b/mace/kernels/depthwise_conv2d.h index ef78e17e..141119a9 100644 --- a/mace/kernels/depthwise_conv2d.h +++ b/mace/kernels/depthwise_conv2d.h @@ -241,15 +241,13 @@ struct DepthwiseConv2dFunctorBase { const std::vector &paddings, const int *dilations, const ActivationType activation, - const float relux_max_limit, - const float prelu_alpha) + const float relux_max_limit) : strides_(strides), padding_type_(padding_type), paddings_(paddings), dilations_(dilations), activation_(activation), - relux_max_limit_(relux_max_limit), - prelu_alpha_(prelu_alpha) {} + relux_max_limit_(relux_max_limit){} const int *strides_; // [stride_h, stride_w] const Padding padding_type_; @@ -257,7 +255,6 @@ struct DepthwiseConv2dFunctorBase { const int *dilations_; // [dilation_h, dilation_w] const ActivationType activation_; const float relux_max_limit_; - const float prelu_alpha_; }; template @@ -267,15 +264,13 @@ struct DepthwiseConv2dFunctor : public DepthwiseConv2dFunctorBase { const std::vector &paddings, const int *dilations, const ActivationType activation, - const float relux_max_limit, - const float prelu_alpha) + const float relux_max_limit) : DepthwiseConv2dFunctorBase(strides, padding_type, paddings, dilations, activation, - relux_max_limit, - prelu_alpha) {} + relux_max_limit) {} void operator()(const Tensor *input, // NHWC const Tensor *filter, // HWIM @@ -408,7 +403,7 @@ struct DepthwiseConv2dFunctor : public DepthwiseConv2dFunctorBase { output_ptr = output->mutable_data(); DoActivation(output_ptr, output_ptr, output->NumElements(), activation_, - relux_max_limit_, prelu_alpha_); + relux_max_limit_); } }; @@ -428,15 +423,13 @@ struct DepthwiseConv2dFunctor const std::vector &paddings, const int *dilations, const ActivationType activation, - const float relux_max_limit, - const float prelu_alpha) + const float relux_max_limit) : DepthwiseConv2dFunctorBase(strides, padding_type, paddings, dilations, activation, - relux_max_limit, - prelu_alpha) {} + relux_max_limit) {} void operator()(const Tensor *input, const Tensor *filter, diff --git a/mace/kernels/fully_connected.h b/mace/kernels/fully_connected.h index c3b66222..7d072016 100644 --- a/mace/kernels/fully_connected.h +++ b/mace/kernels/fully_connected.h @@ -15,23 +15,19 @@ namespace kernels { struct FullyConnectedBase { FullyConnectedBase(const ActivationType activation, - const float relux_max_limit, - const float prelu_alpha) + const float relux_max_limit) : activation_(activation), - relux_max_limit_(relux_max_limit), - prelu_alpha_(prelu_alpha) {} + relux_max_limit_(relux_max_limit){} const ActivationType activation_; const float relux_max_limit_; - const float prelu_alpha_; }; template struct FullyConnectedFunctor : FullyConnectedBase { FullyConnectedFunctor(const ActivationType activation, - const float relux_max_limit, - const float prelu_alpha) : - FullyConnectedBase(activation, relux_max_limit, prelu_alpha) {} + const float relux_max_limit) : + FullyConnectedBase(activation, relux_max_limit) {} void operator()(const Tensor *input, const Tensor *weight, @@ -70,16 +66,15 @@ struct FullyConnectedFunctor : FullyConnectedBase { } DoActivation(output_ptr, output_ptr, output->NumElements(), activation_, - relux_max_limit_, prelu_alpha_); + relux_max_limit_); } }; template struct FullyConnectedFunctor : FullyConnectedBase { FullyConnectedFunctor(const ActivationType activation, - const float relux_max_limit, - const float prelu_alpha) : - FullyConnectedBase(activation, relux_max_limit, prelu_alpha) {} + const float relux_max_limit) : + FullyConnectedBase(activation, relux_max_limit) {} void operator()(const Tensor *input, const Tensor *weight, diff --git a/mace/kernels/opencl/activation_opencl.cc b/mace/kernels/opencl/activation_opencl.cc index 934ce8da..dee01087 100644 --- a/mace/kernels/opencl/activation_opencl.cc +++ b/mace/kernels/opencl/activation_opencl.cc @@ -14,6 +14,7 @@ namespace kernels { template void ActivationFunctor::operator()(const Tensor *input, + const Tensor *alpha, Tensor *output, StatsFuture *future) { const index_t batch = input->dim(0); @@ -60,8 +61,10 @@ void ActivationFunctor::operator()(const Tensor *input, runtime->BuildKernel("activation", kernel_name, built_options); int idx = 0; kernel_.setArg(idx++, *(static_cast(input->buffer()))); + if (activation_ == PRELU) { + kernel_.setArg(idx++, *(static_cast(alpha->buffer()))); + } kernel_.setArg(idx++, static_cast(relux_max_limit_)); - kernel_.setArg(idx++, static_cast(prelu_alpha_)); kernel_.setArg(idx++, *(static_cast(output->buffer()))); } diff --git a/mace/kernels/opencl/batch_norm_opencl.cc b/mace/kernels/opencl/batch_norm_opencl.cc index d88fed51..7696e875 100644 --- a/mace/kernels/opencl/batch_norm_opencl.cc +++ b/mace/kernels/opencl/batch_norm_opencl.cc @@ -50,9 +50,6 @@ void BatchNormFunctor::operator()(const Tensor *input, case RELUX: built_options.emplace("-DUSE_RELUX"); break; - case PRELU: - built_options.emplace("-DUSE_PRELU"); - break; case TANH: built_options.emplace("-DUSE_TANH"); break; @@ -79,7 +76,6 @@ void BatchNormFunctor::operator()(const Tensor *input, } kernel_.setArg(idx++, *(static_cast(output->buffer()))); kernel_.setArg(idx++, relux_max_limit_); - kernel_.setArg(idx++, prelu_alpha_); } const uint32_t gws[3] = {static_cast(channel_blocks), diff --git a/mace/kernels/opencl/cl/activation.cl b/mace/kernels/opencl/cl/activation.cl index fe8619e2..bee0b0e3 100644 --- a/mace/kernels/opencl/cl/activation.cl +++ b/mace/kernels/opencl/cl/activation.cl @@ -1,8 +1,10 @@ #include __kernel void activation(__read_only image2d_t input, +#ifdef USE_PRELU + __read_only image2d_t alpha, +#endif __private const float relux_max_limit, - __private const float prelu_alpha, __write_only image2d_t output) { const int ch_blk = get_global_id(0); const int w = get_global_id(1); @@ -11,7 +13,12 @@ __kernel void activation(__read_only image2d_t input, const int pos = mad24(ch_blk, width, w); DATA_TYPE4 in = READ_IMAGET(input, SAMPLER, (int2)(pos, hb)); - DATA_TYPE4 out = do_activation(in, relux_max_limit, prelu_alpha); +#ifdef USE_PRELU + DATA_TYPE4 prelu_alpha = READ_IMAGET(alpha, SAMPLER, (int2)(ch_blk, 0)); + DATA_TYPE4 out = do_activation(in, prelu_alpha, relux_max_limit); +#else + DATA_TYPE4 out = do_activation(in, relux_max_limit); +#endif WRITE_IMAGET(output, (int2)(pos, hb), out); } diff --git a/mace/kernels/opencl/cl/batch_norm.cl b/mace/kernels/opencl/cl/batch_norm.cl index 99c00fab..773b59c4 100644 --- a/mace/kernels/opencl/cl/batch_norm.cl +++ b/mace/kernels/opencl/cl/batch_norm.cl @@ -9,8 +9,7 @@ __kernel void batch_norm(__read_only image2d_t input, __private const float epsilon, #endif __write_only image2d_t output, - __private const float relux_max_limit, - __private const float prelu_alpha) { + __private const float relux_max_limit) { const int ch_blk = get_global_id(0); const int w = get_global_id(1); const int hb = get_global_id(2); @@ -35,8 +34,8 @@ __kernel void batch_norm(__read_only image2d_t input, DATA_TYPE4 in = READ_IMAGET(input, SAMPLER, (int2)(pos, hb)); DATA_TYPE4 out = mad(in, bn_scale, bn_offset); -#if defined(USE_RELU) || defined(USE_RELUX) || defined(USE_PRELU) || defined(USE_TANH) || defined(USE_SIGMOID) - out = do_activation(out, relux_max_limit, prelu_alpha); +#if defined(USE_RELU) || defined(USE_RELUX) || defined(USE_TANH) || defined(USE_SIGMOID) + out = do_activation(out, relux_max_limit); #endif WRITE_IMAGET(output, (int2)(pos, hb), out); diff --git a/mace/kernels/opencl/cl/common.h b/mace/kernels/opencl/cl/common.h index 13b20e05..28b9addd 100644 --- a/mace/kernels/opencl/cl/common.h +++ b/mace/kernels/opencl/cl/common.h @@ -22,8 +22,10 @@ __constant sampler_t SAMPLER = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | inline DATA_TYPE4 do_activation(DATA_TYPE4 in, - __private const float relux_max_limit, - __private const float prelu_alpha) { +#ifdef USE_PRELU + DATA_TYPE4 prelu_alpha, +#endif + __private const float relux_max_limit) { DATA_TYPE4 out; #ifdef USE_RELU out = fmax(in, 0); diff --git a/mace/kernels/opencl/cl/conv_2d.cl b/mace/kernels/opencl/cl/conv_2d.cl index 35e17da8..11253d69 100644 --- a/mace/kernels/opencl/cl/conv_2d.cl +++ b/mace/kernels/opencl/cl/conv_2d.cl @@ -7,7 +7,6 @@ __kernel void conv_2d(__read_only image2d_t input, /* [c%4 * w * c/4, h * b] */ #endif __write_only image2d_t output, __private const float relux_max_limit, - __private const float prelu_alpha, __private const int in_height, __private const int in_width, __private const int in_ch_blks, @@ -112,11 +111,11 @@ __kernel void conv_2d(__read_only image2d_t input, /* [c%4 * w * c/4, h * b] */ } } -#if defined(USE_RELU) || defined(USE_RELUX) || defined(USE_PRELU) || defined(USE_TANH) || defined(USE_SIGMOID) - out0 = do_activation(out0, relux_max_limit, prelu_alpha); - out1 = do_activation(out1, relux_max_limit, prelu_alpha); - out2 = do_activation(out2, relux_max_limit, prelu_alpha); - out3 = do_activation(out3, relux_max_limit, prelu_alpha); +#if defined(USE_RELU) || defined(USE_RELUX) || defined(USE_TANH) || defined(USE_SIGMOID) + out0 = do_activation(out0, relux_max_limit); + out1 = do_activation(out1, relux_max_limit); + out2 = do_activation(out2, relux_max_limit); + out3 = do_activation(out3, relux_max_limit); #endif const int out_x_base = mul24(out_ch_blk, out_width); diff --git a/mace/kernels/opencl/cl/conv_2d_1x1.cl b/mace/kernels/opencl/cl/conv_2d_1x1.cl index ebcbb231..b695165e 100644 --- a/mace/kernels/opencl/cl/conv_2d_1x1.cl +++ b/mace/kernels/opencl/cl/conv_2d_1x1.cl @@ -7,7 +7,6 @@ __kernel void conv_2d_1x1(__read_only image2d_t input, /* [c%4 * w * c/4, h * b] #endif __write_only image2d_t output, __private const float relux_max_limit, - __private const float prelu_alpha, __private const int in_height, __private const int in_width, __private const int in_ch_blks, @@ -86,11 +85,11 @@ __kernel void conv_2d_1x1(__read_only image2d_t input, /* [c%4 * w * c/4, h * b] filter_x_base += 4; } -#if defined(USE_RELU) || defined(USE_RELUX) || defined(USE_PRELU) || defined(USE_TANH) || defined(USE_SIGMOID) - out0 = do_activation(out0, relux_max_limit, prelu_alpha); - out1 = do_activation(out1, relux_max_limit, prelu_alpha); - out2 = do_activation(out2, relux_max_limit, prelu_alpha); - out3 = do_activation(out3, relux_max_limit, prelu_alpha); +#if defined(USE_RELU) || defined(USE_RELUX) || defined(USE_TANH) || defined(USE_SIGMOID) + out0 = do_activation(out0, relux_max_limit); + out1 = do_activation(out1, relux_max_limit); + out2 = do_activation(out2, relux_max_limit); + out3 = do_activation(out3, relux_max_limit); #endif const int out_x_base = mul24(out_ch_blk, width); diff --git a/mace/kernels/opencl/cl/conv_2d_3x3.cl b/mace/kernels/opencl/cl/conv_2d_3x3.cl index 0fce8236..fad561aa 100644 --- a/mace/kernels/opencl/cl/conv_2d_3x3.cl +++ b/mace/kernels/opencl/cl/conv_2d_3x3.cl @@ -7,7 +7,6 @@ __kernel void conv_2d_3x3(__read_only image2d_t input, /* [c%4 * w * c/4, h * b] #endif __write_only image2d_t output, __private const float relux_max_limit, - __private const float prelu_alpha, __private const int in_height, __private const int in_width, __private const int in_ch_blks, @@ -120,12 +119,12 @@ __kernel void conv_2d_3x3(__read_only image2d_t input, /* [c%4 * w * c/4, h * b] } } -#if defined(USE_RELU) || defined(USE_RELUX) || defined(USE_PRELU) || defined(USE_TANH) || defined(USE_SIGMOID) - out0 = do_activation(out0, relux_max_limit, prelu_alpha); - out1 = do_activation(out1, relux_max_limit, prelu_alpha); - out2 = do_activation(out2, relux_max_limit, prelu_alpha); - out3 = do_activation(out3, relux_max_limit, prelu_alpha); - out4 = do_activation(out4, relux_max_limit, prelu_alpha); +#if defined(USE_RELU) || defined(USE_RELUX) || defined(USE_TANH) || defined(USE_SIGMOID) + out0 = do_activation(out0, relux_max_limit); + out1 = do_activation(out1, relux_max_limit); + out2 = do_activation(out2, relux_max_limit); + out3 = do_activation(out3, relux_max_limit); + out4 = do_activation(out4, relux_max_limit); #endif const int out_x_base = mul24(out_ch_blk, out_width); diff --git a/mace/kernels/opencl/cl/depthwise_conv2d.cl b/mace/kernels/opencl/cl/depthwise_conv2d.cl index 5ba07d73..792c0934 100644 --- a/mace/kernels/opencl/cl/depthwise_conv2d.cl +++ b/mace/kernels/opencl/cl/depthwise_conv2d.cl @@ -8,7 +8,6 @@ __kernel void depthwise_conv2d(__read_only image2d_t input, /* [c%4 * w * c/4, h #endif __write_only image2d_t output, __private const float relux_max_limit, - __private const float prelu_alpha, __private const short in_height, __private const short in_width, __private const short in_ch_blks, @@ -103,11 +102,11 @@ __kernel void depthwise_conv2d(__read_only image2d_t input, /* [c%4 * w * c/4, h in_hb_idx += dilation_h; } -#if defined(USE_RELU) || defined(USE_RELUX) || defined(USE_PRELU) || defined(USE_TANH) || defined(USE_SIGMOID) - out0 = do_activation(out0, relux_max_limit, prelu_alpha); - out1 = do_activation(out1, relux_max_limit, prelu_alpha); - out2 = do_activation(out2, relux_max_limit, prelu_alpha); - out3 = do_activation(out3, relux_max_limit, prelu_alpha); +#if defined(USE_RELU) || defined(USE_RELUX) || defined(USE_TANH) || defined(USE_SIGMOID) + out0 = do_activation(out0, relux_max_limit); + out1 = do_activation(out1, relux_max_limit); + out2 = do_activation(out2, relux_max_limit); + out3 = do_activation(out3, relux_max_limit); #endif const short out_x_base = mul24(out_ch_blk, out_width); @@ -134,7 +133,6 @@ __kernel void depthwise_conv2d_s1(__read_only image2d_t input, /* [c%4 * w * c/4 #endif __write_only image2d_t output, __private const DATA_TYPE relux_max_limit, - __private const DATA_TYPE prelu_alpha, __private const short in_height, __private const short in_width, __private const short in_ch_blks, @@ -220,11 +218,11 @@ __kernel void depthwise_conv2d_s1(__read_only image2d_t input, /* [c%4 * w * c/4 in_hb_idx += 1; } -#if defined(USE_RELU) || defined(USE_RELUX) || defined(USE_PRELU) || defined(USE_TANH) || defined(USE_SIGMOID) - out0 = do_activation(out0, relux_max_limit, prelu_alpha); - out1 = do_activation(out1, relux_max_limit, prelu_alpha); - out2 = do_activation(out2, relux_max_limit, prelu_alpha); - out3 = do_activation(out3, relux_max_limit, prelu_alpha); +#if defined(USE_RELU) || defined(USE_RELUX) || defined(USE_TANH) || defined(USE_SIGMOID) + out0 = do_activation(out0, relux_max_limit); + out1 = do_activation(out1, relux_max_limit); + out2 = do_activation(out2, relux_max_limit); + out3 = do_activation(out3, relux_max_limit); #endif const short out_x_base = mul24(out_ch_blk, out_width); diff --git a/mace/kernels/opencl/cl/fully_connected.cl b/mace/kernels/opencl/cl/fully_connected.cl index 021012ff..89264d82 100644 --- a/mace/kernels/opencl/cl/fully_connected.cl +++ b/mace/kernels/opencl/cl/fully_connected.cl @@ -10,8 +10,7 @@ __kernel void fully_connected(__read_only image2d_t input, __private const int input_height, __private const int input_width, __private const int input_channel, - __private const float relux_max_limit, - __private const float prelu_alpha) { + __private const float relux_max_limit) { const int batch_idx = get_global_id(0); const int out_blk_idx = get_global_id(1); const int input_chan_blk = (input_channel + 3) >> 2; @@ -51,8 +50,8 @@ __kernel void fully_connected(__read_only image2d_t input, input_coord.y++; } -#if defined(USE_RELU) || defined(USE_RELUX) || defined(USE_PRELU) || defined(USE_TANH) || defined(USE_SIGMOID) - result = do_activation(result, relux_max_limit, prelu_alpha); +#if defined(USE_RELU) || defined(USE_RELUX) || defined(USE_TANH) || defined(USE_SIGMOID) + result = do_activation(result, relux_max_limit); #endif WRITE_IMAGET(output, (int2)(out_blk_idx, batch_idx), result); } diff --git a/mace/kernels/opencl/cl/winograd_transform.cl b/mace/kernels/opencl/cl/winograd_transform.cl index e4b31598..cbcd3b19 100644 --- a/mace/kernels/opencl/cl/winograd_transform.cl +++ b/mace/kernels/opencl/cl/winograd_transform.cl @@ -115,8 +115,7 @@ __kernel void winograd_inverse_transform_2x2(__read_only image2d_t input, __private const int out_width, __private const int round_hw, __private const int round_w, - __private const float relux_max_limit, - __private const float prelu_alpha) { + __private const float relux_max_limit) { const int width_idx = get_global_id(0); const int height_idx = get_global_id(1); const int out_channel = get_global_size(1); @@ -183,11 +182,11 @@ __kernel void winograd_inverse_transform_2x2(__read_only image2d_t input, #endif -#if defined(USE_RELU) || defined(USE_RELUX) || defined(USE_PRELU) || defined(USE_TANH) || defined(USE_SIGMOID) - in0[0] = do_activation(in0[0], relux_max_limit, prelu_alpha); - in0[1] = do_activation(in0[1], relux_max_limit, prelu_alpha); - in1[0] = do_activation(in1[0], relux_max_limit, prelu_alpha); - in1[1] = do_activation(in1[1], relux_max_limit, prelu_alpha); +#if defined(USE_RELU) || defined(USE_RELUX) || defined(USE_TANH) || defined(USE_SIGMOID) + in0[0] = do_activation(in0[0], relux_max_limit); + in0[1] = do_activation(in0[1], relux_max_limit); + in1[0] = do_activation(in1[0], relux_max_limit); + in1[1] = do_activation(in1[1], relux_max_limit); #endif WRITE_IMAGET(output, (int2)(coord_x, coord_y), in0[0]); @@ -205,6 +204,4 @@ __kernel void winograd_inverse_transform_2x2(__read_only image2d_t input, WRITE_IMAGET(output, (int2)(coord_x + 1, coord_y + 1), in1[1]); } - - } diff --git a/mace/kernels/opencl/conv_2d_opencl.cc b/mace/kernels/opencl/conv_2d_opencl.cc index 58c833cb..94aa7838 100644 --- a/mace/kernels/opencl/conv_2d_opencl.cc +++ b/mace/kernels/opencl/conv_2d_opencl.cc @@ -17,7 +17,6 @@ extern void Conv2dOpenclK1x1(cl::Kernel *kernel, const int *dilations, const ActivationType activation, const float relux_max_limit, - const float prelu_alpha, const DataType dt, Tensor *output, StatsFuture *future); @@ -31,7 +30,6 @@ extern void Conv2dOpenclK3x3(cl::Kernel *kernel, const int *dilations, const ActivationType activation, const float relux_max_limit, - const float prelu_alpha, const DataType dt, Tensor *output, StatsFuture *future); @@ -45,7 +43,6 @@ extern void Conv2dOpencl(cl::Kernel *kernel, const int *dilations, const ActivationType activation, const float relux_max_limit, - const float prelu_alpha, const DataType dt, Tensor *output, StatsFuture *future); @@ -60,7 +57,7 @@ void Conv2dFunctor::operator()(const Tensor *input, cl::Kernel *kernel, const Tensor *input, const Tensor *filter, const Tensor *bias, const int stride, const int *padding, const int *dilations, const ActivationType activation, - const float relux_max_limit, const float prelu_alpha, const DataType dt, + const float relux_max_limit, const DataType dt, Tensor *output, StatsFuture *future); // Selection matrix: kernel_size x stride_size static const Conv2dOpenclFunction selector[5] = @@ -99,12 +96,10 @@ void Conv2dFunctor::operator()(const Tensor *input, selector[kernel_h - 1] != nullptr) { auto conv2d_func = selector[kernel_h - 1]; conv2d_func(&kernel_, input, filter, bias, strides_[0], paddings.data(), dilations_, activation_, - relux_max_limit_, prelu_alpha_, DataTypeToEnum::value, - output, future); + relux_max_limit_, DataTypeToEnum::value, output, future); } else { Conv2dOpencl(&kernel_, input, filter, bias, strides_[0], paddings.data(), dilations_, - activation_, relux_max_limit_, prelu_alpha_, - DataTypeToEnum::value, output, future); + activation_, relux_max_limit_, DataTypeToEnum::value, output, future); } } diff --git a/mace/kernels/opencl/conv_2d_opencl_1x1.cc b/mace/kernels/opencl/conv_2d_opencl_1x1.cc index a8dff646..bee0e12a 100644 --- a/mace/kernels/opencl/conv_2d_opencl_1x1.cc +++ b/mace/kernels/opencl/conv_2d_opencl_1x1.cc @@ -19,7 +19,6 @@ extern void Conv2dOpenclK1x1(cl::Kernel *kernel, const int *dilations, const ActivationType activation, const float relux_max_limit, - const float prelu_alpha, const DataType dt, Tensor *output, StatsFuture *future) { @@ -56,9 +55,6 @@ extern void Conv2dOpenclK1x1(cl::Kernel *kernel, case RELUX: built_options.emplace("-DUSE_RELUX"); break; - case PRELU: - built_options.emplace("-DUSE_PRELU"); - break; case TANH: built_options.emplace("-DUSE_TANH"); break; @@ -86,7 +82,6 @@ extern void Conv2dOpenclK1x1(cl::Kernel *kernel, *(static_cast(output->buffer()))); // FIXME handle flexable data type: half not supported kernel->setArg(idx++, relux_max_limit); - kernel->setArg(idx++, prelu_alpha); kernel->setArg(idx++, static_cast(input_height)); kernel->setArg(idx++, static_cast(input_width)); kernel->setArg(idx++, static_cast(input_channel_blocks)); diff --git a/mace/kernels/opencl/conv_2d_opencl_3x3.cc b/mace/kernels/opencl/conv_2d_opencl_3x3.cc index af2a2bc2..bb677177 100644 --- a/mace/kernels/opencl/conv_2d_opencl_3x3.cc +++ b/mace/kernels/opencl/conv_2d_opencl_3x3.cc @@ -21,7 +21,6 @@ extern void Conv2dOpenclK3x3(cl::Kernel *kernel, const int *dilations, const ActivationType activation, const float relux_max_limit, - const float prelu_alpha, const DataType dt, Tensor *output, StatsFuture *future) { @@ -51,9 +50,6 @@ extern void Conv2dOpenclK3x3(cl::Kernel *kernel, case RELUX: built_options.emplace("-DUSE_RELUX"); break; - case PRELU: - built_options.emplace("-DUSE_PRELU"); - break; case TANH: built_options.emplace("-DUSE_TANH"); break; @@ -80,7 +76,6 @@ extern void Conv2dOpenclK3x3(cl::Kernel *kernel, kernel->setArg(idx++, *(static_cast(output->buffer()))); kernel->setArg(idx++, relux_max_limit); - kernel->setArg(idx++, prelu_alpha); kernel->setArg(idx++, static_cast(input->dim(1))); kernel->setArg(idx++, static_cast(input->dim(2))); kernel->setArg(idx++, static_cast(input_channel_blocks)); diff --git a/mace/kernels/opencl/conv_2d_opencl_general.cc b/mace/kernels/opencl/conv_2d_opencl_general.cc index f9dd6d1d..af344c28 100644 --- a/mace/kernels/opencl/conv_2d_opencl_general.cc +++ b/mace/kernels/opencl/conv_2d_opencl_general.cc @@ -21,7 +21,6 @@ extern void Conv2dOpencl(cl::Kernel *kernel, const int *dilations, const ActivationType activation, const float relux_max_limit, - const float prelu_alpha, const DataType dt, Tensor *output, StatsFuture *future) { @@ -51,9 +50,6 @@ extern void Conv2dOpencl(cl::Kernel *kernel, case RELUX: built_options.emplace("-DUSE_RELUX"); break; - case PRELU: - built_options.emplace("-DUSE_PRELU"); - break; case TANH: built_options.emplace("-DUSE_TANH"); break; @@ -80,7 +76,6 @@ extern void Conv2dOpencl(cl::Kernel *kernel, kernel->setArg(idx++, *(static_cast(output->buffer()))); kernel->setArg(idx++, relux_max_limit); - kernel->setArg(idx++, prelu_alpha); kernel->setArg(idx++, static_cast(input->dim(1))); kernel->setArg(idx++, static_cast(input->dim(2))); kernel->setArg(idx++, static_cast(input_channel_blocks)); diff --git a/mace/kernels/opencl/depthwise_conv_opencl.cc b/mace/kernels/opencl/depthwise_conv_opencl.cc index 063220e4..2942c5d0 100644 --- a/mace/kernels/opencl/depthwise_conv_opencl.cc +++ b/mace/kernels/opencl/depthwise_conv_opencl.cc @@ -20,7 +20,6 @@ void DepthwiseConv2d(cl::Kernel *kernel, const int *dilations, const ActivationType activation, const float relux_max_limit, - const float prelu_alpha, const DataType dt, Tensor *output, StatsFuture *future) { @@ -69,9 +68,6 @@ void DepthwiseConv2d(cl::Kernel *kernel, case RELUX: built_options.emplace("-DUSE_RELUX"); break; - case PRELU: - built_options.emplace("-DUSE_PRELU"); - break; case TANH: built_options.emplace("-DUSE_TANH"); break; @@ -96,7 +92,6 @@ void DepthwiseConv2d(cl::Kernel *kernel, kernel->setArg( idx++, *(static_cast(output->buffer()))); kernel->setArg(idx++, relux_max_limit); - kernel->setArg(idx++, prelu_alpha); kernel->setArg(idx++, static_cast(input_height)); kernel->setArg(idx++, static_cast(input_width)); kernel->setArg(idx++, static_cast(input_channel_blocks)); @@ -140,8 +135,8 @@ void DepthwiseConv2dFunctor::operator()( << " is not implemented yet, using slow version"; // TODO(heliangliang) The CPU/NEON kernel should map the buffer DepthwiseConv2dFunctor( - strides_, padding_type_, paddings_, dilations_, activation_, relux_max_limit_, - prelu_alpha_)(input, filter, bias, output, future); + strides_, padding_type_, paddings_, dilations_, activation_, + relux_max_limit_)(input, filter, bias, output, future); return; } @@ -169,7 +164,7 @@ void DepthwiseConv2dFunctor::operator()( output->ResizeImage(output_shape, output_image_shape); DepthwiseConv2d(&kernel_, input, filter, bias, strides_[0], paddings.data(), dilations_, - activation_, relux_max_limit_, prelu_alpha_, + activation_, relux_max_limit_, DataTypeToEnum::value, output, future); } diff --git a/mace/kernels/opencl/fully_connected_opencl.cc b/mace/kernels/opencl/fully_connected_opencl.cc index f719bb01..d0f4c1d1 100644 --- a/mace/kernels/opencl/fully_connected_opencl.cc +++ b/mace/kernels/opencl/fully_connected_opencl.cc @@ -48,9 +48,6 @@ void FullyConnectedFunctor::operator()( case RELUX: built_options.emplace("-DUSE_RELUX"); break; - case PRELU: - built_options.emplace("-DUSE_PRELU"); - break; case TANH: built_options.emplace("-DUSE_TANH"); break; @@ -78,7 +75,6 @@ void FullyConnectedFunctor::operator()( kernel_.setArg(idx++, static_cast(input->dim(3))); // FIXME handle flexable data type: half not supported kernel_.setArg(idx++, relux_max_limit_); - kernel_.setArg(idx++, prelu_alpha_); } const uint32_t gws[2] = { diff --git a/mace/kernels/opencl/winograd_transform.cc b/mace/kernels/opencl/winograd_transform.cc index ab8704a1..54511220 100644 --- a/mace/kernels/opencl/winograd_transform.cc +++ b/mace/kernels/opencl/winograd_transform.cc @@ -129,7 +129,6 @@ void WinogradInverseTransformFunctor::operator()(const Te kernel_.setArg(idx++, static_cast(round_h * round_w)); kernel_.setArg(idx++, static_cast(round_w)); kernel_.setArg(idx++, relux_max_limit_); - kernel_.setArg(idx++, prelu_alpha_); } const uint32_t gws[2] = {static_cast(input_tensor->dim(2)), diff --git a/mace/kernels/winograd_transform.h b/mace/kernels/winograd_transform.h index 639138ee..fdab5c7c 100644 --- a/mace/kernels/winograd_transform.h +++ b/mace/kernels/winograd_transform.h @@ -58,21 +58,18 @@ struct WinogradInverseTransformFunctorBase { const int height, const int width, const ActivationType activation, - const float relux_max_limit, - const float prelu_alpha) + const float relux_max_limit) : batch_(batch), height_(height), width_(width), activation_(activation), - relux_max_limit_(relux_max_limit), - prelu_alpha_(prelu_alpha) {} + relux_max_limit_(relux_max_limit) {} const int batch_; const int height_; const int width_; const ActivationType activation_; const float relux_max_limit_; - const float prelu_alpha_; }; template @@ -81,9 +78,8 @@ struct WinogradInverseTransformFunctor : WinogradInverseTransformFunctorBase { const int height, const int width, const ActivationType activation, - const float relux_max_limit, - const float prelu_alpha) - : WinogradInverseTransformFunctorBase(batch, height, width, activation, relux_max_limit, prelu_alpha) {} + const float relux_max_limit) + : WinogradInverseTransformFunctorBase(batch, height, width, activation, relux_max_limit) {} void operator()(const Tensor *input, const Tensor *bias, @@ -100,9 +96,8 @@ struct WinogradInverseTransformFunctor : WinogradInverseT const int height, const int width, const ActivationType activation, - const float relux_max_limit, - const float prelu_alpha) - : WinogradInverseTransformFunctorBase(batch, height, width, activation, relux_max_limit, prelu_alpha) {} + const float relux_max_limit) + : WinogradInverseTransformFunctorBase(batch, height, width, activation, relux_max_limit) {} void operator()(const Tensor *input, const Tensor *bias, diff --git a/mace/ops/activation.h b/mace/ops/activation.h index 04ca0249..a55dfe1a 100644 --- a/mace/ops/activation.h +++ b/mace/ops/activation.h @@ -18,15 +18,15 @@ class ActivationOp : public Operator { functor_(kernels::StringToActivationType( OperatorBase::GetSingleArgument("activation", "NOOP")), - OperatorBase::GetSingleArgument("max_limit", 0.0f), - OperatorBase::GetSingleArgument("alpha", 0.0f)) {} + OperatorBase::GetSingleArgument("max_limit", 0.0f)) {} bool Run(StatsFuture *future) override { - const Tensor *input_tensor = this->inputs_[0]; + const Tensor *input_tensor = this->Input(0); + const Tensor *alpha_tensor = this->InputSize() >= 2 ? this->Input(1) : nullptr; Tensor *output_tensor = this->outputs_[0]; output_tensor->ResizeLike(input_tensor); - functor_(input_tensor, output_tensor, future); + functor_(input_tensor, alpha_tensor, output_tensor, future); return true; } diff --git a/mace/ops/activation_test.cc b/mace/ops/activation_test.cc index 02e16108..ce5ddd45 100644 --- a/mace/ops/activation_test.cc +++ b/mace/ops/activation_test.cc @@ -213,17 +213,22 @@ void TestSimplePrelu() { // Add input data net.AddInputFromArray( "Input", {2, 2, 2, 2}, - {-7, 7, -6, 6, -5, 5, -4, 4, -3, 3, -2, 2, -1, 1, 0, 0}); + {-7, 7, -6, 6, -5, -5, -4, -4, -3, 3, -2, 2, -1, -1, 0, 0}); + net.AddInputFromArray( + "Alpha", {2}, + {2.0, 3.0}); if (D == DeviceType::OPENCL) { BufferToImage(net, "Input", "InputImage", kernels::BufferType::IN_OUT_CHANNEL); + BufferToImage(net, "Alpha", "AlphaImage", + kernels::BufferType::ARGUMENT); OpDefBuilder("Activation", "PreluTest") .Input("InputImage") + .Input("AlphaImage") .Output("OutputImage") .AddStringArg("activation", "PRELU") - .AddFloatArg("alpha", 2.0) .Finalize(net.NewOperatorDef()); // Run @@ -235,9 +240,9 @@ void TestSimplePrelu() { } else { OpDefBuilder("Activation", "PreluTest") .Input("Input") + .Input("Alpha") .Output("Output") .AddStringArg("activation", "PRELU") - .AddFloatArg("alpha", 2.0) .Finalize(net.NewOperatorDef()); // Run @@ -245,7 +250,7 @@ void TestSimplePrelu() { } auto expected = CreateTensor( - {2, 2, 2, 2}, {-14, 7, -12, 6, -10, 5, -8, 4, -6, 3, -4, 2, -2, 1, 0, 0}); + {2, 2, 2, 2}, {-14, 7, -12, 6, -10, -15, -8, -12, -6, 3, -4, 2, -2, -3, 0, 0}); ExpectTensorNear(*expected, *net.GetOutput("Output"), 1e-5); } diff --git a/mace/ops/batch_norm.h b/mace/ops/batch_norm.h index a3a0c136..eb2bfd25 100644 --- a/mace/ops/batch_norm.h +++ b/mace/ops/batch_norm.h @@ -16,7 +16,7 @@ class BatchNormOp : public Operator { public: BatchNormOp(const OperatorDef &operator_def, Workspace *ws) : Operator(operator_def, ws), - functor_(false, kernels::ActivationType::NOOP, 0.0f, 0.0f) { + functor_(false, kernels::ActivationType::NOOP, 0.0f) { epsilon_ = OperatorBase::GetSingleArgument("epsilon", static_cast(1e-4)); } diff --git a/mace/ops/conv_2d.h b/mace/ops/conv_2d.h index b7347b5e..c441b0b4 100644 --- a/mace/ops/conv_2d.h +++ b/mace/ops/conv_2d.h @@ -23,7 +23,6 @@ class Conv2dOp : public ConvPool2dOpBase { this->paddings_, this->dilations_.data(), kernels::ActivationType::NOOP, - 0.0f, 0.0f) {} bool Run(StatsFuture *future) override { diff --git a/mace/ops/conv_2d_test.cc b/mace/ops/conv_2d_test.cc index 8beba453..184772c4 100644 --- a/mace/ops/conv_2d_test.cc +++ b/mace/ops/conv_2d_test.cc @@ -412,9 +412,9 @@ static void TestComplexConvNxNS12(const std::vector &shape, const int s } TEST_F(Conv2dOpTest, OPENCLAlignedConvNxNS12) { - TestComplexConvNxNS12({32, 16, 16, 32}, + TestComplexConvNxNS12({32, 16, 16, 32}, 1); - TestComplexConvNxNS12({32, 16, 16, 32}, + TestComplexConvNxNS12({32, 16, 16, 32}, 2); } diff --git a/mace/ops/depthwise_conv2d.h b/mace/ops/depthwise_conv2d.h index 82d6e899..0678ba07 100644 --- a/mace/ops/depthwise_conv2d.h +++ b/mace/ops/depthwise_conv2d.h @@ -26,8 +26,7 @@ class DepthwiseConv2dOp : public ConvPool2dOpBase { kernels::StringToActivationType( OperatorBase::GetSingleArgument("activation", "NOOP")), - OperatorBase::GetSingleArgument("max_limit", 0.0f), - OperatorBase::GetSingleArgument("alpha", 0.0f)) {} + OperatorBase::GetSingleArgument("max_limit", 0.0f)) {} bool Run(StatsFuture *future) override { const Tensor *input = this->Input(INPUT); diff --git a/mace/ops/folded_batch_norm.h b/mace/ops/folded_batch_norm.h index 7e7c7e5d..3420d13e 100644 --- a/mace/ops/folded_batch_norm.h +++ b/mace/ops/folded_batch_norm.h @@ -19,8 +19,7 @@ class FoldedBatchNormOp : public Operator { kernels::StringToActivationType( OperatorBase::GetSingleArgument("activation", "NOOP")), - OperatorBase::GetSingleArgument("max_limit", 0.0f), - OperatorBase::GetSingleArgument("alpha", 0.0f)) {} + OperatorBase::GetSingleArgument("max_limit", 0.0f)) {} bool Run(StatsFuture *future) override { const Tensor *input = this->Input(INPUT); diff --git a/mace/ops/fully_connected.h b/mace/ops/fully_connected.h index 0ee90e2b..c65947af 100644 --- a/mace/ops/fully_connected.h +++ b/mace/ops/fully_connected.h @@ -19,8 +19,7 @@ class FullyConnectedOp : public Operator { kernels::StringToActivationType( OperatorBase::GetSingleArgument("activation", "NOOP")), - OperatorBase::GetSingleArgument("max_limit", 0.0f), - OperatorBase::GetSingleArgument("alpha", 0.0f)) {} + OperatorBase::GetSingleArgument("max_limit", 0.0f)) {} bool Run(StatsFuture *future) override { const Tensor *input = this->Input(INPUT); diff --git a/mace/ops/fused_conv_2d.h b/mace/ops/fused_conv_2d.h index 5c5957f1..0184f92f 100644 --- a/mace/ops/fused_conv_2d.h +++ b/mace/ops/fused_conv_2d.h @@ -25,8 +25,7 @@ class FusedConv2dOp : public ConvPool2dOpBase { kernels::StringToActivationType( OperatorBase::GetSingleArgument("activation", "NOOP")), - OperatorBase::GetSingleArgument("max_limit", 0.0f), - OperatorBase::GetSingleArgument("alpha", 0.0f)) {} + OperatorBase::GetSingleArgument("max_limit", 0.0f)) {} bool Run(StatsFuture *future) override { const Tensor *input = this->Input(INPUT); diff --git a/mace/ops/winograd_inverse_transform.h b/mace/ops/winograd_inverse_transform.h index 4c20769f..aef37473 100644 --- a/mace/ops/winograd_inverse_transform.h +++ b/mace/ops/winograd_inverse_transform.h @@ -24,8 +24,7 @@ class WinogradInverseTransformOp : public Operator { kernels::StringToActivationType( OperatorBase::GetSingleArgument("activation", "NOOP")), - OperatorBase::GetSingleArgument("max_limit", 0.0f), - OperatorBase::GetSingleArgument("alpha", 0.0f)) {} + OperatorBase::GetSingleArgument("max_limit", 0.0f)) {} bool Run(StatsFuture *future) override { const Tensor *input_tensor = this->Input(INPUT); -- GitLab