From b77d113c7d9cb9b3ee9a83c7d8be9201c6e7d795 Mon Sep 17 00:00:00 2001 From: Liangliang He Date: Thu, 22 Mar 2018 12:03:45 +0800 Subject: [PATCH] Diasble OpenCL symbol loading error checking and fix kernels syntax error --- mace/core/runtime/opencl/opencl_wrapper.cc | 19 +++++++++---------- mace/kernels/opencl/cl/common.h | 8 ++++---- mace/kernels/opencl/cl/fully_connected.cl | 5 +++-- mace/kernels/opencl/cl/softmax.cl | 2 +- 4 files changed, 17 insertions(+), 17 deletions(-) diff --git a/mace/core/runtime/opencl/opencl_wrapper.cc b/mace/core/runtime/opencl/opencl_wrapper.cc index cc79b561..b8cd7097 100644 --- a/mace/core/runtime/opencl/opencl_wrapper.cc +++ b/mace/core/runtime/opencl/opencl_wrapper.cc @@ -290,16 +290,15 @@ void *OpenCLLibraryImpl::LoadFromPath(const std::string &path) { return nullptr; } -#define MACE_CL_ASSIGN_FROM_DLSYM(func) \ - do { \ - void *ptr = dlsym(handle, #func); \ - if (ptr == nullptr) { \ - LOG(ERROR) << "Failed to load " << #func << " from " << path; \ - dlclose(handle); \ - return nullptr; \ - } \ - func = reinterpret_cast(ptr); \ - VLOG(2) << "Loaded " << #func << " from " << path; \ +#define MACE_CL_ASSIGN_FROM_DLSYM(func) \ + do { \ + void *ptr = dlsym(handle, #func); \ + if (ptr == nullptr) { \ + VLOG(1) << "Failed to load " << #func << " from " << path; \ + continue; \ + } \ + func = reinterpret_cast(ptr); \ + VLOG(2) << "Loaded " << #func << " from " << path; \ } while (false) MACE_CL_ASSIGN_FROM_DLSYM(clGetPlatformIDs); diff --git a/mace/kernels/opencl/cl/common.h b/mace/kernels/opencl/cl/common.h index ac870bd3..dd61012c 100644 --- a/mace/kernels/opencl/cl/common.h +++ b/mace/kernels/opencl/cl/common.h @@ -28,19 +28,19 @@ inline DATA_TYPE4 do_activation(DATA_TYPE4 in, __private const float relux_max_limit) { DATA_TYPE4 out; #ifdef USE_RELU - out = fmax(in, 0); + out = fmax(in, (DATA_TYPE)0); #endif #ifdef USE_RELUX - out = clamp(in, 0, relux_max_limit); + out = clamp(in, (DATA_TYPE4)0, relux_max_limit); #endif #ifdef USE_PRELU - out = select(prelu_alpha * in, in, in >= 0); + out = select(prelu_alpha * in, in, in >= (DATA_TYPE)0); #endif #ifdef USE_TANH out = tanh(in); #endif #ifdef USE_SIGMOID - out = native_recip(1.0 + native_exp(-in)); + out = native_recip((DATA_TYPE)1 + native_exp(-in)); #endif return out; } diff --git a/mace/kernels/opencl/cl/fully_connected.cl b/mace/kernels/opencl/cl/fully_connected.cl index 878c110d..057a66a4 100644 --- a/mace/kernels/opencl/cl/fully_connected.cl +++ b/mace/kernels/opencl/cl/fully_connected.cl @@ -81,7 +81,7 @@ __kernel void fully_connected_width(__read_only image2d_t input, int2 input_coord, weight_coord; DATA_TYPE4 in, w; - DATA_TYPE sum = 0.0; + DATA_TYPE sum = 0; input_coord = (int2)(0, mul24(batch_idx, input_height)); @@ -107,7 +107,8 @@ __kernel void fully_connected_width(__read_only image2d_t input, input_coord.y++; } - const short inter_out_offset = mad24(get_local_id(1), 4, get_local_id(0)); + const short inter_out_offset = mad24((short)get_local_id(1), (short)4, + (short)get_local_id(0)); const short local_width_blk_size = (short)get_local_size(1); const short local_size = mul24((short)get_local_size(0), local_width_blk_size); diff --git a/mace/kernels/opencl/cl/softmax.cl b/mace/kernels/opencl/cl/softmax.cl index 0188d767..6830b508 100644 --- a/mace/kernels/opencl/cl/softmax.cl +++ b/mace/kernels/opencl/cl/softmax.cl @@ -12,7 +12,7 @@ __kernel void softmax(__read_only image2d_t input, int pos = width_idx; DATA_TYPE max_value = -FLT_MAX; - DATA_TYPE sum = 0.0; + DATA_TYPE sum = 0; DATA_TYPE4 data; for (short i = 0; i < chan_blks; ++i) { data = READ_IMAGET(input, SAMPLER, (int2)(pos, hb_idx)); -- GitLab