diff --git a/mace/core/runtime/opencl/opencl_wrapper.cc b/mace/core/runtime/opencl/opencl_wrapper.cc
index cc79b56143c0c033409235c10a8f1039e3546222..b8cd709730d16aec1715e415bbeff4c552cd8787 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<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);
diff --git a/mace/kernels/opencl/cl/common.h b/mace/kernels/opencl/cl/common.h
index ac870bd38441b64c955cdc5465d8f595bbf84aa1..dd61012cebdb887639f132d63e48d9521898e443 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 878c110d7903a616727f632a39c87f9c92e83a91..057a66a4ca6514fb991870f65c813ffcedb75623 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 0188d7679153c18b347e7691a3b30d9e350e6ef5..6830b50801aff517f0dfeda9868c983721df65dc 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));