From 1f6919fd40de3bf151e45554053a9b0cc07019ea Mon Sep 17 00:00:00 2001 From: Hao Zeng Date: Sun, 18 Aug 2019 18:37:57 +0800 Subject: [PATCH] Implemented hard swish layer --- src/CMakeLists.txt | 1 + src/layer/arm/hardswish_arm.cpp | 75 +++++++++++++++++ src/layer/arm/hardswish_arm.h | 30 +++++++ src/layer/hardswish.cpp | 63 +++++++++++++++ src/layer/hardswish.h | 37 +++++++++ src/layer/vulkan/hardswish_vulkan.cpp | 84 ++++++++++++++++++++ src/layer/vulkan/hardswish_vulkan.h | 39 +++++++++ src/layer/vulkan/shader/hardswish.comp | 58 ++++++++++++++ src/layer/vulkan/shader/hardswish_pack4.comp | 58 ++++++++++++++ 9 files changed, 445 insertions(+) create mode 100644 src/layer/arm/hardswish_arm.cpp create mode 100644 src/layer/arm/hardswish_arm.h create mode 100644 src/layer/hardswish.cpp create mode 100644 src/layer/hardswish.h create mode 100644 src/layer/vulkan/hardswish_vulkan.cpp create mode 100644 src/layer/vulkan/hardswish_vulkan.h create mode 100644 src/layer/vulkan/shader/hardswish.comp create mode 100644 src/layer/vulkan/shader/hardswish_pack4.comp diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 672db143d..1aa23e955 100644 --- a/src/CMakeLists.txt +++ b/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}) diff --git a/src/layer/arm/hardswish_arm.cpp b/src/layer/arm/hardswish_arm.cpp new file mode 100644 index 000000000..1f6b78f7a --- /dev/null +++ b/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 +#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> 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 diff --git a/src/layer/arm/hardswish_arm.h b/src/layer/arm/hardswish_arm.h new file mode 100644 index 000000000..1743d0952 --- /dev/null +++ b/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 diff --git a/src/layer/hardswish.cpp b/src/layer/hardswish.cpp new file mode 100644 index 000000000..a19f8a163 --- /dev/null +++ b/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 + +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 upper) ; + else + ptr[i] = ptr[i] * (ptr[i] * alpha + beta); + } + } + + return 0; +} + +} // namespace ncnn diff --git a/src/layer/hardswish.h b/src/layer/hardswish.h new file mode 100644 index 000000000..021db2f8d --- /dev/null +++ b/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 diff --git a/src/layer/vulkan/hardswish_vulkan.cpp b/src/layer/vulkan/hardswish_vulkan.cpp new file mode 100644 index 000000000..9bafee016 --- /dev/null +++ b/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 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 bindings(1); + bindings[0] = bottom_top_blob; + + std::vector 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 diff --git a/src/layer/vulkan/hardswish_vulkan.h b/src/layer/vulkan/hardswish_vulkan.h new file mode 100644 index 000000000..be2d53801 --- /dev/null +++ b/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 diff --git a/src/layer/vulkan/shader/hardswish.comp b/src/layer/vulkan/shader/hardswish.comp new file mode 100644 index 000000000..a2143fe3e --- /dev/null +++ b/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); +} diff --git a/src/layer/vulkan/shader/hardswish_pack4.comp b/src/layer/vulkan/shader/hardswish_pack4.comp new file mode 100644 index 000000000..52d8f9baa --- /dev/null +++ b/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); +}