diff --git a/src/framework/cl/cl_engine.h b/src/framework/cl/cl_engine.h
index 15af651136b79a8034609596b6454410cb4e5d33..d8a4c771a5da9aff96cbb86abec63762b9cd4dce 100644
--- a/src/framework/cl/cl_engine.h
+++ b/src/framework/cl/cl_engine.h
@@ -88,6 +88,8 @@ class CLEngine {
     return true;
   }
 
+  cl_device_id DeviceID(int index = 0) { return devices_[index]; }
+
  private:
   CLEngine() { initialized_ = false; }
 
diff --git a/src/framework/cl/cl_scope.h b/src/framework/cl/cl_scope.h
index 93a74a36452f42a019954d200af093b699d17344..58357632be8b6257de5a9ecf1e1e950f9f422f77 100644
--- a/src/framework/cl/cl_scope.h
+++ b/src/framework/cl/cl_scope.h
@@ -58,11 +58,24 @@ class CLScope {
     auto program = CLEngine::Instance()->CreateProgramWith(
         context_.get(), "./cl_kernel/" + file_name);
 
+    DLOG << " --- begin build program -> " << file_name << " --- ";
     status_ =
         clBuildProgram(program.get(), 0, 0, "-cl-fast-relaxed-math", 0, 0);
 
     CL_CHECK_ERRORS(status_);
 
+    if (status_ == CL_BUILD_PROGRAM_FAILURE) {
+      size_t log_size;
+      clGetProgramBuildInfo(program.get(), CLEngine::Instance()->DeviceID(),
+                            CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
+      char *log = (char *)malloc(log_size);
+      clGetProgramBuildInfo(program.get(), CLEngine::Instance()->DeviceID(),
+                            CL_PROGRAM_BUILD_LOG, log_size, log, NULL);
+      DLOG << " program build error: " << log;
+    }
+
+    DLOG << " --- end build program -> " << file_name << " --- ";
+
     programs_[file_name] = std::move(program);
 
     return programs_[file_name].get();
diff --git a/src/io/api_paddle_mobile.cc b/src/io/api_paddle_mobile.cc
index 4609438ec9fbdb5b5030b56a4bf18b9437bf7c2e..5dce367b525214e38b78d043410da62c195f3365 100644
--- a/src/io/api_paddle_mobile.cc
+++ b/src/io/api_paddle_mobile.cc
@@ -114,6 +114,8 @@ CreatePaddlePredictor<PaddleMobileConfig, PaddleEngineKind::kPaddleMobile>(
       x.reset(new PaddleMobilePredictor<FPGA, Precision::FP32>(config));
     } else if (config.device == PaddleMobileConfig::kGPU_MALI) {
       x.reset(new PaddleMobilePredictor<GPU_MALI, Precision::FP32>(config));
+    } else if (config.device == PaddleMobileConfig::kGPU_CL) {
+      x.reset(new PaddleMobilePredictor<GPU_CL, Precision::FP32>(config));
     } else {
       LOG(kLOG_ERROR) << "unsupport device type!";
       return nullptr;
diff --git a/src/io/paddle_inference_api.h b/src/io/paddle_inference_api.h
index 97564f4132d2e43cf736c2eb4a95d437584be24f..6c3ce7ef77ee81bf4ce31ecd60efd65380256bbc 100644
--- a/src/io/paddle_inference_api.h
+++ b/src/io/paddle_inference_api.h
@@ -111,7 +111,7 @@ class PaddlePredictor {
 
 struct PaddleMobileConfig : public PaddlePredictor::Config {
   enum Precision { FP32 = 0 };
-  enum Device { kCPU = 0, kFPGA = 1, kGPU_MALI = 2 };
+  enum Device { kCPU = 0, kFPGA = 1, kGPU_MALI = 2, kGPU_CL = 3 };
 
   enum Precision precision;
   enum Device device;
diff --git a/src/operators/kernel/cl/cl_kernel/reshape.cl b/src/operators/kernel/cl/cl_kernel/reshape.cl
index 4055445d1576b2ca54919ed03ad187d08cff14c2..767a7105d29b02b4ced95f3b9834e048225d538a 100644
--- a/src/operators/kernel/cl/cl_kernel/reshape.cl
+++ b/src/operators/kernel/cl/cl_kernel/reshape.cl
@@ -43,7 +43,7 @@ __kernel void reshape(__read_only image2d_t input,
     int i0 = oindex;
     int ix = (i1 / 4) * d3 + i3;
     int iy = i0 * d2 + i2;
-    r[i] = read_imageh(input, sampler, int2(ix, iy))[i1%4];
+    r[i] = read_imageh(input, sampler, (int2)(ix, iy))[i1%4];
   }
-  write_imageh(output, int2(x, y), r);
+  write_imageh(output, (int2)(x, y), r);
 }
\ No newline at end of file