change scalar math to cwise

mace/core/operator.cc

@@ -89,7 +89,6 @@ extern void Register_PSROIAlign(OperatorRegistry *op_registry);
 extern void Register_ReOrganize(OperatorRegistry *op_registry);
 extern void Register_Reshape(OperatorRegistry *op_registry);
 extern void Register_ResizeBilinear(OperatorRegistry *op_registry);
-extern void Register_ScalarMath(OperatorRegistry *op_registry);
 extern void Register_Slice(OperatorRegistry *op_registry);
 extern void Register_Softmax(OperatorRegistry *op_registry);
 extern void Register_SpaceToBatchND(OperatorRegistry *op_registry);
@@ -127,7 +126,6 @@ OperatorRegistry::OperatorRegistry() {
   ops::Register_ReOrganize(this);
   ops::Register_Reshape(this);
   ops::Register_ResizeBilinear(this);
-  ops::Register_ScalarMath(this);
   ops::Register_Slice(this);
   ops::Register_Softmax(this);
   ops::Register_SpaceToBatchND(this);

mace/kernels/opencl/cl/scalar_math.cl

-#include <common.h>
-
-__kernel void scalar_math(__read_only image2d_t input, /* [c%4 * w * c/4, h * b] */
-                      __private const float scalar,
-                      __write_only image2d_t output) {
-  const int w = get_global_id(0);
-  const int hb = get_global_id(1);
-
-  DATA_TYPE4 in0 = READ_IMAGET(input, SAMPLER, (int2)(w, hb));
-  DATA_TYPE4 in1;
-  in1.x = scalar;
-  in1.y = scalar;
-  in1.z = scalar;
-  in1.w = scalar;  
-  DATA_TYPE4 out;
-#if SCALAR_MATH_TYPE == 1
-  out = in0 + in1;
-#elif SCALAR_MATH_TYPE == 4
-  out = in0 - in1;
-#elif SCALAR_MATH_TYPE == 0
-  out = in0 * in1;
-#elif SCALAR_MATH_TYPE == 5
-  out = in0 / in1;
-#endif
-
-  WRITE_IMAGET(output, (int2)(w, hb), out);
-}

mace/kernels/opencl/scalar_math_opencl.cc

-//
-// Copyright (c) 2017 XiaoMi All rights reserved.
-//
-
-#include "mace/kernels/scalar_math.h"
-#include "mace/core/runtime/opencl/opencl_runtime.h"
-#include "mace/kernels/opencl/helper.h"
-#include "mace/utils/tuner.h"
-
-namespace mace {
-namespace kernels {
-
-template <typename T>
-void ScalarMathFunctor<DeviceType::OPENCL, T>::operator()(const Tensor *input,
-                                                       Tensor *output,
-                                                       StatsFuture *future) {
-  const index_t batch = input->dim(0);
-  const index_t height = input->dim(1);
-  const index_t width = input->dim(2);
-  const index_t channels = input->dim(3);
-
-  const index_t channel_blocks = RoundUpDiv4(channels);
-  const index_t width_pixels = channel_blocks * width;
-  const index_t batch_height_pixels = batch * height;
-
-  if (kernel_.get() == nullptr) {
-    auto runtime = OpenCLRuntime::Global();
-    std::set<std::string> built_options;
-    auto dt = DataTypeToEnum<T>::value;
-    std::string kernel_name = MACE_OBFUSCATE_SYMBOL("scalar_math");
-    built_options.emplace("-Dscalar_math=" + kernel_name);
-    built_options.emplace("-DDATA_TYPE=" + DtToUpstreamCLDt(dt));
-    built_options.emplace("-DCMD_DATA_TYPE=" + DtToUpstreamCLCMDDt(dt));
-    built_options.emplace(MakeString("-DSCALAR_MATH_TYPE=", type_));
-    kernel_ = runtime->BuildKernel("scalar_math", kernel_name, built_options);
-  }
-  if (!IsVecEqual(input_shape_, input->shape())) {
-    uint32_t idx = 0;
-    kernel_.setArg(idx++, *(input->opencl_image()));
-    kernel_.setArg(idx++, static_cast<float>(coeff_));
-    kernel_.setArg(idx++, *(output->opencl_image()));
-    input_shape_ = input->shape();
-  }
-
-  const uint32_t gws[2] = {static_cast<uint32_t>(width_pixels),
-                           static_cast<uint32_t>(batch_height_pixels)};
-  const std::vector<uint32_t> lws = {64, 16, 1};
-  std::stringstream ss;
-  ss << "eltwise_opencl_kernel_" << output->dim(0) << "_" << output->dim(1)
-     << "_" << output->dim(2) << "_" << output->dim(3);
-  TuningOrRun2DKernel(kernel_, ss.str(), gws, lws, future);
-}
-
-template struct ScalarMathFunctor<DeviceType::OPENCL, float>;
-template struct ScalarMathFunctor<DeviceType::OPENCL, half>;
-}  // namespace kernels
-}  // namespace mace

mace/kernels/scalar_math.h

-//
-// Copyright (c) 2017 XiaoMi All rights reserved.
-//
-#ifndef MACE_KERNELS_SCALAR_MATH_H_
-#define MACE_KERNELS_SCALAR_MATH_H_
-
-#include <algorithm>
-#include <vector>
-
-#include "mace/core/future.h"
-#include "mace/core/runtime/opencl/cl2_header.h"
-#include "mace/core/tensor.h"
-
-namespace mace {
-namespace kernels {
-
-enum ScalarMathType {
-  MUL = 0,
-  ADD = 1,
-  MAX = 2,
-  MIN = 3,
-  SUB = 4,
-  DIV = 5,
-};
-
-struct ScalarMathFunctorBase {
-  ScalarMathFunctorBase(const ScalarMathType type, const float coeff)
-      : type_(type), coeff_(coeff) {}
-
-  ScalarMathType type_;
-  float coeff_;
-};
-
-template <DeviceType D, typename T>
-struct ScalarMathFunctor : ScalarMathFunctorBase {
-  ScalarMathFunctor(const ScalarMathType type, const float coeff)
-      : ScalarMathFunctorBase(type, coeff) {}
-
-  void operator()(const Tensor *input,
-                  Tensor *output,
-                  StatsFuture *future) {
-    Tensor::MappingGuard input_guard(input);
-    Tensor::MappingGuard output_guard(output);
-
-    const T *input_ptr = input->data<T>();
-    T *output_ptr = output->mutable_data<T>();
-    const index_t size = input->size();
-
-    switch (type_) {
-      case MUL:
-#pragma omp parallel for
-        for (index_t i = 0; i < size; ++i) {
-          output_ptr[i] =  coeff_ * input_ptr[i];
-        }
-        break;
-      case ADD:
-#pragma omp parallel for
-        for (index_t i = 0; i < size; ++i) {
-          output_ptr[i] =  coeff_ + input_ptr[i];
-        }
-        break;
-      case SUB:
-#pragma omp parallel for
-        for (index_t i = 0; i < size; ++i) {
-          output_ptr[i] =  input_ptr[i] - coeff_;
-        }
-        break;
-      case DIV:
-#pragma omp parallel for
-        for (index_t i = 0; i < size; ++i) {
-          output_ptr[i] =  input_ptr[i] / coeff_;
-        }
-        break;
-      default:
-        LOG(FATAL) << "ScalarMath op not support type " << type_;
-    }
-  }
-};
-
-template <typename T>
-struct ScalarMathFunctor<DeviceType::OPENCL, T> : ScalarMathFunctorBase {
-  ScalarMathFunctor(const ScalarMathType type, const float coeff)
-      : ScalarMathFunctorBase(type, coeff) {}
-
-  void operator()(const Tensor *input,
-                  Tensor *output,
-                  StatsFuture *future);
-
-  cl::Kernel kernel_;
-  std::vector<index_t> input_shape_;
-};
-
-}  // namespace kernels
-}  // namespace mace
-
-#endif  // MACE_KERNELS_SCALAR_MATH_H_

mace/ops/scalar_math.cc

-//
-// Copyright (c) 2017 XiaoMi All rights reserved.
-//
-
-#include "mace/ops/scalar_math.h"
-
-namespace mace {
-namespace ops {
-
-void Register_ScalarMath(OperatorRegistry *op_registry) {
-  REGISTER_OPERATOR(op_registry, OpKeyBuilder("ScalarMath")
-                                     .Device(DeviceType::CPU)
-                                     .TypeConstraint<float>("T")
-                                     .Build(),
-                    ScalarMathOp<DeviceType::CPU, float>);
-
-  REGISTER_OPERATOR(op_registry, OpKeyBuilder("ScalarMath")
-                                     .Device(DeviceType::OPENCL)
-                                     .TypeConstraint<float>("T")
-                                     .Build(),
-                    ScalarMathOp<DeviceType::OPENCL, float>);
-
-  REGISTER_OPERATOR(op_registry, OpKeyBuilder("ScalarMath")
-                                     .Device(DeviceType::OPENCL)
-                                     .TypeConstraint<half>("T")
-                                     .Build(),
-                    ScalarMathOp<DeviceType::OPENCL, half>);
-}
-
-}  // namespace ops
-}  // namespace mace

mace/ops/scalar_math.h

-//
-// Copyright (c) 2017 XiaoMi All rights reserved.
-//
-
-#ifndef MACE_OPS_SCALAR_MATH_H_
-#define MACE_OPS_SCALAR_MATH_H_
-
-#include <string>
-
-#include "mace/core/operator.h"
-#include "mace/kernels/scalar_math.h"
-
-namespace mace {
-namespace ops {
-
-template <DeviceType D, class T>
-class ScalarMathOp : public Operator<D, T> {
- public:
-  ScalarMathOp(const OperatorDef &operator_def, Workspace *ws)
-      : Operator<D, T>(operator_def, ws),
-        x_(OperatorBase::GetSingleArgument<float>("x", 1.0)),
-        functor_(static_cast<kernels::ScalarMathType>(
-                     OperatorBase::GetSingleArgument<int>(
-                         "type", static_cast<int>(
-                             kernels::ScalarMathType::ADD))),
-                 this->x_) {}
-
-  bool Run(StatsFuture *future) override {
-    const Tensor *input_tensor = this->Input(INPUT);
-    Tensor *output_tensor = this->Output(OUTPUT);
-    output_tensor->ResizeLike(input_tensor);
-
-    functor_(input_tensor, output_tensor, future);
-    return true;
-  }
-
- protected:
-  const float x_;
-  OP_INPUT_TAGS(INPUT);
-  OP_OUTPUT_TAGS(OUTPUT);
-
- private:
-  kernels::ScalarMathFunctor<D, T> functor_;
-};
-
-}  // namespace ops
-}  // namespace mace
-
-#endif  // MACE_OPS_SCALAR_MATH_H_

mace/ops/scalar_math_benchmark.cc

-//
-// Copyright (c) 2017 XiaoMi All rights reserved.
-//
-
-#include "mace/core/operator.h"
-#include "mace/core/runtime/opencl/opencl_runtime.h"
-#include "mace/core/testing/test_benchmark.h"
-#include "mace/ops/ops_test_util.h"
-
-namespace mace {
-namespace ops {
-namespace test {
-
-template <DeviceType D, typename T>
-static void ScalarMath(int iters, int batch, int channels,
-                       int height, int width, float x, int type) {
-  mace::testing::StopTiming();
-
-  OpsTestNet net;
-
-  // Add input data
-  net.AddRandomInput<D, T>("Input", {batch, height, width, channels});
-
-  if (D == DeviceType::OPENCL) {
-    BufferToImage<D, T>(&net, "Input", "InputImage",
-                        kernels::BufferType::IN_OUT_CHANNEL);
-    OpDefBuilder("ScalarMath", "ScalarMathBM")
-        .Input("InputImage")
-        .Output("Output")
-        .AddIntArg("type", type)
-        .AddFloatArg("x", x)
-        .Finalize(net.NewOperatorDef());
-  } else {
-    OpDefBuilder("ScalarMath", "ScalarMathBM")
-        .Input("Input")
-        .Output("Output")
-        .AddIntArg("type", type)
-        .AddFloatArg("x", x)
-        .Finalize(net.NewOperatorDef());
-  }
-
-  // Warm-up
-  for (int i = 0; i < 5; ++i) {
-    net.RunOp(D);
-  }
-  net.Sync();
-
-  mace::testing::StartTiming();
-  while (iters--) {
-    net.RunOp(D);
-  }
-  net.Sync();
-}
-
-#define BM_SCALAR_MATH_MACRO(N, C, H, W, X, G, TYPE, DEVICE)              \
-  static void                                                             \
-    BM_SCALAR_MATH_##N##_##C##_##H##_##W##_##X##_##G##_##TYPE##_##DEVICE( \
-      int iters) {                                                   \
-    const int64_t tot = static_cast<int64_t>(iters) * N * C * H * W; \
-    mace::testing::MaccProcessed(tot);                               \
-    mace::testing::BytesProcessed(tot *(sizeof(TYPE)));              \
-    ScalarMath<DEVICE, TYPE>(iters, N, C, H, W, X, G);               \
-  }                                                                  \
-  BENCHMARK(                                                              \
-    BM_SCALAR_MATH_##N##_##C##_##H##_##W##_##X##_##G##_##TYPE##_##DEVICE)
-
-#define BM_SCALAR_MATH(N, C, H, W, X, G)                 \
-  BM_SCALAR_MATH_MACRO(N, C, H, W, X, G, float, CPU);    \
-  BM_SCALAR_MATH_MACRO(N, C, H, W, X, G, float, OPENCL); \
-  BM_SCALAR_MATH_MACRO(N, C, H, W, X, G, half, OPENCL);
-
-BM_SCALAR_MATH(1, 1, 512, 512, 2, 0);
-BM_SCALAR_MATH(1, 3, 128, 128, 2, 1);
-BM_SCALAR_MATH(1, 3, 512, 512, 2, 4);
-BM_SCALAR_MATH(1, 32, 112, 112, 2, 5);
-BM_SCALAR_MATH(1, 64, 256, 256, 3, 0);
-BM_SCALAR_MATH(1, 64, 512, 512, 3, 1);
-BM_SCALAR_MATH(1, 128, 56, 56, 3, 4);
-BM_SCALAR_MATH(1, 128, 256, 256, 3, 5);
-BM_SCALAR_MATH(1, 256, 14, 14, 3, 0);
-BM_SCALAR_MATH(1, 512, 14, 14, 3, 1);
-BM_SCALAR_MATH(1, 1024, 7, 7, 3, 4);
-BM_SCALAR_MATH(32, 1, 256, 256, 3, 5);
-
-}  // namespace test
-}  // namespace ops
-}  // namespace mace

mace/ops/scalar_math_test.cc

-//
-// Copyright (c) 2017 XiaoMi All rights reserved.
-//
-
-#include "mace/core/operator.h"
-#include "mace/ops/ops_test_util.h"
-#include "../kernels/scalar_math.h"
-
-namespace mace {
-namespace ops {
-namespace test {
-
-class ScalarMathOpTest : public OpsTestBase {};
-
-
-template <DeviceType D>
-void Simple(const kernels::ScalarMathType type,
-            const std::vector<index_t> &shape,
-            const std::vector<float> &input0,
-            const float x,
-            const std::vector<float> &output) {
-  // Construct graph
-  OpsTestNet net;
-
-  // Add input data
-  net.AddInputFromArray<D, float>("Input1", shape, input0);
-
-  if (D == DeviceType::CPU) {
-    OpDefBuilder("ScalarMath", "ScalarMathTest")
-        .Input("Input1")
-        .AddIntArg("type", static_cast<int>(type))
-        .AddFloatArg("x", x)
-        .Output("Output")
-        .Finalize(net.NewOperatorDef());
-
-    // Run
-    net.RunOp(D);
-  } else {
-    BufferToImage<D, half>(&net, "Input1", "InputImg1",
-                           kernels::BufferType::IN_OUT_CHANNEL);
-    OpDefBuilder("ScalarMath", "ScalarMathTest")
-        .Input("InputImg1")
-        .AddIntArg("type", static_cast<int>(type))
-        .AddFloatArg("x", x)
-        .Output("OutputImg")
-        .Finalize(net.NewOperatorDef());
-
-    // Run
-    net.RunOp(D);
-
-    ImageToBuffer<D, float>(&net, "OutputImg", "Output",
-                            kernels::BufferType::IN_OUT_CHANNEL);
-  }
-
-  auto expected = CreateTensor<float>(shape, output);
-
-  ExpectTensorNear<float>(*expected, *net.GetOutput("Output"), 1e-3);
-}
-
-TEST_F(ScalarMathOpTest, CPUSimple) {
-  Simple<DeviceType::CPU>(kernels::ScalarMathType::MUL, {1, 1, 2, 3},
-                          {1, 2, 3, 4, 5, 6}, 0.1, {0.1, 0.2, .3, .4, .5, .6});
-
-  Simple<DeviceType::CPU>(kernels::ScalarMathType::ADD, {1, 1, 2, 3},
-                          {1, 2, 3, 4, 5, 6}, 2.0, {3, 4, 5, 6, 7, 8});
-
-  Simple<DeviceType::CPU>(kernels::ScalarMathType::DIV, {1, 1, 2, 3},
-                          {1, 2, 3, 4, 5, 6}, 0.1, {10, 20, 30, 40, 50, 60});
-
-  Simple<DeviceType::CPU>(kernels::ScalarMathType::SUB, {1, 1, 2, 3},
-                          {1, 2, 3, 4, 5, 6}, 2.0, {-1, 0, 1, 2, 3, 4});
-}
-
-TEST_F(ScalarMathOpTest, GPUSimple) {
-  Simple<DeviceType::OPENCL>(kernels::ScalarMathType::MUL, {1, 1, 2, 3},
-                          {1, 2, 3, 4, 5, 6}, 0.1, {0.1, 0.2, .3, .4, .5, .6});
-
-  Simple<DeviceType::OPENCL>(kernels::ScalarMathType::ADD, {1, 1, 2, 3},
-                          {1, 2, 3, 4, 5, 6}, 2.0, {3, 4, 5, 6, 7, 8});
-
-  Simple<DeviceType::OPENCL>(kernels::ScalarMathType::DIV, {1, 1, 2, 3},
-                          {1, 2, 3, 4, 5, 6}, 0.1, {10, 20, 30, 40, 50, 60});
-
-  Simple<DeviceType::OPENCL>(kernels::ScalarMathType::SUB, {1, 1, 2, 3},
-                          {1, 2, 3, 4, 5, 6}, 2.0, {-1, 0, 1, 2, 3, 4});
-}
-
-template <DeviceType D, typename T>
-void RandomTest(const kernels::ScalarMathType type,
-                const std::vector<index_t> &shape) {
-  testing::internal::LogToStderr();
-  srand(time(NULL));
-
-  // Construct graph
-  OpsTestNet net;
-
-  // Add input data
-  net.AddRandomInput<D, float>("Input1", shape);
-
-  OpDefBuilder("ScalarMath", "ScalarMathTest")
-      .Input("Input1")
-      .AddIntArg("type", static_cast<int>(type))
-      .AddFloatArg("x", 1.2)
-      .Output("Output")
-      .Finalize(net.NewOperatorDef());
-
-  // Run
-  net.RunOp();
-
-  BufferToImage<D, T>(&net, "Input1", "InputImg1",
-                      kernels::BufferType::IN_OUT_CHANNEL);
-
-  OpDefBuilder("ScalarMath", "ScalarMathTest")
-      .Input("InputImg1")
-      .AddIntArg("type", static_cast<int>(type))
-      .AddFloatArg("x", 1.2)
-      .AddIntArg("T", static_cast<int>(DataTypeToEnum<T>::value))
-      .Output("OutputImg")
-      .Finalize(net.NewOperatorDef());
-
-  // Run
-  net.RunOp(D);
-
-  ImageToBuffer<D, float>(&net, "OutputImg", "OPENCLOutput",
-                          kernels::BufferType::IN_OUT_CHANNEL);
-
-  if (DataTypeToEnum<T>::value == DT_FLOAT) {
-    ExpectTensorNear<float>(*net.GetTensor("Output"),
-                            *net.GetOutput("OPENCLOutput"), 1e-3);
-  } else {
-    ExpectTensorNear<float>(*net.GetTensor("Output"),
-                            *net.GetOutput("OPENCLOutput"), 1e-1);
-  }
-}
-
-TEST_F(ScalarMathOpTest, OPENCLRandomFloat) {
-  RandomTest<DeviceType::OPENCL, float>(kernels::ScalarMathType::MUL,
-                                        {3, 23, 37, 19});
-  RandomTest<DeviceType::OPENCL, float>(kernels::ScalarMathType::ADD,
-                                        {13, 32, 32, 64});
-  RandomTest<DeviceType::OPENCL, float>(kernels::ScalarMathType::SUB,
-                                        {3, 32, 32, 64});
-  RandomTest<DeviceType::OPENCL, float>(kernels::ScalarMathType::DIV,
-                                        {13, 32, 32, 64});
-}
-
-TEST_F(ScalarMathOpTest, OPENCLRandomHalf) {
-  RandomTest<DeviceType::OPENCL, half>(kernels::ScalarMathType::MUL,
-                                       {3, 23, 37, 19});
-  RandomTest<DeviceType::OPENCL, half>(kernels::ScalarMathType::ADD,
-                                       {13, 32, 32, 64});
-  RandomTest<DeviceType::OPENCL, half>(kernels::ScalarMathType::SUB,
-                                       {3, 32, 32, 64});
-  RandomTest<DeviceType::OPENCL, half>(kernels::ScalarMathType::DIV,
-                                       {13, 32, 32, 64});
-}
-
-}  // namespace test
-}  // namespace ops
-}  // namespace mace