diff --git a/mace/core/operator.cc b/mace/core/operator.cc index dcbd9f7be8c2e252f0e16003a185cff2888fe499..60eabfdaa0a8de900690e2ef87e9719ead07d639 100644 --- a/mace/core/operator.cc +++ b/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); diff --git a/mace/kernels/opencl/cl/scalar_math.cl b/mace/kernels/opencl/cl/scalar_math.cl deleted file mode 100644 index 19678b08e7f8d5a3fb63156f06a100aee1bd80b4..0000000000000000000000000000000000000000 --- a/mace/kernels/opencl/cl/scalar_math.cl +++ /dev/null @@ -1,27 +0,0 @@ -#include - -__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); -} diff --git a/mace/kernels/opencl/scalar_math_opencl.cc b/mace/kernels/opencl/scalar_math_opencl.cc deleted file mode 100644 index 42ad751829d7f1810e5bdb793adc5439cd50f6b3..0000000000000000000000000000000000000000 --- a/mace/kernels/opencl/scalar_math_opencl.cc +++ /dev/null @@ -1,57 +0,0 @@ -// -// 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 -void ScalarMathFunctor::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 built_options; - auto dt = DataTypeToEnum::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(coeff_)); - kernel_.setArg(idx++, *(output->opencl_image())); - input_shape_ = input->shape(); - } - - const uint32_t gws[2] = {static_cast(width_pixels), - static_cast(batch_height_pixels)}; - const std::vector 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; -template struct ScalarMathFunctor; -} // namespace kernels -} // namespace mace diff --git a/mace/kernels/scalar_math.h b/mace/kernels/scalar_math.h deleted file mode 100644 index 75c5286757173124df730e94ef7d0d7517b9605a..0000000000000000000000000000000000000000 --- a/mace/kernels/scalar_math.h +++ /dev/null @@ -1,96 +0,0 @@ -// -// Copyright (c) 2017 XiaoMi All rights reserved. -// -#ifndef MACE_KERNELS_SCALAR_MATH_H_ -#define MACE_KERNELS_SCALAR_MATH_H_ - -#include -#include - -#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 -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 *output_ptr = output->mutable_data(); - 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 -struct ScalarMathFunctor : 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 input_shape_; -}; - -} // namespace kernels -} // namespace mace - -#endif // MACE_KERNELS_SCALAR_MATH_H_ diff --git a/mace/ops/scalar_math.cc b/mace/ops/scalar_math.cc deleted file mode 100644 index 9891994f5fa2e86b0db0b6f8590fd33a89752d7d..0000000000000000000000000000000000000000 --- a/mace/ops/scalar_math.cc +++ /dev/null @@ -1,31 +0,0 @@ -// -// 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("T") - .Build(), - ScalarMathOp); - - REGISTER_OPERATOR(op_registry, OpKeyBuilder("ScalarMath") - .Device(DeviceType::OPENCL) - .TypeConstraint("T") - .Build(), - ScalarMathOp); - - REGISTER_OPERATOR(op_registry, OpKeyBuilder("ScalarMath") - .Device(DeviceType::OPENCL) - .TypeConstraint("T") - .Build(), - ScalarMathOp); -} - -} // namespace ops -} // namespace mace diff --git a/mace/ops/scalar_math.h b/mace/ops/scalar_math.h deleted file mode 100644 index 2f0f4394e92b1f1d8abe15a7201e68ddf3b88fc1..0000000000000000000000000000000000000000 --- a/mace/ops/scalar_math.h +++ /dev/null @@ -1,49 +0,0 @@ -// -// Copyright (c) 2017 XiaoMi All rights reserved. -// - -#ifndef MACE_OPS_SCALAR_MATH_H_ -#define MACE_OPS_SCALAR_MATH_H_ - -#include - -#include "mace/core/operator.h" -#include "mace/kernels/scalar_math.h" - -namespace mace { -namespace ops { - -template -class ScalarMathOp : public Operator { - public: - ScalarMathOp(const OperatorDef &operator_def, Workspace *ws) - : Operator(operator_def, ws), - x_(OperatorBase::GetSingleArgument("x", 1.0)), - functor_(static_cast( - OperatorBase::GetSingleArgument( - "type", static_cast( - 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 functor_; -}; - -} // namespace ops -} // namespace mace - -#endif // MACE_OPS_SCALAR_MATH_H_ diff --git a/mace/ops/scalar_math_benchmark.cc b/mace/ops/scalar_math_benchmark.cc deleted file mode 100644 index e01d39e03943adea2dd6f0a7642e9b184a49db91..0000000000000000000000000000000000000000 --- a/mace/ops/scalar_math_benchmark.cc +++ /dev/null @@ -1,87 +0,0 @@ -// -// 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 -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("Input", {batch, height, width, channels}); - - if (D == DeviceType::OPENCL) { - BufferToImage(&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(iters) * N * C * H * W; \ - mace::testing::MaccProcessed(tot); \ - mace::testing::BytesProcessed(tot *(sizeof(TYPE))); \ - ScalarMath(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 diff --git a/mace/ops/scalar_math_test.cc b/mace/ops/scalar_math_test.cc deleted file mode 100644 index 3da6176d82e5b5e183615548a5000dffeab2462b..0000000000000000000000000000000000000000 --- a/mace/ops/scalar_math_test.cc +++ /dev/null @@ -1,160 +0,0 @@ -// -// 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 -void Simple(const kernels::ScalarMathType type, - const std::vector &shape, - const std::vector &input0, - const float x, - const std::vector &output) { - // Construct graph - OpsTestNet net; - - // Add input data - net.AddInputFromArray("Input1", shape, input0); - - if (D == DeviceType::CPU) { - OpDefBuilder("ScalarMath", "ScalarMathTest") - .Input("Input1") - .AddIntArg("type", static_cast(type)) - .AddFloatArg("x", x) - .Output("Output") - .Finalize(net.NewOperatorDef()); - - // Run - net.RunOp(D); - } else { - BufferToImage(&net, "Input1", "InputImg1", - kernels::BufferType::IN_OUT_CHANNEL); - OpDefBuilder("ScalarMath", "ScalarMathTest") - .Input("InputImg1") - .AddIntArg("type", static_cast(type)) - .AddFloatArg("x", x) - .Output("OutputImg") - .Finalize(net.NewOperatorDef()); - - // Run - net.RunOp(D); - - ImageToBuffer(&net, "OutputImg", "Output", - kernels::BufferType::IN_OUT_CHANNEL); - } - - auto expected = CreateTensor(shape, output); - - ExpectTensorNear(*expected, *net.GetOutput("Output"), 1e-3); -} - -TEST_F(ScalarMathOpTest, CPUSimple) { - Simple(kernels::ScalarMathType::MUL, {1, 1, 2, 3}, - {1, 2, 3, 4, 5, 6}, 0.1, {0.1, 0.2, .3, .4, .5, .6}); - - Simple(kernels::ScalarMathType::ADD, {1, 1, 2, 3}, - {1, 2, 3, 4, 5, 6}, 2.0, {3, 4, 5, 6, 7, 8}); - - Simple(kernels::ScalarMathType::DIV, {1, 1, 2, 3}, - {1, 2, 3, 4, 5, 6}, 0.1, {10, 20, 30, 40, 50, 60}); - - Simple(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(kernels::ScalarMathType::MUL, {1, 1, 2, 3}, - {1, 2, 3, 4, 5, 6}, 0.1, {0.1, 0.2, .3, .4, .5, .6}); - - Simple(kernels::ScalarMathType::ADD, {1, 1, 2, 3}, - {1, 2, 3, 4, 5, 6}, 2.0, {3, 4, 5, 6, 7, 8}); - - Simple(kernels::ScalarMathType::DIV, {1, 1, 2, 3}, - {1, 2, 3, 4, 5, 6}, 0.1, {10, 20, 30, 40, 50, 60}); - - Simple(kernels::ScalarMathType::SUB, {1, 1, 2, 3}, - {1, 2, 3, 4, 5, 6}, 2.0, {-1, 0, 1, 2, 3, 4}); -} - -template -void RandomTest(const kernels::ScalarMathType type, - const std::vector &shape) { - testing::internal::LogToStderr(); - srand(time(NULL)); - - // Construct graph - OpsTestNet net; - - // Add input data - net.AddRandomInput("Input1", shape); - - OpDefBuilder("ScalarMath", "ScalarMathTest") - .Input("Input1") - .AddIntArg("type", static_cast(type)) - .AddFloatArg("x", 1.2) - .Output("Output") - .Finalize(net.NewOperatorDef()); - - // Run - net.RunOp(); - - BufferToImage(&net, "Input1", "InputImg1", - kernels::BufferType::IN_OUT_CHANNEL); - - OpDefBuilder("ScalarMath", "ScalarMathTest") - .Input("InputImg1") - .AddIntArg("type", static_cast(type)) - .AddFloatArg("x", 1.2) - .AddIntArg("T", static_cast(DataTypeToEnum::value)) - .Output("OutputImg") - .Finalize(net.NewOperatorDef()); - - // Run - net.RunOp(D); - - ImageToBuffer(&net, "OutputImg", "OPENCLOutput", - kernels::BufferType::IN_OUT_CHANNEL); - - if (DataTypeToEnum::value == DT_FLOAT) { - ExpectTensorNear(*net.GetTensor("Output"), - *net.GetOutput("OPENCLOutput"), 1e-3); - } else { - ExpectTensorNear(*net.GetTensor("Output"), - *net.GetOutput("OPENCLOutput"), 1e-1); - } -} - -TEST_F(ScalarMathOpTest, OPENCLRandomFloat) { - RandomTest(kernels::ScalarMathType::MUL, - {3, 23, 37, 19}); - RandomTest(kernels::ScalarMathType::ADD, - {13, 32, 32, 64}); - RandomTest(kernels::ScalarMathType::SUB, - {3, 32, 32, 64}); - RandomTest(kernels::ScalarMathType::DIV, - {13, 32, 32, 64}); -} - -TEST_F(ScalarMathOpTest, OPENCLRandomHalf) { - RandomTest(kernels::ScalarMathType::MUL, - {3, 23, 37, 19}); - RandomTest(kernels::ScalarMathType::ADD, - {13, 32, 32, 64}); - RandomTest(kernels::ScalarMathType::SUB, - {3, 32, 32, 64}); - RandomTest(kernels::ScalarMathType::DIV, - {13, 32, 32, 64}); -} - -} // namespace test -} // namespace ops -} // namespace mace