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
2018-09-28 12:48:51 +03:00
parent 29e88e50ff 4b895a4d1f
commit a8b0db4e5d
77 changed files with 586 additions and 443 deletions
Download Patch File Download Diff File Expand all files Collapse all files
modules/dnn/CMakeLists.txt
-1
View File
@@ -36,7 +36,6 @@ else()
-Wunused-parameter -Wunused-local-typedefs -Wsign-compare -Wsign-promo
-Wundef -Wtautological-undefined-compare -Wignored-qualifiers -Wextra
-Wunused-function -Wunused-const-variable -Wdeprecated-declarations
-Werror=non-virtual-dtor
)
endif()
modules/dnn/include/opencv2/dnn/dnn.hpp
+5
View File
@@ -528,6 +528,11 @@ CV__DNN_INLINE_NS_BEGIN
/** @brief Returns indexes of layers with unconnected outputs.
*/
CV_WRAP std::vector<int> getUnconnectedOutLayers() const;
/** @brief Returns names of layers with unconnected outputs.
*/
CV_WRAP std::vector<String> getUnconnectedOutLayersNames() const;
/** @brief Returns input and output shapes for all layers in loaded model;
* preliminary inferencing isn't necessary.
* @param netInputShapes shapes for all input blobs in net input layer.
modules/dnn/src/dnn.cpp
+27 -5
View File
@@ -1078,12 +1078,22 @@ struct Net::Impl
}
#else
{
if (!DNN_OPENCL_ALLOW_ALL_DEVICES
&& !(ocl::Device::getDefault().isIntel() && ocl::Device::getDefault().type() == ocl::Device::TYPE_GPU) // Current implementation is only valid for Intel GPU (#11494)
)
if (!DNN_OPENCL_ALLOW_ALL_DEVICES)
{
CV_LOG_WARNING(NULL, "DNN: OpenCL target is not supported with current OpenCL device (tested with Intel GPUs only), switching to CPU.");
preferableTarget = DNN_TARGET_CPU;
// Current implementation is only valid for GPU (#11494)
if (ocl::Device::getDefault().type() != ocl::Device::TYPE_GPU)
{
CV_LOG_WARNING(NULL, "DNN: OpenCL target is not supported with current OpenCL device (tested with GPUs only), switching to CPU.");
preferableTarget = DNN_TARGET_CPU;
}
else if (preferableTarget == DNN_TARGET_OPENCL_FP16 && !ocl::Device::getDefault().isIntel())
{
CV_LOG_WARNING(NULL,
"DNN: OpenCL target with fp16 precision is not supported "
"with current OpenCL device (tested with Intel GPUs only), "
"switching to OpenCL with fp32 precision.");
preferableTarget = DNN_TARGET_OPENCL;
}
}
}
#endif
@@ -2789,6 +2799,18 @@ std::vector<int> Net::getUnconnectedOutLayers() const
return layersIds;
}
std::vector<String> Net::getUnconnectedOutLayersNames() const
{
std::vector<int> ids = getUnconnectedOutLayers();
const size_t n = ids.size();
std::vector<String> names(n);
for (size_t i = 0; i < n; ++i)
{
names[i] = impl->layers[ids[i]].name;
}
return names;
}
void Net::getLayersShapes(const ShapesVec& netInputShapes,
std::vector<int>& layersIds,
std::vector<ShapesVec>& inLayersShapes,
modules/dnn/src/layers/batch_norm_layer.cpp
+1 -2
View File
@@ -230,8 +230,7 @@ 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()),
CV_OCL_RUN(IS_DNN_OPENCL_TARGET(preferableTarget),
forward_ocl(inputs_arr, outputs_arr, internals_arr))
if (inputs_arr.depth() == CV_16S)
modules/dnn/src/layers/blank_layer.cpp
+1 -8
View File
@@ -95,16 +95,9 @@ 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()),
CV_OCL_RUN(IS_DNN_OPENCL_TARGET(preferableTarget),
forward_ocl(inputs_arr, outputs_arr, internals_arr))
if (inputs_arr.depth() == CV_16S)
{
forward_fallback(inputs_arr, outputs_arr, internals_arr);
return;
}
std::vector<Mat> inputs, outputs;
inputs_arr.getMatVector(inputs);
outputs_arr.getMatVector(outputs);
modules/dnn/src/layers/concat_layer.cpp
+1 -8
View File
@@ -237,16 +237,9 @@ 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()),
CV_OCL_RUN(IS_DNN_OPENCL_TARGET(preferableTarget),
forward_ocl(inputs_arr, outputs_arr, internals_arr))
if (inputs_arr.depth() == CV_16S)
{
forward_fallback(inputs_arr, outputs_arr, internals_arr);
return;
}
std::vector<Mat> inputs, outputs;
inputs_arr.getMatVector(inputs);
outputs_arr.getMatVector(outputs);
modules/dnn/src/layers/convolution_layer.cpp
+1 -2
View File
@@ -1529,8 +1529,7 @@ 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()),
CV_OCL_RUN(IS_DNN_OPENCL_TARGET(preferableTarget),
forward_ocl(inputs_arr, outputs_arr, internals_arr));
if (inputs_arr.depth() == CV_16S)
modules/dnn/src/layers/crop_layer.cpp
-6
View File
@@ -137,12 +137,6 @@ public:
CV_TRACE_FUNCTION();
CV_TRACE_ARG_VALUE(name, "name", name.c_str());
if (inputs_arr.depth() == CV_16S)
{
forward_fallback(inputs_arr, outputs_arr, internals_arr);
return;
}
std::vector<Mat> inputs, outputs;
inputs_arr.getMatVector(inputs);
outputs_arr.getMatVector(outputs);
modules/dnn/src/layers/detection_output_layer.cpp
+1 -2
View File
@@ -415,8 +415,7 @@ public:
if (_bboxesNormalized)
{
CV_OCL_RUN(IS_DNN_OPENCL_TARGET(preferableTarget) &&
OCL_PERFORMANCE_CHECK(ocl::Device::getDefault().isIntel()),
CV_OCL_RUN(IS_DNN_OPENCL_TARGET(preferableTarget),
forward_ocl(inputs_arr, outputs_arr, internals_arr))
}
if (inputs_arr.depth() == CV_16S)
modules/dnn/src/layers/eltwise_layer.cpp
+1 -2
View File
@@ -354,8 +354,7 @@ 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()),
CV_OCL_RUN(IS_DNN_OPENCL_TARGET(preferableTarget),
forward_ocl(inputs_arr, outputs_arr, internals_arr))
if (inputs_arr.depth() == CV_16S)
modules/dnn/src/layers/flatten_layer.cpp
+1 -8
View File
@@ -135,16 +135,9 @@ public:
CV_TRACE_ARG_VALUE(name, "name", name.c_str());
CV_OCL_RUN(IS_DNN_OPENCL_TARGET(preferableTarget) &&
outputs_arr.isUMatVector() &&
OCL_PERFORMANCE_CHECK(ocl::Device::getDefault().isIntel()),
outputs_arr.isUMatVector(),
forward_ocl(inputs_arr, outputs_arr, internals_arr))
if (inputs_arr.depth() == CV_16S)
{
forward_fallback(inputs_arr, outputs_arr, internals_arr);
return;
}
std::vector<Mat> inputs, outputs;
inputs_arr.getMatVector(inputs);
outputs_arr.getMatVector(outputs);
modules/dnn/src/layers/fully_connected_layer.cpp
+1 -2
View File
@@ -389,8 +389,7 @@ 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()),
CV_OCL_RUN(IS_DNN_OPENCL_TARGET(preferableTarget),
forward_ocl(inputs_arr, outputs_arr, internals_arr))
if (inputs_arr.depth() == CV_16S)
modules/dnn/src/layers/lrn_layer.cpp
+1 -2
View File
@@ -148,8 +148,7 @@ public:
CV_Assert(inputs_arr.total() == outputs_arr.total());
CV_OCL_RUN(IS_DNN_OPENCL_TARGET(preferableTarget) &&
OCL_PERFORMANCE_CHECK(ocl::Device::getDefault().isIntel()),
CV_OCL_RUN(IS_DNN_OPENCL_TARGET(preferableTarget),
forward_ocl(inputs_arr, outputs_arr, internals_arr))
if (inputs_arr.depth() == CV_16S)
modules/dnn/src/layers/normalize_bbox_layer.cpp
+1 -2
View File
@@ -184,8 +184,7 @@ 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()),
CV_OCL_RUN(IS_DNN_OPENCL_TARGET(preferableTarget),
forward_ocl(inputs_arr, outputs_arr, internals_arr))
if (inputs_arr.depth() == CV_16S)
modules/dnn/src/layers/padding_layer.cpp
+9 -7
View File
@@ -99,19 +99,21 @@ public:
CV_TRACE_FUNCTION();
CV_TRACE_ARG_VALUE(name, "name", name.c_str());
if (inputs_arr.depth() == CV_16S)
{
forward_fallback(inputs_arr, outputs_arr, internals_arr);
return;
}
std::vector<Mat> inputs, outputs;
inputs_arr.getMatVector(inputs);
outputs_arr.getMatVector(outputs);
if (paddingType == "constant")
{
outputs[0].setTo(paddingValue);
if (inputs_arr.depth() == CV_16S)
{
std::vector<float> paddingValue_fp32(1, paddingValue);
std::vector<int16_t> paddingValue_fp16(1);
convertFp16(paddingValue_fp32, paddingValue_fp16);
outputs[0].setTo(paddingValue_fp16[0]);
}
else
outputs[0].setTo(paddingValue);
inputs[0].copyTo(outputs[0](dstRanges));
}
else if (paddingType == "reflect")
modules/dnn/src/layers/permute_layer.cpp
+1 -2
View File
@@ -304,8 +304,7 @@ 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()),
CV_OCL_RUN(IS_DNN_OPENCL_TARGET(preferableTarget),
forward_ocl(inputs_arr, outputs_arr, internals_arr))
if (inputs_arr.depth() == CV_16S)
modules/dnn/src/layers/prior_box_layer.cpp
+1 -2
View File
@@ -402,8 +402,7 @@ 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()),
CV_OCL_RUN(IS_DNN_OPENCL_TARGET(preferableTarget),
forward_ocl(inputs_arr, outputs_arr, internals_arr))
if (inputs_arr.depth() == CV_16S)
modules/dnn/src/layers/region_layer.cpp
+1 -2
View File
@@ -196,8 +196,7 @@ 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()),
CV_OCL_RUN(IS_DNN_OPENCL_TARGET(preferableTarget),
forward_ocl(inputs_arr, outputs_arr, internals_arr))
if (inputs_arr.depth() == CV_16S)
modules/dnn/src/layers/reorg_layer.cpp
+1 -2
View File
@@ -160,8 +160,7 @@ 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()),
CV_OCL_RUN(IS_DNN_OPENCL_TARGET(preferableTarget),
forward_ocl(inputs_arr, outputs_arr, internals_arr))
if (inputs_arr.depth() == CV_16S)
modules/dnn/src/layers/reshape_layer.cpp
+1 -8
View File
@@ -233,16 +233,9 @@ 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()),
CV_OCL_RUN(IS_DNN_OPENCL_TARGET(preferableTarget),
forward_ocl(inputs_arr, outputs_arr, internals_arr))
if (inputs_arr.depth() == CV_16S)
{
forward_fallback(inputs_arr, outputs_arr, internals_arr);
return;
}
std::vector<Mat> inputs, outputs;
inputs_arr.getMatVector(inputs);
outputs_arr.getMatVector(outputs);
modules/dnn/src/layers/shuffle_channel_layer.cpp
+1 -2
View File
@@ -92,8 +92,7 @@ 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()),
CV_OCL_RUN(IS_DNN_OPENCL_TARGET(preferableTarget),
forward_ocl(inputs_arr, outputs_arr, internals_arr))
if (inputs_arr.depth() == CV_16S)
modules/dnn/src/layers/slice_layer.cpp
+1 -8
View File
@@ -239,16 +239,9 @@ 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()),
CV_OCL_RUN(IS_DNN_OPENCL_TARGET(preferableTarget),
forward_ocl(inputs_arr, outputs_arr, internals_arr))
if (inputs_arr.depth() == CV_16S)
{
forward_fallback(inputs_arr, outputs_arr, internals_arr);
return;
}
std::vector<Mat> inputs, outputs;
inputs_arr.getMatVector(inputs);
outputs_arr.getMatVector(outputs);
modules/dnn/src/layers/softmax_layer.cpp
+1 -2
View File
@@ -187,8 +187,7 @@ 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()),
CV_OCL_RUN(IS_DNN_OPENCL_TARGET(preferableTarget),
forward_ocl(inputs_arr, outputs_arr, internals_arr))
if (inputs_arr.depth() == CV_16S)
modules/dnn/src/layers/split_layer.cpp
-6
View File
@@ -83,12 +83,6 @@ public:
CV_TRACE_FUNCTION();
CV_TRACE_ARG_VALUE(name, "name", name.c_str());
if (inputs_arr.depth() == CV_16S)
{
forward_fallback(inputs_arr, outputs_arr, internals_arr);
return;
}
std::vector<Mat> inputs, outputs;
inputs_arr.getMatVector(inputs);
outputs_arr.getMatVector(outputs);
modules/dnn/src/ocl4dnn/src/ocl4dnn_conv_spatial.cpp
+183 -69
View File
@@ -60,6 +60,8 @@
#if defined WIN32 || defined _WIN32
#include <windows.h>
#include <direct.h>
#undef min
#undef max
#endif
namespace cv { namespace dnn { namespace ocl4dnn {
@@ -68,6 +70,30 @@ typedef std::map<std::string, std::string> kernel_hash_t;
static kernel_hash_t kernelConfigMap;
static bool defaultConfigLoaded = false;
static bool enableWorkaroundIDLF()
{
static bool param = utils::getConfigurationParameterSizeT("OPENCV_OCL4DNN_WORKAROUND_IDLF", true);
return param;
}
static bool dumpFailedResult()
{
static bool param = utils::getConfigurationParameterSizeT("OPENCV_OCL4DNN_DUMP_FAILED_RESULT", false);
return param;
}
static size_t testAllKernels()
{
static size_t param = utils::getConfigurationParameterSizeT("OPENCV_OCL4DNN_TEST_ALL_KERNELS", 0);
return param;
}
static bool raiseOnCheckError()
{
static bool param = utils::getConfigurationParameterBool("OPENCV_OCL4DNN_TUNING_RAISE_CHECK_ERROR", false);
return param;
}
static std::string sanitize(const std::string& s)
{
std::string s_ = s;
@@ -1221,9 +1247,6 @@ bool OCL4DNNConvSpatial<float>::verifyResult(const UMat &bottom,
kernelConfig* config,
UMat &verifyTop)
{
uint32_t verificationFail = 0;
if (config->verified)
return true;
else if (config->tested)
@@ -1236,6 +1259,8 @@ bool OCL4DNNConvSpatial<float>::verifyResult(const UMat &bottom,
convolve(bottom, top, weight, bias, numImages, config);
tuned_ = saved_tuned;
config->tested = true;
UMat new_top, new_verify_top;
Mat mat_top, mat_verify_top;
if (use_half_)
@@ -1254,41 +1279,88 @@ bool OCL4DNNConvSpatial<float>::verifyResult(const UMat &bottom,
const float* data = mat_top.ptr<float>();
const float* verify_data = mat_verify_top.ptr<float>();
for (int32_t n = 0; n < num_; ++n) {
for (int32_t g = 0; g < group_; ++g) {
int32_t output_image_offset = n * top_dim_ + output_w_ * output_h_ * M_ * g;
for (int out_ch = 0; out_ch < M_ && !verificationFail; out_ch++)
for (int h = 0; h < output_h_ && !verificationFail; h++)
for (int w = 0; w < output_w_; w++) {
size_t offset = output_image_offset + out_ch * output_w_ * output_h_ + h * output_w_ + w;
int error_slice_offset = 0;
int error_slice = 0;
float relative_eps = use_half_ ? 0.1f : 0.01f;
float error_factor = fabs(data[offset] - verify_data[offset]);
if (use_half_ && error_factor > 0.1 * fabs(verify_data[offset]) &&
error_factor > 0.04 && !(fabs(verify_data[offset]) < 1.e-3 && error_factor < 1.e-4))
{
CV_LOG_ERROR(NULL, "test verification failed @ image " << n << " group " << g
<< " out_ch " << out_ch << " h " << h << " w " << w
<< " got " << data[offset] << " expected " << verify_data[offset]);
verificationFail = 1;
goto out;
size_t errors = 0;
double rel_err = norm(mat_top.reshape(1, 1), mat_verify_top.reshape(1, 1), NORM_L1 | NORM_RELATIVE);
if (rel_err >= relative_eps)
{
for (int32_t n = 0; n < num_; ++n) {
for (int32_t g = 0; g < group_; ++g) {
int32_t output_image_offset = n * top_dim_ + output_w_ * output_h_ * M_ * g;
for (int out_ch = 0; out_ch < M_; out_ch++)
for (int h = 0; h < output_h_; h++)
for (int w = 0; w < output_w_; w++) {
size_t offset = output_image_offset + out_ch * output_w_ * output_h_ + h * output_w_ + w;
bool has_error = !(data[offset] == data[offset]); // is NaN
if (!has_error)
{
float error_factor = std::fabs(data[offset] - verify_data[offset]);
float base_value_abs = std::max(1e-3f, std::fabs(verify_data[offset]));
has_error = error_factor > relative_eps * base_value_abs;
}
if (has_error)
{
if (errors == 0)
{
error_slice = (int)(offset / (output_w_ * output_h_));
error_slice_offset = (int)(offset % (output_w_ * output_h_));
CV_LOG_ERROR(NULL, "Kernel: " << config->kernelName);
}
if (errors < 10)
CV_LOG_ERROR(NULL, "test verification failed @ image " << n << " group " << g
<< " out_ch " << out_ch << " h " << h << " w " << w
<< " (offset: " << offset << ")"
<< " got " << data[offset] << " expected " << verify_data[offset]);
errors++;
}
}
else if (!use_half_ && error_factor > 0.1 * fabs(verify_data[offset]) &&
!(fabs(verify_data[offset]) < 1.e-3 && error_factor < 1.e-4))
{
CV_LOG_ERROR(NULL, "test verification failed @ image " << n << " group " << g
<< " out_ch " << out_ch << " h " << h << " w " << w
<< " got " << data[offset] << " expected " << verify_data[offset]);
verificationFail = 1;
goto out;
}
}
}
}
}
out:
if (verificationFail == 1)
if (errors)
{
if (dumpFailedResult())
{
try
{
int n_outputs = (int)(mat_top.size[0]*mat_top.size[1]);
int slice_size = (int)(mat_top.total() / n_outputs);
Rect roi(0, 0, slice_size, n_outputs);
roi.width = std::min(roi.width, 32);
roi.height = std::min(roi.height, 16);
roi.x = std::max(0, std::min(slice_size - roi.width, error_slice_offset - roi.width/2));
roi.y = std::max(0, std::min(n_outputs - roi.height, error_slice - roi.height/2));
std::cout << "roi = " << roi << " errors=" << errors << std::endl;
std::cout << "mat_top = " << shape(mat_top) << std::endl
<< mat_top.reshape(1, 1).reshape(1, n_outputs)(roi) << std::endl;
std::cout << "verify_top = " << shape(mat_verify_top) << std::endl
<< mat_verify_top.reshape(1, 1).reshape(1, n_outputs)(roi) << std::endl;
}
catch (const std::exception& e)
{
CV_LOG_ERROR(NULL, "Results dump failed: " << e.what());
}
catch (...)
{
CV_LOG_ERROR(NULL, "Results dump failed")
}
}
if (raiseOnCheckError())
CV_Error_(Error::StsError, ("ocl4dnn tuning verification failed: %s (errors %lld)", config->kernelName.c_str(), (long long int)errors));
return false;
}
else
{
config->verified = true;
return true;
}
}
template<typename Dtype>
@@ -1408,6 +1480,17 @@ bool OCL4DNNConvSpatial<float>::createIDLFKernel(int32_t blockWidth,
setupKernel();
if (enableWorkaroundIDLF() && ocl::Device::getDefault().intelSubgroupsSupport())
{
// Issues are observed with these kernels: 3x1 (covered by tests), 2x1, 4x1, 5x1, 3x2
// kernels 1x3, 3x3, 2x3 are good
if (pad_h_ != 0 && kernel_w_ <= simd_size && kernel_h_ <= 2)
{
CV_LOG_INFO(NULL, "DNN(workaround): skip IDLF kernel: " << kernel_name_);
return false;
}
}
ocl::Program program = compileKernel();
if (program.ptr())
{
@@ -1623,13 +1706,38 @@ void OCL4DNNConvSpatial<float>::useFirstAvailable(const UMat &bottom,
generateTunerItems(tunerItems);
tunerItems.push_back(makePtr<tunerParam>(KERNEL_TYPE_BASIC, 1, 1, 1));
for (int i = 0; i < tunerItems.size(); i++) {
for (int i = 0; i < tunerItems.size(); i++)
{
if (createConvolutionKernel(tunerItems[i]->kernelType,
tunerItems[i]->blockWidth,
tunerItems[i]->blockHeight,
tunerItems[i]->blockDepth)) {
tunerItems[i]->blockDepth))
{
int kernelIdx = kernelQueue.size() - 1;
if (verifyResult(bottom, top, weight, bias, numImages, kernelQueue[kernelIdx], verifyTop)) {
kernelConfig* config = kernelQueue[kernelIdx].get();
bool failed = false;
const size_t testCount = testAllKernels();
for(int t = 0; t < testCount; t++)
{
try
{
config->tested = false;
config->verified = false;
if (!verifyResult(bottom, top, weight, bias, numImages, config, verifyTop))
{
CV_LOG_ERROR(NULL, "Failed on test iteration: " << t);
failed = true;
break;
}
}
catch (...)
{
CV_LOG_ERROR(NULL, "Failed on test iteration: " << t);
throw;
}
}
if (!failed && verifyResult(bottom, top, weight, bias, numImages, config, verifyTop))
{
bestKernelConfig = kernelQueue[kernelIdx];
if (bestKernelConfig->kernelType != KERNEL_TYPE_INTEL_IDLF &&
bestKernelConfig->kernelType != KERNEL_TYPE_GEMM_LIKE)
@@ -1685,42 +1793,50 @@ void OCL4DNNConvSpatial<float>::setupConvolution(const UMat &bottom,
tunerItems[i]->blockHeight,
tunerItems[i]->blockDepth);
for (int32_t x = 0; x < kernelQueue.size(); x++) {
kernelQueue[x]->executionTime = timedConvolve(bottom, top, weight, bias, numImages,
kernelQueue[x]);
#ifdef TEST_ALL_KERNELS
if (kernelQueue[x]->tested == false) {
bool verified = verifyResult(bottom, top, weight, bias, numImages, kernelQueue[x], verifyTop);
if (verified == false) {
CV_LOG_ERROR(NULL, "Kernel " << kernelQueue[x]->kernelName << " failed verification");
CV_LOG_ERROR(NULL, "kernelQueue[x]->workItem_output[0]: "
<< kernelQueue[x]->workItem_output[0] << " "
<< "kernelQueue[x]->workItem_output[1]: "
<< kernelQueue[x]->workItem_output[1] << " "
<< "kernelQueue[x]->workItem_output[2]: "
<< kernelQueue[x]->workItem_output[2] << " "
<< "kernelQueue[x]->kernelType: "
<< kernelQueue[x]->kernelType << " "
<< "kernelQueue[x]->global_work_size[0]: "
<< kernelQueue[x]->global_work_size[0] << " "
<< "kernelQueue[x]->global_work_size[1]: "
<< kernelQueue[x]->global_work_size[1] << " "
<< "kernelQueue[x]->global_work_size[2]: "
<< kernelQueue[x]->global_work_size[2] << " "
<< "kernelQueue[x]->local_work_size[0]: "
<< kernelQueue[x]->local_work_size[0] << " "
<< "kernelQueue[x]->local_work_size[1]: "
<< kernelQueue[x]->local_work_size[1] << " "
<< "kernelQueue[x]->local_work_size[2]: "
<< kernelQueue[x]->local_work_size[2] << " "
<< kernelQueue[x]->swizzle_weights << " "
<< kernelQueue[x]->use_null_local);
} else {
CV_LOG_INFO(NULL, "Kernel " << kernelQueue[x]->kernelName << " pass verification");
const size_t testCount = testAllKernels();
for (int32_t x = 0; x < kernelQueue.size(); x++)
{
kernelConfig* config = kernelQueue[x];
config->executionTime = timedConvolve(bottom, top, weight, bias, numImages, config);
for(int t = 0; t < testCount; t++)
{
try
{
config->tested = false;
config->verified = false;
bool verified = verifyResult(bottom, top, weight, bias, numImages, config, verifyTop);
if (verified == false)
{
CV_LOG_ERROR(NULL, "Kernel " << config->kernelName << " failed verification");
CV_LOG_ERROR(NULL, "workItem="
<< config->workItem_output[0] << ","
<< config->workItem_output[1] << ","
<< config->workItem_output[2] << " "
<< "kernelType: " << config->kernelType << " "
<< "global_work_size="
<< config->global_work_size[0] << ","
<< config->global_work_size[1] << ","
<< config->global_work_size[2] << " "
<< "local_work_size="
<< config->local_work_size[0] << ","
<< config->local_work_size[1] << ","
<< config->local_work_size[2] << " "
<< config->swizzle_weights << " "
<< config->use_null_local);
}
else
{
CV_LOG_VERBOSE(NULL, "Kernel " << config->kernelName << " pass verification");
}
}
catch (...)
{
CV_LOG_ERROR(NULL, "Failed on test iteration: " << t);
throw;
}
}
#endif
}
int32_t failures = 0;
bool verification = false;
if (kernelQueue.size()) {
@@ -1739,12 +1855,10 @@ void OCL4DNNConvSpatial<float>::setupConvolution(const UMat &bottom,
// Test fastest kernel
bool verified = verifyResult(bottom, top, weight, bias, numImages, kernelQueue[fastestKernel], verifyTop);
if (verified == true) {
kernelQueue[fastestKernel]->verified = true;
kernel_index_ = fastestKernel;
verification = true;
break;
} else {
kernelQueue[fastestKernel]->tested = true;
CV_LOG_ERROR(NULL, "Kernel " << kernelQueue[fastestKernel]->kernelName <<
" failed verification");
failures++;
modules/dnn/src/ocl4dnn/src/ocl4dnn_lrn.cpp
-3
View File
@@ -69,9 +69,6 @@ bool OCL4DNNLRN<Dtype>::Forward(const UMat& bottom, UMat& top)
{
bool ret = true;
if (!ocl::Device::getDefault().intelSubgroupsSupport())
return false;
switch (lrn_type_)
{
case LRNParameter_NormRegion_ACROSS_CHANNELS:
modules/dnn/src/onnx/onnx_importer.cpp
+1 -1
View File
@@ -213,7 +213,7 @@ LayerParams ONNXImporter::getLayerParams(const opencv_onnx::NodeProto& node_prot
else if (attribute_proto.floats_size() > 0)
{
lp.set(attribute_name, DictValue::arrayReal(
(float*)attribute_proto.mutable_floats(), attribute_proto.floats_size()));
attribute_proto.floats().data(), attribute_proto.floats_size()));
}
else if (attribute_proto.ints_size() > 0)
{
modules/dnn/src/opencl/prior_box.cl
+1 -1
View File
@@ -114,6 +114,6 @@ __kernel void clip(const int nthreads,
for (int index = get_global_id(0); index < nthreads; index += get_global_size(0))
{
Dtype4 vec = vload4(index, dst);
vstore4(clamp(vec, 0, 1), index, dst);
vstore4(clamp(vec, 0.0f, 1.0f), index, dst);
}
}
modules/dnn/src/tensorflow/tf_graph_simplifier.cpp
+2
View File
@@ -20,6 +20,8 @@ using ::google::protobuf::MapPair;
class Subgraph // Interface to match and replace TensorFlow subgraphs.
{
public:
virtual ~Subgraph() {}
// Add a node to be matched in the origin graph. Specify ids of nodes that
// are expected to be inputs. Returns id of a newly added node.
// TODO: Replace inputs to std::vector<int> in C++11
modules/dnn/test/test_common.hpp
+2
View File
@@ -276,6 +276,8 @@ static testing::internal::ParamGenerator<tuple<Backend, Target> > dnnBackendsAnd
targets.push_back(make_tuple(DNN_BACKEND_OPENCV, DNN_TARGET_OPENCL_FP16));
}
#endif
if (targets.empty()) // validate at least CPU mode
targets.push_back(make_tuple(DNN_BACKEND_OPENCV, DNN_TARGET_CPU));
return testing::ValuesIn(targets);
}
modules/dnn/test/test_halide_layers.cpp
-8
View File
@@ -99,14 +99,6 @@ TEST_P(Convolution, Accuracy)
#endif
bool skipCheck = false;
if (cvtest::skipUnstableTests && backendId == DNN_BACKEND_OPENCV &&
(targetId == DNN_TARGET_OPENCL || targetId == DNN_TARGET_OPENCL_FP16) &&
(
(kernel == Size(3, 1) && stride == Size(1, 1) && pad == Size(0, 1)) ||
(stride.area() > 1 && !(pad.width == 0 && pad.height == 0))
)
)
skipCheck = true;
int sz[] = {outChannels, inChannels / group, kernel.height, kernel.width};
Mat weights(4, &sz[0], CV_32F);
modules/dnn/test/test_onnx_importer.cpp
+2 -2
View File
@@ -295,7 +295,7 @@ TEST_P(Test_ONNX_nets, TinyYolov2)
TEST_P(Test_ONNX_nets, CNN_MNIST)
{
// output range: [-1952; 6574]
const double l1 = (target == DNN_TARGET_OPENCL_FP16 || target == DNN_TARGET_MYRIAD) ? 3.82 : 4.3e-4;
const double l1 = (target == DNN_TARGET_OPENCL_FP16 || target == DNN_TARGET_MYRIAD) ? 3.82 : 4.4e-4;
const double lInf = (target == DNN_TARGET_OPENCL_FP16 || target == DNN_TARGET_MYRIAD) ? 13.5 : 2e-3;
testONNXModels("cnn_mnist", pb, l1, lInf);
@@ -341,7 +341,7 @@ TEST_P(Test_ONNX_nets, Inception_v2)
TEST_P(Test_ONNX_nets, DenseNet121)
{
// output range: [-87; 138]
const double l1 = (target == DNN_TARGET_OPENCL_FP16 || target == DNN_TARGET_MYRIAD) ? 0.12 : 1.88e-5;
const double l1 = (target == DNN_TARGET_OPENCL_FP16 || target == DNN_TARGET_MYRIAD) ? 0.12 : 2.2e-5;
const double lInf = (target == DNN_TARGET_OPENCL_FP16 || target == DNN_TARGET_MYRIAD) ? 0.74 : 1.23e-4;
testONNXModels("densenet121", pb, l1, lInf);
}