OpenVINO version bump

* Enable Javascript bindings for photo module.

1. Enable the build flag in build_js.py.
2. Append js into WRAP list of photo’s CMakefiles.txt
3. Add photo module's API into JS API whitelist (embindgen.py)

Exposing the HDR imaging part of photo module.

[TODO]
Add tests
Fix opencv/doc/js_tutorials/

* [WIP] TODO: Add tests

* Remove TonemapDurand: algorithm patented in US, so moved to opencv_contrib

* Fix ningxin's comment: expose the base class.

* Add some more simple binding tests.

Also expose process function

3rdparty/carotene/hal/CMakeLists.txt

@@ -53,7 +53,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()
 

CMakeLists.txt

@@ -219,71 +219,200 @@ 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" ON
+  VISIBLE_IF WIN32
+  VERIFY HAVE_MSMF)
+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 +452,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 +472,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) )
@@ -1479,7 +1608,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 +1621,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 +1656,10 @@ status("")
 
 ocv_finalize_status()
 
+if(ENABLE_CONFIG_VERIFICATION)
+  ocv_verify_config()
+endif()
+
 ocv_cmake_hook(POST_FINALIZE)
 
 # ----------------------------------------------------------------------------

cmake/OpenCVCompilerOptimizations.cmake

@@ -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()

cmake/OpenCVGenABI.cmake

@@ -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")

cmake/OpenCVModule.cmake

@@ -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()

cmake/OpenCVUtils.cmake

@@ -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)

cmake/OpenCVVersion.cmake

@@ -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")

cmake/checks/OpenCVDetectCudaArch.cu

@@ -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;

cmake/checks/cpu_vsx.cpp

@@ -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()
 {

cmake/checks/cpu_vsx3.cpp

@@ -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;
+}

cmake/templates/opencv_abi.xml.in

@@ -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>
 

doc/py_tutorials/py_photo/py_hdr/py_hdr.markdown

@@ -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.

doc/tutorials/photo/hdr_imaging/hdr_imaging.markdown

@@ -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

modules/core/CMakeLists.txt

@@ -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)

modules/core/include/opencv2/core/cv_cpu_dispatch.h

@@ -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

modules/core/include/opencv2/core/cv_cpu_helper.h

@@ -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 */

modules/core/include/opencv2/core/cvdef.h

@@ -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
 

modules/core/include/opencv2/core/hal/intrin_avx.hpp

@@ -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)
@@ -1128,6 +1133,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)                \

modules/core/include/opencv2/core/hal/intrin_cpp.hpp

@@ -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.

modules/core/include/opencv2/core/hal/intrin_neon.hpp

@@ -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)
@@ -992,6 +999,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) \
 { \

modules/core/include/opencv2/core/hal/intrin_sse.hpp

@@ -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)
@@ -1472,6 +1477,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 +1965,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)

modules/core/include/opencv2/core/hal/intrin_vsx.hpp

@@ -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)
@@ -734,6 +739,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)

modules/core/include/opencv2/core/persistence.hpp

@@ -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;
 

modules/core/include/opencv2/core/sse_utils.hpp

@@ -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);

modules/core/include/opencv2/core/types.hpp

@@ -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)
 {

modules/core/include/opencv2/core/types_c.h

@@ -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();

modules/core/include/opencv2/core/version.hpp

@@ -7,8 +7,8 @@
 
 #define CV_VERSION_MAJOR    4
 #define CV_VERSION_MINOR    0
-#define CV_VERSION_REVISION 0
-#define CV_VERSION_STATUS   ""
+#define CV_VERSION_REVISION 1
+#define CV_VERSION_STATUS   "-openvino"
 
 #define CVAUX_STR_EXP(__A)  #__A
 #define CVAUX_STR(__A)      CVAUX_STR_EXP(__A)

modules/core/perf/perf_norm.cpp

@@ -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

modules/core/src/arithm.cpp

@@ -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;
     }
 };

modules/core/src/norm.cpp

@@ -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;

modules/core/src/ocl.cpp

@@ -2772,6 +2772,7 @@ struct Kernel::Impl
         for( int i = 0; i < MAX_ARRS; i++ )
             u[i] = 0;
         haveTempDstUMats = false;
+        haveTempSrcUMats = false;
     }
 
     void cleanupUMats()
@@ -2788,6 +2789,7 @@ struct Kernel::Impl
             }
         nu = 0;
         haveTempDstUMats = false;
+        haveTempSrcUMats = false;
     }
 
     void addUMat(const UMat& m, bool dst)
@@ -2798,6 +2800,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 +2839,7 @@ struct Kernel::Impl
     int nu;
     std::list<Image2D> images;
     bool haveTempDstUMats;
+    bool haveTempSrcUMats;
 };
 
 }} // namespace cv::ocl
@@ -3108,6 +3113,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;

modules/core/src/parallel.cpp

@@ -454,7 +454,7 @@ static inline int _initMaxThreads()
     {
         omp_set_dynamic(maxThreads);
     }
-    return numThreads;
+    return maxThreads;
 }
 static int numThreadsMax = _initMaxThreads();
 #elif defined HAVE_GCD

modules/core/src/persistence.cpp

@@ -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); }

modules/core/src/persistence.hpp

@@ -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;
 };

modules/core/src/system.cpp

@@ -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 };

modules/core/src/utils/datafile.cpp

@@ -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

modules/core/src/utils/filesystem.cpp

@@ -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 }

modules/core/src/va_intel.cpp

@@ -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];

modules/core/test/test_operations.cpp

@@ -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);

modules/dnn/include/opencv2/dnn/all_layers.hpp

@@ -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
     {

modules/dnn/include/opencv2/dnn/dnn.hpp

@@ -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);
 

modules/dnn/include/opencv2/dnn/version.hpp

@@ -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 20181205
 
 #if !defined CV_DOXYGEN && !defined CV_DNN_DONT_ADD_INLINE_NS
 #define CV__DNN_INLINE_NS __CV_CAT(dnn4_v, OPENCV_DNN_API_VERSION)

modules/dnn/perf/perf_net.cpp

@@ -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);

modules/dnn/src/dnn.cpp

@@ -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)

modules/dnn/src/init.cpp

@@ -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);

modules/dnn/src/layers/blank_layer.cpp

@@ -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

modules/dnn/src/layers/const_layer.cpp

@@ -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

modules/dnn/src/layers/convolution_layer.cpp

@@ -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());

modules/dnn/src/layers/elementwise_layers.cpp

@@ -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
 

modules/dnn/src/layers/eltwise_layer.cpp

@@ -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)

modules/dnn/src/onnx/onnx_importer.cpp

@@ -358,7 +358,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);
                 }
             }
@@ -375,10 +376,28 @@ void ONNXImporter::populateNet(Net dstNet)
                 layerParams.set("shift", blob.at<float>(0));
             }
             else {
-                layerParams.type = "Shift";
+                layerParams.type = "Scale";
+                layerParams.set("has_bias", true);
                 layerParams.blobs.push_back(blob);
             }
         }
+        else if (layer_type == "Div")
+        {
+            Mat blob = getBlob(node_proto, constBlobs, 1);
+            CV_Assert_N(blob.type() == CV_32F, blob.total());
+            divide(1.0, blob, blob);
+            if (blob.total() == 1)
+            {
+                layerParams.set("scale", blob.at<float>(0));
+                layerParams.type = "Power";
+            }
+            else
+            {
+                layerParams.type = "Scale";
+                layerParams.blobs.push_back(blob);
+                layerParams.set("bias_term", false);
+            }
+        }
         else if (layer_type == "Constant")
         {
             CV_Assert(node_proto.input_size() == 0);

modules/dnn/src/op_inf_engine.cpp

@@ -233,6 +233,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 +308,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 +321,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 +424,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 +476,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 +504,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[] = {

modules/dnn/src/op_inf_engine.hpp

@@ -68,6 +68,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;

modules/dnn/src/tensorflow/tf_importer.cpp

@@ -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));

modules/dnn/src/vkcom/include/op_base.hpp

@@ -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_;

modules/dnn/src/vkcom/include/vkcom.hpp

@@ -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

modules/dnn/src/vkcom/src/buffer.cpp

@@ -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)) &&

modules/dnn/src/vkcom/src/common.hpp

@@ -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"); \
         } \
 }

modules/dnn/src/vkcom/src/context.cpp

@@ -0,0 +1,297 @@
+// 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) {
+            throw std::runtime_error("Could not load vkDestroyDebugReportCallbackEXT");
+        }
+        func(kInstance, kDebugReportCallback, NULL);
+    }
+    kShaders.clear();
+    vkDestroyInstance(kInstance, NULL);
+
+    return;
+}
+
+#endif // HAVE_VULKAN
+
+}}} // namespace cv::dnn::vkcom

modules/dnn/src/vkcom/src/context.hpp

@@ -7,7 +7,6 @@
 
 #ifndef OPENCV_DNN_VKCOM_CONTEXT_HPP
 #define OPENCV_DNN_VKCOM_CONTEXT_HPP
-#include "common.hpp"
 
 namespace cv { namespace dnn { namespace vkcom {
 
@@ -15,13 +14,12 @@ namespace cv { namespace dnn { namespace vkcom {
 
 struct Context
 {
-    VkDevice device;
-    VkQueue queue;
-    VkCommandPool cmd_pool;
-    std::map<std::string, VkShaderModule> shader_modules;
-    int ref;
+    Context();
+    ~Context();
 };
 
+void createContext();
+
 #endif // HAVE_VULKAN
 
 }}} // namespace cv::dnn::vkcom

modules/dnn/src/vkcom/src/op_base.cpp

@@ -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;
@@ -139,7 +139,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 +150,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 +177,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);
 }

modules/dnn/src/vkcom/src/tensor.cpp

@@ -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);
 }
 

modules/dnn/src/vkcom/src/vkcom.cpp

@@ -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

modules/dnn/src/vkcom/vulkan/vk_loader.hpp

@@ -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

modules/dnn/test/test_backends.cpp

@@ -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
     if (backend == DNN_BACKEND_INFERENCE_ENGINE && target == DNN_TARGET_MYRIAD)
         throw SkipTestException("Test is enabled starts from OpenVINO 2018R3");
+#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), "");
 }

modules/dnn/test/test_caffe_importer.cpp

@@ -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)
 {

modules/dnn/test/test_common.hpp

@@ -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)
            {

modules/dnn/test/test_darknet_importer.cpp

@@ -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");
 }
 

modules/dnn/test/test_googlenet.cpp

@@ -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

modules/dnn/test/test_halide_layers.cpp

@@ -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;

modules/dnn/test/test_ie_models.cpp

@@ -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[] = {
@@ -202,7 +203,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 +246,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

modules/dnn/test/test_layers.cpp

@@ -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
@@ -349,12 +354,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());
 

modules/dnn/test/test_misc.cpp

@@ -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)

modules/dnn/test/test_tf_importer.cpp

@@ -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)

modules/features2d/include/opencv2/features2d.hpp

@@ -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

modules/features2d/src/brisk.cpp

@@ -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

modules/features2d/src/draw.cpp

@@ -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 );

modules/features2d/test/test_brisk.cpp

@@ -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);

modules/flann/include/opencv2/flann/dist.h

@@ -462,10 +462,9 @@ struct Hamming
             }
         }
 #else // NO NEON and NOT GNUC
-        typedef unsigned long long pop_t;
         HammingLUT lut;
         result = lut(reinterpret_cast<const unsigned char*> (a),
-                     reinterpret_cast<const unsigned char*> (b), size * sizeof(pop_t));
+                     reinterpret_cast<const unsigned char*> (b), size);
 #endif
         return result;
     }

modules/gapi/CMakeLists.txt

@@ -23,6 +23,7 @@ file(GLOB gapi_ext_hdrs
     "${CMAKE_CURRENT_LIST_DIR}/include/opencv2/${name}/util/*.hpp"
     "${CMAKE_CURRENT_LIST_DIR}/include/opencv2/${name}/cpu/*.hpp"
     "${CMAKE_CURRENT_LIST_DIR}/include/opencv2/${name}/gpu/*.hpp"
+    "${CMAKE_CURRENT_LIST_DIR}/include/opencv2/${name}/ocl/*.hpp"
     "${CMAKE_CURRENT_LIST_DIR}/include/opencv2/${name}/fluid/*.hpp"
     "${CMAKE_CURRENT_LIST_DIR}/include/opencv2/${name}/own/*.hpp"
     )
@@ -72,11 +73,11 @@ set(gapi_srcs
     src/backends/fluid/gfluidimgproc_func.dispatch.cpp
     src/backends/fluid/gfluidcore.cpp
 
-    # GPU Backend (currently built-in)
-    src/backends/gpu/ggpubackend.cpp
-    src/backends/gpu/ggpukernel.cpp
-    src/backends/gpu/ggpuimgproc.cpp
-    src/backends/gpu/ggpucore.cpp
+    # OCL Backend (currently built-in)
+    src/backends/ocl/goclbackend.cpp
+    src/backends/ocl/goclkernel.cpp
+    src/backends/ocl/goclimgproc.cpp
+    src/backends/ocl/goclcore.cpp
 
     # Compound
     src/backends/common/gcompoundbackend.cpp

modules/gapi/cmake/init.cmake

@@ -1,3 +1,9 @@
+OCV_OPTION(WITH_ADE "Enable ADE framework (required for Graph API module)" ON)
+
+if(NOT WITH_ADE)
+  return()
+endif()
+
 if (ade_DIR)
   # if ade_DIR is set, use ADE-supplied CMake script
   # to set up variables to the prebuilt ADE

modules/gapi/include/opencv2/gapi/fluid/gfluidbuffer.hpp

@@ -45,21 +45,37 @@ class GAPI_EXPORTS Buffer;
 class GAPI_EXPORTS View
 {
 public:
+    struct Cache
+    {
+        std::vector<const uint8_t*> m_linePtrs;
+        GMatDesc m_desc;
+        int m_border_size = 0;
+
+        inline const uint8_t* linePtr(int index) const
+        {
+            return m_linePtrs[index + m_border_size];
+        }
+    };
+
     View() = default;
 
-    const uint8_t* InLineB(int index) const; // -(w-1)/2...0...+(w-1)/2 for Filters
+    const inline uint8_t* InLineB(int index) const // -(w-1)/2...0...+(w-1)/2 for Filters
+    {
+        return m_cache->linePtr(index);
+    }
+
     template<typename T> const inline T* InLine(int i) const
     {
         const uint8_t* ptr = this->InLineB(i);
         return reinterpret_cast<const T*>(ptr);
     }
 
-    operator bool() const;
+    inline operator bool() const { return m_priv != nullptr; }
     bool ready() const;
-    int length() const;
+    inline int length() const { return m_cache->m_desc.size.width; }
     int y() const;
 
-    const GMatDesc& meta() const;
+    inline const GMatDesc& meta() const { return m_cache->m_desc; }
 
     class GAPI_EXPORTS Priv;      // internal use only
     Priv& priv();               // internal use only
@@ -69,11 +85,18 @@ public:
 
 private:
     std::shared_ptr<Priv> m_priv;
+    const Cache* m_cache;
 };
 
 class GAPI_EXPORTS Buffer
 {
 public:
+    struct Cache
+    {
+        std::vector<uint8_t*> m_linePtrs;
+        GMatDesc m_desc;
+    };
+
     // Default constructor (executable creation stage,
     // all following initialization performed in Priv::init())
     Buffer();
@@ -89,7 +112,11 @@ public:
     // Constructor for in/out buffers (for tests)
     Buffer(const cv::gapi::own::Mat &data, bool is_input);
 
-    uint8_t* OutLineB(int index = 0);
+    inline uint8_t* OutLineB(int index = 0)
+    {
+        return m_cache->m_linePtrs[index];
+    }
+
     template<typename T> inline T* OutLine(int index = 0)
     {
         uint8_t* ptr = this->OutLineB(index);
@@ -100,10 +127,10 @@ public:
 
     int linesReady() const;
     void debug(std::ostream &os) const;
-    int length() const;
+    inline int length() const { return m_cache->m_desc.size.width; }
     int lpi() const;  // LPI for WRITER
 
-    const GMatDesc& meta() const;
+    inline const GMatDesc& meta() const { return m_cache->m_desc; }
 
     View mkView(int borderSize, bool ownStorage);
 
@@ -113,7 +140,7 @@ public:
 
 private:
     std::shared_ptr<Priv> m_priv;
-
+    const Cache* m_cache;
 };
 
 } // namespace cv::gapi::fluid

modules/gapi/include/opencv2/gapi/gpu/core.hpp

@@ -7,17 +7,17 @@
 
 #ifndef OPENCV_GAPI_GPU_CORE_API_HPP
 #define OPENCV_GAPI_GPU_CORE_API_HPP
+/** @file
+* @deprecated Use "opencv2/gapi/ocl/core.hpp" instead.
+*/
 
-#include <opencv2/core/cvdef.h>     // GAPI_EXPORTS
-#include <opencv2/gapi/gkernel.hpp> // GKernelPackage
+#include "opencv2/gapi/ocl/core.hpp"
 
 namespace cv {
 namespace gapi {
 namespace core {
 namespace gpu {
-
-GAPI_EXPORTS GKernelPackage kernels();
-
+    using namespace ocl;
 } // namespace gpu
 } // namespace core
 } // namespace gapi

modules/gapi/include/opencv2/gapi/gpu/ggpukernel.hpp

@@ -7,238 +7,12 @@
 
 #ifndef OPENCV_GAPI_GGPUKERNEL_HPP
 #define OPENCV_GAPI_GGPUKERNEL_HPP
+/** @file
+* @deprecated Use "opencv2/gapi/ocl/goclkernel.hpp" instead.
+*/
 
-#include <vector>
-#include <functional>
-#include <map>
-#include <unordered_map>
+#include "opencv2/gapi/ocl/goclkernel.hpp"
+#define GAPI_GPU_KERNEL GAPI_OCL_KERNEL
 
-#include <opencv2/core/mat.hpp>
-#include <opencv2/gapi/gcommon.hpp>
-#include <opencv2/gapi/gkernel.hpp>
-#include <opencv2/gapi/garg.hpp>
-
-// FIXME: namespace scheme for backends?
-namespace cv {
-
-namespace gimpl
-{
-    // Forward-declare an internal class
-    class GGPUExecutable;
-} // namespace gimpl
-
-namespace gapi
-{
-namespace gpu
-{
-    /**
-     * \addtogroup gapi_std_backends G-API Standard backends
-     * @{
-     */
-    /**
-     * @brief Get a reference to GPU backend.
-     *
-     * At the moment, the GPU backend is built atop of OpenCV
-     * "Transparent API" (T-API), see cv::UMat for details.
-     *
-     * @sa gapi_std_backends
-     */
-    GAPI_EXPORTS cv::gapi::GBackend backend();
-    /** @} */
-} // namespace gpu
-} // namespace gapi
-
-
-// Represents arguments which are passed to a wrapped GPU function
-// FIXME: put into detail?
-class GAPI_EXPORTS GGPUContext
-{
-public:
-    // Generic accessor API
-    template<typename T>
-    const T& inArg(int input) { return m_args.at(input).get<T>(); }
-
-    // Syntax sugar
-    const cv::UMat&  inMat(int input);
-    cv::UMat&  outMatR(int output); // FIXME: Avoid cv::Mat m = ctx.outMatR()
-
-    const cv::gapi::own::Scalar& inVal(int input);
-    cv::gapi::own::Scalar& outValR(int output); // FIXME: Avoid cv::gapi::own::Scalar s = ctx.outValR()
-    template<typename T> std::vector<T>& outVecR(int output) // FIXME: the same issue
-    {
-        return outVecRef(output).wref<T>();
-    }
-
-protected:
-    detail::VectorRef& outVecRef(int output);
-
-    std::vector<GArg> m_args;
-    std::unordered_map<std::size_t, GRunArgP> m_results;
-
-
-    friend class gimpl::GGPUExecutable;
-};
-
-class GAPI_EXPORTS GGPUKernel
-{
-public:
-    // This function is kernel's execution entry point (does the processing work)
-    using F = std::function<void(GGPUContext &)>;
-
-    GGPUKernel();
-    explicit GGPUKernel(const F& f);
-
-    void apply(GGPUContext &ctx);
-
-protected:
-    F m_f;
-};
-
-// FIXME: This is an ugly ad-hoc imlpementation. TODO: refactor
-
-namespace detail
-{
-template<class T> struct gpu_get_in;
-template<> struct gpu_get_in<cv::GMat>
-{
-    static cv::UMat    get(GGPUContext &ctx, int idx) { return ctx.inMat(idx); }
-};
-template<> struct gpu_get_in<cv::GScalar>
-{
-    static cv::Scalar get(GGPUContext &ctx, int idx) { return to_ocv(ctx.inVal(idx)); }
-};
-template<typename U> struct gpu_get_in<cv::GArray<U> >
-{
-    static const std::vector<U>& get(GGPUContext &ctx, int idx) { return ctx.inArg<VectorRef>(idx).rref<U>(); }
-};
-template<class T> struct gpu_get_in
-{
-    static T get(GGPUContext &ctx, int idx) { return ctx.inArg<T>(idx); }
-};
-
-struct tracked_cv_umat{
-    //TODO Think if T - API could reallocate UMat to a proper size - how do we handle this ?
-    //tracked_cv_umat(cv::UMat& m) : r{(m)}, original_data{m.getMat(ACCESS_RW).data} {}
-    tracked_cv_umat(cv::UMat& m) : r{ (m) }, original_data{ nullptr } {}
-    cv::UMat r;
-    uchar* original_data;
-
-    operator cv::UMat& (){ return r;}
-    void validate() const{
-        //if (r.getMat(ACCESS_RW).data != original_data)
-        //{
-        //    util::throw_error
-        //        (std::logic_error
-        //         ("OpenCV kernel output parameter was reallocated. \n"
-        //          "Incorrect meta data was provided ?"));
-        //}
-
-    }
-};
-
-struct scalar_wrapper_gpu
-{
-    //FIXME reuse CPU (OpenCV) plugin code
-    scalar_wrapper_gpu(cv::gapi::own::Scalar& s) : m_s{cv::gapi::own::to_ocv(s)}, m_org_s(s) {};
-    operator cv::Scalar& () { return m_s; }
-    void writeBack() const  { m_org_s = to_own(m_s); }
-
-    cv::Scalar m_s;
-    cv::gapi::own::Scalar& m_org_s;
-};
-
-template<typename... Outputs>
-void postprocess_gpu(Outputs&... outs)
-{
-    struct
-    {
-        void operator()(tracked_cv_umat* bm) { bm->validate(); }
-        void operator()(scalar_wrapper_gpu* sw) { sw->writeBack(); }
-        void operator()(...) {                  }
-
-    } validate;
-    //dummy array to unfold parameter pack
-    int dummy[] = { 0, (validate(&outs), 0)... };
-    cv::util::suppress_unused_warning(dummy);
-}
-
-template<class T> struct gpu_get_out;
-template<> struct gpu_get_out<cv::GMat>
-{
-    static tracked_cv_umat get(GGPUContext &ctx, int idx)
-    {
-        auto& r = ctx.outMatR(idx);
-        return{ r };
-    }
-};
-template<> struct gpu_get_out<cv::GScalar>
-{
-    static scalar_wrapper_gpu get(GGPUContext &ctx, int idx)
-    {
-        auto& s = ctx.outValR(idx);
-        return{ s };
-    }
-};
-template<typename U> struct gpu_get_out<cv::GArray<U> >
-{
-    static std::vector<U>& get(GGPUContext &ctx, int idx) { return ctx.outVecR<U>(idx);  }
-};
-
-template<typename, typename, typename>
-struct GPUCallHelper;
-
-// FIXME: probably can be simplified with std::apply or analogue.
-template<typename Impl, typename... Ins, typename... Outs>
-struct GPUCallHelper<Impl, std::tuple<Ins...>, std::tuple<Outs...> >
-{
-    template<typename... Inputs>
-    struct call_and_postprocess
-    {
-        template<typename... Outputs>
-        static void call(Inputs&&... ins, Outputs&&... outs)
-        {
-            //not using a std::forward on outs is deliberate in order to
-            //cause compilation error, by tring to bind rvalue references to lvalue references
-            Impl::run(std::forward<Inputs>(ins)..., outs...);
-
-            postprocess_gpu(outs...);
-        }
-    };
-
-    template<int... IIs, int... OIs>
-    static void call_impl(GGPUContext &ctx, detail::Seq<IIs...>, detail::Seq<OIs...>)
-    {
-        //TODO: Make sure that OpenCV kernels do not reallocate memory for output parameters
-        //by comparing it's state (data ptr) before and after the call.
-        //Convert own::Scalar to cv::Scalar before call kernel and run kernel
-        //convert cv::Scalar to own::Scalar after call kernel and write back results
-        call_and_postprocess<decltype(gpu_get_in<Ins>::get(ctx, IIs))...>::call(gpu_get_in<Ins>::get(ctx, IIs)..., gpu_get_out<Outs>::get(ctx, OIs)...);
-    }
-
-    static void call(GGPUContext &ctx)
-    {
-        call_impl(ctx,
-            typename detail::MkSeq<sizeof...(Ins)>::type(),
-            typename detail::MkSeq<sizeof...(Outs)>::type());
-    }
-};
-
-} // namespace detail
-
-template<class Impl, class K>
-class GGPUKernelImpl: public detail::GPUCallHelper<Impl, typename K::InArgs, typename K::OutArgs>
-{
-    using P = detail::GPUCallHelper<Impl, typename K::InArgs, typename K::OutArgs>;
-
-public:
-    using API = K;
-
-    static cv::gapi::GBackend backend()  { return cv::gapi::gpu::backend(); }
-    static cv::GGPUKernel     kernel()   { return GGPUKernel(&P::call);     }
-};
-
-#define GAPI_GPU_KERNEL(Name, API) struct Name: public cv::GGPUKernelImpl<Name, API>
-
-} // namespace cv
 
 #endif // OPENCV_GAPI_GGPUKERNEL_HPP

modules/gapi/include/opencv2/gapi/gpu/imgproc.hpp

@@ -7,17 +7,18 @@
 
 #ifndef OPENCV_GAPI_GPU_IMGPROC_API_HPP
 #define OPENCV_GAPI_GPU_IMGPROC_API_HPP
+/** @file
+* @deprecated Use "opencv2/gapi/ocl/imgproc.hpp" instead.
+*/
+
+#include "opencv2/gapi/ocl/imgproc.hpp"
 
-#include <opencv2/core/cvdef.h>     // GAPI_EXPORTS
-#include <opencv2/gapi/gkernel.hpp> // GKernelPackage
 
 namespace cv {
 namespace gapi {
 namespace imgproc {
 namespace gpu {
-
-GAPI_EXPORTS GKernelPackage kernels();
-
+    using namespace ocl;
 } // namespace gpu
 } // namespace imgproc
 } // namespace gapi

modules/gapi/include/opencv2/gapi/ocl/core.hpp

@@ -0,0 +1,27 @@
+// 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
+
+
+#ifndef OPENCV_GAPI_OCL_CORE_API_HPP
+#define OPENCV_GAPI_OCL_CORE_API_HPP
+
+#include <opencv2/core/cvdef.h>     // GAPI_EXPORTS
+#include <opencv2/gapi/gkernel.hpp> // GKernelPackage
+
+namespace cv {
+namespace gapi {
+namespace core {
+namespace ocl {
+
+        GAPI_EXPORTS GKernelPackage kernels();
+
+} // namespace ocl
+} // namespace core
+} // namespace gapi
+} // namespace cv
+
+
+#endif // OPENCV_GAPI_OCL_CORE_API_HPP

modules/gapi/include/opencv2/gapi/ocl/goclkernel.hpp

@@ -0,0 +1,244 @@
+// 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
+
+
+#ifndef OPENCV_GAPI_GOCLKERNEL_HPP
+#define OPENCV_GAPI_GOCLKERNEL_HPP
+
+#include <vector>
+#include <functional>
+#include <map>
+#include <unordered_map>
+
+#include <opencv2/core/mat.hpp>
+#include <opencv2/gapi/gcommon.hpp>
+#include <opencv2/gapi/gkernel.hpp>
+#include <opencv2/gapi/garg.hpp>
+
+// FIXME: namespace scheme for backends?
+namespace cv {
+
+namespace gimpl
+{
+    // Forward-declare an internal class
+    class GOCLExecutable;
+} // namespace gimpl
+
+namespace gapi
+{
+namespace ocl
+{
+    /**
+     * \addtogroup gapi_std_backends G-API Standard backends
+     * @{
+     */
+    /**
+     * @brief Get a reference to OCL backend.
+     *
+     * At the moment, the OCL backend is built atop of OpenCV
+     * "Transparent API" (T-API), see cv::UMat for details.
+     *
+     * @sa gapi_std_backends
+     */
+    GAPI_EXPORTS cv::gapi::GBackend backend();
+    /** @} */
+} // namespace ocl
+} // namespace gapi
+
+
+// Represents arguments which are passed to a wrapped OCL function
+// FIXME: put into detail?
+class GAPI_EXPORTS GOCLContext
+{
+public:
+    // Generic accessor API
+    template<typename T>
+    const T& inArg(int input) { return m_args.at(input).get<T>(); }
+
+    // Syntax sugar
+    const cv::UMat&  inMat(int input);
+    cv::UMat&  outMatR(int output); // FIXME: Avoid cv::Mat m = ctx.outMatR()
+
+    const cv::gapi::own::Scalar& inVal(int input);
+    cv::gapi::own::Scalar& outValR(int output); // FIXME: Avoid cv::gapi::own::Scalar s = ctx.outValR()
+    template<typename T> std::vector<T>& outVecR(int output) // FIXME: the same issue
+    {
+        return outVecRef(output).wref<T>();
+    }
+
+protected:
+    detail::VectorRef& outVecRef(int output);
+
+    std::vector<GArg> m_args;
+    std::unordered_map<std::size_t, GRunArgP> m_results;
+
+
+    friend class gimpl::GOCLExecutable;
+};
+
+class GAPI_EXPORTS GOCLKernel
+{
+public:
+    // This function is kernel's execution entry point (does the processing work)
+    using F = std::function<void(GOCLContext &)>;
+
+    GOCLKernel();
+    explicit GOCLKernel(const F& f);
+
+    void apply(GOCLContext &ctx);
+
+protected:
+    F m_f;
+};
+
+// FIXME: This is an ugly ad-hoc imlpementation. TODO: refactor
+
+namespace detail
+{
+template<class T> struct ocl_get_in;
+template<> struct ocl_get_in<cv::GMat>
+{
+    static cv::UMat    get(GOCLContext &ctx, int idx) { return ctx.inMat(idx); }
+};
+template<> struct ocl_get_in<cv::GScalar>
+{
+    static cv::Scalar get(GOCLContext &ctx, int idx) { return to_ocv(ctx.inVal(idx)); }
+};
+template<typename U> struct ocl_get_in<cv::GArray<U> >
+{
+    static const std::vector<U>& get(GOCLContext &ctx, int idx) { return ctx.inArg<VectorRef>(idx).rref<U>(); }
+};
+template<class T> struct ocl_get_in
+{
+    static T get(GOCLContext &ctx, int idx) { return ctx.inArg<T>(idx); }
+};
+
+struct tracked_cv_umat{
+    //TODO Think if T - API could reallocate UMat to a proper size - how do we handle this ?
+    //tracked_cv_umat(cv::UMat& m) : r{(m)}, original_data{m.getMat(ACCESS_RW).data} {}
+    tracked_cv_umat(cv::UMat& m) : r{ (m) }, original_data{ nullptr } {}
+    cv::UMat r;
+    uchar* original_data;
+
+    operator cv::UMat& (){ return r;}
+    void validate() const{
+        //if (r.getMat(ACCESS_RW).data != original_data)
+        //{
+        //    util::throw_error
+        //        (std::logic_error
+        //         ("OpenCV kernel output parameter was reallocated. \n"
+        //          "Incorrect meta data was provided ?"));
+        //}
+
+    }
+};
+
+struct scalar_wrapper_ocl
+{
+    //FIXME reuse CPU (OpenCV) plugin code
+    scalar_wrapper_ocl(cv::gapi::own::Scalar& s) : m_s{cv::gapi::own::to_ocv(s)}, m_org_s(s) {};
+    operator cv::Scalar& () { return m_s; }
+    void writeBack() const  { m_org_s = to_own(m_s); }
+
+    cv::Scalar m_s;
+    cv::gapi::own::Scalar& m_org_s;
+};
+
+template<typename... Outputs>
+void postprocess_ocl(Outputs&... outs)
+{
+    struct
+    {
+        void operator()(tracked_cv_umat* bm) { bm->validate(); }
+        void operator()(scalar_wrapper_ocl* sw) { sw->writeBack(); }
+        void operator()(...) {                  }
+
+    } validate;
+    //dummy array to unfold parameter pack
+    int dummy[] = { 0, (validate(&outs), 0)... };
+    cv::util::suppress_unused_warning(dummy);
+}
+
+template<class T> struct ocl_get_out;
+template<> struct ocl_get_out<cv::GMat>
+{
+    static tracked_cv_umat get(GOCLContext &ctx, int idx)
+    {
+        auto& r = ctx.outMatR(idx);
+        return{ r };
+    }
+};
+template<> struct ocl_get_out<cv::GScalar>
+{
+    static scalar_wrapper_ocl get(GOCLContext &ctx, int idx)
+    {
+        auto& s = ctx.outValR(idx);
+        return{ s };
+    }
+};
+template<typename U> struct ocl_get_out<cv::GArray<U> >
+{
+    static std::vector<U>& get(GOCLContext &ctx, int idx) { return ctx.outVecR<U>(idx);  }
+};
+
+template<typename, typename, typename>
+struct OCLCallHelper;
+
+// FIXME: probably can be simplified with std::apply or analogue.
+template<typename Impl, typename... Ins, typename... Outs>
+struct OCLCallHelper<Impl, std::tuple<Ins...>, std::tuple<Outs...> >
+{
+    template<typename... Inputs>
+    struct call_and_postprocess
+    {
+        template<typename... Outputs>
+        static void call(Inputs&&... ins, Outputs&&... outs)
+        {
+            //not using a std::forward on outs is deliberate in order to
+            //cause compilation error, by tring to bind rvalue references to lvalue references
+            Impl::run(std::forward<Inputs>(ins)..., outs...);
+
+            postprocess_ocl(outs...);
+        }
+    };
+
+    template<int... IIs, int... OIs>
+    static void call_impl(GOCLContext &ctx, detail::Seq<IIs...>, detail::Seq<OIs...>)
+    {
+        //TODO: Make sure that OpenCV kernels do not reallocate memory for output parameters
+        //by comparing it's state (data ptr) before and after the call.
+        //Convert own::Scalar to cv::Scalar before call kernel and run kernel
+        //convert cv::Scalar to own::Scalar after call kernel and write back results
+        call_and_postprocess<decltype(ocl_get_in<Ins>::get(ctx, IIs))...>::call(ocl_get_in<Ins>::get(ctx, IIs)..., ocl_get_out<Outs>::get(ctx, OIs)...);
+    }
+
+    static void call(GOCLContext &ctx)
+    {
+        call_impl(ctx,
+            typename detail::MkSeq<sizeof...(Ins)>::type(),
+            typename detail::MkSeq<sizeof...(Outs)>::type());
+    }
+};
+
+} // namespace detail
+
+template<class Impl, class K>
+class GOCLKernelImpl: public detail::OCLCallHelper<Impl, typename K::InArgs, typename K::OutArgs>
+{
+    using P = detail::OCLCallHelper<Impl, typename K::InArgs, typename K::OutArgs>;
+
+public:
+    using API = K;
+
+    static cv::gapi::GBackend backend()  { return cv::gapi::ocl::backend(); }
+    static cv::GOCLKernel     kernel()   { return GOCLKernel(&P::call);     }
+};
+
+#define GAPI_OCL_KERNEL(Name, API) struct Name: public cv::GOCLKernelImpl<Name, API>
+
+} // namespace cv
+
+#endif // OPENCV_GAPI_GOCLKERNEL_HPP

modules/gapi/include/opencv2/gapi/ocl/imgproc.hpp

@@ -0,0 +1,27 @@
+// 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
+
+
+#ifndef OPENCV_GAPI_OCL_IMGPROC_API_HPP
+#define OPENCV_GAPI_OCL_IMGPROC_API_HPP
+
+#include <opencv2/core/cvdef.h>     // GAPI_EXPORTS
+#include <opencv2/gapi/gkernel.hpp> // GKernelPackage
+
+namespace cv {
+namespace gapi {
+namespace imgproc {
+namespace ocl {
+
+    GAPI_EXPORTS GKernelPackage kernels();
+
+} // namespace ocl
+} // namespace imgproc
+} // namespace gapi
+} // namespace cv
+
+
+#endif // OPENCV_GAPI_OCL_IMGPROC_API_HPP

modules/gapi/include/opencv2/gapi/own/assert.hpp

@@ -8,9 +8,10 @@
 #ifndef OPENCV_GAPI_OWN_ASSERT_HPP
 #define OPENCV_GAPI_OWN_ASSERT_HPP
 
-#if 0
+#if !defined(GAPI_STANDALONE)
 #include <opencv2/core/base.hpp>
-#define GAPI_Assert(expr) CV_Assert(expr)
+#define GAPI_Assert CV_Assert
+#define GAPI_DbgAssert CV_DbgAssert
 
 #else
 #include <stdexcept>
@@ -30,12 +31,13 @@ namespace detail
 #define GAPI_Assert(expr) \
 { if (!(expr)) ::detail::assert_abort(#expr, __LINE__, __FILE__, __func__); }
 
+
+#ifdef NDEBUG
+#  define GAPI_DbgAssert(expr)
+#else
+#  define GAPI_DbgAssert(expr) GAPI_Assert(expr)
 #endif
 
-#ifdef _DEBUG
-#  define GAPI_DbgAssert(expr) GAPI_Assert(expr)
-#else
-#  define GAPI_DbgAssert(expr)
-#endif
+#endif // GAPI_STANDALONE
 
 #endif // OPENCV_GAPI_OWN_ASSERT_HPP

modules/gapi/perf/common/gapi_core_perf_tests.hpp

@@ -50,9 +50,9 @@ namespace opencv_test
     class MaxPerfTest : public TestPerfParams<tuple<cv::Size, MatType, cv::GCompileArgs>> {};
     class AbsDiffPerfTest : public TestPerfParams<tuple<cv::Size, MatType, cv::GCompileArgs>> {};
     class AbsDiffCPerfTest : public TestPerfParams<tuple<cv::Size, MatType, cv::GCompileArgs>> {};
-    class SumPerfTest : public TestPerfParams<tuple<cv::Size, MatType, double, cv::GCompileArgs>> {};
-    class AddWeightedPerfTest : public TestPerfParams<tuple<cv::Size, MatType, int, double, cv::GCompileArgs>> {};
-    class NormPerfTest : public TestPerfParams<tuple<NormTypes, cv::Size, MatType, double, cv::GCompileArgs>> {};
+    class SumPerfTest : public TestPerfParams<tuple<compare_scalar_f, cv::Size, MatType, cv::GCompileArgs>> {};
+    class AddWeightedPerfTest : public TestPerfParams<tuple<compare_f, cv::Size, MatType, int, cv::GCompileArgs>> {};
+    class NormPerfTest : public TestPerfParams<tuple<compare_scalar_f, NormTypes, cv::Size, MatType, cv::GCompileArgs>> {};
     class IntegralPerfTest : public TestPerfParams<tuple<cv::Size, MatType, cv::GCompileArgs>> {};
     class ThresholdPerfTest : public TestPerfParams<tuple<cv::Size, MatType, int, cv::GCompileArgs>> {};
     class ThresholdOTPerfTest : public TestPerfParams<tuple<cv::Size, MatType, int, cv::GCompileArgs>> {};

modules/gapi/perf/common/gapi_core_perf_tests_inl.hpp

@@ -900,13 +900,13 @@ PERF_TEST_P_(AbsDiffCPerfTest, TestPerformance)
 
 PERF_TEST_P_(SumPerfTest, TestPerformance)
 {
-    cv::Size sz_in = get<0>(GetParam());
-    MatType type = get<1>(GetParam());
-    double tolerance = get<2>(GetParam());
+    compare_scalar_f cmpF = get<0>(GetParam());
+    cv::Size sz_in = get<1>(GetParam());
+    MatType type = get<2>(GetParam());
     cv::GCompileArgs compile_args = get<3>(GetParam());
 
 
-    initMatrixRandU(type, sz_in, false);
+    initMatrixRandU(type, sz_in, type, false);
     cv::Scalar out_sum;
     cv::Scalar out_sum_ocv;
 
@@ -928,8 +928,7 @@ PERF_TEST_P_(SumPerfTest, TestPerformance)
 
     // Comparison ////////////////////////////////////////////////////////////
     {
-        EXPECT_LE(std::abs(out_sum[0] - out_sum_ocv[0]) / std::max(1.0, std::abs(out_sum_ocv[0])), tolerance)
-            << "OCV=" << out_sum_ocv[0] << "   GAPI=" << out_sum[0];
+        EXPECT_TRUE(cmpF(out_sum, out_sum_ocv));
     }
 
     SANITY_CHECK_NOTHING();
@@ -939,10 +938,10 @@ PERF_TEST_P_(SumPerfTest, TestPerformance)
 
 PERF_TEST_P_(AddWeightedPerfTest, TestPerformance)
 {
-    cv::Size sz_in = get<0>(GetParam());
-    MatType type = get<1>(GetParam());
-    int dtype = get<2>(GetParam());
-    double tolerance = get<3>(GetParam());
+    compare_f cmpF = get<0>(GetParam());
+    cv::Size sz_in = get<1>(GetParam());
+    MatType type = get<2>(GetParam());
+    int dtype = get<3>(GetParam());
     cv::GCompileArgs compile_args = get<4>(GetParam());
 
     auto& rng = cv::theRNG();
@@ -968,45 +967,9 @@ PERF_TEST_P_(AddWeightedPerfTest, TestPerformance)
     }
 
     // Comparison ////////////////////////////////////////////////////////////
-    // FIXIT unrealiable check
-    if (0)
-    {
-        // Note, that we cannot expect bitwise results for add-weighted:
-        //
-        //    tmp = src1*alpha + src2*beta + gamma;
-        //    dst = saturate<DST>( round(tmp) );
-        //
-        // Because tmp is floating-point, dst depends on compiler optimizations
-        //
-        // However, we must expect good accuracy of tmp, and rounding correctly
+    EXPECT_TRUE(cmpF(out_mat_gapi, out_mat_ocv));
+    EXPECT_EQ(out_mat_gapi.size(), sz_in);
 
-        cv::Mat failures;
-
-        if (out_mat_ocv.type() == CV_32FC1)
-        {
-            // result: float - may vary in 7th decimal digit
-            failures = abs(out_mat_gapi - out_mat_ocv) > abs(out_mat_ocv) * 1e-6;
-        }
-        else
-        {
-            // result: integral - rounding may vary if fractional part of tmp
-            //                    is nearly 0.5
-
-            cv::Mat inexact, incorrect, diff, tmp;
-
-            inexact = out_mat_gapi != out_mat_ocv;
-
-            // even if rounded differently, check if still rounded correctly
-            cv::addWeighted(in_mat1, alpha, in_mat2, beta, gamma, tmp, CV_32F);
-            cv::subtract(out_mat_gapi, tmp, diff, cv::noArray(), CV_32F);
-            incorrect = abs(diff) >= tolerance;// 0.5000005f; // relative to 6 digits
-
-            failures = inexact & incorrect;
-        }
-
-        EXPECT_EQ(0, cv::countNonZero(failures));
-        EXPECT_EQ(out_mat_gapi.size(), sz_in);
-    }
 
     SANITY_CHECK_NOTHING();
 }
@@ -1015,10 +978,10 @@ PERF_TEST_P_(AddWeightedPerfTest, TestPerformance)
 
 PERF_TEST_P_(NormPerfTest, TestPerformance)
 {
-    NormTypes opType = get<0>(GetParam());
-    cv::Size sz = get<1>(GetParam());
-    MatType type = get<2>(GetParam());
-    double tolerance = get<3>(GetParam());
+    compare_scalar_f cmpF = get<0>(GetParam());
+    NormTypes opType = get<1>(GetParam());
+    cv::Size sz = get<2>(GetParam());
+    MatType type = get<3>(GetParam());
     cv::GCompileArgs compile_args = get<4>(GetParam());
 
 
@@ -1051,8 +1014,7 @@ PERF_TEST_P_(NormPerfTest, TestPerformance)
 
     // Comparison ////////////////////////////////////////////////////////////
     {
-        EXPECT_LE(std::abs(out_norm[0] - out_norm_ocv[0]) / std::max(1.0, std::abs(out_norm_ocv[0])), tolerance)
-            << "OCV=" << out_norm_ocv[0] << "   GAPI=" << out_norm[0];
+        EXPECT_TRUE(cmpF(out_norm, out_norm_ocv));
     }
 
     SANITY_CHECK_NOTHING();

modules/gapi/perf/common/gapi_imgproc_perf_tests_inl.hpp

@@ -52,7 +52,7 @@ PERF_TEST_P_(SepFilterPerfTest, TestPerformance)
 
     TEST_CYCLE()
     {
-      c.apply(in_mat1, out_mat_gapi, std::move(compile_args));
+      c.apply(in_mat1, out_mat_gapi);
     }
 
     // Comparison //////////////////////////////////////////////////////////////
@@ -100,7 +100,7 @@ PERF_TEST_P_(Filter2DPerfTest, TestPerformance)
 
     TEST_CYCLE()
     {
-        c.apply(in_mat1, out_mat_gapi, std::move(compile_args));
+        c.apply(in_mat1, out_mat_gapi);
     }
 
     // Comparison //////////////////////////////////////////////////////////////
@@ -145,7 +145,7 @@ PERF_TEST_P_(BoxFilterPerfTest, TestPerformance)
 
     TEST_CYCLE()
     {
-        c.apply(in_mat1, out_mat_gapi, std::move(compile_args));
+        c.apply(in_mat1, out_mat_gapi);
     }
 
     // Comparison //////////////////////////////////////////////////////////////
@@ -188,7 +188,7 @@ PERF_TEST_P_(BlurPerfTest, TestPerformance)
 
     TEST_CYCLE()
     {
-        c.apply(in_mat1, out_mat_gapi, std::move(compile_args));
+        c.apply(in_mat1, out_mat_gapi);
     }
 
     // Comparison //////////////////////////////////////////////////////////////
@@ -230,7 +230,7 @@ PERF_TEST_P_(GaussianBlurPerfTest, TestPerformance)
 
     TEST_CYCLE()
     {
-        c.apply(in_mat1, out_mat_gapi, std::move(compile_args));
+        c.apply(in_mat1, out_mat_gapi);
     }
 
     // Comparison //////////////////////////////////////////////////////////////
@@ -271,7 +271,7 @@ PERF_TEST_P_(MedianBlurPerfTest, TestPerformance)
 
     TEST_CYCLE()
     {
-        c.apply(in_mat1, out_mat_gapi, std::move(compile_args));
+        c.apply(in_mat1, out_mat_gapi);
     }
 
     // Comparison //////////////////////////////////////////////////////////////
@@ -314,7 +314,7 @@ PERF_TEST_P_(ErodePerfTest, TestPerformance)
 
     TEST_CYCLE()
     {
-        c.apply(in_mat1, out_mat_gapi, std::move(compile_args));
+        c.apply(in_mat1, out_mat_gapi);
     }
 
     // Comparison //////////////////////////////////////////////////////////////
@@ -357,7 +357,7 @@ PERF_TEST_P_(Erode3x3PerfTest, TestPerformance)
 
     TEST_CYCLE()
     {
-        c.apply(in_mat1, out_mat_gapi, std::move(compile_args));
+        c.apply(in_mat1, out_mat_gapi);
     }
 
     // Comparison //////////////////////////////////////////////////////////////
@@ -400,7 +400,7 @@ PERF_TEST_P_(DilatePerfTest, TestPerformance)
 
     TEST_CYCLE()
     {
-        c.apply(in_mat1, out_mat_gapi, std::move(compile_args));
+        c.apply(in_mat1, out_mat_gapi);
     }
 
     // Comparison //////////////////////////////////////////////////////////////
@@ -443,7 +443,7 @@ PERF_TEST_P_(Dilate3x3PerfTest, TestPerformance)
 
     TEST_CYCLE()
     {
-        c.apply(in_mat1, out_mat_gapi, std::move(compile_args));
+        c.apply(in_mat1, out_mat_gapi);
     }
 
     // Comparison //////////////////////////////////////////////////////////////
@@ -526,7 +526,7 @@ PERF_TEST_P_(CannyPerfTest, TestPerformance)
 
     TEST_CYCLE()
     {
-        c.apply(in_mat1, out_mat_gapi, std::move(compile_args));
+        c.apply(in_mat1, out_mat_gapi);
     }
 
     // Comparison //////////////////////////////////////////////////////////////
@@ -564,7 +564,7 @@ PERF_TEST_P_(EqHistPerfTest, TestPerformance)
 
     TEST_CYCLE()
     {
-        c.apply(in_mat1, out_mat_gapi, std::move(compile_args));
+        c.apply(in_mat1, out_mat_gapi);
     }
 
     // Comparison //////////////////////////////////////////////////////////////
@@ -678,7 +678,7 @@ PERF_TEST_P_(RGB2YUVPerfTest, TestPerformance)
 
     TEST_CYCLE()
     {
-        c.apply(in_mat1, out_mat_gapi, std::move(compile_args));
+        c.apply(in_mat1, out_mat_gapi);
     }
 
     // Comparison //////////////////////////////////////////////////////////////
@@ -716,7 +716,7 @@ PERF_TEST_P_(YUV2RGBPerfTest, TestPerformance)
 
     TEST_CYCLE()
     {
-        c.apply(in_mat1, out_mat_gapi, std::move(compile_args));
+        c.apply(in_mat1, out_mat_gapi);
     }
 
     // Comparison //////////////////////////////////////////////////////////////
@@ -754,7 +754,7 @@ PERF_TEST_P_(RGB2LabPerfTest, TestPerformance)
 
     TEST_CYCLE()
     {
-        c.apply(in_mat1, out_mat_gapi, std::move(compile_args));
+        c.apply(in_mat1, out_mat_gapi);
     }
 
     // Comparison //////////////////////////////////////////////////////////////
@@ -792,7 +792,7 @@ PERF_TEST_P_(BGR2LUVPerfTest, TestPerformance)
 
     TEST_CYCLE()
     {
-        c.apply(in_mat1, out_mat_gapi, std::move(compile_args));
+        c.apply(in_mat1, out_mat_gapi);
     }
 
     // Comparison //////////////////////////////////////////////////////////////
@@ -830,7 +830,7 @@ PERF_TEST_P_(LUV2BGRPerfTest, TestPerformance)
 
     TEST_CYCLE()
     {
-        c.apply(in_mat1, out_mat_gapi, std::move(compile_args));
+        c.apply(in_mat1, out_mat_gapi);
     }
 
     // Comparison //////////////////////////////////////////////////////////////
@@ -864,7 +864,7 @@ PERF_TEST_P_(BGR2YUVPerfTest, TestPerformance)
 
     TEST_CYCLE()
     {
-        c.apply(in_mat1, out_mat_gapi, std::move(compile_args));
+        c.apply(in_mat1, out_mat_gapi);
     }
 
     EXPECT_TRUE(cmpF(out_mat_gapi, out_mat_ocv));
@@ -894,7 +894,7 @@ PERF_TEST_P_(YUV2BGRPerfTest, TestPerformance)
 
     TEST_CYCLE()
     {
-        c.apply(in_mat1, out_mat_gapi, std::move(compile_args));
+        c.apply(in_mat1, out_mat_gapi);
     }
 
     EXPECT_TRUE(cmpF(out_mat_gapi, out_mat_ocv));

modules/gapi/perf/cpu/gapi_core_perf_tests_cpu.cpp

@@ -152,24 +152,24 @@ INSTANTIATE_TEST_CASE_P(AbsDiffCPerfTestCPU, AbsDiffCPerfTest,
         Values(cv::compile_args(CORE_CPU))));
 
 INSTANTIATE_TEST_CASE_P(SumPerfTestCPU, SumPerfTest,
-    Combine(Values(szSmall128, szVGA, sz720p, sz1080p),
+    Combine(Values(AbsToleranceScalar(0.0).to_compare_f()),
+        Values(szSmall128, szVGA, sz720p, sz1080p),
         Values(CV_8UC1, CV_8UC3, CV_16UC1, CV_16SC1, CV_32FC1),
-        Values(0.0),
+        //Values(0.0),
         Values(cv::compile_args(CORE_CPU))));
 
-// FIXME: Comparison introduced by YL doesn't work with C3
 INSTANTIATE_TEST_CASE_P(AddWeightedPerfTestCPU, AddWeightedPerfTest,
-    Combine(Values(szSmall128, szVGA, sz720p, sz1080p),
-        Values(CV_8UC1, /*CV_8UC3,*/ CV_16UC1, CV_16SC1, CV_32FC1),
+    Combine(Values(Tolerance_FloatRel_IntAbs(1e-6, 1).to_compare_f()),
+        Values(szSmall128, szVGA, sz720p, sz1080p),
+        Values(CV_8UC1, CV_8UC3, CV_16UC1, CV_16SC1, CV_32FC1),
         Values(-1, CV_8U, CV_16U, CV_32F),
-        Values(0.5000005),
         Values(cv::compile_args(CORE_CPU))));
 
 INSTANTIATE_TEST_CASE_P(NormPerfTestCPU, NormPerfTest,
-    Combine(Values(NORM_INF, NORM_L1, NORM_L2),
+    Combine(Values(AbsToleranceScalar(0.0).to_compare_f()),
+        Values(NORM_INF, NORM_L1, NORM_L2),
         Values(szSmall128, szVGA, sz720p, sz1080p),
         Values(CV_8UC1, CV_8UC3, CV_16UC1, CV_16SC1, CV_32FC1),
-        Values(0.0),
         Values(cv::compile_args(CORE_CPU))));
 
 INSTANTIATE_TEST_CASE_P(IntegralPerfTestCPU, IntegralPerfTest,

modules/gapi/perf/cpu/gapi_imgproc_perf_tests_fluid.cpp

@@ -13,9 +13,101 @@
 namespace opencv_test
 {
 
-    INSTANTIATE_TEST_CASE_P(SobelPerfTestFluid, SobelPerfTest,
-        Combine(Values(AbsExact().to_compare_f()),
-            Values(CV_8UC1, CV_8UC3, CV_16UC1, CV_16SC1),  // add CV_32FC1 when ready
+INSTANTIATE_TEST_CASE_P(SepFilterPerfTestFluid_8U, SepFilterPerfTest,
+    Combine(Values(ToleranceFilter(1e-4f, 0.01).to_compare_f()),
+            Values(CV_8UC1, CV_8UC3),
+            Values(3),
+            Values(szVGA, sz720p, sz1080p),
+            Values(-1, CV_16S, CV_32F),
+            Values(cv::compile_args(IMGPROC_FLUID))));
+
+INSTANTIATE_TEST_CASE_P(SepFilterPerfTestFluid_other, SepFilterPerfTest,
+    Combine(Values(ToleranceFilter(1e-4f, 0.01).to_compare_f()),
+            Values(CV_16UC1, CV_16SC1, CV_32FC1),
+            Values(3),
+            Values(szVGA, sz720p, sz1080p),
+            Values(-1, CV_32F),
+            Values(cv::compile_args(IMGPROC_FLUID))));
+
+INSTANTIATE_TEST_CASE_P(Filter2DPerfTestFluid, Filter2DPerfTest,
+    Combine(Values(ToleranceFilter(1e-4f, 0.01).to_compare_f()),
+            Values(CV_8UC1, CV_8UC3, CV_16UC1, CV_16SC1, CV_32FC1),
+            Values(3),                                     // add 4, 5, 7 when kernel is ready
+            Values(szVGA, sz720p, sz1080p),
+            Values(cv::BORDER_DEFAULT),
+            Values(-1, CV_32F),
+            Values(cv::compile_args(IMGPROC_FLUID))));
+
+INSTANTIATE_TEST_CASE_P(BoxFilterPerfTestFluid, BoxFilterPerfTest,
+    Combine(Values(ToleranceFilter(1e-4f, 0.01).to_compare_f()),
+            Values(CV_8UC1, CV_8UC3, CV_16UC1, CV_16SC1, CV_32FC1),
+            Values(3),                                     // add size=5, when kernel is ready
+            Values(szVGA, sz720p, sz1080p),
+            Values(cv::BORDER_DEFAULT),
+            Values(-1, CV_32F),
+            Values(cv::compile_args(IMGPROC_FLUID))));
+
+INSTANTIATE_TEST_CASE_P(BlurPerfTestFluid, BlurPerfTest,
+    Combine(Values(ToleranceFilter(1e-4f, 0.01).to_compare_f()),
+            Values(CV_8UC1, CV_8UC3, CV_16UC1, CV_16SC1, CV_32FC1),
+            Values(3),                                     // add size=5, when kernel is ready
+            Values(szVGA, sz720p, sz1080p),
+            Values(cv::BORDER_DEFAULT),
+            Values(cv::compile_args(IMGPROC_FLUID))));
+
+INSTANTIATE_TEST_CASE_P(GaussianBlurPerfTestFluid, GaussianBlurPerfTest,
+    Combine(Values(ToleranceFilter(1e-3f, 0.01).to_compare_f()),
+            Values(CV_8UC1, CV_8UC3, CV_16UC1, CV_16SC1, CV_32FC1),
+            Values(3),                                     // add size=5, when kernel is ready
+            Values(szVGA, sz720p, sz1080p),
+            Values(cv::compile_args(IMGPROC_FLUID))));
+
+INSTANTIATE_TEST_CASE_P(MedianBlurPerfTestFluid, MedianBlurPerfTest,
+    Combine(Values(AbsExact().to_compare_f()),
+            Values(CV_8UC1, CV_8UC3, CV_16UC1, CV_16SC1, CV_32FC1),
+            Values(3),                                     // add size=5, when kernel is ready
+            Values(szVGA, sz720p, sz1080p),
+            Values(cv::compile_args(IMGPROC_FLUID))));
+
+INSTANTIATE_TEST_CASE_P(ErodePerfTestFluid, ErodePerfTest,
+    Combine(Values(AbsExact().to_compare_f()),
+            Values(CV_8UC1, CV_8UC3, CV_16UC1, CV_16SC1, CV_32FC1),
+            Values(3),                                     // add size=5, when kernel is ready
+            Values(szVGA, sz720p, sz1080p),
+            Values(cv::MorphShapes::MORPH_RECT,
+                   cv::MorphShapes::MORPH_CROSS,
+                   cv::MorphShapes::MORPH_ELLIPSE),
+            Values(cv::compile_args(IMGPROC_FLUID))));
+
+// GAPI/fluid does not support iterations parameter for the Erode kernel
+INSTANTIATE_TEST_CASE_P(DISABLED_Erode3x3PerfTestFluid, Erode3x3PerfTest,
+    Combine(Values(AbsExact().to_compare_f()),
+            Values(CV_8UC1, CV_8UC3, CV_16UC1, CV_16SC1, CV_32FC1),
+            Values(szVGA, sz720p, sz1080p),
+            Values(1, 2, 4),
+            Values(cv::compile_args(IMGPROC_FLUID))));
+
+INSTANTIATE_TEST_CASE_P(DilatePerfTestFluid, DilatePerfTest,
+    Combine(Values(AbsExact().to_compare_f()),
+            Values(CV_8UC1, CV_8UC3, CV_16UC1, CV_16SC1, CV_32FC1),
+            Values(3),                                     // add size=5, when kernel is ready
+            Values(szVGA, sz720p, sz1080p),
+            Values(cv::MorphShapes::MORPH_RECT,
+                   cv::MorphShapes::MORPH_CROSS,
+                   cv::MorphShapes::MORPH_ELLIPSE),
+            Values(cv::compile_args(IMGPROC_FLUID))));
+
+// GAPI/fluid does not support iterations parameter for the Dilate kernel
+INSTANTIATE_TEST_CASE_P(DISABLED_Dilate3x3PerfTestFluid, Dilate3x3PerfTest,
+    Combine(Values(AbsExact().to_compare_f()),
+            Values(CV_8UC1, CV_8UC3, CV_16UC1, CV_16SC1, CV_32FC1),
+            Values(szVGA, sz720p, sz1080p),
+            Values(1, 2, 4),
+            Values(cv::compile_args(IMGPROC_FLUID))));
+
+INSTANTIATE_TEST_CASE_P(SobelPerfTestFluid, SobelPerfTest,
+    Combine(Values(AbsExact().to_compare_f()),
+            Values(CV_8UC1, CV_8UC3, CV_16UC1, CV_16SC1),
             Values(3),                                     // add 5x5 once supported
             Values(szVGA, sz720p, sz1080p),
             Values(-1, CV_16S, CV_32F),
@@ -23,8 +115,8 @@ namespace opencv_test
             Values(1, 2),
             Values(cv::compile_args(IMGPROC_FLUID))));
 
-    INSTANTIATE_TEST_CASE_P(SobelPerfTestFluid32F, SobelPerfTest,
-        Combine(Values(ToleranceFilter(1e-3f, 0.0).to_compare_f()),
+INSTANTIATE_TEST_CASE_P(SobelPerfTestFluid32F, SobelPerfTest,
+    Combine(Values(ToleranceFilter(1e-3f, 0.0).to_compare_f()),
             Values(CV_32FC1),
             Values(3),                                     // add 5x5 once supported
             Values(szVGA, sz720p, sz1080p),
@@ -33,14 +125,44 @@ namespace opencv_test
             Values(1, 2),
             Values(cv::compile_args(IMGPROC_FLUID))));
 
-    INSTANTIATE_TEST_CASE_P(RGB2GrayPerfTestFluid, RGB2GrayPerfTest,
-        Combine(Values(ToleranceColor(1e-3).to_compare_f()),
-                Values(szVGA, sz720p, sz1080p),
-                Values(cv::compile_args(IMGPROC_FLUID))));
+INSTANTIATE_TEST_CASE_P(RGB2GrayPerfTestFluid, RGB2GrayPerfTest,
+    Combine(Values(ToleranceColor(1e-3).to_compare_f()),
+            Values(szVGA, sz720p, sz1080p),
+            Values(cv::compile_args(IMGPROC_FLUID))));
 
-    INSTANTIATE_TEST_CASE_P(BGR2GrayPerfTestFluid, BGR2GrayPerfTest,
-        Combine(Values(ToleranceColor(1e-3).to_compare_f()),
-                Values(szVGA, sz720p, sz1080p),
-                Values(cv::compile_args(IMGPROC_FLUID))));
+INSTANTIATE_TEST_CASE_P(BGR2GrayPerfTestFluid, BGR2GrayPerfTest,
+    Combine(Values(ToleranceColor(1e-3).to_compare_f()),
+            Values(szVGA, sz720p, sz1080p),
+            Values(cv::compile_args(IMGPROC_FLUID))));
+
+INSTANTIATE_TEST_CASE_P(RGB2YUVPerfTestFluid, RGB2YUVPerfTest,
+    Combine(Values(ToleranceColor(1e-3).to_compare_f()),
+            Values(szVGA, sz720p, sz1080p),
+            Values(cv::compile_args(IMGPROC_FLUID))));
+
+INSTANTIATE_TEST_CASE_P(YUV2RGBPerfTestFluid, YUV2RGBPerfTest,
+    Combine(Values(ToleranceColor(1e-3).to_compare_f()),
+            Values(szVGA, sz720p, sz1080p),
+            Values(cv::compile_args(IMGPROC_FLUID))));
+
+INSTANTIATE_TEST_CASE_P(BGR2YUVPerfTestFluid, BGR2YUVPerfTest,
+    Combine(Values(ToleranceColor(1e-3).to_compare_f()),
+            Values(szVGA, sz720p, sz1080p),
+            Values(cv::compile_args(IMGPROC_FLUID))));
+
+INSTANTIATE_TEST_CASE_P(YUV2BGRPerfTestFluid, YUV2BGRPerfTest,
+    Combine(Values(ToleranceColor(1e-3).to_compare_f()),
+            Values(szVGA, sz720p, sz1080p),
+            Values(cv::compile_args(IMGPROC_FLUID))));
+
+INSTANTIATE_TEST_CASE_P(BGR2LUVPerfTestFluid, BGR2LUVPerfTest,
+    Combine(Values(AbsSimilarPoints(1, 0.05).to_compare_f()),
+            Values(szVGA, sz720p, sz1080p),
+            Values(cv::compile_args(IMGPROC_FLUID))));
+
+INSTANTIATE_TEST_CASE_P(RGB2LabPerfTestFluid, RGB2LabPerfTest,
+    Combine(Values(AbsSimilarPoints(1, 0.05).to_compare_f()),
+            Values(szVGA, sz720p, sz1080p),
+            Values(cv::compile_args(IMGPROC_FLUID))));
 
 }

modules/gapi/perf/gpu/gapi_core_perf_tests_gpu.cpp

@@ -7,7 +7,6 @@
 
 #include "../perf_precomp.hpp"
 #include "../common/gapi_core_perf_tests.hpp"
-#include "opencv2/gapi/gpu/core.hpp"
 
 #define CORE_GPU cv::gapi::core::gpu::kernels()
 
@@ -153,24 +152,23 @@ INSTANTIATE_TEST_CASE_P(AbsDiffCPerfTestGPU, AbsDiffCPerfTest,
                                 Values(cv::compile_args(CORE_GPU))));
 
 INSTANTIATE_TEST_CASE_P(SumPerfTestGPU, SumPerfTest,
-                        Combine(Values( szSmall128, szVGA, sz720p, sz1080p ),
+                        Combine(Values(AbsToleranceScalar(1e-5).to_compare_f()),
+                                Values( szSmall128, szVGA, sz720p, sz1080p ),
                                 Values( CV_8UC1, CV_8UC3, CV_16UC1, CV_16SC1, CV_32FC1 ),
-                                Values(4.0), //TODO: too relaxed?
                                 Values(cv::compile_args(CORE_GPU))));
 
-// FIXME: Comparison introduced by YL doesn't work with C3
 INSTANTIATE_TEST_CASE_P(AddWeightedPerfTestGPU, AddWeightedPerfTest,
-                        Combine(Values( szSmall128, szVGA, sz720p, sz1080p ),
-                                Values( CV_8UC1, /*CV_8UC3,*/ CV_16UC1, CV_16SC1, CV_32FC1 ),
+                        Combine(Values(Tolerance_FloatRel_IntAbs(1e-6, 1).to_compare_f()),
+                                Values( szSmall128, szVGA, sz720p, sz1080p ),
+                                Values( CV_8UC1, CV_8UC3, CV_16UC1, CV_16SC1, CV_32FC1 ),
                                 Values( -1, CV_8U, CV_16U, CV_32F ),
-                                Values(0.50005),
                                 Values(cv::compile_args(CORE_GPU))));
 
 INSTANTIATE_TEST_CASE_P(NormPerfTestGPU, NormPerfTest,
-                        Combine(Values(NORM_INF, NORM_L1, NORM_L2),
+                        Combine(Values(AbsToleranceScalar(1e-5).to_compare_f()),
+                                Values(NORM_INF, NORM_L1, NORM_L2),
                                 Values( szSmall128, szVGA, sz720p, sz1080p ),
                                 Values( CV_8UC1, CV_8UC3, CV_16UC1, CV_16SC1, CV_32FC1 ),
-                                Values(4.0), //TODO: too relaxed?
                                 Values(cv::compile_args(CORE_GPU))));
 
 INSTANTIATE_TEST_CASE_P(IntegralPerfTestGPU, IntegralPerfTest,

modules/gapi/perf/gpu/gapi_imgproc_perf_tests_gpu.cpp

@@ -7,7 +7,6 @@
 
 #include "../perf_precomp.hpp"
 #include "../common/gapi_imgproc_perf_tests.hpp"
-#include "opencv2/gapi/gpu/imgproc.hpp"
 
 #define IMGPROC_GPU cv::gapi::imgproc::gpu::kernels()
 
@@ -109,10 +108,20 @@ INSTANTIATE_TEST_CASE_P(Dilate3x3PerfTestGPU, Dilate3x3PerfTest,
 
 INSTANTIATE_TEST_CASE_P(SobelPerfTestGPU, SobelPerfTest,
                         Combine(Values(ToleranceFilter(1e-4f, 0.01).to_compare_f()),
-                                Values(CV_8UC1, CV_8UC3, CV_16UC1, CV_16SC1/*, CV_32FC1*/), //TODO: CV_32FC1 fails accuracy
+                                Values(CV_8UC1, CV_8UC3, CV_16UC1, CV_16SC1),
                                 Values(3, 5),
                                 Values(szVGA, sz720p, sz1080p),
-                                Values(-1, CV_32F),
+                                Values(-1, CV_16S, CV_32F),
+                                Values(0, 1),
+                                Values(1, 2),
+                                Values(cv::compile_args(IMGPROC_GPU))));
+
+INSTANTIATE_TEST_CASE_P(SobelPerfTestGPU32F, SobelPerfTest,
+                        Combine(Values(ToleranceFilter(1e-4f, 0.01).to_compare_f()),
+                                Values(CV_32FC1),
+                                Values(3, 5),
+                                Values(szVGA, sz720p, sz1080p),
+                                Values(CV_32F),
                                 Values(0, 1),
                                 Values(1, 2),
                                 Values(cv::compile_args(IMGPROC_GPU))));

modules/gapi/perf/perf_precomp.hpp

@@ -17,6 +17,8 @@
 #include "opencv2/gapi/core.hpp"
 #include "opencv2/gapi/cpu/gcpukernel.hpp"
 #include "opencv2/gapi/gpu/ggpukernel.hpp"
+#include "opencv2/gapi/gpu/imgproc.hpp"
+#include "opencv2/gapi/gpu/core.hpp"
 #include "opencv2/gapi/operators.hpp"
 
 #include "opencv2/gapi/fluid/core.hpp"

modules/gapi/src/api/gbackend.cpp

@@ -94,8 +94,12 @@ void bindInArg(Mag& mag, const RcDesc &rc, const GRunArg &arg, bool is_umat)
         case GRunArg::index_of<cv::gapi::own::Mat>() :
             if (is_umat)
             {
+#if !defined(GAPI_STANDALONE)
                 auto& mag_umat = mag.template slot<cv::UMat>()[rc.id];
                 mag_umat = to_ocv(util::get<cv::gapi::own::Mat>(arg)).getUMat(ACCESS_READ);
+#else
+                util::throw_error(std::logic_error("UMat is not supported in stadnalone build"));
+#endif // !defined(GAPI_STANDALONE)
             }
             else
             {
@@ -157,8 +161,12 @@ void bindOutArg(Mag& mag, const RcDesc &rc, const GRunArgP &arg, bool is_umat)
         case GRunArgP::index_of<cv::gapi::own::Mat*>() :
             if (is_umat)
             {
+#if !defined(GAPI_STANDALONE)
                 auto& mag_umat = mag.template slot<cv::UMat>()[rc.id];
                 mag_umat = to_ocv(*(util::get<cv::gapi::own::Mat*>(arg))).getUMat(ACCESS_RW);
+#else
+                util::throw_error(std::logic_error("UMat is not supported in standalone build"));
+#endif // !defined(GAPI_STANDALONE)
             }
             else
             {
@@ -256,7 +264,13 @@ cv::GRunArgP getObjPtr(Mag& mag, const RcDesc &rc, bool is_umat)
     {
     case GShape::GMAT:
         if (is_umat)
+        {
+#if !defined(GAPI_STANDALONE)
             return GRunArgP(&mag.template slot<cv::UMat>()[rc.id]);
+#else
+            util::throw_error(std::logic_error("UMat is not supported in standalone build"));
+#endif //  !defined(GAPI_STANDALONE)
+        }
         else
             return GRunArgP(&mag.template slot<cv::gapi::own::Mat>()[rc.id]);
     case GShape::GSCALAR: return GRunArgP(&mag.template slot<cv::gapi::own::Scalar>()[rc.id]);
@@ -300,8 +314,12 @@ void writeBack(const Mag& mag, const RcDesc &rc, GRunArgP &g_arg, bool is_umat)
         }
         if (is_umat)
         {
+#if !defined(GAPI_STANDALONE)
             auto& in_mag = mag.template slot<cv::UMat>().at(rc.id);
             GAPI_Assert((out_arg_data == (in_mag.getMat(ACCESS_RW).data)) && " data for output parameters was reallocated ?");
+#else
+            util::throw_error(std::logic_error("UMat is not supported in standalone build"));
+#endif // !defined(GAPI_STANDALONE)
         }
         else
         {

modules/gapi/src/backends/fluid/gfluidbackend.cpp

@@ -107,19 +107,29 @@ cv::gapi::GBackend cv::gapi::fluid::backend()
 // FluidAgent implementation ///////////////////////////////////////////////////
 
 namespace cv { namespace gimpl {
+struct FluidMapper
+{
+    FluidMapper(double ratio, int lpi) : m_ratio(ratio), m_lpi(lpi) {}
+    virtual ~FluidMapper() = default;
+    virtual int firstWindow(int outCoord, int lpi) const = 0;
+    virtual std::pair<int,int> linesReadAndNextWindow(int outCoord, int lpi) const = 0;
+
+protected:
+    double m_ratio = 0.0;
+    int    m_lpi   = 0;
+};
+
 struct FluidDownscaleMapper : public FluidMapper
 {
     virtual int firstWindow(int outCoord, int lpi) const override;
-    virtual int nextWindow(int outCoord, int lpi) const override;
-    virtual int linesRead(int outCoord) const override;
+    virtual std::pair<int,int> linesReadAndNextWindow(int outCoord, int lpi) const override;
     using FluidMapper::FluidMapper;
 };
 
 struct FluidUpscaleMapper : public FluidMapper
 {
     virtual int firstWindow(int outCoord, int lpi) const override;
-    virtual int nextWindow(int outCoord, int lpi) const override;
-    virtual int linesRead(int outCoord) const override;
+    virtual std::pair<int,int> linesReadAndNextWindow(int outCoord, int lpi) const override;
     FluidUpscaleMapper(double ratio, int lpi, int inHeight) : FluidMapper(ratio, lpi), m_inHeight(inHeight) {}
 private:
     int m_inHeight = 0;
@@ -129,8 +139,7 @@ struct FluidFilterAgent : public FluidAgent
 {
 private:
     virtual int firstWindow() const override;
-    virtual int nextWindow() const override;
-    virtual int linesRead() const override;
+    virtual std::pair<int,int> linesReadAndnextWindow() const override;
     virtual void setRatio(double) override { /* nothing */ }
 public:
     using FluidAgent::FluidAgent;
@@ -140,8 +149,7 @@ struct FluidResizeAgent : public FluidAgent
 {
 private:
     virtual int firstWindow() const override;
-    virtual int nextWindow() const override;
-    virtual int linesRead() const override;
+    virtual std::pair<int,int> linesReadAndnextWindow() const override;
     virtual void setRatio(double ratio) override;
 
     std::unique_ptr<FluidMapper> m_mapper;
@@ -267,35 +275,35 @@ static int borderSize(const cv::GFluidKernel& k)
     }
 }
 
-double inCoord(int outIdx, double ratio)
+inline double inCoord(int outIdx, double ratio)
 {
     return outIdx * ratio;
 }
 
-int windowStart(int outIdx, double ratio)
+inline int windowStart(int outIdx, double ratio)
 {
     return static_cast<int>(inCoord(outIdx, ratio) + 1e-3);
 }
 
-int windowEnd(int outIdx, double ratio)
+inline int windowEnd(int outIdx, double ratio)
 {
     return static_cast<int>(std::ceil(inCoord(outIdx + 1, ratio) - 1e-3));
 }
 
-double inCoordUpscale(int outCoord, double ratio)
+inline double inCoordUpscale(int outCoord, double ratio)
 {
     // Calculate the projection of output pixel's center
     return (outCoord + 0.5) * ratio - 0.5;
 }
 
-int upscaleWindowStart(int outCoord, double ratio)
+inline int upscaleWindowStart(int outCoord, double ratio)
 {
     int start = static_cast<int>(inCoordUpscale(outCoord, ratio));
     GAPI_DbgAssert(start >= 0);
     return start;
 }
 
-int upscaleWindowEnd(int outCoord, double ratio, int inSz)
+inline int upscaleWindowEnd(int outCoord, double ratio, int inSz)
 {
     int end = static_cast<int>(std::ceil(inCoordUpscale(outCoord, ratio)) + 1);
     if (end > inSz)
@@ -311,16 +319,19 @@ int cv::gimpl::FluidDownscaleMapper::firstWindow(int outCoord, int lpi) const
     return windowEnd(outCoord + lpi - 1, m_ratio) - windowStart(outCoord, m_ratio);
 }
 
-int cv::gimpl::FluidDownscaleMapper::nextWindow(int outCoord, int lpi) const
+std::pair<int,int> cv::gimpl::FluidDownscaleMapper::linesReadAndNextWindow(int outCoord, int lpi) const
 {
     auto nextStartIdx = outCoord + 1 + m_lpi - 1;
     auto nextEndIdx   = nextStartIdx + lpi - 1;
-    return windowEnd(nextEndIdx, m_ratio) - windowStart(nextStartIdx, m_ratio);
-}
 
-int cv::gimpl::FluidDownscaleMapper::linesRead(int outCoord) const
-{
-    return windowStart(outCoord + 1 + m_lpi - 1, m_ratio) - windowStart(outCoord, m_ratio);
+    auto currStart = windowStart(outCoord, m_ratio);
+    auto nextStart = windowStart(nextStartIdx, m_ratio);
+    auto nextEnd   = windowEnd(nextEndIdx, m_ratio);
+
+    auto lines_read = nextStart - currStart;
+    auto next_window = nextEnd - nextStart;
+
+    return std::make_pair(lines_read, next_window);
 }
 
 int cv::gimpl::FluidUpscaleMapper::firstWindow(int outCoord, int lpi) const
@@ -328,16 +339,19 @@ int cv::gimpl::FluidUpscaleMapper::firstWindow(int outCoord, int lpi) const
     return upscaleWindowEnd(outCoord + lpi - 1, m_ratio, m_inHeight) - upscaleWindowStart(outCoord, m_ratio);
 }
 
-int cv::gimpl::FluidUpscaleMapper::nextWindow(int outCoord, int lpi) const
+std::pair<int,int> cv::gimpl::FluidUpscaleMapper::linesReadAndNextWindow(int outCoord, int lpi) const
 {
     auto nextStartIdx = outCoord + 1 + m_lpi - 1;
     auto nextEndIdx   = nextStartIdx + lpi - 1;
-    return upscaleWindowEnd(nextEndIdx, m_ratio, m_inHeight) - upscaleWindowStart(nextStartIdx, m_ratio);
-}
 
-int cv::gimpl::FluidUpscaleMapper::linesRead(int outCoord) const
-{
-    return upscaleWindowStart(outCoord + 1 + m_lpi - 1, m_ratio) - upscaleWindowStart(outCoord, m_ratio);
+    auto currStart = upscaleWindowStart(outCoord, m_ratio);
+    auto nextStart = upscaleWindowStart(nextStartIdx, m_ratio);
+    auto nextEnd   = upscaleWindowEnd(nextEndIdx, m_ratio, m_inHeight);
+
+    auto lines_read = nextStart - currStart;
+    auto next_window = nextEnd - nextStart;
+
+    return std::make_pair(lines_read, next_window);
 }
 
 int cv::gimpl::FluidFilterAgent::firstWindow() const
@@ -345,15 +359,10 @@ int cv::gimpl::FluidFilterAgent::firstWindow() const
     return k.m_window + k.m_lpi - 1;
 }
 
-int cv::gimpl::FluidFilterAgent::nextWindow() const
+std::pair<int,int> cv::gimpl::FluidFilterAgent::linesReadAndnextWindow() const
 {
     int lpi = std::min(k.m_lpi, m_outputLines - m_producedLines - k.m_lpi);
-    return k.m_window - 1 + lpi;
-}
-
-int cv::gimpl::FluidFilterAgent::linesRead() const
-{
-    return k.m_lpi;
+    return std::make_pair(k.m_lpi, k.m_window - 1 + lpi);
 }
 
 int cv::gimpl::FluidResizeAgent::firstWindow() const
@@ -363,17 +372,11 @@ int cv::gimpl::FluidResizeAgent::firstWindow() const
     return m_mapper->firstWindow(outIdx, lpi);
 }
 
-int cv::gimpl::FluidResizeAgent::nextWindow() const
+std::pair<int,int> cv::gimpl::FluidResizeAgent::linesReadAndnextWindow() const
 {
     auto outIdx = out_buffers[0]->priv().y();
     auto lpi = std::min(m_outputLines - m_producedLines - k.m_lpi, k.m_lpi);
-    return m_mapper->nextWindow(outIdx, lpi);
-}
-
-int cv::gimpl::FluidResizeAgent::linesRead() const
-{
-    auto outIdx = out_buffers[0]->priv().y();
-    return m_mapper->linesRead(outIdx);
+    return m_mapper->linesReadAndNextWindow(outIdx, lpi);
 }
 
 void cv::gimpl::FluidResizeAgent::setRatio(double ratio)
@@ -437,7 +440,11 @@ void cv::gimpl::FluidAgent::doWork()
 
     for (auto& in_view : in_views)
     {
-        if (in_view) in_view.priv().readDone(linesRead(), nextWindow());
+        if (in_view)
+        {
+            auto pair = linesReadAndnextWindow();
+            in_view.priv().readDone(pair.first, pair.second);
+        };
     }
 
     for (auto out_buf : out_buffers)
@@ -1030,6 +1037,10 @@ void cv::gimpl::GFluidExecutable::makeReshape(const std::vector<gapi::own::Rect>
             GAPI_LOG_INFO(NULL, stream.str());
         }
     }
+
+    // FIXME: calculate the size (lpi * ..)
+    m_script.clear();
+    m_script.reserve(10000);
 }
 
 void cv::gimpl::GFluidExecutable::reshape(ade::Graph &g, const GCompileArgs &args)
@@ -1134,24 +1145,36 @@ void cv::gimpl::GFluidExecutable::run(std::vector<InObj>  &&input_objs,
     //       and output buffers get "writeDone()"
     //   - if there's not enough data, Agent is skipped
     // Yes, THAT easy!
-    bool complete = true;
-    do {
-        complete = true;
-        bool work_done=false;
-        for (auto &agent : m_agents)
-        {
-            // agent->debug(std::cout);
-            if (!agent->done())
+
+    if (m_script.empty())
+    {
+        bool complete = true;
+        do {
+            complete = true;
+            bool work_done=false;
+            for (auto &agent : m_agents)
             {
-                if (agent->canWork())
+                // agent->debug(std::cout);
+                if (!agent->done())
                 {
-                    agent->doWork(); work_done=true;
+                    if (agent->canWork())
+                    {
+                        agent->doWork(); work_done=true;
+                        m_script.push_back(agent.get());
+                    }
+                    if (!agent->done())   complete = false;
                 }
-                if (!agent->done())   complete = false;
             }
+            GAPI_Assert(work_done || complete);
+        } while (!complete); // FIXME: number of iterations can be calculated statically
+    }
+    else
+    {
+        for (auto &agent : m_script)
+        {
+            agent->doWork();
         }
-        GAPI_Assert(work_done || complete);
-    } while (!complete); // FIXME: number of iterations can be calculated statically
+    }
 }
 
 // FIXME: these passes operate on graph global level!!!