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

Merge remote-tracking branch 'upstream/3.4' into merge-3.4

Browse Source
This commit is contained in:
Alexander Alekhin
2020-08-06 14:15:52 +00:00
parent dbab8d8c38 1067cd0649
commit fa25faa2d2
52 changed files with 1498 additions and 319 deletions
Download Patch File Download Diff File Expand all files Collapse all files
modules/dnn/include/opencv2/dnn/all_layers.hpp
+8
View File
@@ -600,6 +600,14 @@ CV__DNN_INLINE_NS_BEGIN
static Ptr<RegionLayer> create(const LayerParams& params);
};
/**
* @brief Detection output layer.
*
* The layer size is: @f$ (1 \times 1 \times N \times 7) @f$
* where N is [keep_top_k] parameter multiplied by batch size. Each row is:
* [image_id, label, confidence, xmin, ymin, xmax, ymax]
* where image_id is the index of image input in the batch.
*/
class CV_EXPORTS DetectionOutputLayer : public Layer
{
public:
modules/dnn/src/darknet/darknet_io.cpp
+6 -18
View File
@@ -221,6 +221,10 @@ namespace cv {
{
cv::dnn::LayerParams activation_param;
if (type == "relu")
{
activation_param.type = "ReLU";
}
else if (type == "leaky")
{
activation_param.set<float>("negative_slope", 0.1f);
activation_param.type = "ReLU";
@@ -862,24 +866,8 @@ namespace cv {
}
std::string activation = getParam<std::string>(layer_params, "activation", "linear");
if (activation == "leaky")
{
setParams.setActivation("relu");
}
else if (activation == "swish")
{
setParams.setActivation("swish");
}
else if (activation == "mish")
{
setParams.setActivation("mish");
}
else if (activation == "logistic")
{
setParams.setActivation("logistic");
}
else if (activation != "linear")
CV_Error(cv::Error::StsParseError, "Unsupported activation: " + activation);
if (activation != "linear")
setParams.setActivation(activation);
net->out_channels_vec[layers_counter] = tensor_shape[0];
}
modules/dnn/src/layers/convolution_layer.cpp
+116 -50
View File
@@ -114,18 +114,19 @@ public:
inputs_arr.getMatVector(inputs);
outputs_arr.getMatVector(outputs);
CV_Assert(inputs.size() > 0);
CV_Assert((inputs.size() > outputs.size() && blobs.empty()) ||
(!inputs.empty() && (blobs.size() == 1 || blobs.size() == 2)));
MatSize weightShape = blobs.empty() ? inputs[1].size : blobs[0].size;
CV_Assert(blobs.size() == 1 || blobs.size() == 2);
CV_Assert(inputs[0].dims == outputs[0].dims);
CV_Assert(blobs[0].dims == kernel_size.size() + 2);
CV_Assert(weightShape.dims() == kernel_size.size() + 2);
for (int i = 0; i < kernel_size.size(); i++) {
CV_Assert(blobs[0].size[i + 2] == kernel_size[i]);
CV_Assert(weightShape[i + 2] == kernel_size[i]);
}
const Mat &input = inputs[0];
CV_Assert((input.dims == 4 || input.dims == 5) && (input.type() == CV_32F || input.type() == CV_16S));
for (size_t i = 0; i < inputs.size(); i++)
for (size_t i = 0; i < outputs.size(); i++)
{
CV_Assert(inputs[i].type() == input.type());
CV_Assert((inputs[i].dims == 4 || inputs[i].dims == 5) && inputs[i].size[1] == input.size[1]);
@@ -270,6 +271,7 @@ public:
MatShape computeColRowShape(const MatShape &inpShape, const MatShape &outShape) const CV_OVERRIDE
{
CV_Assert(!blobs.empty());
int dims = inpShape.size();
int inpD = dims == 5 ? inpShape[2] : 1;
int inpH = inpShape[dims - 2];
@@ -296,6 +298,8 @@ public:
{
if (kernel_size.size() == 3)
return preferableTarget == DNN_TARGET_CPU;
if ((backendId == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 || preferableTarget != DNN_TARGET_MYRIAD) && blobs.empty())
return false;
return (preferableTarget != DNN_TARGET_MYRIAD || dilation.width == dilation.height);
}
else
@@ -305,7 +309,7 @@ public:
return (preferableTarget == DNN_TARGET_CPU && backendId == DNN_BACKEND_OPENCV);
else if (kernel_size.size() == 2)
return backendId == DNN_BACKEND_OPENCV ||
backendId == DNN_BACKEND_HALIDE ||
(backendId == DNN_BACKEND_HALIDE && !blobs.empty()) ||
(backendId == DNN_BACKEND_VKCOM && haveVulkan());
else
return false;
@@ -317,16 +321,16 @@ public:
std::vector<MatShape> &outputs,
std::vector<MatShape> &internals) const CV_OVERRIDE
{
CV_Assert(blobs.size() != 0);
CV_Assert(!hasBias() || blobs[1].total() == (size_t)blobs[0].size[0]);
CV_Assert(inputs.size() == (size_t)1);
CV_Assert(!blobs.empty() || inputs.size() > 1);
const int* weightShape = blobs.empty() ? &inputs[1][0] : blobs[0].size.p;
CV_Assert(!hasBias() || blobs[1].total() == (size_t)weightShape[0]);
internals.clear();
CV_Assert(inputs.size() != 0);
std::vector<int> inpShape(inputs[0].begin() + 2, inputs[0].end());
int outCn = blobs[0].size[0];
int outCn = weightShape[0];
std::vector<int> outShape;
outShape.push_back(inputs[0][0]);
outShape.push_back(outCn);
@@ -342,10 +346,10 @@ public:
getConvPoolOutParams(inpShape, kernel_size, strides, padMode, dilations, outShape);
}
int ngroups = inpCn / blobs[0].size[1];
if (ngroups == 0 || ngroups * blobs[0].size[1] != inpCn)
int ngroups = inpCn / weightShape[1];
if (ngroups == 0 || ngroups * weightShape[1] != inpCn)
CV_Error(Error::StsError, format("Number of input channels should "
"be multiple of %d but got %d", blobs[0].size[1], inpCn));
"be multiple of %d but got %d", weightShape[1], inpCn));
CV_Assert(ngroups > 0 && inpCn % ngroups == 0 && outCn % ngroups == 0);
outputs.resize(1, outShape);
@@ -357,15 +361,15 @@ public:
{
BaseConvolutionLayerImpl::finalize(inputs_arr, outputs_arr);
CV_Assert(!blobs.empty());
const int outCn = blobs[0].size[0];
std::vector<Mat> inputs;
inputs_arr.getMatVector(inputs);
// prepare weightsMat where each row is aligned and has enough zero padding on the right to
// use vectorized (i.e. with intrinsics) loops without tail processing
Mat wm = blobs[0].reshape(1, outCn);
Mat wm = blobs.empty() ? inputs[1].reshape(1, numOutput) : blobs[0].reshape(1, numOutput);
if( wm.step1() % VEC_ALIGN != 0 )
{
int newcols = (int)alignSize(wm.step1(), VEC_ALIGN);
Mat wm_buffer = Mat(outCn, newcols, wm.type());
Mat wm_buffer = Mat(numOutput, newcols, wm.type());
Mat wm_padding = wm_buffer.colRange(wm.cols, newcols);
wm_padding.setTo(Scalar::all(0.));
Mat wm_aligned = wm_buffer.colRange(0, wm.cols);
@@ -373,18 +377,18 @@ public:
wm = wm_aligned;
}
weightsMat = wm;
weightsMultipliers.assign(outCn, 1.0);
weightsMultipliers.assign(numOutput, 1.0);
Mat biasMat = hasBias() ? blobs[1].reshape(1, outCn) : Mat();
biasvec.resize(outCn+2);
Mat biasMat = hasBias() ? blobs[1].reshape(1, numOutput) : Mat();
biasvec.resize(numOutput+2);
if( biasMat.empty() )
{
for(int i = 0; i < outCn; i++ )
for(int i = 0; i < numOutput; i++ )
biasvec[i] = 0.f;
}
else
{
for(int i = 0; i < outCn; i++ )
for(int i = 0; i < numOutput; i++ )
biasvec[i] = biasMat.at<float>(i);
}
#ifdef HAVE_OPENCL
@@ -394,7 +398,7 @@ public:
bool setActivation(const Ptr<ActivationLayer>& layer) CV_OVERRIDE
{
if (!activ.empty() && !layer.empty())
if ((!activ.empty() && !layer.empty()) || blobs.empty())
return false;
activ = layer;
@@ -743,37 +747,48 @@ public:
virtual Ptr<BackendNode> initNgraph(const std::vector<Ptr<BackendWrapper> > &inputs,
const std::vector<Ptr<BackendNode> >& nodes) CV_OVERRIDE
{
CV_Assert_N(inputs.size() == 1, nodes.size() == 1);
CV_Assert_N(inputs.size() >= 1, nodes.size() >= 1);
auto& ieInpNode = nodes[0].dynamicCast<InfEngineNgraphNode>()->node;
std::vector<size_t> dims = ieInpNode->get_shape();
CV_Assert(dims.size() == 4 || dims.size() == 5);
std::shared_ptr<ngraph::Node> ieWeights = nodes.size() > 1 ? nodes[1].dynamicCast<InfEngineNgraphNode>()->node : nullptr;
const int inpCn = dims[1];
const int outCn = blobs[0].size[0];
const int inpGroupCn = blobs[0].size[1];
const int inpGroupCn = nodes.size() > 1 ? ieWeights->get_shape()[1] : blobs[0].size[1];
const int group = inpCn / inpGroupCn;
std::vector<size_t> kernel_shape = getShape<size_t>(blobs[0]);
std::vector<size_t> kernel_shape;
if (group != 1)
{
kernel_shape[0] /= group;
kernel_shape.insert(kernel_shape.begin(), group);
kernel_shape.push_back(group);
}
kernel_shape.push_back(numOutput / group);
kernel_shape.push_back(inpCn / group);
std::copy(kernel_size.begin(), kernel_size.end(), back_inserter(kernel_shape));
auto ieWeights = std::make_shared<ngraph::op::Constant>(ngraph::element::f32, kernel_shape, blobs[0].data);
if (fusedWeights)
if (nodes.size() == 1)
{
if (weightsMat.isContinuous())
ieWeights = std::make_shared<ngraph::op::Constant>(ngraph::element::f32, kernel_shape, blobs[0].data);
if (fusedWeights)
{
ieWeights = std::make_shared<ngraph::op::Constant>(ngraph::element::f32, kernel_shape, weightsMat.data);
}
else
{
Mat newWeights;
Mat cvWeights = weightsMat.colRange(0, blobs[0].total() / outCn);
cvWeights.copyTo(newWeights);
ieWeights = std::make_shared<ngraph::op::Constant>(ngraph::element::f32, kernel_shape, newWeights.data);
if (weightsMat.isContinuous())
{
ieWeights = std::make_shared<ngraph::op::Constant>(ngraph::element::f32, kernel_shape, weightsMat.data);
}
else
{
Mat newWeights;
Mat cvWeights = weightsMat.colRange(0, blobs[0].total() / numOutput);
cvWeights.copyTo(newWeights);
ieWeights = std::make_shared<ngraph::op::Constant>(ngraph::element::f32, kernel_shape, newWeights.data);
}
}
}
else
{
auto shape = std::make_shared<ngraph::op::Constant>(ngraph::element::i64,
ngraph::Shape{kernel_shape.size()}, kernel_shape.data());
ieWeights = std::make_shared<ngraph::op::v1::Reshape>(ieWeights, shape, true);
}
ngraph::op::PadType pad_type = ngraph::op::PadType::EXPLICIT;
if (!padMode.empty())
@@ -798,11 +813,21 @@ public:
pad_type);
}
if (hasBias() || fusedBias)
if (hasBias() || fusedBias || nodes.size() == 3)
{
std::vector<size_t> shape(conv_node->get_shape().size(), 1);
shape[1] = outCn;
auto bias = std::make_shared<ngraph::op::Constant>(ngraph::element::f32, ngraph::Shape(shape), biasvec.data());
shape[1] = conv_node->get_shape()[1];
std::shared_ptr<ngraph::Node> bias;
if (nodes.size() == 3)
{
auto bias_shape = std::make_shared<ngraph::op::Constant>(ngraph::element::i64,
ngraph::Shape{shape.size()}, shape.data());
bias = std::make_shared<ngraph::op::v1::Reshape>(nodes[2].dynamicCast<InfEngineNgraphNode>()->node, bias_shape, true);
}
else
{
bias = std::make_shared<ngraph::op::Constant>(ngraph::element::f32, ngraph::Shape(shape), biasvec.data());
}
auto conv_bias = std::make_shared<ngraph::op::v1::Add>(conv_node, bias, ngraph::op::AutoBroadcastType::NUMPY);
return Ptr<BackendNode>(new InfEngineNgraphNode(conv_bias));
}
@@ -1516,6 +1541,26 @@ public:
for (int i = 0; i < inputs.size(); ++i)
CV_Assert(inputs[i].u != outputs[0].u);
if (blobs.empty())
{
size_t n = inputs.size() - 1;
umat_blobs.resize(n);
for (size_t i = 0; i < n; i++)
{
if (use_half)
{
Mat matFP32;
convertFp16(inputs[i + 1], matFP32);
matFP32.copyTo(umat_blobs[i]);
}
else
{
inputs[i + 1].copyTo(umat_blobs[i]);
}
}
inputs.resize(1);
}
if (umat_blobs.empty())
{
size_t n = blobs.size();
@@ -1526,7 +1571,7 @@ public:
}
}
if (convolutionOp.empty())
if (convolutionOp.empty() || blobs.empty())
{
OCL4DNNConvConfig config;
config.in_shape = shape(inputs[0]);
@@ -1536,7 +1581,7 @@ public:
config.stride = stride;
config.dilation = dilation;
config.group = inputs[0].size[1] / umat_blobs[0].size[1];
config.bias_term = (hasBias()) ? true : false;
config.bias_term = umat_blobs.size() == 2;
config.use_half = use_half;
convolutionOp = Ptr<OCL4DNNConvSpatial<float> >(new OCL4DNNConvSpatial<float>(config));
@@ -1663,16 +1708,37 @@ public:
inputs_arr.getMatVector(inputs);
outputs_arr.getMatVector(outputs);
int outCn = blobs.empty() ? inputs[1].size[0] : blobs[0].size[0];
// Need to align non-const blobs
if (blobs.empty())
{
Mat wm = inputs[1].reshape(1, outCn);
if( wm.step1() % VEC_ALIGN != 0 )
{
wm.copyTo(weightsMat);
if (inputs.size() > 2)
{
Mat biasMat = inputs[2].reshape(1, outCn);
biasMat.col(0).copyTo(biasvec);
biasvec.resize(outCn + 2);
}
else
{
biasvec.resize(outCn + 2, 0);
}
}
}
/*printf("conv %s: input (%d x %d x %d x %d), kernel (%d x %d), pad (%d x %d), stride (%d x %d), dilation (%d x %d)\n",
name.c_str(), inputs[0].size[0], inputs[0].size[1], inputs[0].size[2], inputs[0].size[3],
kernel.width, kernel.height, pad.width, pad.height,
stride.width, stride.height, dilation.width, dilation.height);*/
CV_Assert_N(inputs.size() == (size_t)1, inputs[0].size[1] % blobs[0].size[1] == 0,
int inpGroupCn = blobs.empty() ? inputs[1].size[1] : blobs[0].size[1];
CV_Assert_N(inputs.size() >= (size_t)1, inputs[0].size[1] % inpGroupCn == 0,
outputs.size() == 1, inputs[0].data != outputs[0].data);
int ngroups = inputs[0].size[1]/blobs[0].size[1];
int ngroups = inputs[0].size[1] / inpGroupCn;
CV_Assert(outputs[0].size[1] % ngroups == 0);
int outCn = blobs[0].size[0];
reluslope.clear();
if( activ )
@@ -1810,11 +1876,11 @@ public:
virtual int64 getFLOPS(const std::vector<MatShape> &inputs,
const std::vector<MatShape> &outputs) const CV_OVERRIDE
{
CV_Assert(inputs.size() == outputs.size());
CV_Assert(inputs.size() == outputs.size() || inputs.size() == outputs.size() + blobs.size());
int64 flops = 0;
int karea = std::accumulate(kernel_size.begin(), kernel_size.end(), 1, std::multiplies<size_t>());
for (int i = 0; i < inputs.size(); i++)
for (int i = 0; i < outputs.size(); i++)
{
flops += total(outputs[i])*(CV_BIG_INT(2)*karea*inputs[i][1] + 1);
}
modules/dnn/src/layers/fully_connected_layer.cpp
+1 -1
View File
@@ -587,7 +587,7 @@ public:
}
else
{
std::vector<size_t> data = {(size_t)ieInpNode->get_shape()[0], (size_t)blobs[0].size[1]};
std::vector<int64_t> data = {(int64_t)ieInpNode->get_shape()[0], (int64_t)blobs[0].size[1]};
auto new_shape = std::make_shared<ngraph::op::Constant>(ngraph::element::i64, ngraph::Shape{2}, data.data());
auto inp = std::make_shared<ngraph::op::v1::Reshape>(ieInpNode, new_shape, true);
modules/dnn/src/layers/permute_layer.cpp
+2 -1
View File
@@ -397,8 +397,9 @@ public:
const std::vector<Ptr<BackendNode> >& nodes) CV_OVERRIDE
{
auto& ieInpNode = nodes[0].dynamicCast<InfEngineNgraphNode>()->node;
std::vector<int64_t> order(_order.begin(), _order.end());
auto tr_axes = std::make_shared<ngraph::op::Constant>(ngraph::element::i64,
ngraph::Shape({_order.size()}), _order.data());
ngraph::Shape({order.size()}), order.data());
auto transpose = std::make_shared<ngraph::op::Transpose>(ieInpNode, tr_axes);
return Ptr<BackendNode>(new InfEngineNgraphNode(transpose));
}
modules/dnn/src/layers/slice_layer.cpp
+110 -19
View File
@@ -167,6 +167,10 @@ public:
void finalize(InputArrayOfArrays inputs_arr, OutputArrayOfArrays outputs_arr) CV_OVERRIDE
{
#ifdef HAVE_OPENCL
ocl_exec_cache.clear();
#endif
std::vector<Mat> inputs, outputs;
inputs_arr.getMatVector(inputs);
outputs_arr.getMatVector(outputs);
@@ -221,26 +225,33 @@ public:
}
#ifdef HAVE_OPENCL
bool forward_ocl(InputArrayOfArrays inputs_, OutputArrayOfArrays outputs_, OutputArrayOfArrays internals_)
struct OpenCLExecInfo
{
std::vector<UMat> inputs;
std::vector<UMat> outputs;
std::string kernel_name;
std::string build_opts;
size_t local_size[2];
size_t global_size[2];
inputs_.getUMatVector(inputs);
outputs_.getUMatVector(outputs);
OpenCLExecInfo()
{
local_size[0] = local_size[1] = 0;
global_size[0] = global_size[1] = 0;
}
};
std::vector<OpenCLExecInfo> ocl_exec_cache;
void ocl_prepare(const std::vector<UMat>& inputs, const std::vector<UMat>& outputs)
{
CV_TRACE_FUNCTION();
CV_Assert(outputs.size() == finalSliceRanges.size());
ocl_exec_cache.resize(outputs.size());
const UMat& input = inputs[0];
if (input.dims > 5)
{
CV_LOG_INFO(NULL, "DNN/OpenCL/Slice: implementation doesn't support dims=" << input.dims << ". Fallback to CPU");
return false;
}
const int dims = input.dims;
size_t WSZ = 128;
const int dims = input.dims;
const int elemSize = (int)input.elemSize();
String opts0 = cv::format(
"-DDIMS=%d -DELEMSIZE=%d",
@@ -250,10 +261,11 @@ public:
{
opts0 += cv::format(" -DSRC_STEP_%d=%d", d, (int)input.step[dims - 1 - d]);
}
String kname = cv::format("slice_%d", dims);
for (size_t i = 0; i < outputs.size(); i++)
{
UMat& output = outputs[i];
OpenCLExecInfo& ocl = ocl_exec_cache[i];
const UMat& output = outputs[i];
const std::vector<Range>& range = finalSliceRanges[i];
String opts = opts0;
@@ -269,6 +281,8 @@ public:
CV_CheckEQ(range[d].size(), (int)output.size[d], "");
}
const size_t param_LIMIT_BLOCK_SIZE_PER_WG = WSZ * 64;
int block_dims = 0;
size_t block_size = elemSize;
for (int i = dims - 1; i >= 0; --i)
@@ -277,12 +291,14 @@ public:
break;
block_size *= output.size[i];
block_dims++;
if (block_size >= param_LIMIT_BLOCK_SIZE_PER_WG)
break;
}
const size_t total = output.total() * elemSize;
size_t num_blocks = total / block_size;
if ((num_blocks <= 8 && block_size >= WSZ * 4) || (block_size >= WSZ * 64))
if ((num_blocks <= 8 && block_size >= WSZ * 4) || (block_size >= param_LIMIT_BLOCK_SIZE_PER_WG))
{
// use 1D copy mode
opts += cv::format(" -DUSE_COPY_1D=1");
@@ -352,23 +368,98 @@ public:
opts += cv::format(" -DWSZ=%d", (int)WSZ);
size_t local[] = { WSZ, 1 };
size_t global[] = { WSZ, num_blocks };
std::ostringstream kernel_suffix;
kernel_suffix << dims << 'x' << elemSize << "_bsz" << block_size;
kernel_suffix << "__src_";
for (int d = 0; d < dims; d++)
{
kernel_suffix << input.size[dims - 1 - d] << '_';
}
kernel_suffix << '_';
/*for (int d = 0; d < dims; d++)
{
kernel_suffix << input.step[dims - 1 - d] << '_';
}
kernel_suffix << '_';*/
ocl::Kernel kernel(kname.c_str(), ocl::dnn::slice_oclsrc, opts);
kernel_suffix << "dst_";
for (int d = 0; d < dims; d++)
{
kernel_suffix << output.size[dims - 1 - d] << '_';
}
/*kernel_suffix << '_';
for (int d = 0; d < dims; d++)
{
kernel_suffix << output.step[dims - 1 - d] << '_';
}*/
kernel_suffix << "_slice_";
for (int d = 0; d < dims; d++)
{
kernel_suffix << range[dims - 1 - d].start << '_';
}
for (int d = 0; d < dims; d++)
{
kernel_suffix << '_' << range[dims - 1 - d].end;
}
std::string kernel_suffix_str = kernel_suffix.str();
opts += cv::format(" -DSLICE_KERNEL_SUFFIX=%s", kernel_suffix_str.c_str());
ocl.kernel_name = cv::format("slice_%s", kernel_suffix_str.c_str());
ocl.build_opts = opts;
ocl.local_size[0] = WSZ;
ocl.local_size[1] = 1;
ocl.global_size[0] = WSZ;
ocl.global_size[1] = num_blocks;
} // for outputs.size()
} // ocl_prepare
bool forward_ocl(InputArrayOfArrays inputs_, OutputArrayOfArrays outputs_, OutputArrayOfArrays internals_)
{
CV_TRACE_FUNCTION();
std::vector<UMat> inputs;
std::vector<UMat> outputs;
inputs_.getUMatVector(inputs);
outputs_.getUMatVector(outputs);
CV_Assert(outputs.size() == finalSliceRanges.size());
const UMat& input = inputs[0];
const int dims = input.dims;
if (dims > 5)
{
CV_LOG_INFO(NULL, "DNN/OpenCL/Slice: implementation doesn't support dims=" << dims << ". Fallback to CPU");
return false;
}
if (ocl_exec_cache.empty())
{
ocl_prepare(inputs, outputs);
}
CV_CheckEQ(ocl_exec_cache.size(), outputs.size(), "");
for (size_t i = 0; i < outputs.size(); i++)
{
const OpenCLExecInfo& ocl = ocl_exec_cache[i];
UMat& output = outputs[i];
ocl::Kernel kernel(ocl.kernel_name.c_str(), ocl::dnn::slice_oclsrc, ocl.build_opts);
if (kernel.empty())
return false;
bool ret = kernel.args(
ocl::KernelArg::PtrReadOnly(input),
ocl::KernelArg::PtrWriteOnly(output)
)
.run(2, global, local, false);
.run(2, (size_t*)ocl.global_size, (size_t*)ocl.local_size, false);
if (!ret)
return false;
} // for outputs.size()
return true;
}
} // forward_ocl
#endif
void forward(InputArrayOfArrays inputs_arr, OutputArrayOfArrays outputs_arr, OutputArrayOfArrays internals_arr) CV_OVERRIDE
modules/dnn/src/onnx/onnx_importer.cpp
+6 -3
View File
@@ -1003,10 +1003,13 @@ void ONNXImporter::populateNet(Net dstNet)
CV_Assert(node_proto.input_size() >= 2);
layerParams.type = "Convolution";
for (int j = 1; j < node_proto.input_size(); j++) {
layerParams.blobs.push_back(getBlob(node_proto, constBlobs, j));
if (constBlobs.find(node_proto.input(j)) != constBlobs.end())
{
layerParams.blobs.push_back(getBlob(node_proto, constBlobs, j));
}
}
layerParams.set("num_output", layerParams.blobs[0].size[0]);
layerParams.set("bias_term", node_proto.input_size() == 3);
int outCn = layerParams.blobs.empty() ? outShapes[node_proto.input(1)][0] : layerParams.blobs[0].size[0];
layerParams.set("num_output", outCn);
}
else if (layer_type == "ConvTranspose")
{
modules/dnn/src/opencl/slice.cl
+86 -105
View File
@@ -48,19 +48,85 @@ global: <WSZ, number_of_copy_blocks, 1>
#define BLOCK_COLS_X4 (BLOCK_COLS / 4)
#define BLOCK_COLS_X16 (BLOCK_COLS / 16)
#ifdef USE_COPY_1D
static inline
__attribute__((always_inline))
void copy_block_1d(
__attribute__((reqd_work_group_size(WSZ, 1, 1)))
__kernel void
CONCAT(slice_, SLICE_KERNEL_SUFFIX)(
__global const uchar* src0,
const uint src_offset,
__global uchar* dst0,
const uint dst_offset
__global uchar* dst0
)
{
__global const uchar* src = src0 + src_offset;
__global uchar* dst = dst0 + dst_offset;
uint block_id = get_global_id(1);
uint dst_offset0 = block_id * BLOCK_SIZE;
uint src_offset0 = 0;
{ // calculate src_offset0
#define CALC_SRC_INDEX(dim) \
{ \
uint plane_sz = CONCAT(DST_STEP_, dim) / BLOCK_SIZE; \
CONCAT(idx_, dim) = block_id / plane_sz; \
block_id = block_id - CONCAT(idx_, dim) * plane_sz; \
}
#define UPDATE_SRC_OFFSET(dim) \
src_offset0 = mad24((uint)(CONCAT(idx_, dim) + CONCAT(SRC_START_, dim)), (uint)CONCAT(SRC_STEP_, dim), (uint)src_offset0);
/*
if (get_global_id(0) == 0 && get_global_id(1) == 0) \
printf("(%d, %d): @%d src_offset0=%d idx_dim=%d block_id=%d\n", \
get_global_id(0), get_global_id(1), \
dim, src_offset0, CONCAT(idx_, dim), block_id \
);
*/
#if DIMS > 5
#error "invalid configuration"
#endif
#if DIMS > 4
uint idx_4 = 0;
#if BLOCK_DIMS <= 4
CALC_SRC_INDEX(4)
#endif
UPDATE_SRC_OFFSET(4)
#endif
#if DIMS > 3
uint idx_3 = 0;
#if BLOCK_DIMS <= 3
CALC_SRC_INDEX(3)
#endif
UPDATE_SRC_OFFSET(3)
#endif
#if DIMS > 2
uint idx_2 = 0;
#if BLOCK_DIMS <= 2
CALC_SRC_INDEX(2)
#endif
UPDATE_SRC_OFFSET(2)
#endif
#if DIMS > 1
uint idx_1 = 0;
#if BLOCK_DIMS <= 1
CALC_SRC_INDEX(1)
#endif
UPDATE_SRC_OFFSET(1)
#endif
#if DIMS > 0
uint idx_0 = 0;
UPDATE_SRC_OFFSET(0)
#endif
/*
if (get_global_id(0) == 0)
printf("(%d, %d): src_offset0=%d dst_offset0=%d\n",
get_global_id(0), get_global_id(1),
src_offset0, dst_offset0
);
*/
} // calculate src_offset0
#ifdef USE_COPY_1D
{ // copy_block_1d
__global const uchar* src = src0 + src_offset0;
__global uchar* dst = dst0 + dst_offset0;
uint processed = 0;
@@ -70,8 +136,9 @@ void copy_block_1d(
uint i = get_local_id(0) * 16; // uchar16
while (i < BLOCK_COLS_X16 * 16)
{
uint4 idx = (uint4)(i, i + 16 * WSZ, i + 32 * WSZ, i + 48 * WSZ);
idx = select((uint4)i, idx, idx < (BLOCK_COLS_X16 * 16));
uint4 idx0 = (uint4)i;
uint4 idx = idx0 + (uint4)(0, 16 * WSZ, 32 * WSZ, 48 * WSZ);
idx = select(idx0, idx, idx < (BLOCK_COLS_X16 * 16));
uchar16 a0 = vload16(0, src + idx.s0);
uchar16 a1 = vload16(0, src + idx.s1);
@@ -97,8 +164,9 @@ void copy_block_1d(
uint i = get_local_id(0) * 4 + processed; // uchar4
while (i < BLOCK_COLS_X4 * 4)
{
uint4 idx = (uint4)(i, i + 4 * WSZ, i + 8 * WSZ, i + 12 * WSZ);
idx = select((uint4)i, idx, idx < (BLOCK_COLS_X4 * 4));
uint4 idx0 = (uint4)i;
uint4 idx = idx0 + (uint4)(0, 4 * WSZ, 8 * WSZ, 12 * WSZ);
idx = select(idx0, idx, idx < (BLOCK_COLS_X4 * 4));
uchar4 a0 = vload4(0, src + idx.s0);
uchar4 a1 = vload4(0, src + idx.s1);
@@ -130,19 +198,11 @@ void copy_block_1d(
}
}
#endif
}
} // copy_block_1d
#else // USE_COPY_1D
#else
static inline
__attribute__((always_inline))
void copy_block_2d(
__global const uchar* src0,
const uint src_offset0,
__global uchar* dst0,
const uint dst_offset0
)
{
{ // copy_block_2d
__global const uchar* src = src0 + src_offset0;
__global uchar* dst = dst0 + dst_offset0;
@@ -199,85 +259,6 @@ void copy_block_2d(
#endif // BLOCK_COLS_FILL_X4 != BLOCK_COLS
i += WSZ * 4;
}
}
#endif // USE_COPY_1D
__kernel void
CONCAT(slice_, DIMS)(
__global const uchar* src,
__global uchar* dst
)
{
uint block_id = get_global_id(1);
uint dst_offset = block_id * BLOCK_SIZE;
uint src_offset = 0;
#define CALC_SRC_INDEX(dim) \
{ \
uint plane_sz = CONCAT(DST_STEP_, dim) / BLOCK_SIZE; \
CONCAT(idx_, dim) = block_id / plane_sz; \
block_id = block_id - CONCAT(idx_, dim) * plane_sz; \
}
#define UPDATE_SRC_OFFSET(dim) \
src_offset = mad24((uint)(CONCAT(idx_, dim) + CONCAT(SRC_START_, dim)), (uint)CONCAT(SRC_STEP_, dim), (uint)src_offset);
/*
if (get_global_id(0) == 0 && get_global_id(1) == 0) \
printf("(%d, %d): @%d src_offset=%d idx_dim=%d block_id=%d\n", \
get_global_id(0), get_global_id(1), \
dim, src_offset, CONCAT(idx_, dim), block_id \
);
*/
#if DIMS > 5
#error "invalid configuration"
#endif
#if DIMS > 4
uint idx_4 = 0;
#if BLOCK_DIMS <= 4
CALC_SRC_INDEX(4)
#endif
UPDATE_SRC_OFFSET(4)
#endif
#if DIMS > 3
uint idx_3 = 0;
#if BLOCK_DIMS <= 3
CALC_SRC_INDEX(3)
#endif
UPDATE_SRC_OFFSET(3)
#endif
#if DIMS > 2
uint idx_2 = 0;
#if BLOCK_DIMS <= 2
CALC_SRC_INDEX(2)
#endif
UPDATE_SRC_OFFSET(2)
#endif
#if DIMS > 1
uint idx_1 = 0;
#if BLOCK_DIMS <= 1
CALC_SRC_INDEX(1)
#endif
UPDATE_SRC_OFFSET(1)
#endif
#if DIMS > 0
uint idx_0 = 0;
UPDATE_SRC_OFFSET(0)
#endif
/*
if (get_global_id(0) == 0)
printf("(%d, %d): src_offset=%d dst_offset=%d\n",
get_global_id(0), get_global_id(1),
src_offset, dst_offset
);
*/
#ifdef USE_COPY_1D
copy_block_1d(src, src_offset, dst, dst_offset);
#else
copy_block_2d(src, src_offset, dst, dst_offset);
} // copy_block_2d
#endif
}
modules/dnn/test/test_darknet_importer.cpp
+5
View File
@@ -784,6 +784,11 @@ TEST_P(Test_Darknet_layers, connected)
testDarknetLayer("connected", true);
}
TEST_P(Test_Darknet_layers, relu)
{
testDarknetLayer("relu");
}
INSTANTIATE_TEST_CASE_P(/**/, Test_Darknet_layers, dnnBackendsAndTargets());
}} // namespace
modules/dnn/test/test_layers.cpp
+18 -11
View File
@@ -1133,6 +1133,9 @@ TEST_P(Layer_Test_Convolution_DLDT, Accuracy)
const Backend backendId = get<0>(GetParam());
const Target targetId = get<1>(GetParam());
if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && targetId == DNN_TARGET_MYRIAD)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_MYRIAD, CV_TEST_TAG_DNN_SKIP_IE_NN_BUILDER);
if (backendId != DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && backendId != DNN_BACKEND_INFERENCE_ENGINE_NGRAPH)
throw SkipTestException("No support for async forward");
@@ -1143,9 +1146,8 @@ TEST_P(Layer_Test_Convolution_DLDT, Accuracy)
else
FAIL() << "Unknown backendId";
std::string suffix = (targetId == DNN_TARGET_OPENCL_FP16 || targetId == DNN_TARGET_MYRIAD) ? "_fp16" : "";
Net netDefault = readNet(_tf("layer_convolution.caffemodel"), _tf("layer_convolution.prototxt"));
Net net = readNet(_tf("layer_convolution" + suffix + ".xml"), _tf("layer_convolution" + suffix + ".bin"));
Net net = readNet(_tf("layer_convolution.xml"), _tf("layer_convolution.bin"));
Mat inp = blobFromNPY(_tf("blob.npy"));
@@ -1165,7 +1167,10 @@ TEST_P(Layer_Test_Convolution_DLDT, Accuracy)
std::vector<int> outLayers = net.getUnconnectedOutLayers();
ASSERT_EQ(net.getLayer(outLayers[0])->name, "output");
ASSERT_EQ(net.getLayer(outLayers[0])->type, "Convolution");
if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019)
ASSERT_EQ(net.getLayer(outLayers[0])->type, "Convolution");
else
ASSERT_EQ(net.getLayer(outLayers[0])->type, "Add");
}
TEST_P(Layer_Test_Convolution_DLDT, setInput_uint8)
@@ -1173,6 +1178,9 @@ TEST_P(Layer_Test_Convolution_DLDT, setInput_uint8)
const Backend backendId = get<0>(GetParam());
const Target targetId = get<1>(GetParam());
if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && targetId == DNN_TARGET_MYRIAD)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_MYRIAD, CV_TEST_TAG_DNN_SKIP_IE_NN_BUILDER);
if (backendId != DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && backendId != DNN_BACKEND_INFERENCE_ENGINE_NGRAPH)
throw SkipTestException("No support for async forward");
@@ -1189,12 +1197,10 @@ TEST_P(Layer_Test_Convolution_DLDT, setInput_uint8)
randu(inputs[0], 0, 255);
inputs[0].convertTo(inputs[1], CV_32F);
std::string suffix = (targetId == DNN_TARGET_OPENCL_FP16 || targetId == DNN_TARGET_MYRIAD) ? "_fp16" : "";
Mat outs[2];
for (int i = 0; i < 2; ++i)
{
Net net = readNet(_tf("layer_convolution" + suffix + ".xml"), _tf("layer_convolution" + suffix + ".bin"));
Net net = readNet(_tf("layer_convolution.xml"), _tf("layer_convolution.bin"));
net.setPreferableBackend(backendId);
net.setPreferableTarget(targetId);
net.setInput(inputs[i]);
@@ -1210,6 +1216,9 @@ TEST_P(Layer_Test_Convolution_DLDT, multithreading)
const Backend backendId = get<0>(GetParam());
const Target targetId = get<1>(GetParam());
if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && targetId == DNN_TARGET_MYRIAD)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_MYRIAD, CV_TEST_TAG_DNN_SKIP_IE_NN_BUILDER);
if (backendId != DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && backendId != DNN_BACKEND_INFERENCE_ENGINE_NGRAPH)
throw SkipTestException("No support for async forward");
@@ -1220,9 +1229,8 @@ TEST_P(Layer_Test_Convolution_DLDT, multithreading)
else
FAIL() << "Unknown backendId";
std::string suffix = (targetId == DNN_TARGET_OPENCL_FP16 || targetId == DNN_TARGET_MYRIAD) ? "_fp16" : "";
std::string xmlPath = _tf("layer_convolution" + suffix + ".xml");
std::string binPath = _tf("layer_convolution" + suffix + ".bin");
std::string xmlPath = _tf("layer_convolution.xml");
std::string binPath = _tf("layer_convolution.bin");
Net firstNet = readNet(xmlPath, binPath);
Net secondNet = readNet(xmlPath, binPath);
Mat inp = blobFromNPY(_tf("blob.npy"));
@@ -1281,8 +1289,7 @@ TEST_P(Test_DLDT_two_inputs_3dim, as_IR)
int secondInpType = get<1>(GetParam());
Target targetId = get<2>(GetParam());
std::string suffix = (targetId == DNN_TARGET_OPENCL_FP16 || targetId == DNN_TARGET_MYRIAD) ? "_fp16" : "";
Net net = readNet(_tf("net_two_inputs" + suffix + ".xml"), _tf("net_two_inputs.bin"));
Net net = readNet(_tf("net_two_inputs.xml"), _tf("net_two_inputs.bin"));
std::vector<int> inpSize = get<3>(GetParam());
Mat firstInp(3, inpSize.data(), firstInpType);
Mat secondInp(3, inpSize.data(), secondInpType);
modules/dnn/test/test_misc.cpp
+25 -14
View File
@@ -444,12 +444,14 @@ TEST_P(Async, model_optimizer_pipeline_set_and_forward_single)
const Backend backendId = get<0>(get<1>(GetParam()));
const Target targetId = get<1>(get<1>(GetParam()));
if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && targetId == DNN_TARGET_MYRIAD)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_MYRIAD, CV_TEST_TAG_DNN_SKIP_IE_NN_BUILDER);
if (backendId != DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && backendId != DNN_BACKEND_INFERENCE_ENGINE_NGRAPH)
throw SkipTestException("No support for async forward");
const std::string suffix = (targetId == DNN_TARGET_OPENCL_FP16 || targetId == DNN_TARGET_MYRIAD) ? "_fp16" : "";
const std::string& model = findDataFile("dnn/layers/layer_convolution" + suffix + ".bin");
const std::string& proto = findDataFile("dnn/layers/layer_convolution" + suffix + ".xml");
const std::string& model = findDataFile("dnn/layers/layer_convolution.bin");
const std::string& proto = findDataFile("dnn/layers/layer_convolution.xml");
if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019)
setInferenceEngineBackendType(CV_DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_API);
@@ -503,12 +505,14 @@ TEST_P(Async, model_optimizer_pipeline_set_and_forward_all)
const Backend backendId = get<0>(get<1>(GetParam()));
const Target targetId = get<1>(get<1>(GetParam()));
if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && targetId == DNN_TARGET_MYRIAD)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_MYRIAD, CV_TEST_TAG_DNN_SKIP_IE_NN_BUILDER);
if (backendId != DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && backendId != DNN_BACKEND_INFERENCE_ENGINE_NGRAPH)
throw SkipTestException("No support for async forward");
const std::string suffix = (targetId == DNN_TARGET_OPENCL_FP16 || targetId == DNN_TARGET_MYRIAD) ? "_fp16" : "";
const std::string& model = findDataFile("dnn/layers/layer_convolution" + suffix + ".bin");
const std::string& proto = findDataFile("dnn/layers/layer_convolution" + suffix + ".xml");
const std::string& model = findDataFile("dnn/layers/layer_convolution.bin");
const std::string& proto = findDataFile("dnn/layers/layer_convolution.xml");
if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019)
setInferenceEngineBackendType(CV_DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_API);
@@ -677,9 +681,11 @@ TEST_P(Test_Model_Optimizer, forward_two_nets)
const Backend backendId = get<0>(GetParam());
const Target targetId = get<1>(GetParam());
const std::string suffix = (targetId == DNN_TARGET_OPENCL_FP16 || targetId == DNN_TARGET_MYRIAD) ? "_fp16" : "";
const std::string& model = findDataFile("dnn/layers/layer_convolution" + suffix + ".bin");
const std::string& proto = findDataFile("dnn/layers/layer_convolution" + suffix + ".xml");
if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && targetId == DNN_TARGET_MYRIAD)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_MYRIAD, CV_TEST_TAG_DNN_SKIP_IE_NN_BUILDER);
const std::string& model = findDataFile("dnn/layers/layer_convolution.bin");
const std::string& proto = findDataFile("dnn/layers/layer_convolution.xml");
if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019)
setInferenceEngineBackendType(CV_DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_API);
@@ -716,12 +722,14 @@ TEST_P(Test_Model_Optimizer, readFromBuffer)
const Backend backendId = get<0>(GetParam());
const Target targetId = get<1>(GetParam());
if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && targetId == DNN_TARGET_MYRIAD)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_MYRIAD, CV_TEST_TAG_DNN_SKIP_IE_NN_BUILDER);
if (backendId != DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && backendId != DNN_BACKEND_INFERENCE_ENGINE_NGRAPH)
throw SkipTestException("No support for async forward");
const std::string suffix = (targetId == DNN_TARGET_OPENCL_FP16 || targetId == DNN_TARGET_MYRIAD) ? "_fp16" : "";
const std::string& weightsFile = findDataFile("dnn/layers/layer_convolution" + suffix + ".bin");
const std::string& modelFile = findDataFile("dnn/layers/layer_convolution" + suffix + ".xml");
const std::string& weightsFile = findDataFile("dnn/layers/layer_convolution.bin");
const std::string& modelFile = findDataFile("dnn/layers/layer_convolution.xml");
if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019)
setInferenceEngineBackendType(CV_DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_API);
@@ -769,8 +777,11 @@ TEST_P(Test_Model_Optimizer, flexible_inputs)
const Backend backendId = get<0>(GetParam());
const Target targetId = get<1>(GetParam());
const std::string& model = findDataFile("dnn/layers/layer_convolution_fp16.bin");
const std::string& proto = findDataFile("dnn/layers/layer_convolution_fp16.xml");
if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && targetId == DNN_TARGET_MYRIAD)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_MYRIAD, CV_TEST_TAG_DNN_SKIP_IE_NN_BUILDER);
const std::string& model = findDataFile("dnn/layers/layer_convolution.bin");
const std::string& proto = findDataFile("dnn/layers/layer_convolution.xml");
if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019)
setInferenceEngineBackendType(CV_DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_API);
modules/dnn/test/test_onnx_importer.cpp
+56
View File
@@ -114,6 +114,62 @@ TEST_P(Test_ONNX_layers, Convolution)
testONNXModels("convolution");
}
TEST_P(Test_ONNX_layers, Convolution_variable_weight)
{
if ((backend == DNN_BACKEND_INFERENCE_ENGINE_NGRAPH ||
backend == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019) && target == DNN_TARGET_MYRIAD)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_MYRIAD, CV_TEST_TAG_DNN_SKIP_IE_NN_BUILDER, CV_TEST_TAG_DNN_SKIP_IE_NGRAPH);
String basename = "conv_variable_w";
Net net = readNetFromONNX(_tf("models/" + basename + ".onnx"));
ASSERT_FALSE(net.empty());
net.setPreferableBackend(backend);
net.setPreferableTarget(target);
for (int i = 0; i < 2; i++)
{
Mat input = blobFromNPY(_tf("data/input_" + basename + format("_%d", i) + "_0.npy"));
Mat weights = blobFromNPY(_tf("data/input_" + basename + format("_%d", i) + "_1.npy"));
Mat ref = blobFromNPY(_tf("data/output_" + basename + format("_%d", i) + ".npy"));
net.setInput(input, "0");
net.setInput(weights, "1");
Mat out = net.forward();
normAssert(ref, out, "", default_l1, default_lInf);
}
}
TEST_P(Test_ONNX_layers, Convolution_variable_weight_bias)
{
if ((backend == DNN_BACKEND_INFERENCE_ENGINE_NGRAPH ||
backend == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019) && target == DNN_TARGET_MYRIAD)
applyTestTag(CV_TEST_TAG_DNN_SKIP_IE_MYRIAD, CV_TEST_TAG_DNN_SKIP_IE_NN_BUILDER, CV_TEST_TAG_DNN_SKIP_IE_NGRAPH);
String basename = "conv_variable_wb";
Net net = readNetFromONNX(_tf("models/" + basename + ".onnx"));
ASSERT_FALSE(net.empty());
net.setPreferableBackend(backend);
net.setPreferableTarget(target);
for (int i = 0; i < 2; i++)
{
Mat input = blobFromNPY(_tf("data/input_" + basename + format("_%d", i) + "_0.npy"));
Mat weights = blobFromNPY(_tf("data/input_" + basename + format("_%d", i) + "_1.npy"));
Mat bias = blobFromNPY(_tf("data/input_" + basename + format("_%d", i) + "_2.npy"));
Mat ref = blobFromNPY(_tf("data/output_" + basename + format("_%d", i) + ".npy"));
net.setInput(input, "0");
net.setInput(weights, "1");
net.setInput(bias, "bias");
Mat out = net.forward();
normAssert(ref, out, "", default_l1, default_lInf);
}
}
TEST_P(Test_ONNX_layers, Gather)
{
if (backend == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019 && target == DNN_TARGET_MYRIAD)