1
0
mirror of https://github.com/opencv/opencv.git synced 2026-07-29 15:23:05 +04:00

Merge pull request #12052 from dkurt:dnn_ie_torch_tests

This commit is contained in:
Vadim Pisarevsky
2018-07-26 09:09:35 +00:00
12 changed files with 309 additions and 159 deletions
+8
View File
@@ -3071,6 +3071,14 @@ std::vector<Mat> Layer::finalize(const std::vector<Mat> &inputs)
return outputs;
}
void Layer::forward(InputArrayOfArrays inputs, OutputArrayOfArrays outputs, OutputArrayOfArrays internals)
{
CV_TRACE_FUNCTION();
CV_TRACE_ARG_VALUE(name, "name", name.c_str());
Layer::forward_fallback(inputs, outputs, internals);
}
void Layer::forward_fallback(InputArrayOfArrays inputs_arr, OutputArrayOfArrays outputs_arr, OutputArrayOfArrays internals_arr)
{
CV_TRACE_FUNCTION();
@@ -196,7 +196,7 @@ public:
virtual bool supportBackend(int backendId) CV_OVERRIDE
{
return backendId == DNN_BACKEND_OPENCV ||
backendId == DNN_BACKEND_INFERENCE_ENGINE && haveInfEngine() && !_locPredTransposed;
backendId == DNN_BACKEND_INFERENCE_ENGINE && !_locPredTransposed && _bboxesNormalized;
}
bool getMemoryShapes(const std::vector<MatShape> &inputs,
@@ -411,9 +411,12 @@ public:
CV_TRACE_FUNCTION();
CV_TRACE_ARG_VALUE(name, "name", name.c_str());
CV_OCL_RUN(IS_DNN_OPENCL_TARGET(preferableTarget) &&
OCL_PERFORMANCE_CHECK(ocl::Device::getDefault().isIntel()),
forward_ocl(inputs_arr, outputs_arr, internals_arr))
if (_bboxesNormalized)
{
CV_OCL_RUN(IS_DNN_OPENCL_TARGET(preferableTarget) &&
OCL_PERFORMANCE_CHECK(ocl::Device::getDefault().isIntel()),
forward_ocl(inputs_arr, outputs_arr, internals_arr))
}
Layer::forward_fallback(inputs_arr, outputs_arr, internals_arr);
}
+49 -20
View File
@@ -135,10 +135,17 @@ public:
virtual bool supportBackend(int backendId) CV_OVERRIDE
{
return backendId == DNN_BACKEND_OPENCV ||
backendId == DNN_BACKEND_HALIDE && haveHalide() &&
(type == MAX || type == AVE && !pad.width && !pad.height) ||
backendId == DNN_BACKEND_INFERENCE_ENGINE && haveInfEngine() && (type == MAX || type == AVE);
if (backendId == DNN_BACKEND_INFERENCE_ENGINE)
{
if (preferableTarget == DNN_TARGET_MYRIAD)
return type == MAX || type == AVE;
else
return type != STOCHASTIC;
}
else
return backendId == DNN_BACKEND_OPENCV ||
backendId == DNN_BACKEND_HALIDE && haveHalide() &&
(type == MAX || type == AVE && !pad.width && !pad.height);
}
#ifdef HAVE_OPENCL
@@ -192,8 +199,11 @@ public:
CV_TRACE_FUNCTION();
CV_TRACE_ARG_VALUE(name, "name", name.c_str());
CV_OCL_RUN(IS_DNN_OPENCL_TARGET(preferableTarget),
forward_ocl(inputs_arr, outputs_arr, internals_arr))
if (type == MAX || type == AVE || type == STOCHASTIC)
{
CV_OCL_RUN(IS_DNN_OPENCL_TARGET(preferableTarget),
forward_ocl(inputs_arr, outputs_arr, internals_arr))
}
Layer::forward_fallback(inputs_arr, outputs_arr, internals_arr);
}
@@ -238,22 +248,41 @@ public:
#ifdef HAVE_INF_ENGINE
InferenceEngine::LayerParams lp;
lp.name = name;
lp.type = "Pooling";
lp.precision = InferenceEngine::Precision::FP32;
std::shared_ptr<InferenceEngine::PoolingLayer> ieLayer(new InferenceEngine::PoolingLayer(lp));
ieLayer->_kernel_x = kernel.width;
ieLayer->_kernel_y = kernel.height;
ieLayer->_stride_x = stride.width;
ieLayer->_stride_y = stride.height;
ieLayer->_padding_x = pad.width;
ieLayer->_padding_y = pad.height;
ieLayer->_exclude_pad = type == AVE && padMode == "SAME";
ieLayer->params["rounding-type"] = ceilMode ? "ceil" : "floor";
if (type == MAX)
ieLayer->_type = InferenceEngine::PoolingLayer::PoolType::MAX;
else if (type == AVE)
ieLayer->_type = InferenceEngine::PoolingLayer::PoolType::AVG;
std::shared_ptr<InferenceEngine::CNNLayer> ieLayer;
if (type == MAX || type == AVE)
{
lp.type = "Pooling";
InferenceEngine::PoolingLayer* poolLayer = new InferenceEngine::PoolingLayer(lp);
poolLayer->_kernel_x = kernel.width;
poolLayer->_kernel_y = kernel.height;
poolLayer->_stride_x = stride.width;
poolLayer->_stride_y = stride.height;
poolLayer->_padding_x = pad.width;
poolLayer->_padding_y = pad.height;
poolLayer->_exclude_pad = type == AVE && padMode == "SAME";
poolLayer->params["rounding-type"] = ceilMode ? "ceil" : "floor";
poolLayer->_type = type == MAX ? InferenceEngine::PoolingLayer::PoolType::MAX :
InferenceEngine::PoolingLayer::PoolType::AVG;
ieLayer = std::shared_ptr<InferenceEngine::CNNLayer>(poolLayer);
}
else if (type == ROI)
{
lp.type = "ROIPooling";
ieLayer = std::shared_ptr<InferenceEngine::CNNLayer>(new InferenceEngine::CNNLayer(lp));
ieLayer->params["pooled_w"] = format("%d", pooledSize.width);
ieLayer->params["pooled_h"] = format("%d", pooledSize.height);
ieLayer->params["spatial_scale"] = format("%f", spatialScale);
}
else if (type == PSROI)
{
lp.type = "PSROIPooling";
ieLayer = std::shared_ptr<InferenceEngine::CNNLayer>(new InferenceEngine::CNNLayer(lp));
ieLayer->params["output_dim"] = format("%d", psRoiOutChannels);
ieLayer->params["group_size"] = format("%d", pooledSize.width);
ieLayer->params["spatial_scale"] = format("%f", spatialScale);
}
else
CV_Error(Error::StsNotImplemented, "Unsupported pooling type");
+48 -7
View File
@@ -6,6 +6,7 @@
// Third party copyrights are property of their respective owners.
#include "../precomp.hpp"
#include "layers_common.hpp"
#include "../op_inf_engine.hpp"
namespace cv { namespace dnn {
@@ -16,14 +17,14 @@ public:
{
setParamsFrom(params);
uint32_t featStride = params.get<uint32_t>("feat_stride", 16);
uint32_t baseSize = params.get<uint32_t>("base_size", 16);
featStride = params.get<uint32_t>("feat_stride", 16);
baseSize = params.get<uint32_t>("base_size", 16);
// uint32_t minSize = params.get<uint32_t>("min_size", 16);
uint32_t keepTopBeforeNMS = params.get<uint32_t>("pre_nms_topn", 6000);
keepTopBeforeNMS = params.get<uint32_t>("pre_nms_topn", 6000);
keepTopAfterNMS = params.get<uint32_t>("post_nms_topn", 300);
float nmsThreshold = params.get<float>("nms_thresh", 0.7);
DictValue ratios = params.get("ratio");
DictValue scales = params.get("scale");
nmsThreshold = params.get<float>("nms_thresh", 0.7);
ratios = params.get("ratio");
scales = params.get("scale");
{
LayerParams lp;
@@ -83,6 +84,12 @@ public:
}
}
virtual bool supportBackend(int backendId) CV_OVERRIDE
{
return backendId == DNN_BACKEND_OPENCV ||
backendId == DNN_BACKEND_INFERENCE_ENGINE && preferableTarget != DNN_TARGET_MYRIAD;
}
bool getMemoryShapes(const std::vector<MatShape> &inputs,
const int requiredOutputs,
std::vector<MatShape> &outputs,
@@ -312,6 +319,38 @@ public:
outputs[i].rowRange(numDets, keepTopAfterNMS).setTo(0);
}
virtual Ptr<BackendNode> initInfEngine(const std::vector<Ptr<BackendWrapper> >&) CV_OVERRIDE
{
#ifdef HAVE_INF_ENGINE
InferenceEngine::LayerParams lp;
lp.name = name;
lp.type = "Proposal";
lp.precision = InferenceEngine::Precision::FP32;
std::shared_ptr<InferenceEngine::CNNLayer> ieLayer(new InferenceEngine::CNNLayer(lp));
ieLayer->params["base_size"] = format("%d", baseSize);
ieLayer->params["feat_stride"] = format("%d", featStride);
ieLayer->params["min_size"] = "16";
ieLayer->params["nms_thresh"] = format("%f", nmsThreshold);
ieLayer->params["post_nms_topn"] = format("%d", keepTopAfterNMS);
ieLayer->params["pre_nms_topn"] = format("%d", keepTopBeforeNMS);
if (ratios.size())
{
ieLayer->params["ratio"] = format("%f", ratios.get<float>(0));
for (int i = 1; i < ratios.size(); ++i)
ieLayer->params["ratio"] += format(",%f", ratios.get<float>(i));
}
if (scales.size())
{
ieLayer->params["scale"] = format("%f", scales.get<float>(0));
for (int i = 1; i < scales.size(); ++i)
ieLayer->params["scale"] += format(",%f", scales.get<float>(i));
}
return Ptr<BackendNode>(new InfEngineBackendNode(ieLayer));
#endif // HAVE_INF_ENGINE
return Ptr<BackendNode>();
}
private:
// A first half of channels are background scores. We need only a second one.
static Mat getObjectScores(const Mat& m)
@@ -342,8 +381,10 @@ private:
Ptr<PermuteLayer> deltasPermute;
Ptr<PermuteLayer> scoresPermute;
uint32_t keepTopAfterNMS;
uint32_t keepTopBeforeNMS, keepTopAfterNMS, featStride, baseSize;
Mat fakeImageBlob;
float nmsThreshold;
DictValue ratios, scales;
#ifdef HAVE_OPENCL
UMat umat_fakeImageBlob;
#endif
+2 -1
View File
@@ -183,8 +183,9 @@ bool OCL4DNNPool<Dtype>::Forward(const UMat& bottom,
ocl::Kernel oclk_sto_pool_forward(
kname.c_str(),
ocl::dnn::ocl4dnn_pooling_oclsrc,
format("-D KERNEL_STO_POOL=1 -D KERNEL_W=%d -D KERNEL_H=%d"
format(" -D Dtype=%s -D KERNEL_STO_POOL=1 -D KERNEL_W=%d -D KERNEL_H=%d"
" -D STRIDE_W=%d -D STRIDE_H=%d",
(use_half) ? "half" : "float",
kernel_w_, kernel_h_,
stride_w_, stride_h_
));
+2 -2
View File
@@ -104,7 +104,7 @@ __kernel void
#elif defined KERNEL_AVE_POOL
__kernel void TEMPLATE(ave_pool_forward, Dtype)(
const int nthreads, __global const Dtype* const bottom_data,
const int nthreads, __global const Dtype* bottom_data,
const int channels, const int height, const int width,
const int pooled_height, const int pooled_width,
__global Dtype* top_data)
@@ -150,7 +150,7 @@ __kernel void TEMPLATE(ave_pool_forward, Dtype)(
#elif defined KERNEL_STO_POOL
__kernel void TEMPLATE(sto_pool_forward_test,Dtype)(
const int nthreads, __global const Dtype* const bottom_data,
const int nthreads, __global const Dtype* bottom_data,
const int channels, const int height, const int width,
const int pooled_height, const int pooled_width,
__global Dtype* top_data)
+10
View File
@@ -938,6 +938,16 @@ struct TorchImporter
layerParams.set("end", DictValue::arrayInt<int*>(&ends[0], 4));
curModule->modules.push_back(newModule);
}
else if (nnName == "SpatialUpSamplingNearest")
{
readTorchTable(scalarParams, tensorParams);
CV_Assert(scalarParams.has("scale_factor"));
int scale_factor = scalarParams.get<int>("scale_factor");
newModule->apiType = "Resize";
layerParams.set("interpolation", "nearest");
layerParams.set("zoom_factor", scale_factor);
curModule->modules.push_back(newModule);
}
else
{
// Importer does not know how to map Torch's layer type to an OpenCV's one.