Add sync function for opencl test and benchmark.

mace/core/runtime/opencl/opencl_runtime.cc

@@ -8,11 +8,9 @@
 #include <mutex>
 
 #include <dirent.h>
-#include <errno.h>
 
 #include "mace/core/logging.h"
 #include "mace/core/runtime/opencl/opencl_runtime.h"
-#include "mace/core/runtime/opencl/opencl_wrapper.h"
 
 namespace mace {
 namespace {
@@ -66,7 +64,7 @@ bool BuildProgram(OpenCLRuntime *runtime,
   };
 
   *program = cl::Program(runtime->context(), sources);
-  std::string build_options = "-Werror -cl-mad-enable -I" + path;
+  std::string build_options = "-Werror -cl-mad-enable -cl-fast-relaxed-math -I" + path;
   // TODO(heliangliang) -cl-unsafe-math-optimizations -cl-fast-relaxed-math
   if (program->build({runtime->device()}, build_options.c_str()) != CL_SUCCESS) {
     if (program->getBuildInfo<CL_PROGRAM_BUILD_STATUS>(runtime->device()) ==

mace/core/runtime/opencl/opencl_runtime.h

@@ -20,15 +20,18 @@ namespace mace {
 class OpenCLRuntime {
  public:
   static OpenCLRuntime *Get();
-  OpenCLRuntime(cl::Context context,
-                cl::Device device,
-                cl::CommandQueue command_queue);
-  ~OpenCLRuntime();
 
   cl::Context &context();
   cl::Device &device();
   cl::CommandQueue &command_queue();
   cl::Program &program();
+ private:
+  OpenCLRuntime(cl::Context context,
+                cl::Device device,
+                cl::CommandQueue command_queue);
+  ~OpenCLRuntime();
+  OpenCLRuntime(const OpenCLRuntime&) = delete;
+  OpenCLRuntime &operator=(const OpenCLRuntime&) = delete;
 
  private:
   cl::Context context_;

mace/kernels/opencl/batch_norm_opencl.cc

@@ -24,25 +24,21 @@ void BatchNormFunctor<DeviceType::OPENCL, float>::operator()(
 
   auto runtime = OpenCLRuntime::Get();
   auto program = runtime->program();
-  auto batch_norm_kernel =
-      cl::KernelFunctor<cl::Buffer, cl::Buffer, cl::Buffer,
-                        cl::Buffer, cl::Buffer, cl::Buffer,
-                        int, int, cl::Buffer>(program, "batch_norm");
-  cl_int error;
-  auto res_event = batch_norm_kernel(cl::EnqueueArgs(runtime->command_queue(),
-                              cl::NDRange(n * channel * sample_size),
-                              cl::NDRange(128)),
-                    *(static_cast<const cl::Buffer *>(input->buffer())),
-                    *(static_cast<cl::Buffer *>(scale->buffer())),
-                    *(static_cast<cl::Buffer *>(offset->buffer())),
-                    *(static_cast<cl::Buffer *>(mean->buffer())),
-                    *(static_cast<cl::Buffer *>(var->buffer())),
-                    *(static_cast<cl::Buffer *>(epsilon->buffer())),
-                    static_cast<int>(channel),
-                    static_cast<int>(sample_size),
-                    *(static_cast<cl::Buffer *>(output->buffer())),
-                    error);
-  res_event.wait();
+  auto _kernel = cl::Kernel(program, "batch_norm");
+  _kernel.setArg(0, *(static_cast<const cl::Buffer *>(input->buffer())));
+  _kernel.setArg(1, *(static_cast<cl::Buffer *>(scale->buffer())));
+  _kernel.setArg(2, *(static_cast<cl::Buffer *>(offset->buffer())));
+  _kernel.setArg(3, *(static_cast<cl::Buffer *>(mean->buffer())));
+  _kernel.setArg(4, *(static_cast<cl::Buffer *>(var->buffer())));
+  _kernel.setArg(5, *(static_cast<cl::Buffer *>(epsilon->buffer())));
+  _kernel.setArg(6, static_cast<int>(sample_size));
+  _kernel.setArg(7, *(static_cast<cl::Buffer *>(output->buffer())));
+  _kernel.setArg(8, 32u, nullptr);
+  _kernel.setArg(9, 32u, nullptr);
+  cl_int error = runtime->command_queue().enqueueNDRangeKernel(
+      _kernel, cl::NullRange,
+      cl::NDRange(n, channel, sample_size),
+      cl::NDRange(1, 1, 128));
   MACE_CHECK(error == CL_SUCCESS);
 }
 

mace/kernels/opencl/cl/batch_norm.cl

@@ -4,16 +4,28 @@ void kernel batch_norm(global const float *input,
                        global const float *mean,
                        global const float *var,
                        global const float *epsilon,
-                       private const int channels,
                        private const int pixels,
-                       global float *output) {
-  int idx = get_global_id(0);
-  int channel = (idx % (channels * pixels)) / pixels;
+                       global float *output,
+                       __local float *new_scale,
+                       __local float *new_offset) {
+  const int batch = get_global_id(0);
+  const int channel = get_global_id(1);
+  const int channels = get_global_size(1);
+  const int pixel_offset = get_global_id(2);
+  const unsigned int local_channel = get_local_id(1);
+  const int local_pixel_idx = get_local_id(2);
 
-  const float *input_ptr = input + idx;
-  const float new_scale = scale[channel] * rsqrt(var[channel] + *epsilon);
-  const float new_offset = offset[channel] - mean[channel] * new_scale;
-  float *output_ptr = output + idx;
-  *output_ptr = new_scale * *input_ptr + new_offset;
+  if(local_pixel_idx == 0) {
+    new_scale[local_channel] = scale[channel] * rsqrt(var[channel] + *epsilon);
+    new_offset[local_channel] = offset[channel] - mean[channel] * new_scale[local_channel];
+  }
+
+  barrier(CLK_LOCAL_MEM_FENCE);
+
+  const int sample_offset = (batch * channels + channel) * pixels + pixel_offset;
+
+  const float *input_ptr = input + sample_offset;
+  float *output_ptr = output + sample_offset;
+  *output_ptr = new_scale[local_channel] * *input_ptr + new_offset[local_channel];
 }
 

mace/ops/BUILD

@@ -17,6 +17,7 @@ cc_library(
     ],
     deps = [
         "//mace/core",
+        "//mace/core:opencl_runtime",
         "@gtest//:gtest",
     ],
 )
@@ -39,7 +40,6 @@ cc_library(
         "-fopenmp",
     ],
     deps = [
-        "//mace/core",
         "//mace/kernels",
         "//mace/proto:cc_proto",
     ],
@@ -72,7 +72,6 @@ cc_test(
     deps = [
         ":ops",
         ":test",
-        "//mace/core",
         "//mace/core:test_benchmark_main",
     ],
 )

mace/ops/batch_norm_benchmark.cc

@@ -34,11 +34,13 @@ static void BatchNorm(
   // Warm-up
   for (int i = 0; i < 5; ++i) {
     net.RunOp(D);
+    net.Sync();
   }
 
   mace::testing::StartTiming();
   while (iters--) {
     net.RunOp(D);
+    net.Sync();
   }
 }
 

mace/ops/batch_norm_test.cc

@@ -208,6 +208,7 @@ TEST_F(BatchNormOpTest, ComplexRandomOPENCL) {
 
   // Run NEON
   net.RunOp(DeviceType::OPENCL);
+  net.Sync();
 
   // Check
   Tensor expected;

mace/ops/ops_test_util.h

@@ -11,6 +11,7 @@
 #include "mace/core/common.h"
 #include "mace/core/net.h"
 #include "mace/core/tensor.h"
+#include "mace/core/runtime/opencl/opencl_runtime.h"
 
 namespace mace {
 
@@ -152,6 +153,12 @@ class OpsTestNet {
     return ws_.GetTensor(output_name);
   }
 
+  void Sync() {
+    if (net_) {
+      OpenCLRuntime::Get()->command_queue().finish();
+    }
+  }
+
  public:
   Workspace ws_;
   OperatorDef op_def_;