add sigmoid for cl add add op unit test fixed #1411

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

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

src/operators/kernel/cl/cl_kernel/sigmoid.cl

+/* 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

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<GPU_CL, float>;
 
 }  // namespace operators
 }  // namespace paddle_mobile
+#endif

src/operators/kernel/cl/sigmoid_kernel.cpp

+/* 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<GPU_CL, float>::Init(SigmoidParam<GPU_CL>* param) {
+  this->cl_helper_.AddKernel("sigmoid", "sigmoid.cl");
+  return true;
+}
+
+template <>
+void SigmoidKernel<GPU_CL, float>::Compute(const SigmoidParam<GPU_CL>& 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<GPU_CL, float>;
+
+}  // namespace operators
+}  // namespace paddle_mobile
+#endif

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)

test/operators/cl/test_sigmoid_cl_op.cpp

+/* 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 <CL/cl.h>
+#include <cmath>
+#include <iostream>
+#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>();
+  float *y = Y->mutable_data<float>();
+
+  for (int i = 0; i < X->numel(); ++i) {
+    y[i] = 1.f / (1.f + exp(-x[i]));
+  }
+}
+
+int TestSigmoidOp(const std::vector<int> input_shape) {
+  paddle_mobile::PaddleMobile<paddle_mobile::GPU_CL> 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<float>(&input_tensor, framework::make_ddim(input_shape),
+                     static_cast<float>(0), static_cast<float>(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<float>(output->dims());
+
+  Sigmoid(&input_tensor2, &output_cmp);
+
+  const float *output_data = output->data<float>();
+  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<float>()[i] << "   "
+                     << input_tensor2.data<float>()[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;
+}

tools/build.sh

 #!/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

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--")