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 3282556bb1afeff1085f35cfccd49301445a58a8..eac8446f572bb3398461fa386dd8c94c39597179 100644 --- a/src/operators/kernel/cl/feed_kernel.cpp +++ b/src/operators/kernel/cl/feed_kernel.cpp @@ -61,11 +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); -<<<<<<< HEAD - for (int i = 0; i < numel; i++) DLOG << Half2Float(out[i])<<","<>>>>>> 289b739de8517c21872107c16790b9cb2e7042d7 + } template class FeedKernel;