From 4782257aa5fc74defded6272c0cc3a65e225cacc Mon Sep 17 00:00:00 2001 From: xiebaiyuan Date: Wed, 23 Jan 2019 10:52:06 +0800 Subject: [PATCH] add sigmoid for cl add add op unit test fixed #1411 --- CMakeLists.txt | 2 +- src/operators/activation_op.cpp | 3 + src/operators/kernel/cl/cl_kernel/sigmoid.cl | 30 +++++++ src/operators/kernel/cl/reshape_kernel.cpp | 2 + src/operators/kernel/cl/sigmoid_kernel.cpp | 46 +++++++++++ test/CMakeLists.txt | 17 ++++ test/operators/cl/test_sigmoid_cl_op.cpp | 85 ++++++++++++++++++++ tools/build.sh | 2 +- tools/op.cmake | 6 ++ 9 files changed, 191 insertions(+), 2 deletions(-) create mode 100644 src/operators/kernel/cl/cl_kernel/sigmoid.cl create mode 100644 src/operators/kernel/cl/sigmoid_kernel.cpp create mode 100644 test/operators/cl/test_sigmoid_cl_op.cpp diff --git a/CMakeLists.txt b/CMakeLists.txt index e669f1b1c6..3d2d45b1ac 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -187,7 +187,7 @@ else() set(NET "default" CACHE STRING "select net type") endif() -set_property(CACHE NET PROPERTY STRINGS "default" "googlenet" "mobilenet" "yolo" "squeezenet" "FPGA_NET_V1" "FPGA_NET_V2" "NLP") +set_property(CACHE NET PROPERTY STRINGS "default" "googlenet" "mobilenet" "yolo" "squeezenet" "FPGA_NET_V1" "FPGA_NET_V2" "NLP" "op") include("${CMAKE_CURRENT_LIST_DIR}/tools/op.cmake") # build library diff --git a/src/operators/activation_op.cpp b/src/operators/activation_op.cpp index 6cb1c3c040..7bc78ef77f 100644 --- a/src/operators/activation_op.cpp +++ b/src/operators/activation_op.cpp @@ -66,6 +66,9 @@ REGISTER_OPERATOR_CL(relu, ops::ReluOp); #ifdef PADDLE_MOBILE_CPU REGISTER_OPERATOR_CPU(sigmoid, ops::SigmoidOp); #endif +#ifdef PADDLE_MOBILE_CL +REGISTER_OPERATOR_CL(sigmoid, ops::SigmoidOp); +#endif #endif // SIGMOID_OP #ifdef TANH_OP diff --git a/src/operators/kernel/cl/cl_kernel/sigmoid.cl b/src/operators/kernel/cl/cl_kernel/sigmoid.cl new file mode 100644 index 0000000000..d6dd9cfae2 --- /dev/null +++ b/src/operators/kernel/cl/cl_kernel/sigmoid.cl @@ -0,0 +1,30 @@ +/* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. + +Licensed under the Apache License, Version 2.0 (the "License"); +you may not use this file except in compliance with the License. +You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + +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. */ + +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +__kernel void sigmoid(__read_only image2d_t input, + __write_only image2d_t output){ + + const int x = get_global_id(0); + const int y = get_global_id(1); + + const sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE | + CLK_ADDRESS_CLAMP | + CLK_FILTER_NEAREST; + + half4 in = read_imageh(input, sampler, (int2)(x, y)); + in = 1.0f / (1 + exp(-in)); + write_imageh(output, (int2)(x, y), in); +} \ No newline at end of file diff --git a/src/operators/kernel/cl/reshape_kernel.cpp b/src/operators/kernel/cl/reshape_kernel.cpp index 4e8d3e1d60..18d98b0ff9 100644 --- a/src/operators/kernel/cl/reshape_kernel.cpp +++ b/src/operators/kernel/cl/reshape_kernel.cpp @@ -11,6 +11,7 @@ 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. */ +#ifdef RESHAPE_OP #include "operators/kernel/reshape_kernel.h" @@ -102,3 +103,4 @@ template class ReshapeKernel; } // namespace operators } // namespace paddle_mobile +#endif diff --git a/src/operators/kernel/cl/sigmoid_kernel.cpp b/src/operators/kernel/cl/sigmoid_kernel.cpp new file mode 100644 index 0000000000..4ac8b54ba9 --- /dev/null +++ b/src/operators/kernel/cl/sigmoid_kernel.cpp @@ -0,0 +1,46 @@ +/* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. + +Licensed under the Apache License, Version 2.0 (the "License"); +you may not use this file except in compliance with the License. +You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + +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. */ +#ifdef SIGMOID_OP + +#include "operators/kernel/activation_kernel.h" + +namespace paddle_mobile { +namespace operators { + +template <> +bool SigmoidKernel::Init(SigmoidParam* param) { + this->cl_helper_.AddKernel("sigmoid", "sigmoid.cl"); + return true; +} + +template <> +void SigmoidKernel::Compute(const SigmoidParam& param) { + auto kernel = this->cl_helper_.KernelAt(0); + const auto* input = param.InputX(); + auto* output = param.Out(); + auto default_work_size = this->cl_helper_.DefaultWorkSize(*output); + auto inputImage = input->GetCLImage(); + auto outputImage = output->GetCLImage(); + clSetKernelArg(kernel, 0, sizeof(cl_mem), &inputImage); + clSetKernelArg(kernel, 1, sizeof(cl_mem), &outputImage); + const size_t work_size[2] = {input->ImageWidth(), input->ImageHeight()}; + clEnqueueNDRangeKernel(this->cl_helper_.CLCommandQueue(), kernel, 2, NULL, + work_size, NULL, 0, NULL, NULL); +} + +template class SigmoidKernel; + +} // namespace operators +} // namespace paddle_mobile +#endif diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 167b374de9..28f83418be 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -154,6 +154,19 @@ if (CON GREATER -1) endif () +list(FIND NET "op" CON) +if (CON GREATER -1) + # gen test + ADD_EXECUTABLE(test-sigmoid operators/test_sigmoid_op.cpp test_include.h) + target_link_libraries(test-sigmoid paddle-mobile) + + # gen test + ADD_EXECUTABLE(test-sigmoid-cl operators/cl/test_sigmoid_cl_op.cpp test_helper.h test_include.h) + target_link_libraries(test-sigmoid-cl paddle-mobile) + + set(FOUND_MATCH ON) +endif () + if (NOT FOUND_MATCH) # gen test ADD_EXECUTABLE(test-resnet net/test_resnet.cpp test_helper.h test_include.h executor_for_test.h) @@ -352,6 +365,10 @@ if (NOT FOUND_MATCH) ADD_EXECUTABLE(test-sigmoid-op operators/test_sigmoid_op.cpp test_include.h) target_link_libraries(test-sigmoid-op paddle-mobile) + # gen test + ADD_EXECUTABLE(test-sigmoid-cl operators/cl/test_sigmoid_cl_op.cpp test_helper.h test_include.h) + target_link_libraries(test-sigmoid-cl paddle-mobile) + # gen test ADD_EXECUTABLE(test-depthwise-conv-op operators/test_depthwise_conv_op.cpp test_helper.h test_include.h executor_for_test.h) target_link_libraries(test-depthwise-conv-op paddle-mobile) diff --git a/test/operators/cl/test_sigmoid_cl_op.cpp b/test/operators/cl/test_sigmoid_cl_op.cpp new file mode 100644 index 0000000000..9909adc6c8 --- /dev/null +++ b/test/operators/cl/test_sigmoid_cl_op.cpp @@ -0,0 +1,85 @@ +/* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved. + +Licensed under the Apache License, Version 2.0 (the "License"); +you may not use this file except in compliance with the License. +You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + +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. */ +#ifdef PADDLE_MOBILE_CL +#include +#include +#include +#include "../../test_helper.h" +#include "../../test_include.h" +#include "framework/cl/cl_tensor.h" +#include "operators/activation_op.h" + +namespace paddle_mobile { + +void Sigmoid(const framework::Tensor *X, framework::Tensor *Y) { + const float *x = X->data(); + float *y = Y->mutable_data(); + + for (int i = 0; i < X->numel(); ++i) { + y[i] = 1.f / (1.f + exp(-x[i])); + } +} + +int TestSigmoidOp(const std::vector input_shape) { + paddle_mobile::PaddleMobile paddle_mobile; + // paddle_mobile.SetThreadNum(4); + auto time1 = paddle_mobile::time(); +#ifdef PADDLE_MOBILE_CL + paddle_mobile.SetCLPath("/data/local/tmp/bin"); +#endif + auto isok = paddle_mobile.Load(std::string(g_sigmoid), true, false, 1, true); + if (!isok) { + exit(1); + } + framework::DDim dims = framework::make_ddim(input_shape); + Tensor input_tensor; + SetupTensor(&input_tensor, framework::make_ddim(input_shape), + static_cast(0), static_cast(1)); + Tensor input_tensor2 = input_tensor; + + // paddle_mobile.Feed(input_tensor, "feed"); + paddle_mobile.Predict(input_tensor); + auto output = paddle_mobile.Fetch(); + + framework::Tensor output_cmp; + float *output_cmp_data = output_cmp.mutable_data(output->dims()); + + Sigmoid(&input_tensor2, &output_cmp); + + const float *output_data = output->data(); + for (int i = 0; i < output->numel(); ++i) { + float gap = output_data[i] - output_cmp_data[i]; + if (std::abs(gap / (output_data[i] + 1e-5)) > 1e-2) { + LOG(kLOG_INFO) << "input_tensor[" << i + << "]= " << input_tensor.data()[i] << " " + << input_tensor2.data()[i] << "output_data[" << i + << "] = " << output_data[i] << ", output_cmp_data[" << i + << "] = " << output_cmp_data[i]; + exit(1); + } + } + return 0; +} + +} // namespace paddle_mobile +#endif +int main() { +#ifdef PADDLE_MOBILE_CL + paddle_mobile::TestSigmoidOp({1, 1, 1000, 1000}); + paddle_mobile::TestSigmoidOp({1, 3, 11, 22}); + paddle_mobile::TestSigmoidOp({1, 32, 112, 112}); + std::cout << "test sigmoid op cl pass." << std::endl; +#endif + return 0; +} diff --git a/tools/build.sh b/tools/build.sh index 117e74adea..cac18c79d7 100755 --- a/tools/build.sh +++ b/tools/build.sh @@ -1,6 +1,6 @@ #!/usr/bin/env bash NETS="" -declare -a supportedNets=("googlenet" "mobilenet" "yolo" "squeezenet" "resnet" "mobilenetssd" "nlp" "mobilenetfssd" "genet" "super") +declare -a supportedNets=("googlenet" "mobilenet" "yolo" "squeezenet" "resnet" "mobilenetssd" "nlp" "mobilenetfssd" "genet" "super" "op") build_for_mac() { if [ ! `which brew` ]; then diff --git a/tools/op.cmake b/tools/op.cmake index 84b5bb6ef0..94dc0b4215 100755 --- a/tools/op.cmake +++ b/tools/op.cmake @@ -228,6 +228,12 @@ if (CON GREATER -1) set(FOUND_MATCH ON) endif() +list(FIND NET "op" CON) +if (CON GREATER -1) + message("op enabled") + set(SIGMOID_OP ON) + set(FOUND_MATCH ON) +endif() if(NOT FOUND_MATCH) message("--default--") -- GitLab