Merge pull request #1092 from codeWorm2015/opencl

 update relu kernel

src/framework/executor.cpp

@@ -37,7 +37,7 @@ limitations under the License. */
 #include "framework/cl/cl_image.h"
 #endif
 
-int debug_to = 4;
+int debug_to = 5;
 
 namespace paddle_mobile {
 namespace framework {

src/operators/kernel/cl/cl_kernel/relu.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 relu(__read_only image2d_t input,
-                   __write_only image2d_t output)
+                   __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 r = read_imageh(input, sampler, int2(x, y));
-  r = max(half4(0, 0, 0, 0), r);
-  write_imageh(output, int2(x, y), r);
+
+  half4 in = read_imageh(input, sampler, (int2)(x, y));
+  in = max((half4)(0.0), in);
+  write_imageh(output, (int2)(x, y), in);
 }
\ No newline at end of file

src/operators/kernel/cl/feed_kernel.cpp

@@ -61,7 +61,7 @@ void FeedKernel<GPU_CL, float>::Compute(const FeedParam<GPU_CL> &param) {
   size_t region[3] = {height, width, 1};
   clEnqueueReadImage(commandQueue, cl_image, CL_TRUE, origin, region, 0, 0, out,
                      0, NULL, NULL);
-//  for (int i = 0; i < numel; i++) DLOG << Half2Float(out[i]);
+  //  for (int i = 0; i < numel; i++) DLOG << Half2Float(out[i]);
 }
 
 template class FeedKernel<GPU_CL, float>;