format code, add success check in bias_add

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>> {

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

mace/kernels/opencl/cl/bias_add.cl

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