diff --git a/src/layer/celu.cpp b/src/layer/celu.cpp index 4bddfc368..58782f877 100644 --- a/src/layer/celu.cpp +++ b/src/layer/celu.cpp @@ -35,8 +35,9 @@ 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; + int size = w * h * d; #pragma omp parallel for num_threads(opt.num_threads) for (int q = 0; q < channels; q++) diff --git a/src/layer/vulkan/celu_vulkan.cpp b/src/layer/vulkan/celu_vulkan.cpp new file mode 100644 index 000000000..3a45619d6 --- /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 000000000..b5e25e19b --- /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 000000000..6f70a53b2 --- /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.0f), 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 000000000..697deab47 --- /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.0f), 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 000000000..cfd659812 --- /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.0f), 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 +} diff --git a/tests/test_celu.cpp b/tests/test_celu.cpp index 79b0bbe18..703864eab 100644 --- a/tests/test_celu.cpp +++ b/tests/test_celu.cpp @@ -25,7 +25,7 @@ static int test_celu(const ncnn::Mat& a, float alpha) int ret = test_layer("CELU", pd, weights, a); if (ret != 0) { - fprintf(stderr, "test_celu failed a.dims=%d a=(%d %d %d) alpha=%f\n", a.dims, a.w, a.h, a.c, alpha); + fprintf(stderr, "test_celu failed a.dims=%d a=(%d %d %d %d) alpha=%f\n", a.dims, a.w, a.h, a.d, a.c, alpha); } return ret; @@ -34,25 +34,45 @@ static int test_celu(const ncnn::Mat& a, float alpha) static int test_celu_0() { return 0 - || test_celu(RandomMat(5, 7, 24), 1.f) - || test_celu(RandomMat(7, 9, 12), 0.5f) - || test_celu(RandomMat(3, 5, 13), 0.2f); + || test_celu(RandomMat(3, 8, 12, 18), 1.f) + || test_celu(RandomMat(4, 7, 9, 16), 0.1f) + || test_celu(RandomMat(3, 5, 12, 16), 1.f) + || test_celu(RandomMat(9, 6, 7, 14), 0.1f) + || test_celu(RandomMat(5, 6, 9, 10), 1.f) + || test_celu(RandomMat(6, 8, 2, 15), 0.1f); } static int test_celu_1() { return 0 - || test_celu(RandomMat(15, 24), 1.f) - || test_celu(RandomMat(17, 12), 0.5f) - || test_celu(RandomMat(19, 15), 0.2f); + || test_celu(RandomMat(7, 6, 18), 1.f) + || test_celu(RandomMat(9, 6, 15), 0.1f) + || test_celu(RandomMat(9, 7, 16), 1.f) + || test_celu(RandomMat(6, 10, 15), 0.1f) + || test_celu(RandomMat(2, 7, 11), 1.f) + || test_celu(RandomMat(6, 10, 7), 0.1f); } static int test_celu_2() { return 0 + || test_celu(RandomMat(12, 18), 1.f) + || test_celu(RandomMat(18, 12), 0.1f) + || test_celu(RandomMat(23, 27), 1.f) + || test_celu(RandomMat(18, 16), 0.1f) + || test_celu(RandomMat(18, 16), 1.f) + || test_celu(RandomMat(20, 16), 0.1f); +} + +static int test_celu_3() +{ + return 0 + || test_celu(RandomMat(256), 1.f) + || test_celu(RandomMat(64), 0.1f) + || test_celu(RandomMat(128), 1.f) + || test_celu(RandomMat(96), 0.1f) || test_celu(RandomMat(128), 1.f) - || test_celu(RandomMat(124), 0.5f) - || test_celu(RandomMat(127), 0.2f); + || test_celu(RandomMat(128), 0.1f); } int main() @@ -62,5 +82,6 @@ int main() return 0 || test_celu_0() || test_celu_1() - || test_celu_2(); + || test_celu_2() + || test_celu_3(); } diff --git a/tools/onnx/onnx2ncnn.cpp b/tools/onnx/onnx2ncnn.cpp index b3ae40115..137ecf7bf 100644 --- a/tools/onnx/onnx2ncnn.cpp +++ b/tools/onnx/onnx2ncnn.cpp @@ -3641,6 +3641,10 @@ int main(int argc, char** argv) { fprintf(pp, "%-16s", "UnaryOp"); } + else if (op == "Celu") + { + fprintf(pp, "%-16s", "CELU"); + } else if (op == "Clip") { fprintf(pp, "%-16s", "Clip"); @@ -4154,6 +4158,12 @@ int main(int argc, char** argv) int op_type = 3; fprintf(pp, " 0=%d", op_type); } + else if (op == "CeLU") + { + float alpha = get_node_attr_f(node, "alpha", 1.0f); + + fprintf(pp, " 0=%e", alpha); + } else if (op == "Clip") { float min;