Skip to content

Mace
Mace

Xiaomi / Mace

通知 107
Star 40
Fork 27
Mace
Mace
提交 190385bb 编写于 作者: 刘 刘琦
浏览文件
操作

Revert "Merge branch 'resize_bicubic' into 'master'" 

This reverts merge request !748
上级 e46c5a39
隐藏空白更改
内联 并排
Showing with 0 addition and 897 deletion
mace/kernels/opencl/cl/resize_bicubic.cl 已删除 100644 → 0
浏览文件 @ e46c5a39
#include <common.h>
//#include <cstdio>
const int kTableSize = (1 << 10);
inline float ComputeCoeffs(int i) {
const float A = -0.75;
float x = (i / 2) * 1.0 / kTableSize;
if (i % 2 == 0){
float coeff = ((A + 2) * x - (A + 3)) * x * x + 1;
return coeff;
}
else {
x += 1.0;
float coeff = ((A * x - 5 * A) * x + 8 * A) * x - 4 * A;
return coeff;
}
}
#define BOUND(val, limit) min(limit - 1, max(0, val))
__kernel void resize_bicubic_nocache(KERNEL_ERROR_PARAMS
GLOBAL_WORK_GROUP_SIZE_DIM3
__read_only image2d_t input,
__write_only image2d_t output,
__private const float height_scale,
__private const float width_scale,
__private const int in_height,
__private const int in_width,
__private const int out_height) {
const int ch_blk = get_global_id(0);
const int w = get_global_id(1);
const int hb = get_global_id(2);
#ifndef NON_UNIFORM_WORK_GROUP
if (ch_blk >= global_size_dim0 || w >= global_size_dim1
|| hb >= global_size_dim2) {
return;
}
const int ch_blks = global_size_dim0;
const int out_width = global_size_dim1;
#else
const int ch_blks = get_global_size(0);
const int out_width = get_global_size(1);
#endif
const int b = hb / out_height;
const int h = hb % out_height;
const float h_in = h * height_scale;
const float w_in = w * width_scale;
const int in_w_offset = mul24(ch_blk, in_width);
const int in_h_offset = mul24(b, in_height);
const int h_in_loc = height_scale * h;
const float h_delta = height_scale * h - h_in_loc;
const int h_offset = h_delta * kTableSize + 0.5;
const int w_in_loc = width_scale * w;
const float w_delta = width_scale * w - w_in_loc;
const int w_offset = w_delta * kTableSize + 0.5;
float4 y_weights = {ComputeCoeffs(h_offset * 2 + 1),
ComputeCoeffs(h_offset * 2),
ComputeCoeffs((kTableSize - h_offset) * 2),
ComputeCoeffs((kTableSize - h_offset) * 2 + 1)};
int4 y_indices = {BOUND(h_in_loc - 1, in_height),
BOUND(h_in_loc, in_height),
BOUND(h_in_loc + 1, in_height),
BOUND(h_in_loc + 2, in_height)};
float4 x_weights = {ComputeCoeffs(w_offset * 2 + 1),
ComputeCoeffs(w_offset * 2),
ComputeCoeffs((kTableSize - w_offset) * 2),
ComputeCoeffs((kTableSize - w_offset) * 2 + 1)};
int4 x_indices = {BOUND(w_in_loc - 1, in_width),
BOUND(w_in_loc, in_width),
BOUND(w_in_loc + 1, in_width),
BOUND(w_in_loc + 2, in_width)};
float4 coeffs0 = {0, 0, 0, 0};
float4 coeffs1 = {0, 0, 0, 0};
float4 coeffs2 = {0, 0, 0, 0};
float4 coeffs3 = {0, 0, 0, 0};
for (int i = 0; i < 4; ++i) {
int y_index = y_indices.s0;
if ( i == 1 ) { y_index = y_indices.s1; }
if ( i == 2 ) { y_index = y_indices.s2; }
if ( i == 3 ) { y_index = y_indices.s3; }
DATA_TYPE4 data0 = READ_IMAGET(input, SAMPLER,
(int2)(in_w_offset + x_indices.s0, in_h_offset + y_index));
DATA_TYPE4 data1 = READ_IMAGET(input, SAMPLER,
(int2)(in_w_offset + x_indices.s1, in_h_offset + y_index));
DATA_TYPE4 data2 = READ_IMAGET(input, SAMPLER,
(int2)(in_w_offset + x_indices.s2, in_h_offset + y_index));
DATA_TYPE4 data3 = READ_IMAGET(input, SAMPLER,
(int2)(in_w_offset + x_indices.s3, in_h_offset + y_index));
float4 xw0 = { x_weights.s0, x_weights.s0, x_weights.s0, x_weights.s0 };
float4 xw1 = { x_weights.s1, x_weights.s1, x_weights.s1, x_weights.s1 };
float4 xw2 = { x_weights.s2, x_weights.s2, x_weights.s2, x_weights.s2 };
float4 xw3 = { x_weights.s3, x_weights.s3, x_weights.s3, x_weights.s3 };
float4 res = { 0, 0, 0, 0 };
res = mad(xw0, data0, res);
res = mad(xw1, data1, res);
res = mad(xw2, data2, res);
res = mad(xw3, data3, res);
if ( i == 0 ) { coeffs0 = res; }
if ( i == 1 ) { coeffs1 = res; }
if ( i == 2 ) { coeffs2 = res; }
if ( i == 3 ) { coeffs3 = res; }
}
float4 yw0 = { y_weights.s0, y_weights.s0, y_weights.s0, y_weights.s0 };
float4 yw1 = { y_weights.s1, y_weights.s1, y_weights.s1, y_weights.s1 };
float4 yw2 = { y_weights.s2, y_weights.s2, y_weights.s2, y_weights.s2 };
float4 yw3 = { y_weights.s3, y_weights.s3, y_weights.s3, y_weights.s3 };
DATA_TYPE4 outdata = { 0, 0, 0, 0 };
outdata = mad(yw0, coeffs0, outdata);
outdata = mad(yw1, coeffs1, outdata);
outdata = mad(yw2, coeffs2, outdata);
outdata = mad(yw3, coeffs3, outdata);
const int out_w_offset = mul24(ch_blk, out_width);
const int out_h_offset = mul24(b, out_height);
WRITE_IMAGET(output, (int2)(out_w_offset + w, out_h_offset + h), outdata);
}
mace/kernels/opencl/resize_bicubic.cc 已删除 100644 → 0
浏览文件 @ e46c5a39
// Copyright 2018 Xiaomi, Inc. All rights reserved.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include "mace/kernels/resize_bicubic.h"
#include "mace/core/runtime/opencl/opencl_runtime.h"
#include "mace/core/tensor.h"
#include "mace/kernels/opencl/helper.h"
#include "mace/utils/tuner.h"
#include "mace/utils/utils.h"
namespace mace {
namespace kernels {
namespace {
std::vector<uint32_t> LocalWS(const uint32_t *gws, const uint32_t kwg_size) {
std::vector<uint32_t> lws(4, 0);
uint64_t cache_size = OpenCLRuntime::Global()->device_global_mem_cache_size();
uint32_t base = std::max<uint32_t>(cache_size / kBaseGPUMemCacheSize, 1);
lws[1] = std::min<uint32_t>(gws[1], kwg_size);
if (lws[1] >= base) {
lws[0] = std::min<uint32_t>(gws[0], base);
} else {
lws[0] = gws[0] / 8;
if (lws[0] == 0) {
lws[0] = gws[0];
}
}
lws[0] = std::min<uint32_t>(lws[0], kwg_size / lws[1]);
const uint32_t lws_size = lws[0] * lws[1];
lws[2] = gws[2] / 8;
if (lws[2] == 0) {
lws[2] = gws[2];
}
lws[2] = std::max<uint32_t>(std::min<uint32_t>(lws[2], kwg_size / lws_size),
1);
return lws;
}
} // namespace
template <typename T>
MaceStatus ResizeBicubicFunctor<DeviceType::GPU, T>::operator()(
const Tensor *input, Tensor *output, StatsFuture *future) {
const index_t batch = input->dim(0);
const index_t in_height = input->dim(1);
const index_t in_width = input->dim(2);
const index_t channels = input->dim(3);
const index_t channel_blocks = RoundUpDiv4(channels);
const index_t out_height = out_height_;
const index_t out_width = out_width_;
const uint32_t gws[3] = {static_cast<uint32_t>(channel_blocks),
static_cast<uint32_t>(out_width),
static_cast<uint32_t>(out_height * batch)};
auto runtime = OpenCLRuntime::Global();
if (kernel_.get() == nullptr) {
std::set<std::string> built_options;
OUT_OF_RANGE_CONFIG(kernel_error_);
NON_UNIFORM_WG_CONFIG;
std::string kernel_name = MACE_OBFUSCATE_SYMBOL("resize_bicubic_nocache");
built_options.emplace("-Dresize_bicubic_nocache=" + kernel_name);
auto dt = DataTypeToEnum<T>::value;
built_options.emplace("-DDATA_TYPE=" + DtToUpCompatibleCLDt(dt));
built_options.emplace("-DCMD_DATA_TYPE=" + DtToUpCompatibleCLCMDDt(dt));
MACE_RETURN_IF_ERROR(
runtime->BuildKernel("resize_bicubic",
kernel_name,
built_options,
&kernel_));
kwg_size_ =
static_cast<uint32_t>(runtime->GetKernelMaxWorkGroupSize(kernel_));
}
if (!IsVecEqual(input_shape_, input->shape())) {
MACE_CHECK(out_height > 0 && out_width > 0);
std::vector<index_t> output_shape{batch, out_height, out_width, channels};
std::vector<size_t> output_image_shape;
CalImage2DShape(output_shape, BufferType::IN_OUT_CHANNEL,
&output_image_shape);
MACE_RETURN_IF_ERROR(output->ResizeImage(output_shape, output_image_shape));
float height_scale =
CalculateResizeScale(in_height, out_height, align_corners_);
float width_scale =
CalculateResizeScale(in_width, out_width, align_corners_);
uint32_t idx = 0;
OUT_OF_RANGE_SET_ARG;
SET_3D_GWS_ARGS(kernel_);
kernel_.setArg(idx++, *(input->opencl_image()));
kernel_.setArg(idx++, *(output->opencl_image()));
kernel_.setArg(idx++, height_scale);
kernel_.setArg(idx++, width_scale);
kernel_.setArg(idx++, static_cast<int32_t>(in_height));
kernel_.setArg(idx++, static_cast<int32_t>(in_width));
kernel_.setArg(idx++, static_cast<int32_t>(out_height));
input_shape_ = input->shape();
}
const std::vector<uint32_t> lws = LocalWS(gws, kwg_size_);
std::string tuning_key =
Concat("resize_bicubic_opencl_kernel", output->dim(0), output->dim(1),
output->dim(2), output->dim(3));
MACE_RETURN_IF_ERROR(TuningOrRun3DKernel(kernel_, tuning_key,
gws, lws, future));
OUT_OF_RANGE_VALIDATION(kernel_error_);
return MACE_SUCCESS;
}
template struct ResizeBicubicFunctor<DeviceType::GPU, float>;
template struct ResizeBicubicFunctor<DeviceType::GPU, half>;
} // namespace kernels
} // namespace mace
mace/kernels/resize_bicubic.h 已删除 100644 → 0
浏览文件 @ e46c5a39
// Copyright 2018 Xiaomi, Inc. All rights reserved.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#ifndef MACE_KERNELS_RESIZE_BICUBIC_H_
#define MACE_KERNELS_RESIZE_BICUBIC_H_
#include <algorithm>
#include <memory>
#include <vector>
#include "mace/core/future.h"
#include "mace/core/tensor.h"
#include "mace/utils/logging.h"
#ifdef MACE_ENABLE_OPENCL
#include "mace/core/runtime/opencl/cl2_header.h"
#endif // MACE_ENABLE_OPENCL
namespace mace {
namespace kernels {
static const int64_t kTableSize = (1 << 10);
inline const float* InitCoeffsTable() {
// Allocate and initialize coefficients table using Bicubic
// convolution algorithm.
// https://en.wikipedia.org/wiki/Bicubic_interpolation
float* coeffs_tab = new float[(kTableSize + 1) * 2];
static const double A = -0.75;
for (int i = 0; i <= kTableSize; ++i) {
float x = i * 1.0 / kTableSize;
coeffs_tab[i * 2] = ((A + 2) * x - (A + 3)) * x * x + 1;
x += 1.0;
coeffs_tab[i * 2 + 1] = ((A * x - 5 * A) * x + 8 * A) * x - 4 * A;
}
return coeffs_tab;
}
inline const float* GetCoeffsTable() {
// Static so that we initialize it on first use
static const float* coeffs_tab = InitCoeffsTable();
return coeffs_tab;
}
inline int64_t Bound(int64_t val, int64_t limit) {
return std::min<int64_t>(limit - 1ll, std::max<int64_t>(0ll, val));
}
inline void GetWeightsAndIndices(float scale, int64_t out_loc, int64_t limit,
std::array<float, 4>* weights,
std::array<int64_t, 4>* indices) {
const int64_t in_loc = scale * out_loc;
const float delta = scale * out_loc - in_loc;
const int64_t offset = lrintf(delta * kTableSize);
const float* coeffs_tab = GetCoeffsTable();
*weights = {{coeffs_tab[offset * 2 + 1], coeffs_tab[offset * 2],
coeffs_tab[(kTableSize - offset) * 2],
coeffs_tab[(kTableSize - offset) * 2 + 1]}};
*indices = {{Bound(in_loc - 1, limit), Bound(in_loc, limit),
Bound(in_loc + 1, limit), Bound(in_loc + 2, limit)}};
}
inline float Interpolate1D(const std::array<float, 4>& weights,
const std::array<float, 4>& values) {
return values[0] * weights[0] + values[1] * weights[1] +
values[2] * weights[2] + values[3] * weights[3];
}
inline float CalculateResizeScale(index_t in_size,
index_t out_size,
bool align_corners) {
return (align_corners && out_size > 1)
? (in_size - 1) / static_cast<float>(out_size - 1)
: in_size / static_cast<float>(out_size);
}
inline void ResizeImage(const float *images,
const index_t batch_size,
const index_t in_height,
const index_t in_width,
const index_t out_height,
const index_t out_width,
const index_t channels,
const float height_scale,
const float width_scale,
float *output) {
std::array<float, 4> coeff = {{0.0, 0.0, 0.0, 0.0}};
#pragma omp parallel for collapse(2)
for (index_t b = 0; b < batch_size; ++b) {
for (index_t y = 0; y < out_height; ++y) {
std::array<float, 4> y_weights;
std::array<index_t, 4> y_indices;
GetWeightsAndIndices(height_scale, y, in_height, &y_weights,
&y_indices);
std::stringstream ss;
for (index_t x = 0; x < out_width; ++x) {
std::array<float, 4> x_weights;
std::array<index_t, 4> x_indices;
GetWeightsAndIndices(width_scale, x, in_width, &x_weights,
&x_indices);
for (index_t c = 0; c < channels; ++c) {
// Use a 4x4 patch to compute the interpolated output value at
// (b, y, x, c).
const float *channel_input_ptr =
images + (b * channels + c) * in_height * in_width;
float *channel_output_ptr =
output + (b * channels + c) * out_height * out_width;
for (index_t i = 0; i < 4; ++i) {
const std::array<float, 4> values = {
{static_cast<float>(channel_input_ptr
[y_indices[i] * in_width + x_indices[0]]),
static_cast<float>(channel_input_ptr
[y_indices[i] * in_width + x_indices[1]]),
static_cast<float>(channel_input_ptr
[y_indices[i] * in_width + x_indices[2]]),
static_cast<float>(channel_input_ptr
[y_indices[i] * in_width + x_indices[3]])}};
coeff[i] = Interpolate1D(x_weights, values);
}
channel_output_ptr[y * out_width + x] =
Interpolate1D(y_weights, coeff);
}
}
}
}
}
struct ResizeBicubicFunctorBase {
ResizeBicubicFunctorBase(const std::vector<index_t> &size,
bool align_corners)
: align_corners_(align_corners) {
MACE_CHECK(size.size() == 2);
out_height_ = size[0];
out_width_ = size[1];
}
protected:
bool align_corners_;
index_t out_height_;
index_t out_width_;
};
template<DeviceType D, typename T>
struct ResizeBicubicFunctor;
template<>
struct ResizeBicubicFunctor<DeviceType::CPU, float>
: ResizeBicubicFunctorBase {
ResizeBicubicFunctor(const std::vector<index_t> &size, bool align_corners)
: ResizeBicubicFunctorBase(size, align_corners) {}
MaceStatus operator()(const Tensor *input,
Tensor *output,
StatsFuture *future) {
MACE_UNUSED(future);
const index_t batch = input->dim(0);
const index_t channels = input->dim(1);
const index_t in_height = input->dim(2);
const index_t in_width = input->dim(3);
index_t out_height = out_height_;
index_t out_width = out_width_;
MACE_CHECK(out_height > 0 && out_width > 0);
std::vector<index_t> out_shape{batch, channels, out_height, out_width};
MACE_RETURN_IF_ERROR(output->Resize(out_shape));
Tensor::MappingGuard input_mapper(input);
Tensor::MappingGuard output_mapper(output);
const float *input_data = input->data<float>();
float *output_data = output->mutable_data<float>();
if (out_height == in_height && out_width == in_width) {
std::copy(input_data,
input_data + batch * channels * in_height * in_width,
output_data);
return MACE_SUCCESS;
}
float height_scale =
CalculateResizeScale(in_height, out_height, align_corners_);
float width_scale =
CalculateResizeScale(in_width, out_width, align_corners_);
ResizeImage(input_data, batch, in_height, in_width, out_height, out_width,
channels, height_scale, width_scale, output_data);
return MACE_SUCCESS;
}
};
#ifdef MACE_ENABLE_OPENCL
template<typename T>
struct ResizeBicubicFunctor<DeviceType::GPU, T>
: ResizeBicubicFunctorBase {
ResizeBicubicFunctor(const std::vector<index_t> &size, bool align_corners)
: ResizeBicubicFunctorBase(size, align_corners) {}
MaceStatus operator()(const Tensor *input,
Tensor *output,
StatsFuture *future);
cl::Kernel kernel_;
uint32_t kwg_size_;
std::unique_ptr<BufferBase> kernel_error_;
std::vector<index_t> input_shape_;
};
#endif // MACE_ENABLE_OPENCL
} // namespace kernels
} // namespace mace
#endif // MACE_KERNELS_RESIZE_BICUBIC_H_
mace/ops/ops_register.cc
浏览文件 @ 190385bb
...@@ -47,7 +47,6 @@ extern void Register_Proposal(OperatorRegistryBase *op_registry); ...@@ -47,7 +47,6 @@ extern void Register_Proposal(OperatorRegistryBase *op_registry);
extern void Register_Quantize(OperatorRegistryBase *op_registry); extern void Register_Quantize(OperatorRegistryBase *op_registry);
extern void Register_ReduceMean(OperatorRegistryBase *op_registry); extern void Register_ReduceMean(OperatorRegistryBase *op_registry);
extern void Register_Reshape(OperatorRegistryBase *op_registry); extern void Register_Reshape(OperatorRegistryBase *op_registry);
extern void Register_ResizeBicubic(OperatorRegistryBase *op_registry);
extern void Register_ResizeBilinear(OperatorRegistryBase *op_registry); extern void Register_ResizeBilinear(OperatorRegistryBase *op_registry);
extern void Register_ScalarMath(OperatorRegistryBase *op_registry); extern void Register_ScalarMath(OperatorRegistryBase *op_registry);
extern void Register_Shape(OperatorRegistryBase *op_registry); extern void Register_Shape(OperatorRegistryBase *op_registry);
...@@ -100,7 +99,6 @@ OperatorRegistry::OperatorRegistry() : OperatorRegistryBase() { ...@@ -100,7 +99,6 @@ OperatorRegistry::OperatorRegistry() : OperatorRegistryBase() {
ops::Register_Quantize(this); ops::Register_Quantize(this);
ops::Register_ReduceMean(this); ops::Register_ReduceMean(this);
ops::Register_Reshape(this); ops::Register_Reshape(this);
ops::Register_ResizeBicubic(this);
ops::Register_ResizeBilinear(this); ops::Register_ResizeBilinear(this);
ops::Register_ScalarMath(this); ops::Register_ScalarMath(this);
ops::Register_Shape(this); ops::Register_Shape(this);
......
mace/ops/resize_bicubic.cc 已删除 100644 → 0
浏览文件 @ e46c5a39
// Copyright 2018 Xiaomi, Inc. All rights reserved.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include "mace/ops/resize_bicubic.h"
namespace mace {
namespace ops {
void Register_ResizeBicubic(OperatorRegistryBase *op_registry) {
MACE_REGISTER_OPERATOR(op_registry, OpKeyBuilder("ResizeBicubic")
.Device(DeviceType::CPU)
.TypeConstraint<float>("T")
.Build(),
ResizeBicubicOp<DeviceType::CPU, float>);
#ifdef MACE_ENABLE_OPENCL
MACE_REGISTER_OPERATOR(op_registry, OpKeyBuilder("ResizeBicubic")
.Device(DeviceType::GPU)
.TypeConstraint<float>("T")
.Build(),
ResizeBicubicOp<DeviceType::GPU, float>);
MACE_REGISTER_OPERATOR(op_registry, OpKeyBuilder("ResizeBicubic")
.Device(DeviceType::GPU)
.TypeConstraint<half>("T")
.Build(),
ResizeBicubicOp<DeviceType::GPU, half>);
#endif // MACE_ENABLE_OPENCL
}
} // namespace ops
} // namespace mace
mace/ops/resize_bicubic.h 已删除 100644 → 0
浏览文件 @ e46c5a39
// Copyright 2018 Xiaomi, Inc. All rights reserved.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#ifndef MACE_OPS_RESIZE_BICUBIC_H_
#define MACE_OPS_RESIZE_BICUBIC_H_
#include "mace/core/operator.h"
#include "mace/kernels/resize_bicubic.h"
namespace mace {
namespace ops {
template <DeviceType D, class T>
class ResizeBicubicOp : public Operator<D, T> {
public:
ResizeBicubicOp(const OperatorDef &operator_def, Workspace *ws)
: Operator<D, T>(operator_def, ws),
functor_(OperatorBase::GetRepeatedArgs<index_t>("size", {-1, -1}),
OperatorBase::GetOptionalArg<bool>("align_corners", false)) {}
MaceStatus Run(StatsFuture *future) override {
const Tensor *input = this->Input(0);
Tensor *output = this->Output(0);
MACE_CHECK(input->dim_size() == 4, "input must be 4-dimensional.",
input->dim_size());
return functor_(input, output, future);
}
private:
kernels::ResizeBicubicFunctor<D, T> functor_;
};
} // namespace ops
} // namespace mace
#endif // MACE_OPS_RESIZE_BICUBIC_H_
mace/ops/resize_bicubic_benchmark.cc 已删除 100644 → 0
浏览文件 @ e46c5a39
// Copyright 2018 Xiaomi, Inc. All rights reserved.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include <string>
#include "mace/core/operator.h"
#include "mace/core/testing/test_benchmark.h"
#include "mace/ops/ops_test_util.h"
namespace mace {
namespace ops {
namespace test {
namespace {
template <DeviceType D, typename T>
void ResizeBicubicBenchmark(int iters,
int batch,
int channels,
int input_height,
int input_width,
int output_height,
int output_width) {
mace::testing::StopTiming();
OpsTestNet net;
// Add input data
if (D == DeviceType::CPU) {
net.AddRandomInput<D, float>("Input",
{batch, channels, input_height, input_width});
} else if (D == DeviceType::GPU) {
net.AddRandomInput<D, float>("Input",
{batch, input_height, input_width, channels});
} else {
MACE_NOT_IMPLEMENTED;
}
net.AddInputFromArray<D, int>("OutSize", {2},
{output_height, output_width});
if (D == DeviceType::CPU) {
OpDefBuilder("ResizeBicubic", "ResizeBicubicBenchmark")
.Input("Input")
.Input("OutSize")
.Output("Output")
.AddIntsArg("size", {output_height, output_width})
.AddIntArg("T", static_cast<int>(DataTypeToEnum<T>::value))
.Finalize(net.NewOperatorDef());
} else if (D == DeviceType::GPU) {
BufferToImage<D, T>(&net, "Input", "InputImage",
kernels::BufferType::IN_OUT_CHANNEL);
OpDefBuilder("ResizeBicubic", "ResizeBicubicBenchmark")
.Input("InputImage")
.Input("OutSize")
.Output("OutputImage")
.AddIntsArg("size", {output_height, output_width})
.AddIntArg("T", static_cast<int>(DataTypeToEnum<T>::value))
.Finalize(net.NewOperatorDef());
} else {
MACE_NOT_IMPLEMENTED;
}
// Warm-up
for (int i = 0; i < 5; ++i) {
net.RunOp(D);
}
mace::testing::StartTiming();
while (iters--) {
net.RunOp(D);
}
net.Sync();
}
} // namespace
#define MACE_BM_RESIZE_BICUBIC_MACRO(N, C, H0, W0, H1, W1, TYPE, DEVICE) \
static void \
MACE_BM_RESIZE_BICUBIC_##N##_##C##_##H0##_##W0##_##H1##_##W1##_##TYPE##_\
##DEVICE( \
int iters) { \
const int64_t macc = static_cast<int64_t>(iters) * N * C * H1 * W1 * 3; \
const int64_t tot = static_cast<int64_t>(iters) * N * C * H0 * W0; \
mace::testing::MaccProcessed(macc); \
mace::testing::BytesProcessed(tot *(sizeof(TYPE))); \
ResizeBicubicBenchmark<DEVICE, TYPE>(iters, N, C, H0, W0, H1, W1); \
} \
MACE_BENCHMARK( \
MACE_BM_RESIZE_BICUBIC_##N##_##C##_##H0##_##W0##_##H1##_##W1##_##TYPE##_\
##DEVICE)
#define MACE_BM_RESIZE_BICUBIC(N, C, H0, W0, H1, W1) \
MACE_BM_RESIZE_BICUBIC_MACRO(N, C, H0, W0, H1, W1, float, CPU); \
MACE_BM_RESIZE_BICUBIC_MACRO(N, C, H0, W0, H1, W1, float, GPU); \
MACE_BM_RESIZE_BICUBIC_MACRO(N, C, H0, W0, H1, W1, half, GPU);
MACE_BM_RESIZE_BICUBIC(1, 128, 120, 120, 480, 480);
MACE_BM_RESIZE_BICUBIC(1, 256, 7, 7, 15, 15);
MACE_BM_RESIZE_BICUBIC(1, 256, 15, 15, 30, 30);
MACE_BM_RESIZE_BICUBIC(1, 128, 30, 30, 60, 60);
MACE_BM_RESIZE_BICUBIC(1, 128, 240, 240, 480, 480);
MACE_BM_RESIZE_BICUBIC(1, 3, 4032, 3016, 480, 480);
MACE_BM_RESIZE_BICUBIC(1, 3, 480, 480, 4032, 3016);
} // namespace test
} // namespace ops
} // namespace mace
mace/ops/resize_bicubic_test.cc 已删除 100644 → 0
浏览文件 @ e46c5a39
// Copyright 2018 Xiaomi, Inc. All rights reserved.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include <vector>
#include "mace/core/operator.h"
#include "mace/ops/ops_test_util.h"
#include "mace/ops/resize_bicubic.h"
namespace mace {
namespace ops {
namespace test {
class ResizeBicubicTest : public OpsTestBase {};
TEST_F(ResizeBicubicTest, CPUResizeBicubicWOAlignCorners) {
testing::internal::LogToStderr();
// Construct graph
OpsTestNet net;
// Add input data
std::vector<float> input(24);
std::iota(begin(input), end(input), 0);
net.AddInputFromArray<DeviceType::CPU, float>("Input", {1, 2, 4, 3}, input);
net.TransformDataFormat<DeviceType::CPU, float>("Input", NHWC, "InputNCHW",
NCHW);
OpDefBuilder("ResizeBicubic", "ResizeBicubicTest")
.Input("InputNCHW")
.Output("OutputNCHW")
.AddIntsArg("size", {1, 2})
.Finalize(net.NewOperatorDef());
// Run
net.RunOp();
net.TransformDataFormat<DeviceType::CPU, float>("OutputNCHW", NCHW, "Output",
NHWC);
// Check
auto expected = CreateTensor<float>({1, 1, 2, 3}, {0, 1, 2, 6, 7, 8});
ExpectTensorNear<float>(*expected, *net.GetOutput("Output"), 1e-5);
}
TEST_F(ResizeBicubicTest, CPUResizeBicubicWOAlignCorners1) {
testing::internal::LogToStderr();
// Construct graph
OpsTestNet net;
// Add input data
std::vector<float> input(48);
std::iota(begin(input), end(input), 0);
net.AddInputFromArray<DeviceType::CPU, float>("Input", {1, 4, 4, 3}, input);
net.TransformDataFormat<DeviceType::CPU, float>("Input", NHWC, "InputNCHW",
NCHW);
OpDefBuilder("ResizeBicubic", "ResizeBicubicTest")
.Input("InputNCHW")
.Output("OutputNCHW")
.AddIntsArg("size", {2, 3})
.Finalize(net.NewOperatorDef());
// Run
net.RunOp();
net.TransformDataFormat<DeviceType::CPU, float>("OutputNCHW", NCHW, "Output",
NHWC);
// Check
auto expected = CreateTensor<float>({1, 2, 3, 3},
{0., 1., 2., 4.110297, 5.110297, 6.110297,
8.223037, 9.223036, 10.223037, 24., 25., 26.,
28.110298, 29.1103, 30.110298, 32.223038, 33.223038, 34.223038});
ExpectTensorNear<float>(*expected, *net.GetOutput("Output"), 1e-5);
}
TEST_F(ResizeBicubicTest, ResizeBicubicWAlignCorners) {
testing::internal::LogToStderr();
// Construct graph
OpsTestNet net;
// Add input data
std::vector<float> input(24);
std::iota(begin(input), end(input), 0);
net.AddInputFromArray<DeviceType::CPU, float>("Input", {1, 2, 4, 3}, input);
net.TransformDataFormat<DeviceType::CPU, float>("Input", NHWC, "InputNCHW",
NCHW);
OpDefBuilder("ResizeBicubic", "ResizeBicubicTest")
.Input("InputNCHW")
.Output("OutputNCHW")
.AddIntArg("align_corners", 1)
.AddIntsArg("size", {1, 2})
.Finalize(net.NewOperatorDef());
// Run
net.RunOp();
net.TransformDataFormat<DeviceType::CPU, float>("OutputNCHW", NCHW, "Output",
NHWC);
// Check
auto expected = CreateTensor<float>({1, 1, 2, 3}, {0, 1, 2, 9, 10, 11});
ExpectTensorNear<float>(*expected, *net.GetOutput("Output"), 1e-5);
}
namespace {
template <DeviceType D>
void TestRandomResizeBicubic() {
testing::internal::LogToStderr();
static unsigned int seed = time(NULL);
for (int round = 0; round < 10; ++round) {
int batch = 1 + rand_r(&seed) % 5;
int channels = 1 + rand_r(&seed) % 100;
int height = 1 + rand_r(&seed) % 100;
int width = 1 + rand_r(&seed) % 100;
int in_height = 1 + rand_r(&seed) % 100;
int in_width = 1 + rand_r(&seed) % 100;
int align_corners = rand_r(&seed) % 1;
// Construct graph
OpsTestNet net;
// Add input data
net.AddRandomInput<D, float>("Input",
{batch, in_height, in_width, channels});
net.TransformDataFormat<DeviceType::CPU, float>("Input", NHWC, "InputNCHW",
NCHW);
OpDefBuilder("ResizeBicubic", "ResizeBicubicTest")
.Input("InputNCHW")
.Output("OutputNCHW")
.AddIntArg("align_corners", align_corners)
.AddIntsArg("size", {height, width})
.Finalize(net.NewOperatorDef());
// Run on CPU
net.RunOp(DeviceType::CPU);
net.TransformDataFormat<DeviceType::CPU, float>("OutputNCHW", NCHW,
"Output", NHWC);
Tensor expected;
expected.Copy(*net.GetOutput("Output"));
if (D == DeviceType::GPU) {
BufferToImage<D, float>(&net, "Input", "InputImage",
kernels::BufferType::IN_OUT_CHANNEL);
OpDefBuilder("ResizeBicubic", "ResizeBicubicTest")
.Input("InputImage")
.Output("OutputImage")
.AddIntArg("align_corners", align_corners)
.AddIntsArg("size", {height, width})
.Finalize(net.NewOperatorDef());
// Run
net.RunOp(D);
ImageToBuffer<D, float>(&net, "OutputImage", "DeviceOutput",
kernels::BufferType::IN_OUT_CHANNEL);
}
// Check
ExpectTensorNear<float>(expected, *net.GetOutput("DeviceOutput"), 1e-5,
1e-4);
}
}
} // namespace
TEST_F(ResizeBicubicTest, OPENCLRandomResizeBicubic) {
TestRandomResizeBicubic<DeviceType::GPU>();
}
} // namespace test
} // namespace ops
} // namespace mace
mace/python/tools/converter_tool/base_converter.py
浏览文件 @ 190385bb
...@@ -100,7 +100,6 @@ MaceSupportedOps = [ ...@@ -100,7 +100,6 @@ MaceSupportedOps = [
'Quantize', 'Quantize',
'ReduceMean', 'ReduceMean',
'Reshape', 'Reshape',
'ResizeBicubic',
'ResizeBilinear', 'ResizeBilinear',
'ScalarMath', 'ScalarMath',
'Slice', 'Slice',
......
mace/python/tools/converter_tool/tensorflow_converter.py
浏览文件 @ 190385bb
...@@ -80,7 +80,6 @@ TFSupportedOps = [ ...@@ -80,7 +80,6 @@ TFSupportedOps = [
'Shape', 'Shape',
'Transpose', 'Transpose',
'Softmax', 'Softmax',
'ResizeBicubic',
'ResizeBilinear', 'ResizeBilinear',
'Placeholder', 'Placeholder',
'SpaceToBatchND', 'SpaceToBatchND',
...@@ -179,7 +178,6 @@ class TensorflowConverter(base_converter.ConverterInterface): ...@@ -179,7 +178,6 @@ class TensorflowConverter(base_converter.ConverterInterface):
TFOpType.Squeeze.name: self.convert_squeeze, TFOpType.Squeeze.name: self.convert_squeeze,
TFOpType.Transpose.name: self.convert_transpose, TFOpType.Transpose.name: self.convert_transpose,
TFOpType.Softmax.name: self.convert_softmax, TFOpType.Softmax.name: self.convert_softmax,
TFOpType.ResizeBicubic.name: self.convert_resize_bicubic,
TFOpType.ResizeBilinear.name: self.convert_resize_bilinear, TFOpType.ResizeBilinear.name: self.convert_resize_bilinear,
TFOpType.Placeholder.name: self.convert_nop, TFOpType.Placeholder.name: self.convert_nop,
TFOpType.SpaceToBatchND.name: self.convert_space_batch, TFOpType.SpaceToBatchND.name: self.convert_space_batch,
...@@ -536,20 +534,6 @@ class TensorflowConverter(base_converter.ConverterInterface): ...@@ -536,20 +534,6 @@ class TensorflowConverter(base_converter.ConverterInterface):
op = self.convert_general_op(tf_op) op = self.convert_general_op(tf_op)
op.type = MaceOp.Softmax.name op.type = MaceOp.Softmax.name
def convert_resize_bicubic(self, tf_op):
op = self.convert_general_op(tf_op)
op.type = MaceOp.ResizeBicubic.name
del op.input[1:]
size_arg = op.arg.add()
size_arg.name = MaceKeyword.mace_resize_size_str
size_value = tf_op.inputs[1].eval().astype(np.int32)
size_arg.ints.extend(size_value)
self._skip_tensor.add(tf_op.inputs[1].name)
align_corners_arg = op.arg.add()
align_corners_arg.name = MaceKeyword.mace_align_corners_str
align_corners_arg.i = tf_op.get_attr(tf_align_corners)
def convert_resize_bilinear(self, tf_op): def convert_resize_bilinear(self, tf_op):
op = self.convert_general_op(tf_op) op = self.convert_general_op(tf_op)
op.type = MaceOp.ResizeBilinear.name op.type = MaceOp.ResizeBilinear.name
......
repository/opencl-kernel/opencl_kernel_configure.bzl
浏览文件 @ 190385bb
...@@ -41,7 +41,6 @@ def _opencl_encrypt_kernel_impl(repository_ctx): ...@@ -41,7 +41,6 @@ def _opencl_encrypt_kernel_impl(repository_ctx):
unused_var = repository_ctx.path(Label("//:mace/kernels/opencl/cl/pad.cl")) unused_var = repository_ctx.path(Label("//:mace/kernels/opencl/cl/pad.cl"))
unused_var = repository_ctx.path(Label("//:mace/kernels/opencl/cl/pooling.cl")) unused_var = repository_ctx.path(Label("//:mace/kernels/opencl/cl/pooling.cl"))
unused_var = repository_ctx.path(Label("//:mace/kernels/opencl/cl/reduce_mean.cl")) unused_var = repository_ctx.path(Label("//:mace/kernels/opencl/cl/reduce_mean.cl"))
unused_var = repository_ctx.path(Label("//:mace/kernels/opencl/cl/resize_bicubic.cl"))
unused_var = repository_ctx.path(Label("//:mace/kernels/opencl/cl/resize_bilinear.cl")) unused_var = repository_ctx.path(Label("//:mace/kernels/opencl/cl/resize_bilinear.cl"))
unused_var = repository_ctx.path(Label("//:mace/kernels/opencl/cl/split.cl")) unused_var = repository_ctx.path(Label("//:mace/kernels/opencl/cl/split.cl"))
unused_var = repository_ctx.path(Label("//:mace/kernels/opencl/cl/softmax.cl")) unused_var = repository_ctx.path(Label("//:mace/kernels/opencl/cl/softmax.cl"))
......
Markdown is supported
0% 或 .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑！
想要评论请 注册 或