diff --git a/src/framework/executor.cpp b/src/framework/executor.cpp index 609e069d48bac3347eb74d3f9573bdb2d45ab10e..bb856fcf549393df1205a1b815f40f653f491fae 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 e773d1c2577461abb35fabfa752ffc272970492b..20714202a6745bcb4f1810c20665a20e387e665d 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 4f691342446d3fef4af920477fd661b6623d8567..9b6ff736b96d7c5265f5caee2dfe0f2dc63a5edb 100644 --- a/src/operators/kernel/cl/feed_kernel.cpp +++ b/src/operators/kernel/cl/feed_kernel.cpp @@ -61,7 +61,7 @@ void FeedKernel::Compute(const FeedParam ¶m) { 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;