From 23d4246ceda73cecda063e44051a55ce5f32851f Mon Sep 17 00:00:00 2001
From: yejianwu <yejianwu@xiaomi.com>
Date: Tue, 5 Dec 2017 10:16:05 +0800
Subject: [PATCH] format code, add success check in bias_add

---
 mace/kernels/opencl/batch_norm_opencl.cc | 10 +++++-----
 mace/kernels/opencl/bias_add_opencl.cc   |  3 ++-
 mace/kernels/opencl/cl/bias_add.cl       |  4 ++--
 3 files changed, 9 insertions(+), 8 deletions(-)

diff --git a/mace/kernels/opencl/batch_norm_opencl.cc b/mace/kernels/opencl/batch_norm_opencl.cc
index c1728689..da0a7556 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 70c8f05d..b778c6b7 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 1663c7ea..3dc0eabe 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);
-- 
GitLab