Skip to content

Mace
Mace

Xiaomi / Mace

通知 106
Star 40
Fork 27
Mace
Mace

前往新版Gitcode，体验更适合开发者的 AI 搜索 >>

提交 b90f0ad4 编写于 作者: U Unknown 提交者： liutuo
浏览文件
操作

delete neg

上级 7d301f24
隐藏空白更改
内联 并排
Showing with 0 addition and 363 deletion
mace/core/operator.cc
浏览文件 @ b90f0ad4
......@@ -83,7 +83,6 @@ extern void Register_FusedConv2D(OperatorRegistry *op_registry);
extern void Register_GlobalAvgPooling(OperatorRegistry *op_registry);
extern void Register_ImageToBuffer(OperatorRegistry *op_registry);
extern void Register_MatMul(OperatorRegistry *op_registry);
extern void Register_Neg(OperatorRegistry *op_registry);
extern void Register_Pooling(OperatorRegistry *op_registry);
extern void Register_Proposal(OperatorRegistry *op_registry);
extern void Register_PSROIAlign(OperatorRegistry *op_registry);
......@@ -122,7 +121,6 @@ OperatorRegistry::OperatorRegistry() {
ops::Register_GlobalAvgPooling(this);
ops::Register_ImageToBuffer(this);
ops::Register_MatMul(this);
ops::Register_Neg(this);
ops::Register_Pooling(this);
ops::Register_Proposal(this);
ops::Register_PSROIAlign(this);
......
mace/kernels/negative.h 已删除 100644 → 0
浏览文件 @ 7d301f24
//
// Copyright (c) 2017 XiaoMi All rights reserved.
//
#ifndef MACE_KERNELS_NEGATIVE_H_
#define MACE_KERNELS_NEGATIVE_H_
#include <vector>
#include "mace/core/future.h"
#include "mace/core/runtime/opencl/cl2_header.h"
#include "mace/core/tensor.h"
#include "mace/public/mace.h"
namespace mace {
namespace kernels {
template <DeviceType D, typename T>
struct NegFunctor {
void 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);
Tensor::MappingGuard input_mapper(input);
Tensor::MappingGuard output_mapper(output);
const T *input_ptr = input->data<T>();
T *output_ptr = output->mutable_data<T>();
#pragma omp parallel for collapse(4)
for (index_t n = 0; n < batch; ++n) {
for (index_t h = 0; h < height; ++h) {
for (index_t w = 0; w < width; ++w) {
for (index_t c = 0; c < channels; ++c) {
index_t pos = (((n * height) + h) * width + w) * channels + c;
output_ptr[pos] = 0 - input_ptr[pos];
}
}
}
}
}
};
/*
template <>
void NegFunctor<DeviceType::NEON, float>::operator()(
const Tensor *input,
const Tensor *bias,
Tensor *output,
StatsFuture *future);
*/
template <typename T>
struct NegFunctor<DeviceType::OPENCL, T> {
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_NEGATIVE_H_
mace/kernels/opencl/cl/neg.cl 已删除 100644 → 0
浏览文件 @ 7d301f24
#include <common.h>
// Supported data types: half/float
__kernel void neg(__read_only image2d_t input,
__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);
const int width = get_global_size(1);
const int pos = mad24(ch_blk, width, w);
DATA_TYPE4 in = READ_IMAGET(input, SAMPLER, (int2)(pos, hb));
DATA_TYPE4 out = 0 - in;
WRITE_IMAGET(output, (int2)(pos, hb), out);
}
mace/kernels/opencl/negative_opencl.cc 已删除 100644 → 0
浏览文件 @ 7d301f24
//
// Copyright (c) 2017 XiaoMi All rights reserved.
//
#include "mace/kernels/negative.h"
#include "mace/core/runtime/opencl/cl2_header.h"
#include "mace/core/runtime/opencl/opencl_runtime.h"
#include "mace/kernels/opencl/helper.h"
#include "mace/utils/utils.h"
namespace mace {
namespace kernels {
template <typename T>
void NegFunctor<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);
auto runtime = OpenCLRuntime::Global();
if (kernel_.get() == nullptr) {
std::set<std::string> built_options;
auto dt = DataTypeToEnum<T>::value;
std::string kernel_name = MACE_OBFUSCATE_SYMBOL("neg");
built_options.emplace("-Dneg=" + kernel_name);
built_options.emplace("-DDATA_TYPE=" + DtToUpstreamCLDt(dt));
built_options.emplace("-DCMD_DATA_TYPE=" + DtToUpstreamCLCMDDt(dt));
kernel_ = runtime->BuildKernel("neg", kernel_name, built_options);
}
if (!IsVecEqual(input_shape_, input->shape())) {
uint32_t idx = 0;
kernel_.setArg(idx++, *(input->opencl_image()));
kernel_.setArg(idx++, *(output->opencl_image()));
input_shape_ = input->shape();
}
const uint32_t gws[3] = {static_cast<uint32_t>(channel_blocks),
static_cast<uint32_t>(width),
static_cast<uint32_t>(height * batch)};
const std::vector<uint32_t> lws = {8, 16, 8};
cl::Event event;
cl_int error = runtime->command_queue().enqueueNDRangeKernel(
kernel_, cl::NullRange, cl::NDRange(gws[0], gws[1], gws[2]),
cl::NDRange(lws[0], lws[1], lws[2]), nullptr, &event);
MACE_CHECK(error == CL_SUCCESS);
if (future != nullptr) {
future->wait_fn = [runtime, event](CallStats *stats) {
event.wait();
if (stats != nullptr) {
runtime->GetCallStats(event, stats);
}
};
}
}
template struct NegFunctor<DeviceType::OPENCL, float>;
template struct NegFunctor<DeviceType::OPENCL, half>;
} // namespace kernels
} // namespace mace
mace/ops/neg.cc 已删除 100644 → 0
浏览文件 @ 7d301f24
//
// Copyright (c) 2017 XiaoMi All rights reserved.
//
#include "mace/ops/neg.h"
namespace mace {
namespace ops {
void Register_Neg(OperatorRegistry *op_registry) {
REGISTER_OPERATOR(op_registry, OpKeyBuilder("Neg")
.Device(DeviceType::CPU)
.TypeConstraint<float>("T")
.Build(),
NegOp<DeviceType::CPU, float>);
REGISTER_OPERATOR(op_registry, OpKeyBuilder("Neg")
.Device(DeviceType::OPENCL)
.TypeConstraint<float>("T")
.Build(),
NegOp<DeviceType::OPENCL, float>);
REGISTER_OPERATOR(op_registry, OpKeyBuilder("Neg")
.Device(DeviceType::OPENCL)
.TypeConstraint<half>("T")
.Build(),
NegOp<DeviceType::OPENCL, half>);
}
} // namespace ops
} // namespace mace
mace/ops/neg.h 已删除 100644 → 0
浏览文件 @ 7d301f24
//
// Copyright (c) 2017 XiaoMi All rights reserved.
//
#ifndef MACE_OPS_NEG_H_
#define MACE_OPS_NEG_H_
#include <string>
#include "mace/core/operator.h"
#include "mace/kernels/negative.h"
namespace mace {
namespace ops {
template <DeviceType D, class T>
class NegOp : public Operator<D, T> {
public:
NegOp(const OperatorDef &operator_def, Workspace *ws)
: Operator<D, T>(operator_def, ws),
functor_() {}
bool Run(StatsFuture *future) override {
const Tensor *input_tensor = this->Input(0);
Tensor *output_tensor = this->outputs_[0];
output_tensor->ResizeLike(input_tensor);
functor_(input_tensor, output_tensor, future);
return true;
}
private:
kernels::NegFunctor<D, T> functor_;
};
} // namespace ops
} // namespace mace
#endif // MACE_OPS_NEGATIVE_H_
mace/ops/neg_benchmark.cc 已删除 100644 → 0
浏览文件 @ 7d301f24
//
// 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 Neg(int iters, int batch, int channels, int height, int width) {
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("Neg", "NegBM")
.Input("InputImage")
.Output("Output")
.Finalize(net.NewOperatorDef());
} else {
OpDefBuilder("Neg", "NegBM")
.Input("Input")
.Output("Output")
.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_NEG_MACRO(N, C, H, W, TYPE, DEVICE) \
static void BM_NEG_##N##_##C##_##H##_##W##_##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))); \
Neg<DEVICE, TYPE>(iters, N, C, H, W); \
} \
BENCHMARK(BM_NEG_##N##_##C##_##H##_##W##_##TYPE##_##DEVICE)
#define BM_NEG(N, C, H, W) \
BM_NEG_MACRO(N, C, H, W, float, CPU); \
BM_NEG_MACRO(N, C, H, W, float, OPENCL); \
BM_NEG_MACRO(N, C, H, W, half, OPENCL);
BM_NEG(1, 1, 512, 512);
BM_NEG(1, 3, 128, 128);
BM_NEG(1, 3, 512, 512);
BM_NEG(1, 32, 112, 112);
BM_NEG(1, 64, 256, 256);
BM_NEG(1, 64, 512, 512);
BM_NEG(1, 128, 56, 56);
BM_NEG(1, 128, 256, 256);
BM_NEG(1, 256, 14, 14);
BM_NEG(1, 512, 14, 14);
BM_NEG(1, 1024, 7, 7);
BM_NEG(32, 1, 256, 256);
BM_NEG(32, 3, 256, 256);
} // namespace test
} // namespace ops
} // namespace mace
mace/ops/neg_test.cc 已删除 100644 → 0
浏览文件 @ 7d301f24
//
// Copyright (c) 2017 XiaoMi All rights reserved.
//
#include "mace/core/operator.h"
#include "mace/ops/ops_test_util.h"
namespace mace {
namespace ops {
namespace test {
class NegOpTest : public OpsTestBase {};
template <DeviceType D>
void NegSimple() {
OpsTestNet net;
// Add input data
net.AddInputFromArray<D, float>("Input", {1, 6, 2, 1},
{5, 5, 7, 7, 9, 9, 11, 11, 13, 13, 15, 15});
if (D == DeviceType::OPENCL) {
BufferToImage<D, float>(&net, "Input", "InputImage",
kernels::BufferType::IN_OUT_CHANNEL);
OpDefBuilder("Neg", "NegTest")
.Input("InputImage")
.Output("OutputImage")
.Finalize(net.NewOperatorDef());
// Run
net.RunOp(D);
// Transfer output
ImageToBuffer<D, float>(&net, "OutputImage", "Output",
kernels::BufferType::IN_OUT_CHANNEL);
} else {
OpDefBuilder("Neg", "NegTest")
.Input("Input")
.Output("Output")
.Finalize(net.NewOperatorDef());
// Run
net.RunOp(D);
}
// Check
auto expected = CreateTensor<float>(
{1, 6, 2, 1},
{-5, -5, -7, -7, -9, -9, -11, -11, -13, -13, -15, -15});
ExpectTensorNear<float>(*expected, *net.GetOutput("Output"), 1e-8);
}
TEST_F(NegOpTest, NegSimpleCPU) { NegSimple<DeviceType::CPU>(); }
TEST_F(NegOpTest, NegSimpleOPENCL) {
NegSimple<DeviceType::OPENCL>();
}
} // namespace test
} // namespace ops
} // namespace mace
Markdown is supported
0% 或 .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑！
想要评论请 注册 或