potok/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2026-07-30 15:53:03 +04:00
Code Packages Releases Wiki Activity

Merge branch 4.x

Browse Source
This commit is contained in:
Alexander Smorkalov
2023-07-27 18:38:25 +03:00
parent 0e8748746f 677a28fd2a
commit 47188b7c7e
151 changed files with 5199 additions and 1111 deletions
Download Patch File Download Diff File Expand all files Collapse all files
modules/dnn/include/opencv2/dnn/dnn.hpp
-1
View File
@@ -894,7 +894,6 @@ CV__DNN_INLINE_NS_BEGIN
* @param cfgFile path to the .cfg file with text description of the network architecture.
* @param darknetModel path to the .weights file with learned network.
* @returns Network object that ready to do forward, throw an exception in failure cases.
* @returns Net object.
*/
CV_EXPORTS_W Net readNetFromDarknet(const String &cfgFile, const String &darknetModel = String());
modules/dnn/include/opencv2/dnn/version.hpp
+1 -1
View File
@@ -6,7 +6,7 @@
#define OPENCV_DNN_VERSION_HPP
/// Use with major OpenCV version only.
#define OPENCV_DNN_API_VERSION 20221220
#define OPENCV_DNN_API_VERSION 20230620
#if !defined CV_DOXYGEN && !defined CV_STATIC_ANALYSIS && !defined CV_DNN_DONT_ADD_INLINE_NS
#define CV__DNN_INLINE_NS __CV_CAT(dnn5_v, OPENCV_DNN_API_VERSION)
modules/dnn/perf/perf_convolution.cpp
+2 -1
View File
@@ -744,7 +744,8 @@ static const ConvParam_t testConvolutionConfigs[] = {
/* GFLOPS 0.000 x 1 = 0.000 */ {{1, 1}, {{1, 4, 1, 1}}, 96, 1, {1, 1}, {1, 1}, {0, 0}, {0, 0}, "SAME", true, 864.},
/* GFLOPS 0.000 x 1 = 0.000 */ {{1, 1}, {{1, 96, 1, 1}}, 4, 1, {1, 1}, {1, 1}, {0, 0}, {0, 0}, "SAME", true, 772.},
/* GFLOPS 0.000 x 1 = 0.000 */ {{1, 1}, {{1, 8, 1, 1}}, 32, 1, {1, 1}, {1, 1}, {0, 0}, {0, 0}, "SAME", true, 544.},
/* GFLOPS 0.000 x 1 = 0.000 */ {{1, 1}, {{1, 32, 1, 1}}, 8, 1, {1, 1}, {1, 1}, {0, 0}, {0, 0}, "SAME", true, 520.}
/* GFLOPS 0.000 x 1 = 0.000 */ {{1, 1}, {{1, 32, 1, 1}}, 8, 1, {1, 1}, {1, 1}, {0, 0}, {0, 0}, "SAME", true, 520.},
/* GFLOPS 0.472 x 1 = 0.472 */ {{5, 5}, {{1, 32, 96, 96}}, 32, 32, {1, 1}, {1, 1}, {2, 2}, {0, 0}, "", true, 472154112.}
};
struct ConvParamID
{
modules/dnn/src/layers/cpu_kernels/conv_block.simd.hpp
+7 -7
View File
@@ -452,13 +452,13 @@ void convBlockMR1_F32(int np, const float * a, const float * b, float *c, const
if (init_c)
{
c0 += vld1q_f32(c);
c1 += vld1q_f32(c + 4);
c2 += vld1q_f32(c + 8);
c3 += vld1q_f32(c + 12);
c4 += vld1q_f32(c + 16);
c5 += vld1q_f32(c + 20);
c6 += vld1q_f32(c + 24);
c0 = vaddq_f32(c0, vld1q_f32(c));
c1 = vaddq_f32(c1, vld1q_f32(c + 4));
c2 = vaddq_f32(c2, vld1q_f32(c + 8));
c3 = vaddq_f32(c3, vld1q_f32(c + 12));
c4 = vaddq_f32(c4, vld1q_f32(c + 16));
c5 = vaddq_f32(c5, vld1q_f32(c + 20));
c6 = vaddq_f32(c6, vld1q_f32(c + 24));
}
if (ifMinMaxAct)
modules/dnn/src/layers/cpu_kernels/convolution.cpp
+1 -1
View File
@@ -1290,7 +1290,7 @@ void runFastConv(InputArray _input, OutputArray _output, const Ptr<FastConv>& co
else
Kg_nblocks = 1;
bool separateIm2col = fast_1x1 || stripes_per_plane == 1;
bool separateIm2col = (fast_1x1 || stripes_per_plane == 1) && conv->conv_type != CONV_TYPE_DEPTHWISE_REMAIN;
int Kstripes = Kg_nblocks * stripes_per_plane;
int nsubtasks = N * ngroups * Kstripes;
modules/dnn/src/layers/elementwise_layers.cpp
+107 -12
View File
@@ -138,7 +138,7 @@ public:
{
const float* srcptr = src_->ptr<float>(i) + stripeStart;
float* dstptr = dst_->ptr<float>(i) + stripeStart;
func_->apply(srcptr, dstptr, (int)(stripeEnd - stripeStart), planeSize, 0, outCn);
func_->apply(srcptr, dstptr, stripeStart, (int)(stripeEnd - stripeStart), planeSize, 0, outCn);
}
}
};
@@ -268,7 +268,7 @@ public:
void forwardSlice(const float* src, float* dst, int len, size_t planeSize, int cn0, int cn1) const CV_OVERRIDE
{
func.apply(src, dst, len, planeSize, cn0, cn1);
func.apply(src, dst, -1, len, planeSize, cn0, cn1);
}
#ifdef HAVE_CUDA
@@ -355,8 +355,9 @@ struct ReLUFunctor : public BaseFunctor
backendId == DNN_BACKEND_CANN;
}
void apply(const float* srcptr, float* dstptr, int len, size_t planeSize, int cn0, int cn1) const
void apply(const float* srcptr, float* dstptr, int stripeStart, int len, size_t planeSize, int cn0, int cn1) const
{
CV_UNUSED(stripeStart);
float s = slope;
for( int cn = cn0; cn < cn1; cn++, srcptr += planeSize, dstptr += planeSize )
{
@@ -559,8 +560,9 @@ struct ReLU6Functor : public BaseFunctor
backendId == DNN_BACKEND_CANN;
}
void apply(const float* srcptr, float* dstptr, int len, size_t planeSize, int cn0, int cn1) const
void apply(const float* srcptr, float* dstptr, int stripeStart, int len, size_t planeSize, int cn0, int cn1) const
{
CV_UNUSED(stripeStart);
for( int cn = cn0; cn < cn1; cn++, srcptr += planeSize, dstptr += planeSize )
{
int i = 0;
@@ -704,8 +706,9 @@ struct ReLU6Functor : public BaseFunctor
template <class T>
struct BaseDefaultFunctor : public BaseFunctor
{
void apply(const float* srcptr, float* dstptr, int len, size_t planeSize, int cn0, int cn1) const
void apply(const float* srcptr, float* dstptr, int stripeStart, int len, size_t planeSize, int cn0, int cn1) const
{
CV_UNUSED(stripeStart);
for( int cn = cn0; cn < cn1; cn++, srcptr += planeSize, dstptr += planeSize )
{
for( int i = 0; i < len; i++ )
@@ -1100,7 +1103,14 @@ struct SigmoidFunctor : public BaseDefaultFunctor<SigmoidFunctor>
inline float calculate(float x) const
{
return 1.f / (1.f + exp(-x));
float y;
if (x >= 0)
y = 1.f / (1.f + exp(-x));
else {
y = exp(x);
y = y / (1 + y);
}
return y;
}
#ifdef HAVE_CUDA
@@ -2219,8 +2229,9 @@ struct PowerFunctor : public BaseFunctor
shift = originShift;
}
void apply(const float* srcptr, float* dstptr, int len, size_t planeSize, int cn0, int cn1) const
void apply(const float* srcptr, float* dstptr, int stripeStart, int len, size_t planeSize, int cn0, int cn1) const
{
CV_UNUSED(stripeStart);
float a = scale, b = shift, p = power;
if( p == 1.f )
{
@@ -2445,6 +2456,7 @@ struct ChannelsPReLUFunctor : public BaseFunctor
Mat scale;
#ifdef HAVE_OPENCL
UMat scale_umat;
std::string oclKernelName = "ChannelsPReLUForward";
#endif
explicit ChannelsPReLUFunctor(const Mat& scale_=Mat()) : scale(scale_)
@@ -2463,8 +2475,9 @@ struct ChannelsPReLUFunctor : public BaseFunctor
backendId == DNN_BACKEND_CANN;
}
void apply(const float* srcptr, float* dstptr, int len, size_t planeSize, int cn0, int cn1) const
void apply(const float* srcptr, float* dstptr, int stripeStart, int len, size_t planeSize, int cn0, int cn1) const
{
CV_UNUSED(stripeStart);
CV_Assert(scale.isContinuous() && scale.type() == CV_32F);
const float* scaleptr = scale.ptr<float>();
@@ -2518,7 +2531,7 @@ struct ChannelsPReLUFunctor : public BaseFunctor
UMat& src = inputs[i];
UMat& dst = outputs[i];
ocl::Kernel kernel("PReLUForward", ocl::dnn::activations_oclsrc, buildopt);
ocl::Kernel kernel(oclKernelName.c_str(), ocl::dnn::activations_oclsrc, buildopt);
kernel.set(0, (int)src.total());
kernel.set(1, (int)src.size[1]);
kernel.set(2, (int)total(shape(src), 2));
@@ -2598,6 +2611,75 @@ struct ChannelsPReLUFunctor : public BaseFunctor
int64 getFLOPSPerElement() const { return 1; }
};
struct PReLUFunctor : public ChannelsPReLUFunctor
{
explicit PReLUFunctor(const Mat& scale_=Mat()) : ChannelsPReLUFunctor(scale_)
{
#ifdef HAVE_OPENCL
oclKernelName = "PReLUForward";
#endif
}
bool supportBackend(int backendId, int)
{
return backendId == DNN_BACKEND_OPENCV ||
backendId == DNN_BACKEND_CANN ||
backendId == DNN_BACKEND_INFERENCE_ENGINE_NGRAPH;
}
void apply(const float* srcptr, float* dstptr, int stripeStart, int len, size_t planeSize, int cn0, int cn1) const
{
CV_UNUSED(stripeStart);
CV_Assert(scale.isContinuous() && scale.type() == CV_32F);
if (stripeStart < 0)
CV_Error(Error::StsNotImplemented, "PReLUFunctor requires stripe offset parameter");
const float* scaleptr = scale.ptr<float>() + cn0 * planeSize + stripeStart;
for( int cn = cn0; cn < cn1; cn++, srcptr += planeSize, dstptr += planeSize, scaleptr += planeSize )
{
int i = 0;
#if CV_SIMD128
v_float32x4 z = v_setzero_f32();
for( ; i <= len - 16; i += 16 )
{
v_float32x4 x0 = v_load(srcptr + i);
v_float32x4 x1 = v_load(srcptr + i + 4);
v_float32x4 x2 = v_load(srcptr + i + 8);
v_float32x4 x3 = v_load(srcptr + i + 12);
v_float32x4 s0 = v_load(scaleptr + i);
v_float32x4 s1 = v_load(scaleptr + i + 4);
v_float32x4 s2 = v_load(scaleptr + i + 8);
v_float32x4 s3 = v_load(scaleptr + i + 12);
x0 = v_select(x0 >= z, x0, x0*s0);
x1 = v_select(x1 >= z, x1, x1*s1);
x2 = v_select(x2 >= z, x2, x2*s2);
x3 = v_select(x3 >= z, x3, x3*s3);
v_store(dstptr + i, x0);
v_store(dstptr + i + 4, x1);
v_store(dstptr + i + 8, x2);
v_store(dstptr + i + 12, x3);
}
#endif
for( ; i < len; i++ )
{
float x = srcptr[i];
float s = scaleptr[i];
dstptr[i] = x >= 0.f ? x : s*x;
}
}
}
#ifdef HAVE_DNN_NGRAPH
std::shared_ptr<ngraph::Node> initNgraphAPI(const std::shared_ptr<ngraph::Node>& node)
{
auto shape = getShape<size_t>(scale);
auto slope = std::make_shared<ngraph::op::Constant>(ngraph::element::f32, shape, scale.ptr<float>());
return std::make_shared<ngraph::op::PRelu>(node, slope);
}
#endif // HAVE_DNN_NGRAPH
};
struct SignFunctor : public BaseDefaultFunctor<SignFunctor>
{
typedef SignLayer Layer;
@@ -3033,13 +3115,26 @@ Ptr<ExpLayer> ExpLayer::create(const LayerParams& params)
Ptr<Layer> ChannelsPReLULayer::create(const LayerParams& params)
{
CV_Assert(params.blobs.size() == 1);
if (params.blobs[0].total() == 1)
Mat scale = params.blobs[0];
float slope = *scale.ptr<float>();
if (scale.total() == 1 || countNonZero(scale != slope) == 0)
{
LayerParams reluParams = params;
reluParams.set("negative_slope", *params.blobs[0].ptr<float>());
reluParams.set("negative_slope", slope);
return ReLULayer::create(reluParams);
}
Ptr<ChannelsPReLULayer> l(new ElementWiseLayer<ChannelsPReLUFunctor>(ChannelsPReLUFunctor(params.blobs[0])));
Ptr<Layer> l;
// Check first two dimensions of scale (batch, channels)
MatShape scaleShape = shape(scale);
if (std::count_if(scaleShape.begin(), scaleShape.end(), [](int d){ return d != 1;}) > 1)
{
l = new ElementWiseLayer<PReLUFunctor>(PReLUFunctor(scale));
}
else
{
l = new ElementWiseLayer<ChannelsPReLUFunctor>(ChannelsPReLUFunctor(scale));
}
l->setParamsFrom(params);
return l;
modules/dnn/src/layers/max_unpooling_layer.cpp
+7 -11
View File
@@ -14,6 +14,7 @@ Implementation of Batch Normalization layer.
#include "../op_cuda.hpp"
#include "../op_halide.hpp"
#include <opencv2/dnn/shape_utils.hpp>
#include <opencv2/core/utils/logger.hpp>
#ifdef HAVE_CUDA
#include "../cuda4dnn/primitives/max_unpooling.hpp"
@@ -110,17 +111,12 @@ public:
int index = idxptr[i_wh];
if (!(0 <= index && index < outPlaneTotal))
{
std::cerr
<< "i_n=" << i_n << std::endl
<< "i_c=" << i_c << std::endl
<< "i_wh=" << i_wh << std::endl
<< "index=" << index << std::endl
<< "maxval=" << inptr[i_wh] << std::endl
<< "outPlaneTotal=" << outPlaneTotal << std::endl
<< "input.size=" << input.size << std::endl
<< "indices.size=" << indices.size << std::endl
<< "outBlob=" << outBlob.size << std::endl
;
CV_LOG_ERROR(NULL, cv::format(
"i_n=%d\ni_c=%d\ni_wh=%d\nindex=%d\nmaxval=%lf\noutPlaneTotal=%d\n",
i_n, i_c, i_wh, index, inptr[i_wh], outPlaneTotal));
CV_LOG_ERROR(NULL, "input.size=" << input.size);
CV_LOG_ERROR(NULL, "indices.size=" << indices.size);
CV_LOG_ERROR(NULL, "outBlob=" << outBlob.size);
CV_Assert(0 <= index && index < outPlaneTotal);
}
outptr[index] = inptr[i_wh];
modules/dnn/src/net_cann.cpp
+2 -2
View File
@@ -304,9 +304,9 @@ std::shared_ptr<ge::ModelBufferData> compileCannGraph(std::shared_ptr<ge::Graph>
bool ok;
if ((child=fork()) == 0)
{
// initialize engine
// initialize engine Ascend310/Ascend310P3/Ascend910B/Ascend310B
std::map<ge::AscendString, ge::AscendString> options = {
{ge::AscendString(ge::ir_option::SOC_VERSION), ge::AscendString("Ascend310")},
{ge::AscendString(ge::ir_option::SOC_VERSION), ge::AscendString(aclrtGetSocName())},
};
ACL_CHECK_GRAPH_RET(ge::aclgrphBuildInitialize(options));
modules/dnn/src/net_impl.cpp
+8 -8
View File
@@ -662,14 +662,14 @@ void Net::Impl::forwardLayer(LayerData& ld)
m = u.getMat(ACCESS_READ);
if (!checkRange(m))
{
std::cerr << "WARNING: NaN detected in layer output: id=" << ld.id << " name=" << layer->name << std::endl;
std::cerr << "output id=" << i << " output shape=" << shape(m) << std::endl;
CV_LOG_WARNING(NULL, "NaN detected in layer output: id=" << ld.id << " name=" << layer->name
<< " output id=" << i << " output shape=" << shape(m));
fail = true;
}
else if (!checkRange(m, true, NULL, -1e6, 1e6))
{
std::cerr << "WARNING: Inf detected in layer output: id=" << ld.id << " name=" << layer->name << std::endl;
std::cerr << "output id=" << i << " output shape=" << shape(m) << std::endl;
CV_LOG_WARNING(NULL, "Inf detected in layer output: id=" << ld.id << " name=" << layer->name
<< " output id=" << i << " output shape=" << shape(m));
fail = true;
}
}
@@ -738,14 +738,14 @@ void Net::Impl::forwardLayer(LayerData& ld)
const Mat& m = ld.outputBlobs[i];
if (!checkRange(m))
{
std::cerr << "WARNING: NaN detected in layer output: id=" << ld.id << " name=" << layer->name << std::endl;
std::cerr << "output id=" << i << " output shape=" << shape(m) << std::endl;
CV_LOG_WARNING(NULL, "NaN detected in layer output: "
<< cv::format("id=%d name=%s output id=%zu output shape=", ld.id, layer->name.c_str(), i) << shape(m));
fail = true;
}
else if (!checkRange(m, true, NULL, -1e6, 1e6))
{
std::cerr << "WARNING: Inf detected in layer output: id=" << ld.id << " name=" << layer->name << std::endl;
std::cerr << "output id=" << i << " output shape=" << shape(m) << std::endl;
CV_LOG_WARNING(NULL, "Inf detected in layer output: "
<< cv::format("id=%d name=%s output id=%zu output shape=", ld.id, layer->name.c_str(), i) << shape(m));
fail = true;
}
}
modules/dnn/src/onnx/onnx_graph_simplifier.cpp
+1 -1
View File
@@ -1125,7 +1125,7 @@ Mat getMatFromTensor(const opencv_onnx::TensorProto& tensor_proto)
else if (datatype == opencv_onnx::TensorProto_DataType_FLOAT16)
{
// FIXME, for now, we only load FP16 Tensor as FP32 Mat, full support for FP16 is required in the future.
CV_LOG_ONCE_WARNING(NULL, "DNN: load FP16 model as FP32 model, and it takes twice the FP16 RAM requirement.");
CV_LOG_ONCE_INFO(NULL, "DNN: load FP16 model as FP32 model, and it takes twice the FP16 RAM requirement.");
// ONNX saves float 16 data in two format: int32 and raw_data.
// Link: https://github.com/onnx/onnx/issues/4460#issuecomment-1224373746
modules/dnn/src/op_cann.hpp
+5 -1
View File
@@ -10,7 +10,11 @@
#include "graph/graph.h" // ge::Graph; ge::Operator from operator.h
#include "graph/ge_error_codes.h" // GRAPH_SUCCESS, ...
#include "op_proto/built-in/inc/all_ops.h" // ge::Conv2D, ...
#ifdef CANN_VERSION_BELOW_6_3_ALPHA002
#include "op_proto/built-in/inc/all_ops.h" // ge::Conv2D, ...
#else
#include "built-in/op_proto/inc/all_ops.h" // ge::Conv2D, ...
#endif
#include "graph/tensor.h" // ge::Shape, ge::Tensor, ge::TensorDesc
#include "graph/types.h" // DT_FLOAT, ... ; FORMAT_NCHW, ...
modules/dnn/src/opencl/activations.cl
+12 -3
View File
@@ -73,14 +73,23 @@ __kernel void ReLU6Forward(const int count, __global const T* in, __global T* ou
}
}
__kernel void ChannelsPReLUForward(const int count, const int channels, const int plane_size,
__global const T* in, __global T* out,
__global const KERNEL_ARG_DTYPE* slope_data)
{
int index = get_global_id(0);
int c = (index / plane_size) % channels;
if(index < count)
out[index] = in[index] > 0 ? in[index] : in[index] * slope_data[c];
}
__kernel void PReLUForward(const int count, const int channels, const int plane_size,
__global const T* in, __global T* out,
__global const KERNEL_ARG_DTYPE* slope_data)
{
int index = get_global_id(0);
int c = (index / plane_size) % channels;
if(index < count)
out[index] = in[index] > 0 ? in[index] : in[index] * slope_data[c];
out[index] = in[index] > 0 ? in[index] : in[index] * slope_data[index];
}
__kernel void TanHForward(const int count, __global T* in, __global T* out) {
@@ -352,4 +361,4 @@ __kernel void ReciprocalForward(const int n, __global T* in, __global T* out)
int index = get_global_id(0);
if(index < n)
out[index] = 1.0f/in[index];
}
}
modules/dnn/src/tensorflow/tf_importer.cpp
+23
View File
@@ -589,6 +589,7 @@ private:
void parsePack (tensorflow::GraphDef& net, const tensorflow::NodeDef& layer, LayerParams& layerParams);
void parseClipByValue (tensorflow::GraphDef& net, const tensorflow::NodeDef& layer, LayerParams& layerParams);
void parseLeakyRelu (tensorflow::GraphDef& net, const tensorflow::NodeDef& layer, LayerParams& layerParams);
void parsePReLU (tensorflow::GraphDef& net, const tensorflow::NodeDef& layer, LayerParams& layerParams);
void parseActivation (tensorflow::GraphDef& net, const tensorflow::NodeDef& layer, LayerParams& layerParams);
void parseExpandDims (tensorflow::GraphDef& net, const tensorflow::NodeDef& layer, LayerParams& layerParams);
void parseSquare (tensorflow::GraphDef& net, const tensorflow::NodeDef& layer, LayerParams& layerParams);
@@ -668,6 +669,7 @@ TFImporter::DispatchMap TFImporter::buildDispatchMap()
dispatch["Pack"] = &TFImporter::parsePack;
dispatch["ClipByValue"] = &TFImporter::parseClipByValue;
dispatch["LeakyRelu"] = &TFImporter::parseLeakyRelu;
dispatch["PReLU"] = &TFImporter::parsePReLU;
dispatch["Abs"] = dispatch["Tanh"] = dispatch["Sigmoid"] = dispatch["Relu"] =
dispatch["Elu"] = dispatch["Exp"] = dispatch["Identity"] = dispatch["Relu6"] = &TFImporter::parseActivation;
dispatch["ExpandDims"] = &TFImporter::parseExpandDims;
@@ -2622,6 +2624,27 @@ void TFImporter::parseLeakyRelu(tensorflow::GraphDef& net, const tensorflow::Nod
connectToAllBlobs(layer_id, dstNet, parsePin(layer.input(0)), id, num_inputs);
}
void TFImporter::parsePReLU(tensorflow::GraphDef& net, const tensorflow::NodeDef& layer, LayerParams& layerParams)
{
const std::string& name = layer.name();
Mat scales;
blobFromTensor(getConstBlob(layer, value_id, 1), scales);
layerParams.blobs.resize(1);
if (scales.dims == 3) {
// Considering scales from Keras wih HWC layout;
transposeND(scales, {2, 0, 1}, layerParams.blobs[0]);
} else {
layerParams.blobs[0] = scales;
}
int id = dstNet.addLayer(name, "PReLU", layerParams);
layer_id[name] = id;
connect(layer_id, dstNet, parsePin(layer.input(0)), id, 0);
}
// "Abs" "Tanh" "Sigmoid" "Relu" "Elu" "Exp" "Identity" "Relu6"
void TFImporter::parseActivation(tensorflow::GraphDef& net, const tensorflow::NodeDef& layer, LayerParams& layerParams)
{
modules/dnn/src/tflite/tflite_importer.cpp
+93 -12
View File
@@ -59,6 +59,7 @@ private:
void parseUnpooling(const Operator& op, const std::string& opcode, LayerParams& layerParams);
void parseReshape(const Operator& op, const std::string& opcode, LayerParams& layerParams);
void parseConcat(const Operator& op, const std::string& opcode, LayerParams& layerParams);
void parsePack(const Operator& op, const std::string& opcode, LayerParams& layerParams);
void parseResize(const Operator& op, const std::string& opcode, LayerParams& layerParams);
void parseDeconvolution(const Operator& op, const std::string& opcode, LayerParams& layerParams);
void parseQuantize(const Operator& op, const std::string& opcode, LayerParams& layerParams);
@@ -70,6 +71,8 @@ private:
void parseActivation(const Operator& op, const std::string& opcode, LayerParams& layerParams, bool isFused);
void addLayer(LayerParams& layerParams, const Operator& op);
int addPermuteLayer(const std::vector<int>& order, const std::string& permName, const std::pair<int, int>& inpId, int dtype);
int addReshapeLayer(const std::vector<int>& shape, int axis, int num_axes,
const std::string& name, const std::pair<int, int>& inpId, int dtype);
inline bool isInt8(const Operator& op);
inline void getQuantParams(const Operator& op, float& inpScale, int& inpZero, float& outScale, int& outZero);
};
@@ -267,6 +270,7 @@ TFLiteImporter::DispatchMap TFLiteImporter::buildDispatchMap()
dispatch["PAD"] = &TFLiteImporter::parsePadding;
dispatch["RESHAPE"] = &TFLiteImporter::parseReshape;
dispatch["CONCATENATION"] = &TFLiteImporter::parseConcat;
dispatch["PACK"] = &TFLiteImporter::parsePack;
dispatch["RESIZE_BILINEAR"] = dispatch["RESIZE_NEAREST_NEIGHBOR"] = &TFLiteImporter::parseResize;
dispatch["Convolution2DTransposeBias"] = &TFLiteImporter::parseDeconvolution;
dispatch["QUANTIZE"] = &TFLiteImporter::parseQuantize;
@@ -596,16 +600,6 @@ void TFLiteImporter::parseUnpooling(const Operator& op, const std::string& opcod
void TFLiteImporter::parseReshape(const Operator& op, const std::string& opcode, LayerParams& layerParams) {
DataLayout inpLayout = layouts[op.inputs()->Get(0)];
if (inpLayout == DNN_LAYOUT_NHWC) {
// Permute to NCHW
std::vector<int> order = {0, 2, 3, 1};
const std::string name = layerParams.name + "/permute";
auto inpId = layerIds[op.inputs()->Get(0)];
int permId = addPermuteLayer(order, name, inpId, isInt8(op) ? CV_8S : CV_32F); // NCHW -> NHWC
layerIds[op.inputs()->Get(0)] = std::make_pair(permId, 0);
layouts[op.outputs()->Get(0)] = DNN_LAYOUT_NCHW;
}
layerParams.type = "Reshape";
std::vector<int> shape;
if (op.inputs()->size() > 1) {
@@ -615,6 +609,22 @@ void TFLiteImporter::parseReshape(const Operator& op, const std::string& opcode,
CV_Assert(options);
shape.assign(options->new_shape()->begin(), options->new_shape()->end());
}
if (inpLayout == DNN_LAYOUT_NHWC) {
if (shape.size() == 4) {
// Keep data but change a shape to OpenCV's NCHW order
std::swap(shape[2], shape[3]);
std::swap(shape[1], shape[2]);
} else {
// Permute to NCHW entire data and reshape to given a shape
std::vector<int> order = {0, 2, 3, 1};
const std::string name = layerParams.name + "/permute";
auto inpId = layerIds[op.inputs()->Get(0)];
int permId = addPermuteLayer(order, name, inpId, isInt8(op) ? CV_8S : CV_32F); // NCHW -> NHWC
layerIds[op.inputs()->Get(0)] = std::make_pair(permId, 0);
layouts[op.outputs()->Get(0)] = DNN_LAYOUT_NCHW;
}
}
layerParams.set("dim", DictValue::arrayInt<int*>(shape.data(), shape.size()));
addLayer(layerParams, op);
}
@@ -636,6 +646,47 @@ void TFLiteImporter::parseConcat(const Operator& op, const std::string& opcode,
parseFusedActivation(op, options->fused_activation_function());
}
void TFLiteImporter::parsePack(const Operator& op, const std::string& opcode, LayerParams& layerParams) {
auto options = reinterpret_cast<const PackOptions*>(op.builtin_options());
int axis = options->axis();
DataLayout inpLayout = layouts[op.inputs()->Get(0)];
if (inpLayout == DNN_LAYOUT_NHWC) {
// OpenCV works in NCHW data layout. So change the axis correspondingly.
axis = normalize_axis(axis, 5); // 5 because Pack adds a new axis so -1 would mean 4
static const int remap[] = {0, 1, 3, 4, 2};
axis = remap[axis];
}
// Replace Pack layer to Reshape + Concat
// Use a set because there are models which replicate single layer data by Pack.
std::set<int> op_inputs(op.inputs()->begin(), op.inputs()->end());
std::map<int, std::pair<int, int> > originLayerIds;
for (int inp : op_inputs) {
auto inpId = layerIds[inp];
int dims = modelTensors->Get(inp)->shape()->size();
std::vector<int> shape{1, -1};
if (axis == dims) {
std::swap(shape[0], shape[1]);
}
const auto name = modelTensors->Get(inp)->name()->str() + "/reshape";
int reshapeId = addReshapeLayer(shape, axis == dims ? dims - 1 : axis, 1,
name, inpId, isInt8(op) ? CV_8S : CV_32F);
originLayerIds[inp] = layerIds[inp];
layerIds[inp] = std::make_pair(reshapeId, 0);
}
layerParams.type = "Concat";
layerParams.set("axis", axis);
addLayer(layerParams, op);
// Restore origin layer inputs
for (const auto& ids : originLayerIds) {
layerIds[ids.first] = ids.second;
}
}
void TFLiteImporter::parseResize(const Operator& op, const std::string& opcode, LayerParams& layerParams) {
layerParams.type = "Resize";
@@ -666,6 +717,18 @@ int TFLiteImporter::addPermuteLayer(const std::vector<int>& order, const std::st
return permId;
}
int TFLiteImporter::addReshapeLayer(const std::vector<int>& shape, int axis, int num_axes,
const std::string& name, const std::pair<int, int>& inpId, int dtype)
{
LayerParams lp;
lp.set("axis", axis);
lp.set("dim", DictValue::arrayInt<const int*>(shape.data(), shape.size()));
lp.set("num_axes", num_axes);
int id = dstNet.addLayer(name, "Reshape", dtype, lp);
dstNet.connect(inpId.first, inpId.second, id, 0);
return id;
}
void TFLiteImporter::parseDeconvolution(const Operator& op, const std::string& opcode, LayerParams& layerParams) {
layerParams.type = "Deconvolution";
@@ -771,6 +834,8 @@ void TFLiteImporter::parseDetectionPostProcess(const Operator& op, const std::st
parameters[keys[i]] = *reinterpret_cast<const uint32_t*>(data + offset + i * 4);
}
parameters["num_classes"] = modelTensors->Get(op.inputs()->Get(1))->shape()->Get(2);
layerParams.type = "DetectionOutput";
layerParams.set("num_classes", parameters["num_classes"]);
layerParams.set("share_location", true);
@@ -780,7 +845,6 @@ void TFLiteImporter::parseDetectionPostProcess(const Operator& op, const std::st
layerParams.set("top_k", parameters["max_detections"]);
layerParams.set("keep_top_k", parameters["max_detections"]);
layerParams.set("code_type", "CENTER_SIZE");
layerParams.set("variance_encoded_in_target", true);
layerParams.set("loc_pred_transposed", true);
// Replace third input from tensor to Const layer with the priors
@@ -796,10 +860,27 @@ void TFLiteImporter::parseDetectionPostProcess(const Operator& op, const std::st
priors.col(2) = priors.col(0) + priors.col(3);
priors.col(3) = priors.col(1) + tmp;
float x_scale = *(float*)&parameters["x_scale"];
float y_scale = *(float*)&parameters["y_scale"];
float w_scale = *(float*)&parameters["w_scale"];
float h_scale = *(float*)&parameters["h_scale"];
if (x_scale != 1.0f || y_scale != 1.0f || w_scale != 1.0f || h_scale != 1.0f) {
int numPriors = priors.rows;
priors.resize(numPriors * 2);
Mat_<float> scales({1, 4}, {1.f / x_scale, 1.f / y_scale,
1.f / w_scale, 1.f / h_scale});
repeat(scales, numPriors, 1, priors.rowRange(numPriors, priors.rows));
priors = priors.reshape(1, {1, 2, (int)priors.total() / 2});
layerParams.set("variance_encoded_in_target", false);
} else {
priors = priors.reshape(1, {1, 1, (int)priors.total()});
layerParams.set("variance_encoded_in_target", true);
}
LayerParams priorsLP;
priorsLP.name = layerParams.name + "/priors";
priorsLP.type = "Const";
priorsLP.blobs.resize(1, priors.reshape(1, {1, 1, (int)priors.total()}));
priorsLP.blobs.resize(1, priors);
int priorsId = dstNet.addLayer(priorsLP.name, priorsLP.type, priorsLP);
layerIds[op.inputs()->Get(2)] = std::make_pair(priorsId, 0);
modules/dnn/test/test_tf_importer.cpp
+7 -1
View File
@@ -1675,6 +1675,7 @@ TEST_P(Test_TensorFlow_layers, clip_by_value)
TEST_P(Test_TensorFlow_layers, tf2_prelu)
{
double l1 = 0, lInf = 0;
if (backend == DNN_BACKEND_CUDA)
applyTestTag(CV_TEST_TAG_DNN_SKIP_CUDA); // not supported; only across channels is supported
#if defined(INF_ENGINE_RELEASE) && INF_ENGINE_VER_MAJOR_EQ(2022010000)
@@ -1686,6 +1687,11 @@ TEST_P(Test_TensorFlow_layers, tf2_prelu)
applyTestTag(target == DNN_TARGET_OPENCL ? CV_TEST_TAG_DNN_SKIP_IE_OPENCL : CV_TEST_TAG_DNN_SKIP_IE_OPENCL_FP16,
CV_TEST_TAG_DNN_SKIP_IE_NGRAPH, CV_TEST_TAG_DNN_SKIP_IE_VERSION
);
#elif defined(INF_ENGINE_RELEASE) && INF_ENGINE_VER_MAJOR_GE(2023000000)
if (backend == DNN_BACKEND_INFERENCE_ENGINE_NGRAPH && target == DNN_TARGET_OPENCL) {
l1 = 1e-4;
lInf = 1e-3;
}
#elif defined(INF_ENGINE_RELEASE) && INF_ENGINE_VER_MAJOR_EQ(2021040000)
if (backend == DNN_BACKEND_INFERENCE_ENGINE_NGRAPH)
{
@@ -1705,7 +1711,7 @@ TEST_P(Test_TensorFlow_layers, tf2_prelu)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_NGRAPH);
#endif
runTensorFlowNet("tf2_prelu");
runTensorFlowNet("tf2_prelu", false, l1, lInf);
}
TEST_P(Test_TensorFlow_layers, tf2_permute_nhwc_ncwh)
modules/dnn/test/test_tflite_importer.cpp
+18 -2
View File
@@ -31,9 +31,8 @@ void testInputShapes(const Net& net, const std::vector<Mat>& inps) {
}
}
void testModel(const std::string& modelName, const Mat& input, double l1 = 1e-5, double lInf = 1e-4)
void testModel(Net& net, const std::string& modelName, const Mat& input, double l1 = 1e-5, double lInf = 1e-4)
{
Net net = readNet(findDataFile("dnn/tflite/" + modelName + ".tflite", false));
testInputShapes(net, {input});
net.setInput(input);
@@ -49,6 +48,12 @@ void testModel(const std::string& modelName, const Mat& input, double l1 = 1e-5,
}
}
void testModel(const std::string& modelName, const Mat& input, double l1 = 1e-5, double lInf = 1e-4)
{
Net net = readNet(findDataFile("dnn/tflite/" + modelName + ".tflite", false));
testModel(net, modelName, input, l1, lInf);
}
void testModel(const std::string& modelName, const Size& inpSize, double l1 = 1e-5, double lInf = 1e-4)
{
Mat input = imread(findDataFile("cv/shared/lena.png"));
@@ -56,6 +61,13 @@ void testModel(const std::string& modelName, const Size& inpSize, double l1 = 1e
testModel(modelName, input, l1, lInf);
}
void testLayer(const std::string& modelName, double l1 = 1e-5, double lInf = 1e-4)
{
Mat inp = blobFromNPY(findDataFile("dnn/tflite/" + modelName + "_inp.npy"));
Net net = readNet(findDataFile("dnn/tflite/" + modelName + ".tflite"));
testModel(net, modelName, inp, l1, lInf);
}
// https://google.github.io/mediapipe/solutions/face_mesh
TEST(Test_TFLite, face_landmark)
{
@@ -146,6 +158,10 @@ TEST(Test_TFLite, EfficientDet_int8) {
normAssertDetections(ref, out, "", 0.5, 0.05, 0.1);
}
TEST(Test_TFLite, replicate_by_pack) {
testLayer("replicate_by_pack");
}
}} // namespace
#endif // OPENCV_TEST_DNN_TFLITE
modules/dnn/test/test_torch_importer.cpp
+2 -2
View File
@@ -566,14 +566,14 @@ TEST_P(Test_Torch_nets, FastNeuralStyle_accuracy)
}
else if(target == DNN_TARGET_CUDA_FP16)
{
normAssert(out, refBlob, "", 0.6, 25);
normAssert(out, refBlob, "", 0.6, 26);
}
else if (target == DNN_TARGET_CPU_FP16)
{
normAssert(out, refBlob, "", 0.62, 25);
}
else
normAssert(out, refBlob, "", 0.5, 1.1);
normAssert(out, refBlob, "", 0.5, 1.11);
}
}