Implemented hard swish layer

src/CMakeLists.txt

@@ -261,6 +261,7 @@ ncnn_add_layer(Requantize)
 ncnn_add_layer(Cast)
 ncnn_add_layer(HardSigmoid)
 ncnn_add_layer(SELU)
+ncnn_add_layer(HardSwish)
 
 add_custom_target(generate-spirv DEPENDS ${SHADER_SPV_HEX_FILES})
 

src/layer/arm/hardswish_arm.cpp

@@ -0,0 +1,75 @@
+// Tencent is pleased to support the open source community by making ncnn available.
+//
+// Copyright (C) 2017 THL A29 Limited, a Tencent company. All rights reserved.
+//
+// Licensed under the BSD 3-Clause License (the "License"); you may not use this file except
+// in compliance with the License. You may obtain a copy of the License at
+//
+// https://opensource.org/licenses/BSD-3-Clause
+//
+// Unless required by applicable law or agreed to in writing, software distributed
+// under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR
+// CONDITIONS OF ANY KIND, either express or implied. See the License for the
+// specific language governing permissions and limitations under the License.
+
+#include "hardswish_arm.h"
+
+#if __ARM_NEON
+#include <arm_neon.h>
+#endif // __ARM_NEON
+
+namespace ncnn {
+
+DEFINE_LAYER_CREATOR(HardSwish_arm)
+
+int HardSwish_arm::forward_inplace(Mat& bottom_top_blob, const Option& opt) const
+{
+    int w = bottom_top_blob.w;
+    int h = bottom_top_blob.h;
+    int channels = bottom_top_blob.c;
+    int size = w * h;
+
+    #pragma omp parallel for num_threads(opt.num_threads)
+    for (int q=0; q<channels; q++)
+    {
+        float* ptr = bottom_top_blob.channel(q);
+
+#if __ARM_NEON
+        int nn = size >> 2;
+        int remain = size - (nn << 2);
+#else
+        int remain = size;
+#endif // __ARM_NEON
+
+#if __ARM_NEON
+        float32x4_t _zero = vdupq_n_f32(0.f);
+        float32x4_t _one = vdupq_n_f32(1.f);
+        while (nn--)
+        {
+            float32x4_t _p = vld1q_f32(ptr);
+            float32x4_t _ans = vdupq_n_f32(beta);
+            _ans = vmlaq_n_f32(_ans, _p, alpha);
+            _ans = vmaxq_f32(_ans, _zero);
+            _ans = vminq_f32(_ans, _one);
+            _ans = vmulq_f32(_ans, _p);
+            vst1q_f32(ptr, _ans);
+
+            ptr += 4;
+        }
+#endif // __ARM_NEON
+        for (; remain>0; remain--)
+        {
+            if (*ptr < lower)
+                *ptr = 0.f;
+            else if (*ptr > upper)
+                ;
+            else
+                *ptr = *ptr * (*ptr * alpha + beta);
+            ++ptr;
+        }
+    }
+
+    return 0;
+}
+
+} // namespace ncnn

src/layer/arm/hardswish_arm.h

@@ -0,0 +1,30 @@
+// Tencent is pleased to support the open source community by making ncnn available.
+//
+// Copyright (C) 2017 THL A29 Limited, a Tencent company. All rights reserved.
+//
+// Licensed under the BSD 3-Clause License (the "License"); you may not use this file except
+// in compliance with the License. You may obtain a copy of the License at
+//
+// https://opensource.org/licenses/BSD-3-Clause
+//
+// Unless required by applicable law or agreed to in writing, software distributed
+// under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR
+// CONDITIONS OF ANY KIND, either express or implied. See the License for the
+// specific language governing permissions and limitations under the License.
+
+#ifndef LAYER_HARDSWISH_ARM_H
+#define LAYER_HARDSWISH_ARM_H
+
+#include "hardswish.h"
+
+namespace ncnn {
+
+class HardSwish_arm : virtual public HardSwish
+{
+public:
+    virtual int forward_inplace(Mat& bottom_top_blob, const Option& opt) const;
+};
+
+} // namespace ncnn
+
+#endif // LAYER_HARDSWISH_ARM_H

src/layer/hardswish.cpp

@@ -0,0 +1,63 @@
+// Tencent is pleased to support the open source community by making ncnn available.
+//
+// Copyright (C) 2017 THL A29 Limited, a Tencent company. All rights reserved.
+//
+// Licensed under the BSD 3-Clause License (the "License"); you may not use this file except
+// in compliance with the License. You may obtain a copy of the License at
+//
+// https://opensource.org/licenses/BSD-3-Clause
+//
+// Unless required by applicable law or agreed to in writing, software distributed
+// under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR
+// CONDITIONS OF ANY KIND, either express or implied. See the License for the
+// specific language governing permissions and limitations under the License.
+
+#include "hardswish.h"
+#include <algorithm>
+
+namespace ncnn {
+
+DEFINE_LAYER_CREATOR(HardSwish)
+
+HardSwish::HardSwish()
+{
+    one_blob_only = true;
+    support_inplace = true;
+}
+
+int HardSwish::load_param(const ParamDict& pd)
+{
+    alpha = pd.get(0, 0.2f);
+    beta = pd.get(1, 0.5f);
+    lower = -beta / alpha;
+    upper = (1.f / alpha) + lower;
+
+    return 0;
+}
+
+int HardSwish::forward_inplace(Mat& bottom_top_blob, const Option& opt) const
+{
+    int w = bottom_top_blob.w;
+    int h = bottom_top_blob.h;
+    int channels = bottom_top_blob.c;
+    int size = w * h;
+
+    #pragma omp parallel for num_threads(opt.num_threads)
+    for (int q=0; q<channels; q++)
+    {
+        float* ptr = bottom_top_blob.channel(q);
+
+        for (int i=0; i<size; i++)
+        {
+            if (ptr[i] < lower)
+                ptr[i] = 0.f;
+            else if (ptr[i] > upper) ;
+            else
+                ptr[i] = ptr[i] * (ptr[i] * alpha + beta);
+        }
+    }
+
+    return 0;
+}
+
+} // namespace ncnn

src/layer/hardswish.h

@@ -0,0 +1,37 @@
+// Tencent is pleased to support the open source community by making ncnn available.
+//
+// Copyright (C) 2017 THL A29 Limited, a Tencent company. All rights reserved.
+//
+// Licensed under the BSD 3-Clause License (the "License"); you may not use this file except
+// in compliance with the License. You may obtain a copy of the License at
+//
+// https://opensource.org/licenses/BSD-3-Clause
+//
+// Unless required by applicable law or agreed to in writing, software distributed
+// under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR
+// CONDITIONS OF ANY KIND, either express or implied. See the License for the
+// specific language governing permissions and limitations under the License.
+
+#ifndef LAYER_HARDSWISH_H
+#define LAYER_HARDSWISH_H
+
+#include "layer.h"
+
+namespace ncnn {
+
+class HardSwish : public Layer
+{
+public:
+    HardSwish();
+
+    virtual int load_param(const ParamDict& pd);
+
+    virtual int forward_inplace(Mat& bottom_top_blob, const Option& opt) const;
+
+public:
+    float alpha, beta, lower, upper;
+};
+
+} // namespace ncnn
+
+#endif // LAYER_HARDSWISH_H

src/layer/vulkan/hardswish_vulkan.cpp

@@ -0,0 +1,84 @@
+// Tencent is pleased to support the open source community by making ncnn available.
+//
+// Copyright (C) 2019 THL A29 Limited, a Tencent company. All rights reserved.
+//
+// Licensed under the BSD 3-Clause License (the "License"); you may not use this file except
+// in compliance with the License. You may obtain a copy of the License at
+//
+// https://opensource.org/licenses/BSD-3-Clause
+//
+// Unless required by applicable law or agreed to in writing, software distributed
+// under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR
+// CONDITIONS OF ANY KIND, either express or implied. See the License for the
+// specific language governing permissions and limitations under the License.
+
+#include "hardswish_vulkan.h"
+
+namespace ncnn {
+
+DEFINE_LAYER_CREATOR(HardSwish_vulkan)
+
+HardSwish_vulkan::HardSwish_vulkan()
+{
+    support_vulkan = true;
+
+    pipeline_hardswish = 0;
+    pipeline_hardswish_pack4 = 0;
+}
+
+int HardSwish_vulkan::create_pipeline(const Option& opt)
+{
+    std::vector<vk_specialization_type> specializations(2);
+    specializations[0].f = alpha;
+    specializations[1].f = beta;
+
+    // pack1
+    {
+        pipeline_hardswish = new Pipeline(vkdev);
+        pipeline_hardswish->set_optimal_local_size_xyz();
+        pipeline_hardswish->create("hardswish", opt, specializations, 1, 5);
+    }
+
+    // pack4
+    {
+        pipeline_hardswish_pack4 = new Pipeline(vkdev);
+        pipeline_hardswish_pack4->set_optimal_local_size_xyz();
+        pipeline_hardswish_pack4->create("hardswish_pack4", opt, specializations, 1, 5);
+    }
+
+    return 0;
+}
+
+int HardSwish_vulkan::destroy_pipeline(const Option& opt)
+{
+    delete pipeline_hardswish;
+    pipeline_hardswish = 0;
+
+    delete pipeline_hardswish_pack4;
+    pipeline_hardswish_pack4 = 0;
+
+    return 0;
+}
+
+int HardSwish_vulkan::forward_inplace(VkMat& bottom_top_blob, VkCompute& cmd, const Option& opt) const
+{
+    int elempack = bottom_top_blob.elempack;
+
+    std::vector<VkMat> bindings(1);
+    bindings[0] = bottom_top_blob;
+
+    std::vector<vk_constant_type> constants(5);
+    constants[0].i = bottom_top_blob.dims;
+    constants[1].i = bottom_top_blob.w;
+    constants[2].i = bottom_top_blob.h;
+    constants[3].i = bottom_top_blob.c;
+    constants[4].i = bottom_top_blob.cstep;
+
+    const Pipeline* pipeline = elempack == 4 ? pipeline_hardswish_pack4 : pipeline_hardswish;
+
+    cmd.record_pipeline(pipeline, bindings, constants, bottom_top_blob);
+
+    return 0;
+}
+
+} // namespace ncnn

src/layer/vulkan/hardswish_vulkan.h

@@ -0,0 +1,39 @@
+// Tencent is pleased to support the open source community by making ncnn available.
+//
+// Copyright (C) 2019 THL A29 Limited, a Tencent company. All rights reserved.
+//
+// Licensed under the BSD 3-Clause License (the "License"); you may not use this file except
+// in compliance with the License. You may obtain a copy of the License at
+//
+// https://opensource.org/licenses/BSD-3-Clause
+//
+// Unless required by applicable law or agreed to in writing, software distributed
+// under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR
+// CONDITIONS OF ANY KIND, either express or implied. See the License for the
+// specific language governing permissions and limitations under the License.
+
+#ifndef LAYER_HARDSWISH_VULKAN_H
+#define LAYER_HARDSWISH_VULKAN_H
+
+#include "hardswish.h"
+
+namespace ncnn {
+
+class HardSwish_vulkan : virtual public HardSwish
+{
+public:
+    HardSwish_vulkan();
+
+    virtual int create_pipeline(const Option& opt);
+    virtual int destroy_pipeline(const Option& opt);
+
+    virtual int forward_inplace(VkMat& bottom_top_blob, VkCompute& cmd, const Option& opt) const;
+
+public:
+    Pipeline* pipeline_hardswish;
+    Pipeline* pipeline_hardswish_pack4;
+};
+
+} // namespace ncnn
+
+#endif // LAYER_HARDSWISH_VULKAN_H

src/layer/vulkan/shader/hardswish.comp

@@ -0,0 +1,58 @@
+// Tencent is pleased to support the open source community by making ncnn available.
+//
+// Copyright (C) 2018 THL A29 Limited, a Tencent company. All rights reserved.
+//
+// Licensed under the BSD 3-Clause License (the "License"); you may not use this file except
+// in compliance with the License. You may obtain a copy of the License at
+//
+// https://opensource.org/licenses/BSD-3-Clause
+//
+// Unless required by applicable law or agreed to in writing, software distributed
+// under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR
+// CONDITIONS OF ANY KIND, either express or implied. See the License for the
+// specific language governing permissions and limitations under the License.
+
+#version 450
+
+#if NCNN_fp16_storage
+#extension GL_EXT_shader_16bit_storage: require
+#endif
+#if NCNN_fp16_arithmetic
+#extension GL_AMD_gpu_shader_half_float: require
+#endif
+
+layout (constant_id = 0) const float alpha = 0;
+layout (constant_id = 1) const float beta = 0;
+
+layout (local_size_x_id = 233) in;
+layout (local_size_y_id = 234) in;
+layout (local_size_z_id = 235) in;
+
+layout (binding = 0) buffer bottom_top_blob { sfp bottom_top_blob_data[]; };
+
+layout (push_constant) uniform parameter
+{
+    int dims;
+    int w;
+    int h;
+    int c;
+    int cstep;
+} p;
+
+void main()
+{
+    int gx = int(gl_GlobalInvocationID.x);
+    int gy = int(gl_GlobalInvocationID.y);
+    int gz = int(gl_GlobalInvocationID.z);
+
+    if (gx >= p.w || gy >= p.h || gz >= p.c)
+        return;
+
+    const int gi = gz * p.cstep + gy * p.w + gx;
+
+    afp v = sfp2afp(bottom_top_blob_data[gi]);
+
+    v = v * clamp(v * afp(alpha) + afp(beta), afp(0.f), afp(1.f));
+
+    bottom_top_blob_data[gi] = afp2sfp(v);
+}

src/layer/vulkan/shader/hardswish_pack4.comp

@@ -0,0 +1,58 @@
+// Tencent is pleased to support the open source community by making ncnn available.
+//
+// Copyright (C) 2019 THL A29 Limited, a Tencent company. All rights reserved.
+//
+// Licensed under the BSD 3-Clause License (the "License"); you may not use this file except
+// in compliance with the License. You may obtain a copy of the License at
+//
+// https://opensource.org/licenses/BSD-3-Clause
+//
+// Unless required by applicable law or agreed to in writing, software distributed
+// under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR
+// CONDITIONS OF ANY KIND, either express or implied. See the License for the
+// specific language governing permissions and limitations under the License.
+
+#version 450
+
+#if NCNN_fp16_storage
+#extension GL_EXT_shader_16bit_storage: require
+#endif
+#if NCNN_fp16_arithmetic
+#extension GL_AMD_gpu_shader_half_float: require
+#endif
+
+layout (constant_id = 0) const float alpha = 0;
+layout (constant_id = 1) const float beta = 0;
+
+layout (local_size_x_id = 233) in;
+layout (local_size_y_id = 234) in;
+layout (local_size_z_id = 235) in;
+
+layout (binding = 0) buffer bottom_top_blob { sfpvec4 bottom_top_blob_data[]; };
+
+layout (push_constant) uniform parameter
+{
+    int dims;
+    int w;
+    int h;
+    int c;
+    int cstep;
+} p;
+
+void main()
+{
+    int gx = int(gl_GlobalInvocationID.x);
+    int gy = int(gl_GlobalInvocationID.y);
+    int gz = int(gl_GlobalInvocationID.z);
+
+    if (gx >= p.w || gy >= p.h || gz >= p.c)
+        return;
+
+    const int gi = gz * p.cstep + gy * p.w + gx;
+
+    afpvec4 v = sfp2afpvec4(bottom_top_blob_data[gi]);
+
+    v = v * clamp(v * afp(alpha) + afp(beta), afp(0.f), afp(1.f));
+
+    bottom_top_blob_data[gi] = afp2sfpvec4(v);
+}