From b83becca691b530481f80a91a20a1d999b2420a8 Mon Sep 17 00:00:00 2001
From: liuruilong <liuruilong@baidu.com>
Date: Mon, 15 Oct 2018 23:19:08 +0800
Subject: [PATCH]  update relu kernel

---
 src/framework/executor.cpp                |  2 +-
 src/operators/kernel/cl/cl_kernel/relu.cl | 24 ++++++++---------------
 src/operators/kernel/cl/feed_kernel.cpp   |  2 +-
 3 files changed, 10 insertions(+), 18 deletions(-)

diff --git a/src/framework/executor.cpp b/src/framework/executor.cpp
index 609e069d48..bb856fcf54 100644
--- a/src/framework/executor.cpp
+++ b/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 {
diff --git a/src/operators/kernel/cl/cl_kernel/relu.cl b/src/operators/kernel/cl/cl_kernel/relu.cl
index e773d1c257..20714202a6 100644
--- a/src/operators/kernel/cl/cl_kernel/relu.cl
+++ b/src/operators/kernel/cl/cl_kernel/relu.cl
@@ -1,25 +1,17 @@
-/* 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
diff --git a/src/operators/kernel/cl/feed_kernel.cpp b/src/operators/kernel/cl/feed_kernel.cpp
index 4f69134244..9b6ff736b9 100644
--- a/src/operators/kernel/cl/feed_kernel.cpp
+++ b/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>;
-- 
GitLab