diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 255f20b72c53..5e6dfc111aea 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -164,6 +164,7 @@ ncnn_add_layer(CumulativeSum) ncnn_add_layer(CopyTo) ncnn_add_layer(Erf) ncnn_add_layer(Diag) +ncnn_add_layer(CeLU) if(NCNN_VULKAN) ncnn_add_shader(${CMAKE_CURRENT_SOURCE_DIR}/convert_ycbcr.comp) diff --git a/src/layer/celu.cpp b/src/layer/celu.cpp new file mode 100644 index 000000000000..536d31542ad2 --- /dev/null +++ b/src/layer/celu.cpp @@ -0,0 +1,54 @@ +// 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 "relu.h" + +namespace ncnn { + +CeLU::CeLU() +{ + one_blob_only = true; + support_inplace = true; +} + +int CeLU::load_param(const ParamDict& pd) +{ + alpha = pd.get(0, 1.f); + + return 0; +} + +int CeLU::forward_inplace(Mat& bottom_top_blob, const Option& opt) const +{ + int w = bottom_top_blob.w; + int h = bottom_top_blob.h; + int d = bottom_top_blob.d; + int channels = bottom_top_blob.c; + int size = w * h * d; + + #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++) + { + ptr[i] = std::max(0,ptr[i]) + std::min(0,alpha*(expf(x/alpha)-1)); + } + } + + return 0; +} + +} // namespace ncnn diff --git a/src/layer/celu.h b/src/layer/celu.h new file mode 100644 index 000000000000..78cace95e8f3 --- /dev/null +++ b/src/layer/celu.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_CELU_H +#define LAYER_CELU_H + +#include "layer.h" + +namespace ncnn { + +class CeLU : public Layer +{ +public: + CeLU(); + + virtual int load_param(const ParamDict& pd); + + virtual int forward_inplace(Mat& bottom_top_blob, const Option& opt) const; + +public: + float alpha; +}; + +} // namespace ncnn + +#endif // LAYER_CELU_H diff --git a/src/layer/vulkan/celu_vulkan.cpp b/src/layer/vulkan/celu_vulkan.cpp new file mode 100644 index 000000000000..7bf3b0349294 --- /dev/null +++ b/src/layer/vulkan/celu_vulkan.cpp @@ -0,0 +1,182 @@ +// 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 "celu_vulkan.h" + +#include "layer_shader_type.h" + +namespace ncnn { + +CeLU_vulkan::CeLU_vulkan() +{ + support_vulkan = true; + support_image_storage = true; + + pipeline_celu = 0; + pipeline_celu_pack4 = 0; + pipeline_celu_pack8 = 0; +} + +int CeLU_vulkan::create_pipeline(const Option& opt) +{ + const Mat& shape = top_shapes.empty() ? Mat() : top_shapes[0]; + + int elempack = 1; + if (shape.dims == 1) elempack = opt.use_shader_pack8 && shape.w % 8 == 0 ? 8 : shape.w % 4 == 0 ? 4 : 1; + if (shape.dims == 2) elempack = opt.use_shader_pack8 && shape.h % 8 == 0 ? 8 : shape.h % 4 == 0 ? 4 : 1; + if (shape.dims == 3 || shape.dims == 4) elempack = opt.use_shader_pack8 && shape.c % 8 == 0 ? 8 : shape.c % 4 == 0 ? 4 : 1; + + size_t elemsize; + if (opt.use_fp16_storage) + { + elemsize = elempack * 2u; + } + else if (opt.use_fp16_packed) + { + elemsize = elempack == 1 ? 4u : elempack * 2u; + } + else + { + elemsize = elempack * 4u; + } + + Mat shape_packed; + if (shape.dims == 1) shape_packed = Mat(shape.w / elempack, (void*)0, elemsize, elempack); + if (shape.dims == 2) shape_packed = Mat(shape.w, shape.h / elempack, (void*)0, elemsize, elempack); + if (shape.dims == 3) shape_packed = Mat(shape.w, shape.h, shape.c / elempack, (void*)0, elemsize, elempack); + if (shape.dims == 4) shape_packed = Mat(shape.w, shape.h, shape.d, shape.c / elempack, (void*)0, elemsize, elempack); + + std::vector specializations(1 + 5); + specializations[0].f = alpha; + specializations[1 + 0].i = shape_packed.dims; + specializations[1 + 1].i = shape_packed.w; + specializations[1 + 2].i = shape_packed.h * shape_packed.d; + specializations[1 + 3].i = shape_packed.c; + specializations[1 + 4].i = shape_packed.cstep; + + Mat local_size_xyz; + if (shape_packed.dims == 1) + { + local_size_xyz.w = std::min(64, shape_packed.w); + local_size_xyz.h = 1; + local_size_xyz.c = 1; + } + if (shape_packed.dims == 2) + { + local_size_xyz.w = std::min(8, shape_packed.w); + local_size_xyz.h = std::min(8, shape_packed.h); + local_size_xyz.c = 1; + } + if (shape_packed.dims == 3) + { + local_size_xyz.w = std::min(4, shape_packed.w); + local_size_xyz.h = std::min(4, shape_packed.h); + local_size_xyz.c = std::min(4, shape_packed.c); + } + if (shape_packed.dims == 4) + { + local_size_xyz.w = std::min(4, shape_packed.w); + local_size_xyz.h = std::min(4, shape_packed.h * shape_packed.d); + local_size_xyz.c = std::min(4, shape_packed.c); + } + + // pack1 + if (shape.dims == 0 || elempack == 1) + { + pipeline_celu = new Pipeline(vkdev); + pipeline_celu->set_optimal_local_size_xyz(local_size_xyz); + pipeline_celu->create(LayerShaderType::celu, opt, specializations); + } + + // pack4 + if (shape.dims == 0 || elempack == 4) + { + pipeline_celu_pack4 = new Pipeline(vkdev); + pipeline_celu_pack4->set_optimal_local_size_xyz(local_size_xyz); + pipeline_celu_pack4->create(LayerShaderType::celu_pack4, opt, specializations); + } + + // pack8 + if ((opt.use_shader_pack8 && shape.dims == 0) || elempack == 8) + { + pipeline_celu_pack8 = new Pipeline(vkdev); + pipeline_celu_pack8->set_optimal_local_size_xyz(local_size_xyz); + pipeline_celu_pack8->create(LayerShaderType::celu_pack8, opt, specializations); + } + + return 0; +} + +int CeLU_vulkan::destroy_pipeline(const Option& /*opt*/) +{ + delete pipeline_celu; + pipeline_celu = 0; + + delete pipeline_celu_pack4; + pipeline_celu_pack4 = 0; + + delete pipeline_celu_pack8; + pipeline_celu_pack8 = 0; + + return 0; +} + +int CeLU_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 * bottom_top_blob.d; + constants[3].i = bottom_top_blob.c; + constants[4].i = bottom_top_blob.cstep; + + const Pipeline* pipeline = elempack == 8 ? pipeline_celu_pack8 + : elempack == 4 ? pipeline_celu_pack4 + : pipeline_celu; + + cmd.record_pipeline(pipeline, bindings, constants, bottom_top_blob); + + return 0; +} + +int CeLU_vulkan::forward_inplace(VkImageMat& bottom_top_blob, VkCompute& cmd, const Option& /*opt*/) const +{ + int elempack = bottom_top_blob.elempack; + + std::vector bindings(2); + bindings[0] = bottom_top_blob; + bindings[1] = 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 * bottom_top_blob.d; + constants[3].i = bottom_top_blob.c; + constants[4].i = 0; //bottom_top_blob.cstep; + + const Pipeline* pipeline = elempack == 8 ? pipeline_celu_pack8 + : elempack == 4 ? pipeline_celu_pack4 + : pipeline_celu; + + cmd.record_pipeline(pipeline, bindings, constants, bottom_top_blob); + + return 0; +} + +} // namespace ncnn diff --git a/src/layer/vulkan/celu_vulkan.h b/src/layer/vulkan/celu_vulkan.h new file mode 100644 index 000000000000..7f49131986c0 --- /dev/null +++ b/src/layer/vulkan/celu_vulkan.h @@ -0,0 +1,42 @@ +// 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_CELU_VULKAN_H +#define LAYER_CELU_VULKAN_H + +#include "celu.h" + +namespace ncnn { + +class CeLU_vulkan : virtual public CeLU +{ +public: + CeLU_vulkan(); + + virtual int create_pipeline(const Option& opt); + virtual int destroy_pipeline(const Option& opt); + + using CeLU::forward_inplace; + virtual int forward_inplace(VkMat& bottom_top_blob, VkCompute& cmd, const Option& opt) const; + virtual int forward_inplace(VkImageMat& bottom_top_blob, VkCompute& cmd, const Option& opt) const; + +public: + Pipeline* pipeline_celu; + Pipeline* pipeline_celu_pack4; + Pipeline* pipeline_celu_pack8; +}; + +} // namespace ncnn + +#endif // LAYER_CELU_VULKAN_H diff --git a/src/layer/vulkan/shader/celu.comp b/src/layer/vulkan/shader/celu.comp new file mode 100644 index 000000000000..a41dac3e7955 --- /dev/null +++ b/src/layer/vulkan/shader/celu.comp @@ -0,0 +1,73 @@ +// 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_EXT_shader_explicit_arithmetic_types_float16: require +#endif + +layout (constant_id = 0) const float alpha = 0; + +#define shape_constant_id_offset 1 +layout (constant_id = shape_constant_id_offset + 0) const int dims = 0; +layout (constant_id = shape_constant_id_offset + 1) const int w = 0; +layout (constant_id = shape_constant_id_offset + 2) const int h = 0; +layout (constant_id = shape_constant_id_offset + 3) const int c = 0; +layout (constant_id = shape_constant_id_offset + 4) const int cstep = 0; + +#if NCNN_image_shader +layout (binding = 0) uniform unfp sampler3D bottom_blob_3d; +layout (binding = 1, imfmtc1) writeonly uniform unfp image3D top_blob_3d; +#else +layout (binding = 0) buffer bottom_top_blob { sfp bottom_top_blob_data[]; }; +#endif + +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 >= psc(w) || gy >= psc(h) || gz >= psc(c)) + return; + +#if NCNN_image_shader + afp v = image3d_ld1(bottom_blob_3d, ivec3(gx, gy, gz)); +#else + const int gi = gz * psc(cstep) + gy * psc(w) + gx; + + afp v = buffer_ld1(bottom_top_blob_data, gi); +#endif + + v = max(v, afp(0.0f)) + min(alpha*(exp(v/alpha)-1), afp(0.0f)); + +#if NCNN_image_shader + image3d_st1(top_blob_3d, ivec3(gx, gy, gz), v); +#else + buffer_st1(bottom_top_blob_data, gi, v); +#endif +} diff --git a/src/layer/vulkan/shader/celu_pack4.comp b/src/layer/vulkan/shader/celu_pack4.comp new file mode 100644 index 000000000000..4d12a81d7ba5 --- /dev/null +++ b/src/layer/vulkan/shader/celu_pack4.comp @@ -0,0 +1,73 @@ +// 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_EXT_shader_explicit_arithmetic_types_float16: require +#endif + +layout (constant_id = 0) const float alpha = 0; + +#define shape_constant_id_offset 1 +layout (constant_id = shape_constant_id_offset + 0) const int dims = 0; +layout (constant_id = shape_constant_id_offset + 1) const int w = 0; +layout (constant_id = shape_constant_id_offset + 2) const int h = 0; +layout (constant_id = shape_constant_id_offset + 3) const int c = 0; +layout (constant_id = shape_constant_id_offset + 4) const int cstep = 0; + +#if NCNN_image_shader +layout (binding = 0) uniform unfp sampler3D bottom_blob_3d; +layout (binding = 1, imfmtc4) writeonly uniform unfp image3D top_blob_3d; +#else +layout (binding = 0) buffer bottom_top_blob { sfpvec4 bottom_top_blob_data[]; }; +#endif + +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 >= psc(w) || gy >= psc(h) || gz >= psc(c)) + return; + +#if NCNN_image_shader + afpvec4 v = image3d_ld4(bottom_blob_3d, ivec3(gx, gy, gz)); +#else + const int gi = gz * psc(cstep) + gy * psc(w) + gx; + + afpvec4 v = buffer_ld4(bottom_top_blob_data, gi); +#endif + + v = max(v, afp(0.0f)) + min(alpha*(exp(v/alpha)-1), afp(0.0f)); + +#if NCNN_image_shader + image3d_st4(top_blob_3d, ivec3(gx, gy, gz), v); +#else + buffer_st4(bottom_top_blob_data, gi, v); +#endif +} diff --git a/src/layer/vulkan/shader/celu_pack8.comp b/src/layer/vulkan/shader/celu_pack8.comp new file mode 100644 index 000000000000..98b957ea6910 --- /dev/null +++ b/src/layer/vulkan/shader/celu_pack8.comp @@ -0,0 +1,75 @@ +// Tencent is pleased to support the open source community by making ncnn available. +// +// Copyright (C) 2020 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 +struct sfpvec8 { f16vec4 abcd; f16vec4 efgh; }; +#endif +#if NCNN_fp16_arithmetic +#extension GL_EXT_shader_explicit_arithmetic_types_float16: require +#endif + +layout (constant_id = 0) const float alpha = 0; + +#define shape_constant_id_offset 1 +layout (constant_id = shape_constant_id_offset + 0) const int dims = 0; +layout (constant_id = shape_constant_id_offset + 1) const int w = 0; +layout (constant_id = shape_constant_id_offset + 2) const int h = 0; +layout (constant_id = shape_constant_id_offset + 3) const int c = 0; +layout (constant_id = shape_constant_id_offset + 4) const int cstep = 0; + +#if NCNN_image_shader +layout (binding = 0) uniform unfp sampler3D bottom_blob_3d; +layout (binding = 1, imfmtc4) writeonly uniform unfp image3D top_blob_3d; +#else +layout (binding = 0) buffer bottom_top_blob { sfpvec8 bottom_top_blob_data[]; }; +#endif + +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 >= psc(w) || gy >= psc(h) || gz >= psc(c)) + return; + +#if NCNN_image_shader + afpvec8 v = image3d_ld8(bottom_blob_3d, ivec3(gx, gy, gz)); +#else + const int gi = gz * psc(cstep) + gy * psc(w) + gx; + + afpvec8 v = buffer_ld8(bottom_top_blob_data, gi); +#endif + + v[0] = max(v[0], afp(0.0f)) + min(alpha*(exp(v[0]/alpha)-1), afp(0.0f)); + v[1] = max(v[1], afp(0.0f)) + min(alpha*(exp(v[1]/alpha)-1), afp(0.0f)); + +#if NCNN_image_shader + image3d_st8(top_blob_3d, ivec3(gx, gy, gz), v); +#else + buffer_st8(bottom_top_blob_data, gi, v); +#endif +}