diff --git a/mace/kernels/activation.h b/mace/kernels/activation.h index 83acf4fb70311c710bcde1d7b08c1e6c4630879f..b768eb28c93a52502af8dd51f0e6aa13a9a145a8 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 bf5035ded2d67a0469347a0a001e6bc66f1d437f..0d489f40af6398240772cbabd8ff99f5b72be96c 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 91e1471ff662aa2e82e46a5103aa5d5214cbca77..e6c22cd9f927364c80733f7b98f46535582eaaf7 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 ef78e17ee53e328372e7c92cfbb6b8f14edb2ecc..141119a9ba8a02002d5e37db46d74654973add11 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 c3b662222a58d8dea1180195e2bdb8aae0c56eca..7d072016f284a839895ab408cb51852da54bb389 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 934ce8da666b8e9910815b9f56de250a3511fd68..dee010875e853ba192cf05b90abfcdf5ee2cb48f 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 d88fed51cee468073dcf5521c40a8187f892dd5c..7696e875e538b0f06aefb4e30c4032fcba56a538 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 fe8619e20b7f567bc36dfa4bc5d6c53bd5f792fb..bee0b0e35313b4129fe6741cd9575f88b60e1431 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 99c00fabe4d7a988d8388c0d0fa61d45301baf09..773b59c44e0021ab68a4d621514056d5327b5427 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 13b20e05c07e5ebddb78bc94c8decd7694bbbc25..28b9addd3dfe67f2caac81be7890ecd6624f0e90 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 35e17da8ff7b3337fb7cc30f73dcd92a793fc647..11253d69f02485c9c712c7abd7e56dbae2ad9414 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 ebcbb23177924555b76b89b12111f45cb0f156f3..b695165e1c3398ad333f2e52f307cd91e3eb4f59 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 0fce82366e16bdbffab48b73c0038cba02dd48b8..fad561aaca4aa8f6fe862f314177221214264053 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 5ba07d73dfe0c73617879a884a8a7310536cca32..792c0934a4f7af5774b3065ecd349300a1f18854 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 021012ffe39eb5f9a3744267bdb088001123f1aa..89264d82b890ca0effd9de53ea935a8821506f19 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 e4b315984ed040d5ea088e22237665b14abe4271..cbcd3b193a92e8e135a55014ad5e62b5545ed57e 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 58c833cb206ff110854f34090a08db0c7bb23bdb..94aa783886c8eab5214f5a1edbf75ec6b42c6501 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 a8dff64699ea3952d53b7026c6b1a66dc659ccda..bee0e12a8826c2cd4d7bbe28ba3e70c9fe42f259 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 af2a2bc29d74196a7788c76f3c317c6a9ab897ff..bb67717791fd05f820ad92f734af545fe2b99e1e 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 f9dd6d1db5010a8575492ae93357ff9e4c5ce062..af344c284fe04836d1d2ac23b4014ffdf76ac22b 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 063220e44404a8e85e056043b6ed8a34a030a9e6..2942c5d060b9a240c0e9c3aa47cf6e2a82a6fdfd 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 f719bb0100e763514360d4bed837ba12e2ee9e0d..d0f4c1d1257c123f1a19b545e0277b24acf34a4c 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 ab8704a19e77eb8cedaa9ed8989dbeae90f62d66..54511220fdc4ce1cec32f8e2a38f0fbf38b35519 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 639138ee31485c540eb0b066fdf7ec9bdc8fc662..fdab5c7c8d42ee815ca69b37fab34775a60047e2 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 04ca0249c18a19fb2c47529d97535f4cd8663073..a55dfe1a4dea9ae7bff92c43ee0133def76d7c64 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 02e16108eaacc8f91608457717bfeb7e55260dac..ce5ddd4598d154b760849bee078d944973c60ac8 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 a3a0c136ba12105fa121dbe18372c1d41112e75e..eb2bfd25cdc1af2ea7c34e236f71c9384055f464 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 b7347b5eeb1f222c9fb04f127c1cf68669b058cb..c441b0b45b0b00619fe0a36554dcef307227b2d9 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 8beba453cc20b421581d16270be9dedeb82b6ddb..184772c47c7a7183f40db324b6a519fa4d820e34 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 82d6e899649d1a04192725fe62a56c8e776ee8f3..0678ba0757d17a50fa1b2ef2cdcc5c4b2635a46d 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 7e7c7e5d855ffcaf4a7fea7b5feb332fc8695523..3420d13e5814052b1902aee4dc98d9ecfce35129 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 0ee90e2b3fd86d20d77d0f20d10b75c214cbb46e..c65947af7148d1ace34e9fe8b899a1044b3c3265 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 5c5957f1c1975212f8e51d9cddd82cf36171cf30..0184f92f80700fd686e20a8c5918c3ce3c6ecb28 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 4c20769f1fd461f393c1c57e58bc5f089197ed7c..aef374731408885aa19f4c88f5ffc15fe389bd2c 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);