release: OpenCV 4.0.1 (version++)

Merge tag '4.0.1-openvino'
OpenCV 4.0.1 for Intel(R) OpenVINO(TM) toolkit
2026-07-25 13:23:02 +04:00 · 2018-12-22 07:03:30 +00:00 · 2018-12-22 07:01:13 +00:00 · 2018-12-22 05:40:15 +00:00 · 2018-12-22 05:37:29 +00:00 · 2018-12-21 21:48:03 +03:00
303 changed files with 12037 additions and 6673 deletions
@@ -12,6 +12,8 @@ elseif (CMAKE_SYSTEM_PROCESSOR MATCHES "aarch64.*|AARCH64.*")
  set(AARCH64 TRUE)
 endif()

+ocv_warnings_disable(CMAKE_CXX_FLAGS -Wunused-function)
+
 set(TEGRA_COMPILER_FLAGS "")

 if(CV_GCC OR CV_CLANG)
@@ -53,7 +55,7 @@ endif()
 set(CAROTENE_NS "carotene_o4t" CACHE STRING "" FORCE)

 function(compile_carotene)
-  if(ENABLE_NEON)
+  if(";${CPU_BASELINE_FINAL};" MATCHES ";NEON;")
    set(WITH_NEON ON)
  endif()

@@ -30,6 +30,7 @@

 #include <sys/cdefs.h>
 #include <stdint.h>
+#include <string.h>

 __BEGIN_DECLS

@@ -219,71 +219,203 @@ OCV_OPTION(BUILD_ITT                "Build Intel ITT from source"        (NOT MI

 # Optional 3rd party components
 # ===================================================
-OCV_OPTION(WITH_1394           "Include IEEE1394 support"                    ON   IF (NOT ANDROID AND NOT IOS AND NOT WINRT) )
-OCV_OPTION(WITH_AVFOUNDATION   "Use AVFoundation for Video I/O (iOS/Mac)"    ON   IF APPLE)
-OCV_OPTION(WITH_CAROTENE       "Use NVidia carotene acceleration library for ARM platform"                   ON  IF (ARM OR AARCH64) AND NOT IOS AND NOT (CMAKE_VERSION VERSION_LESS "2.8.11"))
-OCV_OPTION(WITH_CPUFEATURES    "Use cpufeatures Android library"             ON   IF ANDROID)
-OCV_OPTION(WITH_VTK            "Include VTK library support (and build opencv_viz module eiher)"             ON  IF (NOT ANDROID AND NOT IOS AND NOT WINRT AND NOT CMAKE_CROSSCOMPILING) )
-OCV_OPTION(WITH_CUDA           "Include NVidia Cuda Runtime support"                                         OFF IF (NOT IOS AND NOT WINRT) )
-OCV_OPTION(WITH_CUFFT          "Include NVidia Cuda Fast Fourier Transform (FFT) library support"            ON  IF (NOT IOS AND NOT WINRT) )
-OCV_OPTION(WITH_CUBLAS         "Include NVidia Cuda Basic Linear Algebra Subprograms (BLAS) library support" ON IF (NOT IOS AND NOT WINRT) )
-OCV_OPTION(WITH_NVCUVID        "Include NVidia Video Decoding library support"                               ON IF (NOT IOS AND NOT APPLE) )
-OCV_OPTION(WITH_EIGEN          "Include Eigen2/Eigen3 support"               (NOT CV_DISABLE_OPTIMIZATION)       IF (NOT WINRT AND NOT CMAKE_CROSSCOMPILING) )
-OCV_OPTION(WITH_FFMPEG         "Include FFMPEG support"                      ON   IF (NOT ANDROID AND NOT IOS AND NOT WINRT) )
-OCV_OPTION(WITH_GSTREAMER      "Include Gstreamer support"                   ON   IF (NOT ANDROID AND NOT IOS AND NOT WINRT) )
-OCV_OPTION(WITH_GSTREAMER_0_10 "Enable Gstreamer 0.10 support (instead of 1.x)"                              OFF )
-OCV_OPTION(WITH_GTK            "Include GTK support"                         ON   IF (UNIX AND NOT APPLE AND NOT ANDROID) )
-OCV_OPTION(WITH_GTK_2_X        "Use GTK version 2"                           OFF  IF (UNIX AND NOT APPLE AND NOT ANDROID) )
-OCV_OPTION(WITH_IPP            "Include Intel IPP support"                   (NOT MINGW AND NOT CV_DISABLE_OPTIMIZATION)  IF (X86_64 OR X86) AND NOT WINRT AND NOT IOS )
-OCV_OPTION(WITH_HALIDE         "Include Halide support"                      OFF)
-OCV_OPTION(WITH_VULKAN         "Include Vulkan support"                      OFF)
-OCV_OPTION(WITH_INF_ENGINE     "Include Intel Inference Engine support"      OFF)
-OCV_OPTION(WITH_JASPER         "Include JPEG2K support"                      ON   IF (NOT IOS) )
-OCV_OPTION(WITH_JPEG           "Include JPEG support"                        ON)
-OCV_OPTION(WITH_WEBP           "Include WebP support"                        ON   IF (NOT WINRT) )
-OCV_OPTION(WITH_OPENEXR        "Include ILM support via OpenEXR"             ON   IF (NOT IOS AND NOT WINRT) )
-OCV_OPTION(WITH_OPENGL         "Include OpenGL support"                      OFF  IF (NOT ANDROID AND NOT WINRT) )
-OCV_OPTION(WITH_OPENVX         "Include OpenVX support"                      OFF)
-OCV_OPTION(WITH_OPENNI         "Include OpenNI support"                      OFF  IF (NOT ANDROID AND NOT IOS AND NOT WINRT) )
-OCV_OPTION(WITH_OPENNI2        "Include OpenNI2 support"                     OFF  IF (NOT ANDROID AND NOT IOS AND NOT WINRT) )
-OCV_OPTION(WITH_PNG            "Include PNG support"                         ON)
-OCV_OPTION(WITH_GDCM           "Include DICOM support"                       OFF)
-OCV_OPTION(WITH_PVAPI          "Include Prosilica GigE support"              OFF   IF (NOT ANDROID AND NOT IOS AND NOT WINRT) )
-OCV_OPTION(WITH_GIGEAPI        "Include Smartek GigE support"                OFF   IF (NOT ANDROID AND NOT IOS AND NOT WINRT) )
-OCV_OPTION(WITH_ARAVIS         "Include Aravis GigE support"                 OFF   IF (NOT ANDROID AND NOT IOS AND NOT WINRT AND NOT WIN32) )
-OCV_OPTION(WITH_QT             "Build with Qt Backend support"               OFF  IF (NOT ANDROID AND NOT IOS AND NOT WINRT) )
-OCV_OPTION(WITH_WIN32UI        "Build with Win32 UI Backend support"         ON   IF WIN32 AND NOT WINRT)
-OCV_OPTION(WITH_TBB            "Include Intel TBB support"                   OFF  IF (NOT IOS AND NOT WINRT) )
-OCV_OPTION(WITH_HPX            "Include Ste||ar Group HPX support"           OFF)
-OCV_OPTION(WITH_OPENMP         "Include OpenMP support"                      OFF)
-OCV_OPTION(WITH_PTHREADS_PF    "Use pthreads-based parallel_for"             ON   IF (NOT WIN32 OR MINGW) )
-OCV_OPTION(WITH_TIFF           "Include TIFF support"                        ON   IF (NOT IOS) )
-OCV_OPTION(WITH_V4L            "Include Video 4 Linux support"               ON   IF (UNIX AND NOT ANDROID AND NOT APPLE) )
-OCV_OPTION(WITH_DSHOW          "Build VideoIO with DirectShow support"       ON   IF (WIN32 AND NOT ARM AND NOT WINRT) )
-OCV_OPTION(WITH_MSMF           "Build VideoIO with Media Foundation support" ON  IF WIN32 )
-OCV_OPTION(WITH_XIMEA          "Include XIMEA cameras support"               OFF  IF (NOT ANDROID AND NOT WINRT) )
-OCV_OPTION(WITH_XINE           "Include Xine support (GPL)"                  OFF  IF (UNIX AND NOT APPLE AND NOT ANDROID) )
-OCV_OPTION(WITH_CLP            "Include Clp support (EPL)"                   OFF)
-OCV_OPTION(WITH_OPENCL         "Include OpenCL Runtime support"              (NOT ANDROID AND NOT CV_DISABLE_OPTIMIZATION)  IF (NOT IOS AND NOT WINRT) )
-OCV_OPTION(WITH_OPENCL_SVM     "Include OpenCL Shared Virtual Memory support" OFF ) # experimental
-OCV_OPTION(WITH_OPENCLAMDFFT   "Include AMD OpenCL FFT library support"      ON   IF (NOT ANDROID AND NOT IOS AND NOT WINRT) )
-OCV_OPTION(WITH_OPENCLAMDBLAS  "Include AMD OpenCL BLAS library support"     ON   IF (NOT ANDROID AND NOT IOS AND NOT WINRT) )
-OCV_OPTION(WITH_DIRECTX        "Include DirectX support"                     ON   IF (WIN32 AND NOT WINRT) )
-OCV_OPTION(WITH_INTELPERC      "Include Intel Perceptual Computing support"  OFF  IF (WIN32 AND NOT WINRT) )
-OCV_OPTION(WITH_LIBREALSENSE   "Include Intel librealsense support"          OFF  IF (NOT WITH_INTELPERC) )
-OCV_OPTION(WITH_VA             "Include VA support"                          OFF  IF (UNIX AND NOT ANDROID) )
-OCV_OPTION(WITH_VA_INTEL       "Include Intel VA-API/OpenCL support"         OFF  IF (UNIX AND NOT ANDROID) )
-OCV_OPTION(WITH_MFX            "Include Intel Media SDK support"             OFF   IF ((UNIX AND NOT ANDROID) OR (WIN32 AND NOT WINRT AND NOT MINGW)) )
-OCV_OPTION(WITH_GDAL           "Include GDAL Support"                        OFF  IF (NOT ANDROID AND NOT IOS AND NOT WINRT) )
-OCV_OPTION(WITH_GPHOTO2        "Include gPhoto2 library support"             OFF  IF (UNIX AND NOT ANDROID AND NOT IOS) )
-OCV_OPTION(WITH_LAPACK         "Include Lapack library support"              (NOT CV_DISABLE_OPTIMIZATION)  IF (NOT ANDROID AND NOT IOS) )
-OCV_OPTION(WITH_ITT            "Include Intel ITT support"                   ON   IF (NOT APPLE_FRAMEWORK) )
-OCV_OPTION(WITH_PROTOBUF       "Enable libprotobuf"                          ON )
-OCV_OPTION(WITH_IMGCODEC_HDR   "Include HDR support"                         ON)
-OCV_OPTION(WITH_IMGCODEC_SUNRASTER "Include SUNRASTER support"               ON)
-OCV_OPTION(WITH_IMGCODEC_PXM   "Include PNM (PBM,PGM,PPM) and PAM formats support" ON)
-OCV_OPTION(WITH_IMGCODEC_PFM   "Include PFM formats support"                 ON)
-OCV_OPTION(WITH_QUIRC          "Include library QR-code decoding"            ON)
+OCV_OPTION(WITH_1394 "Include IEEE1394 support" ON
+  VISIBLE_IF NOT ANDROID AND NOT IOS AND NOT WINRT
+  VERIFY HAVE_DC1394)
+OCV_OPTION(WITH_AVFOUNDATION "Use AVFoundation for Video I/O (iOS/Mac)" ON
+  VISIBLE_IF APPLE
+  VERIFY HAVE_AVFOUNDATION)
+OCV_OPTION(WITH_CAROTENE "Use NVidia carotene acceleration library for ARM platform" ON
+  VISIBLE_IF (ARM OR AARCH64) AND NOT IOS AND NOT (CMAKE_VERSION VERSION_LESS "2.8.11"))
+OCV_OPTION(WITH_CPUFEATURES "Use cpufeatures Android library" ON
+  VISIBLE_IF ANDROID
+  VERIFY HAVE_CPUFEATURES)
+OCV_OPTION(WITH_VTK "Include VTK library support (and build opencv_viz module eiher)" ON
+  VISIBLE_IF NOT ANDROID AND NOT IOS AND NOT WINRT AND NOT CMAKE_CROSSCOMPILING
+  VERIFY HAVE_VTK)
+OCV_OPTION(WITH_CUDA "Include NVidia Cuda Runtime support" OFF
+  VISIBLE_IF NOT IOS AND NOT WINRT
+  VERIFY HAVE_CUDA)
+OCV_OPTION(WITH_CUFFT "Include NVidia Cuda Fast Fourier Transform (FFT) library support" WITH_CUDA
+  VISIBLE_IF WITH_CUDA
+  VERIFY HAVE_CUFFT)
+OCV_OPTION(WITH_CUBLAS "Include NVidia Cuda Basic Linear Algebra Subprograms (BLAS) library support" WITH_CUDA
+  VISIBLE_IF WITH_CUDA
+  VERIFY HAVE_CUBLAS)
+OCV_OPTION(WITH_NVCUVID "Include NVidia Video Decoding library support" WITH_CUDA
+  VISIBLE_IF WITH_CUDA
+  VERIFY HAVE_NVCUVID)
+OCV_OPTION(WITH_EIGEN "Include Eigen2/Eigen3 support" (NOT CV_DISABLE_OPTIMIZATION)
+  VISIBLE_IF NOT WINRT AND NOT CMAKE_CROSSCOMPILING
+  VERIFY HAVE_EIGEN)
+OCV_OPTION(WITH_FFMPEG "Include FFMPEG support" ON
+  VISIBLE_IF NOT ANDROID AND NOT IOS AND NOT WINRT
+  VERIFY HAVE_FFMPEG)
+OCV_OPTION(WITH_GSTREAMER "Include Gstreamer support" ON
+  VISIBLE_IF NOT ANDROID AND NOT IOS AND NOT WINRT
+  VERIFY HAVE_GSTREAMER AND GSTREAMER_BASE_VERSION VERSION_GREATER "0.99")
+OCV_OPTION(WITH_GSTREAMER_0_10 "Enable Gstreamer 0.10 support (instead of 1.x)" OFF
+  VISIBLE_IF TRUE
+  VERIFY HAVE_GSTREAMER AND GSTREAMER_BASE_VERSION VERSION_LESS "1.0")
+OCV_OPTION(WITH_GTK "Include GTK support" ON
+  VISIBLE_IF UNIX AND NOT APPLE AND NOT ANDROID
+  VERIFY HAVE_GTK)
+OCV_OPTION(WITH_GTK_2_X "Use GTK version 2" OFF
+  VISIBLE_IF UNIX AND NOT APPLE AND NOT ANDROID
+  VERIFY HAVE_GTK AND NOT HAVE_GTK3)
+OCV_OPTION(WITH_IPP "Include Intel IPP support" (NOT MINGW AND NOT CV_DISABLE_OPTIMIZATION)
+  VISIBLE_IF (X86_64 OR X86) AND NOT WINRT AND NOT IOS
+  VERIFY HAVE_IPP)
+OCV_OPTION(WITH_HALIDE "Include Halide support" OFF
+  VISIBLE_IF TRUE
+  VERIFY HAVE_HALIDE)
+OCV_OPTION(WITH_VULKAN "Include Vulkan support" OFF
+  VISIBLE_IF TRUE
+  VERIFY HAVE_VULKAN)
+OCV_OPTION(WITH_INF_ENGINE "Include Intel Inference Engine support" OFF
+  VISIBLE_IF TRUE
+  VERIFY INF_ENGINE_TARGET)
+OCV_OPTION(WITH_JASPER "Include JPEG2K support" ON
+  VISIBLE_IF NOT IOS
+  VERIFY HAVE_JASPER)
+OCV_OPTION(WITH_JPEG "Include JPEG support" ON
+  VISIBLE_IF TRUE
+  VERIFY HAVE_JPEG)
+OCV_OPTION(WITH_WEBP "Include WebP support" ON
+  VISIBLE_IF NOT WINRT
+  VERIFY HAVE_WEBP)
+OCV_OPTION(WITH_OPENEXR "Include ILM support via OpenEXR" ON
+  VISIBLE_IF NOT IOS AND NOT WINRT
+  VERIFY HAVE_OPENEXR)
+OCV_OPTION(WITH_OPENGL "Include OpenGL support" OFF
+  VISIBLE_IF NOT ANDROID AND NOT WINRT
+  VERIFY HAVE_OPENGL)
+OCV_OPTION(WITH_OPENVX "Include OpenVX support" OFF
+  VISIBLE_IF TRUE
+  VERIFY HAVE_OPENVX)
+OCV_OPTION(WITH_OPENNI "Include OpenNI support" OFF
+  VISIBLE_IF NOT ANDROID AND NOT IOS AND NOT WINRT
+  VERIFY HAVE_OPENNI)
+OCV_OPTION(WITH_OPENNI2 "Include OpenNI2 support" OFF
+  VISIBLE_IF NOT ANDROID AND NOT IOS AND NOT WINRT
+  VERIFY HAVE_OPENNI2)
+OCV_OPTION(WITH_PNG "Include PNG support" ON
+  VISIBLE_IF TRUE
+  VERIFY HAVE_PNG)
+OCV_OPTION(WITH_GDCM "Include DICOM support" OFF
+  VISIBLE_IF TRUE
+  VERIFY HAVE_GDCM)
+OCV_OPTION(WITH_PVAPI "Include Prosilica GigE support" OFF
+  VISIBLE_IF NOT ANDROID AND NOT IOS AND NOT WINRT
+  VERIFY HAVE_PVAPI)
+OCV_OPTION(WITH_GIGEAPI "Include Smartek GigE support" OFF
+  VISIBLE_IF NOT ANDROID AND NOT IOS AND NOT WINRT
+  VERIFY HAVE_GIGE_API)
+OCV_OPTION(WITH_ARAVIS "Include Aravis GigE support" OFF
+  VISIBLE_IF NOT ANDROID AND NOT IOS AND NOT WINRT AND NOT WIN32
+  VERIFY HAVE_ARAVIS_API)
+OCV_OPTION(WITH_QT "Build with Qt Backend support" OFF
+  VISIBLE_IF NOT ANDROID AND NOT IOS AND NOT WINRT
+  VERIFY HAVE_QT)
+OCV_OPTION(WITH_WIN32UI "Build with Win32 UI Backend support" ON
+  VISIBLE_IF WIN32 AND NOT WINRT
+  VERIFY HAVE_WIN32UI)
+OCV_OPTION(WITH_TBB "Include Intel TBB support" OFF
+  VISIBLE_IF NOT IOS AND NOT WINRT
+  VERIFY HAVE_TBB)
+OCV_OPTION(WITH_HPX "Include Ste||ar Group HPX support" OFF
+  VISIBLE_IF TRUE
+  VERIFY HAVE_HPX)
+OCV_OPTION(WITH_OPENMP "Include OpenMP support" OFF
+  VISIBLE_IF TRUE
+  VERIFY HAVE_OPENMP)
+OCV_OPTION(WITH_PTHREADS_PF "Use pthreads-based parallel_for" ON
+  VISIBLE_IF NOT WIN32 OR MINGW
+  VERIFY HAVE_PTHREADS_PF)
+OCV_OPTION(WITH_TIFF "Include TIFF support" ON
+  VISIBLE_IF NOT IOS
+  VERIFY HAVE_TIFF)
+OCV_OPTION(WITH_V4L "Include Video 4 Linux support" ON
+  VISIBLE_IF UNIX AND NOT ANDROID AND NOT APPLE
+  VERIFY HAVE_CAMV4L OR HAVE_CAMV4L2 OR HAVE_VIDEOIO)
+OCV_OPTION(WITH_DSHOW "Build VideoIO with DirectShow support" ON
+  VISIBLE_IF WIN32 AND NOT ARM AND NOT WINRT
+  VERIFY HAVE_DSHOW)
+OCV_OPTION(WITH_MSMF "Build VideoIO with Media Foundation support" NOT MINGW
+  VISIBLE_IF WIN32
+  VERIFY HAVE_MSMF)
+OCV_OPTION(WITH_MSMF_DXVA "Enable hardware acceleration in Media Foundation backend" WITH_MSMF
+  VISIBLE_IF WIN32
+  VERIFY HAVE_MSMF_DXVA)
+OCV_OPTION(WITH_XIMEA "Include XIMEA cameras support" OFF
+  VISIBLE_IF NOT ANDROID AND NOT WINRT
+  VERIFY HAVE_XIMEA)
+OCV_OPTION(WITH_XINE "Include Xine support (GPL)" OFF
+  VISIBLE_IF UNIX AND NOT APPLE AND NOT ANDROID
+  VERIFY HAVE_XINE)
+OCV_OPTION(WITH_CLP "Include Clp support (EPL)" OFF
+  VISIBLE_IF TRUE
+  VERIFY HAVE_CLP)
+OCV_OPTION(WITH_OPENCL "Include OpenCL Runtime support" (NOT ANDROID AND NOT CV_DISABLE_OPTIMIZATION)
+  VISIBLE_IF NOT IOS AND NOT WINRT
+  VERIFY HAVE_OPENCL)
+OCV_OPTION(WITH_OPENCL_SVM "Include OpenCL Shared Virtual Memory support" OFF
+  VISIBLE_IF TRUE
+  VERIFY HAVE_OPENCL_SVM) # experimental
+OCV_OPTION(WITH_OPENCLAMDFFT "Include AMD OpenCL FFT library support" ON
+  VISIBLE_IF NOT ANDROID AND NOT IOS AND NOT WINRT
+  VERIFY HAVE_CLAMDFFT)
+OCV_OPTION(WITH_OPENCLAMDBLAS "Include AMD OpenCL BLAS library support" ON
+  VISIBLE_IF NOT ANDROID AND NOT IOS AND NOT WINRT
+  VERIFY HAVE_CLAMDBLAS)
+OCV_OPTION(WITH_DIRECTX "Include DirectX support" ON
+  VISIBLE_IF WIN32 AND NOT WINRT
+  VERIFY HAVE_DIRECTX)
+OCV_OPTION(WITH_INTELPERC "Include Intel Perceptual Computing support" OFF
+  VISIBLE_IF WIN32 AND NOT WINRT
+  VERIFY HAVE_INTELPERC)
+OCV_OPTION(WITH_LIBREALSENSE "Include Intel librealsense support" OFF
+  VISIBLE_IF NOT WITH_INTELPERC
+  VERIFY HAVE_LIBREALSENSE)
+OCV_OPTION(WITH_VA "Include VA support" OFF
+  VISIBLE_IF UNIX AND NOT ANDROID
+  VERIFY HAVE_VA)
+OCV_OPTION(WITH_VA_INTEL "Include Intel VA-API/OpenCL support" OFF
+  VISIBLE_IF UNIX AND NOT ANDROID
+  VERIFY HAVE_VA_INTEL)
+OCV_OPTION(WITH_MFX "Include Intel Media SDK support" OFF
+  VISIBLE_IF (UNIX AND NOT ANDROID) OR (WIN32 AND NOT WINRT AND NOT MINGW)
+  VERIFY HAVE_MFX)
+OCV_OPTION(WITH_GDAL "Include GDAL Support" OFF
+  VISIBLE_IF NOT ANDROID AND NOT IOS AND NOT WINRT
+  VERIFY HAVE_GDAL)
+OCV_OPTION(WITH_GPHOTO2 "Include gPhoto2 library support" OFF
+  VISIBLE_IF UNIX AND NOT ANDROID AND NOT IOS
+  VERIFY HAVE_GPHOTO2)
+OCV_OPTION(WITH_LAPACK "Include Lapack library support" (NOT CV_DISABLE_OPTIMIZATION)
+  VISIBLE_IF NOT ANDROID AND NOT IOS
+  VERIFY HAVE_LAPACK)
+OCV_OPTION(WITH_ITT "Include Intel ITT support" ON
+  VISIBLE_IF NOT APPLE_FRAMEWORK
+  VERIFY HAVE_ITT)
+OCV_OPTION(WITH_PROTOBUF "Enable libprotobuf" ON
+  VISIBLE_IF TRUE
+  VERIFY HAVE_PROTOBUF)
+OCV_OPTION(WITH_IMGCODEC_HDR "Include HDR support" ON
+  VISIBLE_IF TRUE
+  VERIFY HAVE_IMGCODEC_HDR)
+OCV_OPTION(WITH_IMGCODEC_SUNRASTER "Include SUNRASTER support" ON
+  VISIBLE_IF TRUE
+  VERIFY HAVE_IMGCODEC_SUNRASTER)
+OCV_OPTION(WITH_IMGCODEC_PXM "Include PNM (PBM,PGM,PPM) and PAM formats support" ON
+  VISIBLE_IF TRUE
+  VERIFY HAVE_IMGCODEC_PXM)
+OCV_OPTION(WITH_IMGCODEC_PFM "Include PFM formats support" ON
+  VISIBLE_IF TRUE
+  VERIFY HAVE_IMGCODEC_PFM)
+OCV_OPTION(WITH_QUIRC "Include library QR-code decoding" ON
+  VISIBLE_IF TRUE
+  VERIFY HAVE_QUIRC)

 # OpenCV build components
 # ===================================================
@@ -323,7 +455,6 @@ OCV_OPTION(ENABLE_PROFILING           "Enable profiling in the GCC compiler (Add
 OCV_OPTION(ENABLE_COVERAGE            "Enable coverage collection with  GCov"                    OFF  IF CV_GCC )
 OCV_OPTION(ENABLE_OMIT_FRAME_POINTER  "Enable -fomit-frame-pointer for GCC"                      ON   IF CV_GCC )
 OCV_OPTION(ENABLE_POWERPC             "Enable PowerPC for GCC"                                   ON   IF (CV_GCC AND CMAKE_SYSTEM_PROCESSOR MATCHES powerpc.*) )
-OCV_OPTION(ENABLE_VSX                 "Enable POWER8 and above VSX (64-bit little-endian)"       ON   IF ((CV_GCC OR CV_CLANG) AND PPC64LE) )
 OCV_OPTION(ENABLE_FAST_MATH           "Enable -ffast-math (not recommended for GCC 4.6.x)"       OFF  IF (CV_GCC AND (X86 OR X86_64)) )
 if(NOT IOS AND (NOT ANDROID OR OPENCV_ANDROID_USE_LEGACY_FLAGS))  # Use CPU_BASELINE instead
 OCV_OPTION(ENABLE_NEON                "Enable NEON instructions"                                 (NEON OR ANDROID_ARM_NEON OR AARCH64) IF (CV_GCC OR CV_CLANG) AND (ARM OR AARCH64 OR IOS) )
@@ -344,6 +475,7 @@ OCV_OPTION(CV_ENABLE_INTRINSICS       "Use intrinsic-based optimized code" ON )
 OCV_OPTION(CV_DISABLE_OPTIMIZATION    "Disable explicit optimized code (dispatched code/intrinsics/loop unrolling/etc)" OFF )
 OCV_OPTION(CV_TRACE                   "Enable OpenCV code trace" ON)
 OCV_OPTION(OPENCV_GENERATE_SETUPVARS  "Generate setup_vars* scripts" ON IF (NOT ANDROID AND NOT APPLE_FRAMEWORK) )
+OCV_OPTION(ENABLE_CONFIG_VERIFICATION "Fail build if actual configuration doesn't match requested (WITH_XXX != HAVE_XXX)" OFF)

 OCV_OPTION(ENABLE_PYLINT              "Add target with Pylint checks"                            (BUILD_DOCS OR BUILD_EXAMPLES) IF (NOT CMAKE_CROSSCOMPILING AND NOT APPLE_FRAMEWORK) )
 OCV_OPTION(ENABLE_FLAKE8              "Add target with Python flake8 checker"                    (BUILD_DOCS OR BUILD_EXAMPLES) IF (NOT CMAKE_CROSSCOMPILING AND NOT APPLE_FRAMEWORK) )
@@ -461,7 +593,7 @@ else()
    ocv_update(OPENCV_OTHER_INSTALL_PATH         "${CMAKE_INSTALL_DATAROOTDIR}/opencv4")
    ocv_update(OPENCV_LICENSES_INSTALL_PATH      "${CMAKE_INSTALL_DATAROOTDIR}/licenses/opencv4")
  endif()
-  ocv_update(OPENCV_PYTHON_INSTALL_PATH  "python")
+  #ocv_update(OPENCV_PYTHON_INSTALL_PATH  "python")  # no default value, see https://github.com/opencv/opencv/issues/13202
 endif()

 ocv_update(CMAKE_INSTALL_RPATH "${CMAKE_INSTALL_PREFIX}/${OPENCV_LIB_INSTALL_PATH}")
@@ -1318,6 +1450,7 @@ endif()

 if(WITH_MSMF OR HAVE_MSMF)
  status("    Media Foundation:" HAVE_MSMF    THEN YES                                        ELSE NO)
+  status("      DXVA:"           HAVE_MSMF_DXVA    THEN YES    ELSE NO)
 endif()

 if(WITH_XIMEA OR HAVE_XIMEA)
@@ -1479,7 +1612,7 @@ if(BUILD_opencv_python2)
    status("    Libraries:"   HAVE_opencv_python2  THEN  "${PYTHON2_LIBRARIES}"                                      ELSE NO)
  endif()
  status("    numpy:"         PYTHON2_NUMPY_INCLUDE_DIRS THEN "${PYTHON2_NUMPY_INCLUDE_DIRS} (ver ${PYTHON2_NUMPY_VERSION})" ELSE "NO (Python wrappers can not be generated)")
-  status("    packages path:" PYTHON2_EXECUTABLE         THEN "${PYTHON2_PACKAGES_PATH}"                                    ELSE "-")
+  status("    install path:"  HAVE_opencv_python2  THEN "${__INSTALL_PATH_PYTHON2}"                            ELSE "-")
 endif()

 if(BUILD_opencv_python3)
@@ -1492,7 +1625,7 @@ if(BUILD_opencv_python3)
    status("    Libraries:"   HAVE_opencv_python3  THEN  "${PYTHON3_LIBRARIES}"                                      ELSE NO)
  endif()
  status("    numpy:"         PYTHON3_NUMPY_INCLUDE_DIRS THEN "${PYTHON3_NUMPY_INCLUDE_DIRS} (ver ${PYTHON3_NUMPY_VERSION})" ELSE "NO (Python3 wrappers can not be generated)")
-  status("    packages path:" PYTHON3_EXECUTABLE         THEN "${PYTHON3_PACKAGES_PATH}"                                    ELSE "-")
+  status("    install path:"  HAVE_opencv_python3  THEN "${__INSTALL_PATH_PYTHON3}"                            ELSE "-")
 endif()

 status("")
@@ -1527,6 +1660,10 @@ status("")

 ocv_finalize_status()

+if(ENABLE_CONFIG_VERIFICATION)
+  ocv_verify_config()
+endif()
+
 ocv_cmake_hook(POST_FINALIZE)

 # ----------------------------------------------------------------------------
@@ -5,6 +5,10 @@
 # AVX / AVX2 / AVX_512F
 # FMA3

+# ppc64le arch:
+# VSX  (always available on Power8)
+# VSX3 (always available on Power9)
+
 # CPU_{opt}_SUPPORTED=ON/OFF - compiler support (possibly with additional flag)
 # CPU_{opt}_IMPLIES=<list>
 # CPU_{opt}_FORCE=<list> - subset of "implies" list
@@ -26,10 +30,12 @@
 # CPU_DISPATCH_FINAL=<list> - final list of dispatched optimizations
 #
 # CPU_DISPATCH_FLAGS_${opt} - flags for source files compiled separately (<name>.avx2.cpp)
+#
+# CPU_{opt}_ENABLED_DEFAULT=ON/OFF - has compiler support without additional flag (CPU_BASELINE_DETECT=ON only)

 set(CPU_ALL_OPTIMIZATIONS "SSE;SSE2;SSE3;SSSE3;SSE4_1;SSE4_2;POPCNT;AVX;FP16;AVX2;FMA3;AVX_512F;AVX512_SKX")
 list(APPEND CPU_ALL_OPTIMIZATIONS NEON VFPV3 FP16)
-list(APPEND CPU_ALL_OPTIMIZATIONS VSX)
+list(APPEND CPU_ALL_OPTIMIZATIONS VSX VSX3)
 list(REMOVE_DUPLICATES CPU_ALL_OPTIMIZATIONS)

 ocv_update(CPU_VFPV3_FEATURE_ALIAS "")
@@ -81,7 +87,7 @@ ocv_optimization_process_obsolete_option(ENABLE_FMA3 FMA3 ON)
 ocv_optimization_process_obsolete_option(ENABLE_VFPV3 VFPV3 OFF)
 ocv_optimization_process_obsolete_option(ENABLE_NEON NEON OFF)

-ocv_optimization_process_obsolete_option(ENABLE_VSX VSX OFF)
+ocv_optimization_process_obsolete_option(ENABLE_VSX VSX ON)

 macro(ocv_is_optimization_in_list resultvar check_opt)
  set(__checked "")
@@ -289,14 +295,24 @@ elseif(ARM OR AARCH64)
    set(CPU_BASELINE "NEON;FP16" CACHE STRING "${HELP_CPU_BASELINE}")
  endif()
 elseif(PPC64LE)
-  ocv_update(CPU_KNOWN_OPTIMIZATIONS "VSX")
+  ocv_update(CPU_KNOWN_OPTIMIZATIONS "VSX;VSX3")
  ocv_update(CPU_VSX_TEST_FILE "${OpenCV_SOURCE_DIR}/cmake/checks/cpu_vsx.cpp")
+  ocv_update(CPU_VSX3_TEST_FILE "${OpenCV_SOURCE_DIR}/cmake/checks/cpu_vsx3.cpp")
+
+  if(NOT OPENCV_CPU_OPT_IMPLIES_IGNORE)
+    ocv_update(CPU_VSX3_IMPLIES "VSX")
+  endif()

  if(CV_CLANG AND (NOT ${CMAKE_CXX_COMPILER} MATCHES "xlc"))
    ocv_update(CPU_VSX_FLAGS_ON "-mvsx -maltivec")
+    ocv_update(CPU_VSX3_FLAGS_ON "-mpower9-vector")
  else()
    ocv_update(CPU_VSX_FLAGS_ON "-mcpu=power8")
+    ocv_update(CPU_VSX3_FLAGS_ON "-mcpu=power9 -mtune=power9")
  endif()
+
+  set(CPU_DISPATCH "VSX3" CACHE STRING "${HELP_CPU_DISPATCH}")
+  set(CPU_BASELINE "VSX" CACHE STRING "${HELP_CPU_BASELINE}")
 endif()

 # Helper values for cmake-gui
@@ -331,6 +347,7 @@ macro(ocv_check_compiler_optimization OPT)
          ocv_check_compiler_flag(CXX "${CPU_BASELINE_FLAGS}" "${_varname}" "${CPU_${OPT}_TEST_FILE}")
          if(${_varname})
            list(APPEND CPU_BASELINE_FINAL ${OPT})
+            set(CPU_${OPT}_ENABLED_DEFAULT ON)
            set(__available 1)
          endif()
        endif()
@@ -448,7 +465,7 @@ foreach(OPT ${CPU_KNOWN_OPTIMIZATIONS})
      if(NOT ";${CPU_BASELINE_FINAL};" MATCHES ";${OPT};")
        list(APPEND CPU_BASELINE_FINAL ${OPT})
      endif()
-      if(NOT CPU_BASELINE_DETECT)  # Don't change compiler flags in 'detection' mode
+      if(NOT CPU_${OPT}_ENABLED_DEFAULT)  # Don't change compiler flags in 'detection' mode
        ocv_append_optimization_flag(CPU_BASELINE_FLAGS ${OPT})
      endif()
    endif()
@@ -1,3 +1,6 @@
+set(OPENCV_JAVA_SOURCE_VERSION "" CACHE STRING "Java source version (javac Ant target)")
+set(OPENCV_JAVA_TARGET_VERSION "" CACHE STRING "Java target version (javac Ant target)")
+
 file(TO_CMAKE_PATH "$ENV{ANT_DIR}" ANT_DIR_ENV_PATH)
 file(TO_CMAKE_PATH "$ENV{ProgramFiles}" ProgramFiles_ENV_PATH)

@@ -52,7 +52,7 @@ if(CUDA_FOUND)

  message(STATUS "CUDA detected: " ${CUDA_VERSION})

-  set(_generations "Fermi" "Kepler" "Maxwell" "Pascal" "Volta")
+  set(_generations "Fermi" "Kepler" "Maxwell" "Pascal" "Volta" "Turing")
  if(NOT CMAKE_CROSSCOMPILING)
    list(APPEND _generations "Auto")
  endif()
@@ -115,7 +115,7 @@ if(CUDA_FOUND)
      string(REGEX REPLACE ".*\n" "" _nvcc_out "${_nvcc_out}") #Strip leading warning messages, if any
      if(NOT _nvcc_res EQUAL 0)
        message(STATUS "Automatic detection of CUDA generation failed. Going to build for all known architectures.")
-        set(__cuda_arch_bin "5.3 6.2 7.0 7.5")
+        set(__cuda_arch_bin "5.3 6.2 7.2")
      else()
        set(__cuda_arch_bin "${_nvcc_out}")
        string(REPLACE "2.1" "2.1(2.0)" __cuda_arch_bin "${__cuda_arch_bin}")
@@ -78,9 +78,9 @@ endif()

 if(INF_ENGINE_TARGET)
  if(NOT INF_ENGINE_RELEASE)
-    message(WARNING "InferenceEngine version have not been set, 2018R4 will be used by default. Set INF_ENGINE_RELEASE variable if you experience build errors.")
+    message(WARNING "InferenceEngine version have not been set, 2018R5 will be used by default. Set INF_ENGINE_RELEASE variable if you experience build errors.")
  endif()
-  set(INF_ENGINE_RELEASE "2018040000" CACHE STRING "Force IE version, should be in form YYYYAABBCC (e.g. 2018R2.0.2 -> 2018020002)")
+  set(INF_ENGINE_RELEASE "2018050000" CACHE STRING "Force IE version, should be in form YYYYAABBCC (e.g. 2018R2.0.2 -> 2018020002)")
  set_target_properties(${INF_ENGINE_TARGET} PROPERTIES
    INTERFACE_COMPILE_DEFINITIONS "HAVE_INF_ENGINE=1;INF_ENGINE_RELEASE=${INF_ENGINE_RELEASE}"
  )
@@ -246,14 +246,15 @@ endif(WITH_DSHOW)
 ocv_clear_vars(HAVE_MSMF)
 if(WITH_MSMF)
  check_include_file(Mfapi.h HAVE_MSMF)
-  check_include_file(D3D11.h D3D11_found)
-  check_include_file(D3d11_4.h D3D11_4_found)
-  if(D3D11_found AND D3D11_4_found)
-    set(HAVE_DXVA YES)
-  else()
-    set(HAVE_DXVA NO)
+  set(HAVE_MSMF_DXVA "")
+  if(WITH_MSMF_DXVA)
+    check_include_file(D3D11.h D3D11_found)
+    check_include_file(D3d11_4.h D3D11_4_found)
+    if(D3D11_found AND D3D11_4_found)
+      set(HAVE_MSMF_DXVA YES)
+    endif()
  endif()
-endif(WITH_MSMF)
+endif()

 # --- Extra HighGUI and VideoIO libs on Windows ---
 if(WIN32)
@@ -22,8 +22,6 @@ set(OPENCV_ABI_HEADERS "{RELPATH}/${OPENCV_INCLUDE_INSTALL_PATH}")
 # Libraries
 set(OPENCV_ABI_LIBRARIES "{RELPATH}/${OPENCV_LIB_INSTALL_PATH}")

-set(OPENCV_ABI_SKIP_HEADERS "")
-set(OPENCV_ABI_SKIP_LIBRARIES "")
 foreach(mod ${OPENCV_MODULES_BUILD})
  string(REGEX REPLACE "^opencv_" "" mod "${mod}")
  if(NOT OPENCV_MODULE_opencv_${mod}_CLASS STREQUAL "PUBLIC"
@@ -44,7 +42,7 @@ string(REPLACE ";" "\n    " OPENCV_ABI_SKIP_HEADERS "${OPENCV_ABI_SKIP_HEADERS}"
 string(REPLACE ";" "\n    " OPENCV_ABI_SKIP_LIBRARIES "${OPENCV_ABI_SKIP_LIBRARIES}")

 # Options
-set(OPENCV_ABI_GCC_OPTIONS "${CMAKE_CXX_FLAGS} ${CMAKE_CXX_FLAGS_RELEASE} -DOPENCV_ABI_CHECK=1")
+set(OPENCV_ABI_GCC_OPTIONS "${CMAKE_CXX_FLAGS} ${CMAKE_CXX_FLAGS_RELEASE} -DOPENCV_ABI_CHECK=1 -DCV_DNN_DONT_ADD_INLINE_NS=1")
 string(REGEX REPLACE "([^ ]) +([^ ])" "\\1\\n    \\2" OPENCV_ABI_GCC_OPTIONS "${OPENCV_ABI_GCC_OPTIONS}")

 configure_file("${CMAKE_CURRENT_SOURCE_DIR}/cmake/templates/opencv_abi.xml.in" "${path1}.base")
@@ -43,11 +43,24 @@ else()
 endif()
 file(RELATIVE_PATH OPENCV_PYTHON_DIR_RELATIVE_CMAKECONFIG
    "${CMAKE_INSTALL_PREFIX}/${OPENCV_SETUPVARS_INSTALL_PATH}/" "${CMAKE_INSTALL_PREFIX}/")
-if(IS_ABSOLUTE "${OPENCV_PYTHON_INSTALL_PATH}")
-  set(OPENCV_PYTHON_DIR_RELATIVE_CMAKECONFIG "${OPENCV_PYTHON_INSTALL_PATH}")
-  message(WARNING "CONFIGURATION IS NOT SUPPORTED: validate setupvars script in install directory")
+if(DEFINED OPENCV_PYTHON_INSTALL_PATH)
+  set(__python_path "${OPENCV_PYTHON_INSTALL_PATH}")
+elseif(DEFINED OPENCV_PYTHON_INSTALL_PATH_SETUPVARS)
+  set(__python_path "${OPENCV_PYTHON_INSTALL_PATH_SETUPVARS}")
+endif()
+if(DEFINED __python_path)
+  if(IS_ABSOLUTE "${__python_path}")
+    set(OPENCV_PYTHON_DIR_RELATIVE_CMAKECONFIG "${__python_path}")
+    message(WARNING "CONFIGURATION IS NOT SUPPORTED: validate setupvars script in install directory")
+  else()
+    ocv_path_join(OPENCV_PYTHON_DIR_RELATIVE_CMAKECONFIG "${OPENCV_PYTHON_DIR_RELATIVE_CMAKECONFIG}" "${__python_path}")
+  endif()
 else()
-  ocv_path_join(OPENCV_PYTHON_DIR_RELATIVE_CMAKECONFIG "${OPENCV_PYTHON_DIR_RELATIVE_CMAKECONFIG}" "${OPENCV_PYTHON_INSTALL_PATH}")
+  if(DEFINED OPENCV_PYTHON3_INSTALL_PATH)
+    ocv_path_join(OPENCV_PYTHON_DIR_RELATIVE_CMAKECONFIG "${OPENCV_PYTHON_DIR_RELATIVE_CMAKECONFIG}" "${OPENCV_PYTHON3_INSTALL_PATH}")
+  else()
+    set(OPENCV_PYTHON_DIR_RELATIVE_CMAKECONFIG "python_loader_is_not_installed")
+  endif()
 endif()
 configure_file("${OpenCV_SOURCE_DIR}/cmake/templates/${OPENCV_SETUPVARS_TEMPLATE}" "${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/install/${OPENCV_SETUPVARS_FILENAME}" @ONLY)
 install(FILES "${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/install/${OPENCV_SETUPVARS_FILENAME}"
@@ -779,6 +779,7 @@ macro(ocv_glob_module_sources)
       "${CMAKE_CURRENT_LIST_DIR}/include/opencv2/${name}/hal/*.h"
       "${CMAKE_CURRENT_LIST_DIR}/include/opencv2/${name}/utils/*.hpp"
       "${CMAKE_CURRENT_LIST_DIR}/include/opencv2/${name}/utils/*.h"
+       "${CMAKE_CURRENT_LIST_DIR}/include/opencv2/${name}/legacy/*.h"
  )
  file(GLOB lib_hdrs_detail
       "${CMAKE_CURRENT_LIST_DIR}/include/opencv2/${name}/detail/*.hpp"
@@ -909,7 +910,11 @@ macro(_ocv_create_module)
      source_group("Src" FILES "${_VS_VERSION_FILE}")
    endif()
  endif()
-  if(WIN32 AND NOT ("${the_module}" STREQUAL "opencv_core" OR "${the_module}" STREQUAL "opencv_world")
+  if(WIN32 AND NOT (
+          "${the_module}" STREQUAL "opencv_core" OR
+          "${the_module}" STREQUAL "opencv_world" OR
+          "${the_module}" STREQUAL "opencv_cudev"
+      )
      AND (BUILD_SHARED_LIBS AND NOT "x${OPENCV_MODULE_TYPE}" STREQUAL "xSTATIC")
      AND NOT OPENCV_SKIP_DLLMAIN_GENERATION
  )
@@ -1011,6 +1016,8 @@ macro(_ocv_create_module)
      string(REGEX REPLACE "^.*opencv2/" "opencv2/" hdr2 "${hdr}")
      if(NOT hdr2 MATCHES "private" AND hdr2 MATCHES "^(opencv2/?.*)/[^/]+.h(..)?$" )
        install(FILES ${hdr} OPTIONAL DESTINATION "${OPENCV_INCLUDE_INSTALL_PATH}/${CMAKE_MATCH_1}" COMPONENT dev)
+      else()
+        #message("Header file will be NOT installed: ${hdr}")
      endif()
    endforeach()
  endif()
@@ -508,7 +508,7 @@ macro(ocv_warnings_disable)
      foreach(var ${_flag_vars})
        foreach(warning ${_gxx_warnings})
          if(NOT warning MATCHES "^-Wno-")
-            string(REGEX REPLACE "${warning}(=[^ ]*)?" "" ${var} "${${var}}")
+            string(REGEX REPLACE "(^|[ ]+)${warning}(=[^ ]*)?([ ]+|$)" " " ${var} "${${var}}")
            string(REPLACE "-W" "-Wno-" warning "${warning}")
          endif()
          ocv_check_flag_support(${var} "${warning}" _varname "")
@@ -571,14 +571,21 @@ endmacro()
 # Provides an option that the user can optionally select.
 # Can accept condition to control when option is available for user.
 # Usage:
-#   option(<option_variable> "help string describing the option" <initial value or boolean expression> [IF <condition>])
+#   option(<option_variable>
+#          "help string describing the option"
+#          <initial value or boolean expression>
+#          [VISIBLE_IF <condition>]
+#          [VERIFY <condition>])
 macro(OCV_OPTION variable description value)
  set(__value ${value})
  set(__condition "")
+  set(__verification)
  set(__varname "__value")
  foreach(arg ${ARGN})
-    if(arg STREQUAL "IF" OR arg STREQUAL "if")
+    if(arg STREQUAL "IF" OR arg STREQUAL "if" OR arg STREQUAL "VISIBLE_IF")
      set(__varname "__condition")
+    elseif(arg STREQUAL "VERIFY")
+      set(__varname "__verification")
    else()
      list(APPEND ${__varname} ${arg})
    endif()
@@ -614,10 +621,46 @@ macro(OCV_OPTION variable description value)
      unset(${variable} CACHE)
    endif()
  endif()
+  if(__verification)
+    set(OPENCV_VERIFY_${variable} "${__verification}") # variable containing condition to verify
+    list(APPEND OPENCV_VERIFICATIONS "${variable}") # list of variable names (WITH_XXX;WITH_YYY;...)
+  endif()
  unset(__condition)
  unset(__value)
 endmacro()

+
+# Check that each variable stored in OPENCV_VERIFICATIONS list
+# is consistent with actual detection result (stored as condition in OPENCV_VERIFY_...) variables
+function(ocv_verify_config)
+  set(broken_options)
+  foreach(var ${OPENCV_VERIFICATIONS})
+    set(evaluated FALSE)
+    if(${OPENCV_VERIFY_${var}})
+      set(evaluated TRUE)
+    endif()
+    status("Verifying ${var}=${${var}} => '${OPENCV_VERIFY_${var}}'=${evaluated}")
+    if (${var} AND NOT evaluated)
+      list(APPEND broken_options ${var})
+      message(WARNING
+        "Option ${var} is enabled but corresponding dependency "
+        "have not been found: \"${OPENCV_VERIFY_${var}}\" is FALSE")
+    elseif(NOT ${var} AND evaluated)
+      list(APPEND broken_options ${var})
+      message(WARNING
+        "Option ${var} is disabled or unset but corresponding dependency "
+        "have been explicitly turned on: \"${OPENCV_VERIFY_${var}}\" is TRUE")
+    endif()
+  endforeach()
+  if(broken_options)
+    string(REPLACE ";" "\n" broken_options "${broken_options}")
+    message(FATAL_ERROR
+      "Some dependencies have not been found or have been forced, "
+      "unset ENABLE_CONFIG_VERIFICATION option to ignore these failures "
+      "or change following options:\n${broken_options}")
+  endif()
+endfunction()
+
 # Usage: ocv_append_build_options(HIGHGUI FFMPEG)
 macro(ocv_append_build_options var_prefix pkg_prefix)
  foreach(suffix INCLUDE_DIRS LIBRARIES LIBRARY_DIRS)
@@ -15,7 +15,7 @@ set(OPENCV_LIBVERSION "${OPENCV_VERSION_MAJOR}.${OPENCV_VERSION_MINOR}.${OPENCV_

 # create a dependency on the version file
 # we never use the output of the following command but cmake will rerun automatically if the version file changes
-configure_file("${OPENCV_VERSION_FILE}" "${CMAKE_BINARY_DIR}/junk/version.junk" COPYONLY)
+configure_file("${OPENCV_VERSION_FILE}" "${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/opencv_junk/version.junk" COPYONLY)

 ocv_update(OPENCV_VS_VER_FILEVERSION_QUAD "${OPENCV_VERSION_MAJOR},${OPENCV_VERSION_MINOR},${OPENCV_VERSION_PATCH},0")
 ocv_update(OPENCV_VS_VER_PRODUCTVERSION_QUAD "${OPENCV_VERSION_MAJOR},${OPENCV_VERSION_MINOR},${OPENCV_VERSION_PATCH},0")
@@ -8,6 +8,24 @@ if(DEFINED ANDROID_NDK_REVISION AND ANDROID_NDK_REVISION MATCHES "(1[56])([0-9]+
  set(ANDROID_NDK_REVISION "${ANDROID_NDK_REVISION}" CACHE INTERNAL "Android NDK revision")
 endif()

+# fixup -g option: https://github.com/opencv/opencv/issues/8460#issuecomment-434249750
+if((INSTALL_CREATE_DISTRIB OR CMAKE_BUILD_TYPE STREQUAL "Release")
+  AND NOT OPENCV_SKIP_ANDROID_G_OPTION_FIX
+)
+  if(" ${CMAKE_CXX_FLAGS} " MATCHES " -g ")
+    message(STATUS "Android: fixup -g compiler option from Android toolchain")
+  endif()
+  string(REPLACE " -g " " " CMAKE_CXX_FLAGS " ${CMAKE_CXX_FLAGS} ")
+  string(REPLACE " -g " " " CMAKE_C_FLAGS " ${CMAKE_C_FLAGS} ")
+  string(REPLACE " -g " " " CMAKE_ASM_FLAGS " ${CMAKE_ASM_FLAGS} ")
+  if(NOT " ${CMAKE_CXX_FLAGS_DEBUG}" MATCHES " -g")
+    set(CMAKE_CXX_FLAGS_DEBUG "${CMAKE_CXX_FLAGS_DEBUG} -g")
+  endif()
+  if(NOT " ${CMAKE_C_FLAGS_DEBUG}" MATCHES " -g")
+    set(CMAKE_C_FLAGS_DEBUG "${CMAKE_C_FLAGS_DEBUG} -g")
+  endif()
+endif()
+
 # https://developer.android.com/studio/command-line/variables.html
 ocv_check_environment_variables(ANDROID_SDK_ROOT ANDROID_HOME ANDROID_SDK)

@@ -1,5 +1,151 @@
-message(FATAL_ERROR "
-Android gradle-based build/projects are not supported in this version of OpenCV.
-You need to downgrade Android SDK Tools to version 25.2.5.
-Details: https://github.com/opencv/opencv/issues/8460
+# https://developer.android.com/studio/releases/gradle-plugin
+set(ANDROID_GRADLE_PLUGIN_VERSION "3.2.1" CACHE STRING "Android Gradle Plugin version (3.0+)")
+message(STATUS "Android Gradle Plugin version: ${ANDROID_GRADLE_PLUGIN_VERSION}")
+
+set(ANDROID_COMPILE_SDK_VERSION "26" CACHE STRING "Android compileSdkVersion")
+set(ANDROID_MIN_SDK_VERSION "21" CACHE STRING "Android minSdkVersion")
+set(ANDROID_TARGET_SDK_VERSION "26" CACHE STRING "Android minSdkVersion")
+
+set(ANDROID_BUILD_BASE_DIR "${OpenCV_BINARY_DIR}/opencv_android" CACHE INTERNAL "")
+set(ANDROID_TMP_INSTALL_BASE_DIR "${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/install/opencv_android")
+
+set(ANDROID_INSTALL_SAMPLES_DIR "samples")
+
+set(ANDROID_BUILD_ABI_FILTER "
+reset()
+include '${ANDROID_ABI}'
 ")
+
+set(ANDROID_INSTALL_ABI_FILTER "
+//reset()
+//include 'armeabi-v7a'
+//include 'arm64-v8a'
+//include 'x86'
+//include 'x86_64'
+")
+if(NOT INSTALL_CREATE_DISTRIB)
+  set(ANDROID_INSTALL_ABI_FILTER "${ANDROID_BUILD_ABI_FILTER}")
+endif()
+
+# BUG: Ninja generator generates broken targets with ANDROID_ABI_FILTER name (CMake 3.11.2)
+#set(__spaces "                        ")
+#string(REPLACE "\n" "\n${__spaces}" ANDROID_ABI_FILTER "${__spaces}${ANDROID_BUILD_ABI_FILTER}")
+#string(REPLACE REGEX "[ ]+$" "" ANDROID_ABI_FILTER "${ANDROID_ABI_FILTER}")
+set(ANDROID_ABI_FILTER "${ANDROID_BUILD_ABI_FILTER}")
+configure_file("${OpenCV_SOURCE_DIR}/samples/android/build.gradle.in" "${ANDROID_BUILD_BASE_DIR}/build.gradle" @ONLY)
+
+set(ANDROID_ABI_FILTER "${ANDROID_INSTALL_ABI_FILTER}")
+configure_file("${OpenCV_SOURCE_DIR}/samples/android/build.gradle.in" "${ANDROID_TMP_INSTALL_BASE_DIR}/${ANDROID_INSTALL_SAMPLES_DIR}/build.gradle" @ONLY)
+install(FILES "${ANDROID_TMP_INSTALL_BASE_DIR}/${ANDROID_INSTALL_SAMPLES_DIR}/build.gradle" DESTINATION "${ANDROID_INSTALL_SAMPLES_DIR}" COMPONENT samples)
+
+set(GRADLE_WRAPPER_FILES
+    "gradle/wrapper/gradle-wrapper.jar"
+    "gradle/wrapper/gradle-wrapper.properties"
+    "gradlew.bat"
+    "gradlew"
+    "gradle.properties"
+)
+foreach(fname ${GRADLE_WRAPPER_FILES})
+  get_filename_component(__dir "${fname}" DIRECTORY)
+  set(__permissions "")
+  set(__permissions_prefix "")
+  if(fname STREQUAL "gradlew")
+    set(__permissions FILE_PERMISSIONS OWNER_READ OWNER_WRITE OWNER_EXECUTE GROUP_READ GROUP_EXECUTE WORLD_READ WORLD_EXECUTE)
+  endif()
+  file(COPY "${OpenCV_SOURCE_DIR}/platforms/android/gradle-wrapper/${fname}" DESTINATION "${ANDROID_BUILD_BASE_DIR}/${__dir}" ${__permissions})
+  string(REPLACE "FILE_PERMISSIONS" "PERMISSIONS" __permissions "${__permissions}")
+  if("${__dir}" STREQUAL "")
+    set(__dir ".")
+  endif()
+  install(FILES "${OpenCV_SOURCE_DIR}/platforms/android/gradle-wrapper/${fname}" DESTINATION "${ANDROID_INSTALL_SAMPLES_DIR}/${__dir}" COMPONENT samples ${__permissions})
+endforeach()
+
+file(WRITE "${ANDROID_BUILD_BASE_DIR}/settings.gradle" "
+include ':opencv'
+")
+
+file(WRITE "${ANDROID_TMP_INSTALL_BASE_DIR}/settings.gradle" "
+rootProject.name = 'opencv_samples'
+
+def opencvsdk='../'
+//def opencvsdk='/<path to OpenCV-android-sdk>'
+//println opencvsdk
+include ':opencv'
+project(':opencv').projectDir = new File(opencvsdk + '/sdk')
+")
+
+
+macro(add_android_project target path)
+  get_filename_component(__dir "${path}" NAME)
+
+  set(OPENCV_ANDROID_CMAKE_EXTRA_ARGS "")
+  if(DEFINED ANDROID_TOOLCHAIN)
+    set(OPENCV_ANDROID_CMAKE_EXTRA_ARGS "${OPENCV_ANDROID_CMAKE_EXTRA_ARGS},\n\"-DANDROID_TOOLCHAIN=${ANDROID_TOOLCHAIN}\"")
+  endif()
+  if(DEFINED ANDROID_STL)
+    set(OPENCV_ANDROID_CMAKE_EXTRA_ARGS "${OPENCV_ANDROID_CMAKE_EXTRA_ARGS},\n\"-DANDROID_STL=${ANDROID_STL}\"")
+  endif()
+
+  #
+  # Build
+  #
+  set(ANDROID_SAMPLE_JNI_PATH "${path}/jni")
+  set(ANDROID_SAMPLE_JAVA_PATH "${path}/src")
+  set(ANDROID_SAMPLE_RES_PATH "${path}/res")
+  set(ANDROID_SAMPLE_MANIFEST_PATH "${path}/gradle/AndroidManifest.xml")
+
+  set(ANDROID_ABI_FILTER "${ANDROID_BUILD_ABI_FILTER}")
+  set(ANDROID_PROJECT_JNI_PATH "../../")
+
+  string(REPLACE ";" "', '" ANDROID_SAMPLE_JAVA_PATH "['${ANDROID_SAMPLE_JAVA_PATH}']")
+  string(REPLACE ";" "', '" ANDROID_SAMPLE_RES_PATH "['${ANDROID_SAMPLE_RES_PATH}']")
+  configure_file("${path}/build.gradle.in" "${ANDROID_BUILD_BASE_DIR}/${__dir}/build.gradle" @ONLY)
+
+  file(APPEND "${ANDROID_BUILD_BASE_DIR}/settings.gradle" "
+include ':${__dir}'
+")
+
+  # build apk
+  set(APK_FILE "${ANDROID_BUILD_BASE_DIR}/${__dir}/build/outputs/apk/release/${__dir}-${ANDROID_ABI}-release-unsigned.apk")
+  ocv_update(OPENCV_GRADLE_VERBOSE_OPTIONS "-i")
+  add_custom_command(
+      OUTPUT "${APK_FILE}" "${OPENCV_DEPHELPER}/android_sample_${__dir}"
+      COMMAND ./gradlew ${OPENCV_GRADLE_VERBOSE_OPTIONS} "${__dir}:assemble"
+      COMMAND ${CMAKE_COMMAND} -E touch "${OPENCV_DEPHELPER}/android_sample_${__dir}"
+      WORKING_DIRECTORY "${ANDROID_BUILD_BASE_DIR}"
+      DEPENDS ${depends} opencv_java_android
+      COMMENT "Building OpenCV Android sample project: ${__dir}"
+  )
+  file(REMOVE "${OPENCV_DEPHELPER}/android_sample_${__dir}")  # force rebuild after CMake run
+
+  add_custom_target(android_sample_${__dir} ALL DEPENDS "${OPENCV_DEPHELPER}/android_sample_${__dir}" SOURCES "${ANDROID_SAMPLE_MANIFEST_PATH}")
+
+  #
+  # Install
+  #
+  set(ANDROID_SAMPLE_JNI_PATH "jni")
+  set(ANDROID_SAMPLE_JAVA_PATH "src")
+  set(ANDROID_SAMPLE_RES_PATH "res")
+  set(ANDROID_SAMPLE_MANIFEST_PATH "AndroidManifest.xml")
+
+  install(DIRECTORY "${path}/res" DESTINATION "${ANDROID_INSTALL_SAMPLES_DIR}/${__dir}" COMPONENT samples OPTIONAL)
+  install(DIRECTORY "${path}/src" DESTINATION "${ANDROID_INSTALL_SAMPLES_DIR}/${__dir}" COMPONENT samples)
+  install(DIRECTORY "${path}/jni" DESTINATION "${ANDROID_INSTALL_SAMPLES_DIR}/${__dir}" COMPONENT samples OPTIONAL)
+
+  install(FILES "${path}/gradle/AndroidManifest.xml" DESTINATION "${ANDROID_INSTALL_SAMPLES_DIR}/${__dir}" COMPONENT samples)
+
+  set(ANDROID_ABI_FILTER "${ANDROID_INSTALL_ABI_FILTER}")
+  set(ANDROID_PROJECT_JNI_PATH "native/jni")
+
+  string(REPLACE ";" "', '" ANDROID_SAMPLE_JAVA_PATH "['${ANDROID_SAMPLE_JAVA_PATH}']")
+  string(REPLACE ";" "', '" ANDROID_SAMPLE_RES_PATH "['${ANDROID_SAMPLE_RES_PATH}']")
+  configure_file("${path}/build.gradle.in" "${ANDROID_TMP_INSTALL_BASE_DIR}/${__dir}/build.gradle" @ONLY)
+  install(FILES "${ANDROID_TMP_INSTALL_BASE_DIR}/${__dir}/build.gradle" DESTINATION "${ANDROID_INSTALL_SAMPLES_DIR}/${__dir}" COMPONENT samples)
+
+  file(APPEND "${ANDROID_TMP_INSTALL_BASE_DIR}/settings.gradle" "
+include ':${__dir}'
+")
+
+endmacro()
+
+install(FILES "${ANDROID_TMP_INSTALL_BASE_DIR}/settings.gradle" DESTINATION "${ANDROID_INSTALL_SAMPLES_DIR}" COMPONENT samples)
@@ -18,7 +18,7 @@ int main()
        archs.push_back(arch.str());
        arch.str("");
    }
-    archs.unique(); #Some devices might have the same arch
+    archs.unique(); // Some devices might have the same arch
    for (std::list<std::string>::iterator it=archs.begin(); it!=archs.end(); ++it)
        std::cout << *it << " ";
    return 0;
@@ -1,8 +1,12 @@
-# if defined(__VSX__)
-#   include <altivec.h>
-# else
-#   error "VSX is not supported"
-# endif
+#if defined(__VSX__)
+    #if defined(__PPC64__) && defined(__LITTLE_ENDIAN__)
+        #include <altivec.h>
+    #else
+        #error "OpenCV only supports little-endian mode"
+    #endif
+#else
+    #error "VSX is not supported"
+#endif

 int main()
 {
@@ -0,0 +1,17 @@
+#if defined(__VSX__)
+    #if defined(__PPC64__) && defined(__LITTLE_ENDIAN__)
+        #include <altivec.h>
+    #else
+        #error "OpenCV only supports little-endian mode"
+    #endif
+#else
+    #error "VSX3 is not supported"
+#endif
+
+int main()
+{
+    __vector unsigned char a = vec_splats((unsigned char)1);
+    __vector unsigned char b = vec_splats((unsigned char)2);
+    __vector unsigned char r = vec_absd(a, b);
+    return 0;
+}
@@ -28,7 +28,9 @@
    opencv/cxeigen.hpp
    opencv2/core/eigen.hpp
    opencv2/flann/hdf5.h
+    opencv2/imgcodecs/imgcodecs_c.h
    opencv2/imgcodecs/ios.h
+    opencv2/videoio/videoio_c.h
    opencv2/videoio/cap_ios.h
    opencv2/xobjdetect/private.hpp
    @OPENCV_ABI_SKIP_HEADERS@
@@ -39,6 +41,7 @@
 </skip_libs>

 <gcc_options>
+ -std=c++11
 @OPENCV_ABI_GCC_OPTIONS@
 </gcc_options>

@@ -85,10 +85,8 @@ we will later have to clip the data in order to avoid overflow.

@code{.py}
 # Tonemap HDR image
-tonemap1 = cv.createTonemapDurand(gamma=2.2)
+tonemap1 = cv.createTonemap(gamma=2.2)
 res_debevec = tonemap1.process(hdr_debevec.copy())
-tonemap2 = cv.createTonemapDurand(gamma=1.3)
-res_robertson = tonemap2.process(hdr_robertson.copy())
@endcode

 ### 4. Merge exposures using Mertens fusion
@@ -173,5 +171,5 @@ Additional Resources

 Exercises
 ---------
-1. Try all tonemap algorithms: cv::TonemapDrago, cv::TonemapDurand, cv::TonemapMantiuk and cv::TonemapReinhard
+1. Try all tonemap algorithms: cv::TonemapDrago, cv::TonemapMantiuk and cv::TonemapReinhard
 2. Try changing the parameters in the HDR calibration and tonemap methods.
@@ -12,7 +12,7 @@ Tutorial was written for the following versions of corresponding software:

 - Download and install Android Studio from https://developer.android.com/studio.

- Get the latest pre-built OpenCV for Android release from https://github.com/opencv/opencv/releases and unpack it (for example, `opencv-3.4.4-android-sdk.zip`).
+- Get the latest pre-built OpenCV for Android release from https://github.com/opencv/opencv/releases and unpack it (for example, `opencv-3.4.5-android-sdk.zip`).

 - Download MobileNet object detection model from https://github.com/chuanqi305/MobileNet-SSD. We need a configuration file `MobileNetSSD_deploy.prototxt` and weights `MobileNetSSD_deploy.caffemodel`.

@@ -65,7 +65,7 @@ that should be used to find the match.
    -#  **Mask image (M):** The mask, a grayscale image that masks the template


-   Only two matching methods currently accept a mask: CV_TM_SQDIFF and CV_TM_CCORR_NORMED (see
+-   Only two matching methods currently accept a mask: TM_SQDIFF and TM_CCORR_NORMED (see
    below for explanation of all the matching methods available in opencv).


@@ -86,23 +86,23 @@ that should be used to find the match.
 Good question. OpenCV implements Template matching in the function **matchTemplate()**. The
 available methods are 6:

-#  **method=CV_TM_SQDIFF**
+-#  **method=TM_SQDIFF**

    \f[R(x,y)= \sum _{x',y'} (T(x',y')-I(x+x',y+y'))^2\f]

-#  **method=CV_TM_SQDIFF_NORMED**
+-#  **method=TM_SQDIFF_NORMED**

    \f[R(x,y)= \frac{\sum_{x',y'} (T(x',y')-I(x+x',y+y'))^2}{\sqrt{\sum_{x',y'}T(x',y')^2 \cdot \sum_{x',y'} I(x+x',y+y')^2}}\f]

-#  **method=CV_TM_CCORR**
+-#  **method=TM_CCORR**

    \f[R(x,y)= \sum _{x',y'} (T(x',y')  \cdot I(x+x',y+y'))\f]

-#  **method=CV_TM_CCORR_NORMED**
+-#  **method=TM_CCORR_NORMED**

    \f[R(x,y)= \frac{\sum_{x',y'} (T(x',y') \cdot I(x+x',y+y'))}{\sqrt{\sum_{x',y'}T(x',y')^2 \cdot \sum_{x',y'} I(x+x',y+y')^2}}\f]

-#  **method=CV_TM_CCOEFF**
+-#  **method=TM_CCOEFF**

    \f[R(x,y)= \sum _{x',y'} (T'(x',y')  \cdot I'(x+x',y+y'))\f]

@@ -110,7 +110,7 @@ available methods are 6:

    \f[\begin{array}{l} T'(x',y')=T(x',y') - 1/(w  \cdot h)  \cdot \sum _{x'',y''} T(x'',y'') \\ I'(x+x',y+y')=I(x+x',y+y') - 1/(w  \cdot h)  \cdot \sum _{x'',y''} I(x+x'',y+y'') \end{array}\f]

-#  **method=CV_TM_CCOEFF_NORMED**
+-#  **method=TM_CCOEFF_NORMED**

    \f[R(x,y)= \frac{ \sum_{x',y'} (T'(x',y') \cdot I'(x+x',y+y')) }{ \sqrt{\sum_{x',y'}T'(x',y')^2 \cdot \sum_{x',y'} I'(x+x',y+y')^2} }\f]

@@ -36,14 +36,14 @@ Open your Doxyfile using your favorite text editor and search for the key
 `TAGFILES`. Change it as follows:

@code
-TAGFILES = ./docs/doxygen-tags/opencv.tag=http://docs.opencv.org/4.0.0
+TAGFILES = ./docs/doxygen-tags/opencv.tag=http://docs.opencv.org/4.0.1
@endcode

 If you had other definitions already, you can append the line using a `\`:

@code
 TAGFILES = ./docs/doxygen-tags/libstdc++.tag=https://gcc.gnu.org/onlinedocs/libstdc++/latest-doxygen \
-           ./docs/doxygen-tags/opencv.tag=http://docs.opencv.org/4.0.0
+           ./docs/doxygen-tags/opencv.tag=http://docs.opencv.org/4.0.1
@endcode

 Doxygen can now use the information from the tag file to link to the OpenCV
@@ -171,7 +171,7 @@ Now it's time to look at the results. Note that HDR image can't be stored in one
 formats, so we save it to Radiance image (.hdr). Also all HDR imaging functions return results in
 [0, 1] range so we should multiply result by 255.

-You can try other tonemap algorithms: cv::TonemapDrago, cv::TonemapDurand, cv::TonemapMantiuk and cv::TonemapReinhard
+You can try other tonemap algorithms: cv::TonemapDrago, cv::TonemapMantiuk and cv::TonemapReinhard
 You can also adjust the parameters in the HDR calibration and tonemap methods for your own photos.

 Results
@@ -2,7 +2,7 @@ set(the_description "The Core Functionality")

 ocv_add_dispatched_file(mathfuncs_core SSE2 AVX AVX2)
 ocv_add_dispatched_file(stat SSE4_2 AVX2)
-ocv_add_dispatched_file(arithm SSE2 SSE4_1 AVX2)
+ocv_add_dispatched_file(arithm SSE2 SSE4_1 AVX2 VSX3)

 # dispatching for accuracy tests
 ocv_add_dispatched_file_force_all(test_intrin128 TEST SSE2 SSE3 SSSE3 SSE4_1 SSE4_2 AVX FP16 AVX2)
@@ -92,13 +92,51 @@ namespace cv { namespace cuda { namespace device
            return vec.w;
        }

+        //constants for conversion from/to RGB and Gray, YUV, YCrCb according to BT.601
+        const float B2YF = 0.114f;
+        const float G2YF = 0.587f;
+        const float R2YF = 0.299f;
+
+        //to YCbCr
+        const float YCBF = 0.564f; // == 1/2/(1-B2YF)
+        const float YCRF = 0.713f; // == 1/2/(1-R2YF)
+        const int YCBI = 9241;  // == YCBF*16384
+        const int YCRI = 11682; // == YCRF*16384
+        //to YUV
+        const float B2UF = 0.492f;
+        const float R2VF = 0.877f;
+        const int B2UI = 8061;  // == B2UF*16384
+        const int R2VI = 14369; // == R2VF*16384
+        //from YUV
+        const float U2BF = 2.032f;
+        const float U2GF = -0.395f;
+        const float V2GF = -0.581f;
+        const float V2RF = 1.140f;
+        const int U2BI = 33292;
+        const int U2GI = -6472;
+        const int V2GI = -9519;
+        const int V2RI = 18678;
+        //from YCrCb
+        const float CB2BF = 1.773f;
+        const float CB2GF = -0.344f;
+        const float CR2GF = -0.714f;
+        const float CR2RF = 1.403f;
+        const int CB2BI = 29049;
+        const int CB2GI = -5636;
+        const int CR2GI = -11698;
+        const int CR2RI = 22987;
+
        enum
        {
            yuv_shift  = 14,
            xyz_shift  = 12,
+            gray_shift = 15,
            R2Y        = 4899,
            G2Y        = 9617,
            B2Y        = 1868,
+            RY15 =  9798, // == R2YF*32768 + 0.5
+            GY15 = 19235, // == G2YF*32768 + 0.5
+            BY15 =  3735, // == B2YF*32768 + 0.5
            BLOCK_SIZE = 256
        };
    }
@@ -406,7 +444,7 @@ namespace cv { namespace cuda { namespace device
        {
            static __device__ __forceinline__ uchar cvt(uint t)
            {
-                return (uchar)CV_DESCALE(((t << 3) & 0xf8) * B2Y + ((t >> 3) & 0xfc) * G2Y + ((t >> 8) & 0xf8) * R2Y, yuv_shift);
+                return (uchar)CV_DESCALE(((t << 3) & 0xf8) * BY15 + ((t >> 3) & 0xfc) * GY15 + ((t >> 8) & 0xf8) * RY15, gray_shift);
            }
        };

@@ -414,7 +452,7 @@ namespace cv { namespace cuda { namespace device
        {
            static __device__ __forceinline__ uchar cvt(uint t)
            {
-                return (uchar)CV_DESCALE(((t << 3) & 0xf8) * B2Y + ((t >> 2) & 0xf8) * G2Y + ((t >> 7) & 0xf8) * R2Y, yuv_shift);
+                return (uchar)CV_DESCALE(((t << 3) & 0xf8) * BY15 + ((t >> 2) & 0xf8) * GY15 + ((t >> 7) & 0xf8) * RY15, gray_shift);
            }
        };

@@ -443,7 +481,7 @@ namespace cv { namespace cuda { namespace device
    {
        template <int bidx, typename T> static __device__ __forceinline__ T RGB2GrayConvert(const T* src)
        {
-            return (T)CV_DESCALE((unsigned)(src[bidx] * B2Y + src[1] * G2Y + src[bidx^2] * R2Y), yuv_shift);
+            return (T)CV_DESCALE((unsigned)(src[bidx] * BY15 + src[1] * GY15 + src[bidx^2] * RY15), gray_shift);
        }

        template <int bidx> static __device__ __forceinline__ uchar RGB2GrayConvert(uint src)
@@ -451,12 +489,12 @@ namespace cv { namespace cuda { namespace device
            uint b = 0xffu & (src >> (bidx * 8));
            uint g = 0xffu & (src >> 8);
            uint r = 0xffu & (src >> ((bidx ^ 2) * 8));
-            return CV_DESCALE((uint)(b * B2Y + g * G2Y + r * R2Y), yuv_shift);
+            return CV_DESCALE((uint)(b * BY15 + g * GY15 + r * RY15), gray_shift);
        }

        template <int bidx> static __device__ __forceinline__ float RGB2GrayConvert(const float* src)
        {
-            return src[bidx] * 0.114f + src[1] * 0.587f + src[bidx^2] * 0.299f;
+            return src[bidx] * B2YF + src[1] * G2YF + src[bidx^2] * R2YF;
        }

        template <typename T, int scn, int bidx> struct RGB2Gray : unary_function<typename TypeVec<T, scn>::vec_type, T>
@@ -494,8 +532,8 @@ namespace cv { namespace cuda { namespace device

    namespace color_detail
    {
-        __constant__ float c_RGB2YUVCoeffs_f[5] = { 0.114f, 0.587f, 0.299f, 0.492f, 0.877f };
-        __constant__ int   c_RGB2YUVCoeffs_i[5] = { B2Y, G2Y, R2Y, 8061, 14369 };
+        __constant__ float c_RGB2YUVCoeffs_f[5] = { B2YF, G2YF, R2YF, B2UF, R2VF };
+        __constant__ int   c_RGB2YUVCoeffs_i[5] = { B2Y, G2Y, R2Y, B2UI, R2VI };

        template <int bidx, typename T, typename D> static __device__ void RGB2YUVConvert(const T* src, D& dst)
        {
@@ -543,8 +581,8 @@ namespace cv { namespace cuda { namespace device

    namespace color_detail
    {
-        __constant__ float c_YUV2RGBCoeffs_f[5] = { 2.032f, -0.395f, -0.581f, 1.140f };
-        __constant__ int   c_YUV2RGBCoeffs_i[5] = { 33292, -6472, -9519, 18678 };
+        __constant__ float c_YUV2RGBCoeffs_f[5] = { U2BF, U2GF, V2GF, V2RF };
+        __constant__ int   c_YUV2RGBCoeffs_i[5] = { U2BI, U2GI, V2GI, V2RI };

        template <int bidx, typename T, typename D> static __device__ void YUV2RGBConvert(const T& src, D* dst)
        {
@@ -633,8 +671,8 @@ namespace cv { namespace cuda { namespace device

    namespace color_detail
    {
-        __constant__ float c_RGB2YCrCbCoeffs_f[5] = {0.299f, 0.587f, 0.114f, 0.713f, 0.564f};
-        __constant__ int   c_RGB2YCrCbCoeffs_i[5] = {R2Y, G2Y, B2Y, 11682, 9241};
+        __constant__ float c_RGB2YCrCbCoeffs_f[5] = {R2YF, G2YF, B2YF, YCRF, YCBF};
+        __constant__ int   c_RGB2YCrCbCoeffs_i[5] = {R2Y, G2Y, B2Y, YCRI, YCBI};

        template <int bidx, typename T, typename D> static __device__ void RGB2YCrCbConvert(const T* src, D& dst)
        {
@@ -710,8 +748,8 @@ namespace cv { namespace cuda { namespace device

    namespace color_detail
    {
-        __constant__ float c_YCrCb2RGBCoeffs_f[5] = {1.403f, -0.714f, -0.344f, 1.773f};
-        __constant__ int   c_YCrCb2RGBCoeffs_i[5] = {22987, -11698, -5636, 29049};
+        __constant__ float c_YCrCb2RGBCoeffs_f[5] = {CR2RF, CR2GF, CB2GF, CB2BF};
+        __constant__ int   c_YCrCb2RGBCoeffs_i[5] = {CR2RI, CR2GI, CB2GI, CB2BI};

        template <int bidx, typename T, typename D> static __device__ void YCrCb2RGBConvert(const T& src, D* dst)
        {
@@ -107,7 +107,7 @@
 #  include <arm_neon.h>
 #endif

-#if defined(__VSX__) && defined(__PPC64__) && defined(__LITTLE_ENDIAN__)
+#ifdef CV_CPU_COMPILE_VSX
 #  include <altivec.h>
 #  undef vector
 #  undef pixel
@@ -115,6 +115,10 @@
 #  define CV_VSX 1
 #endif

+#ifdef CV_CPU_COMPILE_VSX3
+#  define CV_VSX3 1
+#endif
+
 #endif // CV_ENABLE_INTRINSICS && !CV_DISABLE_OPTIMIZATION && !__CUDACC__

 #if defined CV_CPU_COMPILE_AVX && !defined CV_CPU_BASELINE_COMPILE_AVX
@@ -237,3 +241,7 @@ struct VZeroUpperGuard {
 #ifndef CV_VSX
 #  define CV_VSX 0
 #endif
+
+#ifndef CV_VSX3
+#  define CV_VSX3 0
+#endif
@@ -315,5 +315,26 @@
 #endif
 #define __CV_CPU_DISPATCH_CHAIN_VSX(fn, args, mode, ...)  CV_CPU_CALL_VSX(fn, args); __CV_EXPAND(__CV_CPU_DISPATCH_CHAIN_ ## mode(fn, args, __VA_ARGS__))

+#if !defined CV_DISABLE_OPTIMIZATION && defined CV_ENABLE_INTRINSICS && defined CV_CPU_COMPILE_VSX3
+#  define CV_TRY_VSX3 1
+#  define CV_CPU_FORCE_VSX3 1
+#  define CV_CPU_HAS_SUPPORT_VSX3 1
+#  define CV_CPU_CALL_VSX3(fn, args) return (cpu_baseline::fn args)
+#  define CV_CPU_CALL_VSX3_(fn, args) return (opt_VSX3::fn args)
+#elif !defined CV_DISABLE_OPTIMIZATION && defined CV_ENABLE_INTRINSICS && defined CV_CPU_DISPATCH_COMPILE_VSX3
+#  define CV_TRY_VSX3 1
+#  define CV_CPU_FORCE_VSX3 0
+#  define CV_CPU_HAS_SUPPORT_VSX3 (cv::checkHardwareSupport(CV_CPU_VSX3))
+#  define CV_CPU_CALL_VSX3(fn, args) if (CV_CPU_HAS_SUPPORT_VSX3) return (opt_VSX3::fn args)
+#  define CV_CPU_CALL_VSX3_(fn, args) if (CV_CPU_HAS_SUPPORT_VSX3) return (opt_VSX3::fn args)
+#else
+#  define CV_TRY_VSX3 0
+#  define CV_CPU_FORCE_VSX3 0
+#  define CV_CPU_HAS_SUPPORT_VSX3 0
+#  define CV_CPU_CALL_VSX3(fn, args)
+#  define CV_CPU_CALL_VSX3_(fn, args)
+#endif
+#define __CV_CPU_DISPATCH_CHAIN_VSX3(fn, args, mode, ...)  CV_CPU_CALL_VSX3(fn, args); __CV_EXPAND(__CV_CPU_DISPATCH_CHAIN_ ## mode(fn, args, __VA_ARGS__))
+
 #define CV_CPU_CALL_BASELINE(fn, args) return (cpu_baseline::fn args)
 #define __CV_CPU_DISPATCH_CHAIN_BASELINE(fn, args, mode, ...)  CV_CPU_CALL_BASELINE(fn, args) /* last in sequence */
@@ -226,9 +226,10 @@ namespace cv { namespace debug_build_guard { } using namespace debug_build_guard
 #define CV_CPU_AVX_512VBMI      20
 #define CV_CPU_AVX_512VL        21

-#define CV_CPU_NEON   100
+#define CV_CPU_NEON             100

-#define CV_CPU_VSX 200
+#define CV_CPU_VSX              200
+#define CV_CPU_VSX3             201

 // CPU features groups
 #define CV_CPU_AVX512_SKX       256
@@ -266,6 +267,7 @@ enum CpuFeatures {
    CPU_NEON            = 100,

    CPU_VSX             = 200,
+    CPU_VSX3            = 201,

    CPU_AVX512_SKX      = 256, //!< Skylake-X with AVX-512F/CD/BW/DQ/VL

@@ -60,7 +60,7 @@ namespace cv
 //! @{

 template<typename _Tp, int _rows, int _cols, int _options, int _maxRows, int _maxCols> static inline
-void eigen2cv( const Eigen::Matrix<_Tp, _rows, _cols, _options, _maxRows, _maxCols>& src, Mat& dst )
+void eigen2cv( const Eigen::Matrix<_Tp, _rows, _cols, _options, _maxRows, _maxCols>& src, OutputArray dst )
 {
    if( !(src.Flags & Eigen::RowMajorBit) )
    {
@@ -905,6 +905,11 @@ OPENCV_HAL_IMPL_AVX_CMP_OP_64BIT(v_int64x4)
 OPENCV_HAL_IMPL_AVX_CMP_OP_FLT(v_float32x8, ps)
 OPENCV_HAL_IMPL_AVX_CMP_OP_FLT(v_float64x4, pd)

+inline v_float32x8 v_not_nan(const v_float32x8& a)
+{ return v_float32x8(_mm256_cmp_ps(a.val, a.val, _CMP_ORD_Q)); }
+inline v_float64x4 v_not_nan(const v_float64x4& a)
+{ return v_float64x4(_mm256_cmp_pd(a.val, a.val, _CMP_ORD_Q)); }
+
 /** min/max **/
 OPENCV_HAL_IMPL_AVX_BIN_FUNC(v_min, v_uint8x32,  _mm256_min_epu8)
 OPENCV_HAL_IMPL_AVX_BIN_FUNC(v_max, v_uint8x32,  _mm256_max_epu8)
@@ -1120,6 +1125,12 @@ inline float v_reduce_sum(const v_float32x8& a)
    return _mm_cvtss_f32(s1);
 }

+inline double v_reduce_sum(const v_float64x4& a)
+{
+    __m256d s0 = _mm256_hadd_pd(a.val, a.val);
+    return _mm_cvtsd_f64(_mm_add_pd(_v256_extract_low(s0), _v256_extract_high(s0)));
+}
+
 inline v_float32x8 v_reduce_sum4(const v_float32x8& a, const v_float32x8& b,
                                 const v_float32x8& c, const v_float32x8& d)
 {
@@ -1128,6 +1139,41 @@ inline v_float32x8 v_reduce_sum4(const v_float32x8& a, const v_float32x8& b,
    return v_float32x8(_mm256_hadd_ps(ab, cd));
 }

+inline unsigned v_reduce_sad(const v_uint8x32& a, const v_uint8x32& b)
+{
+    return (unsigned)_v_cvtsi256_si32(_mm256_sad_epu8(a.val, b.val));
+}
+inline unsigned v_reduce_sad(const v_int8x32& a, const v_int8x32& b)
+{
+    __m256i half = _mm256_set1_epi8(0x7f);
+    return (unsigned)_v_cvtsi256_si32(_mm256_sad_epu8(_mm256_add_epi8(a.val, half), _mm256_add_epi8(b.val, half)));
+}
+inline unsigned v_reduce_sad(const v_uint16x16& a, const v_uint16x16& b)
+{
+    v_uint32x8 l, h;
+    v_expand(v_add_wrap(a - b, b - a), l, h);
+    return v_reduce_sum(l + h);
+}
+inline unsigned v_reduce_sad(const v_int16x16& a, const v_int16x16& b)
+{
+    v_uint32x8 l, h;
+    v_expand(v_reinterpret_as_u16(v_sub_wrap(v_max(a, b), v_min(a, b))), l, h);
+    return v_reduce_sum(l + h);
+}
+inline unsigned v_reduce_sad(const v_uint32x8& a, const v_uint32x8& b)
+{
+    return v_reduce_sum(v_max(a, b) - v_min(a, b));
+}
+inline unsigned v_reduce_sad(const v_int32x8& a, const v_int32x8& b)
+{
+    v_int32x8 m = a < b;
+    return v_reduce_sum(v_reinterpret_as_u32(((a - b) ^ m) - m));
+}
+inline float v_reduce_sad(const v_float32x8& a, const v_float32x8& b)
+{
+    return v_reduce_sum((a - b) & v_float32x8(_mm256_castsi256_ps(_mm256_set1_epi32(0x7fffffff))));
+}
+
 /** Popcount **/
 #define OPENCV_HAL_IMPL_AVX_POPCOUNT(_Tpvec)                     \
    inline v_uint32x8 v_popcount(const _Tpvec& a)                \
@@ -1232,6 +1278,16 @@ OPENCV_HAL_IMPL_AVX_CHECK_FLT(v_float64x4, 15)
 OPENCV_HAL_IMPL_AVX_MULADD(v_float32x8, ps)
 OPENCV_HAL_IMPL_AVX_MULADD(v_float64x4, pd)

+inline v_int32x8 v_fma(const v_int32x8& a, const v_int32x8& b, const v_int32x8& c)
+{
+    return a * b + c;
+}
+
+inline v_int32x8 v_muladd(const v_int32x8& a, const v_int32x8& b, const v_int32x8& c)
+{
+    return v_fma(a, b, c);
+}
+
 inline v_float32x8 v_invsqrt(const v_float32x8& x)
 {
    v_float32x8 half = x * v256_setall_f32(0.5);
@@ -683,6 +683,25 @@ OPENCV_HAL_IMPL_CMP_OP(==)
 For all types except 64-bit integer values. */
 OPENCV_HAL_IMPL_CMP_OP(!=)

+template<int n>
+inline v_reg<float, n> v_not_nan(const v_reg<float, n>& a)
+{
+    typedef typename V_TypeTraits<float>::int_type itype;
+    v_reg<float, n> c;
+    for (int i = 0; i < n; i++)
+        c.s[i] = V_TypeTraits<float>::reinterpret_from_int((itype)-(int)(a.s[i] == a.s[i]));
+    return c;
+}
+template<int n>
+inline v_reg<double, n> v_not_nan(const v_reg<double, n>& a)
+{
+    typedef typename V_TypeTraits<double>::int_type itype;
+    v_reg<double, n> c;
+    for (int i = 0; i < n; i++)
+        c.s[i] = V_TypeTraits<double>::reinterpret_from_int((itype)-(int)(a.s[i] == a.s[i]));
+    return c;
+}
+
 //! @brief Helper macro
 //! @ingroup core_hal_intrin_impl
 #define OPENCV_HAL_IMPL_ARITHM_OP(func, bin_op, cast_op, _Tp2) \
@@ -1044,6 +1063,21 @@ inline v_float32x4 v_reduce_sum4(const v_float32x4& a, const v_float32x4& b,
    return r;
 }

+/** @brief Sum absolute differences of values
+
+Scheme:
+@code
+{A1 A2 A3 ...} {B1 B2 B3 ...} => sum{ABS(A1-B1),abs(A2-B2),abs(A3-B3),...}
+@endcode
+For all types except 64-bit types.*/
+template<typename _Tp, int n> inline typename V_TypeTraits< typename V_TypeTraits<_Tp>::abs_type >::sum_type v_reduce_sad(const v_reg<_Tp, n>& a, const v_reg<_Tp, n>& b)
+{
+    typename V_TypeTraits< typename V_TypeTraits<_Tp>::abs_type >::sum_type c = _absdiff(a.s[0], b.s[0]);
+    for (int i = 1; i < n; i++)
+        c += _absdiff(a.s[i], b.s[i]);
+    return c;
+}
+
 /** @brief Get negative values mask

 Returned value is a bit mask with bits set to 1 on places corresponding to negative packed values indexes.
@@ -764,6 +764,13 @@ OPENCV_HAL_IMPL_NEON_INT_CMP_OP(v_int64x2, vreinterpretq_s64_u64, s64, u64)
 OPENCV_HAL_IMPL_NEON_INT_CMP_OP(v_float64x2, vreinterpretq_f64_u64, f64, u64)
 #endif

+inline v_float32x4 v_not_nan(const v_float32x4& a)
+{ return v_float32x4(vreinterpretq_f32_u32(vceqq_f32(a.val, a.val))); }
+#if CV_SIMD128_64F
+inline v_float64x2 v_not_nan(const v_float64x2& a)
+{ return v_float64x2(vreinterpretq_f64_u64(vceqq_f64(a.val, a.val))); }
+#endif
+
 OPENCV_HAL_IMPL_NEON_BIN_FUNC(v_uint8x16, v_add_wrap, vaddq_u8)
 OPENCV_HAL_IMPL_NEON_BIN_FUNC(v_int8x16, v_add_wrap, vaddq_s8)
 OPENCV_HAL_IMPL_NEON_BIN_FUNC(v_uint16x8, v_add_wrap, vaddq_u16)
@@ -977,6 +984,13 @@ OPENCV_HAL_IMPL_NEON_REDUCE_OP_4(v_float32x4, float32x2, float, sum, add, f32)
 OPENCV_HAL_IMPL_NEON_REDUCE_OP_4(v_float32x4, float32x2, float, max, max, f32)
 OPENCV_HAL_IMPL_NEON_REDUCE_OP_4(v_float32x4, float32x2, float, min, min, f32)

+#if CV_SIMD128_64F
+inline double v_reduce_sum(const v_float64x2& a)
+{
+    return vgetq_lane_f64(a.val, 0) + vgetq_lane_f64(a.val, 1);
+}
+#endif
+
 inline v_float32x4 v_reduce_sum4(const v_float32x4& a, const v_float32x4& b,
                                 const v_float32x4& c, const v_float32x4& d)
 {
@@ -992,6 +1006,49 @@ inline v_float32x4 v_reduce_sum4(const v_float32x4& a, const v_float32x4& b,
    return v_float32x4(vaddq_f32(v0, v1));
 }

+inline unsigned v_reduce_sad(const v_uint8x16& a, const v_uint8x16& b)
+{
+    uint32x4_t t0 = vpaddlq_u16(vpaddlq_u8(vabdq_u8(a.val, b.val)));
+    uint32x2_t t1 = vpadd_u32(vget_low_u32(t0), vget_high_u32(t0));
+    return vget_lane_u32(vpadd_u32(t1, t1), 0);
+}
+inline unsigned v_reduce_sad(const v_int8x16& a, const v_int8x16& b)
+{
+    uint32x4_t t0 = vpaddlq_u16(vpaddlq_u8(vreinterpretq_u8_s8(vabdq_s8(a.val, b.val))));
+    uint32x2_t t1 = vpadd_u32(vget_low_u32(t0), vget_high_u32(t0));
+    return vget_lane_u32(vpadd_u32(t1, t1), 0);
+}
+inline unsigned v_reduce_sad(const v_uint16x8& a, const v_uint16x8& b)
+{
+    uint32x4_t t0 = vpaddlq_u16(vabdq_u16(a.val, b.val));
+    uint32x2_t t1 = vpadd_u32(vget_low_u32(t0), vget_high_u32(t0));
+    return vget_lane_u32(vpadd_u32(t1, t1), 0);
+}
+inline unsigned v_reduce_sad(const v_int16x8& a, const v_int16x8& b)
+{
+    uint32x4_t t0 = vpaddlq_u16(vreinterpretq_u16_s16(vabdq_s16(a.val, b.val)));
+    uint32x2_t t1 = vpadd_u32(vget_low_u32(t0), vget_high_u32(t0));
+    return vget_lane_u32(vpadd_u32(t1, t1), 0);
+}
+inline unsigned v_reduce_sad(const v_uint32x4& a, const v_uint32x4& b)
+{
+    uint32x4_t t0 = vabdq_u32(a.val, b.val);
+    uint32x2_t t1 = vpadd_u32(vget_low_u32(t0), vget_high_u32(t0));
+    return vget_lane_u32(vpadd_u32(t1, t1), 0);
+}
+inline unsigned v_reduce_sad(const v_int32x4& a, const v_int32x4& b)
+{
+    uint32x4_t t0 = vreinterpretq_u32_s32(vabdq_s32(a.val, b.val));
+    uint32x2_t t1 = vpadd_u32(vget_low_u32(t0), vget_high_u32(t0));
+    return vget_lane_u32(vpadd_u32(t1, t1), 0);
+}
+inline float v_reduce_sad(const v_float32x4& a, const v_float32x4& b)
+{
+    float32x4_t t0 = vabdq_f32(a.val, b.val);
+    float32x2_t t1 = vpadd_f32(vget_low_f32(t0), vget_high_f32(t0));
+    return vget_lane_f32(vpadd_f32(t1, t1), 0);
+}
+
 #define OPENCV_HAL_IMPL_NEON_POPCOUNT(_Tpvec, cast) \
 inline v_uint32x4 v_popcount(const _Tpvec& a) \
 { \
@@ -1041,6 +1041,11 @@ inline _Tpvec operator != (const _Tpvec& a, const _Tpvec& b) \
 OPENCV_HAL_IMPL_SSE_64BIT_CMP_OP(v_uint64x2, v_reinterpret_as_u64)
 OPENCV_HAL_IMPL_SSE_64BIT_CMP_OP(v_int64x2, v_reinterpret_as_s64)

+inline v_float32x4 v_not_nan(const v_float32x4& a)
+{ return v_float32x4(_mm_cmpord_ps(a.val, a.val)); }
+inline v_float64x2 v_not_nan(const v_float64x2& a)
+{ return v_float64x2(_mm_cmpord_pd(a.val, a.val)); }
+
 OPENCV_HAL_IMPL_SSE_BIN_FUNC(v_uint8x16, v_add_wrap, _mm_add_epi8)
 OPENCV_HAL_IMPL_SSE_BIN_FUNC(v_int8x16, v_add_wrap, _mm_add_epi8)
 OPENCV_HAL_IMPL_SSE_BIN_FUNC(v_uint16x8, v_add_wrap, _mm_add_epi16)
@@ -1451,6 +1456,13 @@ OPENCV_HAL_IMPL_SSE_REDUCE_OP_4_SUM(v_uint32x4, unsigned, __m128i, epi32, OPENCV
 OPENCV_HAL_IMPL_SSE_REDUCE_OP_4_SUM(v_int32x4, int, __m128i, epi32, OPENCV_HAL_NOP, OPENCV_HAL_NOP, si128_si32)
 OPENCV_HAL_IMPL_SSE_REDUCE_OP_4_SUM(v_float32x4, float, __m128, ps, _mm_castps_si128, _mm_castsi128_ps, ss_f32)

+inline double v_reduce_sum(const v_float64x2& a)
+{
+    double CV_DECL_ALIGNED(32) idx[2];
+    v_store_aligned(idx, a);
+    return idx[0] + idx[1];
+}
+
 inline v_float32x4 v_reduce_sum4(const v_float32x4& a, const v_float32x4& b,
                                 const v_float32x4& c, const v_float32x4& d)
 {
@@ -1472,6 +1484,41 @@ OPENCV_HAL_IMPL_SSE_REDUCE_OP_4(v_int32x4, int, min, std::min)
 OPENCV_HAL_IMPL_SSE_REDUCE_OP_4(v_float32x4, float, max, std::max)
 OPENCV_HAL_IMPL_SSE_REDUCE_OP_4(v_float32x4, float, min, std::min)

+inline unsigned v_reduce_sad(const v_uint8x16& a, const v_uint8x16& b)
+{
+    return (unsigned)_mm_cvtsi128_si32(_mm_sad_epu8(a.val, b.val));
+}
+inline unsigned v_reduce_sad(const v_int8x16& a, const v_int8x16& b)
+{
+    __m128i half = _mm_set1_epi8(0x7f);
+    return (unsigned)_mm_cvtsi128_si32(_mm_sad_epu8(_mm_add_epi8(a.val, half),
+                                                    _mm_add_epi8(b.val, half)));
+}
+inline unsigned v_reduce_sad(const v_uint16x8& a, const v_uint16x8& b)
+{
+    v_uint32x4 l, h;
+    v_expand(v_absdiff(a, b), l, h);
+    return v_reduce_sum(l + h);
+}
+inline unsigned v_reduce_sad(const v_int16x8& a, const v_int16x8& b)
+{
+    v_uint32x4 l, h;
+    v_expand(v_absdiff(a, b), l, h);
+    return v_reduce_sum(l + h);
+}
+inline unsigned v_reduce_sad(const v_uint32x4& a, const v_uint32x4& b)
+{
+    return v_reduce_sum(v_absdiff(a, b));
+}
+inline unsigned v_reduce_sad(const v_int32x4& a, const v_int32x4& b)
+{
+    return v_reduce_sum(v_absdiff(a, b));
+}
+inline float v_reduce_sad(const v_float32x4& a, const v_float32x4& b)
+{
+    return v_reduce_sum(v_absdiff(a, b));
+}
+
 #define OPENCV_HAL_IMPL_SSE_POPCOUNT(_Tpvec) \
 inline v_uint32x4 v_popcount(const _Tpvec& a) \
 { \
@@ -1925,13 +1972,11 @@ inline void v_load_deinterleave(const unsigned* ptr, v_uint32x4& a, v_uint32x4&

 inline void v_load_deinterleave(const float* ptr, v_float32x4& a, v_float32x4& b)
 {
-    const int mask_lo = _MM_SHUFFLE(2, 0, 2, 0), mask_hi = _MM_SHUFFLE(3, 1, 3, 1);
-
    __m128 u0 = _mm_loadu_ps(ptr);       // a0 b0 a1 b1
    __m128 u1 = _mm_loadu_ps((ptr + 4)); // a2 b2 a3 b3

-    a.val = _mm_shuffle_ps(u0, u1, mask_lo); // a0 a1 a2 a3
-    b.val = _mm_shuffle_ps(u0, u1, mask_hi); // b0 b1 ab b3
+    a.val = _mm_shuffle_ps(u0, u1, _MM_SHUFFLE(2, 0, 2, 0)); // a0 a1 a2 a3
+    b.val = _mm_shuffle_ps(u0, u1, _MM_SHUFFLE(3, 1, 3, 1)); // b0 b1 ab b3
 }

 inline void v_load_deinterleave(const float* ptr, v_float32x4& a, v_float32x4& b, v_float32x4& c)
@@ -607,6 +607,11 @@ OPENCV_HAL_IMPL_VSX_INT_CMP_OP(v_float64x2)
 OPENCV_HAL_IMPL_VSX_INT_CMP_OP(v_uint64x2)
 OPENCV_HAL_IMPL_VSX_INT_CMP_OP(v_int64x2)

+inline v_float32x4 v_not_nan(const v_float32x4& a)
+{ return v_float32x4(vec_cmpeq(a.val, a.val)); }
+inline v_float64x2 v_not_nan(const v_float64x2& a)
+{ return v_float64x2(vec_cmpeq(a.val, a.val)); }
+
 /** min/max **/
 OPENCV_HAL_IMPL_VSX_BIN_FUNC(v_min, vec_min)
 OPENCV_HAL_IMPL_VSX_BIN_FUNC(v_max, vec_max)
@@ -711,6 +716,11 @@ OPENCV_HAL_IMPL_VSX_REDUCE_OP_4(v_float32x4, vec_float4, float, sum, vec_add)
 OPENCV_HAL_IMPL_VSX_REDUCE_OP_4(v_float32x4, vec_float4, float, max, vec_max)
 OPENCV_HAL_IMPL_VSX_REDUCE_OP_4(v_float32x4, vec_float4, float, min, vec_min)

+inline double v_reduce_sum(const v_float64x2& a)
+{
+    return vec_extract(vec_add(a.val, vec_permi(a.val, a.val, 3)), 0);
+}
+
 #define OPENCV_HAL_IMPL_VSX_REDUCE_OP_8(_Tpvec, _Tpvec2, scalartype, suffix, func) \
 inline scalartype v_reduce_##suffix(const _Tpvec& a)                               \
 {                                                                                  \
@@ -734,6 +744,50 @@ inline v_float32x4 v_reduce_sum4(const v_float32x4& a, const v_float32x4& b,
    return v_float32x4(vec_mergeh(ac, bd));
 }

+inline unsigned v_reduce_sad(const v_uint8x16& a, const v_uint8x16& b)
+{
+    const vec_uint4 zero4 = vec_uint4_z;
+    vec_uint4 sum4 = vec_sum4s(vec_absd(a.val, b.val), zero4);
+    return (unsigned)vec_extract(vec_sums(vec_int4_c(sum4), vec_int4_c(zero4)), 3);
+}
+inline unsigned v_reduce_sad(const v_int8x16& a, const v_int8x16& b)
+{
+    const vec_int4 zero4 = vec_int4_z;
+    vec_char16 ad = vec_abss(vec_subs(a.val, b.val));
+    vec_int4 sum4 = vec_sum4s(ad, zero4);
+    return (unsigned)vec_extract(vec_sums(sum4, zero4), 3);
+}
+inline unsigned v_reduce_sad(const v_uint16x8& a, const v_uint16x8& b)
+{
+    vec_ushort8 ad = vec_absd(a.val, b.val);
+    VSX_UNUSED(vec_int4) sum = vec_sums(vec_int4_c(vec_unpackhu(ad)), vec_int4_c(vec_unpacklu(ad)));
+    return (unsigned)vec_extract(sum, 3);
+}
+inline unsigned v_reduce_sad(const v_int16x8& a, const v_int16x8& b)
+{
+    const vec_int4 zero4 = vec_int4_z;
+    vec_short8 ad = vec_abss(vec_subs(a.val, b.val));
+    vec_int4 sum4 = vec_sum4s(ad, zero4);
+    return (unsigned)vec_extract(vec_sums(sum4, zero4), 3);
+}
+inline unsigned v_reduce_sad(const v_uint32x4& a, const v_uint32x4& b)
+{
+    const vec_uint4 ad = vec_absd(a.val, b.val);
+    const vec_uint4 rd = vec_add(ad, vec_sld(ad, ad, 8));
+    return vec_extract(vec_add(rd, vec_sld(rd, rd, 4)), 0);
+}
+inline unsigned v_reduce_sad(const v_int32x4& a, const v_int32x4& b)
+{
+    vec_int4 ad = vec_abss(vec_sub(a.val, b.val));
+    return (unsigned)vec_extract(vec_sums(ad, vec_int4_z), 3);
+}
+inline float v_reduce_sad(const v_float32x4& a, const v_float32x4& b)
+{
+    const vec_float4 ad = vec_abs(vec_sub(a.val, b.val));
+    const vec_float4 rd = vec_add(ad, vec_sld(ad, ad, 8));
+    return vec_extract(vec_add(rd, vec_sld(rd, rd, 4)), 0);
+}
+
 /** Popcount **/
 template<typename _Tpvec>
 inline v_uint32x4 v_popcount(const _Tpvec& a)
@@ -565,7 +565,7 @@ public:
    //! returns the node content as double
    operator double() const;
    //! returns the node content as text string
-    operator std::string() const;
+    inline operator std::string() const { return this->string(); }

    static bool isMap(int flags);
    static bool isSeq(int flags);
@@ -599,7 +599,7 @@ public:
    //! Simplified reading API to use with bindings.
    CV_WRAP double real() const;
    //! Simplified reading API to use with bindings.
-    CV_WRAP String string() const;
+    CV_WRAP std::string string() const;
    //! Simplified reading API to use with bindings.
    CV_WRAP Mat mat() const;

@@ -567,7 +567,7 @@ inline void _mm_deinterleave_ps(__m128 & v_r0, __m128 & v_r1, __m128 & v_g0, __m

 inline void _mm_interleave_ps(__m128 & v_r0, __m128 & v_r1, __m128 & v_g0, __m128 & v_g1)
 {
-    const int mask_lo = _MM_SHUFFLE(2, 0, 2, 0), mask_hi = _MM_SHUFFLE(3, 1, 3, 1);
+    enum { mask_lo = _MM_SHUFFLE(2, 0, 2, 0), mask_hi = _MM_SHUFFLE(3, 1, 3, 1) };

    __m128 layer2_chunk0 = _mm_shuffle_ps(v_r0, v_r1, mask_lo);
    __m128 layer2_chunk2 = _mm_shuffle_ps(v_r0, v_r1, mask_hi);
@@ -588,7 +588,7 @@ inline void _mm_interleave_ps(__m128 & v_r0, __m128 & v_r1, __m128 & v_g0, __m12
 inline void _mm_interleave_ps(__m128 & v_r0, __m128 & v_r1, __m128 & v_g0,
                              __m128 & v_g1, __m128 & v_b0, __m128 & v_b1)
 {
-    const int mask_lo = _MM_SHUFFLE(2, 0, 2, 0), mask_hi = _MM_SHUFFLE(3, 1, 3, 1);
+    enum { mask_lo = _MM_SHUFFLE(2, 0, 2, 0), mask_hi = _MM_SHUFFLE(3, 1, 3, 1) };

    __m128 layer2_chunk0 = _mm_shuffle_ps(v_r0, v_r1, mask_lo);
    __m128 layer2_chunk3 = _mm_shuffle_ps(v_r0, v_r1, mask_hi);
@@ -615,7 +615,7 @@ inline void _mm_interleave_ps(__m128 & v_r0, __m128 & v_r1, __m128 & v_g0,
 inline void _mm_interleave_ps(__m128 & v_r0, __m128 & v_r1, __m128 & v_g0, __m128 & v_g1,
                              __m128 & v_b0, __m128 & v_b1, __m128 & v_a0, __m128 & v_a1)
 {
-    const int mask_lo = _MM_SHUFFLE(2, 0, 2, 0), mask_hi = _MM_SHUFFLE(3, 1, 3, 1);
+    enum { mask_lo = _MM_SHUFFLE(2, 0, 2, 0), mask_hi = _MM_SHUFFLE(3, 1, 3, 1) };

    __m128 layer2_chunk0 = _mm_shuffle_ps(v_r0, v_r1, mask_lo);
    __m128 layer2_chunk4 = _mm_shuffle_ps(v_r0, v_r1, mask_hi);
@@ -1941,8 +1941,11 @@ Rect_<_Tp>& operator += ( Rect_<_Tp>& a, const Size_<_Tp>& b )
 template<typename _Tp> static inline
 Rect_<_Tp>& operator -= ( Rect_<_Tp>& a, const Size_<_Tp>& b )
 {
-    a.width -= b.width;
-    a.height -= b.height;
+    const _Tp width = a.width - b.width;
+    const _Tp height = a.height - b.height;
+    CV_DbgAssert(width >= 0 && height >= 0);
+    a.width = width;
+    a.height = height;
    return a;
 }

@@ -2007,6 +2010,15 @@ Rect_<_Tp> operator + (const Rect_<_Tp>& a, const Size_<_Tp>& b)
    return Rect_<_Tp>( a.x, a.y, a.width + b.width, a.height + b.height );
 }

+template<typename _Tp> static inline
+Rect_<_Tp> operator - (const Rect_<_Tp>& a, const Size_<_Tp>& b)
+{
+    const _Tp width = a.width - b.width;
+    const _Tp height = a.height - b.height;
+    CV_DbgAssert(width >= 0 && height >= 0);
+    return Rect_<_Tp>( a.x, a.y, width, height );
+}
+
 template<typename _Tp> static inline
 Rect_<_Tp> operator & (const Rect_<_Tp>& a, const Rect_<_Tp>& b)
 {
@@ -364,7 +364,7 @@ IplImage;

 CV_INLINE IplImage cvIplImage()
 {
-#if !defined(CV__ENABLE_C_API_CTORS)
+#if !(defined(CV__ENABLE_C_API_CTORS) && defined(__cplusplus))
    IplImage self = CV_STRUCT_INITIALIZER; self.nSize = sizeof(IplImage); return self;
 #else
    return _IplImage();
@@ -7,7 +7,7 @@

 #define CV_VERSION_MAJOR    4
 #define CV_VERSION_MINOR    0
-#define CV_VERSION_REVISION 0
+#define CV_VERSION_REVISION 1
 #define CV_VERSION_STATUS   ""

 #define CVAUX_STR_EXP(__A)  #__A
@@ -253,4 +253,53 @@ PERF_TEST_P( Size_MatType, normalize_minmax, TYPICAL_MATS )
    SANITY_CHECK(dst, 1e-6, ERROR_RELATIVE);
 }

+typedef TestBaseWithParam< int > test_len;
+PERF_TEST_P(test_len, hal_normL1_u8,
+            testing::Values(300000, 2000000)
+           )
+{
+    int len = GetParam();
+
+    Mat src1(1, len, CV_8UC1);
+    Mat src2(1, len, CV_8UC1);
+
+    declare.in(src1, src2, WARMUP_RNG);
+    double n;
+    TEST_CYCLE() n = hal::normL1_(src1.ptr<uchar>(0), src2.ptr<uchar>(0), len);
+    CV_UNUSED(n);
+    SANITY_CHECK_NOTHING();
+}
+
+PERF_TEST_P(test_len, hal_normL1_f32,
+            testing::Values(300000, 2000000)
+           )
+{
+    int len = GetParam();
+
+    Mat src1(1, len, CV_32FC1);
+    Mat src2(1, len, CV_32FC1);
+
+    declare.in(src1, src2, WARMUP_RNG);
+    double n;
+    TEST_CYCLE() n = hal::normL1_(src1.ptr<float>(0), src2.ptr<float>(0), len);
+    CV_UNUSED(n);
+    SANITY_CHECK_NOTHING();
+}
+
+PERF_TEST_P(test_len, hal_normL2Sqr,
+            testing::Values(300000, 2000000)
+           )
+{
+    int len = GetParam();
+
+    Mat src1(1, len, CV_32FC1);
+    Mat src2(1, len, CV_32FC1);
+
+    declare.in(src1, src2, WARMUP_RNG);
+    double n;
+    TEST_CYCLE() n = hal::normL2Sqr_(src1.ptr<float>(0), src2.ptr<float>(0), len);
+    CV_UNUSED(n);
+    SANITY_CHECK_NOTHING();
+}
+
 } // namespace
@@ -1333,7 +1333,7 @@ struct InRange_SIMD
    }
 };

-#if CV_SIMD128
+#if CV_SIMD

 template <>
 struct InRange_SIMD<uchar>
@@ -1342,16 +1342,17 @@ struct InRange_SIMD<uchar>
        uchar * dst, int len) const
    {
        int x = 0;
-        const int width = v_uint8x16::nlanes;
+        const int width = v_uint8::nlanes;

        for (; x <= len - width; x += width)
        {
-            v_uint8x16 values = v_load(src1 + x);
-            v_uint8x16 low = v_load(src2 + x);
-            v_uint8x16 high = v_load(src3 + x);
+            v_uint8 values = vx_load(src1 + x);
+            v_uint8 low = vx_load(src2 + x);
+            v_uint8 high = vx_load(src3 + x);

            v_store(dst + x, (values >= low) & (high >= values));
        }
+        vx_cleanup();
        return x;
    }
 };
@@ -1363,16 +1364,17 @@ struct InRange_SIMD<schar>
        uchar * dst, int len) const
    {
        int x = 0;
-        const int width = v_int8x16::nlanes;
+        const int width = v_int8::nlanes;

        for (; x <= len - width; x += width)
        {
-            v_int8x16 values = v_load(src1 + x);
-            v_int8x16 low = v_load(src2 + x);
-            v_int8x16 high = v_load(src3 + x);
+            v_int8 values = vx_load(src1 + x);
+            v_int8 low = vx_load(src2 + x);
+            v_int8 high = vx_load(src3 + x);

            v_store((schar*)(dst + x), (values >= low) & (high >= values));
        }
+        vx_cleanup();
        return x;
    }
 };
@@ -1384,20 +1386,21 @@ struct InRange_SIMD<ushort>
        uchar * dst, int len) const
    {
        int x = 0;
-        const int width = v_uint16x8::nlanes * 2;
+        const int width = v_uint16::nlanes * 2;

        for (; x <= len - width; x += width)
        {
-            v_uint16x8 values1 = v_load(src1 + x);
-            v_uint16x8 low1 = v_load(src2 + x);
-            v_uint16x8 high1 = v_load(src3 + x);
+            v_uint16 values1 = vx_load(src1 + x);
+            v_uint16 low1 = vx_load(src2 + x);
+            v_uint16 high1 = vx_load(src3 + x);

-            v_uint16x8 values2 = v_load(src1 + x + v_uint16x8::nlanes);
-            v_uint16x8 low2 = v_load(src2 + x + v_uint16x8::nlanes);
-            v_uint16x8 high2 = v_load(src3 + x + v_uint16x8::nlanes);
+            v_uint16 values2 = vx_load(src1 + x + v_uint16::nlanes);
+            v_uint16 low2 = vx_load(src2 + x + v_uint16::nlanes);
+            v_uint16 high2 = vx_load(src3 + x + v_uint16::nlanes);

            v_store(dst + x, v_pack((values1 >= low1) & (high1 >= values1), (values2 >= low2) & (high2 >= values2)));
        }
+        vx_cleanup();
        return x;
    }
 };
@@ -1409,20 +1412,21 @@ struct InRange_SIMD<short>
        uchar * dst, int len) const
    {
        int x = 0;
-        const int width = (int)v_int16x8::nlanes * 2;
+        const int width = (int)v_int16::nlanes * 2;

        for (; x <= len - width; x += width)
        {
-            v_int16x8 values1 = v_load(src1 + x);
-            v_int16x8 low1 = v_load(src2 + x);
-            v_int16x8 high1 = v_load(src3 + x);
+            v_int16 values1 = vx_load(src1 + x);
+            v_int16 low1 = vx_load(src2 + x);
+            v_int16 high1 = vx_load(src3 + x);

-            v_int16x8 values2 = v_load(src1 + x + v_int16x8::nlanes);
-            v_int16x8 low2 = v_load(src2 + x + v_int16x8::nlanes);
-            v_int16x8 high2 = v_load(src3 + x + v_int16x8::nlanes);
+            v_int16 values2 = vx_load(src1 + x + v_int16::nlanes);
+            v_int16 low2 = vx_load(src2 + x + v_int16::nlanes);
+            v_int16 high2 = vx_load(src3 + x + v_int16::nlanes);

            v_store((schar*)(dst + x), v_pack((values1 >= low1) & (high1 >= values1), (values2 >= low2) & (high2 >= values2)));
        }
+        vx_cleanup();
        return x;
    }
 };
@@ -1434,20 +1438,21 @@ struct InRange_SIMD<int>
        uchar * dst, int len) const
    {
        int x = 0;
-        const int width = (int)v_int32x4::nlanes * 2;
+        const int width = (int)v_int32::nlanes * 2;

        for (; x <= len - width; x += width)
        {
-            v_int32x4 values1 = v_load(src1 + x);
-            v_int32x4 low1 = v_load(src2 + x);
-            v_int32x4 high1 = v_load(src3 + x);
+            v_int32 values1 = vx_load(src1 + x);
+            v_int32 low1 = vx_load(src2 + x);
+            v_int32 high1 = vx_load(src3 + x);

-            v_int32x4 values2 = v_load(src1 + x + v_int32x4::nlanes);
-            v_int32x4 low2 = v_load(src2 + x + v_int32x4::nlanes);
-            v_int32x4 high2 = v_load(src3 + x + v_int32x4::nlanes);
+            v_int32 values2 = vx_load(src1 + x + v_int32::nlanes);
+            v_int32 low2 = vx_load(src2 + x + v_int32::nlanes);
+            v_int32 high2 = vx_load(src3 + x + v_int32::nlanes);

            v_pack_store(dst + x, v_reinterpret_as_u16(v_pack((values1 >= low1) & (high1 >= values1), (values2 >= low2) & (high2 >= values2))));
        }
+        vx_cleanup();
        return x;
    }
 };
@@ -1459,20 +1464,21 @@ struct InRange_SIMD<float>
        uchar * dst, int len) const
    {
        int x = 0;
-        const int width = (int)v_float32x4::nlanes * 2;
+        const int width = (int)v_float32::nlanes * 2;

        for (; x <= len - width; x += width)
        {
-            v_float32x4 values1 = v_load(src1 + x);
-            v_float32x4 low1 = v_load(src2 + x);
-            v_float32x4 high1 = v_load(src3 + x);
+            v_float32 values1 = vx_load(src1 + x);
+            v_float32 low1 = vx_load(src2 + x);
+            v_float32 high1 = vx_load(src3 + x);

-            v_float32x4 values2 = v_load(src1 + x + v_float32x4::nlanes);
-            v_float32x4 low2 = v_load(src2 + x + v_float32x4::nlanes);
-            v_float32x4 high2 = v_load(src3 + x + v_float32x4::nlanes);
+            v_float32 values2 = vx_load(src1 + x + v_float32::nlanes);
+            v_float32 low2 = vx_load(src2 + x + v_float32::nlanes);
+            v_float32 high2 = vx_load(src3 + x + v_float32::nlanes);

            v_pack_store(dst + x, v_pack(v_reinterpret_as_u32((values1 >= low1) & (high1 >= values1)), v_reinterpret_as_u32((values2 >= low2) & (high2 >= values2))));
        }
+        vx_cleanup();
        return x;
    }
 };
@@ -98,43 +98,15 @@ int normHamming(const uchar* a, const uchar* b, int n, int cellSize)
 float normL2Sqr_(const float* a, const float* b, int n)
 {
    int j = 0; float d = 0.f;
-#if CV_AVX2
-    float CV_DECL_ALIGNED(32) buf[8];
-    __m256 d0 = _mm256_setzero_ps();
-
-    for( ; j <= n - 8; j += 8 )
+#if CV_SIMD
+    v_float32 v_d = vx_setzero_f32();
+    for (; j <= n - v_float32::nlanes; j += v_float32::nlanes)
    {
-        __m256 t0 = _mm256_sub_ps(_mm256_loadu_ps(a + j), _mm256_loadu_ps(b + j));
-#if CV_FMA3
-        d0 = _mm256_fmadd_ps(t0, t0, d0);
-#else
-        d0 = _mm256_add_ps(d0, _mm256_mul_ps(t0, t0));
+        v_float32 t = vx_load(a + j) - vx_load(b + j);
+        v_d = v_muladd(t, t, v_d);
+    }
+    d = v_reduce_sum(v_d);
 #endif
-    }
-    _mm256_store_ps(buf, d0);
-    d = buf[0] + buf[1] + buf[2] + buf[3] + buf[4] + buf[5] + buf[6] + buf[7];
-#elif CV_SSE
-    float CV_DECL_ALIGNED(16) buf[4];
-    __m128 d0 = _mm_setzero_ps(), d1 = _mm_setzero_ps();
-
-    for( ; j <= n - 8; j += 8 )
-    {
-        __m128 t0 = _mm_sub_ps(_mm_loadu_ps(a + j), _mm_loadu_ps(b + j));
-        __m128 t1 = _mm_sub_ps(_mm_loadu_ps(a + j + 4), _mm_loadu_ps(b + j + 4));
-        d0 = _mm_add_ps(d0, _mm_mul_ps(t0, t0));
-        d1 = _mm_add_ps(d1, _mm_mul_ps(t1, t1));
-    }
-    _mm_store_ps(buf, _mm_add_ps(d0, d1));
-    d = buf[0] + buf[1] + buf[2] + buf[3];
-#endif
-    {
-        for( ; j <= n - 4; j += 4 )
-        {
-            float t0 = a[j] - b[j], t1 = a[j+1] - b[j+1], t2 = a[j+2] - b[j+2], t3 = a[j+3] - b[j+3];
-            d += t0*t0 + t1*t1 + t2*t2 + t3*t3;
-        }
-    }
-
    for( ; j < n; j++ )
    {
        float t = a[j] - b[j];
@@ -147,38 +119,12 @@ float normL2Sqr_(const float* a, const float* b, int n)
 float normL1_(const float* a, const float* b, int n)
 {
    int j = 0; float d = 0.f;
-#if CV_SSE
-    float CV_DECL_ALIGNED(16) buf[4];
-    static const int CV_DECL_ALIGNED(16) absbuf[4] = {0x7fffffff, 0x7fffffff, 0x7fffffff, 0x7fffffff};
-    __m128 d0 = _mm_setzero_ps(), d1 = _mm_setzero_ps();
-    __m128 absmask = _mm_load_ps((const float*)absbuf);
-
-    for( ; j <= n - 8; j += 8 )
-    {
-        __m128 t0 = _mm_sub_ps(_mm_loadu_ps(a + j), _mm_loadu_ps(b + j));
-        __m128 t1 = _mm_sub_ps(_mm_loadu_ps(a + j + 4), _mm_loadu_ps(b + j + 4));
-        d0 = _mm_add_ps(d0, _mm_and_ps(t0, absmask));
-        d1 = _mm_add_ps(d1, _mm_and_ps(t1, absmask));
-    }
-    _mm_store_ps(buf, _mm_add_ps(d0, d1));
-    d = buf[0] + buf[1] + buf[2] + buf[3];
-#elif CV_NEON
-    float32x4_t v_sum = vdupq_n_f32(0.0f);
-    for ( ; j <= n - 4; j += 4)
-        v_sum = vaddq_f32(v_sum, vabdq_f32(vld1q_f32(a + j), vld1q_f32(b + j)));
-
-    float CV_DECL_ALIGNED(16) buf[4];
-    vst1q_f32(buf, v_sum);
-    d = buf[0] + buf[1] + buf[2] + buf[3];
+#if CV_SIMD
+    v_float32 v_d = vx_setzero_f32();
+    for (; j <= n - v_float32::nlanes; j += v_float32::nlanes)
+        v_d += v_absdiff(vx_load(a + j), vx_load(b + j));
+    d = v_reduce_sum(v_d);
 #endif
-    {
-        for( ; j <= n - 4; j += 4 )
-        {
-            d += std::abs(a[j] - b[j]) + std::abs(a[j+1] - b[j+1]) +
-            std::abs(a[j+2] - b[j+2]) + std::abs(a[j+3] - b[j+3]);
-        }
-    }
-
    for( ; j < n; j++ )
        d += std::abs(a[j] - b[j]);
    return d;
@@ -187,46 +133,10 @@ float normL1_(const float* a, const float* b, int n)
 int normL1_(const uchar* a, const uchar* b, int n)
 {
    int j = 0, d = 0;
-#if CV_SSE
-    __m128i d0 = _mm_setzero_si128();
-
-    for( ; j <= n - 16; j += 16 )
-    {
-        __m128i t0 = _mm_loadu_si128((const __m128i*)(a + j));
-        __m128i t1 = _mm_loadu_si128((const __m128i*)(b + j));
-
-        d0 = _mm_add_epi32(d0, _mm_sad_epu8(t0, t1));
-    }
-
-    for( ; j <= n - 4; j += 4 )
-    {
-        __m128i t0 = _mm_cvtsi32_si128(*(const int*)(a + j));
-        __m128i t1 = _mm_cvtsi32_si128(*(const int*)(b + j));
-
-        d0 = _mm_add_epi32(d0, _mm_sad_epu8(t0, t1));
-    }
-    d = _mm_cvtsi128_si32(_mm_add_epi32(d0, _mm_unpackhi_epi64(d0, d0)));
-#elif CV_NEON
-    uint32x4_t v_sum = vdupq_n_u32(0.0f);
-    for ( ; j <= n - 16; j += 16)
-    {
-        uint8x16_t v_dst = vabdq_u8(vld1q_u8(a + j), vld1q_u8(b + j));
-        uint16x8_t v_low = vmovl_u8(vget_low_u8(v_dst)), v_high = vmovl_u8(vget_high_u8(v_dst));
-        v_sum = vaddq_u32(v_sum, vaddl_u16(vget_low_u16(v_low), vget_low_u16(v_high)));
-        v_sum = vaddq_u32(v_sum, vaddl_u16(vget_high_u16(v_low), vget_high_u16(v_high)));
-    }
-
-    uint CV_DECL_ALIGNED(16) buf[4];
-    vst1q_u32(buf, v_sum);
-    d = buf[0] + buf[1] + buf[2] + buf[3];
+#if CV_SIMD
+    for (; j <= n - v_uint8::nlanes; j += v_uint8::nlanes)
+        d += v_reduce_sad(vx_load(a + j), vx_load(b + j));
 #endif
-    {
-        for( ; j <= n - 4; j += 4 )
-        {
-            d += std::abs(a[j] - b[j]) + std::abs(a[j+1] - b[j+1]) +
-            std::abs(a[j+2] - b[j+2]) + std::abs(a[j+3] - b[j+3]);
-        }
-    }
    for( ; j < n; j++ )
        d += std::abs(a[j] - b[j]);
    return d;
@@ -1257,6 +1257,14 @@ struct Device::Impl
        else
            vendorID_ = UNKNOWN_VENDOR;

+        const size_t CV_OPENCL_DEVICE_MAX_WORK_GROUP_SIZE = utils::getConfigurationParameterSizeT("OPENCV_OPENCL_DEVICE_MAX_WORK_GROUP_SIZE", 0);
+        if (CV_OPENCL_DEVICE_MAX_WORK_GROUP_SIZE > 0)
+        {
+            const size_t new_maxWorkGroupSize = std::min(maxWorkGroupSize_, CV_OPENCL_DEVICE_MAX_WORK_GROUP_SIZE);
+            if (new_maxWorkGroupSize != maxWorkGroupSize_)
+                CV_LOG_WARNING(NULL, "OpenCL: using workgroup size: " << new_maxWorkGroupSize << " (was " << maxWorkGroupSize_ << ")");
+            maxWorkGroupSize_ = new_maxWorkGroupSize;
+        }
 #if 0
        if (isExtensionSupported("cl_khr_spir"))
        {
@@ -2772,6 +2780,7 @@ struct Kernel::Impl
        for( int i = 0; i < MAX_ARRS; i++ )
            u[i] = 0;
        haveTempDstUMats = false;
+        haveTempSrcUMats = false;
    }

    void cleanupUMats()
@@ -2788,6 +2797,7 @@ struct Kernel::Impl
            }
        nu = 0;
        haveTempDstUMats = false;
+        haveTempSrcUMats = false;
    }

    void addUMat(const UMat& m, bool dst)
@@ -2798,6 +2808,8 @@ struct Kernel::Impl
        nu++;
        if(dst && m.u->tempUMat())
            haveTempDstUMats = true;
+        if(m.u->originalUMatData == NULL && m.u->tempUMat())
+            haveTempSrcUMats = true;  // UMat is created on RAW memory (without proper lifetime management, even from Mat)
    }

    void addImage(const Image2D& image)
@@ -2835,6 +2847,7 @@ struct Kernel::Impl
    int nu;
    std::list<Image2D> images;
    bool haveTempDstUMats;
+    bool haveTempSrcUMats;
 };

 }} // namespace cv::ocl
@@ -3108,6 +3121,8 @@ bool Kernel::Impl::run(int dims, size_t globalsize[], size_t localsize[],
    cl_command_queue qq = getQueue(q);
    if (haveTempDstUMats)
        sync = true;
+    if (haveTempSrcUMats)
+        sync = true;
    if (timeNS)
        sync = true;
    cl_event asyncEvent = 0;
@@ -454,7 +454,7 @@ static inline int _initMaxThreads()
    {
        omp_set_dynamic(maxThreads);
    }
-    return numThreads;
+    return maxThreads;
 }
 static int numThreadsMax = _initMaxThreads();
 #elif defined HAVE_GCD
@@ -114,9 +114,11 @@ char* floatToString( char* buf, float value, bool halfprecision, bool explicitZe
        }
        else
        {
-            static const char* fmt = halfprecision ? "%.4e" : "%.8e";
            char* ptr = buf;
-            sprintf( buf, fmt, value );
+            if (halfprecision)
+                sprintf(buf, "%.4e", value);
+            else
+                sprintf(buf, "%.8e", value);
            if( *ptr == '+' || *ptr == '-' )
                ptr++;
            for( ; cv_isdigit(*ptr); ptr++ )
@@ -350,6 +352,7 @@ public:

    void init()
    {
+        flags = 0;
        buffer.clear();
        bufofs = 0;
        state = UNDEFINED;
@@ -358,6 +361,7 @@ public:
        write_mode = false;
        mem_mode = false;
        space = 0;
+        wrap_margin = 71;
        fmt = 0;
        file = 0;
        gzfile = 0;
@@ -615,7 +619,8 @@ public:
                    for(;;)
                    {
                        int line_offset = (int)ftell( file );
-                        char* ptr0 = gets( &xml_buf_[0], xml_buf_size ), *ptr;
+                        const char* ptr0 = this->gets(&xml_buf_[0], xml_buf_size );
+                        const char* ptr = NULL;
                        if( !ptr0 )
                            break;
                        ptr = ptr0;
@@ -708,7 +713,7 @@ public:
            const char* json_signature = "{";
            const char* xml_signature  = "<?xml";
            char buf[16];
-            gets( buf, sizeof(buf)-2 );
+            this->gets( buf, sizeof(buf)-2 );
            char* bufPtr = cv_skip_BOM(buf);
            size_t bufOffset = bufPtr - buf;

@@ -861,7 +866,7 @@ public:

    char* gets()
    {
-        char* ptr = gets(bufferStart(), (int)(bufferEnd() - bufferStart()));
+        char* ptr = this->gets(bufferStart(), (int)(bufferEnd() - bufferStart()));
        if( !ptr )
        {
            ptr = bufferStart();  // FIXIT Why do we need this hack? What is about other parsers JSON/YAML?
@@ -1766,11 +1771,13 @@ public:
 };

 FileStorage::FileStorage()
+    : state(0)
 {
    p = makePtr<FileStorage::Impl>(this);
 }

 FileStorage::FileStorage(const String& filename, int flags, const String& encoding)
+    : state(0)
 {
    p = makePtr<FileStorage::Impl>(this);
    bool ok = p->open(filename.c_str(), flags, encoding.c_str());
@@ -2197,7 +2204,8 @@ FileNode::operator double() const
        return DBL_MAX;
 }

-FileNode::operator std::string() const
+double FileNode::real() const  { return double(*this); }
+std::string FileNode::string() const
 {
    const uchar* p = ptr();
    if( !p || (*p & TYPE_MASK) != STRING )
@@ -2206,9 +2214,6 @@ FileNode::operator std::string() const
    size_t sz = (size_t)(unsigned)readInt(p);
    return std::string((const char*)(p + 4), sz - 1);
 }
-
-double FileNode::real() const  { return double(*this); }
-std::string FileNode::string() const { return std::string(*this); }
 Mat FileNode::mat() const { Mat value; read(*this, value, Mat()); return value; }

 FileNodeIterator FileNode::begin() const { return FileNodeIterator(*this, false); }
@@ -96,11 +96,20 @@ int decodeFormat( const char* dt, int* fmt_pairs, int max_len );
 int decodeSimpleFormat( const char* dt );
 }

+
+#ifdef CV_STATIC_ANALYSIS
+#define CV_PARSE_ERROR_CPP(errmsg) do { (void)fs; abort(); } while (0)
+#else
 #define CV_PARSE_ERROR_CPP( errmsg ) \
    fs->parseError( CV_Func, (errmsg), __FILE__, __LINE__ )
+#endif
+
+
+#define CV_PERSISTENCE_CHECK_END_OF_BUFFER_BUG_CPP() do { \
+    CV_DbgAssert(ptr); \
+    if((ptr)[0] == 0 && (ptr) == fs->bufferEnd() - 1) CV_PARSE_ERROR_CPP("OpenCV persistence doesn't support very long lines"); \
+} while (0)

-#define CV_PERSISTENCE_CHECK_END_OF_BUFFER_BUG_CPP() \
-    if((ptr)[0] == 0 && (ptr) == fs->bufferEnd() - 1) CV_PARSE_ERROR_CPP("OpenCV persistence doesn't support very long lines")

 class FileStorageParser;
 class FileStorageEmitter;
@@ -151,6 +160,7 @@ public:
    virtual double strtod(char* ptr, char** endptr) = 0;

    virtual char* parseBase64(char* ptr, int indent, FileNode& collection) = 0;
+    CV_NORETURN
    virtual void parseError(const char* funcname, const std::string& msg,
                            const char* filename, int lineno) = 0;
 };
@@ -107,15 +107,14 @@ void* allocSingletonBuffer(size_t size) { return fastMalloc(size); }
 #  include <cpu-features.h>
 #endif

-#ifndef __VSX__
-# if defined __PPC64__ && defined __linux__
-#   include "sys/auxv.h"
-#   ifndef AT_HWCAP2
-#     define AT_HWCAP2 26
-#   endif
-#   ifndef PPC_FEATURE2_ARCH_2_07
-#     define PPC_FEATURE2_ARCH_2_07 0x80000000
-#   endif
+
+#if CV_VSX && defined __linux__
+# include "sys/auxv.h"
+# ifndef AT_HWCAP2
+#   define AT_HWCAP2 26
+# endif
+# ifndef PPC_FEATURE2_ARCH_3_00
+#   define PPC_FEATURE2_ARCH_3_00 0x00800000
 # endif
 #endif

@@ -359,6 +358,7 @@ struct HWFeatures
        g_hwFeatureNames[CPU_NEON] = "NEON";

        g_hwFeatureNames[CPU_VSX] = "VSX";
+        g_hwFeatureNames[CPU_VSX3] = "VSX3";

        g_hwFeatureNames[CPU_AVX512_SKX] = "AVX512-SKX";
    }
@@ -513,14 +513,14 @@ struct HWFeatures
    #endif
    #endif

-    #ifdef __VSX__
-        have[CV_CPU_VSX] = true;
-    #elif (defined __PPC64__ && defined __linux__)
-        uint64 hwcaps = getauxval(AT_HWCAP);
+    // there's no need to check VSX availability in runtime since it's always available on ppc64le CPUs
+    have[CV_CPU_VSX] = (CV_VSX);
+    // TODO: Check VSX3 availability in runtime for other platforms
+    #if CV_VSX && defined __linux__
        uint64 hwcap2 = getauxval(AT_HWCAP2);
-        have[CV_CPU_VSX] = (hwcaps & PPC_FEATURE_PPC_LE && hwcaps & PPC_FEATURE_HAS_VSX && hwcap2 & PPC_FEATURE2_ARCH_2_07);
+        have[CV_CPU_VSX3] = (hwcap2 & PPC_FEATURE2_ARCH_3_00);
    #else
-        have[CV_CPU_VSX] = false;
+        have[CV_CPU_VSX3] = (CV_VSX3);
    #endif

        int baseline_features[] = { CV_CPU_BASELINE_FEATURES };
@@ -24,6 +24,8 @@
 #undef min
 #undef max
 #undef abs
+#elif defined(__linux__)
+#include <dlfcn.h>  // requires -ldl
 #elif defined(__APPLE__)
 #include <TargetConditionals.h>
 #if TARGET_OS_MAC
@@ -123,27 +125,10 @@ static cv::String getModuleLocation(const void* addr)
        }
    }
 #elif defined(__linux__)
-    std::ifstream fs("/proc/self/maps");
-    std::string line;
-    while (std::getline(fs, line, '\n'))
+    Dl_info info;
+    if (0 != dladdr(addr, &info))
    {
-        long long int addr_begin = 0, addr_end = 0;
-        if (2 == sscanf(line.c_str(), "%llx-%llx", &addr_begin, &addr_end))
-        {
-            if ((intptr_t)addr >= (intptr_t)addr_begin && (intptr_t)addr < (intptr_t)addr_end)
-            {
-                size_t pos = line.rfind("  "); // 2 spaces
-                if (pos == cv::String::npos)
-                    pos = line.rfind(' '); // 1 spaces
-                else
-                    pos++;
-                if (pos == cv::String::npos)
-                {
-                    CV_LOG_DEBUG(NULL, "Can't parse module path: '" << line << '\'');
-                }
-                return line.substr(pos + 1);
-            }
-        }
+        return cv::String(info.dli_fname);
    }
 #elif defined(__APPLE__)
 # if TARGET_OS_MAC
@@ -34,7 +34,7 @@
 #include <errno.h>
 #include <io.h>
 #include <stdio.h>
-#elif defined __linux__ || defined __APPLE__ || defined __HAIKU__
+#elif defined __linux__ || defined __APPLE__ || defined __HAIKU__ || defined __FreeBSD__
 #include <sys/types.h>
 #include <sys/stat.h>
 #include <fcntl.h>
@@ -178,7 +178,7 @@ cv::String getcwd()
    sz = GetCurrentDirectoryA((DWORD)buf.size(), buf.data());
    return cv::String(buf.data(), (size_t)sz);
 #endif
-#elif defined __linux__ || defined __APPLE__ || defined __HAIKU__
+#elif defined __linux__ || defined __APPLE__ || defined __HAIKU__ || defined __FreeBSD__
    for(;;)
    {
        char* p = ::getcwd(buf.data(), buf.size());
@@ -212,7 +212,7 @@ bool createDirectory(const cv::String& path)
 #else
    int result = _mkdir(path.c_str());
 #endif
-#elif defined __linux__ || defined __APPLE__ || defined __HAIKU__
+#elif defined __linux__ || defined __APPLE__ || defined __HAIKU__ || defined __FreeBSD__
    int result = mkdir(path.c_str(), 0777);
 #else
    int result = -1;
@@ -327,7 +327,7 @@ private:
    Impl& operator=(const Impl&); // disabled
 };

-#elif defined __linux__ || defined __APPLE__ || defined __HAIKU__
+#elif defined __linux__ || defined __APPLE__ || defined __HAIKU__ || defined __FreeBSD__

 struct FileLock::Impl
 {
@@ -441,7 +441,7 @@ cv::String getCacheDirectory(const char* sub_directory_name, const char* configu
            default_cache_path = "/tmp/";
            CV_LOG_WARNING(NULL, "Using world accessible cache directory. This may be not secure: " << default_cache_path);
        }
-#elif defined __linux__ || defined __HAIKU__
+#elif defined __linux__ || defined __HAIKU__ || defined __FreeBSD__
        // https://specifications.freedesktop.org/basedir-spec/basedir-spec-latest.html
        if (default_cache_path.empty())
        {
@@ -563,6 +563,7 @@ cv::String getCacheDirectory(const char* sub_directory_name, const char* configu
 cv::String canonical(const cv::String& /*path*/) { NOT_IMPLEMENTED }
 bool exists(const cv::String& /*path*/) { NOT_IMPLEMENTED }
 void remove_all(const cv::String& /*path*/) { NOT_IMPLEMENTED }
+cv::String getcwd() { NOT_IMPLEMENTED }
 bool createDirectory(const cv::String& /*path*/) { NOT_IMPLEMENTED }
 bool createDirectories(const cv::String& /*path*/) { NOT_IMPLEMENTED }
 cv::String getCacheDirectory(const char* /*sub_directory_name*/, const char* /*configuration_name = NULL*/) { NOT_IMPLEMENTED }
@@ -340,8 +340,8 @@ static void copy_convert_yv12_to_bgr(const VAImage& image, const unsigned char*
            1.5959997177f
        };

-    CV_CheckEQ(image.format.fourcc, VA_FOURCC_YV12, "Unexpected image format");
-    CV_CheckEQ(image.num_planes, 3, "");
+    CV_CheckEQ((size_t)image.format.fourcc, (size_t)VA_FOURCC_YV12, "Unexpected image format");
+    CV_CheckEQ((size_t)image.num_planes, (size_t)3, "");

    const size_t srcOffsetY = image.offsets[0];
    const size_t srcOffsetV = image.offsets[1];
@@ -417,8 +417,8 @@ static void copy_convert_bgr_to_yv12(const VAImage& image, const Mat& bgr, unsig
            -0.2909994125f, 0.438999176f, -0.3679990768f, -0.0709991455f
        };

-    CV_CheckEQ(image.format.fourcc, VA_FOURCC_YV12, "Unexpected image format");
-    CV_CheckEQ(image.num_planes, 3, "");
+    CV_CheckEQ((size_t)image.format.fourcc, (size_t)VA_FOURCC_YV12, "Unexpected image format");
+    CV_CheckEQ((size_t)image.num_planes, (size_t)3, "");

    const size_t dstOffsetY = image.offsets[0];
    const size_t dstOffsetV = image.offsets[1];
@@ -3,6 +3,12 @@
 // of this distribution and at http://opencv.org/license.html.
 #include "test_precomp.hpp"

+#ifdef HAVE_EIGEN
+#include <Eigen/Core>
+#include <Eigen/Dense>
+#include "opencv2/core/eigen.hpp"
+#endif
+
 namespace opencv_test { namespace {

 class Core_ReduceTest : public cvtest::BaseTest
@@ -1962,4 +1968,22 @@ TEST(Core_Vectors, issue_13078_workaround)
    ASSERT_EQ(7, ints[3]);
 }

+
+#ifdef HAVE_EIGEN
+TEST(Core_Eigen, eigen2cv_check_Mat_type)
+{
+    Mat A(4, 4, CV_32FC1, Scalar::all(0));
+    Eigen::MatrixXf eigen_A;
+    cv2eigen(A, eigen_A);
+
+    Mat_<float> f_mat;
+    EXPECT_NO_THROW(eigen2cv(eigen_A, f_mat));
+    EXPECT_EQ(CV_32FC1, f_mat.type());
+
+    Mat_<double> d_mat;
+    EXPECT_ANY_THROW(eigen2cv(eigen_A, d_mat));
+    //EXPECT_EQ(CV_64FC1, d_mat.type());
+}
+#endif // HAVE_EIGEN
+
 }} // namespace
@@ -972,6 +972,13 @@ bool CV_OperationsTest::operations1()
        if (sz.width != 10 || sz.height != 20) throw test_excep();
        if (cvSize(sz).width != 10 || cvSize(sz).height != 20) throw test_excep();

+        Rect r1(0, 0, 10, 20);
+        Size sz1(5, 10);
+        r1 -= sz1;
+        if (r1.size().width != 5 || r1.size().height != 10) throw test_excep();
+        Rect r2 = r1 - sz1;
+        if (r2.size().width != 0 || r2.size().height != 0) throw test_excep();
+
        Vec<double, 5> v5d(1, 1, 1, 1, 1);
        Vec<double, 6> v6d(1, 1, 1, 1, 1, 1);
        Vec<double, 7> v7d(1, 1, 1, 1, 1, 1, 1);
@@ -77,6 +77,15 @@ CV__DNN_INLINE_NS_BEGIN
        static Ptr<Layer> create(const LayerParams &params);
    };

+    /**
+     * Constant layer produces the same data blob at an every forward pass.
+     */
+    class CV_EXPORTS ConstLayer : public Layer
+    {
+    public:
+        static Ptr<Layer> create(const LayerParams &params);
+    };
+
    //! LSTM recurrent layer
    class CV_EXPORTS LSTMLayer : public Layer
    {
@@ -83,9 +83,14 @@ CV__DNN_INLINE_NS_BEGIN
        DNN_TARGET_OPENCL,
        DNN_TARGET_OPENCL_FP16,
        DNN_TARGET_MYRIAD,
-        DNN_TARGET_VULKAN
+        DNN_TARGET_VULKAN,
+        //! FPGA device with CPU fallbacks using Inference Engine's Heterogeneous plugin.
+        DNN_TARGET_FPGA
    };

+    CV_EXPORTS std::vector< std::pair<Backend, Target> > getAvailableBackends();
+    CV_EXPORTS std::vector<Target> getAvailableTargets(Backend be);
+
    /** @brief This class provides all data needed to initialize layer.
     *
     * It includes dictionary with scalar params (which can be read by using Dict interface),
@@ -497,6 +502,7 @@ CV__DNN_INLINE_NS_BEGIN
         * | DNN_TARGET_OPENCL      |                  + |                            + |                  + |
         * | DNN_TARGET_OPENCL_FP16 |                  + |                            + |                    |
         * | DNN_TARGET_MYRIAD      |                    |                            + |                    |
+         * | DNN_TARGET_FPGA        |                    |                            + |                    |
         */
        CV_WRAP void setPreferableTarget(int targetId);

@@ -744,6 +750,7 @@ CV__DNN_INLINE_NS_BEGIN
     *  @brief Reads a network model stored in <a href="http://torch.ch">Torch7</a> framework's format.
     *  @param model    path to the file, dumped from Torch by using torch.save() function.
     *  @param isBinary specifies whether the network was serialized in ascii mode or binary.
+     *  @param evaluate specifies testing phase of network. If true, it's similar to evaluate() method in Torch.
     *  @returns Net object.
     *
     *  @note Ascii mode of Torch serializer is more preferable, because binary mode extensively use `long` type of C language,
@@ -765,7 +772,7 @@ CV__DNN_INLINE_NS_BEGIN
     *
     * Also some equivalents of these classes from cunn, cudnn, and fbcunn may be successfully imported.
     */
-     CV_EXPORTS_W Net readNetFromTorch(const String &model, bool isBinary = true);
+     CV_EXPORTS_W Net readNetFromTorch(const String &model, bool isBinary = true, bool evaluate = true);

     /**
      * @brief Read deep learning network represented in one of the supported formats.
@@ -6,7 +6,7 @@
 #define OPENCV_DNN_VERSION_HPP

 /// Use with major OpenCV version only.
-#define OPENCV_DNN_API_VERSION 20180917
+#define OPENCV_DNN_API_VERSION 20181221

 #if !defined CV_DOXYGEN && !defined CV_DNN_DONT_ADD_INLINE_NS
 #define CV__DNN_INLINE_NS __CV_CAT(dnn4_v, OPENCV_DNN_API_VERSION)
@@ -31,23 +31,6 @@ public:
    void processNet(std::string weights, std::string proto, std::string halide_scheduler,
                    const Mat& input, const std::string& outputLayer = "")
    {
-        if (backend == DNN_BACKEND_OPENCV && (target == DNN_TARGET_OPENCL || target == DNN_TARGET_OPENCL_FP16))
-        {
-#if defined(HAVE_OPENCL)
-            if (!cv::ocl::useOpenCL())
-#endif
-            {
-                throw cvtest::SkipTestException("OpenCL is not available/disabled in OpenCV");
-            }
-        }
-        if (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_MYRIAD)
-        {
-            if (!checkMyriadTarget())
-            {
-                throw SkipTestException("Myriad is not available/disabled in OpenCV");
-            }
-        }
-
        randu(input, 0.0f, 1.0f);

        weights = findDataFile(weights, false);
@@ -85,6 +85,111 @@ using std::map;
 using std::make_pair;
 using std::set;

+//==================================================================================================
+
+class BackendRegistry
+{
+public:
+    typedef std::vector< std::pair<Backend, Target> > BackendsList;
+    const BackendsList & getBackends() const { return backends; }
+    static BackendRegistry & getRegistry()
+    {
+        static BackendRegistry impl;
+        return impl;
+    }
+private:
+    BackendRegistry()
+    {
+#ifdef HAVE_HALIDE
+        backends.push_back(std::make_pair(DNN_BACKEND_HALIDE, DNN_TARGET_CPU));
+#  ifdef HAVE_OPENCL
+        if (cv::ocl::useOpenCL())
+            backends.push_back(std::make_pair(DNN_BACKEND_HALIDE, DNN_TARGET_OPENCL));
+#  endif
+#endif // HAVE_HALIDE
+
+#ifdef HAVE_INF_ENGINE
+        if (checkIETarget(DNN_TARGET_CPU))
+            backends.push_back(std::make_pair(DNN_BACKEND_INFERENCE_ENGINE, DNN_TARGET_CPU));
+        if (checkIETarget(DNN_TARGET_MYRIAD))
+            backends.push_back(std::make_pair(DNN_BACKEND_INFERENCE_ENGINE, DNN_TARGET_MYRIAD));
+        if (checkIETarget(DNN_TARGET_FPGA))
+            backends.push_back(std::make_pair(DNN_BACKEND_INFERENCE_ENGINE, DNN_TARGET_FPGA));
+#  ifdef HAVE_OPENCL
+        if (cv::ocl::useOpenCL() && ocl::Device::getDefault().isIntel())
+        {
+            if (checkIETarget(DNN_TARGET_OPENCL))
+                backends.push_back(std::make_pair(DNN_BACKEND_INFERENCE_ENGINE, DNN_TARGET_OPENCL));
+            if (checkIETarget(DNN_TARGET_OPENCL_FP16))
+                backends.push_back(std::make_pair(DNN_BACKEND_INFERENCE_ENGINE, DNN_TARGET_OPENCL_FP16));
+        }
+#  endif
+#endif // HAVE_INF_ENGINE
+
+#ifdef HAVE_OPENCL
+        if (cv::ocl::useOpenCL())
+        {
+            backends.push_back(std::make_pair(DNN_BACKEND_OPENCV, DNN_TARGET_OPENCL));
+            backends.push_back(std::make_pair(DNN_BACKEND_OPENCV, DNN_TARGET_OPENCL_FP16));
+        }
+#endif
+
+        backends.push_back(std::make_pair(DNN_BACKEND_OPENCV, DNN_TARGET_CPU));
+
+#ifdef HAVE_VULKAN
+        backends.push_back(std::make_pair(DNN_BACKEND_VKCOM, DNN_TARGET_VULKAN));  // TODO Add device check
+#endif
+    }
+    static inline bool checkIETarget(int target)
+    {
+#ifndef HAVE_INF_ENGINE
+        return false;
+#else
+        cv::dnn::Net net;
+        cv::dnn::LayerParams lp;
+        net.addLayerToPrev("testLayer", "Identity", lp);
+        net.setPreferableBackend(cv::dnn::DNN_BACKEND_INFERENCE_ENGINE);
+        net.setPreferableTarget(target);
+        static int inpDims[] = {1, 2, 3, 4};
+        net.setInput(cv::Mat(4, &inpDims[0], CV_32FC1, cv::Scalar(0)));
+        try
+        {
+            net.forward();
+        }
+        catch(...)
+        {
+            return false;
+        }
+        return true;
+#endif
+    }
+
+    BackendsList backends;
+};
+
+
+std::vector< std::pair<Backend, Target> > getAvailableBackends()
+{
+    return BackendRegistry::getRegistry().getBackends();
+}
+
+std::vector<Target> getAvailableTargets(Backend be)
+{
+    if (be == DNN_BACKEND_DEFAULT)
+        be = (Backend)PARAM_DNN_BACKEND_DEFAULT;
+
+    std::vector<Target> result;
+    const BackendRegistry::BackendsList all_backends = getAvailableBackends();
+    for(BackendRegistry::BackendsList::const_iterator i = all_backends.begin(); i != all_backends.end(); ++i )
+    {
+        if (i->first == be)
+            result.push_back(i->second);
+    }
+    return result;
+}
+
+//==================================================================================================
+
 namespace
 {
    typedef std::vector<MatShape> ShapesVec;
@@ -911,21 +1016,8 @@ struct Net::Impl
    typedef std::map<int, LayerShapes> LayersShapesMap;
    typedef std::map<int, LayerData> MapIdToLayerData;

-    ~Impl()
-    {
-#ifdef HAVE_VULKAN
-        // Vulkan requires explicit releasing the child objects of
-        // VkDevice object prior to releasing VkDevice object itself.
-        layers.clear();
-        backendWrappers.clear();
-        vkcom::deinitPerThread();
-#endif
-    }
    Impl()
    {
-#ifdef HAVE_VULKAN
-        vkcom::initPerThread();
-#endif
        //allocate fake net input layer
        netInputLayer = Ptr<DataLayer>(new DataLayer());
        LayerData &inpl = layers.insert( make_pair(0, LayerData()) ).first->second;
@@ -1104,7 +1196,8 @@ struct Net::Impl
                  preferableTarget == DNN_TARGET_CPU ||
                  preferableTarget == DNN_TARGET_OPENCL ||
                  preferableTarget == DNN_TARGET_OPENCL_FP16 ||
-                  preferableTarget == DNN_TARGET_MYRIAD);
+                  preferableTarget == DNN_TARGET_MYRIAD ||
+                  preferableTarget == DNN_TARGET_FPGA);
        CV_Assert(preferableBackend != DNN_BACKEND_VKCOM ||
                  preferableTarget == DNN_TARGET_VULKAN);
        if (!netWasAllocated || this->blobsToKeep != blobsToKeep_)
@@ -1609,7 +1702,9 @@ struct Net::Impl
            ieNode->net = net;

            auto weightableLayer = std::dynamic_pointer_cast<InferenceEngine::WeightableLayer>(ieNode->layer);
-            if ((preferableTarget == DNN_TARGET_OPENCL_FP16 || preferableTarget == DNN_TARGET_MYRIAD) && !fused)
+            if ((preferableTarget == DNN_TARGET_OPENCL_FP16 ||
+                 preferableTarget == DNN_TARGET_MYRIAD ||
+                 preferableTarget == DNN_TARGET_FPGA) && !fused)
            {
                ieNode->layer->precision = InferenceEngine::Precision::FP16;
                if (weightableLayer)
@@ -112,6 +112,7 @@ void initializeLayerFactory()
    CV_DNN_REGISTER_LAYER_CLASS(Dropout,        BlankLayer);
    CV_DNN_REGISTER_LAYER_CLASS(Identity,       BlankLayer);
    CV_DNN_REGISTER_LAYER_CLASS(Silence,        BlankLayer);
+    CV_DNN_REGISTER_LAYER_CLASS(Const,          ConstLayer);

    CV_DNN_REGISTER_LAYER_CLASS(Crop,           CropLayer);
    CV_DNN_REGISTER_LAYER_CLASS(Eltwise,        EltwiseLayer);
@@ -119,8 +119,8 @@ public:
        lp.precision = InferenceEngine::Precision::FP32;
        std::shared_ptr<InferenceEngine::SplitLayer> ieLayer(new InferenceEngine::SplitLayer(lp));
 #if INF_ENGINE_VER_MAJOR_GT(INF_ENGINE_RELEASE_2018R3)
-        ieLayer->params["axis"] = format("%d", input->dims.size() - 1);
-        ieLayer->params["out_sizes"] = format("%d", input->dims[0]);
+        ieLayer->params["axis"] = format("%d", (int)input->dims.size() - 1);
+        ieLayer->params["out_sizes"] = format("%d", (int)input->dims[0]);
 #endif
        return Ptr<BackendNode>(new InfEngineBackendNode(ieLayer));
 #endif  // HAVE_INF_ENGINE
@@ -0,0 +1,68 @@
+// This file is part of OpenCV project.
+// It is subject to the license terms in the LICENSE file found in the top-level directory
+// of this distribution and at http://opencv.org/license.html.
+
+// Copyright (C) 2018, Intel Corporation, all rights reserved.
+// Third party copyrights are property of their respective owners.
+
+#include "../precomp.hpp"
+#include "layers_common.hpp"
+
+#ifdef HAVE_OPENCL
+#include "opencl_kernels_dnn.hpp"
+#endif
+
+namespace cv { namespace dnn {
+
+class ConstLayerImpl CV_FINAL : public ConstLayer
+{
+public:
+    ConstLayerImpl(const LayerParams& params)
+    {
+        setParamsFrom(params);
+        CV_Assert(blobs.size() == 1);
+    }
+
+    virtual bool getMemoryShapes(const std::vector<MatShape> &inputs,
+                                 const int requiredOutputs,
+                                 std::vector<MatShape> &outputs,
+                                 std::vector<MatShape> &internals) const CV_OVERRIDE
+    {
+        CV_Assert(inputs.empty());
+        outputs.assign(1, shape(blobs[0]));
+        return false;
+    }
+
+#ifdef HAVE_OPENCL
+    bool forward_ocl(InputArrayOfArrays inps, OutputArrayOfArrays outs, OutputArrayOfArrays internals)
+    {
+        std::vector<UMat> outputs;
+        outs.getUMatVector(outputs);
+        if (outs.depth() == CV_16S)
+            convertFp16(blobs[0], outputs[0]);
+        else
+            blobs[0].copyTo(outputs[0]);
+        return true;
+    }
+#endif
+
+    void forward(InputArrayOfArrays inputs_arr, OutputArrayOfArrays outputs_arr, OutputArrayOfArrays internals_arr) CV_OVERRIDE
+    {
+        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))
+
+        std::vector<Mat> outputs;
+        outputs_arr.getMatVector(outputs);
+        blobs[0].copyTo(outputs[0]);
+    }
+};
+
+Ptr<Layer> ConstLayer::create(const LayerParams& params)
+{
+    return Ptr<Layer>(new ConstLayerImpl(params));
+}
+
+}}  // namespace cv::dnn
@@ -220,9 +220,14 @@ public:

    virtual bool supportBackend(int backendId) CV_OVERRIDE
    {
+#ifdef HAVE_INF_ENGINE
        if (backendId == DNN_BACKEND_INFERENCE_ENGINE)
-            return preferableTarget != DNN_TARGET_MYRIAD || dilation.width == dilation.height;
+        {
+            return INF_ENGINE_VER_MAJOR_GE(INF_ENGINE_RELEASE_2018R4) ||
+                   (preferableTarget != DNN_TARGET_MYRIAD || dilation.width == dilation.height);
+        }
        else
+#endif
            return backendId == DNN_BACKEND_OPENCV ||
                   backendId == DNN_BACKEND_HALIDE ||
                   (backendId == DNN_BACKEND_VKCOM && haveVulkan());
@@ -752,7 +752,8 @@ struct AbsValFunctor

    bool supportBackend(int backendId, int)
    {
-        return backendId == DNN_BACKEND_OPENCV || backendId == DNN_BACKEND_HALIDE;
+        return backendId == DNN_BACKEND_OPENCV || backendId == DNN_BACKEND_HALIDE ||
+               backendId == DNN_BACKEND_INFERENCE_ENGINE;
    }

    void apply(const float* srcptr, float* dstptr, int len, size_t planeSize, int cn0, int cn1) const
@@ -806,8 +807,11 @@ struct AbsValFunctor
 #ifdef HAVE_INF_ENGINE
    InferenceEngine::CNNLayerPtr initInfEngine(InferenceEngine::LayerParams& lp)
    {
-        CV_Error(Error::StsNotImplemented, "Abs");
-        return InferenceEngine::CNNLayerPtr();
+        lp.type = "ReLU";
+        std::shared_ptr<InferenceEngine::ReLULayer> ieLayer(new InferenceEngine::ReLULayer(lp));
+        ieLayer->negative_slope = -1;
+        ieLayer->params["negative_slope"] = "-1.0";
+        return ieLayer;
    }
 #endif  // HAVE_INF_ENGINE

@@ -900,7 +904,7 @@ struct PowerFunctor
    bool supportBackend(int backendId, int targetId)
    {
        if (backendId == DNN_BACKEND_INFERENCE_ENGINE)
-            return (targetId != DNN_TARGET_OPENCL && targetId != DNN_TARGET_OPENCL_FP16) || power == 1.0;
+            return (targetId != DNN_TARGET_OPENCL && targetId != DNN_TARGET_OPENCL_FP16) || power == 1.0 || power == 0.5;
        else
            return backendId == DNN_BACKEND_OPENCV || backendId == DNN_BACKEND_HALIDE;
    }
@@ -1054,7 +1058,8 @@ struct ChannelsPReLUFunctor

    bool supportBackend(int backendId, int)
    {
-        return backendId == DNN_BACKEND_OPENCV || backendId == DNN_BACKEND_HALIDE;
+        return backendId == DNN_BACKEND_OPENCV || backendId == DNN_BACKEND_HALIDE ||
+               backendId == DNN_BACKEND_INFERENCE_ENGINE;
    }

    void apply(const float* srcptr, float* dstptr, int len, size_t planeSize, int cn0, int cn1) const
@@ -1140,8 +1145,11 @@ struct ChannelsPReLUFunctor
 #ifdef HAVE_INF_ENGINE
    InferenceEngine::CNNLayerPtr initInfEngine(InferenceEngine::LayerParams& lp)
    {
-        CV_Error(Error::StsNotImplemented, "PReLU");
-        return InferenceEngine::CNNLayerPtr();
+        lp.type = "PReLU";
+        std::shared_ptr<InferenceEngine::PReLULayer> ieLayer(new InferenceEngine::PReLULayer(lp));
+        const size_t numChannels = scale.total();
+        ieLayer->_weights = wrapToInfEngineBlob(scale, {numChannels}, InferenceEngine::Layout::C);
+        return ieLayer;
    }
 #endif  // HAVE_INF_ENGINE

@@ -98,7 +98,8 @@ public:
    {
        return backendId == DNN_BACKEND_OPENCV ||
               backendId == DNN_BACKEND_HALIDE ||
-               (backendId == DNN_BACKEND_INFERENCE_ENGINE && (op != SUM || coeffs.empty()));
+               (backendId == DNN_BACKEND_INFERENCE_ENGINE &&
+                (preferableTarget != DNN_TARGET_MYRIAD || coeffs.empty()));
    }

    bool getMemoryShapes(const std::vector<MatShape> &inputs,
@@ -427,6 +428,7 @@ public:
        lp.type = "Eltwise";
        lp.precision = InferenceEngine::Precision::FP32;
        std::shared_ptr<InferenceEngine::EltwiseLayer> ieLayer(new InferenceEngine::EltwiseLayer(lp));
+        ieLayer->coeff = coeffs;
        if (op == SUM)
            ieLayer->_operation = InferenceEngine::EltwiseLayer::Sum;
        else if (op == PROD)
@@ -116,9 +116,15 @@ public:

    virtual bool supportBackend(int backendId) CV_OVERRIDE
    {
+#ifdef HAVE_INF_ENGINE
        if (backendId == DNN_BACKEND_INFERENCE_ENGINE)
+#if INF_ENGINE_VER_MAJOR_GE(INF_ENGINE_RELEASE_2018R5)
+            return !zeroDev && eps <= 1e-7f;
+#else
            return !zeroDev && (preferableTarget == DNN_TARGET_CPU || eps <= 1e-7f);
+#endif
        else
+#endif  // HAVE_INF_ENGINE
            return backendId == DNN_BACKEND_OPENCV;
    }

@@ -6,6 +6,7 @@
 // Third party copyrights are property of their respective owners.

 #include "../precomp.hpp"
+#include <opencv2/dnn/shape_utils.hpp>

 #ifdef HAVE_PROTOBUF

@@ -134,9 +135,38 @@ Mat getMatFromTensor(opencv_onnx::TensorProto& tensor_proto)
    else
        CV_Error(Error::StsUnsupportedFormat, "Unsupported data type: " +
                        opencv_onnx::TensorProto_DataType_Name(datatype));
+    if (tensor_proto.dims_size() == 0)
+        blob.dims = 1;  // To force 1-dimensional cv::Mat for scalars.
    return blob;
 }

+void runLayer(Ptr<Layer> layer, const std::vector<Mat>& inputs,
+              std::vector<Mat>& outputs)
+{
+    std::vector<MatShape> inpShapes(inputs.size());
+    int ddepth = CV_32F;
+    for (size_t i = 0; i < inputs.size(); ++i)
+    {
+        inpShapes[i] = shape(inputs[i]);
+        if (i > 0 && ddepth != inputs[i].depth())
+            CV_Error(Error::StsNotImplemented, "Mixed input data types.");
+        ddepth = inputs[i].depth();
+    }
+
+    std::vector<MatShape> outShapes, internalShapes;
+    layer->getMemoryShapes(inpShapes, 0, outShapes, internalShapes);
+
+    std::vector<Mat> internals(internalShapes.size());
+    outputs.resize(outShapes.size());
+    for (size_t i = 0; i < outShapes.size(); ++i)
+        outputs[i].create(outShapes[i], ddepth);
+    for (size_t i = 0; i < internalShapes.size(); ++i)
+        internals[i].create(internalShapes[i], ddepth);
+
+    layer->finalize(inputs, outputs);
+    layer->forward(inputs, outputs, internals);
+}
+
 std::map<std::string, Mat> ONNXImporter::getGraphTensors(
                                        const opencv_onnx::GraphProto& graph_proto)
 {
@@ -292,6 +322,26 @@ void ONNXImporter::populateNet(Net dstNet)
    CV_Assert(model_proto.has_graph());
    opencv_onnx::GraphProto graph_proto = model_proto.graph();
    std::map<std::string, Mat> constBlobs = getGraphTensors(graph_proto);
+    // List of internal blobs shapes.
+    std::map<std::string, MatShape> outShapes;
+    // Add all the inputs shapes. It includes as constant blobs as network's inputs shapes.
+    for (int i = 0; i < graph_proto.input_size(); ++i)
+    {
+        opencv_onnx::ValueInfoProto valueInfoProto = graph_proto.input(i);
+        CV_Assert(valueInfoProto.has_type());
+        opencv_onnx::TypeProto typeProto = valueInfoProto.type();
+        CV_Assert(typeProto.has_tensor_type());
+        opencv_onnx::TypeProto::Tensor tensor = typeProto.tensor_type();
+        CV_Assert(tensor.has_shape());
+        opencv_onnx::TensorShapeProto tensorShape = tensor.shape();
+
+        MatShape inpShape(tensorShape.dim_size());
+        for (int j = 0; j < inpShape.size(); ++j)
+        {
+            inpShape[j] = tensorShape.dim(j).dim_value();
+        }
+        outShapes[valueInfoProto.name()] = inpShape;
+    }

    std::string framework_name;
    if (model_proto.has_producer_name()) {
@@ -301,6 +351,7 @@ void ONNXImporter::populateNet(Net dstNet)
    // create map with network inputs (without const blobs)
    std::map<std::string, LayerInfo> layer_id;
    std::map<std::string, LayerInfo>::iterator layerId;
+    std::map<std::string, MatShape>::iterator shapeIt;
    // fill map: push layer name, layer id and output id
    std::vector<String> netInputs;
    for (int j = 0; j < graph_proto.input_size(); j++)
@@ -317,9 +368,9 @@ void ONNXImporter::populateNet(Net dstNet)
    LayerParams layerParams;
    opencv_onnx::NodeProto node_proto;

-    for(int i = 0; i < layersSize; i++)
+    for(int li = 0; li < layersSize; li++)
    {
-        node_proto = graph_proto.node(i);
+        node_proto = graph_proto.node(li);
        layerParams = getLayerParams(node_proto);
        CV_Assert(node_proto.output_size() >= 1);
        layerParams.name = node_proto.output(0);
@@ -358,7 +409,8 @@ void ONNXImporter::populateNet(Net dstNet)
                    layerParams.set("shift", blob.at<float>(0));
                }
                else {
-                    layerParams.type = "Shift";
+                    layerParams.type = "Scale";
+                    layerParams.set("bias_term", true);
                    layerParams.blobs.push_back(blob);
                }
            }
@@ -368,15 +420,32 @@ void ONNXImporter::populateNet(Net dstNet)
        }
        else if (layer_type == "Sub")
        {
-            Mat blob = (-1.0f) * getBlob(node_proto, constBlobs, 1);
-            blob = blob.reshape(1, 1);
+            Mat blob = getBlob(node_proto, constBlobs, 1);
            if (blob.total() == 1) {
                layerParams.type = "Power";
-                layerParams.set("shift", blob.at<float>(0));
+                layerParams.set("shift", -blob.at<float>(0));
            }
            else {
-                layerParams.type = "Shift";
+                layerParams.type = "Scale";
+                layerParams.set("has_bias", true);
+                layerParams.blobs.push_back(-1.0f * blob.reshape(1, 1));
+            }
+        }
+        else if (layer_type == "Div")
+        {
+            Mat blob = getBlob(node_proto, constBlobs, 1);
+            CV_Assert_N(blob.type() == CV_32F, blob.total());
+            if (blob.total() == 1)
+            {
+                layerParams.set("scale", 1.0f / blob.at<float>(0));
+                layerParams.type = "Power";
+            }
+            else
+            {
+                layerParams.type = "Scale";
+                divide(1.0, blob, blob);
                layerParams.blobs.push_back(blob);
+                layerParams.set("bias_term", false);
            }
        }
        else if (layer_type == "Constant")
@@ -579,6 +648,65 @@ void ONNXImporter::populateNet(Net dstNet)
        {
            layerParams.type = "Padding";
        }
+        else if (layer_type == "Shape")
+        {
+            CV_Assert(node_proto.input_size() == 1);
+            shapeIt = outShapes.find(node_proto.input(0));
+            CV_Assert(shapeIt != outShapes.end());
+            MatShape inpShape = shapeIt->second;
+
+            Mat shapeMat(inpShape.size(), 1, CV_32S);
+            for (int j = 0; j < inpShape.size(); ++j)
+                shapeMat.at<int>(j) = inpShape[j];
+            shapeMat.dims = 1;
+
+            constBlobs.insert(std::make_pair(layerParams.name, shapeMat));
+            continue;
+        }
+        else if (layer_type == "Gather")
+        {
+            CV_Assert(node_proto.input_size() == 2);
+            CV_Assert(layerParams.has("axis"));
+            Mat input = getBlob(node_proto, constBlobs, 0);
+            Mat indexMat = getBlob(node_proto, constBlobs, 1);
+            CV_Assert_N(indexMat.type() == CV_32S, indexMat.total() == 1);
+            int index = indexMat.at<int>(0);
+            int axis = layerParams.get<int>("axis");
+
+            std::vector<cv::Range> ranges(input.dims, Range::all());
+            ranges[axis] = Range(index, index + 1);
+
+            Mat out = input(ranges);
+            constBlobs.insert(std::make_pair(layerParams.name, out));
+            continue;
+        }
+        else if (layer_type == "Concat")
+        {
+            bool hasVariableInps = false;
+            for (int i = 0; i < node_proto.input_size(); ++i)
+            {
+                if (layer_id.find(node_proto.input(i)) != layer_id.end())
+                {
+                    hasVariableInps = true;
+                    break;
+                }
+            }
+
+            if (!hasVariableInps)
+            {
+                std::vector<Mat> inputs(node_proto.input_size()), concatenated;
+                for (size_t i = 0; i < inputs.size(); ++i)
+                {
+                    inputs[i] = getBlob(node_proto, constBlobs, i);
+                }
+                Ptr<Layer> concat = ConcatLayer::create(layerParams);
+                runLayer(concat, inputs, concatenated);
+
+                CV_Assert(concatenated.size() == 1);
+                constBlobs.insert(std::make_pair(layerParams.name, concatenated[0]));
+                continue;
+            }
+        }
        else
        {
            for (int j = 0; j < node_proto.input_size(); j++) {
@@ -590,12 +718,24 @@ void ONNXImporter::populateNet(Net dstNet)
         int id = dstNet.addLayer(layerParams.name, layerParams.type, layerParams);
         layer_id.insert(std::make_pair(layerParams.name, LayerInfo(id, 0)));

+
+         std::vector<MatShape> layerInpShapes, layerOutShapes, layerInternalShapes;
         for (int j = 0; j < node_proto.input_size(); j++) {
             layerId = layer_id.find(node_proto.input(j));
             if (layerId != layer_id.end()) {
                 dstNet.connect(layerId->second.layerId, layerId->second.outputId, id, j);
+                 // Collect input shapes.
+                 shapeIt = outShapes.find(node_proto.input(j));
+                 CV_Assert(shapeIt != outShapes.end());
+                 layerInpShapes.push_back(shapeIt->second);
             }
         }
+
+         // Compute shape of output blob for this layer.
+         Ptr<Layer> layer = dstNet.getLayer(id);
+         layer->getMemoryShapes(layerInpShapes, 0, layerOutShapes, layerInternalShapes);
+         CV_Assert(!layerOutShapes.empty());
+         outShapes[layerParams.name] = layerOutShapes[0];
     }
 }

@@ -152,6 +152,7 @@ InfEngineBackendNet::InfEngineBackendNet()
 {
    targetDevice = InferenceEngine::TargetDevice::eCPU;
    precision = InferenceEngine::Precision::FP32;
+    hasNetOwner = false;
 }

 InfEngineBackendNet::InfEngineBackendNet(InferenceEngine::CNNNetwork& net)
@@ -162,6 +163,7 @@ InfEngineBackendNet::InfEngineBackendNet(InferenceEngine::CNNNetwork& net)
    outputs = net.getOutputsInfo();
    layers.resize(net.layerCount());  // A hack to execute InfEngineBackendNet::layerCount correctly.
    netOwner = net;
+    hasNetOwner = true;
 }

 void InfEngineBackendNet::Release() CV_NOEXCEPT
@@ -178,12 +180,12 @@ void InfEngineBackendNet::setPrecision(InferenceEngine::Precision p) CV_NOEXCEPT

 InferenceEngine::Precision InfEngineBackendNet::getPrecision() CV_NOEXCEPT
 {
-    return precision;
+    return hasNetOwner ? netOwner.getPrecision() : precision;
 }

 InferenceEngine::Precision InfEngineBackendNet::getPrecision() const CV_NOEXCEPT
 {
-    return precision;
+    return hasNetOwner ? netOwner.getPrecision() : precision;
 }

 // Assume that outputs of network is unconnected blobs.
@@ -233,6 +235,12 @@ const std::string& InfEngineBackendNet::getName() const CV_NOEXCEPT
    return name;
 }

+InferenceEngine::StatusCode InfEngineBackendNet::serialize(const std::string&, const std::string&, InferenceEngine::ResponseDesc*) const CV_NOEXCEPT
+{
+    CV_Error(Error::StsNotImplemented, "");
+    return InferenceEngine::StatusCode::OK;
+}
+
 size_t InfEngineBackendNet::layerCount() CV_NOEXCEPT
 {
    return const_cast<const InfEngineBackendNet*>(this)->layerCount();
@@ -302,7 +310,8 @@ void InfEngineBackendNet::setTargetDevice(InferenceEngine::TargetDevice device)
 {
    if (device != InferenceEngine::TargetDevice::eCPU &&
        device != InferenceEngine::TargetDevice::eGPU &&
-        device != InferenceEngine::TargetDevice::eMYRIAD)
+        device != InferenceEngine::TargetDevice::eMYRIAD &&
+        device != InferenceEngine::TargetDevice::eFPGA)
        CV_Error(Error::StsNotImplemented, "");
    targetDevice = device;
 }
@@ -314,7 +323,8 @@ InferenceEngine::TargetDevice InfEngineBackendNet::getTargetDevice() CV_NOEXCEPT

 InferenceEngine::TargetDevice InfEngineBackendNet::getTargetDevice() const CV_NOEXCEPT
 {
-    return targetDevice;
+    return targetDevice == InferenceEngine::TargetDevice::eFPGA ?
+           InferenceEngine::TargetDevice::eHETERO : targetDevice;
 }

 InferenceEngine::StatusCode InfEngineBackendNet::setBatchSize(const size_t) CV_NOEXCEPT
@@ -416,6 +426,8 @@ void InfEngineBackendNet::init(int targetId)
        InferenceEngine::OutputsDataMap unconnectedOuts;
        for (const auto& l : layers)
        {
+            if (l->type == "Input")
+                continue;
            // Add all outputs.
            for (const InferenceEngine::DataPtr& out : l->outData)
            {
@@ -466,6 +478,11 @@ void InfEngineBackendNet::init(int targetId)
        setPrecision(InferenceEngine::Precision::FP16);
        setTargetDevice(InferenceEngine::TargetDevice::eMYRIAD); break;
    }
+    case DNN_TARGET_FPGA:
+    {
+        setPrecision(InferenceEngine::Precision::FP16);
+        setTargetDevice(InferenceEngine::TargetDevice::eFPGA); break;
+    }
    default:
        CV_Error(Error::StsError, format("Unknown target identifier: %d", targetId));
    }
@@ -489,10 +506,15 @@ void InfEngineBackendNet::initPlugin(InferenceEngine::ICNNNetwork& net)
        }
        else
        {
-            enginePtr = InferenceEngine::PluginDispatcher({""}).getSuitablePlugin(targetDevice);
+            auto dispatcher = InferenceEngine::PluginDispatcher({""});
+            if (targetDevice == InferenceEngine::TargetDevice::eFPGA)
+                enginePtr = dispatcher.getPluginByDevice("HETERO:FPGA,CPU");
+            else
+                enginePtr = dispatcher.getSuitablePlugin(targetDevice);
            sharedPlugins[targetDevice] = enginePtr;

-            if (targetDevice == InferenceEngine::TargetDevice::eCPU)
+            if (targetDevice == InferenceEngine::TargetDevice::eCPU ||
+                targetDevice == InferenceEngine::TargetDevice::eFPGA)
            {
                std::string suffixes[] = {"_avx2", "_sse4", ""};
                bool haveFeature[] = {
@@ -26,10 +26,11 @@
 #define INF_ENGINE_RELEASE_2018R2 2018020000
 #define INF_ENGINE_RELEASE_2018R3 2018030000
 #define INF_ENGINE_RELEASE_2018R4 2018040000
+#define INF_ENGINE_RELEASE_2018R5 2018050000

 #ifndef INF_ENGINE_RELEASE
-#warning("IE version have not been provided via command-line. Using 2018R4 by default")
-#define INF_ENGINE_RELEASE INF_ENGINE_RELEASE_2018R4
+#warning("IE version have not been provided via command-line. Using 2018R5 by default")
+#define INF_ENGINE_RELEASE INF_ENGINE_RELEASE_2018R5
 #endif

 #define INF_ENGINE_VER_MAJOR_GT(ver) (((INF_ENGINE_RELEASE) / 10000) > ((ver) / 10000))
@@ -68,6 +69,8 @@ public:

    virtual InferenceEngine::InputInfo::Ptr getInput(const std::string &inputName) const CV_NOEXCEPT;

+    virtual InferenceEngine::StatusCode serialize(const std::string &xmlPath, const std::string &binPath, InferenceEngine::ResponseDesc* resp) const CV_NOEXCEPT;
+
    virtual void getName(char *pName, size_t len) CV_NOEXCEPT;

    virtual void getName(char *pName, size_t len) const CV_NOEXCEPT;
@@ -134,6 +137,9 @@ private:
    InferenceEngine::InferRequest infRequest;
    // In case of models from Model Optimizer we need to manage their lifetime.
    InferenceEngine::CNNNetwork netOwner;
+    // There is no way to check if netOwner is initialized or not so we use
+    // a separate flag to determine if the model has been loaded from IR.
+    bool hasNetOwner;

    std::string name;

@@ -939,7 +939,7 @@ void TFImporter::populateNet(Net dstNet)
            if (getDataLayout(name, data_layouts) == DATA_LAYOUT_UNKNOWN)
                data_layouts[name] = DATA_LAYOUT_NHWC;
        }
-        else if (type == "BiasAdd" || type == "Add")
+        else if (type == "BiasAdd" || type == "Add" || type == "Sub")
        {
            bool haveConst = false;
            for(int ii = 0; !haveConst && ii < layer.input_size(); ++ii)
@@ -953,6 +953,8 @@ void TFImporter::populateNet(Net dstNet)
            {
                Mat values = getTensorContent(getConstBlob(layer, value_id));
                CV_Assert(values.type() == CV_32FC1);
+                if (type == "Sub")
+                    values *= -1.0f;

                int id;
                if (values.total() == 1)  // is a scalar.
@@ -973,6 +975,12 @@ void TFImporter::populateNet(Net dstNet)
            else
            {
                layerParams.set("operation", "sum");
+                if (type == "Sub")
+                {
+                    static float subCoeffs[] = {1.f, -1.f};
+                    layerParams.set("coeff", DictValue::arrayReal<float*>(subCoeffs, 2));
+                }
+
                int id = dstNet.addLayer(name, "Eltwise", layerParams);
                layer_id[name] = id;

@@ -985,36 +993,6 @@ void TFImporter::populateNet(Net dstNet)
                }
            }
        }
-        else if (type == "Sub")
-        {
-            bool haveConst = false;
-            for(int ii = 0; !haveConst && ii < layer.input_size(); ++ii)
-            {
-                Pin input = parsePin(layer.input(ii));
-                haveConst = value_id.find(input.name) != value_id.end();
-            }
-            CV_Assert(haveConst);
-
-            Mat values = getTensorContent(getConstBlob(layer, value_id));
-            CV_Assert(values.type() == CV_32FC1);
-            values *= -1.0f;
-
-            int id;
-            if (values.total() == 1)  // is a scalar.
-            {
-                layerParams.set("shift", values.at<float>(0));
-                id = dstNet.addLayer(name, "Power", layerParams);
-            }
-            else  // is a vector
-            {
-                layerParams.blobs.resize(1, values);
-                id = dstNet.addLayer(name, "Shift", layerParams);
-            }
-            layer_id[name] = id;
-
-            // one input only
-            connect(layer_id, dstNet, parsePin(layer.input(0)), id, 0);
-        }
        else if (type == "MatMul")
        {
            CV_Assert(layer.input_size() == 2);
@@ -1266,14 +1244,31 @@ void TFImporter::populateNet(Net dstNet)
                axis = toNCHW(axis);
            layerParams.set("axis", axis);

-            int id = dstNet.addLayer(name, "Concat", layerParams);
-            layer_id[name] = id;
-
-
+            // input(0) or input(n-1) is concat_dim
            int from = (type == "Concat" ? 1 : 0);
            int to = (type == "Concat" ? layer.input_size() : layer.input_size() - 1);

-            // input(0) or input(n-1) is concat_dim
+            for (int ii = from; ii < to; ii++)
+            {
+                Pin inp = parsePin(layer.input(ii));
+                if (layer_id.find(inp.name) == layer_id.end())
+                {
+                    // There are constant inputs.
+                    LayerParams lp;
+                    lp.name = inp.name;
+                    lp.type = "Const";
+                    lp.blobs.resize(1);
+                    blobFromTensor(getConstBlob(layer, value_id, ii), lp.blobs.back());
+                    CV_Assert_N(!lp.blobs[0].empty(), lp.blobs[0].type() == CV_32F);
+
+                    int constInpId = dstNet.addLayer(lp.name, lp.type, lp);
+                    layer_id[lp.name] = constInpId;
+                }
+            }
+
+            int id = dstNet.addLayer(name, "Concat", layerParams);
+            layer_id[name] = id;
+
            for (int ii = from; ii < to; ii++)
            {
                Pin inp = parsePin(layer.input(ii));
@@ -129,13 +129,15 @@ struct TorchImporter
    Module *rootModule;
    Module *curModule;
    int moduleCounter;
+    bool testPhase;

-    TorchImporter(String filename, bool isBinary)
+    TorchImporter(String filename, bool isBinary, bool evaluate)
    {
        CV_TRACE_FUNCTION();

        rootModule = curModule = NULL;
        moduleCounter = 0;
+        testPhase = evaluate;

        file = cv::Ptr<THFile>(THDiskFile_new(filename, "r", 0), THFile_free);
        CV_Assert(file && THFile_isOpened(file));
@@ -680,7 +682,8 @@ struct TorchImporter
                    layerParams.blobs.push_back(tensorParams["bias"].second);
                }

-                if (nnName == "InstanceNormalization")
+                bool trainPhase = scalarParams.get<bool>("train", false);
+                if (nnName == "InstanceNormalization" || (trainPhase && !testPhase))
                {
                    cv::Ptr<Module> mvnModule(new Module(nnName));
                    mvnModule->apiType = "MVN";
@@ -1243,18 +1246,18 @@ struct TorchImporter

 Mat readTorchBlob(const String &filename, bool isBinary)
 {
-    TorchImporter importer(filename, isBinary);
+    TorchImporter importer(filename, isBinary, true);
    importer.readObject();
    CV_Assert(importer.tensors.size() == 1);

    return importer.tensors.begin()->second;
 }

-Net readNetFromTorch(const String &model, bool isBinary)
+Net readNetFromTorch(const String &model, bool isBinary, bool evaluate)
 {
    CV_TRACE_FUNCTION();

-    TorchImporter importer(model, isBinary);
+    TorchImporter importer(model, isBinary, evaluate);
    Net net;
    importer.populateNet(net);
    return net;
@@ -36,7 +36,6 @@ protected:
    void recordCommandBuffer(void* push_constants = NULL, size_t push_constants_size = 0);
    void runCommandBuffer();

-    const Context* ctx_;
    VkPipeline pipeline_;
    VkCommandBuffer cmd_buffer_;
    VkDescriptorPool descriptor_pool_;
@@ -59,7 +59,7 @@ private:
    int avg_pool_padded_area_;
    int need_mask_;
    PaddingMode padding_mode_;
-    int activation_;
+    //int activation_;
    PoolShaderConfig config_;
 };

@@ -39,9 +39,6 @@ enum PaddingMode { kPaddingModeSame, kPaddingModeValid, kPaddingModeCaffe, kPadd
 enum FusedActivationType { kNone, kRelu, kRelu1, kRelu6, kActivationNum };
 typedef std::vector<int> Shape;

-/* context APIs */
-bool initPerThread();
-void deinitPerThread();
 bool isAvailable();

 #endif // HAVE_VULKAN
@@ -18,7 +18,7 @@ static uint32_t findMemoryType(uint32_t memoryTypeBits, VkMemoryPropertyFlags pr
 {
    VkPhysicalDeviceMemoryProperties memoryProperties;

-    vkGetPhysicalDeviceMemoryProperties(getPhysicalDevice(), &memoryProperties);
+    vkGetPhysicalDeviceMemoryProperties(kPhysicalDevice, &memoryProperties);

    for (uint32_t i = 0; i < memoryProperties.memoryTypeCount; ++i) {
        if ((memoryTypeBits & (1 << i)) &&
@@ -29,6 +29,11 @@
 namespace cv { namespace dnn { namespace vkcom {

 #ifdef HAVE_VULKAN
+extern VkPhysicalDevice kPhysicalDevice;
+extern VkDevice kDevice;
+extern VkQueue kQueue;
+extern VkCommandPool kCmdPool;
+extern cv::Mutex kContextMtx;

 enum ShapeIdx
 {
@@ -42,7 +47,7 @@ enum ShapeIdx
 { \
        if (f != VK_SUCCESS) \
        { \
-            CV_LOG_ERROR(NULL, "Vulkan check failed, result = " << f); \
+            CV_LOG_ERROR(NULL, "Vulkan check failed, result = " << (int)f); \
            CV_Error(Error::StsError, "Vulkan check failed"); \
        } \
 }
@@ -0,0 +1,301 @@
+// This file is part of OpenCV project.
+// It is subject to the license terms in the LICENSE file found in the top-level directory
+// of this distribution and at http://opencv.org/license.html.
+//
+// Copyright (C) 2018, Intel Corporation, all rights reserved.
+// Third party copyrights are property of their respective owners.
+
+#include "../../precomp.hpp"
+#include "../vulkan/vk_loader.hpp"
+#include "common.hpp"
+#include "context.hpp"
+
+namespace cv { namespace dnn { namespace vkcom {
+
+#ifdef HAVE_VULKAN
+
+std::shared_ptr<Context> kCtx;
+bool enableValidationLayers = false;
+VkInstance kInstance;
+VkPhysicalDevice kPhysicalDevice;
+VkDevice kDevice;
+VkQueue kQueue;
+VkCommandPool kCmdPool;
+VkDebugReportCallbackEXT kDebugReportCallback;
+uint32_t kQueueFamilyIndex;
+std::vector<const char *> kEnabledLayers;
+std::map<std::string, std::vector<uint32_t>> kShaders;
+cv::Mutex kContextMtx;
+
+static uint32_t getComputeQueueFamilyIndex()
+{
+    uint32_t queueFamilyCount;
+
+    vkGetPhysicalDeviceQueueFamilyProperties(kPhysicalDevice, &queueFamilyCount, NULL);
+
+    std::vector<VkQueueFamilyProperties> queueFamilies(queueFamilyCount);
+    vkGetPhysicalDeviceQueueFamilyProperties(kPhysicalDevice,
+                                             &queueFamilyCount,
+                                             queueFamilies.data());
+
+    uint32_t i = 0;
+    for (; i < queueFamilies.size(); ++i)
+    {
+        VkQueueFamilyProperties props = queueFamilies[i];
+
+        if (props.queueCount > 0 && (props.queueFlags & VK_QUEUE_COMPUTE_BIT))
+        {
+            break;
+        }
+    }
+
+    if (i == queueFamilies.size())
+    {
+        throw std::runtime_error("could not find a queue family that supports operations");
+    }
+
+    return i;
+}
+
+bool checkExtensionAvailability(const char *extension_name,
+                                const std::vector<VkExtensionProperties> &available_extensions)
+{
+    for( size_t i = 0; i < available_extensions.size(); ++i )
+    {
+      if( strcmp( available_extensions[i].extensionName, extension_name ) == 0 )
+      {
+        return true;
+      }
+    }
+    return false;
+}
+
+VKAPI_ATTR VkBool32 VKAPI_CALL debugReportCallbackFn(
+        VkDebugReportFlagsEXT                       flags,
+        VkDebugReportObjectTypeEXT                  objectType,
+        uint64_t                                    object,
+        size_t                                      location,
+        int32_t                                     messageCode,
+        const char*                                 pLayerPrefix,
+        const char*                                 pMessage,
+        void*                                       pUserData)
+{
+        std::cout << "Debug Report: " << pLayerPrefix << ":" << pMessage << std::endl;
+        return VK_FALSE;
+}
+
+// internally used
+void createContext()
+{
+    cv::AutoLock lock(kContextMtx);
+    if (!kCtx)
+    {
+        kCtx.reset(new Context());
+    }
+}
+
+bool isAvailable()
+{
+    try
+    {
+        createContext();
+    }
+    catch (const cv::Exception& e)
+    {
+        CV_LOG_ERROR(NULL, "Failed to init Vulkan environment. " << e.what());
+        return false;
+    }
+
+    return true;
+}
+
+Context::Context()
+{
+    if(!loadVulkanLibrary())
+    {
+        CV_Error(Error::StsError, "loadVulkanLibrary failed");
+        return;
+    }
+    else if (!loadVulkanEntry())
+    {
+        CV_Error(Error::StsError, "loadVulkanEntry failed");
+        return;
+    }
+    else if (!loadVulkanGlobalFunctions())
+    {
+        CV_Error(Error::StsError, "loadVulkanGlobalFunctions failed");
+        return;
+    }
+
+    // create VkInstance, VkPhysicalDevice
+    std::vector<const char *> enabledExtensions;
+    if (enableValidationLayers)
+    {
+        uint32_t layerCount;
+        vkEnumerateInstanceLayerProperties(&layerCount, NULL);
+
+        std::vector<VkLayerProperties> layerProperties(layerCount);
+        vkEnumerateInstanceLayerProperties(&layerCount, layerProperties.data());
+
+        bool foundLayer = false;
+        for (VkLayerProperties prop : layerProperties)
+        {
+            if (strcmp("VK_LAYER_LUNARG_standard_validation", prop.layerName) == 0)
+            {
+                foundLayer = true;
+                break;
+            }
+        }
+
+        if (!foundLayer)
+        {
+            throw std::runtime_error("Layer VK_LAYER_LUNARG_standard_validation not supported\n");
+        }
+        kEnabledLayers.push_back("VK_LAYER_LUNARG_standard_validation");
+
+        uint32_t extensionCount;
+
+        vkEnumerateInstanceExtensionProperties(nullptr, &extensionCount, NULL);
+        std::vector<VkExtensionProperties> extensionProperties(extensionCount);
+        vkEnumerateInstanceExtensionProperties(nullptr, &extensionCount, extensionProperties.data());
+
+        bool foundExtension = false;
+        for (VkExtensionProperties prop : extensionProperties)
+        {
+            if (strcmp(VK_EXT_DEBUG_REPORT_EXTENSION_NAME, prop.extensionName) == 0)
+            {
+                foundExtension = true;
+                break;
+            }
+        }
+
+        if (!foundExtension) {
+            throw std::runtime_error("Extension VK_EXT_DEBUG_REPORT_EXTENSION_NAME not supported\n");
+        }
+        enabledExtensions.push_back(VK_EXT_DEBUG_REPORT_EXTENSION_NAME);
+    }
+
+    VkApplicationInfo applicationInfo = {};
+    applicationInfo.sType = VK_STRUCTURE_TYPE_APPLICATION_INFO;
+    applicationInfo.pApplicationName = "VkCom Library";
+    applicationInfo.applicationVersion = 0;
+    applicationInfo.pEngineName = "vkcom";
+    applicationInfo.engineVersion = 0;
+    applicationInfo.apiVersion = VK_API_VERSION_1_0;;
+
+    VkInstanceCreateInfo createInfo = {};
+    createInfo.sType = VK_STRUCTURE_TYPE_INSTANCE_CREATE_INFO;
+    createInfo.flags = 0;
+    createInfo.pApplicationInfo = &applicationInfo;
+
+    // Give our desired layers and extensions to vulkan.
+    createInfo.enabledLayerCount = kEnabledLayers.size();
+    createInfo.ppEnabledLayerNames = kEnabledLayers.data();
+    createInfo.enabledExtensionCount = enabledExtensions.size();
+    createInfo.ppEnabledExtensionNames = enabledExtensions.data();
+
+    VK_CHECK_RESULT(vkCreateInstance(&createInfo, NULL, &kInstance));
+
+    if (!loadVulkanFunctions(kInstance))
+    {
+        CV_Error(Error::StsError, "loadVulkanFunctions failed");
+        return;
+    }
+
+    if (enableValidationLayers && vkCreateDebugReportCallbackEXT)
+    {
+        VkDebugReportCallbackCreateInfoEXT createInfo = {};
+        createInfo.sType = VK_STRUCTURE_TYPE_DEBUG_REPORT_CALLBACK_CREATE_INFO_EXT;
+        createInfo.flags = VK_DEBUG_REPORT_ERROR_BIT_EXT |
+                           VK_DEBUG_REPORT_WARNING_BIT_EXT |
+                           VK_DEBUG_REPORT_PERFORMANCE_WARNING_BIT_EXT;
+        createInfo.pfnCallback = &debugReportCallbackFn;
+
+        // Create and register callback.
+        VK_CHECK_RESULT(vkCreateDebugReportCallbackEXT(kInstance, &createInfo,
+                                                       NULL, &kDebugReportCallback));
+    }
+
+    // find physical device
+    uint32_t deviceCount;
+    vkEnumeratePhysicalDevices(kInstance, &deviceCount, NULL);
+    if (deviceCount == 0)
+    {
+        throw std::runtime_error("could not find a device with vulkan support");
+    }
+
+    std::vector<VkPhysicalDevice> devices(deviceCount);
+    vkEnumeratePhysicalDevices(kInstance, &deviceCount, devices.data());
+
+    for (VkPhysicalDevice device : devices)
+    {
+        if (true)
+        {
+            kPhysicalDevice = device;
+            break;
+        }
+    }
+
+    kQueueFamilyIndex = getComputeQueueFamilyIndex();
+
+    // create device, queue, command pool
+    VkDeviceQueueCreateInfo queueCreateInfo = {};
+    queueCreateInfo.sType = VK_STRUCTURE_TYPE_DEVICE_QUEUE_CREATE_INFO;
+    queueCreateInfo.queueFamilyIndex = kQueueFamilyIndex;
+    queueCreateInfo.queueCount = 1; // create one queue in this family. We don't need more.
+    float queuePriorities = 1.0;  // we only have one queue, so this is not that imporant.
+    queueCreateInfo.pQueuePriorities = &queuePriorities;
+
+    VkDeviceCreateInfo deviceCreateInfo = {};
+
+    // Specify any desired device features here. We do not need any for this application, though.
+    VkPhysicalDeviceFeatures deviceFeatures = {};
+
+    deviceCreateInfo.sType = VK_STRUCTURE_TYPE_DEVICE_CREATE_INFO;
+    deviceCreateInfo.enabledLayerCount = kEnabledLayers.size();
+    deviceCreateInfo.ppEnabledLayerNames = kEnabledLayers.data();
+    deviceCreateInfo.pQueueCreateInfos = &queueCreateInfo;
+    deviceCreateInfo.queueCreateInfoCount = 1;
+    deviceCreateInfo.pEnabledFeatures = &deviceFeatures;
+
+    VK_CHECK_RESULT(vkCreateDevice(kPhysicalDevice, &deviceCreateInfo, NULL, &kDevice));
+
+    // Get a handle to the only member of the queue family.
+    vkGetDeviceQueue(kDevice, kQueueFamilyIndex, 0, &kQueue);
+
+    // create command pool
+    VkCommandPoolCreateInfo commandPoolCreateInfo = {};
+    commandPoolCreateInfo.sType = VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO;
+    commandPoolCreateInfo.flags = VK_COMMAND_POOL_CREATE_RESET_COMMAND_BUFFER_BIT;
+    // the queue family of this command pool. All command buffers allocated from this command pool,
+    // must be submitted to queues of this family ONLY.
+    commandPoolCreateInfo.queueFamilyIndex = kQueueFamilyIndex;
+    VK_CHECK_RESULT(vkCreateCommandPool(kDevice, &commandPoolCreateInfo, NULL, &kCmdPool));
+}
+
+Context::~Context()
+{
+    vkDestroyCommandPool(kDevice, kCmdPool, NULL);
+    vkDestroyDevice(kDevice, NULL);
+
+    if (enableValidationLayers) {
+        auto func = (PFN_vkDestroyDebugReportCallbackEXT)
+            vkGetInstanceProcAddr(kInstance, "vkDestroyDebugReportCallbackEXT");
+        if (func == nullptr)
+        {
+            CV_LOG_FATAL(NULL, "Could not load vkDestroyDebugReportCallbackEXT");
+        }
+        else
+        {
+            func(kInstance, kDebugReportCallback, NULL);
+        }
+    }
+    kShaders.clear();
+    vkDestroyInstance(kInstance, NULL);
+
+    return;
+}
+
+#endif // HAVE_VULKAN
+
+}}} // namespace cv::dnn::vkcom
@@ -7,21 +7,20 @@

 #ifndef OPENCV_DNN_VKCOM_CONTEXT_HPP
 #define OPENCV_DNN_VKCOM_CONTEXT_HPP
-#include "common.hpp"

 namespace cv { namespace dnn { namespace vkcom {

 #ifdef HAVE_VULKAN

-struct Context
+class Context
 {
-    VkDevice device;
-    VkQueue queue;
-    VkCommandPool cmd_pool;
-    std::map<std::string, VkShaderModule> shader_modules;
-    int ref;
+public:
+    Context();
+    ~Context();
 };

+void createContext();
+
 #endif // HAVE_VULKAN

 }}} // namespace cv::dnn::vkcom
@@ -16,8 +16,8 @@ namespace cv { namespace dnn { namespace vkcom {

 OpBase::OpBase()
 {
-    ctx_ = getContext();
-    device_ = ctx_->device;
+    createContext();
+    device_ = kDevice;
    pipeline_ = VK_NULL_HANDLE;
    cmd_buffer_ = VK_NULL_HANDLE;
    descriptor_pool_ = VK_NULL_HANDLE;
@@ -45,7 +45,9 @@ void OpBase::initVulkanThing(int buffer_num)

 void OpBase::createDescriptorSetLayout(int buffer_num)
 {
-    VkDescriptorSetLayoutBinding bindings[buffer_num] = {};
+    if (buffer_num <= 0)
+        return;
+    std::vector<VkDescriptorSetLayoutBinding> bindings(buffer_num);
    for (int i = 0; i < buffer_num; i++)
    {
        bindings[i].binding = i;
@@ -56,7 +58,7 @@ void OpBase::createDescriptorSetLayout(int buffer_num)
    VkDescriptorSetLayoutCreateInfo info = {};
    info.sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO;
    info.bindingCount = buffer_num;
-    info.pBindings = bindings;
+    info.pBindings = &bindings[0];
    VK_CHECK_RESULT(vkCreateDescriptorSetLayout(device_, &info, NULL, &descriptor_set_layout_));
 }

@@ -139,7 +141,7 @@ void OpBase::createCommandBuffer()
 {
    VkCommandBufferAllocateInfo info = {};
    info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO;
-    info.commandPool = ctx_->cmd_pool;
+    info.commandPool = kCmdPool;
    info.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY;
    info.commandBufferCount = 1;
    VK_CHECK_RESULT(vkAllocateCommandBuffers(device_, &info, &cmd_buffer_));
@@ -150,6 +152,7 @@ void OpBase::recordCommandBuffer(void* push_constants, size_t push_constants_siz
    VkCommandBufferBeginInfo beginInfo = {};
    beginInfo.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO;
    beginInfo.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT;
+    cv::AutoLock lock(kContextMtx);
    VK_CHECK_RESULT(vkBeginCommandBuffer(cmd_buffer_, &beginInfo));
    if (push_constants)
        vkCmdPushConstants(cmd_buffer_, pipeline_layout_,
@@ -176,7 +179,10 @@ void OpBase::runCommandBuffer()
    fence_create_info_.flags = 0;

    VK_CHECK_RESULT(vkCreateFence(device_, &fence_create_info_, NULL, &fence));
-    VK_CHECK_RESULT(vkQueueSubmit(ctx_->queue, 1, &submit_info, fence));
+    {
+        cv::AutoLock lock(kContextMtx);
+        VK_CHECK_RESULT(vkQueueSubmit(kQueue, 1, &submit_info, fence));
+    }
    VK_CHECK_RESULT(vkWaitForFences(device_, 1, &fence, VK_TRUE, 100000000000));
    vkDestroyFence(device_, fence, NULL);
 }
@@ -15,15 +15,15 @@ namespace cv { namespace dnn { namespace vkcom {

 Tensor::Tensor(Format fmt) : size_in_byte_(0), format_(fmt)
 {
-    Context *ctx = getContext();
-    device_ = ctx->device;
+    createContext();
+    device_ = kDevice;
 }

 Tensor::Tensor(const char* data, std::vector<int>& shape, Format fmt)
               : size_in_byte_(0), format_(fmt)
 {
-    Context *ctx = getContext();
-    device_ = ctx->device;
+    createContext();
+    device_ = kDevice;
    reshape(data, shape);
 }

@@ -1,403 +0,0 @@
-// This file is part of OpenCV project.
-// It is subject to the license terms in the LICENSE file found in the top-level directory
-// of this distribution and at http://opencv.org/license.html.
-//
-// Copyright (C) 2018, Intel Corporation, all rights reserved.
-// Third party copyrights are property of their respective owners.
-
-#include "../../precomp.hpp"
-#include "common.hpp"
-#include "internal.hpp"
-#include "../include/op_conv.hpp"
-#include "../include/op_pool.hpp"
-#include "../include/op_lrn.hpp"
-#include "../include/op_concat.hpp"
-#include "../include/op_softmax.hpp"
-#include "../vulkan/vk_loader.hpp"
-
-namespace cv { namespace dnn { namespace vkcom {
-
-#ifdef HAVE_VULKAN
-
-static bool enableValidationLayers = false;
-static VkInstance kInstance;
-static VkPhysicalDevice kPhysicalDevice;
-static VkDebugReportCallbackEXT kDebugReportCallback;
-static uint32_t kQueueFamilyIndex;
-std::vector<const char *> kEnabledLayers;
-typedef std::map<std::thread::id, Context*> IdToContextMap;
-IdToContextMap kThreadResources;
-static std::map<std::string, std::vector<uint32_t>> kShaders;
-static int init_count = 0;
-static bool init();
-static void release();
-static uint32_t getComputeQueueFamilyIndex();
-static bool checkExtensionAvailability(const char *extension_name,
-                                       const std::vector<VkExtensionProperties>
-                                       &available_extensions);
-static VKAPI_ATTR VkBool32 VKAPI_CALL debugReportCallbackFn(
-       VkDebugReportFlagsEXT                       flags,
-       VkDebugReportObjectTypeEXT                  objectType,
-       uint64_t                                    object,
-       size_t                                      location,
-       int32_t                                     messageCode,
-       const char*                                 pLayerPrefix,
-       const char*                                 pMessage,
-       void*                                       pUserData);
-
-static void setContext(Context* ctx)
-{
-    cv::AutoLock lock(getInitializationMutex());
-    std::thread::id tid = std::this_thread::get_id();
-    if (kThreadResources.find(tid) != kThreadResources.end())
-    {
-        return;
-    }
-    kThreadResources.insert(std::pair<std::thread::id, Context*>(tid, ctx));
-}
-
-Context* getContext()
-{
-    Context* ctx = NULL;
-
-    cv::AutoLock lock(getInitializationMutex());
-    std::thread::id tid = std::this_thread::get_id();
-    IdToContextMap::iterator it = kThreadResources.find(tid);
-    if (it != kThreadResources.end())
-    {
-        ctx = it->second;
-    }
-    return ctx;
-}
-
-static void removeContext()
-{
-    cv::AutoLock lock(getInitializationMutex());
-    std::thread::id tid = std::this_thread::get_id();
-    IdToContextMap::iterator it = kThreadResources.find(tid);
-    if (it == kThreadResources.end())
-    {
-        return;
-    }
-    kThreadResources.erase(it);
-}
-
-bool initPerThread()
-{
-    VkDevice device;
-    VkQueue queue;
-    VkCommandPool cmd_pool;
-
-    VKCOM_CHECK_BOOL_RET_VAL(init(), false);
-    Context* ctx = getContext();
-    if (ctx)
-    {
-        ctx->ref++;
-        return true;
-    }
-
-    // create device, queue, command pool
-    VkDeviceQueueCreateInfo queueCreateInfo = {};
-    queueCreateInfo.sType = VK_STRUCTURE_TYPE_DEVICE_QUEUE_CREATE_INFO;
-    queueCreateInfo.queueFamilyIndex = kQueueFamilyIndex;
-    queueCreateInfo.queueCount = 1; // create one queue in this family. We don't need more.
-    float queuePriorities = 1.0;  // we only have one queue, so this is not that imporant.
-    queueCreateInfo.pQueuePriorities = &queuePriorities;
-
-    VkDeviceCreateInfo deviceCreateInfo = {};
-
-    // Specify any desired device features here. We do not need any for this application, though.
-    VkPhysicalDeviceFeatures deviceFeatures = {};
-
-    deviceCreateInfo.sType = VK_STRUCTURE_TYPE_DEVICE_CREATE_INFO;
-    deviceCreateInfo.enabledLayerCount = kEnabledLayers.size();
-    deviceCreateInfo.ppEnabledLayerNames = kEnabledLayers.data();
-    deviceCreateInfo.pQueueCreateInfos = &queueCreateInfo;
-    deviceCreateInfo.queueCreateInfoCount = 1;
-    deviceCreateInfo.pEnabledFeatures = &deviceFeatures;
-
-    VK_CHECK_RESULT(vkCreateDevice(kPhysicalDevice, &deviceCreateInfo, NULL, &device));
-
-    // Get a handle to the only member of the queue family.
-    vkGetDeviceQueue(device, kQueueFamilyIndex, 0, &queue);
-
-    // create command pool
-    VkCommandPoolCreateInfo commandPoolCreateInfo = {};
-    commandPoolCreateInfo.sType = VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO;
-    commandPoolCreateInfo.flags = VK_COMMAND_POOL_CREATE_RESET_COMMAND_BUFFER_BIT;
-    // the queue family of this command pool. All command buffers allocated from this command pool,
-    // must be submitted to queues of this family ONLY.
-    commandPoolCreateInfo.queueFamilyIndex = kQueueFamilyIndex;
-    VK_CHECK_RESULT(vkCreateCommandPool(device, &commandPoolCreateInfo, NULL, &cmd_pool));
-
-    ctx = new Context();
-    ctx->device = device;
-    ctx->queue = queue;
-    ctx->cmd_pool = cmd_pool;
-    ctx->ref = 1;
-    setContext(ctx);
-    return true;
-}
-
-void deinitPerThread()
-{
-    Context* ctx = getContext();
-    if (ctx == NULL)
-    {
-        release();
-        return;
-    }
-
-    if (ctx->ref > 1)
-    {
-        ctx->ref--;
-    }
-    else if (ctx->ref == 1)
-    {
-        for(auto &kv: ctx->shader_modules)
-        {
-            vkDestroyShaderModule(ctx->device, kv.second, NULL);
-        }
-        ctx->shader_modules.clear();
-        vkDestroyCommandPool(ctx->device, ctx->cmd_pool, NULL);
-        vkDestroyDevice(ctx->device, NULL);
-        removeContext();
-        delete ctx;
-    }
-    else
-        CV_Assert(0);
-    release();
-}
-
-static bool init()
-{
-    cv::AutoLock lock(getInitializationMutex());
-
-    if (init_count == 0)
-    {
-        if(!loadVulkanLibrary())
-        {
-            return false;
-        }
-        else if (!loadVulkanEntry())
-        {
-            return false;
-        }
-        else if (!loadVulkanGlobalFunctions())
-        {
-            return false;
-        }
-
-        // create VkInstance, VkPhysicalDevice
-        std::vector<const char *> enabledExtensions;
-        if (enableValidationLayers)
-        {
-            uint32_t layerCount;
-            vkEnumerateInstanceLayerProperties(&layerCount, NULL);
-
-            std::vector<VkLayerProperties> layerProperties(layerCount);
-            vkEnumerateInstanceLayerProperties(&layerCount, layerProperties.data());
-
-            bool foundLayer = false;
-            for (VkLayerProperties prop : layerProperties)
-            {
-                if (strcmp("VK_LAYER_LUNARG_standard_validation", prop.layerName) == 0)
-                {
-                    foundLayer = true;
-                    break;
-                }
-            }
-
-            if (!foundLayer)
-            {
-                throw std::runtime_error("Layer VK_LAYER_LUNARG_standard_validation not supported\n");
-            }
-            kEnabledLayers.push_back("VK_LAYER_LUNARG_standard_validation");
-
-            uint32_t extensionCount;
-
-            vkEnumerateInstanceExtensionProperties(nullptr, &extensionCount, NULL);
-            std::vector<VkExtensionProperties> extensionProperties(extensionCount);
-            vkEnumerateInstanceExtensionProperties(nullptr, &extensionCount, extensionProperties.data());
-
-            bool foundExtension = false;
-            for (VkExtensionProperties prop : extensionProperties)
-            {
-                if (strcmp(VK_EXT_DEBUG_REPORT_EXTENSION_NAME, prop.extensionName) == 0)
-                {
-                    foundExtension = true;
-                    break;
-                }
-            }
-
-            if (!foundExtension) {
-                throw std::runtime_error("Extension VK_EXT_DEBUG_REPORT_EXTENSION_NAME not supported\n");
-            }
-            enabledExtensions.push_back(VK_EXT_DEBUG_REPORT_EXTENSION_NAME);
-        }
-
-        VkApplicationInfo applicationInfo = {};
-        applicationInfo.sType = VK_STRUCTURE_TYPE_APPLICATION_INFO;
-        applicationInfo.pApplicationName = "VkCom Library";
-        applicationInfo.applicationVersion = 0;
-        applicationInfo.pEngineName = "vkcom";
-        applicationInfo.engineVersion = 0;
-        applicationInfo.apiVersion = VK_API_VERSION_1_0;;
-
-        VkInstanceCreateInfo createInfo = {};
-        createInfo.sType = VK_STRUCTURE_TYPE_INSTANCE_CREATE_INFO;
-        createInfo.flags = 0;
-        createInfo.pApplicationInfo = &applicationInfo;
-
-        // Give our desired layers and extensions to vulkan.
-        createInfo.enabledLayerCount = kEnabledLayers.size();
-        createInfo.ppEnabledLayerNames = kEnabledLayers.data();
-        createInfo.enabledExtensionCount = enabledExtensions.size();
-        createInfo.ppEnabledExtensionNames = enabledExtensions.data();
-
-        VK_CHECK_RESULT(vkCreateInstance(&createInfo, NULL, &kInstance));
-
-        if (!loadVulkanFunctions(kInstance))
-        {
-            return false;
-        }
-
-        if (enableValidationLayers && vkCreateDebugReportCallbackEXT)
-        {
-            VkDebugReportCallbackCreateInfoEXT createInfo = {};
-            createInfo.sType = VK_STRUCTURE_TYPE_DEBUG_REPORT_CALLBACK_CREATE_INFO_EXT;
-            createInfo.flags = VK_DEBUG_REPORT_ERROR_BIT_EXT |
-                               VK_DEBUG_REPORT_WARNING_BIT_EXT |
-                               VK_DEBUG_REPORT_PERFORMANCE_WARNING_BIT_EXT;
-            createInfo.pfnCallback = &debugReportCallbackFn;
-
-            // Create and register callback.
-            VK_CHECK_RESULT(vkCreateDebugReportCallbackEXT(kInstance, &createInfo,
-                                                           NULL, &kDebugReportCallback));
-        }
-
-        // find physical device
-        uint32_t deviceCount;
-        vkEnumeratePhysicalDevices(kInstance, &deviceCount, NULL);
-        if (deviceCount == 0)
-        {
-            throw std::runtime_error("could not find a device with vulkan support");
-        }
-
-        std::vector<VkPhysicalDevice> devices(deviceCount);
-        vkEnumeratePhysicalDevices(kInstance, &deviceCount, devices.data());
-
-        for (VkPhysicalDevice device : devices)
-        {
-            if (true)
-            {
-                kPhysicalDevice = device;
-                break;
-            }
-        }
-
-        kQueueFamilyIndex = getComputeQueueFamilyIndex();
-    }
-
-    init_count++;
-    return true;
-}
-
-static void release()
-{
-    cv::AutoLock lock(getInitializationMutex());
-    if (init_count == 0)
-    {
-        return;
-    }
-
-    init_count--;
-    if (init_count == 0)
-    {
-        if (enableValidationLayers) {
-            auto func = (PFN_vkDestroyDebugReportCallbackEXT)
-                vkGetInstanceProcAddr(kInstance, "vkDestroyDebugReportCallbackEXT");
-            if (func == nullptr) {
-                throw std::runtime_error("Could not load vkDestroyDebugReportCallbackEXT");
-            }
-            func(kInstance, kDebugReportCallback, NULL);
-        }
-        kShaders.clear();
-        vkDestroyInstance(kInstance, NULL);
-    }
-
-    return;
-}
-
-// Returns the index of a queue family that supports compute operations.
-static uint32_t getComputeQueueFamilyIndex()
-{
-    uint32_t queueFamilyCount;
-
-    vkGetPhysicalDeviceQueueFamilyProperties(kPhysicalDevice, &queueFamilyCount, NULL);
-
-    std::vector<VkQueueFamilyProperties> queueFamilies(queueFamilyCount);
-    vkGetPhysicalDeviceQueueFamilyProperties(kPhysicalDevice,
-                                             &queueFamilyCount,
-                                             queueFamilies.data());
-
-    uint32_t i = 0;
-    for (; i < queueFamilies.size(); ++i)
-    {
-        VkQueueFamilyProperties props = queueFamilies[i];
-
-        if (props.queueCount > 0 && (props.queueFlags & VK_QUEUE_COMPUTE_BIT))
-        {
-            break;
-        }
-    }
-
-    if (i == queueFamilies.size())
-    {
-        throw std::runtime_error("could not find a queue family that supports operations");
-    }
-
-    return i;
-}
-
-bool checkExtensionAvailability(const char *extension_name,
-                                const std::vector<VkExtensionProperties> &available_extensions)
-{
-    for( size_t i = 0; i < available_extensions.size(); ++i )
-    {
-      if( strcmp( available_extensions[i].extensionName, extension_name ) == 0 )
-      {
-        return true;
-      }
-    }
-    return false;
-}
-
-VKAPI_ATTR VkBool32 VKAPI_CALL debugReportCallbackFn(
-        VkDebugReportFlagsEXT                       flags,
-        VkDebugReportObjectTypeEXT                  objectType,
-        uint64_t                                    object,
-        size_t                                      location,
-        int32_t                                     messageCode,
-        const char*                                 pLayerPrefix,
-        const char*                                 pMessage,
-        void*                                       pUserData)
-{
-        std::cout << "Debug Report: " << pLayerPrefix << ":" << pMessage << std::endl;
-        return VK_FALSE;
-}
-
-// internally used functions
-VkPhysicalDevice getPhysicalDevice()
-{
-    return kPhysicalDevice;
-}
-
-bool isAvailable()
-{
-    return getContext() != NULL;
-}
-
-#endif // HAVE_VULKAN
-
-}}} // namespace cv::dnn::vkcom
@@ -8,6 +8,10 @@
 #ifndef OPENCV_DNN_VKCOM_VULKAN_VK_LOADER_HPP
 #define OPENCV_DNN_VKCOM_VULKAN_VK_LOADER_HPP

+#ifdef HAVE_VULKAN
+#include <vulkan/vulkan.h>
+#endif // HAVE_VULKAN
+
 namespace cv { namespace dnn { namespace vkcom {

 #ifdef HAVE_VULKAN
@@ -128,10 +128,16 @@ TEST_P(DNNTestNetwork, GoogLeNet)

 TEST_P(DNNTestNetwork, Inception_5h)
 {
-    if (backend == DNN_BACKEND_INFERENCE_ENGINE) throw SkipTestException("");
+    double l1 = default_l1, lInf = default_lInf;
+    if (backend == DNN_BACKEND_INFERENCE_ENGINE && (target == DNN_TARGET_CPU || target == DNN_TARGET_OPENCL))
+    {
+        l1 = 1.72e-5;
+        lInf = 8e-4;
+    }
    processNet("dnn/tensorflow_inception_graph.pb", "", Size(224, 224), "softmax2",
               target == DNN_TARGET_OPENCL ? "dnn/halide_scheduler_opencl_inception_5h.yml" :
-                                             "dnn/halide_scheduler_inception_5h.yml");
+                                             "dnn/halide_scheduler_inception_5h.yml",
+               l1, lInf);
 }

 TEST_P(DNNTestNetwork, ENet)
@@ -193,41 +199,42 @@ TEST_P(DNNTestNetwork, SSD_VGG16)

 TEST_P(DNNTestNetwork, OpenPose_pose_coco)
 {
-    if (backend == DNN_BACKEND_HALIDE ||
-        (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_MYRIAD))
+    if (backend == DNN_BACKEND_HALIDE)
        throw SkipTestException("");
    processNet("dnn/openpose_pose_coco.caffemodel", "dnn/openpose_pose_coco.prototxt",
-               Size(368, 368));
+               Size(46, 46));
 }

 TEST_P(DNNTestNetwork, OpenPose_pose_mpi)
 {
-    if (backend == DNN_BACKEND_HALIDE ||
-        (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_MYRIAD))
+    if (backend == DNN_BACKEND_HALIDE)
        throw SkipTestException("");
    processNet("dnn/openpose_pose_mpi.caffemodel", "dnn/openpose_pose_mpi.prototxt",
-               Size(368, 368));
+               Size(46, 46));
 }

 TEST_P(DNNTestNetwork, OpenPose_pose_mpi_faster_4_stages)
 {
-    if (backend == DNN_BACKEND_HALIDE ||
-        (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_MYRIAD))
+    if (backend == DNN_BACKEND_HALIDE)
        throw SkipTestException("");
    // The same .caffemodel but modified .prototxt
    // See https://github.com/CMU-Perceptual-Computing-Lab/openpose/blob/master/src/openpose/pose/poseParameters.cpp
    processNet("dnn/openpose_pose_mpi.caffemodel", "dnn/openpose_pose_mpi_faster_4_stages.prototxt",
-               Size(368, 368));
+               Size(46, 46));
 }

 TEST_P(DNNTestNetwork, OpenFace)
 {
-#if defined(INF_ENGINE_RELEASE) && INF_ENGINE_RELEASE < 2018030000
+#if defined(INF_ENGINE_RELEASE)
+#if (INF_ENGINE_RELEASE < 2018030000 || INF_ENGINE_RELEASE == 2018050000)
    if (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_MYRIAD)
-        throw SkipTestException("Test is enabled starts from OpenVINO 2018R3");
+        throw SkipTestException("");
+#elif INF_ENGINE_RELEASE < 2018040000
+    if (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_OPENCL_FP16)
+        throw SkipTestException("Test is enabled starts from OpenVINO 2018R4");
 #endif
-    if (backend == DNN_BACKEND_HALIDE ||
-        (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_OPENCL_FP16))
+#endif
+    if (backend == DNN_BACKEND_HALIDE)
        throw SkipTestException("");
    processNet("dnn/openface_nn4.small2.v1.t7", "", Size(96, 96), "");
 }
@@ -300,10 +300,11 @@ INSTANTIATE_TEST_CASE_P(/**/, Reproducibility_ResNet50,
 typedef testing::TestWithParam<Target> Reproducibility_SqueezeNet_v1_1;
 TEST_P(Reproducibility_SqueezeNet_v1_1, Accuracy)
 {
+    int targetId = GetParam();
+    if(targetId == DNN_TARGET_OPENCL_FP16)
+        throw SkipTestException("This test does not support FP16");
    Net net = readNetFromCaffe(findDataFile("dnn/squeezenet_v1.1.prototxt", false),
                               findDataFile("dnn/squeezenet_v1.1.caffemodel", false));
-
-    int targetId = GetParam();
    net.setPreferableBackend(DNN_BACKEND_OPENCV);
    net.setPreferableTarget(targetId);

@@ -324,7 +325,8 @@ TEST_P(Reproducibility_SqueezeNet_v1_1, Accuracy)
    Mat ref = blobFromNPY(_tf("squeezenet_v1.1_prob.npy"));
    normAssert(ref, out);
 }
-INSTANTIATE_TEST_CASE_P(/**/, Reproducibility_SqueezeNet_v1_1, availableDnnTargets());
+INSTANTIATE_TEST_CASE_P(/**/, Reproducibility_SqueezeNet_v1_1,
+    testing::ValuesIn(getAvailableTargets(DNN_BACKEND_OPENCV)));

 TEST(Reproducibility_AlexNet_fp16, Accuracy)
 {
@@ -469,6 +471,7 @@ TEST(Test_Caffe, shared_weights)

  net.setInput(blob_1, "input_1");
  net.setInput(blob_2, "input_2");
+  net.setPreferableBackend(DNN_BACKEND_OPENCV);

  Mat sum = net.forward();

@@ -68,6 +68,7 @@ static inline void PrintTo(const cv::dnn::Target& v, std::ostream* os)
    case DNN_TARGET_OPENCL_FP16: *os << "OCL_FP16"; return;
    case DNN_TARGET_MYRIAD: *os << "MYRIAD"; return;
    case DNN_TARGET_VULKAN: *os << "VULKAN"; return;
+    case DNN_TARGET_FPGA: *os << "FPGA"; return;
    } // don't use "default:" to emit compiler warnings
    *os << "DNN_TARGET_UNKNOWN(" << (int)v << ")";
 }
@@ -190,30 +191,6 @@ static inline void normAssertDetections(cv::Mat ref, cv::Mat out, const char *co
                         testBoxes, comment, confThreshold, scores_diff, boxes_iou_diff);
 }

-static inline bool checkMyriadTarget()
-{
-#ifndef HAVE_INF_ENGINE
-    return false;
-#else
-    cv::dnn::Net net;
-    cv::dnn::LayerParams lp;
-    net.addLayerToPrev("testLayer", "Identity", lp);
-    net.setPreferableBackend(cv::dnn::DNN_BACKEND_INFERENCE_ENGINE);
-    net.setPreferableTarget(cv::dnn::DNN_TARGET_MYRIAD);
-    static int inpDims[] = {1, 2, 3, 4};
-    net.setInput(cv::Mat(4, &inpDims[0], CV_32FC1, cv::Scalar(0)));
-    try
-    {
-        net.forward();
-    }
-    catch(...)
-    {
-        return false;
-    }
-    return true;
-#endif
-}
-
 static inline bool readFileInMemory(const std::string& filename, std::string& content)
 {
    std::ios::openmode mode = std::ios::in | std::ios::binary;
@@ -238,52 +215,36 @@ namespace opencv_test {
 using namespace cv::dnn;

 static inline
-testing::internal::ParamGenerator<tuple<Backend, Target> > dnnBackendsAndTargets(
+testing::internal::ParamGenerator< tuple<Backend, Target> > dnnBackendsAndTargets(
        bool withInferenceEngine = true,
        bool withHalide = false,
        bool withCpuOCV = true,
        bool withVkCom = true
 )
 {
-    std::vector<tuple<Backend, Target> > targets;
-#ifdef HAVE_HALIDE
+    std::vector< tuple<Backend, Target> > targets;
+    std::vector< Target > available;
    if (withHalide)
    {
-        targets.push_back(make_tuple(DNN_BACKEND_HALIDE, DNN_TARGET_CPU));
-#ifdef HAVE_OPENCL
-        if (cv::ocl::useOpenCL())
-            targets.push_back(make_tuple(DNN_BACKEND_HALIDE, DNN_TARGET_OPENCL));
-#endif
+        available = getAvailableTargets(DNN_BACKEND_HALIDE);
+        for (std::vector< Target >::const_iterator i = available.begin(); i != available.end(); ++i)
+            targets.push_back(make_tuple(DNN_BACKEND_HALIDE, *i));
    }
-#endif
-#ifdef HAVE_INF_ENGINE
    if (withInferenceEngine)
    {
-        targets.push_back(make_tuple(DNN_BACKEND_INFERENCE_ENGINE, DNN_TARGET_CPU));
-#ifdef HAVE_OPENCL
-        if (cv::ocl::useOpenCL() && ocl::Device::getDefault().isIntel())
-        {
-            targets.push_back(make_tuple(DNN_BACKEND_INFERENCE_ENGINE, DNN_TARGET_OPENCL));
-            targets.push_back(make_tuple(DNN_BACKEND_INFERENCE_ENGINE, DNN_TARGET_OPENCL_FP16));
-        }
-#endif
-        if (checkMyriadTarget())
-            targets.push_back(make_tuple(DNN_BACKEND_INFERENCE_ENGINE, DNN_TARGET_MYRIAD));
+        available = getAvailableTargets(DNN_BACKEND_INFERENCE_ENGINE);
+        for (std::vector< Target >::const_iterator i = available.begin(); i != available.end(); ++i)
+            targets.push_back(make_tuple(DNN_BACKEND_INFERENCE_ENGINE, *i));
    }
-#endif
-    if (withCpuOCV)
-        targets.push_back(make_tuple(DNN_BACKEND_OPENCV, DNN_TARGET_CPU));
-#ifdef HAVE_OPENCL
-    if (cv::ocl::useOpenCL())
    {
-        targets.push_back(make_tuple(DNN_BACKEND_OPENCV, DNN_TARGET_OPENCL));
-        targets.push_back(make_tuple(DNN_BACKEND_OPENCV, DNN_TARGET_OPENCL_FP16));
+        available = getAvailableTargets(DNN_BACKEND_OPENCV);
+        for (std::vector< Target >::const_iterator i = available.begin(); i != available.end(); ++i)
+        {
+            if (!withCpuOCV && *i == DNN_TARGET_CPU)
+                continue;
+            targets.push_back(make_tuple(DNN_BACKEND_OPENCV, *i));
+        }
    }
-#endif
-#ifdef HAVE_VULKAN
-    if (withVkCom)
-        targets.push_back(make_tuple(DNN_BACKEND_VKCOM, DNN_TARGET_VULKAN));
-#endif
    if (targets.empty())  // validate at least CPU mode
        targets.push_back(make_tuple(DNN_BACKEND_OPENCV, DNN_TARGET_CPU));
    return testing::ValuesIn(targets);
@@ -295,21 +256,6 @@ testing::internal::ParamGenerator<tuple<Backend, Target> > dnnBackendsAndTargets
 namespace opencv_test {
 using namespace cv::dnn;

-static inline
-testing::internal::ParamGenerator<Target> availableDnnTargets()
-{
-    static std::vector<Target> targets;
-    if (targets.empty())
-    {
-        targets.push_back(DNN_TARGET_CPU);
-#ifdef HAVE_OPENCL
-        if (cv::ocl::useOpenCL())
-            targets.push_back(DNN_TARGET_OPENCL);
-#endif
-    }
-    return testing::ValuesIn(targets);
-}
-
 class DNNTestLayer : public TestWithParam<tuple<Backend, Target> >
 {
 public:
@@ -338,23 +284,10 @@ public:
       }
   }

-   static void checkBackend(int backend, int target, Mat* inp = 0, Mat* ref = 0)
-   {
-       if (backend == DNN_BACKEND_OPENCV && (target == DNN_TARGET_OPENCL || target == DNN_TARGET_OPENCL_FP16))
-       {
-#ifdef HAVE_OPENCL
-           if (!cv::ocl::useOpenCL())
-#endif
-           {
-               throw SkipTestException("OpenCL is not available/disabled in OpenCV");
-           }
-       }
+    static void checkBackend(int backend, int target, Mat* inp = 0, Mat* ref = 0)
+    {
       if (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_MYRIAD)
       {
-           if (!checkMyriadTarget())
-           {
-               throw SkipTestException("Myriad is not available/disabled in OpenCV");
-           }
 #if defined(INF_ENGINE_RELEASE) && INF_ENGINE_RELEASE < 2018030000
           if (inp && ref && inp->size[0] != 1)
           {
@@ -306,7 +306,7 @@ TEST_P(Test_Darknet_nets, TinyYoloVoc)
    // batch size 1
    testDarknetModel(config_file, weights_file, ref.rowRange(0, 2), scoreDiff, iouDiff);

-#if defined(INF_ENGINE_RELEASE) && INF_ENGINE_RELEASE == 2018040000
+#if defined(INF_ENGINE_RELEASE) && INF_ENGINE_RELEASE >= 2018040000
    if (backend == DNN_BACKEND_INFERENCE_ENGINE && target != DNN_TARGET_MYRIAD)
 #endif
    // batch size 2
@@ -347,8 +347,10 @@ INSTANTIATE_TEST_CASE_P(/**/, Test_Darknet_nets, dnnBackendsAndTargets());

 TEST_P(Test_Darknet_layers, shortcut)
 {
+#if defined(INF_ENGINE_RELEASE) && INF_ENGINE_RELEASE < 2018040000
    if (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_CPU)
-        throw SkipTestException("");
+        throw SkipTestException("Test is enabled starts from OpenVINO 2018R4");
+#endif
    testDarknetLayer("shortcut");
 }

@@ -55,9 +55,11 @@ static std::string _tf(TString filename)
 typedef testing::TestWithParam<Target> Reproducibility_GoogLeNet;
 TEST_P(Reproducibility_GoogLeNet, Batching)
 {
+    const int targetId = GetParam();
+    if(targetId == DNN_TARGET_OPENCL_FP16)
+        throw SkipTestException("This test does not support FP16");
    Net net = readNetFromCaffe(findDataFile("dnn/bvlc_googlenet.prototxt", false),
                               findDataFile("dnn/bvlc_googlenet.caffemodel", false));
-    int targetId = GetParam();
    net.setPreferableBackend(DNN_BACKEND_OPENCV);
    net.setPreferableTarget(targetId);

@@ -84,9 +86,11 @@ TEST_P(Reproducibility_GoogLeNet, Batching)

 TEST_P(Reproducibility_GoogLeNet, IntermediateBlobs)
 {
+    const int targetId = GetParam();
+    if(targetId == DNN_TARGET_OPENCL_FP16)
+        throw SkipTestException("This test does not support FP16");
    Net net = readNetFromCaffe(findDataFile("dnn/bvlc_googlenet.prototxt", false),
                               findDataFile("dnn/bvlc_googlenet.caffemodel", false));
-    int targetId = GetParam();
    net.setPreferableBackend(DNN_BACKEND_OPENCV);
    net.setPreferableTarget(targetId);

@@ -113,9 +117,11 @@ TEST_P(Reproducibility_GoogLeNet, IntermediateBlobs)

 TEST_P(Reproducibility_GoogLeNet, SeveralCalls)
 {
+    const int targetId = GetParam();
+    if(targetId == DNN_TARGET_OPENCL_FP16)
+        throw SkipTestException("This test does not support FP16");
    Net net = readNetFromCaffe(findDataFile("dnn/bvlc_googlenet.prototxt", false),
                               findDataFile("dnn/bvlc_googlenet.caffemodel", false));
-    int targetId = GetParam();
    net.setPreferableBackend(DNN_BACKEND_OPENCV);
    net.setPreferableTarget(targetId);

@@ -143,6 +149,7 @@ TEST_P(Reproducibility_GoogLeNet, SeveralCalls)
    normAssert(outs[0], ref, "", 1E-4, 1E-2);
 }

-INSTANTIATE_TEST_CASE_P(/**/, Reproducibility_GoogLeNet, availableDnnTargets());
+INSTANTIATE_TEST_CASE_P(/**/, Reproducibility_GoogLeNet,
+    testing::ValuesIn(getAvailableTargets(DNN_BACKEND_OPENCV)));

 }} // namespace
@@ -166,7 +166,7 @@ TEST_P(Deconvolution, Accuracy)
    if (backendId == DNN_BACKEND_INFERENCE_ENGINE && targetId == DNN_TARGET_CPU &&
        dilation.width == 2 && dilation.height == 2)
        throw SkipTestException("");
-#if defined(INF_ENGINE_RELEASE) && INF_ENGINE_RELEASE == 2018040000
+#if defined(INF_ENGINE_RELEASE) && INF_ENGINE_RELEASE >= 2018040000
    if (backendId == DNN_BACKEND_INFERENCE_ENGINE && targetId == DNN_TARGET_CPU &&
        hasBias && group != 1)
        throw SkipTestException("Test is disabled for OpenVINO 2018R4");
@@ -273,9 +273,11 @@ TEST_P(AvePooling, Accuracy)
    Size stride = get<3>(GetParam());
    Backend backendId = get<0>(get<4>(GetParam()));
    Target targetId = get<1>(get<4>(GetParam()));
+#if defined(INF_ENGINE_RELEASE) && INF_ENGINE_RELEASE < 2018040000
    if (backendId == DNN_BACKEND_INFERENCE_ENGINE && targetId == DNN_TARGET_MYRIAD &&
        stride == Size(3, 2) && kernel == Size(3, 3) && outSize != Size(1, 1))
-        throw SkipTestException("");
+        throw SkipTestException("Test is enabled starts from OpenVINO 2018R4");
+#endif

    const int inWidth = (outSize.width - 1) * stride.width + kernel.width;
    const int inHeight = (outSize.height - 1) * stride.height + kernel.height;
@@ -57,28 +57,29 @@ void runIE(Target target, const std::string& xmlPath, const std::string& binPath
    InferencePlugin plugin;
    ExecutableNetwork netExec;
    InferRequest infRequest;
-    TargetDevice targetDevice;
-    switch (target)
-    {
-        case DNN_TARGET_CPU:
-            targetDevice = TargetDevice::eCPU;
-            break;
-        case DNN_TARGET_OPENCL:
-        case DNN_TARGET_OPENCL_FP16:
-            targetDevice = TargetDevice::eGPU;
-            break;
-        case DNN_TARGET_MYRIAD:
-            targetDevice = TargetDevice::eMYRIAD;
-            break;
-        default:
-            CV_Error(Error::StsNotImplemented, "Unknown target");
-    };
-
    try
    {
-        enginePtr = PluginDispatcher({""}).getSuitablePlugin(targetDevice);
+        auto dispatcher = InferenceEngine::PluginDispatcher({""});
+        switch (target)
+        {
+            case DNN_TARGET_CPU:
+                enginePtr = dispatcher.getSuitablePlugin(TargetDevice::eCPU);
+                break;
+            case DNN_TARGET_OPENCL:
+            case DNN_TARGET_OPENCL_FP16:
+                enginePtr = dispatcher.getSuitablePlugin(TargetDevice::eGPU);
+                break;
+            case DNN_TARGET_MYRIAD:
+                enginePtr = dispatcher.getSuitablePlugin(TargetDevice::eMYRIAD);
+                break;
+            case DNN_TARGET_FPGA:
+                enginePtr = dispatcher.getPluginByDevice("HETERO:FPGA,CPU");
+                break;
+            default:
+                CV_Error(Error::StsNotImplemented, "Unknown target");
+        };

-        if (targetDevice == TargetDevice::eCPU)
+        if (target == DNN_TARGET_CPU || target == DNN_TARGET_FPGA)
        {
            std::string suffixes[] = {"_avx2", "_sse4", ""};
            bool haveFeature[] = {
@@ -189,6 +190,14 @@ TEST_P(DNNTestOpenVINO, models)
                                         modelName == "landmarks-regression-retail-0009" ||
                                          modelName == "semantic-segmentation-adas-0001")))
        throw SkipTestException("");
+#elif INF_ENGINE_RELEASE == 2018050000
+    if (modelName == "single-image-super-resolution-0063" ||
+        modelName == "single-image-super-resolution-1011" ||
+        modelName == "single-image-super-resolution-1021" ||
+        (target == DNN_TARGET_OPENCL_FP16 && modelName == "face-reidentification-retail-0095") ||
+        (target == DNN_TARGET_MYRIAD && (modelName == "license-plate-recognition-barrier-0001" ||
+                                         modelName == "semantic-segmentation-adas-0001")))
+        throw SkipTestException("");
 #endif
 #endif

@@ -202,7 +211,8 @@ TEST_P(DNNTestOpenVINO, models)
    std::map<std::string, cv::Mat> inputsMap;
    std::map<std::string, cv::Mat> ieOutputsMap, cvOutputsMap;
    // Single Myriad device cannot be shared across multiple processes.
-    resetMyriadDevice();
+    if (target == DNN_TARGET_MYRIAD)
+        resetMyriadDevice();
    runIE(target, xmlPath, binPath, inputsMap, ieOutputsMap);
    runCV(target, xmlPath, binPath, inputsMap, cvOutputsMap);

@@ -244,25 +254,10 @@ static testing::internal::ParamGenerator<String> intelModels()
    return ValuesIn(modelsNames);
 }

-static testing::internal::ParamGenerator<Target> dnnDLIETargets()
-{
-    std::vector<Target> targets;
-    targets.push_back(DNN_TARGET_CPU);
-#ifdef HAVE_OPENCL
-    if (cv::ocl::useOpenCL() && ocl::Device::getDefault().isIntel())
-    {
-        targets.push_back(DNN_TARGET_OPENCL);
-        targets.push_back(DNN_TARGET_OPENCL_FP16);
-    }
-#endif
-    if (checkMyriadTarget())
-        targets.push_back(DNN_TARGET_MYRIAD);
-    return testing::ValuesIn(targets);
-}
-
-INSTANTIATE_TEST_CASE_P(/**/, DNNTestOpenVINO, Combine(
-    dnnDLIETargets(), intelModels()
-));
+INSTANTIATE_TEST_CASE_P(/**/,
+    DNNTestOpenVINO,
+    Combine(testing::ValuesIn(getAvailableTargets(DNN_BACKEND_INFERENCE_ENGINE)), intelModels())
+);

 }}
 #endif  // HAVE_INF_ENGINE
@@ -137,7 +137,7 @@ TEST_P(Test_Caffe_layers, Convolution)

 TEST_P(Test_Caffe_layers, DeConvolution)
 {
-#if defined(INF_ENGINE_RELEASE) && INF_ENGINE_RELEASE == 2018040000
+#if defined(INF_ENGINE_RELEASE) && INF_ENGINE_RELEASE >= 2018040000
    if (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_CPU)
        throw SkipTestException("Test is disabled for OpenVINO 2018R4");
 #endif
@@ -243,9 +243,14 @@ TEST_P(Test_Caffe_layers, Concat)

 TEST_P(Test_Caffe_layers, Fused_Concat)
 {
-    if ((backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_CPU) ||
-        (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_OPENCL))
+#if defined(INF_ENGINE_RELEASE)
+    if (backend == DNN_BACKEND_INFERENCE_ENGINE)
+    {
+        if (target == DNN_TARGET_OPENCL || target == DNN_TARGET_OPENCL_FP16 ||
+            (INF_ENGINE_RELEASE < 2018040000 && target == DNN_TARGET_CPU))
        throw SkipTestException("");
+    }
+#endif
    checkBackend();

    // Test case
@@ -290,6 +295,10 @@ TEST_P(Test_Caffe_layers, Eltwise)
 {
    if (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_MYRIAD)
        throw SkipTestException("");
+#if defined(INF_ENGINE_RELEASE) && INF_ENGINE_RELEASE == 2018050000
+    if (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_OPENCL)
+        throw SkipTestException("Test is disabled for OpenVINO 2018R5");
+#endif
    testLayerUsingCaffeModels("layer_eltwise");
 }

@@ -349,12 +358,6 @@ TEST_P(Test_Caffe_layers, Reshape_Split_Slice)

 TEST_P(Test_Caffe_layers, Conv_Elu)
 {
-    if (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_MYRIAD)
-    {
-        if (!checkMyriadTarget())
-            throw SkipTestException("Myriad is not available/disabled in OpenCV");
-    }
-
    Net net = readNetFromTensorflow(_tf("layer_elu_model.pb"));
    ASSERT_FALSE(net.empty());

@@ -919,8 +922,11 @@ INSTANTIATE_TEST_CASE_P(/**/, Layer_Test_DWconv_Prelu, Combine(Values(3, 6), Val
 // Using Intel's Model Optimizer generate .xml and .bin files:
 // ./ModelOptimizer -w /path/to/caffemodel -d /path/to/prototxt \
 //                  -p FP32 -i -b ${batch_size} -o /path/to/output/folder
-TEST(Layer_Test_Convolution_DLDT, Accuracy)
+typedef testing::TestWithParam<Target> Layer_Test_Convolution_DLDT;
+TEST_P(Layer_Test_Convolution_DLDT, Accuracy)
 {
+    Target targetId = GetParam();
+
    Net netDefault = readNet(_tf("layer_convolution.caffemodel"), _tf("layer_convolution.prototxt"));
    Net net = readNet(_tf("layer_convolution.xml"), _tf("layer_convolution.bin"));

@@ -931,17 +937,29 @@ TEST(Layer_Test_Convolution_DLDT, Accuracy)
    Mat outDefault = netDefault.forward();

    net.setInput(inp);
-    Mat out = net.forward();
+    net.setPreferableTarget(targetId);

-    normAssert(outDefault, out);
+    if (targetId != DNN_TARGET_MYRIAD)
+    {
+        Mat out = net.forward();

-    std::vector<int> outLayers = net.getUnconnectedOutLayers();
-    ASSERT_EQ(net.getLayer(outLayers[0])->name, "output_merge");
-    ASSERT_EQ(net.getLayer(outLayers[0])->type, "Concat");
+        normAssert(outDefault, out);
+
+        std::vector<int> outLayers = net.getUnconnectedOutLayers();
+        ASSERT_EQ(net.getLayer(outLayers[0])->name, "output_merge");
+        ASSERT_EQ(net.getLayer(outLayers[0])->type, "Concat");
+    }
+    else
+    {
+        // An assertion is expected because the model is in FP32 format but
+        // Myriad plugin supports only FP16 models.
+        ASSERT_ANY_THROW(net.forward());
+    }
 }

-TEST(Layer_Test_Convolution_DLDT, setInput_uint8)
+TEST_P(Layer_Test_Convolution_DLDT, setInput_uint8)
 {
+    Target targetId = GetParam();
    Mat inp = blobFromNPY(_tf("blob.npy"));

    Mat inputs[] = {Mat(inp.dims, inp.size, CV_8U), Mat()};
@@ -952,12 +970,25 @@ TEST(Layer_Test_Convolution_DLDT, setInput_uint8)
    for (int i = 0; i < 2; ++i)
    {
        Net net = readNet(_tf("layer_convolution.xml"), _tf("layer_convolution.bin"));
+        net.setPreferableTarget(targetId);
        net.setInput(inputs[i]);
-        outs[i] = net.forward();
-        ASSERT_EQ(outs[i].type(), CV_32F);
+        if (targetId != DNN_TARGET_MYRIAD)
+        {
+            outs[i] = net.forward();
+            ASSERT_EQ(outs[i].type(), CV_32F);
+        }
+        else
+        {
+            // An assertion is expected because the model is in FP32 format but
+            // Myriad plugin supports only FP16 models.
+            ASSERT_ANY_THROW(net.forward());
+        }
    }
-    normAssert(outs[0], outs[1]);
+    if (targetId != DNN_TARGET_MYRIAD)
+        normAssert(outs[0], outs[1]);
 }
+INSTANTIATE_TEST_CASE_P(/**/, Layer_Test_Convolution_DLDT,
+    testing::ValuesIn(getAvailableTargets(DNN_BACKEND_INFERENCE_ENGINE)));

 // 1. Create a .prototxt file with the following network:
 // layer {
@@ -981,14 +1012,17 @@ TEST(Layer_Test_Convolution_DLDT, setInput_uint8)
 // net.save('/path/to/caffemodel')
 //
 // 3. Convert using ModelOptimizer.
-typedef testing::TestWithParam<tuple<int, int> > Test_DLDT_two_inputs;
+typedef testing::TestWithParam<tuple<int, int, Target> > Test_DLDT_two_inputs;
 TEST_P(Test_DLDT_two_inputs, as_IR)
 {
    int firstInpType = get<0>(GetParam());
    int secondInpType = get<1>(GetParam());
-    // TODO: It looks like a bug in Inference Engine.
+    Target targetId = get<2>(GetParam());
+
+#if defined(INF_ENGINE_RELEASE) && INF_ENGINE_RELEASE < 2018040000
    if (secondInpType == CV_8U)
-        throw SkipTestException("");
+        throw SkipTestException("Test is enabled starts from OpenVINO 2018R4");
+#endif

    Net net = readNet(_tf("net_two_inputs.xml"), _tf("net_two_inputs.bin"));
    int inpSize[] = {1, 2, 3};
@@ -999,11 +1033,21 @@ TEST_P(Test_DLDT_two_inputs, as_IR)

    net.setInput(firstInp, "data");
    net.setInput(secondInp, "second_input");
-    Mat out = net.forward();
+    net.setPreferableTarget(targetId);
+    if (targetId != DNN_TARGET_MYRIAD)
+    {
+        Mat out = net.forward();

-    Mat ref;
-    cv::add(firstInp, secondInp, ref, Mat(), CV_32F);
-    normAssert(out, ref);
+        Mat ref;
+        cv::add(firstInp, secondInp, ref, Mat(), CV_32F);
+        normAssert(out, ref);
+    }
+    else
+    {
+        // An assertion is expected because the model is in FP32 format but
+        // Myriad plugin supports only FP16 models.
+        ASSERT_ANY_THROW(net.forward());
+    }
 }

 TEST_P(Test_DLDT_two_inputs, as_backend)
@@ -1011,6 +1055,8 @@ TEST_P(Test_DLDT_two_inputs, as_backend)
    static const float kScale = 0.5f;
    static const float kScaleInv = 1.0f / kScale;

+    Target targetId = get<2>(GetParam());
+
    Net net;
    LayerParams lp;
    lp.type = "Eltwise";
@@ -1019,9 +1065,9 @@ TEST_P(Test_DLDT_two_inputs, as_backend)
    int eltwiseId = net.addLayerToPrev(lp.name, lp.type, lp);  // connect to a first input
    net.connect(0, 1, eltwiseId, 1);  // connect to a second input

-    int inpSize[] = {1, 2, 3};
-    Mat firstInp(3, &inpSize[0], get<0>(GetParam()));
-    Mat secondInp(3, &inpSize[0], get<1>(GetParam()));
+    int inpSize[] = {1, 2, 3, 4};
+    Mat firstInp(4, &inpSize[0], get<0>(GetParam()));
+    Mat secondInp(4, &inpSize[0], get<1>(GetParam()));
    randu(firstInp, 0, 255);
    randu(secondInp, 0, 255);

@@ -1029,15 +1075,20 @@ TEST_P(Test_DLDT_two_inputs, as_backend)
    net.setInput(firstInp, "data", kScale);
    net.setInput(secondInp, "second_input", kScaleInv);
    net.setPreferableBackend(DNN_BACKEND_INFERENCE_ENGINE);
+    net.setPreferableTarget(targetId);
    Mat out = net.forward();

    Mat ref;
    addWeighted(firstInp, kScale, secondInp, kScaleInv, 0, ref, CV_32F);
-    normAssert(out, ref);
+    // Output values are in range [0, 637.5].
+    double l1 = (targetId == DNN_TARGET_OPENCL_FP16 || targetId == DNN_TARGET_MYRIAD) ? 0.06 : 1e-6;
+    double lInf = (targetId == DNN_TARGET_OPENCL_FP16 || targetId == DNN_TARGET_MYRIAD) ? 0.3 : 1e-5;
+    normAssert(out, ref, "", l1, lInf);
 }

 INSTANTIATE_TEST_CASE_P(/*nothing*/, Test_DLDT_two_inputs, Combine(
-  Values(CV_8U, CV_32F), Values(CV_8U, CV_32F)
+  Values(CV_8U, CV_32F), Values(CV_8U, CV_32F),
+  testing::ValuesIn(getAvailableTargets(DNN_BACKEND_INFERENCE_ENGINE))
 ));

 class UnsupportedLayer : public Layer
@@ -157,8 +157,6 @@ TEST_P(setInput, normalization)
    const int target   = get<1>(get<3>(GetParam()));
    const bool kSwapRB = true;

-    if (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_MYRIAD && !checkMyriadTarget())
-        throw SkipTestException("Myriad is not available/disabled in OpenCV");
    if (backend == DNN_BACKEND_OPENCV && target == DNN_TARGET_OPENCL_FP16 && dtype != CV_32F)
        throw SkipTestException("");
    if (backend == DNN_BACKEND_VKCOM && dtype != CV_32F)
@@ -162,6 +162,12 @@ TEST_P(Test_ONNX_layers, MultyInputs)
    normAssert(ref, out, "", default_l1,  default_lInf);
 }

+TEST_P(Test_ONNX_layers, DynamicReshape)
+{
+    if (backend == DNN_BACKEND_INFERENCE_ENGINE && (target == DNN_TARGET_OPENCL || target == DNN_TARGET_OPENCL_FP16))
+        throw SkipTestException("");
+    testONNXModels("dynamic_reshape");
+}

 INSTANTIATE_TEST_CASE_P(/*nothing*/, Test_ONNX_layers, dnnBackendsAndTargets());

@@ -245,6 +251,10 @@ TEST_P(Test_ONNX_nets, VGG16)
    else if (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_OPENCL) {
        lInf = 1.2e-4;
    }
+#if defined(INF_ENGINE_RELEASE) && INF_ENGINE_RELEASE >= 2018050000
+    if (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_OPENCL_FP16)
+        l1 = 0.131;
+#endif
    testONNXModels("vgg16", pb, l1, lInf);
 }

@@ -323,7 +333,7 @@ TEST_P(Test_ONNX_nets, CNN_MNIST)
 TEST_P(Test_ONNX_nets, MobileNet_v2)
 {
    // output range: [-166; 317]
-    const double l1 = (target == DNN_TARGET_OPENCL_FP16 || target == DNN_TARGET_MYRIAD) ? 0.38 : 7e-5;
+    const double l1 = (target == DNN_TARGET_OPENCL_FP16 || target == DNN_TARGET_MYRIAD) ? 0.4 : 7e-5;
    const double lInf = (target == DNN_TARGET_OPENCL_FP16 || target == DNN_TARGET_MYRIAD) ? 2.87 : 5e-4;
    testONNXModels("mobilenetv2", pb, l1, lInf);
 }
@@ -346,7 +356,17 @@ TEST_P(Test_ONNX_nets, LResNet100E_IR)

 TEST_P(Test_ONNX_nets, Emotion_ferplus)
 {
-    testONNXModels("emotion_ferplus", pb);
+    double l1 = default_l1;
+    double lInf = default_lInf;
+    // Output values are in range [-2.01109, 2.11111]
+    if (backend == DNN_BACKEND_OPENCV && target == DNN_TARGET_OPENCL_FP16)
+        l1 = 0.007;
+    else if (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_OPENCL_FP16)
+    {
+        l1 = 0.021;
+        lInf = 0.034;
+    }
+    testONNXModels("emotion_ferplus", pb, l1, lInf);
 }

 TEST_P(Test_ONNX_nets, Inception_v2)
@@ -367,6 +387,10 @@ TEST_P(Test_ONNX_nets, DenseNet121)

 TEST_P(Test_ONNX_nets, Inception_v1)
 {
+#if defined(INF_ENGINE_RELEASE) && INF_ENGINE_RELEASE == 2018050000
+    if (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_MYRIAD)
+        throw SkipTestException("");
+#endif
    testONNXModels("inception_v1", pb);
 }

@@ -136,11 +136,13 @@ TEST_P(Test_TensorFlow_layers, padding)
    runTensorFlowNet("padding_same");
    runTensorFlowNet("padding_valid");
    runTensorFlowNet("spatial_padding");
+    runTensorFlowNet("keras_pad_concat");
 }

-TEST_P(Test_TensorFlow_layers, eltwise_add_mul)
+TEST_P(Test_TensorFlow_layers, eltwise)
 {
    runTensorFlowNet("eltwise_add_mul");
+    runTensorFlowNet("eltwise_sub");
 }

 TEST_P(Test_TensorFlow_layers, pad_and_concat)
@@ -239,6 +241,10 @@ TEST_P(Test_TensorFlow_layers, unfused_flatten)

 TEST_P(Test_TensorFlow_layers, leaky_relu)
 {
+#if defined(INF_ENGINE_RELEASE) && INF_ENGINE_RELEASE == 2018050000
+    if (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_OPENCL)
+        throw SkipTestException("");
+#endif
    runTensorFlowNet("leaky_relu_order1");
    runTensorFlowNet("leaky_relu_order2");
    runTensorFlowNet("leaky_relu_order3");
@@ -381,6 +387,10 @@ TEST_P(Test_TensorFlow_nets, Faster_RCNN)

 TEST_P(Test_TensorFlow_nets, MobileNet_v1_SSD_PPN)
 {
+#if defined(INF_ENGINE_RELEASE) && INF_ENGINE_RELEASE == 2018050000
+    if (backend == DNN_BACKEND_INFERENCE_ENGINE && (target == DNN_TARGET_OPENCL || target == DNN_TARGET_OPENCL_FP16))
+        throw SkipTestException("Unstable test case");
+#endif
    checkBackend();
    std::string proto = findDataFile("dnn/ssd_mobilenet_v1_ppn_coco.pbtxt", false);
    std::string model = findDataFile("dnn/ssd_mobilenet_v1_ppn_coco.pb", false);
@@ -558,6 +568,10 @@ TEST_P(Test_TensorFlow_layers, slice)
    if (backend == DNN_BACKEND_INFERENCE_ENGINE &&
        (target == DNN_TARGET_OPENCL || target == DNN_TARGET_OPENCL_FP16))
        throw SkipTestException("");
+#if defined(INF_ENGINE_RELEASE) && INF_ENGINE_RELEASE == 2018050000
+    if (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_MYRIAD)
+        throw SkipTestException("");
+#endif
    runTensorFlowNet("slice_4d");
 }

@@ -73,7 +73,7 @@ class Test_Torch_layers : public DNNTestLayer
 {
 public:
    void runTorchNet(const String& prefix, String outLayerName = "",
-                     bool check2ndBlob = false, bool isBinary = false,
+                     bool check2ndBlob = false, bool isBinary = false, bool evaluate = true,
                     double l1 = 0.0, double lInf = 0.0)
    {
        String suffix = (isBinary) ? ".dat" : ".txt";
@@ -84,7 +84,7 @@ public:

        checkBackend(backend, target, &inp, &outRef);

-        Net net = readNetFromTorch(_tf(prefix + "_net" + suffix), isBinary);
+        Net net = readNetFromTorch(_tf(prefix + "_net" + suffix), isBinary, evaluate);
        ASSERT_FALSE(net.empty());

        net.setPreferableBackend(backend);
@@ -114,7 +114,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;
-    runTorchNet("net_conv", "", false, true, l1, lInf);
+    runTorchNet("net_conv", "", false, true, true, l1, lInf);
 }

 TEST_P(Test_Torch_layers, run_pool_max)
@@ -136,7 +136,7 @@ TEST_P(Test_Torch_layers, run_reshape_change_batch_size)

 TEST_P(Test_Torch_layers, run_reshape)
 {
-#if defined(INF_ENGINE_RELEASE) && INF_ENGINE_RELEASE == 2018040000
+#if defined(INF_ENGINE_RELEASE) && INF_ENGINE_RELEASE >= 2018040000
    if (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_MYRIAD)
        throw SkipTestException("Test is disabled for OpenVINO 2018R4");
 #endif
@@ -147,7 +147,7 @@ TEST_P(Test_Torch_layers, run_reshape)
 TEST_P(Test_Torch_layers, run_reshape_single_sample)
 {
    // Reference output values in range [14.4586, 18.4492].
-    runTorchNet("net_reshape_single_sample", "", false, false,
+    runTorchNet("net_reshape_single_sample", "", false, false, true,
                (target == DNN_TARGET_MYRIAD || target == DNN_TARGET_OPENCL_FP16) ? 0.0073 : default_l1,
                (target == DNN_TARGET_MYRIAD || target == DNN_TARGET_OPENCL_FP16) ? 0.025 : default_lInf);
 }
@@ -166,13 +166,13 @@ TEST_P(Test_Torch_layers, run_concat)

 TEST_P(Test_Torch_layers, run_depth_concat)
 {
-    runTorchNet("net_depth_concat", "", false, true, 0.0,
+    runTorchNet("net_depth_concat", "", false, true, true, 0.0,
                target == DNN_TARGET_OPENCL_FP16 ? 0.021 : 0.0);
 }

 TEST_P(Test_Torch_layers, run_deconv)
 {
-#if defined(INF_ENGINE_RELEASE) && INF_ENGINE_RELEASE == 2018040000
+#if defined(INF_ENGINE_RELEASE) && INF_ENGINE_RELEASE >= 2018040000
    if (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_MYRIAD)
        throw SkipTestException("Test is disabled for OpenVINO 2018R4");
 #endif
@@ -182,6 +182,7 @@ TEST_P(Test_Torch_layers, run_deconv)
 TEST_P(Test_Torch_layers, run_batch_norm)
 {
    runTorchNet("net_batch_norm", "", false, true);
+    runTorchNet("net_batch_norm_train", "", false, true, false);
 }

 TEST_P(Test_Torch_layers, net_prelu)
@@ -216,7 +217,7 @@ TEST_P(Test_Torch_layers, net_conv_gemm_lrn)
 {
    if (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_MYRIAD)
        throw SkipTestException("");
-    runTorchNet("net_conv_gemm_lrn", "", false, true,
+    runTorchNet("net_conv_gemm_lrn", "", false, true, true,
                target == DNN_TARGET_OPENCL_FP16 ? 0.046 : 0.0,
                target == DNN_TARGET_OPENCL_FP16 ? 0.023 : 0.0);
 }
@@ -266,9 +267,9 @@ class Test_Torch_nets : public DNNTestLayer {};

 TEST_P(Test_Torch_nets, OpenFace_accuracy)
 {
-#if defined(INF_ENGINE_RELEASE) && INF_ENGINE_RELEASE < 2018030000
+#if defined(INF_ENGINE_RELEASE) && (INF_ENGINE_RELEASE < 2018030000 || INF_ENGINE_RELEASE == 2018050000)
    if (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_MYRIAD)
-        throw SkipTestException("Test is enabled starts from OpenVINO 2018R3");
+        throw SkipTestException("");
 #endif
    checkBackend();
    if (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_OPENCL_FP16)
@@ -389,6 +390,10 @@ TEST_P(Test_Torch_nets, ENet_accuracy)
 //   -model models/instance_norm/feathers.t7
 TEST_P(Test_Torch_nets, FastNeuralStyle_accuracy)
 {
+#if defined(INF_ENGINE_RELEASE) && INF_ENGINE_RELEASE == 2018050000
+    if (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_MYRIAD)
+        throw SkipTestException("");
+#endif
    checkBackend();
    std::string models[] = {"dnn/fast_neural_style_eccv16_starry_night.t7",
                            "dnn/fast_neural_style_instance_norm_feathers.t7"};
@@ -285,6 +285,18 @@ public:
        const std::vector<int> &numberList, float dMax=5.85f, float dMin=8.2f,
        const std::vector<int>& indexChange=std::vector<int>());
    CV_WRAP virtual String getDefaultName() const CV_OVERRIDE;
+
+    /** @brief Set detection threshold.
+    @param threshold AGAST detection threshold score.
+    */
+    CV_WRAP virtual void setThreshold(int threshold) { CV_UNUSED(threshold); return; }
+    CV_WRAP virtual int getThreshold() const { return -1; }
+
+    /** @brief Set detection octaves.
+    @param octaves detection octaves. Use 0 to do single scale.
+    */
+    CV_WRAP virtual void setOctaves(int octaves) { CV_UNUSED(octaves); return; }
+    CV_WRAP virtual int getOctaves() const { return -1; }
 };

 /** @brief Class implementing the ORB (*oriented BRIEF*) keypoint detector and descriptor extractor
@@ -476,9 +488,9 @@ FastFeatureDetector::TYPE_5_8

 Detects corners using the FAST algorithm by @cite Rosten06 .

-@note In Python API, types are given as cv2.FAST_FEATURE_DETECTOR_TYPE_5_8,
-cv2.FAST_FEATURE_DETECTOR_TYPE_7_12 and cv2.FAST_FEATURE_DETECTOR_TYPE_9_16. For corner
-detection, use cv2.FAST.detect() method.
+@note In Python API, types are given as cv.FAST_FEATURE_DETECTOR_TYPE_5_8,
+cv.FAST_FEATURE_DETECTOR_TYPE_7_12 and cv.FAST_FEATURE_DETECTOR_TYPE_9_16. For corner
+detection, use cv.FAST.detect() method.
 */
 CV_EXPORTS void FAST( InputArray image, CV_OUT std::vector<KeyPoint>& keypoints,
                      int threshold, bool nonmaxSuppression, FastFeatureDetector::DetectorType type );
@@ -1212,9 +1224,9 @@ output image. See possible flags bit values below.
 DrawMatchesFlags. See details above in drawMatches .

@note
-For Python API, flags are modified as cv2.DRAW_MATCHES_FLAGS_DEFAULT,
-cv2.DRAW_MATCHES_FLAGS_DRAW_RICH_KEYPOINTS, cv2.DRAW_MATCHES_FLAGS_DRAW_OVER_OUTIMG,
-cv2.DRAW_MATCHES_FLAGS_NOT_DRAW_SINGLE_POINTS
+For Python API, flags are modified as cv.DRAW_MATCHES_FLAGS_DEFAULT,
+cv.DRAW_MATCHES_FLAGS_DRAW_RICH_KEYPOINTS, cv.DRAW_MATCHES_FLAGS_DRAW_OVER_OUTIMG,
+cv.DRAW_MATCHES_FLAGS_NOT_DRAW_SINGLE_POINTS
 */
 CV_EXPORTS_W void drawKeypoints( InputArray image, const std::vector<KeyPoint>& keypoints, InputOutputArray outImage,
                               const Scalar& color=Scalar::all(-1), DrawMatchesFlags flags=DrawMatchesFlags::DEFAULT );
@@ -7,11 +7,13 @@ import java.util.List;
 import org.opencv.calib3d.Calib3d;
 import org.opencv.core.CvType;
 import org.opencv.core.Mat;
+import org.opencv.core.MatOfInt;
 import org.opencv.core.MatOfDMatch;
 import org.opencv.core.MatOfKeyPoint;
 import org.opencv.core.MatOfPoint2f;
 import org.opencv.core.Point;
 import org.opencv.core.Range;
+import org.opencv.core.Scalar;
 import org.opencv.core.DMatch;
 import org.opencv.features2d.DescriptorMatcher;
 import org.opencv.features2d.Features2d;
@@ -141,4 +143,30 @@ public class Features2dTest extends OpenCVTestCase {
        Imgcodecs.imwrite(outputPath, outimg);
        // OpenCVTestRunner.Log("Output image is saved to: " + outputPath);
    }
+
+    public void testDrawKeypoints()
+    {
+        Mat outImg = Mat.ones(11, 11, CvType.CV_8U);
+
+        MatOfKeyPoint kps = new MatOfKeyPoint(new KeyPoint(5, 5, 1));  // x, y, size
+        Features2d.drawKeypoints(new Mat(), kps, outImg, new Scalar(255),
+                                 Features2d.DrawMatchesFlags_DRAW_OVER_OUTIMG);
+
+        Mat ref = new MatOfInt(new int[] {
+            1,   1,   1,   1,   1,   1,   1,   1,   1,   1,   1,
+            1,   1,   1,   1,  15,  54,  15,   1,   1,   1,   1,
+            1,   1,   1,  76, 217, 217, 221,  81,   1,   1,   1,
+            1,   1, 100, 224, 111,  57, 115, 225, 101,   1,   1,
+            1,  44, 215, 100,   1,   1,   1, 101, 214,  44,   1,
+            1,  54, 212,  57,   1,   1,   1,  55, 212,  55,   1,
+            1,  40, 215, 104,   1,   1,   1, 105, 215,  40,   1,
+            1,   1, 102, 221, 111,  55, 115, 222, 103,   1,   1,
+            1,   1,   1,  76, 218, 217, 220,  81,   1,   1,   1,
+            1,   1,   1,   1,  15,  55,  15,   1,   1,   1,   1,
+            1,   1,   1,   1,   1,   1,   1,   1,   1,   1,   1
+        }).reshape(1, 11);
+        ref.convertTo(ref, CvType.CV_8U);
+
+        assertMatEqual(ref, outImg);
+    }
 }
@@ -80,6 +80,26 @@ public:
        return NORM_HAMMING;
    }

+    virtual void setThreshold(int threshold_in) CV_OVERRIDE
+    {
+        threshold = threshold_in;
+    }
+
+    virtual int getThreshold() const CV_OVERRIDE
+    {
+        return threshold;
+    }
+
+    virtual void setOctaves(int octaves_in) CV_OVERRIDE
+    {
+        octaves = octaves_in;
+    }
+
+    virtual int getOctaves() const CV_OVERRIDE
+    {
+        return octaves;
+    }
+
    // call this to generate the kernel:
    // circle of radius r (pixels), with n points;
    // short pairings with dMax, long pairings with dMin
@@ -117,7 +117,7 @@ void drawKeypoints( InputArray image, const std::vector<KeyPoint>& keypoints, In
                                     end = keypoints.end();
    for( ; it != end; ++it )
    {
-        Scalar color = isRandColor ? Scalar(rng(256), rng(256), rng(256)) : _color;
+        Scalar color = isRandColor ? Scalar( rng(256), rng(256), rng(256), 255 ) : _color;
        _drawKeypoint( outImage, *it, color, flags );
    }
 }
@@ -173,7 +173,7 @@ static inline void _drawMatch( InputOutputArray outImg, InputOutputArray outImg1
 {
    RNG& rng = theRNG();
    bool isRandMatchColor = matchColor == Scalar::all(-1);
-    Scalar color = isRandMatchColor ? Scalar( rng(256), rng(256), rng(256) ) : matchColor;
+    Scalar color = isRandMatchColor ? Scalar( rng(256), rng(256), rng(256), 255 ) : matchColor;

    _drawKeypoint( outImg1, kp1, color, flags );
    _drawKeypoint( outImg2, kp2, color, flags );
@@ -77,7 +77,7 @@ void KeyPointsFilter::retainBest(std::vector<KeyPoint>& keypoints, int n_points)
            return;
        }
        //first use nth element to partition the keypoints into the best and worst.
-        std::nth_element(keypoints.begin(), keypoints.begin() + n_points, keypoints.end(), KeypointResponseGreater());
+        std::nth_element(keypoints.begin(), keypoints.begin() + n_points - 1, keypoints.end(), KeypointResponseGreater());
        //this is the boundary response, and in the case of FAST may be ambiguous
        float ambiguous_response = keypoints[n_points - 1].response;
        //use std::partition to grab all of the keypoints with the boundary response.
@@ -73,6 +73,18 @@ void CV_BRISKTest::run( int )

  Ptr<FeatureDetector> detector = BRISK::create();

+  // Check parameter get/set functions.
+  BRISK* detectorTyped = dynamic_cast<BRISK*>(detector.get());
+  ASSERT_NE(nullptr, detectorTyped);
+  detectorTyped->setOctaves(3);
+  detectorTyped->setThreshold(30);
+  ASSERT_EQ(detectorTyped->getOctaves(), 3);
+  ASSERT_EQ(detectorTyped->getThreshold(), 30);
+  detectorTyped->setOctaves(4);
+  detectorTyped->setThreshold(29);
+  ASSERT_EQ(detectorTyped->getOctaves(), 4);
+  ASSERT_EQ(detectorTyped->getThreshold(), 29);
+
  vector<KeyPoint> keypoints1;
  vector<KeyPoint> keypoints2;
  detector->detect(image1, keypoints1);
--- a/Show More
+++ b/Show More