From 2f9072efdcf0d72023be31f174bd1bf9abd03757 Mon Sep 17 00:00:00 2001 From: Nathan Godwin Date: Mon, 21 Sep 2020 21:56:28 -0500 Subject: [PATCH 01/10] Fixed assertions on ippe solver --- modules/calib3d/src/ippe.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/modules/calib3d/src/ippe.cpp b/modules/calib3d/src/ippe.cpp index 3f1e6a7add..ec6089596c 100644 --- a/modules/calib3d/src/ippe.cpp +++ b/modules/calib3d/src/ippe.cpp @@ -77,18 +77,18 @@ void PoseSolver::solveGeneric(InputArray _objectPoints, InputArray _normalizedIn OutputArray _Ma, OutputArray _Mb) { //argument checking: - size_t n = static_cast(_objectPoints.rows() * _objectPoints.cols()); //number of points + size_t n = static_cast(_normalizedInputPoints.rows()) * static_cast(_normalizedInputPoints.cols()); //number of points int objType = _objectPoints.type(); int type_input = _normalizedInputPoints.type(); CV_CheckType(objType, objType == CV_32FC3 || objType == CV_64FC3, "Type of _objectPoints must be CV_32FC3 or CV_64FC3" ); CV_CheckType(type_input, type_input == CV_32FC2 || type_input == CV_64FC2, - "Type of _normalizedInputPoints must be CV_32FC3 or CV_64FC3" ); + "Type of _normalizedInputPoints must be CV_32FC2 or CV_64FC2" ); CV_Assert(_objectPoints.rows() == 1 || _objectPoints.cols() == 1); CV_Assert(_objectPoints.rows() >= 4 || _objectPoints.cols() >= 4); CV_Assert(_normalizedInputPoints.rows() == 1 || _normalizedInputPoints.cols() == 1); - CV_Assert(static_cast(_objectPoints.rows() * _objectPoints.cols()) == n); + CV_Assert(static_cast(_objectPoints.rows()) * static_cast(_objectPoints.cols()) == n); Mat normalizedInputPoints; if (type_input == CV_32FC2) @@ -101,7 +101,7 @@ void PoseSolver::solveGeneric(InputArray _objectPoints, InputArray _normalizedIn } Mat objectInputPoints; - if (type_input == CV_32FC3) + if (objType == CV_32FC3) { _objectPoints.getMat().convertTo(objectInputPoints, CV_64F); } From 74c8ccb45b2792499683c84b80422866d8b7054f Mon Sep 17 00:00:00 2001 From: Tomoaki Teshima Date: Wed, 23 Sep 2020 21:38:12 +0900 Subject: [PATCH 02/10] fix build error of kernel on Mali --- modules/dnn/src/opencl/ocl4dnn_lrn.cl | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/modules/dnn/src/opencl/ocl4dnn_lrn.cl b/modules/dnn/src/opencl/ocl4dnn_lrn.cl index 36d9d2ae04..31c9f49451 100644 --- a/modules/dnn/src/opencl/ocl4dnn_lrn.cl +++ b/modules/dnn/src/opencl/ocl4dnn_lrn.cl @@ -83,7 +83,7 @@ __kernel void TEMPLATE(lrn_full_no_scale,Dtype)(const int nthreads, __global con * in_off[(head - size) * step]; } scale_val = k + accum_scale * alpha_over_size; - out_off[(head - post_pad) * step] = in_off[(head - post_pad) * step] * (Dtype)native_powr((Dtype)scale_val, (Dtype)negative_beta); + out_off[(head - post_pad) * step] = in_off[(head - post_pad) * step] * (Dtype)native_powr(scale_val, negative_beta); ++head; } // subtract only @@ -93,7 +93,7 @@ __kernel void TEMPLATE(lrn_full_no_scale,Dtype)(const int nthreads, __global con * in_off[(head - size) * step]; } scale_val = k + accum_scale * alpha_over_size; - out_off[(head - post_pad) * step] = in_off[(head - post_pad) * step] * (Dtype)native_powr((Dtype)scale_val, (Dtype)negative_beta); + out_off[(head - post_pad) * step] = in_off[(head - post_pad) * step] * (Dtype)native_powr(scale_val, negative_beta); ++head; } } From ac58b2f857f469ccadfb2941281659553c43f8dd Mon Sep 17 00:00:00 2001 From: Tomoaki Teshima Date: Fri, 25 Sep 2020 22:33:55 +0900 Subject: [PATCH 03/10] compute capability 8.6 - CC for RTX3090, RTX3080 and RTX3070 --- cmake/OpenCVDetectCUDA.cmake | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cmake/OpenCVDetectCUDA.cmake b/cmake/OpenCVDetectCUDA.cmake index 8a62032167..c55f9f1903 100644 --- a/cmake/OpenCVDetectCUDA.cmake +++ b/cmake/OpenCVDetectCUDA.cmake @@ -90,7 +90,7 @@ if(CUDA_FOUND) set(_arch_pascal "6.0;6.1") set(_arch_volta "7.0") set(_arch_turing "7.5") - set(_arch_ampere "8.0") + set(_arch_ampere "8.0;8.6") if(NOT CMAKE_CROSSCOMPILING) list(APPEND _generations "Auto") endif() From 234117800f6b7e9da9461dd1e2c81dd9e3a3895b Mon Sep 17 00:00:00 2001 From: Tomoaki Teshima Date: Fri, 25 Sep 2020 23:57:15 +0900 Subject: [PATCH 04/10] brush up by following the comments --- modules/cudaarithm/src/core.cpp | 6 ++++-- modules/cudaarithm/test/test_core.cpp | 9 ++++++--- 2 files changed, 10 insertions(+), 5 deletions(-) diff --git a/modules/cudaarithm/src/core.cpp b/modules/cudaarithm/src/core.cpp index aefcfc5af2..368c2fcc41 100644 --- a/modules/cudaarithm/src/core.cpp +++ b/modules/cudaarithm/src/core.cpp @@ -163,10 +163,12 @@ void cv::cuda::flip(InputArray _src, OutputArray _dst, int flipCode, Stream& str _dst.create(src.size(), src.type()); GpuMat dst = getOutputMat(_dst, src.size(), src.type(), stream); - if (src.data == dst.data && ((src.cols & 1) == 1 || (src.rows & 1) == 1)) + bool isInplace = (src.data == dst.data) || (src.refcount == dst.refcount); + bool isSizeOdd = (src.cols & 1) == 1 || (src.rows & 1) == 1; + if (isInplace && isSizeOdd) CV_Error(Error::BadROISize, "In-place version of flip only accepts even width/height"); - if (src.data != dst.data) + if (isInplace == false) funcs[src.depth()][src.channels() - 1](src, dst, flipCode, StreamAccessor::getStream(stream)); else // in-place ifuncs[src.depth()][src.channels() - 1](src, flipCode, StreamAccessor::getStream(stream)); diff --git a/modules/cudaarithm/test/test_core.cpp b/modules/cudaarithm/test/test_core.cpp index 04b144ff37..728ee12b16 100644 --- a/modules/cudaarithm/test/test_core.cpp +++ b/modules/cudaarithm/test/test_core.cpp @@ -281,11 +281,14 @@ CUDA_TEST_P(Flip, Accuracy) CUDA_TEST_P(Flip, AccuracyInplace) { - size.width = (size.width >> 1) << 1; // in-place version only accepts even number - size.height = (size.height >> 1) << 1; // in-place version only accepts even number cv::Mat src = randomMat(size, type); - + bool isSizeOdd = ((size.width & 1) == 1) || ((size.height & 1) == 1); cv::cuda::GpuMat srcDst = loadMat(src, useRoi); + if(isSizeOdd) + { + EXPECT_THROW(cv::cuda::flip(srcDst, srcDst, flip_code), cv::Exception); + return; + } cv::cuda::flip(srcDst, srcDst, flip_code); cv::Mat dst_gold; From 691c655630251b5574d24fa1473b6069ad4ab2aa Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Fri, 25 Sep 2020 22:09:25 +0000 Subject: [PATCH 05/10] ippicv: install third-party-programs.txt file --- cmake/OpenCVFindIPP.cmake | 1 + 1 file changed, 1 insertion(+) diff --git a/cmake/OpenCVFindIPP.cmake b/cmake/OpenCVFindIPP.cmake index 59cd497b95..9bc215f415 100644 --- a/cmake/OpenCVFindIPP.cmake +++ b/cmake/OpenCVFindIPP.cmake @@ -252,6 +252,7 @@ if(NOT DEFINED IPPROOT) else() ocv_install_3rdparty_licenses(ippicv "${ICV_PACKAGE_ROOT}/EULA.txt") endif() + ocv_install_3rdparty_licenses(ippicv "${ICV_PACKAGE_ROOT}/third-party-programs.txt") endif() file(TO_CMAKE_PATH "${IPPROOT}" __IPPROOT) From 48368dc9a1c24d9a6062c4137c949bfc808870a2 Mon Sep 17 00:00:00 2001 From: Tomoaki Teshima Date: Sun, 27 Sep 2020 00:37:52 +0900 Subject: [PATCH 06/10] loosen threshold for Mali --- modules/dnn/test/test_torch_importer.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/modules/dnn/test/test_torch_importer.cpp b/modules/dnn/test/test_torch_importer.cpp index 4b89afc331..3be22d6d25 100644 --- a/modules/dnn/test/test_torch_importer.cpp +++ b/modules/dnn/test/test_torch_importer.cpp @@ -113,7 +113,7 @@ TEST_P(Test_Torch_layers, run_convolution) { // Output reference values are in range [23.4018, 72.0181] double l1 = (target == DNN_TARGET_OPENCL_FP16 || target == DNN_TARGET_MYRIAD) ? 0.08 : default_l1; - double lInf = (target == DNN_TARGET_OPENCL_FP16 || target == DNN_TARGET_MYRIAD) ? 0.42 : default_lInf; + double lInf = (target == DNN_TARGET_OPENCL_FP16 || target == DNN_TARGET_MYRIAD) ? 0.43 : default_lInf; runTorchNet("net_conv", "", false, true, true, l1, lInf); } @@ -165,7 +165,7 @@ TEST_P(Test_Torch_layers, run_concat) TEST_P(Test_Torch_layers, run_depth_concat) { runTorchNet("net_depth_concat", "", false, true, true, 0.0, - target == DNN_TARGET_OPENCL_FP16 ? 0.021 : 0.0); + target == DNN_TARGET_OPENCL_FP16 ? 0.032 : 0.0); } TEST_P(Test_Torch_layers, run_deconv) From 233030e417958a8196c76ebb7883ef9a5302d6d0 Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Sun, 27 Sep 2020 06:37:44 +0000 Subject: [PATCH 07/10] core: force check for string literals are used in the message --- modules/core/include/opencv2/core/check.hpp | 2 +- modules/dnn/src/onnx/onnx_importer.cpp | 6 +++--- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/modules/core/include/opencv2/core/check.hpp b/modules/core/include/opencv2/core/check.hpp index 0e0c7cbf31..d975223cc5 100644 --- a/modules/core/include/opencv2/core/check.hpp +++ b/modules/core/include/opencv2/core/check.hpp @@ -63,7 +63,7 @@ struct CheckContext { #define CV__CHECK_LOCATION_VARNAME(id) CVAUX_CONCAT(CVAUX_CONCAT(__cv_check_, id), __LINE__) #define CV__DEFINE_CHECK_CONTEXT(id, message, testOp, p1_str, p2_str) \ static const cv::detail::CheckContext CV__CHECK_LOCATION_VARNAME(id) = \ - { CV__CHECK_FUNCTION, CV__CHECK_FILENAME, __LINE__, testOp, message, p1_str, p2_str } + { CV__CHECK_FUNCTION, CV__CHECK_FILENAME, __LINE__, testOp, "" message, "" p1_str, "" p2_str } CV_EXPORTS void CV_NORETURN check_failed_auto(const int v1, const int v2, const CheckContext& ctx); CV_EXPORTS void CV_NORETURN check_failed_auto(const size_t v1, const size_t v2, const CheckContext& ctx); diff --git a/modules/dnn/src/onnx/onnx_importer.cpp b/modules/dnn/src/onnx/onnx_importer.cpp index 58e5cf75f5..c6f8b65e16 100644 --- a/modules/dnn/src/onnx/onnx_importer.cpp +++ b/modules/dnn/src/onnx/onnx_importer.cpp @@ -446,8 +446,8 @@ void ONNXImporter::populateNet(Net dstNet) avgLp.set("pool", pool); if (axes.size() == 2) { - CV_CheckEQ(clamp(axes.get(0), inpShape.size()), 1, ("Unsupported " + layer_type + " mode").c_str()); - CV_CheckEQ(clamp(axes.get(1), inpShape.size()), 2, ("Unsupported " + layer_type + " mode").c_str()); + CV_CheckEQ(clamp(axes.get(0), inpShape.size()), 1, "Unsupported mode"); + CV_CheckEQ(clamp(axes.get(1), inpShape.size()), 2, "Unsupported mode"); avgLp.set("global_pooling", true); } else @@ -489,7 +489,7 @@ void ONNXImporter::populateNet(Net dstNet) } else if (!layerParams.has("axes") && (layer_type == "ReduceMean" || layer_type == "ReduceSum" || layer_type == "ReduceMax")) { - CV_CheckEQ(layerParams.get("keepdims"), 0, (layer_type + " layer only supports keepdims = false").c_str()); + CV_CheckEQ(layerParams.get("keepdims"), 0, "layer only supports keepdims = false"); LayerParams reshapeLp; reshapeLp.name = layerParams.name + "/reshape"; reshapeLp.type = "Reshape"; From 97bb91d5fa0df316060deee77b4a1f31e99bfb42 Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Sun, 27 Sep 2020 21:14:55 +0000 Subject: [PATCH 08/10] ml: fix python test --- modules/ml/misc/python/test/test_goodfeatures.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/modules/ml/misc/python/test/test_goodfeatures.py b/modules/ml/misc/python/test/test_goodfeatures.py index 91fc885513..a590ba9fa9 100644 --- a/modules/ml/misc/python/test/test_goodfeatures.py +++ b/modules/ml/misc/python/test/test_goodfeatures.py @@ -11,7 +11,7 @@ from tests_common import NewOpenCVTests class TestGoodFeaturesToTrack_test(NewOpenCVTests): def test_goodFeaturesToTrack(self): arr = self.get_sample('samples/data/lena.jpg', 0) - original = arr.copy(True) + original = arr.copy() threshes = [ x / 100. for x in range(1,10) ] numPoints = 20000 From c08f29c8038f4a9bb927130a8babd0570aa19f8d Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Sun, 27 Sep 2020 23:42:30 +0000 Subject: [PATCH 09/10] dnn(opencl): fix convolution kernel w/o bias with activation --- modules/dnn/src/ocl4dnn/src/ocl4dnn_conv_spatial.cpp | 3 +++ modules/dnn/src/opencl/conv_layer_spatial.cl | 2 +- modules/dnn/test/test_layers.cpp | 7 ------- 3 files changed, 4 insertions(+), 8 deletions(-) diff --git a/modules/dnn/src/ocl4dnn/src/ocl4dnn_conv_spatial.cpp b/modules/dnn/src/ocl4dnn/src/ocl4dnn_conv_spatial.cpp index 3cca52d5bd..3707a31846 100644 --- a/modules/dnn/src/ocl4dnn/src/ocl4dnn_conv_spatial.cpp +++ b/modules/dnn/src/ocl4dnn/src/ocl4dnn_conv_spatial.cpp @@ -607,6 +607,7 @@ void OCL4DNNConvSpatial::calculateBenchmark(const UMat &bottom, UMat &ver { options_.str(""); options_.clear(); // clear contents and state flags createBasicKernel(1, 1, 1); + CV_Assert(!kernelQueue.empty()); // basic kernel must be available kernel_index_ = kernelQueue.size() - 1; convolve(bottom, verifyTop, weight, bias, numImages, kernelQueue[kernel_index_]); CV_Assert(phash.find(kernelQueue[kernel_index_]->kernelName) != phash.end()); @@ -1713,6 +1714,7 @@ void OCL4DNNConvSpatial::useFirstAvailable(const UMat &bottom, tunerItems[i]->blockHeight, tunerItems[i]->blockDepth)) { + CV_Assert(!kernelQueue.empty()); // basic kernel must be available int kernelIdx = kernelQueue.size() - 1; kernelConfig* config = kernelQueue[kernelIdx].get(); bool failed = false; @@ -1883,6 +1885,7 @@ void OCL4DNNConvSpatial::setupConvolution(const UMat &bottom, CV_LOG_INFO(NULL, "fallback to basic kernel"); options_.str(""); options_.clear(); // clear contents and state flags createBasicKernel(1, 1, 1); + CV_Assert(!kernelQueue.empty()); // basic kernel must be available kernel_index_ = kernelQueue.size() - 1; } this->bestKernelConfig = kernelQueue[kernel_index_]; diff --git a/modules/dnn/src/opencl/conv_layer_spatial.cl b/modules/dnn/src/opencl/conv_layer_spatial.cl index 5d4d6f3add..236e8d029a 100644 --- a/modules/dnn/src/opencl/conv_layer_spatial.cl +++ b/modules/dnn/src/opencl/conv_layer_spatial.cl @@ -205,7 +205,7 @@ __kernel void ConvolveBasic( #if APPLY_BIAS ACTIVATION_FUNCTION(convolved_image, offset, sum[kern] + bias[biasIndex + kern], biasIndex + kern); #else - ACTIVATION_FUNCTION(convolved_image, offset, sum[kern], biasIndex + kern); + ACTIVATION_FUNCTION(convolved_image, offset, sum[kern], kernelNum + kern); #endif } } diff --git a/modules/dnn/test/test_layers.cpp b/modules/dnn/test/test_layers.cpp index f9b0f62379..085e5a51b8 100644 --- a/modules/dnn/test/test_layers.cpp +++ b/modules/dnn/test/test_layers.cpp @@ -2223,13 +2223,6 @@ TEST_P(ConvolutionActivationFusion, Accuracy) if (actType == "Power" && backendId == DNN_BACKEND_OPENCV && (targetId == DNN_TARGET_OPENCL || targetId == DNN_TARGET_OPENCL_FP16)) applyTestTag(CV_TEST_TAG_DNN_SKIP_OPENCL); - // bug: https://github.com/opencv/opencv/issues/17953 - if (actType == "ChannelsPReLU" && bias_term == false && - backendId == DNN_BACKEND_OPENCV && (targetId == DNN_TARGET_OPENCL || targetId == DNN_TARGET_OPENCL_FP16)) - { - applyTestTag(CV_TEST_TAG_DNN_SKIP_OPENCL); - } - Net net; int convId = net.addLayer(convParams.name, convParams.type, convParams); int activId = net.addLayerToPrev(activationParams.name, activationParams.type, activationParams); From c16e2e6234b56233781127058f07102f5b246349 Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Mon, 28 Sep 2020 01:11:15 +0000 Subject: [PATCH 10/10] ios: don't force BUILD_opencv_world=OFF in case of excluded modules --- platforms/ios/build_framework.py | 1 - 1 file changed, 1 deletion(-) diff --git a/platforms/ios/build_framework.py b/platforms/ios/build_framework.py index 0ff736159b..223542d9f5 100755 --- a/platforms/ios/build_framework.py +++ b/platforms/ios/build_framework.py @@ -150,7 +150,6 @@ class Builder: ] if self.debug_info else []) if len(self.exclude) > 0: - args += ["-DBUILD_opencv_world=OFF"] if not self.dynamic else [] args += ["-DBUILD_opencv_%s=OFF" % m for m in self.exclude] if len(self.disable) > 0: