diff --git a/mace/kernels/opencl/batch_norm_opencl.cc b/mace/kernels/opencl/batch_norm_opencl.cc
index c17286895a8868732ada5608d9454cae31cdd746..da0a755622339dd44f410991675ad2fa208b7cb8 100644
--- a/mace/kernels/opencl/batch_norm_opencl.cc
+++ b/mace/kernels/opencl/batch_norm_opencl.cc
@@ -44,11 +44,11 @@ void BatchNormFunctor<DeviceType::OPENCL, T>::operator()(
 
   uint32_t idx = 0;
   bm_kernel.setArg(idx++, *(static_cast<const cl::Image2D *>(input->buffer())));
-  bm_kernel.setArg(idx++, *(static_cast<cl::Image2D *>(scale->buffer())));
-  bm_kernel.setArg(idx++, *(static_cast<cl::Image2D *>(offset->buffer())));
-  bm_kernel.setArg(idx++, *(static_cast<cl::Image2D *>(mean->buffer())));
-  bm_kernel.setArg(idx++, *(static_cast<cl::Image2D *>(var->buffer())));
-  bm_kernel.setArg(idx++, *(static_cast<cl::Buffer *>(epsilon->buffer())));
+  bm_kernel.setArg(idx++, *(static_cast<const cl::Image2D *>(scale->buffer())));
+  bm_kernel.setArg(idx++, *(static_cast<const cl::Image2D *>(offset->buffer())));
+  bm_kernel.setArg(idx++, *(static_cast<const cl::Image2D *>(mean->buffer())));
+  bm_kernel.setArg(idx++, *(static_cast<const cl::Image2D *>(var->buffer())));
+  bm_kernel.setArg(idx++, *(static_cast<const cl::Buffer *>(epsilon->buffer())));
   bm_kernel.setArg(idx++, *(static_cast<cl::Image2D *>(output->buffer())));
 
   auto params_generator = [&kwg_size]()->std::vector<std::vector<uint32_t>> {
diff --git a/mace/kernels/opencl/bias_add_opencl.cc b/mace/kernels/opencl/bias_add_opencl.cc
index 70c8f05d0fbe8c7fc0a88651592d90c53bbb98b1..b778c6b72bb2de97e93e1882e7663695408da7f4 100644
--- a/mace/kernels/opencl/bias_add_opencl.cc
+++ b/mace/kernels/opencl/bias_add_opencl.cc
@@ -40,7 +40,7 @@ void BiasAddFunctor<DeviceType::OPENCL, T>::operator()(
 
   uint32_t idx = 0;
   bias_kernel.setArg(idx++, *(static_cast<const cl::Image2D *>(input->buffer())));
-  bias_kernel.setArg(idx++, *(static_cast<cl::Image2D *>(bias->buffer())));
+  bias_kernel.setArg(idx++, *(static_cast<const cl::Image2D *>(bias->buffer())));
   bias_kernel.setArg(idx++, *(static_cast<cl::Image2D *>(output->buffer())));
 
   cl_int error = runtime->command_queue().enqueueNDRangeKernel(
@@ -48,6 +48,7 @@ void BiasAddFunctor<DeviceType::OPENCL, T>::operator()(
       cl::NDRange(gws[0], gws[1], gws[2]),
       cl::NDRange(lws[0], lws[1], lws[2]),
       NULL, OpenCLRuntime::Get()->GetDefaultEvent());
+  MACE_CHECK(error == CL_SUCCESS);
 }
 
 template
diff --git a/mace/kernels/opencl/cl/bias_add.cl b/mace/kernels/opencl/cl/bias_add.cl
index 1663c7ea63249d6981213604a4aabf783d8a8f33..3dc0eabe60f8fce68e03e8c8abe789fb35ace510 100644
--- a/mace/kernels/opencl/cl/bias_add.cl
+++ b/mace/kernels/opencl/cl/bias_add.cl
@@ -1,8 +1,8 @@
 #include <common.h>
 // Supported data types: half/float
 __kernel void bias_add(__read_only image2d_t input,
-                         __read_only image2d_t bias,
-                         __write_only image2d_t output) {
+                       __read_only image2d_t bias,
+                       __write_only image2d_t output) {
   const int ch_blk = get_global_id(0);
   const int w = get_global_id(1);
   const int hb = get_global_id(2);