Merge branch 'disable-opencl-loading-error-check' into 'master'

Diasble OpenCL symbol loading error checking and fix kernels syntax error

See merge request !320

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<func##Func>(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<func##Func>(ptr);                    \
+    VLOG(2) << "Loaded " << #func << " from " << path;           \
   } while (false)
 
   MACE_CL_ASSIGN_FROM_DLSYM(clGetPlatformIDs);

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;
 }

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);

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));