diff --git a/doc/tutorials/imgproc/anisotropic_image_segmentation/anisotropic_image_segmentation.markdown b/doc/tutorials/imgproc/anisotropic_image_segmentation/anisotropic_image_segmentation.markdown index 489cd6bf3a..49fd621909 100644 --- a/doc/tutorials/imgproc/anisotropic_image_segmentation/anisotropic_image_segmentation.markdown +++ b/doc/tutorials/imgproc/anisotropic_image_segmentation/anisotropic_image_segmentation.markdown @@ -37,7 +37,7 @@ J_{12} & J_{22} where \f$J_{11} = M[Z_{x}^{2}]\f$, \f$J_{22} = M[Z_{y}^{2}]\f$, \f$J_{12} = M[Z_{x}Z_{y}]\f$ - components of the tensor, \f$M[]\f$ is a symbol of mathematical expectation (we can consider this operation as averaging in a window w), \f$Z_{x}\f$ and \f$Z_{y}\f$ are partial derivatives of an image \f$Z\f$ with respect to \f$x\f$ and \f$y\f$. The eigenvalues of the tensor can be found in the below formula: -\f[\lambda_{1,2} = J_{11} + J_{22} \pm \sqrt{(J_{11} - J_{22})^{2} + 4J_{12}^{2}}\f] +\f[\lambda_{1,2} = \frac{1}{2} \left [ J_{11} + J_{22} \pm \sqrt{(J_{11} - J_{22})^{2} + 4J_{12}^{2}} \right ] \f] where \f$\lambda_1\f$ - largest eigenvalue, \f$\lambda_2\f$ - smallest eigenvalue. ### How to estimate orientation and coherency of an anisotropic image by gradient structure tensor? diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp index 2db28667dd..c1bbd3316c 100644 --- a/modules/core/src/ocl.cpp +++ b/modules/core/src/ocl.cpp @@ -40,6 +40,11 @@ //M*/ #include "precomp.hpp" + +#ifndef HAVE_OPENCL +#include "ocl_disabled.impl.hpp" +#else // HAVE_OPENCL + #include #include #include @@ -106,23 +111,7 @@ #include "opencv2/core/opencl/runtime/opencl_clamdblas.hpp" #include "opencv2/core/opencl/runtime/opencl_clamdfft.hpp" -#ifdef HAVE_OPENCL #include "opencv2/core/opencl/runtime/opencl_core.hpp" -#else -#if defined(_MSC_VER) - #pragma warning(push) - #pragma warning(disable : 4100) - #pragma warning(disable : 4702) -#elif defined(__clang__) - #pragma clang diagnostic push - #pragma clang diagnostic ignored "-Wunused-parameter" -#elif defined(__GNUC__) - #pragma GCC diagnostic push - #pragma GCC diagnostic ignored "-Wunused-parameter" -#endif -// TODO FIXIT: This file can't be build without OPENCL -#include "ocl_deprecated.hpp" -#endif // HAVE_OPENCL #ifdef HAVE_OPENCL_SVM #include "opencv2/core/opencl/runtime/opencl_svm_20.hpp" @@ -147,31 +136,6 @@ cv::utils::AllocatorStatisticsInterface& getOpenCLAllocatorStatistics() return opencl_allocator_stats; } -#ifndef HAVE_OPENCL -#define CV_OPENCL_NO_SUPPORT() CV_Error(cv::Error::OpenCLApiCallError, "OpenCV build without OpenCL support") -namespace { -struct DummyImpl -{ - DummyImpl() { CV_OPENCL_NO_SUPPORT(); } - ~DummyImpl() { /* do not throw in desctructors */ } - IMPLEMENT_REFCOUNTABLE(); -}; -} // namespace - -// TODO Replace to empty body (without HAVE_OPENCL) -#define CV_OCL_TRACE_CHECK_RESULT(status, message) /* nothing */ -#define CV_OCL_API_ERROR_MSG(check_result, msg) cv::String() -#define CV_OCL_CHECK_RESULT(check_result, msg) (void)check_result -#define CV_OCL_CHECK_(expr, check_result) expr; (void)check_result -#define CV_OCL_CHECK(expr) do { cl_int __cl_result = (expr); CV_OCL_CHECK_RESULT(__cl_result, #expr); } while (0) -#define CV_OCL_DBG_CHECK_RESULT(check_result, msg) (void)check_result -#define CV_OCL_DBG_CHECK_(expr, check_result) expr; (void)check_result -#define CV_OCL_DBG_CHECK(expr) do { cl_int __cl_result = (expr); CV_OCL_CHECK_RESULT(__cl_result, #expr); } while (0) - -static const bool CV_OPENCL_DISABLE_BUFFER_RECT_OPERATIONS = false; - -#else // HAVE_OPENCL - #ifndef _DEBUG static bool isRaiseError() { @@ -270,7 +234,6 @@ static const String getBuildExtraOptions() static const bool CV_OPENCL_ENABLE_MEM_USE_HOST_PTR = utils::getConfigurationParameterBool("OPENCV_OPENCL_ENABLE_MEM_USE_HOST_PTR", true); static const size_t CV_OPENCL_ALIGNMENT_MEM_USE_HOST_PTR = utils::getConfigurationParameterSizeT("OPENCV_OPENCL_ALIGNMENT_MEM_USE_HOST_PTR", 4); -#endif // HAVE_OPENCL struct UMat2D { @@ -331,7 +294,7 @@ static uint64 crc64( const uchar* data, size_t size, uint64 crc0=0 ) return ~crc; } -#if defined HAVE_OPENCL && OPENCV_HAVE_FILESYSTEM_SUPPORT +#if OPENCV_HAVE_FILESYSTEM_SUPPORT struct OpenCLBinaryCacheConfigurator { cv::String cache_path_; @@ -872,7 +835,6 @@ static bool g_isOpenCVActivated = false; bool haveOpenCL() { CV_TRACE_FUNCTION(); -#ifdef HAVE_OPENCL static bool g_isOpenCLInitialized = false; static bool g_isOpenCLAvailable = false; @@ -902,9 +864,6 @@ bool haveOpenCL() g_isOpenCLInitialized = true; } return g_isOpenCLAvailable; -#else - return false; -#endif } bool useOpenCL() @@ -924,14 +883,12 @@ bool useOpenCL() return data.useOpenCL > 0; } -#ifdef HAVE_OPENCL bool isOpenCLActivated() { if (!g_isOpenCVActivated) return false; // prevent unnecessary OpenCL activation via useOpenCL()->haveOpenCL() calls return useOpenCL(); } -#endif void setUseOpenCL(bool flag) { @@ -1953,7 +1910,6 @@ static unsigned int getSVMCapabilitiesMask() } // namespace #endif -#ifdef HAVE_OPENCL static size_t getProgramCountLimit() { static bool initialized = false; @@ -1965,7 +1921,6 @@ static size_t getProgramCountLimit() } return count; } -#endif struct Context::Impl { @@ -3548,8 +3503,6 @@ internal::ProgramEntry::operator ProgramSource&() const /////////////////////////////////////////// Program ///////////////////////////////////////////// -#ifdef HAVE_OPENCL - static cv::String joinBuildOptions(const cv::String& a, const cv::String& b) { @@ -3963,10 +3916,6 @@ struct Program::Impl String sourceName_; }; -#else // HAVE_OPENCL -struct Program::Impl : public DummyImpl {}; -#endif // HAVE_OPENCL - Program::Program() { p = 0; } @@ -4009,7 +3958,6 @@ bool Program::create(const ProgramSource& src, p->release(); p = NULL; } -#ifdef HAVE_OPENCL p = new Impl(src, buildflags, errmsg); if(!p->handle) { @@ -4017,18 +3965,11 @@ bool Program::create(const ProgramSource& src, p = 0; } return p != 0; -#else - CV_OPENCL_NO_SUPPORT(); -#endif } void* Program::ptr() const { -#ifdef HAVE_OPENCL return p ? p->handle : 0; -#else - CV_OPENCL_NO_SUPPORT(); -#endif } #ifndef OPENCV_REMOVE_DEPRECATED_API @@ -4051,44 +3992,30 @@ bool Program::write(String& bin) const String Program::getPrefix() const { -#ifdef HAVE_OPENCL if(!p) return String(); Context::Impl* ctx_ = Context::getDefault().getImpl(); CV_Assert(ctx_); return cv::format("opencl=%s\nbuildflags=%s", ctx_->getPrefixString().c_str(), p->buildflags.c_str()); -#else - CV_OPENCL_NO_SUPPORT(); -#endif } String Program::getPrefix(const String& buildflags) { -#ifdef HAVE_OPENCL Context::Impl* ctx_ = Context::getDefault().getImpl(); CV_Assert(ctx_); return cv::format("opencl=%s\nbuildflags=%s", ctx_->getPrefixString().c_str(), buildflags.c_str()); -#else - CV_OPENCL_NO_SUPPORT(); -#endif } -#endif +#endif // OPENCV_REMOVE_DEPRECATED_API void Program::getBinary(std::vector& binary) const { -#ifdef HAVE_OPENCL CV_Assert(p && "Empty program"); p->getProgramBinary(binary); -#else - binary.clear(); - CV_OPENCL_NO_SUPPORT(); -#endif } Program Context::Impl::getProg(const ProgramSource& src, const String& buildflags, String& errmsg) { -#ifdef HAVE_OPENCL size_t limit = getProgramCountLimit(); const ProgramSource::Impl* src_ = src.getImpl(); CV_Assert(src_); @@ -4140,9 +4067,6 @@ Program Context::Impl::getProg(const ProgramSource& src, cacheList.push_front(key); } return prog; -#else - CV_OPENCL_NO_SUPPORT(); -#endif } @@ -4703,9 +4627,6 @@ public: bool allocate(UMatData* u, AccessFlag accessFlags, UMatUsageFlags usageFlags) const CV_OVERRIDE { -#ifndef HAVE_OPENCL - return false; -#else if(!u) return false; @@ -4825,7 +4746,6 @@ public: u->markHostCopyObsolete(true); opencl_allocator_stats.onAllocate(u->size); return true; -#endif // HAVE_OPENCL } /*void sync(UMatData* u) const @@ -6705,27 +6625,19 @@ struct Timer::Impl void start() { -#ifdef HAVE_OPENCL CV_OCL_DBG_CHECK(clFinish((cl_command_queue)queue.ptr())); timer.start(); -#endif } void stop() { -#ifdef HAVE_OPENCL CV_OCL_DBG_CHECK(clFinish((cl_command_queue)queue.ptr())); timer.stop(); -#endif } uint64 durationNS() const { -#ifdef HAVE_OPENCL return (uint64)(timer.getTimeSec() * 1e9); -#else - return 0; -#endif } TickMeter timer; @@ -6752,13 +6664,6 @@ uint64 Timer::durationNS() const return p->durationNS(); } -#ifndef HAVE_OPENCL -#if defined(_MSC_VER) - #pragma warning(pop) -#elif defined(__clang__) - #pragma clang diagnostic pop -#elif defined(__GNUC__) - #pragma GCC diagnostic pop -#endif -#endif }} // namespace + +#endif // HAVE_OPENCL diff --git a/modules/core/src/ocl_deprecated.hpp b/modules/core/src/ocl_deprecated.hpp deleted file mode 100644 index 753e8c312b..0000000000 --- a/modules/core/src/ocl_deprecated.hpp +++ /dev/null @@ -1,1224 +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. - -/* - Part of the file is an extract from the standard OpenCL headers from Khronos site. - Below is the original copyright. -*/ - -/******************************************************************************* - * Copyright (c) 2008 - 2012 The Khronos Group Inc. - * - * Permission is hereby granted, free of charge, to any person obtaining a - * copy of this software and/or associated documentation files (the - * "Materials"), to deal in the Materials without restriction, including - * without limitation the rights to use, copy, modify, merge, publish, - * distribute, sublicense, and/or sell copies of the Materials, and to - * permit persons to whom the Materials are furnished to do so, subject to - * the following conditions: - * - * The above copyright notice and this permission notice shall be included - * in all copies or substantial portions of the Materials. - * - * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, - * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF - * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. - * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY - * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, - * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE - * MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS. - ******************************************************************************/ - -#if 0 //defined __APPLE__ -#define HAVE_OPENCL 1 -#else -#undef HAVE_OPENCL -#endif - -#define OPENCV_CL_NOT_IMPLEMENTED -1000 - -#ifdef HAVE_OPENCL - -#if defined __APPLE__ -#include -#else -#include -#endif - -static const bool g_haveOpenCL = true; - -#else - -extern "C" { - -struct _cl_platform_id { int dummy; }; -struct _cl_device_id { int dummy; }; -struct _cl_context { int dummy; }; -struct _cl_command_queue { int dummy; }; -struct _cl_mem { int dummy; }; -struct _cl_program { int dummy; }; -struct _cl_kernel { int dummy; }; -struct _cl_event { int dummy; }; -struct _cl_sampler { int dummy; }; - -typedef struct _cl_platform_id * cl_platform_id; -typedef struct _cl_device_id * cl_device_id; -typedef struct _cl_context * cl_context; -typedef struct _cl_command_queue * cl_command_queue; -typedef struct _cl_mem * cl_mem; -typedef struct _cl_program * cl_program; -typedef struct _cl_kernel * cl_kernel; -typedef struct _cl_event * cl_event; -typedef struct _cl_sampler * cl_sampler; - -typedef int cl_int; -typedef unsigned cl_uint; -#if defined (_WIN32) && defined(_MSC_VER) - typedef __int64 cl_long; - typedef unsigned __int64 cl_ulong; -#else - typedef long cl_long; - typedef unsigned long cl_ulong; -#endif - -typedef cl_uint cl_bool; /* WARNING! Unlike cl_ types in cl_platform.h, cl_bool is not guaranteed to be the same size as the bool in kernels. */ -typedef cl_ulong cl_bitfield; -typedef cl_bitfield cl_device_type; -typedef cl_uint cl_platform_info; -typedef cl_uint cl_device_info; -typedef cl_bitfield cl_device_fp_config; -typedef cl_uint cl_device_mem_cache_type; -typedef cl_uint cl_device_local_mem_type; -typedef cl_bitfield cl_device_exec_capabilities; -typedef cl_bitfield cl_command_queue_properties; -typedef intptr_t cl_device_partition_property; -typedef cl_bitfield cl_device_affinity_domain; - -typedef intptr_t cl_context_properties; -typedef cl_uint cl_context_info; -typedef cl_uint cl_command_queue_info; -typedef cl_uint cl_channel_order; -typedef cl_uint cl_channel_type; -typedef cl_bitfield cl_mem_flags; -typedef cl_uint cl_mem_object_type; -typedef cl_uint cl_mem_info; -typedef cl_bitfield cl_mem_migration_flags; -typedef cl_uint cl_image_info; -typedef cl_uint cl_buffer_create_type; -typedef cl_uint cl_addressing_mode; -typedef cl_uint cl_filter_mode; -typedef cl_uint cl_sampler_info; -typedef cl_bitfield cl_map_flags; -typedef cl_uint cl_program_info; -typedef cl_uint cl_program_build_info; -typedef cl_uint cl_program_binary_type; -typedef cl_int cl_build_status; -typedef cl_uint cl_kernel_info; -typedef cl_uint cl_kernel_arg_info; -typedef cl_uint cl_kernel_arg_address_qualifier; -typedef cl_uint cl_kernel_arg_access_qualifier; -typedef cl_bitfield cl_kernel_arg_type_qualifier; -typedef cl_uint cl_kernel_work_group_info; -typedef cl_uint cl_event_info; -typedef cl_uint cl_command_type; -typedef cl_uint cl_profiling_info; - - -typedef struct _cl_image_format { - cl_channel_order image_channel_order; - cl_channel_type image_channel_data_type; -} cl_image_format; - -typedef struct _cl_image_desc { - cl_mem_object_type image_type; - size_t image_width; - size_t image_height; - size_t image_depth; - size_t image_array_size; - size_t image_row_pitch; - size_t image_slice_pitch; - cl_uint num_mip_levels; - cl_uint num_samples; - cl_mem buffer; -} cl_image_desc; - -typedef struct _cl_buffer_region { - size_t origin; - size_t size; -} cl_buffer_region; - - -////////////////////////////////////////////////////////// - -#define CL_SUCCESS 0 -#define CL_DEVICE_NOT_FOUND -1 -#define CL_DEVICE_NOT_AVAILABLE -2 -#define CL_COMPILER_NOT_AVAILABLE -3 -#define CL_MEM_OBJECT_ALLOCATION_FAILURE -4 -#define CL_OUT_OF_RESOURCES -5 -#define CL_OUT_OF_HOST_MEMORY -6 -#define CL_PROFILING_INFO_NOT_AVAILABLE -7 -#define CL_MEM_COPY_OVERLAP -8 -#define CL_IMAGE_FORMAT_MISMATCH -9 -#define CL_IMAGE_FORMAT_NOT_SUPPORTED -10 -#define CL_BUILD_PROGRAM_FAILURE -11 -#define CL_MAP_FAILURE -12 -#define CL_MISALIGNED_SUB_BUFFER_OFFSET -13 -#define CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST -14 -#define CL_COMPILE_PROGRAM_FAILURE -15 -#define CL_LINKER_NOT_AVAILABLE -16 -#define CL_LINK_PROGRAM_FAILURE -17 -#define CL_DEVICE_PARTITION_FAILED -18 -#define CL_KERNEL_ARG_INFO_NOT_AVAILABLE -19 - -#define CL_INVALID_VALUE -30 -#define CL_INVALID_DEVICE_TYPE -31 -#define CL_INVALID_PLATFORM -32 -#define CL_INVALID_DEVICE -33 -#define CL_INVALID_CONTEXT -34 -#define CL_INVALID_QUEUE_PROPERTIES -35 -#define CL_INVALID_COMMAND_QUEUE -36 -#define CL_INVALID_HOST_PTR -37 -#define CL_INVALID_MEM_OBJECT -38 -#define CL_INVALID_IMAGE_FORMAT_DESCRIPTOR -39 -#define CL_INVALID_IMAGE_SIZE -40 -#define CL_INVALID_SAMPLER -41 -#define CL_INVALID_BINARY -42 -#define CL_INVALID_BUILD_OPTIONS -43 -#define CL_INVALID_PROGRAM -44 -#define CL_INVALID_PROGRAM_EXECUTABLE -45 -#define CL_INVALID_KERNEL_NAME -46 -#define CL_INVALID_KERNEL_DEFINITION -47 -#define CL_INVALID_KERNEL -48 -#define CL_INVALID_ARG_INDEX -49 -#define CL_INVALID_ARG_VALUE -50 -#define CL_INVALID_ARG_SIZE -51 -#define CL_INVALID_KERNEL_ARGS -52 -#define CL_INVALID_WORK_DIMENSION -53 -#define CL_INVALID_WORK_GROUP_SIZE -54 -#define CL_INVALID_WORK_ITEM_SIZE -55 -#define CL_INVALID_GLOBAL_OFFSET -56 -#define CL_INVALID_EVENT_WAIT_LIST -57 -#define CL_INVALID_EVENT -58 -#define CL_INVALID_OPERATION -59 -#define CL_INVALID_GL_OBJECT -60 -#define CL_INVALID_BUFFER_SIZE -61 -#define CL_INVALID_MIP_LEVEL -62 -#define CL_INVALID_GLOBAL_WORK_SIZE -63 -#define CL_INVALID_PROPERTY -64 -#define CL_INVALID_IMAGE_DESCRIPTOR -65 -#define CL_INVALID_COMPILER_OPTIONS -66 -#define CL_INVALID_LINKER_OPTIONS -67 -#define CL_INVALID_DEVICE_PARTITION_COUNT -68 - -/*#define CL_VERSION_1_0 1 -#define CL_VERSION_1_1 1 -#define CL_VERSION_1_2 1*/ - -#define CL_FALSE 0 -#define CL_TRUE 1 -#define CL_BLOCKING CL_TRUE -#define CL_NON_BLOCKING CL_FALSE - -#define CL_PLATFORM_PROFILE 0x0900 -#define CL_PLATFORM_VERSION 0x0901 -#define CL_PLATFORM_NAME 0x0902 -#define CL_PLATFORM_VENDOR 0x0903 -#define CL_PLATFORM_EXTENSIONS 0x0904 - -#define CL_DEVICE_TYPE_DEFAULT (1 << 0) -#define CL_DEVICE_TYPE_CPU (1 << 1) -#define CL_DEVICE_TYPE_GPU (1 << 2) -#define CL_DEVICE_TYPE_ACCELERATOR (1 << 3) -#define CL_DEVICE_TYPE_CUSTOM (1 << 4) -#define CL_DEVICE_TYPE_ALL 0xFFFFFFFF -#define CL_DEVICE_TYPE 0x1000 -#define CL_DEVICE_VENDOR_ID 0x1001 -#define CL_DEVICE_MAX_COMPUTE_UNITS 0x1002 -#define CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS 0x1003 -#define CL_DEVICE_MAX_WORK_GROUP_SIZE 0x1004 -#define CL_DEVICE_MAX_WORK_ITEM_SIZES 0x1005 -#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR 0x1006 -#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT 0x1007 -#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT 0x1008 -#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG 0x1009 -#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT 0x100A -#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE 0x100B -#define CL_DEVICE_MAX_CLOCK_FREQUENCY 0x100C -#define CL_DEVICE_ADDRESS_BITS 0x100D -#define CL_DEVICE_MAX_READ_IMAGE_ARGS 0x100E -#define CL_DEVICE_MAX_WRITE_IMAGE_ARGS 0x100F -#define CL_DEVICE_MAX_MEM_ALLOC_SIZE 0x1010 -#define CL_DEVICE_IMAGE2D_MAX_WIDTH 0x1011 -#define CL_DEVICE_IMAGE2D_MAX_HEIGHT 0x1012 -#define CL_DEVICE_IMAGE3D_MAX_WIDTH 0x1013 -#define CL_DEVICE_IMAGE3D_MAX_HEIGHT 0x1014 -#define CL_DEVICE_IMAGE3D_MAX_DEPTH 0x1015 -#define CL_DEVICE_IMAGE_SUPPORT 0x1016 -#define CL_DEVICE_MAX_PARAMETER_SIZE 0x1017 -#define CL_DEVICE_MAX_SAMPLERS 0x1018 -#define CL_DEVICE_MEM_BASE_ADDR_ALIGN 0x1019 -#define CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE 0x101A -#define CL_DEVICE_SINGLE_FP_CONFIG 0x101B -#define CL_DEVICE_GLOBAL_MEM_CACHE_TYPE 0x101C -#define CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE 0x101D -#define CL_DEVICE_GLOBAL_MEM_CACHE_SIZE 0x101E -#define CL_DEVICE_GLOBAL_MEM_SIZE 0x101F -#define CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE 0x1020 -#define CL_DEVICE_MAX_CONSTANT_ARGS 0x1021 -#define CL_DEVICE_LOCAL_MEM_TYPE 0x1022 -#define CL_DEVICE_LOCAL_MEM_SIZE 0x1023 -#define CL_DEVICE_ERROR_CORRECTION_SUPPORT 0x1024 -#define CL_DEVICE_PROFILING_TIMER_RESOLUTION 0x1025 -#define CL_DEVICE_ENDIAN_LITTLE 0x1026 -#define CL_DEVICE_AVAILABLE 0x1027 -#define CL_DEVICE_COMPILER_AVAILABLE 0x1028 -#define CL_DEVICE_EXECUTION_CAPABILITIES 0x1029 -#define CL_DEVICE_QUEUE_PROPERTIES 0x102A -#define CL_DEVICE_NAME 0x102B -#define CL_DEVICE_VENDOR 0x102C -#define CL_DRIVER_VERSION 0x102D -#define CL_DEVICE_PROFILE 0x102E -#define CL_DEVICE_VERSION 0x102F -#define CL_DEVICE_EXTENSIONS 0x1030 -#define CL_DEVICE_PLATFORM 0x1031 -#define CL_DEVICE_DOUBLE_FP_CONFIG 0x1032 -#define CL_DEVICE_HALF_FP_CONFIG 0x1033 -#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF 0x1034 -#define CL_DEVICE_HOST_UNIFIED_MEMORY 0x1035 -#define CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR 0x1036 -#define CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT 0x1037 -#define CL_DEVICE_NATIVE_VECTOR_WIDTH_INT 0x1038 -#define CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG 0x1039 -#define CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT 0x103A -#define CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE 0x103B -#define CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF 0x103C -#define CL_DEVICE_OPENCL_C_VERSION 0x103D -#define CL_DEVICE_LINKER_AVAILABLE 0x103E -#define CL_DEVICE_BUILT_IN_KERNELS 0x103F -#define CL_DEVICE_IMAGE_MAX_BUFFER_SIZE 0x1040 -#define CL_DEVICE_IMAGE_MAX_ARRAY_SIZE 0x1041 -#define CL_DEVICE_PARENT_DEVICE 0x1042 -#define CL_DEVICE_PARTITION_MAX_SUB_DEVICES 0x1043 -#define CL_DEVICE_PARTITION_PROPERTIES 0x1044 -#define CL_DEVICE_PARTITION_AFFINITY_DOMAIN 0x1045 -#define CL_DEVICE_PARTITION_TYPE 0x1046 -#define CL_DEVICE_REFERENCE_COUNT 0x1047 -#define CL_DEVICE_PREFERRED_INTEROP_USER_SYNC 0x1048 -#define CL_DEVICE_PRINTF_BUFFER_SIZE 0x1049 -#define CL_DEVICE_IMAGE_PITCH_ALIGNMENT 0x104A -#define CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT 0x104B - -#define CL_FP_DENORM (1 << 0) -#define CL_FP_INF_NAN (1 << 1) -#define CL_FP_ROUND_TO_NEAREST (1 << 2) -#define CL_FP_ROUND_TO_ZERO (1 << 3) -#define CL_FP_ROUND_TO_INF (1 << 4) -#define CL_FP_FMA (1 << 5) -#define CL_FP_SOFT_FLOAT (1 << 6) -#define CL_FP_CORRECTLY_ROUNDED_DIVIDE_SQRT (1 << 7) - -#define CL_NONE 0x0 -#define CL_READ_ONLY_CACHE 0x1 -#define CL_READ_WRITE_CACHE 0x2 -#define CL_LOCAL 0x1 -#define CL_GLOBAL 0x2 -#define CL_EXEC_KERNEL (1 << 0) -#define CL_EXEC_NATIVE_KERNEL (1 << 1) -#define CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE (1 << 0) -#define CL_QUEUE_PROFILING_ENABLE (1 << 1) - -#define CL_CONTEXT_REFERENCE_COUNT 0x1080 -#define CL_CONTEXT_DEVICES 0x1081 -#define CL_CONTEXT_PROPERTIES 0x1082 -#define CL_CONTEXT_NUM_DEVICES 0x1083 -#define CL_CONTEXT_PLATFORM 0x1084 -#define CL_CONTEXT_INTEROP_USER_SYNC 0x1085 - -#define CL_DEVICE_PARTITION_EQUALLY 0x1086 -#define CL_DEVICE_PARTITION_BY_COUNTS 0x1087 -#define CL_DEVICE_PARTITION_BY_COUNTS_LIST_END 0x0 -#define CL_DEVICE_PARTITION_BY_AFFINITY_DOMAIN 0x1088 -#define CL_DEVICE_AFFINITY_DOMAIN_NUMA (1 << 0) -#define CL_DEVICE_AFFINITY_DOMAIN_L4_CACHE (1 << 1) -#define CL_DEVICE_AFFINITY_DOMAIN_L3_CACHE (1 << 2) -#define CL_DEVICE_AFFINITY_DOMAIN_L2_CACHE (1 << 3) -#define CL_DEVICE_AFFINITY_DOMAIN_L1_CACHE (1 << 4) -#define CL_DEVICE_AFFINITY_DOMAIN_NEXT_PARTITIONABLE (1 << 5) -#define CL_QUEUE_CONTEXT 0x1090 -#define CL_QUEUE_DEVICE 0x1091 -#define CL_QUEUE_REFERENCE_COUNT 0x1092 -#define CL_QUEUE_PROPERTIES 0x1093 -#define CL_MEM_READ_WRITE (1 << 0) -#define CL_MEM_WRITE_ONLY (1 << 1) -#define CL_MEM_READ_ONLY (1 << 2) -#define CL_MEM_USE_HOST_PTR (1 << 3) -#define CL_MEM_ALLOC_HOST_PTR (1 << 4) -#define CL_MEM_COPY_HOST_PTR (1 << 5) -// reserved (1 << 6) -#define CL_MEM_HOST_WRITE_ONLY (1 << 7) -#define CL_MEM_HOST_READ_ONLY (1 << 8) -#define CL_MEM_HOST_NO_ACCESS (1 << 9) -#define CL_MIGRATE_MEM_OBJECT_HOST (1 << 0) -#define CL_MIGRATE_MEM_OBJECT_CONTENT_UNDEFINED (1 << 1) - -#define CL_R 0x10B0 -#define CL_A 0x10B1 -#define CL_RG 0x10B2 -#define CL_RA 0x10B3 -#define CL_RGB 0x10B4 -#define CL_RGBA 0x10B5 -#define CL_BGRA 0x10B6 -#define CL_ARGB 0x10B7 -#define CL_INTENSITY 0x10B8 -#define CL_LUMINANCE 0x10B9 -#define CL_Rx 0x10BA -#define CL_RGx 0x10BB -#define CL_RGBx 0x10BC -#define CL_DEPTH 0x10BD -#define CL_DEPTH_STENCIL 0x10BE - -#define CL_SNORM_INT8 0x10D0 -#define CL_SNORM_INT16 0x10D1 -#define CL_UNORM_INT8 0x10D2 -#define CL_UNORM_INT16 0x10D3 -#define CL_UNORM_SHORT_565 0x10D4 -#define CL_UNORM_SHORT_555 0x10D5 -#define CL_UNORM_INT_101010 0x10D6 -#define CL_SIGNED_INT8 0x10D7 -#define CL_SIGNED_INT16 0x10D8 -#define CL_SIGNED_INT32 0x10D9 -#define CL_UNSIGNED_INT8 0x10DA -#define CL_UNSIGNED_INT16 0x10DB -#define CL_UNSIGNED_INT32 0x10DC -#define CL_HALF_FLOAT 0x10DD -#define CL_FLOAT 0x10DE -#define CL_UNORM_INT24 0x10DF - -#define CL_MEM_OBJECT_BUFFER 0x10F0 -#define CL_MEM_OBJECT_IMAGE2D 0x10F1 -#define CL_MEM_OBJECT_IMAGE3D 0x10F2 -#define CL_MEM_OBJECT_IMAGE2D_ARRAY 0x10F3 -#define CL_MEM_OBJECT_IMAGE1D 0x10F4 -#define CL_MEM_OBJECT_IMAGE1D_ARRAY 0x10F5 -#define CL_MEM_OBJECT_IMAGE1D_BUFFER 0x10F6 - -#define CL_MEM_TYPE 0x1100 -#define CL_MEM_FLAGS 0x1101 -#define CL_MEM_SIZE 0x1102 -#define CL_MEM_HOST_PTR 0x1103 -#define CL_MEM_MAP_COUNT 0x1104 -#define CL_MEM_REFERENCE_COUNT 0x1105 -#define CL_MEM_CONTEXT 0x1106 -#define CL_MEM_ASSOCIATED_MEMOBJECT 0x1107 -#define CL_MEM_OFFSET 0x1108 - -#define CL_IMAGE_FORMAT 0x1110 -#define CL_IMAGE_ELEMENT_SIZE 0x1111 -#define CL_IMAGE_ROW_PITCH 0x1112 -#define CL_IMAGE_SLICE_PITCH 0x1113 -#define CL_IMAGE_WIDTH 0x1114 -#define CL_IMAGE_HEIGHT 0x1115 -#define CL_IMAGE_DEPTH 0x1116 -#define CL_IMAGE_ARRAY_SIZE 0x1117 -#define CL_IMAGE_BUFFER 0x1118 -#define CL_IMAGE_NUM_MIP_LEVELS 0x1119 -#define CL_IMAGE_NUM_SAMPLES 0x111A - -#define CL_ADDRESS_NONE 0x1130 -#define CL_ADDRESS_CLAMP_TO_EDGE 0x1131 -#define CL_ADDRESS_CLAMP 0x1132 -#define CL_ADDRESS_REPEAT 0x1133 -#define CL_ADDRESS_MIRRORED_REPEAT 0x1134 - -#define CL_FILTER_NEAREST 0x1140 -#define CL_FILTER_LINEAR 0x1141 - -#define CL_SAMPLER_REFERENCE_COUNT 0x1150 -#define CL_SAMPLER_CONTEXT 0x1151 -#define CL_SAMPLER_NORMALIZED_COORDS 0x1152 -#define CL_SAMPLER_ADDRESSING_MODE 0x1153 -#define CL_SAMPLER_FILTER_MODE 0x1154 - -#define CL_MAP_READ (1 << 0) -#define CL_MAP_WRITE (1 << 1) -#define CL_MAP_WRITE_INVALIDATE_REGION (1 << 2) - -#define CL_PROGRAM_REFERENCE_COUNT 0x1160 -#define CL_PROGRAM_CONTEXT 0x1161 -#define CL_PROGRAM_NUM_DEVICES 0x1162 -#define CL_PROGRAM_DEVICES 0x1163 -#define CL_PROGRAM_SOURCE 0x1164 -#define CL_PROGRAM_BINARY_SIZES 0x1165 -#define CL_PROGRAM_BINARIES 0x1166 -#define CL_PROGRAM_NUM_KERNELS 0x1167 -#define CL_PROGRAM_KERNEL_NAMES 0x1168 -#define CL_PROGRAM_BUILD_STATUS 0x1181 -#define CL_PROGRAM_BUILD_OPTIONS 0x1182 -#define CL_PROGRAM_BUILD_LOG 0x1183 -#define CL_PROGRAM_BINARY_TYPE 0x1184 -#define CL_PROGRAM_BINARY_TYPE_NONE 0x0 -#define CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT 0x1 -#define CL_PROGRAM_BINARY_TYPE_LIBRARY 0x2 -#define CL_PROGRAM_BINARY_TYPE_EXECUTABLE 0x4 - -#define CL_BUILD_SUCCESS 0 -#define CL_BUILD_NONE -1 -#define CL_BUILD_ERROR -2 -#define CL_BUILD_IN_PROGRESS -3 - -#define CL_KERNEL_FUNCTION_NAME 0x1190 -#define CL_KERNEL_NUM_ARGS 0x1191 -#define CL_KERNEL_REFERENCE_COUNT 0x1192 -#define CL_KERNEL_CONTEXT 0x1193 -#define CL_KERNEL_PROGRAM 0x1194 -#define CL_KERNEL_ATTRIBUTES 0x1195 -#define CL_KERNEL_ARG_ADDRESS_QUALIFIER 0x1196 -#define CL_KERNEL_ARG_ACCESS_QUALIFIER 0x1197 -#define CL_KERNEL_ARG_TYPE_NAME 0x1198 -#define CL_KERNEL_ARG_TYPE_QUALIFIER 0x1199 -#define CL_KERNEL_ARG_NAME 0x119A -#define CL_KERNEL_ARG_ADDRESS_GLOBAL 0x119B -#define CL_KERNEL_ARG_ADDRESS_LOCAL 0x119C -#define CL_KERNEL_ARG_ADDRESS_CONSTANT 0x119D -#define CL_KERNEL_ARG_ADDRESS_PRIVATE 0x119E -#define CL_KERNEL_ARG_ACCESS_READ_ONLY 0x11A0 -#define CL_KERNEL_ARG_ACCESS_WRITE_ONLY 0x11A1 -#define CL_KERNEL_ARG_ACCESS_READ_WRITE 0x11A2 -#define CL_KERNEL_ARG_ACCESS_NONE 0x11A3 -#define CL_KERNEL_ARG_TYPE_NONE 0 -#define CL_KERNEL_ARG_TYPE_CONST (1 << 0) -#define CL_KERNEL_ARG_TYPE_RESTRICT (1 << 1) -#define CL_KERNEL_ARG_TYPE_VOLATILE (1 << 2) -#define CL_KERNEL_WORK_GROUP_SIZE 0x11B0 -#define CL_KERNEL_COMPILE_WORK_GROUP_SIZE 0x11B1 -#define CL_KERNEL_LOCAL_MEM_SIZE 0x11B2 -#define CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE 0x11B3 -#define CL_KERNEL_PRIVATE_MEM_SIZE 0x11B4 -#define CL_KERNEL_GLOBAL_WORK_SIZE 0x11B5 - -#define CL_EVENT_COMMAND_QUEUE 0x11D0 -#define CL_EVENT_COMMAND_TYPE 0x11D1 -#define CL_EVENT_REFERENCE_COUNT 0x11D2 -#define CL_EVENT_COMMAND_EXECUTION_STATUS 0x11D3 -#define CL_EVENT_CONTEXT 0x11D4 - -#define CL_COMMAND_NDRANGE_KERNEL 0x11F0 -#define CL_COMMAND_TASK 0x11F1 -#define CL_COMMAND_NATIVE_KERNEL 0x11F2 -#define CL_COMMAND_READ_BUFFER 0x11F3 -#define CL_COMMAND_WRITE_BUFFER 0x11F4 -#define CL_COMMAND_COPY_BUFFER 0x11F5 -#define CL_COMMAND_READ_IMAGE 0x11F6 -#define CL_COMMAND_WRITE_IMAGE 0x11F7 -#define CL_COMMAND_COPY_IMAGE 0x11F8 -#define CL_COMMAND_COPY_IMAGE_TO_BUFFER 0x11F9 -#define CL_COMMAND_COPY_BUFFER_TO_IMAGE 0x11FA -#define CL_COMMAND_MAP_BUFFER 0x11FB -#define CL_COMMAND_MAP_IMAGE 0x11FC -#define CL_COMMAND_UNMAP_MEM_OBJECT 0x11FD -#define CL_COMMAND_MARKER 0x11FE -#define CL_COMMAND_ACQUIRE_GL_OBJECTS 0x11FF -#define CL_COMMAND_RELEASE_GL_OBJECTS 0x1200 -#define CL_COMMAND_READ_BUFFER_RECT 0x1201 -#define CL_COMMAND_WRITE_BUFFER_RECT 0x1202 -#define CL_COMMAND_COPY_BUFFER_RECT 0x1203 -#define CL_COMMAND_USER 0x1204 -#define CL_COMMAND_BARRIER 0x1205 -#define CL_COMMAND_MIGRATE_MEM_OBJECTS 0x1206 -#define CL_COMMAND_FILL_BUFFER 0x1207 -#define CL_COMMAND_FILL_IMAGE 0x1208 - -#define CL_COMPLETE 0x0 -#define CL_RUNNING 0x1 -#define CL_SUBMITTED 0x2 -#define CL_QUEUED 0x3 -#define CL_BUFFER_CREATE_TYPE_REGION 0x1220 - -#define CL_PROFILING_COMMAND_QUEUED 0x1280 -#define CL_PROFILING_COMMAND_SUBMIT 0x1281 -#define CL_PROFILING_COMMAND_START 0x1282 -#define CL_PROFILING_COMMAND_END 0x1283 - -#define CL_CALLBACK CV_STDCALL - - -#ifdef HAVE_OPENCL -static const char* oclFuncToCheck = "clEnqueueReadBufferRect"; -static volatile bool g_haveOpenCL = false; -#endif - -#if defined(__APPLE__) && defined(HAVE_OPENCL) -#include - -static void* initOpenCLAndLoad(const char* funcname) -{ - static bool initialized = false; - static void* handle = 0; - if (!handle) - { - if(!initialized) - { - const char* oclpath = getenv("OPENCV_OPENCL_RUNTIME"); - oclpath = oclpath && strlen(oclpath) > 0 ? oclpath : - "/System/Library/Frameworks/OpenCL.framework/Versions/Current/OpenCL"; - handle = dlopen(oclpath, RTLD_LAZY); - initialized = true; - g_haveOpenCL = handle != 0 && dlsym(handle, oclFuncToCheck) != 0; - if( g_haveOpenCL ) - fprintf(stderr, "Successfully loaded OpenCL v1.1+ runtime from %s\n", oclpath); - else - fprintf(stderr, "Failed to load OpenCL runtime\n"); - } - if(!handle) - return 0; - } - - return funcname && handle ? dlsym(handle, funcname) : 0; -} - -#elif defined _WIN32 && defined(HAVE_OPENCL) - -#ifndef _WIN32_WINNT // This is needed for the declaration of TryEnterCriticalSection in winbase.h with Visual Studio 2005 (and older?) - #define _WIN32_WINNT 0x0400 // http://msdn.microsoft.com/en-us/library/ms686857(VS.85).aspx -#endif -#include -#if (_WIN32_WINNT >= 0x0602) - #include -#endif -#undef small -#undef min -#undef max -#undef abs - -static void* initOpenCLAndLoad(const char* funcname) -{ - static bool initialized = false; - static HMODULE handle = 0; - if (!handle) - { -#ifndef WINRT - if(!initialized) - { - handle = LoadLibraryA("OpenCL.dll"); - initialized = true; - g_haveOpenCL = handle != 0 && GetProcAddress(handle, oclFuncToCheck) != 0; - } -#endif - if(!handle) - return 0; - } - - return funcname ? (void*)GetProcAddress(handle, funcname) : 0; -} - -#elif defined(__linux) && defined(HAVE_OPENCL) - -#include -#include - -static void* initOpenCLAndLoad(const char* funcname) -{ - static bool initialized = false; - static void* handle = 0; - if (!handle) - { - if(!initialized) - { - handle = dlopen("libOpenCL.so", RTLD_LAZY); - if(!handle) - handle = dlopen("libCL.so", RTLD_LAZY); - initialized = true; - g_haveOpenCL = handle != 0 && dlsym(handle, oclFuncToCheck) != 0; - } - if(!handle) - return 0; - } - - return funcname ? (void*)dlsym(handle, funcname) : 0; -} - -#else - -static void* initOpenCLAndLoad(const char*) -{ - return 0; -} - -#endif - - -#define OCL_FUNC(rettype, funcname, argsdecl, args) \ - typedef rettype (CV_STDCALL * funcname##_t) argsdecl; \ - static rettype funcname argsdecl \ - { \ - static funcname##_t funcname##_p = 0; \ - if( !funcname##_p ) \ - { \ - funcname##_p = (funcname##_t)initOpenCLAndLoad(#funcname); \ - if( !funcname##_p ) \ - return OPENCV_CL_NOT_IMPLEMENTED; \ - } \ - return funcname##_p args; \ - } - - -#define OCL_FUNC_P(rettype, funcname, argsdecl, args) \ - typedef rettype (CV_STDCALL * funcname##_t) argsdecl; \ - static rettype funcname argsdecl \ - { \ - static funcname##_t funcname##_p = 0; \ - if( !funcname##_p ) \ - { \ - funcname##_p = (funcname##_t)initOpenCLAndLoad(#funcname); \ - if( !funcname##_p ) \ - { \ - if( errcode_ret ) \ - *errcode_ret = OPENCV_CL_NOT_IMPLEMENTED; \ - return 0; \ - } \ - } \ - return funcname##_p args; \ - } - -OCL_FUNC(cl_int, clGetPlatformIDs, - (cl_uint num_entries, cl_platform_id* platforms, cl_uint* num_platforms), - (num_entries, platforms, num_platforms)) - -OCL_FUNC(cl_int, clGetPlatformInfo, - (cl_platform_id platform, cl_platform_info param_name, - size_t param_value_size, void * param_value, - size_t * param_value_size_ret), - (platform, param_name, param_value_size, param_value, param_value_size_ret)) - -OCL_FUNC(cl_int, clGetDeviceInfo, - (cl_device_id device, - cl_device_info param_name, - size_t param_value_size, - void * param_value, - size_t * param_value_size_ret), - (device, param_name, param_value_size, param_value, param_value_size_ret)) - - -OCL_FUNC(cl_int, clGetDeviceIDs, - (cl_platform_id platform, - cl_device_type device_type, - cl_uint num_entries, - cl_device_id * devices, - cl_uint * num_devices), - (platform, device_type, num_entries, devices, num_devices)) - -OCL_FUNC_P(cl_context, clCreateContext, - (const cl_context_properties * properties, - cl_uint num_devices, - const cl_device_id * devices, - void (CL_CALLBACK * pfn_notify)(const char *, const void *, size_t, void *), - void * user_data, - cl_int * errcode_ret), - (properties, num_devices, devices, pfn_notify, user_data, errcode_ret)) - -OCL_FUNC(cl_int, clReleaseContext, (cl_context context), (context)) - - -OCL_FUNC(cl_int, clRetainContext, (cl_context context), (context)) -/* -OCL_FUNC_P(cl_context, clCreateContextFromType, - (const cl_context_properties * properties, - cl_device_type device_type, - void (CL_CALLBACK * pfn_notify)(const char *, const void *, size_t, void *), - void * user_data, - cl_int * errcode_ret), - (properties, device_type, pfn_notify, user_data, errcode_ret)) - -OCL_FUNC(cl_int, clGetContextInfo, - (cl_context context, - cl_context_info param_name, - size_t param_value_size, - void * param_value, - size_t * param_value_size_ret), - (context, param_name, param_value_size, - param_value, param_value_size_ret)) -*/ -OCL_FUNC_P(cl_command_queue, clCreateCommandQueue, - (cl_context context, - cl_device_id device, - cl_command_queue_properties properties, - cl_int * errcode_ret), - (context, device, properties, errcode_ret)) - -OCL_FUNC(cl_int, clReleaseCommandQueue, (cl_command_queue command_queue), (command_queue)) - -OCL_FUNC_P(cl_mem, clCreateBuffer, - (cl_context context, - cl_mem_flags flags, - size_t size, - void * host_ptr, - cl_int * errcode_ret), - (context, flags, size, host_ptr, errcode_ret)) - -/* -OCL_FUNC(cl_int, clRetainCommandQueue, (cl_command_queue command_queue), (command_queue)) -*/ -OCL_FUNC(cl_int, clGetCommandQueueInfo, - (cl_command_queue command_queue, - cl_command_queue_info param_name, - size_t param_value_size, - void * param_value, - size_t * param_value_size_ret), - (command_queue, param_name, param_value_size, param_value, param_value_size_ret)) -/* -OCL_FUNC_P(cl_mem, clCreateSubBuffer, - (cl_mem buffer, - cl_mem_flags flags, - cl_buffer_create_type buffer_create_type, - const void * buffer_create_info, - cl_int * errcode_ret), - (buffer, flags, buffer_create_type, buffer_create_info, errcode_ret)) -*/ - -OCL_FUNC_P(cl_mem, clCreateImage, - (cl_context context, - cl_mem_flags flags, - const cl_image_format * image_format, - const cl_image_desc * image_desc, - void * host_ptr, - cl_int * errcode_ret), - (context, flags, image_format, image_desc, host_ptr, errcode_ret)) - -OCL_FUNC_P(cl_mem, clCreateImage2D, - (cl_context context, - cl_mem_flags flags, - const cl_image_format * image_format, - size_t image_width, - size_t image_height, - size_t image_row_pitch, - void * host_ptr, - cl_int *errcode_ret), - (context, flags, image_format, image_width, image_height, image_row_pitch, host_ptr, errcode_ret)) - -OCL_FUNC(cl_int, clGetSupportedImageFormats, - (cl_context context, - cl_mem_flags flags, - cl_mem_object_type image_type, - cl_uint num_entries, - cl_image_format * image_formats, - cl_uint * num_image_formats), - (context, flags, image_type, num_entries, image_formats, num_image_formats)) - - -OCL_FUNC(cl_int, clGetMemObjectInfo, - (cl_mem memobj, - cl_mem_info param_name, - size_t param_value_size, - void * param_value, - size_t * param_value_size_ret), - (memobj, param_name, param_value_size, param_value, param_value_size_ret)) - -OCL_FUNC(cl_int, clGetImageInfo, - (cl_mem image, - cl_image_info param_name, - size_t param_value_size, - void * param_value, - size_t * param_value_size_ret), - (image, param_name, param_value_size, param_value, param_value_size_ret)) - -/* -OCL_FUNC(cl_int, clCreateKernelsInProgram, - (cl_program program, - cl_uint num_kernels, - cl_kernel * kernels, - cl_uint * num_kernels_ret), - (program, num_kernels, kernels, num_kernels_ret)) - -OCL_FUNC(cl_int, clRetainKernel, (cl_kernel kernel), (kernel)) - -OCL_FUNC(cl_int, clGetKernelArgInfo, - (cl_kernel kernel, - cl_uint arg_indx, - cl_kernel_arg_info param_name, - size_t param_value_size, - void * param_value, - size_t * param_value_size_ret), - (kernel, arg_indx, param_name, param_value_size, param_value, param_value_size_ret)) - -OCL_FUNC(cl_int, clEnqueueReadImage, - (cl_command_queue command_queue, - cl_mem image, - cl_bool blocking_read, - const size_t * origin[3], - const size_t * region[3], - size_t row_pitch, - size_t slice_pitch, - void * ptr, - cl_uint num_events_in_wait_list, - const cl_event * event_wait_list, - cl_event * event), - (command_queue, image, blocking_read, origin, region, - row_pitch, slice_pitch, - ptr, - num_events_in_wait_list, - event_wait_list, - event)) - -OCL_FUNC(cl_int, clEnqueueWriteImage, - (cl_command_queue command_queue, - cl_mem image, - cl_bool blocking_write, - const size_t * origin[3], - const size_t * region[3], - size_t input_row_pitch, - size_t input_slice_pitch, - const void * ptr, - cl_uint num_events_in_wait_list, - const cl_event * event_wait_list, - cl_event * event), - (command_queue, image, blocking_write, origin, region, input_row_pitch, - input_slice_pitch, ptr, num_events_in_wait_list, event_wait_list, event)) - -OCL_FUNC(cl_int, clEnqueueFillImage, - (cl_command_queue command_queue, - cl_mem image, - const void * fill_color, - const size_t * origin[3], - const size_t * region[3], - cl_uint num_events_in_wait_list, - const cl_event * event_wait_list, - cl_event * event), - (command_queue, image, fill_color, origin, region, - num_events_in_wait_list, event_wait_list, event)) - -OCL_FUNC(cl_int, clEnqueueCopyImage, - (cl_command_queue command_queue, - cl_mem src_image, - cl_mem dst_image, - const size_t * src_origin[3], - const size_t * dst_origin[3], - const size_t * region[3], - cl_uint num_events_in_wait_list, - const cl_event * event_wait_list, - cl_event * event), - (command_queue, src_image, dst_image, src_origin, dst_origin, - region, num_events_in_wait_list, event_wait_list, event)) -*/ - -OCL_FUNC(cl_int, clEnqueueCopyImageToBuffer, - (cl_command_queue command_queue, - cl_mem src_image, - cl_mem dst_buffer, - const size_t * src_origin, - const size_t * region, - size_t dst_offset, - cl_uint num_events_in_wait_list, - const cl_event * event_wait_list, - cl_event * event), - (command_queue, src_image, dst_buffer, src_origin, region, dst_offset, - num_events_in_wait_list, event_wait_list, event)) - -OCL_FUNC(cl_int, clEnqueueCopyBufferToImage, - (cl_command_queue command_queue, - cl_mem src_buffer, - cl_mem dst_image, - size_t src_offset, - const size_t dst_origin[3], - const size_t region[3], - cl_uint num_events_in_wait_list, - const cl_event * event_wait_list, - cl_event * event), - (command_queue, src_buffer, dst_image, src_offset, dst_origin, - region, num_events_in_wait_list, event_wait_list, event)) - - OCL_FUNC(cl_int, clFlush, - (cl_command_queue command_queue), - (command_queue)) - -/* -OCL_FUNC_P(void*, clEnqueueMapImage, - (cl_command_queue command_queue, - cl_mem image, - cl_bool blocking_map, - cl_map_flags map_flags, - const size_t * origin[3], - const size_t * region[3], - size_t * image_row_pitch, - size_t * image_slice_pitch, - cl_uint num_events_in_wait_list, - const cl_event * event_wait_list, - cl_event * event, - cl_int * errcode_ret), - (command_queue, image, blocking_map, map_flags, origin, region, - image_row_pitch, image_slice_pitch, num_events_in_wait_list, - event_wait_list, event, errcode_ret)) -*/ - -/* -OCL_FUNC(cl_int, clRetainProgram, (cl_program program), (program)) - -OCL_FUNC(cl_int, clGetKernelInfo, - (cl_kernel kernel, - cl_kernel_info param_name, - size_t param_value_size, - void * param_value, - size_t * param_value_size_ret), - (kernel, param_name, param_value_size, param_value, param_value_size_ret)) - -*/ - -OCL_FUNC(cl_int, clRetainMemObject, (cl_mem memobj), (memobj)) - -OCL_FUNC(cl_int, clReleaseMemObject, (cl_mem memobj), (memobj)) - -/* -OCL_FUNC_P(cl_program, clCreateProgramWithSource, - (cl_context context, - cl_uint count, - const char ** strings, - const size_t * lengths, - cl_int * errcode_ret), - (context, count, strings, lengths, errcode_ret)) - -OCL_FUNC_P(cl_program, clCreateProgramWithBinary, - (cl_context context, - cl_uint num_devices, - const cl_device_id * device_list, - const size_t * lengths, - const unsigned char ** binaries, - cl_int * binary_status, - cl_int * errcode_ret), - (context, num_devices, device_list, lengths, binaries, binary_status, errcode_ret)) - -OCL_FUNC(cl_int, clReleaseProgram, (cl_program program), (program)) - -OCL_FUNC(cl_int, clBuildProgram, - (cl_program program, - cl_uint num_devices, - const cl_device_id * device_list, - const char * options, - void (CL_CALLBACK * pfn_notify)(cl_program, void *), - void * user_data), - (program, num_devices, device_list, options, pfn_notify, user_data)) - -OCL_FUNC(cl_int, clGetProgramInfo, - (cl_program program, - cl_program_info param_name, - size_t param_value_size, - void * param_value, - size_t * param_value_size_ret), - (program, param_name, param_value_size, param_value, param_value_size_ret)) - -OCL_FUNC(cl_int, clGetProgramBuildInfo, - (cl_program program, - cl_device_id device, - cl_program_build_info param_name, - size_t param_value_size, - void * param_value, - size_t * param_value_size_ret), - (program, device, param_name, param_value_size, param_value, param_value_size_ret)) -*/ -OCL_FUNC_P(cl_kernel, clCreateKernel, - (cl_program program, - const char * kernel_name, - cl_int * errcode_ret), - (program, kernel_name, errcode_ret)) - -OCL_FUNC(cl_int, clReleaseKernel, (cl_kernel kernel), (kernel)) - -OCL_FUNC(cl_int, clSetKernelArg, - (cl_kernel kernel, - cl_uint arg_index, - size_t arg_size, - const void * arg_value), - (kernel, arg_index, arg_size, arg_value)) - -OCL_FUNC(cl_int, clGetKernelWorkGroupInfo, - (cl_kernel kernel, - cl_device_id device, - cl_kernel_work_group_info param_name, - size_t param_value_size, - void * param_value, - size_t * param_value_size_ret), - (kernel, device, param_name, param_value_size, param_value, param_value_size_ret)) - -OCL_FUNC(cl_int, clFinish, (cl_command_queue command_queue), (command_queue)) - -OCL_FUNC(cl_int, clEnqueueReadBuffer, - (cl_command_queue command_queue, - cl_mem buffer, - cl_bool blocking_read, - size_t offset, - size_t size, - void * ptr, - cl_uint num_events_in_wait_list, - const cl_event * event_wait_list, - cl_event * event), - (command_queue, buffer, blocking_read, offset, size, ptr, - num_events_in_wait_list, event_wait_list, event)) - -OCL_FUNC(cl_int, clEnqueueReadBufferRect, - (cl_command_queue command_queue, - cl_mem buffer, - cl_bool blocking_read, - const size_t * buffer_offset, - const size_t * host_offset, - const size_t * region, - size_t buffer_row_pitch, - size_t buffer_slice_pitch, - size_t host_row_pitch, - size_t host_slice_pitch, - void * ptr, - cl_uint num_events_in_wait_list, - const cl_event * event_wait_list, - cl_event * event), - (command_queue, buffer, blocking_read, buffer_offset, host_offset, region, buffer_row_pitch, - buffer_slice_pitch, host_row_pitch, host_slice_pitch, ptr, num_events_in_wait_list, - event_wait_list, event)) - -OCL_FUNC(cl_int, clEnqueueWriteBuffer, - (cl_command_queue command_queue, - cl_mem buffer, - cl_bool blocking_write, - size_t offset, - size_t size, - const void * ptr, - cl_uint num_events_in_wait_list, - const cl_event * event_wait_list, - cl_event * event), - (command_queue, buffer, blocking_write, offset, size, ptr, - num_events_in_wait_list, event_wait_list, event)) - -OCL_FUNC(cl_int, clEnqueueWriteBufferRect, - (cl_command_queue command_queue, - cl_mem buffer, - cl_bool blocking_write, - const size_t * buffer_offset, - const size_t * host_offset, - const size_t * region, - size_t buffer_row_pitch, - size_t buffer_slice_pitch, - size_t host_row_pitch, - size_t host_slice_pitch, - const void * ptr, - cl_uint num_events_in_wait_list, - const cl_event * event_wait_list, - cl_event * event), - (command_queue, buffer, blocking_write, buffer_offset, host_offset, - region, buffer_row_pitch, buffer_slice_pitch, host_row_pitch, - host_slice_pitch, ptr, num_events_in_wait_list, event_wait_list, event)) - -/*OCL_FUNC(cl_int, clEnqueueFillBuffer, - (cl_command_queue command_queue, - cl_mem buffer, - const void * pattern, - size_t pattern_size, - size_t offset, - size_t size, - cl_uint num_events_in_wait_list, - const cl_event * event_wait_list, - cl_event * event), - (command_queue, buffer, pattern, pattern_size, offset, size, - num_events_in_wait_list, event_wait_list, event))*/ - -OCL_FUNC(cl_int, clEnqueueCopyBuffer, - (cl_command_queue command_queue, - cl_mem src_buffer, - cl_mem dst_buffer, - size_t src_offset, - size_t dst_offset, - size_t size, - cl_uint num_events_in_wait_list, - const cl_event * event_wait_list, - cl_event * event), - (command_queue, src_buffer, dst_buffer, src_offset, dst_offset, - size, num_events_in_wait_list, event_wait_list, event)) - -OCL_FUNC(cl_int, clEnqueueCopyBufferRect, - (cl_command_queue command_queue, - cl_mem src_buffer, - cl_mem dst_buffer, - const size_t * src_origin, - const size_t * dst_origin, - const size_t * region, - size_t src_row_pitch, - size_t src_slice_pitch, - size_t dst_row_pitch, - size_t dst_slice_pitch, - cl_uint num_events_in_wait_list, - const cl_event * event_wait_list, - cl_event * event), - (command_queue, src_buffer, dst_buffer, src_origin, dst_origin, - region, src_row_pitch, src_slice_pitch, dst_row_pitch, dst_slice_pitch, - num_events_in_wait_list, event_wait_list, event)) - -OCL_FUNC_P(void*, clEnqueueMapBuffer, - (cl_command_queue command_queue, - cl_mem buffer, - cl_bool blocking_map, - cl_map_flags map_flags, - size_t offset, - size_t size, - cl_uint num_events_in_wait_list, - const cl_event * event_wait_list, - cl_event * event, - cl_int * errcode_ret), - (command_queue, buffer, blocking_map, map_flags, offset, size, - num_events_in_wait_list, event_wait_list, event, errcode_ret)) - -OCL_FUNC(cl_int, clEnqueueUnmapMemObject, - (cl_command_queue command_queue, - cl_mem memobj, - void * mapped_ptr, - cl_uint num_events_in_wait_list, - const cl_event * event_wait_list, - cl_event * event), - (command_queue, memobj, mapped_ptr, num_events_in_wait_list, event_wait_list, event)) - -OCL_FUNC(cl_int, clEnqueueNDRangeKernel, - (cl_command_queue command_queue, - cl_kernel kernel, - cl_uint work_dim, - const size_t * global_work_offset, - const size_t * global_work_size, - const size_t * local_work_size, - cl_uint num_events_in_wait_list, - const cl_event * event_wait_list, - cl_event * event), - (command_queue, kernel, work_dim, global_work_offset, global_work_size, - local_work_size, num_events_in_wait_list, event_wait_list, event)) - -OCL_FUNC(cl_int, clEnqueueTask, - (cl_command_queue command_queue, - cl_kernel kernel, - cl_uint num_events_in_wait_list, - const cl_event * event_wait_list, - cl_event * event), - (command_queue, kernel, num_events_in_wait_list, event_wait_list, event)) - -OCL_FUNC(cl_int, clSetEventCallback, - (cl_event event, - cl_int command_exec_callback_type , - void (CL_CALLBACK *pfn_event_notify) (cl_event event, cl_int event_command_exec_status, void *user_data), - void *user_data), - (event, command_exec_callback_type, pfn_event_notify, user_data)) - -OCL_FUNC(cl_int, clReleaseEvent, (cl_event event), (event)) - -OCL_FUNC(cl_int, clWaitForEvents, - (cl_uint num_events, const cl_event *event_list), - (num_events, event_list)) - - -OCL_FUNC(cl_int, clGetEventProfilingInfo, ( - cl_event event, - cl_profiling_info param_name, - size_t param_value_size, - void *param_value, - size_t *param_value_size_ret), - (event, param_name, param_value_size, param_value, param_value_size_ret)) - -} - -#endif - -#ifndef CL_VERSION_1_2 -#define CL_VERSION_1_2 -#endif diff --git a/modules/core/src/ocl_disabled.impl.hpp b/modules/core/src/ocl_disabled.impl.hpp new file mode 100644 index 0000000000..a0516476bf --- /dev/null +++ b/modules/core/src/ocl_disabled.impl.hpp @@ -0,0 +1,366 @@ +// 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. + +#include "opencv2/core/ocl_genbase.hpp" + +#if defined(_MSC_VER) + #pragma warning(push) + #pragma warning(disable : 4100) + #pragma warning(disable : 4702) +#elif defined(__clang__) + #pragma clang diagnostic push + #pragma clang diagnostic ignored "-Wunused-parameter" +#elif defined(__GNUC__) + #pragma GCC diagnostic push + #pragma GCC diagnostic ignored "-Wunused-parameter" +#endif + +namespace cv { namespace ocl { + +static +CV_NORETURN void throw_no_ocl() +{ + CV_Error(Error::OpenCLApiCallError, "OpenCV build without OpenCL support"); +} +#define OCL_NOT_AVAILABLE() throw_no_ocl(); + +CV_EXPORTS_W bool haveOpenCL() { return false; } +CV_EXPORTS_W bool useOpenCL() { return false; } +CV_EXPORTS_W bool haveAmdBlas() { return false; } +CV_EXPORTS_W bool haveAmdFft() { return false; } +CV_EXPORTS_W void setUseOpenCL(bool flag) { /* nothing */ } +CV_EXPORTS_W void finish() { /* nothing */ } + +CV_EXPORTS bool haveSVM() { return false; } + +Device::Device() : p(NULL) { } +Device::Device(void* d) : p(NULL) { OCL_NOT_AVAILABLE(); } +Device::Device(const Device& d) : p(NULL) { } +Device& Device::operator=(const Device& d) { return *this; } +Device::~Device() { } + +void Device::set(void* d) { OCL_NOT_AVAILABLE(); } + +String Device::name() const { OCL_NOT_AVAILABLE(); } +String Device::extensions() const { OCL_NOT_AVAILABLE(); } +bool Device::isExtensionSupported(const String& extensionName) const { OCL_NOT_AVAILABLE(); } +String Device::version() const { OCL_NOT_AVAILABLE(); } +String Device::vendorName() const { OCL_NOT_AVAILABLE(); } +String Device::OpenCL_C_Version() const { OCL_NOT_AVAILABLE(); } +String Device::OpenCLVersion() const { OCL_NOT_AVAILABLE(); } +int Device::deviceVersionMajor() const { OCL_NOT_AVAILABLE(); } +int Device::deviceVersionMinor() const { OCL_NOT_AVAILABLE(); } +String Device::driverVersion() const { OCL_NOT_AVAILABLE(); } +void* Device::ptr() const { /*OCL_NOT_AVAILABLE();*/ return NULL; } + +int Device::type() const { OCL_NOT_AVAILABLE(); } + +int Device::addressBits() const { OCL_NOT_AVAILABLE(); } +bool Device::available() const { OCL_NOT_AVAILABLE(); } +bool Device::compilerAvailable() const { OCL_NOT_AVAILABLE(); } +bool Device::linkerAvailable() const { OCL_NOT_AVAILABLE(); } + +int Device::doubleFPConfig() const { OCL_NOT_AVAILABLE(); } +int Device::singleFPConfig() const { OCL_NOT_AVAILABLE(); } +int Device::halfFPConfig() const { OCL_NOT_AVAILABLE(); } + +bool Device::endianLittle() const { OCL_NOT_AVAILABLE(); } +bool Device::errorCorrectionSupport() const { OCL_NOT_AVAILABLE(); } + +int Device::executionCapabilities() const { OCL_NOT_AVAILABLE(); } + +size_t Device::globalMemCacheSize() const { OCL_NOT_AVAILABLE(); } + +int Device::globalMemCacheType() const { OCL_NOT_AVAILABLE(); } +int Device::globalMemCacheLineSize() const { OCL_NOT_AVAILABLE(); } +size_t Device::globalMemSize() const { OCL_NOT_AVAILABLE(); } + +size_t Device::localMemSize() const { OCL_NOT_AVAILABLE(); } +int Device::localMemType() const { return NO_LOCAL_MEM; } +bool Device::hostUnifiedMemory() const { OCL_NOT_AVAILABLE(); } + +bool Device::imageSupport() const { OCL_NOT_AVAILABLE(); } + +bool Device::imageFromBufferSupport() const { OCL_NOT_AVAILABLE(); } +uint Device::imagePitchAlignment() const { OCL_NOT_AVAILABLE(); } +uint Device::imageBaseAddressAlignment() const { OCL_NOT_AVAILABLE(); } + +bool Device::intelSubgroupsSupport() const { OCL_NOT_AVAILABLE(); } + +size_t Device::image2DMaxWidth() const { OCL_NOT_AVAILABLE(); } +size_t Device::image2DMaxHeight() const { OCL_NOT_AVAILABLE(); } + +size_t Device::image3DMaxWidth() const { OCL_NOT_AVAILABLE(); } +size_t Device::image3DMaxHeight() const { OCL_NOT_AVAILABLE(); } +size_t Device::image3DMaxDepth() const { OCL_NOT_AVAILABLE(); } + +size_t Device::imageMaxBufferSize() const { OCL_NOT_AVAILABLE(); } +size_t Device::imageMaxArraySize() const { OCL_NOT_AVAILABLE(); } + +int Device::vendorID() const { OCL_NOT_AVAILABLE(); } + +int Device::maxClockFrequency() const { OCL_NOT_AVAILABLE(); } +int Device::maxComputeUnits() const { OCL_NOT_AVAILABLE(); } +int Device::maxConstantArgs() const { OCL_NOT_AVAILABLE(); } +size_t Device::maxConstantBufferSize() const { OCL_NOT_AVAILABLE(); } + +size_t Device::maxMemAllocSize() const { OCL_NOT_AVAILABLE(); } +size_t Device::maxParameterSize() const { OCL_NOT_AVAILABLE(); } + +int Device::maxReadImageArgs() const { OCL_NOT_AVAILABLE(); } +int Device::maxWriteImageArgs() const { OCL_NOT_AVAILABLE(); } +int Device::maxSamplers() const { OCL_NOT_AVAILABLE(); } + +size_t Device::maxWorkGroupSize() const { OCL_NOT_AVAILABLE(); } +int Device::maxWorkItemDims() const { OCL_NOT_AVAILABLE(); } +void Device::maxWorkItemSizes(size_t*) const { OCL_NOT_AVAILABLE(); } + +int Device::memBaseAddrAlign() const { OCL_NOT_AVAILABLE(); } + +int Device::nativeVectorWidthChar() const { OCL_NOT_AVAILABLE(); } +int Device::nativeVectorWidthShort() const { OCL_NOT_AVAILABLE(); } +int Device::nativeVectorWidthInt() const { OCL_NOT_AVAILABLE(); } +int Device::nativeVectorWidthLong() const { OCL_NOT_AVAILABLE(); } +int Device::nativeVectorWidthFloat() const { OCL_NOT_AVAILABLE(); } +int Device::nativeVectorWidthDouble() const { OCL_NOT_AVAILABLE(); } +int Device::nativeVectorWidthHalf() const { OCL_NOT_AVAILABLE(); } + +int Device::preferredVectorWidthChar() const { OCL_NOT_AVAILABLE(); } +int Device::preferredVectorWidthShort() const { OCL_NOT_AVAILABLE(); } +int Device::preferredVectorWidthInt() const { OCL_NOT_AVAILABLE(); } +int Device::preferredVectorWidthLong() const { OCL_NOT_AVAILABLE(); } +int Device::preferredVectorWidthFloat() const { OCL_NOT_AVAILABLE(); } +int Device::preferredVectorWidthDouble() const { OCL_NOT_AVAILABLE(); } +int Device::preferredVectorWidthHalf() const { OCL_NOT_AVAILABLE(); } + +size_t Device::printfBufferSize() const { OCL_NOT_AVAILABLE(); } +size_t Device::profilingTimerResolution() const { OCL_NOT_AVAILABLE(); } + +/* static */ +const Device& Device::getDefault() +{ + static Device dummy; + return dummy; +} + + +Context::Context() : p(NULL) { } +Context::Context(int dtype) : p(NULL) { } +Context::~Context() { } +Context::Context(const Context& c) : p(NULL) { } +Context& Context::operator=(const Context& c) { return *this; } + +bool Context::create() { return false; } +bool Context::create(int dtype) { return false; } +size_t Context::ndevices() const { return 0; } +const Device& Context::device(size_t idx) const { OCL_NOT_AVAILABLE(); } +Program Context::getProg(const ProgramSource& prog, const String& buildopt, String& errmsg) { OCL_NOT_AVAILABLE(); } +void Context::unloadProg(Program& prog) { } + +/* static */ +Context& Context::getDefault(bool initialize) +{ + static Context dummy; + return dummy; +} +void* Context::ptr() const { return NULL; } + +bool Context::useSVM() const { return false; } +void Context::setUseSVM(bool enabled) { } + +Platform::Platform() : p(NULL) { } +Platform::~Platform() { } +Platform::Platform(const Platform&) : p(NULL) { } +Platform& Platform::operator=(const Platform&) { return *this; } + +void* Platform::ptr() const { return NULL; } + +/* static */ +Platform& Platform::getDefault() +{ + static Platform dummy; + return dummy; +} + +void attachContext(const String& platformName, void* platformID, void* context, void* deviceID) { OCL_NOT_AVAILABLE(); } +void convertFromBuffer(void* cl_mem_buffer, size_t step, int rows, int cols, int type, UMat& dst) { OCL_NOT_AVAILABLE(); } +void convertFromImage(void* cl_mem_image, UMat& dst) { OCL_NOT_AVAILABLE(); } + +void initializeContextFromHandle(Context& ctx, void* platform, void* context, void* device) { OCL_NOT_AVAILABLE(); } + +Queue::Queue() : p(NULL) { } +Queue::Queue(const Context& c, const Device& d) : p(NULL) { OCL_NOT_AVAILABLE(); } +Queue::~Queue() { } +Queue::Queue(const Queue& q) {} +Queue& Queue::operator=(const Queue& q) { return *this; } + +bool Queue::create(const Context& c, const Device& d) { OCL_NOT_AVAILABLE(); } +void Queue::finish() {} +void* Queue::ptr() const { return NULL; } +/* static */ +Queue& Queue::getDefault() +{ + static Queue dummy; + return dummy; +} + +/// @brief Returns OpenCL command queue with enable profiling mode support +const Queue& Queue::getProfilingQueue() const { OCL_NOT_AVAILABLE(); } + + +KernelArg::KernelArg() + : flags(0), m(0), obj(0), sz(0), wscale(1), iwscale(1) +{ +} + +KernelArg::KernelArg(int _flags, UMat* _m, int _wscale, int _iwscale, const void* _obj, size_t _sz) + : flags(_flags), m(_m), obj(_obj), sz(_sz), wscale(_wscale), iwscale(_iwscale) +{ + OCL_NOT_AVAILABLE(); +} + +KernelArg KernelArg::Constant(const Mat& m) +{ + OCL_NOT_AVAILABLE(); +} + + +Kernel::Kernel() : p(NULL) { } +Kernel::Kernel(const char* kname, const Program& prog) : p(NULL) { OCL_NOT_AVAILABLE(); } +Kernel::Kernel(const char* kname, const ProgramSource& prog, const String& buildopts, String* errmsg) : p(NULL) { OCL_NOT_AVAILABLE(); } +Kernel::~Kernel() { } +Kernel::Kernel(const Kernel& k) : p(NULL) { } +Kernel& Kernel::operator=(const Kernel& k) { return *this; } + +bool Kernel::empty() const { return true; } +bool Kernel::create(const char* kname, const Program& prog) { OCL_NOT_AVAILABLE(); } +bool Kernel::create(const char* kname, const ProgramSource& prog, const String& buildopts, String* errmsg) { OCL_NOT_AVAILABLE(); } + +int Kernel::set(int i, const void* value, size_t sz) { OCL_NOT_AVAILABLE(); } +int Kernel::set(int i, const Image2D& image2D) { OCL_NOT_AVAILABLE(); } +int Kernel::set(int i, const UMat& m) { OCL_NOT_AVAILABLE(); } +int Kernel::set(int i, const KernelArg& arg) { OCL_NOT_AVAILABLE(); } + +bool Kernel::run(int dims, size_t globalsize[], size_t localsize[], bool sync, const Queue& q) { OCL_NOT_AVAILABLE(); } +bool Kernel::runTask(bool sync, const Queue& q) { OCL_NOT_AVAILABLE(); } + +int64 Kernel::runProfiling(int dims, size_t globalsize[], size_t localsize[], const Queue& q) { OCL_NOT_AVAILABLE(); } + +size_t Kernel::workGroupSize() const { OCL_NOT_AVAILABLE(); } +size_t Kernel::preferedWorkGroupSizeMultiple() const { OCL_NOT_AVAILABLE(); } +bool Kernel::compileWorkGroupSize(size_t wsz[]) const { OCL_NOT_AVAILABLE(); } +size_t Kernel::localMemSize() const { OCL_NOT_AVAILABLE(); } + +void* Kernel::ptr() const { return NULL; } + + +Program::Program() : p(NULL) { } +Program::Program(const ProgramSource& src, const String& buildflags, String& errmsg) : p(NULL) { OCL_NOT_AVAILABLE(); } +Program::Program(const Program& prog) : p(NULL) { } +Program& Program::operator=(const Program& prog) { return *this; } +Program::~Program() { } + +bool Program::create(const ProgramSource& src, const String& buildflags, String& errmsg) { OCL_NOT_AVAILABLE(); } + +void* Program::ptr() const { return NULL; } + +void Program::getBinary(std::vector& binary) const { OCL_NOT_AVAILABLE(); } + +bool Program::read(const String& buf, const String& buildflags) { OCL_NOT_AVAILABLE(); } +bool Program::write(String& buf) const { OCL_NOT_AVAILABLE(); } +const ProgramSource& Program::source() const { OCL_NOT_AVAILABLE(); } +String Program::getPrefix() const { OCL_NOT_AVAILABLE(); } +/* static */ String Program::getPrefix(const String& buildflags) { OCL_NOT_AVAILABLE(); } + + +ProgramSource::ProgramSource() : p(NULL) { } +ProgramSource::ProgramSource(const String& module, const String& name, const String& codeStr, const String& codeHash) : p(NULL) { } +ProgramSource::ProgramSource(const String& prog) : p(NULL) { } +ProgramSource::ProgramSource(const char* prog) : p(NULL) { } +ProgramSource::~ProgramSource() { } +ProgramSource::ProgramSource(const ProgramSource& prog) : p(NULL) { } +ProgramSource& ProgramSource::operator=(const ProgramSource& prog) { return *this; } + +const String& ProgramSource::source() const { OCL_NOT_AVAILABLE(); } +ProgramSource::hash_t ProgramSource::hash() const { OCL_NOT_AVAILABLE(); } + +/* static */ ProgramSource ProgramSource::fromBinary(const String& module, const String& name, const unsigned char* binary, const size_t size, const cv::String& buildOptions) { OCL_NOT_AVAILABLE(); } +/* static */ ProgramSource ProgramSource::fromSPIR(const String& module, const String& name, const unsigned char* binary, const size_t size, const cv::String& buildOptions) { OCL_NOT_AVAILABLE(); } + + +PlatformInfo::PlatformInfo() : p(NULL) { } +PlatformInfo::PlatformInfo(void* id) : p(NULL) { OCL_NOT_AVAILABLE(); } +PlatformInfo::~PlatformInfo() { } + +PlatformInfo::PlatformInfo(const PlatformInfo& i) : p(NULL) { } +PlatformInfo& PlatformInfo::operator=(const PlatformInfo& i) { return *this; } + +String PlatformInfo::name() const { OCL_NOT_AVAILABLE(); } +String PlatformInfo::vendor() const { OCL_NOT_AVAILABLE(); } +String PlatformInfo::version() const { OCL_NOT_AVAILABLE(); } +int PlatformInfo::deviceNumber() const { OCL_NOT_AVAILABLE(); } +void PlatformInfo::getDevice(Device& device, int d) const { OCL_NOT_AVAILABLE(); } + +const char* convertTypeStr(int sdepth, int ddepth, int cn, char* buf) { OCL_NOT_AVAILABLE(); } +const char* typeToStr(int t) { OCL_NOT_AVAILABLE(); } +const char* memopTypeToStr(int t) { OCL_NOT_AVAILABLE(); } +const char* vecopTypeToStr(int t) { OCL_NOT_AVAILABLE(); } +const char* getOpenCLErrorString(int errorCode) { OCL_NOT_AVAILABLE(); } +String kernelToStr(InputArray _kernel, int ddepth, const char* name) { OCL_NOT_AVAILABLE(); } +void getPlatfomsInfo(std::vector& platform_info) { OCL_NOT_AVAILABLE(); } + + +int predictOptimalVectorWidth(InputArray src1, InputArray src2, InputArray src3, + InputArray src4, InputArray src5, InputArray src6, + InputArray src7, InputArray src8, InputArray src9, + OclVectorStrategy strat) +{ OCL_NOT_AVAILABLE(); } + +int checkOptimalVectorWidth(const int *vectorWidths, + InputArray src1, InputArray src2, InputArray src3, + InputArray src4, InputArray src5, InputArray src6, + InputArray src7, InputArray src8, InputArray src9, + OclVectorStrategy strat) +{ OCL_NOT_AVAILABLE(); } + +int predictOptimalVectorWidthMax(InputArray src1, InputArray src2, InputArray src3, + InputArray src4, InputArray src5, InputArray src6, + InputArray src7, InputArray src8, InputArray src9) +{ OCL_NOT_AVAILABLE(); } + +void buildOptionsAddMatrixDescription(String& buildOptions, const String& name, InputArray _m) { OCL_NOT_AVAILABLE(); } + + +Image2D::Image2D() : p(NULL) { } +Image2D::Image2D(const UMat &src, bool norm, bool alias) { OCL_NOT_AVAILABLE(); } +Image2D::Image2D(const Image2D & i) : p(NULL) { OCL_NOT_AVAILABLE(); } +Image2D::~Image2D() { } +Image2D& Image2D::operator=(const Image2D & i) { return *this; } + +/* static */ bool Image2D::canCreateAlias(const UMat &u) { OCL_NOT_AVAILABLE(); } +/* static */ bool Image2D::isFormatSupported(int depth, int cn, bool norm) { OCL_NOT_AVAILABLE(); } + +void* Image2D::ptr() const { return NULL; } + + +Timer::Timer(const Queue& q) : p(NULL) {} +Timer::~Timer() {} +void Timer::start() { OCL_NOT_AVAILABLE(); } +void Timer::stop() { OCL_NOT_AVAILABLE();} + +uint64 Timer::durationNS() const { OCL_NOT_AVAILABLE(); } + +MatAllocator* getOpenCLAllocator() { return NULL; } + +internal::ProgramEntry::operator ProgramSource&() const { OCL_NOT_AVAILABLE(); } + +}} + +#if defined(_MSC_VER) + #pragma warning(pop) +#elif defined(__clang__) + #pragma clang diagnostic pop +#elif defined(__GNUC__) + #pragma GCC diagnostic pop +#endif diff --git a/modules/features2d/src/orb.cpp b/modules/features2d/src/orb.cpp index a664dcac5d..881fc01516 100644 --- a/modules/features2d/src/orb.cpp +++ b/modules/features2d/src/orb.cpp @@ -983,7 +983,11 @@ void ORB_Impl::detectAndCompute( InputArray _image, InputArray _mask, int descPatchSize = cvCeil(halfPatchSize*sqrt(2.0)); int border = std::max(edgeThreshold, std::max(descPatchSize, HARRIS_BLOCK_SIZE/2))+1; +#ifdef HAVE_OPENCL bool useOCL = ocl::isOpenCLActivated() && OCL_FORCE_CHECK(_image.isUMat() || _descriptors.isUMat()); +#else + bool useOCL = false; +#endif Mat image = _image.getMat(), mask = _mask.getMat(); if( image.type() != CV_8UC1 ) diff --git a/modules/flann/include/opencv2/flann/kmeans_index.h b/modules/flann/include/opencv2/flann/kmeans_index.h index 382d95ea26..cb1a54a6d6 100644 --- a/modules/flann/include/opencv2/flann/kmeans_index.h +++ b/modules/flann/include/opencv2/flann/kmeans_index.h @@ -50,6 +50,9 @@ #include "logger.h" #define BITS_PER_CHAR 8 +#define BITS_PER_BASE 2 // for DNA/RNA sequences +#define BASE_PER_CHAR (BITS_PER_CHAR/BITS_PER_BASE) +#define HISTOS_PER_BASE (1<( ensureSquareDistance( @@ -814,6 +817,73 @@ private: } + void computeDnaNodeStatistics(KMeansNodePtr node, int* indices, + unsigned int indices_length) + { + const unsigned int histos_veclen = static_cast( + veclen_*sizeof(CentersType)*(HISTOS_PER_BASE*BASE_PER_CHAR)); + + unsigned long long variance = 0ull; + unsigned int* histograms = new unsigned int[histos_veclen]; + memset(histograms, 0, sizeof(unsigned int)*histos_veclen); + + for (unsigned int i=0; i( ensureSquareDistance( + distance_(dataset_[indices[i]], ZeroIterator(), veclen_))); + + unsigned char* vec = (unsigned char*)dataset_[indices[i]]; + for (size_t k=0, l=0; k>2) & 0x03)]++; + histograms[k + 8 + ((vec[l]>>4) & 0x03)]++; + histograms[k +12 + ((vec[l]>>6) & 0x03)]++; + } + } + + CentersType* mean = new CentersType[veclen_]; + memoryCounter_ += int(veclen_*sizeof(CentersType)); + unsigned char* char_mean = (unsigned char*)mean; + unsigned int* h = histograms; + for (size_t k=0, l=0; k h[k+1] ? h[k+2] > h[k+3] ? h[k] > h[k+2] ? 0x00 : 0x10 + : h[k] > h[k+3] ? 0x00 : 0x11 + : h[k+2] > h[k+3] ? h[k+1] > h[k+2] ? 0x01 : 0x10 + : h[k+1] > h[k+3] ? 0x01 : 0x11) + | (h[k+4]>h[k+5] ? h[k+6] > h[k+7] ? h[k+4] > h[k+6] ? 0x00 : 0x1000 + : h[k+4] > h[k+7] ? 0x00 : 0x1100 + : h[k+6] > h[k+7] ? h[k+5] > h[k+6] ? 0x0100 : 0x1000 + : h[k+5] > h[k+7] ? 0x0100 : 0x1100) + | (h[k+8]>h[k+9] ? h[k+10]>h[k+11] ? h[k+8] >h[k+10] ? 0x00 : 0x100000 + : h[k+8] >h[k+11] ? 0x00 : 0x110000 + : h[k+10]>h[k+11] ? h[k+9] >h[k+10] ? 0x010000 : 0x100000 + : h[k+9] >h[k+11] ? 0x010000 : 0x110000) + | (h[k+12]>h[k+13] ? h[k+14]>h[k+15] ? h[k+12] >h[k+14] ? 0x00 : 0x10000000 + : h[k+12] >h[k+15] ? 0x00 : 0x11000000 + : h[k+14]>h[k+15] ? h[k+13] >h[k+14] ? 0x01000000 : 0x10000000 + : h[k+13] >h[k+15] ? 0x01000000 : 0x11000000); + } + variance = static_cast( + 0.5 + static_cast(variance) / static_cast(indices_length)); + variance -= static_cast( + ensureSquareDistance( + distance_(mean, ZeroIterator(), veclen_))); + + DistanceType radius = 0; + for (unsigned int i=0; iradius) { + radius = tmp; + } + } + + node->variance = static_cast(variance); + node->radius = radius; + node->pivot = mean; + + delete[] histograms; + } + + template void computeNodeStatistics(KMeansNodePtr node, int* indices, unsigned int indices_length, @@ -847,6 +917,22 @@ private: computeBitfieldNodeStatistics(node, indices, indices_length); } + void computeNodeStatistics(KMeansNodePtr node, int* indices, + unsigned int indices_length, + const cvflann::DNAmmingLUT* identifier) + { + (void)identifier; + computeDnaNodeStatistics(node, indices, indices_length); + } + + void computeNodeStatistics(KMeansNodePtr node, int* indices, + unsigned int indices_length, + const cvflann::DNAmming2* identifier) + { + (void)identifier; + computeDnaNodeStatistics(node, indices, indices_length); + } + void refineClustering(int* indices, int indices_length, int branching, CentersType** centers, std::vector& radiuses, int* belongs_to, int* count) @@ -1040,6 +1126,112 @@ private: } + void refineDnaClustering(int* indices, int indices_length, int branching, CentersType** centers, + std::vector& radiuses, int* belongs_to, int* count) + { + for (int i=0; i( + veclen_*sizeof(CentersType)*(HISTOS_PER_BASE*BASE_PER_CHAR)); + cv::AutoBuffer histos_buf(branching*histos_veclen); + Matrix histos(histos_buf.data(), branching, histos_veclen); + + bool converged = false; + int iteration = 0; + while (!converged && iteration>2) & 0x03)]++; + h[k + 8 + ((vec[l]>>4) & 0x03)]++; + h[k +12 + ((vec[l]>>6) & 0x03)]++; + } + } + for (int i=0; i h[k+1] ? h[k+2] > h[k+3] ? h[k] > h[k+2] ? 0x00 : 0x10 + : h[k] > h[k+3] ? 0x00 : 0x11 + : h[k+2] > h[k+3] ? h[k+1] > h[k+2] ? 0x01 : 0x10 + : h[k+1] > h[k+3] ? 0x01 : 0x11) + | (h[k+4]>h[k+5] ? h[k+6] > h[k+7] ? h[k+4] > h[k+6] ? 0x00 : 0x1000 + : h[k+4] > h[k+7] ? 0x00 : 0x1100 + : h[k+6] > h[k+7] ? h[k+5] > h[k+6] ? 0x0100 : 0x1000 + : h[k+5] > h[k+7] ? 0x0100 : 0x1100) + | (h[k+8]>h[k+9] ? h[k+10]>h[k+11] ? h[k+8] >h[k+10] ? 0x00 : 0x100000 + : h[k+8] >h[k+11] ? 0x00 : 0x110000 + : h[k+10]>h[k+11] ? h[k+9] >h[k+10] ? 0x010000 : 0x100000 + : h[k+9] >h[k+11] ? 0x010000 : 0x110000) + | (h[k+12]>h[k+13] ? h[k+14]>h[k+15] ? h[k+12] >h[k+14] ? 0x00 : 0x10000000 + : h[k+12] >h[k+15] ? 0x00 : 0x11000000 + : h[k+14]>h[k+15] ? h[k+13] >h[k+14] ? 0x01000000 : 0x10000000 + : h[k+13] >h[k+15] ? 0x01000000 : 0x11000000); + } + } + + std::vector new_centroids(indices_length); + std::vector dists(indices_length); + + // reassign points to clusters + KMeansDistanceComputer invoker( + distance_, dataset_, branching, indices, centers, veclen_, new_centroids, dists); + parallel_for_(cv::Range(0, (int)indices_length), invoker); + + for (int i=0; i < indices_length; ++i) { + DistanceType dist(dists[i]); + int new_centroid(new_centroids[i]); + if (dist > radiuses[new_centroid]) { + radiuses[new_centroid] = dist; + } + if (new_centroid != belongs_to[i]) { + count[belongs_to[i]]--; + count[new_centroid]++; + belongs_to[i] = new_centroid; + converged = false; + } + } + + for (int i=0; i& radiuses, int* belongs_to, int* count) @@ -1139,7 +1331,7 @@ private: /** * The methods responsible with doing the recursive hierarchical clustering on * binary vectors. - * As some might have heared that KMeans on binary data doesn't make sense, + * As some might have heard that KMeans on binary data doesn't make sense, * it's worth a little explanation why it actually fairly works. As * with the Hierarchical Clustering algortihm, we seed several centers for the * current node by picking some of its points. Then in a first pass each point @@ -1222,6 +1414,34 @@ private: } + void refineAndSplitClustering( + KMeansNodePtr node, int* indices, int indices_length, int branching, + int level, CentersType** centers, std::vector& radiuses, + int* belongs_to, int* count, const cvflann::DNAmmingLUT* identifier) + { + (void)identifier; + refineDnaClustering( + indices, indices_length, branching, centers, radiuses, belongs_to, count); + + computeAnyBitfieldSubClustering(node, indices, indices_length, branching, + level, centers, radiuses, belongs_to, count); + } + + + void refineAndSplitClustering( + KMeansNodePtr node, int* indices, int indices_length, int branching, + int level, CentersType** centers, std::vector& radiuses, + int* belongs_to, int* count, const cvflann::DNAmming2* identifier) + { + (void)identifier; + refineDnaClustering( + indices, indices_length, branching, centers, radiuses, belongs_to, count); + + computeAnyBitfieldSubClustering(node, indices, indices_length, branching, + level, centers, radiuses, belongs_to, count); + } + + /** * The method responsible with actually doing the recursive hierarchical * clustering diff --git a/modules/flann/src/miniflann.cpp b/modules/flann/src/miniflann.cpp index 5d2d655feb..b56578c17f 100644 --- a/modules/flann/src/miniflann.cpp +++ b/modules/flann/src/miniflann.cpp @@ -366,6 +366,7 @@ typedef ::cvflann::Hamming HammingDistance; #else typedef ::cvflann::HammingLUT HammingDistance; #endif +typedef ::cvflann::DNAmming2 DNAmmingDistance; Index::Index() { @@ -418,6 +419,9 @@ void Index::build(InputArray _data, const IndexParams& params, flann_distance_t buildIndex< ::cvflann::L1 >(index, data, params); break; #if MINIFLANN_SUPPORT_EXOTIC_DISTANCE_TYPES + case FLANN_DIST_DNAMMING: + buildIndex< DNAmmingDistance >(index, data, params); + break; case FLANN_DIST_MAX: buildIndex< ::cvflann::MaxDistance >(index, data, params); break; @@ -473,6 +477,9 @@ void Index::release() deleteIndex< ::cvflann::L1 >(index); break; #if MINIFLANN_SUPPORT_EXOTIC_DISTANCE_TYPES + case FLANN_DIST_DNAMMING: + deleteIndex< DNAmmingDistance >(index); + break; case FLANN_DIST_MAX: deleteIndex< ::cvflann::MaxDistance >(index); break; @@ -594,7 +601,8 @@ void Index::knnSearch(InputArray _query, OutputArray _indices, CV_INSTRUMENT_REGION(); Mat query = _query.getMat(), indices, dists; - int dtype = distType == FLANN_DIST_HAMMING ? CV_32S : CV_32F; + int dtype = (distType == FLANN_DIST_HAMMING) + || (distType == FLANN_DIST_DNAMMING) ? CV_32S : CV_32F; createIndicesDists( _indices, _dists, indices, dists, query.rows, knn, knn, dtype ); @@ -610,6 +618,9 @@ void Index::knnSearch(InputArray _query, OutputArray _indices, runKnnSearch< ::cvflann::L1 >(index, query, indices, dists, knn, params); break; #if MINIFLANN_SUPPORT_EXOTIC_DISTANCE_TYPES + case FLANN_DIST_DNAMMING: + runKnnSearch(index, query, indices, dists, knn, params); + break; case FLANN_DIST_MAX: runKnnSearch< ::cvflann::MaxDistance >(index, query, indices, dists, knn, params); break; @@ -638,7 +649,8 @@ int Index::radiusSearch(InputArray _query, OutputArray _indices, CV_INSTRUMENT_REGION(); Mat query = _query.getMat(), indices, dists; - int dtype = distType == FLANN_DIST_HAMMING ? CV_32S : CV_32F; + int dtype = (distType == FLANN_DIST_HAMMING) + || (distType == FLANN_DIST_DNAMMING) ? CV_32S : CV_32F; CV_Assert( maxResults > 0 ); createIndicesDists( _indices, _dists, indices, dists, query.rows, maxResults, INT_MAX, dtype ); @@ -655,6 +667,8 @@ int Index::radiusSearch(InputArray _query, OutputArray _indices, case FLANN_DIST_L1: return runRadiusSearch< ::cvflann::L1 >(index, query, indices, dists, radius, params); #if MINIFLANN_SUPPORT_EXOTIC_DISTANCE_TYPES + case FLANN_DIST_DNAMMING: + return runRadiusSearch< DNAmmingDistance >(index, query, indices, dists, radius, params); case FLANN_DIST_MAX: return runRadiusSearch< ::cvflann::MaxDistance >(index, query, indices, dists, radius, params); case FLANN_DIST_HIST_INTERSECT: @@ -718,6 +732,9 @@ void Index::save(const String& filename) const saveIndex< ::cvflann::L1 >(this, index, fout); break; #if MINIFLANN_SUPPORT_EXOTIC_DISTANCE_TYPES + case FLANN_DIST_DNAMMING: + saveIndex< DNAmmingDistance >(this, index, fout); + break; case FLANN_DIST_MAX: saveIndex< ::cvflann::MaxDistance >(this, index, fout); break; @@ -799,6 +816,7 @@ bool Index::load(InputArray _data, const String& filename) distType = (flann_distance_t)idistType; if( !((distType == FLANN_DIST_HAMMING && featureType == CV_8U) || + (distType == FLANN_DIST_DNAMMING && featureType == CV_8U) || (distType != FLANN_DIST_HAMMING && featureType == CV_32F)) ) { fprintf(stderr, "Reading FLANN index error: unsupported feature type %d for the index type %d\n", featureType, algo); @@ -818,6 +836,9 @@ bool Index::load(InputArray _data, const String& filename) loadIndex< ::cvflann::L1 >(this, index, data, fin); break; #if MINIFLANN_SUPPORT_EXOTIC_DISTANCE_TYPES + case FLANN_DIST_DNAMMING: + loadIndex< DNAmmingDistance >(this, index, data, fin); + break; case FLANN_DIST_MAX: loadIndex< ::cvflann::MaxDistance >(this, index, data, fin); break; diff --git a/modules/imgproc/include/opencv2/imgproc.hpp b/modules/imgproc/include/opencv2/imgproc.hpp index 156d3498cc..4a41ba6b50 100644 --- a/modules/imgproc/include/opencv2/imgproc.hpp +++ b/modules/imgproc/include/opencv2/imgproc.hpp @@ -256,6 +256,9 @@ enum InterpolationFlags{ INTER_LANCZOS4 = 4, /** Bit exact bilinear interpolation */ INTER_LINEAR_EXACT = 5, + /** Bit exact nearest neighbor interpolation. This will produce same results as + the nearest neighbor method in PIL, scikit-image or Matlab. */ + INTER_NEAREST_EXACT = 6, /** mask for interpolation codes */ INTER_MAX = 7, /** flag, fills all of the destination image pixels. If some of them correspond to outliers in the diff --git a/modules/imgproc/perf/perf_resize.cpp b/modules/imgproc/perf/perf_resize.cpp index 236955bd66..a0ea0804cc 100644 --- a/modules/imgproc/perf/perf_resize.cpp +++ b/modules/imgproc/perf/perf_resize.cpp @@ -254,4 +254,30 @@ PERF_TEST_P(MatInfo_Size_Scale_NN, ResizeNN, SANITY_CHECK_NOTHING(); } +PERF_TEST_P(MatInfo_Size_Scale_NN, ResizeNNExact, + testing::Combine( + testing::Values(CV_8UC1, CV_8UC3, CV_8UC4), + testing::Values(sz720p, sz1080p), + testing::Values(0.25, 0.5, 2.0) + ) +) +{ + int matType = get<0>(GetParam()); + Size from = get<1>(GetParam()); + double scale = get<2>(GetParam()); + + cv::Mat src(from, matType); + + Size to(cvRound(from.width * scale), cvRound(from.height * scale)); + cv::Mat dst(to, matType); + + declare.in(src, WARMUP_RNG).out(dst); + declare.time(100); + + TEST_CYCLE() resize(src, dst, dst.size(), 0, 0, INTER_NEAREST_EXACT); + + EXPECT_GT(countNonZero(dst.reshape(1)), 0); + SANITY_CHECK_NOTHING(); +} + } // namespace diff --git a/modules/imgproc/src/fixedpoint.inl.hpp b/modules/imgproc/src/fixedpoint.inl.hpp index 21a5368526..f5f433fec6 100644 --- a/modules/imgproc/src/fixedpoint.inl.hpp +++ b/modules/imgproc/src/fixedpoint.inl.hpp @@ -14,13 +14,14 @@ namespace { class fixedpoint64 { private: - static const int fixedShift = 32; - int64_t val; fixedpoint64(int64_t _val) : val(_val) {} static CV_ALWAYS_INLINE uint64_t fixedround(const uint64_t& _val) { return (_val + ((1LL << fixedShift) >> 1)); } public: + static const int fixedShift = 32; + typedef fixedpoint64 WT; + typedef int64_t raw_t; CV_ALWAYS_INLINE fixedpoint64() { val = 0; } CV_ALWAYS_INLINE fixedpoint64(const fixedpoint64& v) { val = v.val; } CV_ALWAYS_INLINE fixedpoint64(const int8_t& _val) { val = ((int64_t)_val) << fixedShift; } @@ -97,13 +98,14 @@ public: class ufixedpoint64 { private: - static const int fixedShift = 32; - uint64_t val; ufixedpoint64(uint64_t _val) : val(_val) {} static CV_ALWAYS_INLINE uint64_t fixedround(const uint64_t& _val) { return (_val + ((1LL << fixedShift) >> 1)); } public: + static const int fixedShift = 32; + typedef ufixedpoint64 WT; + typedef uint64_t raw_t; CV_ALWAYS_INLINE ufixedpoint64() { val = 0; } CV_ALWAYS_INLINE ufixedpoint64(const ufixedpoint64& v) { val = v.val; } CV_ALWAYS_INLINE ufixedpoint64(const uint8_t& _val) { val = ((uint64_t)_val) << fixedShift; } @@ -157,19 +159,24 @@ public: CV_ALWAYS_INLINE bool isZero() { return val == 0; } static CV_ALWAYS_INLINE ufixedpoint64 zero() { return ufixedpoint64(); } static CV_ALWAYS_INLINE ufixedpoint64 one() { return ufixedpoint64((uint64_t)(1ULL << fixedShift)); } + + static CV_ALWAYS_INLINE ufixedpoint64 fromRaw(uint64_t v) { return ufixedpoint64(v); } + CV_ALWAYS_INLINE uint64_t raw() { return val; } + CV_ALWAYS_INLINE uint32_t cvFloor() const { return cv::saturate_cast(val >> fixedShift); } friend class ufixedpoint32; }; class fixedpoint32 { private: - static const int fixedShift = 16; - int32_t val; fixedpoint32(int32_t _val) : val(_val) {} static CV_ALWAYS_INLINE uint32_t fixedround(const uint32_t& _val) { return (_val + ((1 << fixedShift) >> 1)); } public: + static const int fixedShift = 16; + typedef fixedpoint64 WT; + typedef int32_t raw_t; CV_ALWAYS_INLINE fixedpoint32() { val = 0; } CV_ALWAYS_INLINE fixedpoint32(const fixedpoint32& v) { val = v.val; } CV_ALWAYS_INLINE fixedpoint32(const int8_t& _val) { val = ((int32_t)_val) << fixedShift; } @@ -217,13 +224,14 @@ public: class ufixedpoint32 { private: - static const int fixedShift = 16; - uint32_t val; ufixedpoint32(uint32_t _val) : val(_val) {} static CV_ALWAYS_INLINE uint32_t fixedround(const uint32_t& _val) { return (_val + ((1 << fixedShift) >> 1)); } public: + static const int fixedShift = 16; + typedef ufixedpoint64 WT; + typedef uint32_t raw_t; CV_ALWAYS_INLINE ufixedpoint32() { val = 0; } CV_ALWAYS_INLINE ufixedpoint32(const ufixedpoint32& v) { val = v.val; } CV_ALWAYS_INLINE ufixedpoint32(const uint8_t& _val) { val = ((uint32_t)_val) << fixedShift; } @@ -261,19 +269,23 @@ public: CV_ALWAYS_INLINE bool isZero() { return val == 0; } static CV_ALWAYS_INLINE ufixedpoint32 zero() { return ufixedpoint32(); } static CV_ALWAYS_INLINE ufixedpoint32 one() { return ufixedpoint32((1U << fixedShift)); } + + static CV_ALWAYS_INLINE ufixedpoint32 fromRaw(uint32_t v) { return ufixedpoint32(v); } + CV_ALWAYS_INLINE uint32_t raw() { return val; } friend class ufixedpoint16; }; class fixedpoint16 { private: - static const int fixedShift = 8; - int16_t val; fixedpoint16(int16_t _val) : val(_val) {} static CV_ALWAYS_INLINE uint16_t fixedround(const uint16_t& _val) { return (_val + ((1 << fixedShift) >> 1)); } public: + static const int fixedShift = 8; + typedef fixedpoint32 WT; + typedef int16_t raw_t; CV_ALWAYS_INLINE fixedpoint16() { val = 0; } CV_ALWAYS_INLINE fixedpoint16(const fixedpoint16& v) { val = v.val; } CV_ALWAYS_INLINE fixedpoint16(const int8_t& _val) { val = ((int16_t)_val) << fixedShift; } @@ -314,13 +326,14 @@ public: class ufixedpoint16 { private: - static const int fixedShift = 8; - uint16_t val; ufixedpoint16(uint16_t _val) : val(_val) {} static CV_ALWAYS_INLINE uint16_t fixedround(const uint16_t& _val) { return (_val + ((1 << fixedShift) >> 1)); } public: + static const int fixedShift = 8; + typedef ufixedpoint32 WT; + typedef uint16_t raw_t; CV_ALWAYS_INLINE ufixedpoint16() { val = 0; } CV_ALWAYS_INLINE ufixedpoint16(const ufixedpoint16& v) { val = v.val; } CV_ALWAYS_INLINE ufixedpoint16(const uint8_t& _val) { val = ((uint16_t)_val) << fixedShift; } @@ -357,7 +370,7 @@ public: static CV_ALWAYS_INLINE ufixedpoint16 one() { return ufixedpoint16((uint16_t)(1 << fixedShift)); } static CV_ALWAYS_INLINE ufixedpoint16 fromRaw(uint16_t v) { return ufixedpoint16(v); } - CV_ALWAYS_INLINE ufixedpoint16 raw() { return val; } + CV_ALWAYS_INLINE uint16_t raw() { return val; } }; } diff --git a/modules/imgproc/src/resize.cpp b/modules/imgproc/src/resize.cpp index 2f060b6ae0..4f82bddfa0 100644 --- a/modules/imgproc/src/resize.cpp +++ b/modules/imgproc/src/resize.cpp @@ -51,6 +51,7 @@ #include "opencl_kernels_imgproc.hpp" #include "hal_replacement.hpp" #include "opencv2/core/hal/intrin.hpp" +#include "opencv2/core/utils/buffer_area.private.hpp" #include "opencv2/core/openvx/ovx_defs.hpp" #include "resize.hpp" @@ -1104,6 +1105,121 @@ resizeNN( const Mat& src, Mat& dst, double fx, double fy ) } } +class resizeNN_bitexactInvoker : public ParallelLoopBody +{ +public: + resizeNN_bitexactInvoker(const Mat& _src, Mat& _dst, int* _x_ofse, int _ify, int _ify0) + : src(_src), dst(_dst), x_ofse(_x_ofse), ify(_ify), ify0(_ify0) {} + + virtual void operator() (const Range& range) const CV_OVERRIDE + { + Size ssize = src.size(), dsize = dst.size(); + int pix_size = (int)src.elemSize(); + for( int y = range.start; y < range.end; y++ ) + { + uchar* D = dst.ptr(y); + int _sy = (ify * y + ify0) >> 16; + int sy = std::min(_sy, ssize.height-1); + const uchar* S = src.ptr(sy); + + int x = 0; + switch( pix_size ) + { + case 1: +#if CV_SIMD + for( ; x <= dsize.width - v_uint8::nlanes; x += v_uint8::nlanes ) + v_store(D + x, vx_lut(S, x_ofse + x)); +#endif + for( ; x < dsize.width; x++ ) + D[x] = S[x_ofse[x]]; + break; + case 2: +#if CV_SIMD + for( ; x <= dsize.width - v_uint16::nlanes; x += v_uint16::nlanes ) + v_store((ushort*)D + x, vx_lut((ushort*)S, x_ofse + x)); +#endif + for( ; x < dsize.width; x++ ) + *((ushort*)D + x) = *((ushort*)S + x_ofse[x]); + break; + case 3: + for( ; x < dsize.width; x++, D += 3 ) + { + const uchar* _tS = S + x_ofse[x] * 3; + D[0] = _tS[0]; D[1] = _tS[1]; D[2] = _tS[2]; + } + break; + case 4: +#if CV_SIMD + for( ; x <= dsize.width - v_uint32::nlanes; x += v_uint32::nlanes ) + v_store((uint32_t*)D + x, vx_lut((uint32_t*)S, x_ofse + x)); +#endif + for( ; x < dsize.width; x++ ) + *((uint32_t*)D + x) = *((uint32_t*)S + x_ofse[x]); + break; + case 6: + for( ; x < dsize.width; x++, D += 6 ) + { + const ushort* _tS = (const ushort*)(S + x_ofse[x]*6); + ushort* _tD = (ushort*)D; + _tD[0] = _tS[0]; _tD[1] = _tS[1]; _tD[2] = _tS[2]; + } + break; + case 8: +#if CV_SIMD + for( ; x <= dsize.width - v_uint64::nlanes; x += v_uint64::nlanes ) + v_store((uint64_t*)D + x, vx_lut((uint64_t*)S, x_ofse + x)); +#endif + for( ; x < dsize.width; x++ ) + *((uint64_t*)D + x) = *((uint64_t*)S + x_ofse[x]); + break; + case 12: + for( ; x < dsize.width; x++, D += 12 ) + { + const int* _tS = (const int*)(S + x_ofse[x]*12); + int* _tD = (int*)D; + _tD[0] = _tS[0]; _tD[1] = _tS[1]; _tD[2] = _tS[2]; + } + break; + default: + for( x = 0; x < dsize.width; x++, D += pix_size ) + { + const uchar* _tS = S + x_ofse[x] * pix_size; + for (int k = 0; k < pix_size; k++) + D[k] = _tS[k]; + } + } + } + } +private: + const Mat& src; + Mat& dst; + int* x_ofse; + const int ify; + const int ify0; +}; + +static void resizeNN_bitexact( const Mat& src, Mat& dst, double /*fx*/, double /*fy*/ ) +{ + Size ssize = src.size(), dsize = dst.size(); + int ifx = ((ssize.width << 16) + dsize.width / 2) / dsize.width; // 16bit fixed-point arithmetic + int ifx0 = ifx / 2 - 1; // This method uses center pixel coordinate as Pillow and scikit-images do. + int ify = ((ssize.height << 16) + dsize.height / 2) / dsize.height; + int ify0 = ify / 2 - 1; + + cv::utils::BufferArea area; + int* x_ofse = 0; + area.allocate(x_ofse, dsize.width, CV_SIMD_WIDTH); + area.commit(); + + for( int x = 0; x < dsize.width; x++ ) + { + int sx = (ifx * x + ifx0) >> 16; + x_ofse[x] = std::min(sx, ssize.width-1); // offset in element (not byte) + } + Range range(0, dsize.height); + resizeNN_bitexactInvoker invoker(src, dst, x_ofse, ify, ify0); + parallel_for_(range, invoker, dst.total()/(double)(1<<16)); +} struct VResizeNoVec { @@ -3723,6 +3839,12 @@ void resize(int src_type, return; } + if( interpolation == INTER_NEAREST_EXACT ) + { + resizeNN_bitexact( src, dst, inv_scale_x, inv_scale_y ); + return; + } + int k, sx, sy, dx, dy; diff --git a/modules/imgproc/src/smooth.dispatch.cpp b/modules/imgproc/src/smooth.dispatch.cpp index 65d1fc8ed6..69d07580f2 100644 --- a/modules/imgproc/src/smooth.dispatch.cpp +++ b/modules/imgproc/src/smooth.dispatch.cpp @@ -258,23 +258,20 @@ softdouble getGaussianKernelFixedPoint_ED(CV_OUT std::vector& result, c } static void getGaussianKernel(int n, double sigma, int ktype, Mat& res) { res = getGaussianKernel(n, sigma, ktype); } -template static void getGaussianKernel(int n, double sigma, int, std::vector& res); -//{ res = getFixedpointGaussianKernel(n, sigma); } - -template<> void getGaussianKernel(int n, double sigma, int, std::vector& res) +template static void getGaussianKernel(int n, double sigma, int, std::vector& res) { std::vector res_sd; softdouble s0 = getGaussianKernelBitExact(res_sd, n, sigma); CV_UNUSED(s0); std::vector fixed_256; - softdouble approx_err = getGaussianKernelFixedPoint_ED(fixed_256, res_sd, 8); + softdouble approx_err = getGaussianKernelFixedPoint_ED(fixed_256, res_sd, FT::fixedShift); CV_UNUSED(approx_err); res.resize(n); for (int i = 0; i < n; i++) { - res[i] = ufixedpoint16::fromRaw((uint16_t)fixed_256[i]); + res[i] = FT::fromRaw((typename FT::raw_t)fixed_256[i]); //printf("%03d: %d\n", i, res[i].raw()); } } @@ -688,6 +685,43 @@ void GaussianBlur(InputArray _src, OutputArray _dst, Size ksize, return; } } + if(sdepth == CV_16U && ((borderType & BORDER_ISOLATED) || !_src.isSubmatrix())) + { + CV_LOG_INFO(NULL, "GaussianBlur: running bit-exact version..."); + + std::vector fkx, fky; + createGaussianKernels(fkx, fky, type, ksize, sigma1, sigma2); + + static bool param_check_gaussian_blur_bitexact_kernels = utils::getConfigurationParameterBool("OPENCV_GAUSSIANBLUR_CHECK_BITEXACT_KERNELS", false); + if (param_check_gaussian_blur_bitexact_kernels && !validateGaussianBlurKernel(fkx)) + { + CV_LOG_INFO(NULL, "GaussianBlur: bit-exact fx kernel can't be applied: ksize=" << ksize << " sigma=" << Size2d(sigma1, sigma2)); + } + else if (param_check_gaussian_blur_bitexact_kernels && !validateGaussianBlurKernel(fky)) + { + CV_LOG_INFO(NULL, "GaussianBlur: bit-exact fy kernel can't be applied: ksize=" << ksize << " sigma=" << Size2d(sigma1, sigma2)); + } + else + { + // TODO: implement ocl_sepFilter2D_BitExact -- how to deal with bdepth? + // CV_OCL_RUN(useOpenCL, + // ocl_sepFilter2D_BitExact(_src, _dst, sdepth, + // ksize, + // (const uint32_t*)&fkx[0], (const uint32_t*)&fky[0], + // Point(-1, -1), 0, borderType, + // 16/*shift_bits*/) + // ); + + Mat src = _src.getMat(); + Mat dst = _dst.getMat(); + + if (src.data == dst.data) + src = src.clone(); + CV_CPU_DISPATCH(GaussianBlurFixedPoint, (src, dst, (const uint32_t*)&fkx[0], (int)fkx.size(), (const uint32_t*)&fky[0], (int)fky.size(), borderType), + CV_CPU_DISPATCH_MODES_ALL); + return; + } + } #ifdef HAVE_OPENCL if (useOpenCL) diff --git a/modules/imgproc/src/smooth.simd.hpp b/modules/imgproc/src/smooth.simd.hpp index 3102b36f74..2a7a8e72bb 100644 --- a/modules/imgproc/src/smooth.simd.hpp +++ b/modules/imgproc/src/smooth.simd.hpp @@ -54,9 +54,10 @@ namespace cv { CV_CPU_OPTIMIZATION_NAMESPACE_BEGIN // forward declarations -void GaussianBlurFixedPoint(const Mat& src, /*const*/ Mat& dst, - const uint16_t/*ufixedpoint16*/* fkx, int fkx_size, - const uint16_t/*ufixedpoint16*/* fky, int fky_size, +template +void GaussianBlurFixedPoint(const Mat& src, Mat& dst, + const RFT* fkx, int fkx_size, + const RFT* fky, int fky_size, int borderType); #ifndef CV_CPU_OPTIMIZATION_DECLARATIONS_ONLY @@ -192,8 +193,9 @@ void hlineSmooth3N(const uint8_t* src, int cn, const ufi } } } -template -void hlineSmooth3N121(const ET* src, int cn, const FT*, int, FT* dst, int len, int borderType) + +template +void hlineSmooth3N121Impl(const ET* src, int cn, const FT*, int, FT* dst, int len, int borderType) { if (len == 1) { @@ -217,7 +219,13 @@ void hlineSmooth3N121(const ET* src, int cn, const FT*, int, FT* dst, int len, i } src += cn; dst += cn; - for (int i = cn; i < (len - 1)*cn; i++, src++, dst++) + int i = cn, lencn = (len - 1)*cn; +#if CV_SIMD + const int VECSZ = VFT::nlanes; + for (; i <= lencn - VECSZ; i += VECSZ, src += VECSZ, dst += VECSZ) + v_store((typename FT::raw_t*)dst, (vx_load_expand(src - cn) + vx_load_expand(src + cn) + (vx_load_expand(src) << 1)) << (FT::fixedShift-2)); +#endif + for (; i < lencn; i++, src++, dst++) *dst = (FT(src[-cn])>>2) + (FT(src[cn])>>2) + (FT(src[0])>>1); // Point that fall right from border @@ -231,51 +239,19 @@ void hlineSmooth3N121(const ET* src, int cn, const FT*, int, FT* dst, int len, i } } } +template +void hlineSmooth3N121(const ET* src, int cn, const FT*, int, FT* dst, int len, int borderType); template <> -void hlineSmooth3N121(const uint8_t* src, int cn, const ufixedpoint16*, int, ufixedpoint16* dst, int len, int borderType) +void hlineSmooth3N121(const uint8_t* src, int cn, const ufixedpoint16* _m, int _n, ufixedpoint16* dst, int len, int borderType) { - if (len == 1) - { - if (borderType != BORDER_CONSTANT) - for (int k = 0; k < cn; k++) - dst[k] = ufixedpoint16(src[k]); - else - for (int k = 0; k < cn; k++) - dst[k] = ufixedpoint16(src[k]) >> 1; - } - else - { - // Point that fall left from border - for (int k = 0; k < cn; k++) - dst[k] = (ufixedpoint16(src[k])>>1) + (ufixedpoint16(src[cn + k])>>2); - if (borderType != BORDER_CONSTANT)// If BORDER_CONSTANT out of border values are equal to zero and could be skipped - { - int src_idx = borderInterpolate(-1, len, borderType); - for (int k = 0; k < cn; k++) - dst[k] = dst[k] + (ufixedpoint16(src[src_idx*cn + k])>>2); - } - - src += cn; dst += cn; - int i = cn, lencn = (len - 1)*cn; -#if CV_SIMD - const int VECSZ = v_uint16::nlanes; - for (; i <= lencn - VECSZ; i += VECSZ, src += VECSZ, dst += VECSZ) - v_store((uint16_t*)dst, (vx_load_expand(src - cn) + vx_load_expand(src + cn) + (vx_load_expand(src) << 1)) << 6); -#endif - for (; i < lencn; i++, src++, dst++) - *((uint16_t*)dst) = (uint16_t(src[-cn]) + uint16_t(src[cn]) + (uint16_t(src[0]) << 1)) << 6; - - // Point that fall right from border - for (int k = 0; k < cn; k++) - dst[k] = (ufixedpoint16(src[k - cn])>>2) + (ufixedpoint16(src[k])>>1); - if (borderType != BORDER_CONSTANT)// If BORDER_CONSTANT out of border values are equal to zero and could be skipped - { - int src_idx = (borderInterpolate(len, len, borderType) - (len - 1))*cn; - for (int k = 0; k < cn; k++) - dst[k] = dst[k] + (ufixedpoint16(src[src_idx + k])>>2); - } - } + hlineSmooth3N121Impl(src, cn, _m, _n, dst, len, borderType); } +template <> +void hlineSmooth3N121(const uint16_t* src, int cn, const ufixedpoint32* _m, int _n, ufixedpoint32* dst, int len, int borderType) +{ + hlineSmooth3N121Impl(src, cn, _m, _n, dst, len, borderType); +} + template void hlineSmooth3Naba(const ET* src, int cn, const FT* m, int, FT* dst, int len, int borderType) { @@ -1376,6 +1352,28 @@ void vlineSmooth3N121(const ufixedpoint16* const * src, for (; i < len; i++) dst[i] = (((uint32_t)(((uint16_t*)(src[0]))[i]) + (uint32_t)(((uint16_t*)(src[2]))[i]) + ((uint32_t)(((uint16_t*)(src[1]))[i]) << 1)) + (1 << 9)) >> 10; } +template <> +void vlineSmooth3N121(const ufixedpoint32* const * src, const ufixedpoint32*, int, uint16_t* dst, int len) +{ + int i = 0; +#if CV_SIMD + const int VECSZ = v_uint32::nlanes; + for (; i <= len - 2*VECSZ; i += 2*VECSZ) + { + v_uint64 v_src00, v_src01, v_src02, v_src03, v_src10, v_src11, v_src12, v_src13, v_src20, v_src21, v_src22, v_src23; + v_expand(vx_load((uint32_t*)(src[0]) + i), v_src00, v_src01); + v_expand(vx_load((uint32_t*)(src[0]) + i + VECSZ), v_src02, v_src03); + v_expand(vx_load((uint32_t*)(src[1]) + i), v_src10, v_src11); + v_expand(vx_load((uint32_t*)(src[1]) + i + VECSZ), v_src12, v_src13); + v_expand(vx_load((uint32_t*)(src[2]) + i), v_src20, v_src21); + v_expand(vx_load((uint32_t*)(src[2]) + i + VECSZ), v_src22, v_src23); + v_store(dst + i, v_pack(v_rshr_pack<18>(v_src00 + v_src20 + (v_src10 + v_src10), v_src01 + v_src21 + (v_src11 + v_src11)), + v_rshr_pack<18>(v_src02 + v_src22 + (v_src12 + v_src12), v_src03 + v_src23 + (v_src13 + v_src13)))); + } +#endif + for (; i < len; i++) + dst[i] = (((uint64_t)((uint32_t*)(src[0]))[i]) + (uint64_t)(((uint32_t*)(src[2]))[i]) + ((uint64_t(((uint32_t*)(src[1]))[i]) << 1)) + (1 << 17)) >> 18; +} template void vlineSmooth5N(const FT* const * src, const FT* m, int, ET* dst, int len) { @@ -1525,6 +1523,39 @@ void vlineSmooth5N14641(const ufixedpoint16* const * src (((uint32_t)(((uint16_t*)(src[1]))[i]) + (uint32_t)(((uint16_t*)(src[3]))[i])) << 2) + (uint32_t)(((uint16_t*)(src[0]))[i]) + (uint32_t)(((uint16_t*)(src[4]))[i]) + (1 << 11)) >> 12; } +template <> +void vlineSmooth5N14641(const ufixedpoint32* const * src, const ufixedpoint32*, int, uint16_t* dst, int len) +{ + int i = 0; +#if CV_SIMD + const int VECSZ = v_uint32::nlanes; + for (; i <= len - 2*VECSZ; i += 2*VECSZ) + { + v_uint64 v_src00, v_src10, v_src20, v_src30, v_src40; + v_uint64 v_src01, v_src11, v_src21, v_src31, v_src41; + v_uint64 v_src02, v_src12, v_src22, v_src32, v_src42; + v_uint64 v_src03, v_src13, v_src23, v_src33, v_src43; + v_expand(vx_load((uint32_t*)(src[0]) + i), v_src00, v_src01); + v_expand(vx_load((uint32_t*)(src[0]) + i + VECSZ), v_src02, v_src03); + v_expand(vx_load((uint32_t*)(src[1]) + i), v_src10, v_src11); + v_expand(vx_load((uint32_t*)(src[1]) + i + VECSZ), v_src12, v_src13); + v_expand(vx_load((uint32_t*)(src[2]) + i), v_src20, v_src21); + v_expand(vx_load((uint32_t*)(src[2]) + i + VECSZ), v_src22, v_src23); + v_expand(vx_load((uint32_t*)(src[3]) + i), v_src30, v_src31); + v_expand(vx_load((uint32_t*)(src[3]) + i + VECSZ), v_src32, v_src33); + v_expand(vx_load((uint32_t*)(src[4]) + i), v_src40, v_src41); + v_expand(vx_load((uint32_t*)(src[4]) + i + VECSZ), v_src42, v_src43); + v_store(dst + i, v_pack(v_rshr_pack<20>((v_src20 << 2) + (v_src20 << 1) + ((v_src10 + v_src30) << 2) + v_src00 + v_src40, + (v_src21 << 2) + (v_src21 << 1) + ((v_src11 + v_src31) << 2) + v_src01 + v_src41), + v_rshr_pack<20>((v_src22 << 2) + (v_src22 << 1) + ((v_src12 + v_src32) << 2) + v_src02 + v_src42, + (v_src23 << 2) + (v_src23 << 1) + ((v_src13 + v_src33) << 2) + v_src03 + v_src43))); + } +#endif + for (; i < len; i++) + dst[i] = ((uint64_t)(((uint32_t*)(src[2]))[i]) * 6 + + (((uint64_t)(((uint32_t*)(src[1]))[i]) + (uint64_t)(((uint32_t*)(src[3]))[i])) << 2) + + (uint64_t)(((uint32_t*)(src[0]))[i]) + (uint64_t)(((uint32_t*)(src[4]))[i]) + (1 << 19)) >> 20; +} template void vlineSmooth(const FT* const * src, const FT* m, int n, ET* dst, int len) { @@ -2029,25 +2060,42 @@ private: } // namespace anon -void GaussianBlurFixedPoint(const Mat& src, /*const*/ Mat& dst, - const uint16_t/*ufixedpoint16*/* fkx, int fkx_size, - const uint16_t/*ufixedpoint16*/* fky, int fky_size, - int borderType) +template +void GaussianBlurFixedPointImpl(const Mat& src, /*const*/ Mat& dst, + const RFT* fkx, int fkx_size, + const RFT* fky, int fky_size, + int borderType) { CV_INSTRUMENT_REGION(); - CV_Assert(src.depth() == CV_8U && ((borderType & BORDER_ISOLATED) || !src.isSubmatrix())); - fixedSmoothInvoker invoker( - src.ptr(), src.step1(), - dst.ptr(), dst.step1(), dst.cols, dst.rows, dst.channels(), - (const ufixedpoint16*)fkx, fkx_size, (const ufixedpoint16*)fky, fky_size, + CV_Assert(src.depth() == DataType::depth && ((borderType & BORDER_ISOLATED) || !src.isSubmatrix())); + fixedSmoothInvoker invoker( + src.ptr(), src.step1(), + dst.ptr(), dst.step1(), dst.cols, dst.rows, dst.channels(), + (const FT*)fkx, fkx_size, (const FT*)fky, fky_size, borderType & ~BORDER_ISOLATED); { // TODO AVX guard (external call) parallel_for_(Range(0, dst.rows), invoker, std::max(1, std::min(getNumThreads(), getNumberOfCPUs()))); } } +template <> +void GaussianBlurFixedPoint(const Mat& src, /*const*/ Mat& dst, + const uint16_t/*ufixedpoint16*/* fkx, int fkx_size, + const uint16_t/*ufixedpoint16*/* fky, int fky_size, + int borderType) +{ + GaussianBlurFixedPointImpl(src, dst, fkx, fkx_size, fky, fky_size, borderType); +} +template <> +void GaussianBlurFixedPoint(const Mat& src, /*const*/ Mat& dst, + const uint32_t/*ufixedpoint32*/* fkx, int fkx_size, + const uint32_t/*ufixedpoint32*/* fky, int fky_size, + int borderType) +{ + GaussianBlurFixedPointImpl(src, dst, fkx, fkx_size, fky, fky_size, borderType); +} #endif CV_CPU_OPTIMIZATION_NAMESPACE_END } // namespace diff --git a/modules/imgproc/test/test_imgwarp.cpp b/modules/imgproc/test/test_imgwarp.cpp index 874cd6b2b5..024c4f33fa 100644 --- a/modules/imgproc/test/test_imgwarp.cpp +++ b/modules/imgproc/test/test_imgwarp.cpp @@ -346,14 +346,24 @@ protected: CV_ResizeExactTest::CV_ResizeExactTest() : CV_ResizeTest() { - max_interpolation = 1; + max_interpolation = 2; } void CV_ResizeExactTest::get_test_array_types_and_sizes(int test_case_idx, vector >& sizes, vector >& types) { CV_ResizeTest::get_test_array_types_and_sizes(test_case_idx, sizes, types); - interpolation = INTER_LINEAR_EXACT; + switch (interpolation) + { + case 0: + interpolation = INTER_LINEAR_EXACT; + break; + case 1: + interpolation = INTER_NEAREST_EXACT; + break; + default: + CV_Assert(interpolation < max_interpolation); + } if (CV_MAT_DEPTH(types[INPUT][0]) == CV_32F || CV_MAT_DEPTH(types[INPUT][0]) == CV_64F) types[INPUT][0] = types[INPUT_OUTPUT][0] = types[REF_INPUT_OUTPUT][0] = CV_MAKETYPE(CV_8U, CV_MAT_CN(types[INPUT][0])); diff --git a/modules/imgproc/test/test_resize_bitexact.cpp b/modules/imgproc/test/test_resize_bitexact.cpp index f76eb6f9d2..78cad71d03 100644 --- a/modules/imgproc/test/test_resize_bitexact.cpp +++ b/modules/imgproc/test/test_resize_bitexact.cpp @@ -152,4 +152,89 @@ TEST(Resize_Bitexact, Linear8U) } } +PARAM_TEST_CASE(Resize_Bitexact, int) +{ +public: + int depth; + + virtual void SetUp() + { + depth = GET_PARAM(0); + } + + double CountDiff(const Mat& src) + { + Mat dstExact; cv::resize(src, dstExact, Size(), 2, 1, INTER_NEAREST_EXACT); + Mat dstNonExact; cv::resize(src, dstNonExact, Size(), 2, 1, INTER_NEAREST); + + return cv::norm(dstExact, dstNonExact, NORM_INF); + } +}; + +TEST_P(Resize_Bitexact, Nearest8U_vsNonExact) +{ + Mat mat_color, mat_gray; + Mat src_color = imread(cvtest::findDataFile("shared/lena.png")); + Mat src_gray; cv::cvtColor(src_color, src_gray, COLOR_BGR2GRAY); + src_color.convertTo(mat_color, depth); + src_gray.convertTo(mat_gray, depth); + + EXPECT_EQ(CountDiff(mat_color), 0) << "color, type: " << depth; + EXPECT_EQ(CountDiff(mat_gray), 0) << "gray, type: " << depth; +} + +// Now INTER_NEAREST's convention and INTER_NEAREST_EXACT's one are different. +INSTANTIATE_TEST_CASE_P(DISABLED_Imgproc, Resize_Bitexact, + testing::Values(CV_8U, CV_16U, CV_32F, CV_64F) +); + +TEST(Resize_Bitexact, Nearest8U) +{ + Mat src[6], dst[6]; + + // 2x decimation + src[0] = (Mat_(1, 6) << 0, 1, 2, 3, 4, 5); + dst[0] = (Mat_(1, 3) << 0, 2, 4); + + // decimation odd to 1 + src[1] = (Mat_(1, 5) << 0, 1, 2, 3, 4); + dst[1] = (Mat_(1, 1) << 2); + + // decimation n*2-1 to n + src[2] = (Mat_(1, 5) << 0, 1, 2, 3, 4); + dst[2] = (Mat_(1, 3) << 0, 2, 4); + + // decimation n*2+1 to n + src[3] = (Mat_(1, 5) << 0, 1, 2, 3, 4); + dst[3] = (Mat_(1, 2) << 1, 3); + + // zoom + src[4] = (Mat_(3, 5) << + 0, 1, 2, 3, 4, + 5, 6, 7, 8, 9, + 10, 11, 12, 13, 14); + dst[4] = (Mat_(5, 7) << + 0, 1, 1, 2, 3, 3, 4, + 0, 1, 1, 2, 3, 3, 4, + 5, 6, 6, 7, 8, 8, 9, + 10, 11, 11, 12, 13, 13, 14, + 10, 11, 11, 12, 13, 13, 14); + + src[5] = (Mat_(2, 3) << + 0, 1, 2, + 3, 4, 5); + dst[5] = (Mat_(4, 6) << + 0, 0, 1, 1, 2, 2, + 0, 0, 1, 1, 2, 2, + 3, 3, 4, 4, 5, 5, + 3, 3, 4, 4, 5, 5); + + for (int i = 0; i < 6; i++) + { + Mat calc; + resize(src[i], calc, dst[i].size(), 0, 0, INTER_NEAREST_EXACT); + EXPECT_EQ(cvtest::norm(calc, dst[i], cv::NORM_L1), 0); + } +} + }} // namespace diff --git a/modules/imgproc/test/test_smooth_bitexact.cpp b/modules/imgproc/test/test_smooth_bitexact.cpp index 8151c48238..f446deb8d8 100644 --- a/modules/imgproc/test/test_smooth_bitexact.cpp +++ b/modules/imgproc/test/test_smooth_bitexact.cpp @@ -7,13 +7,15 @@ namespace opencv_test { namespace { static const int fixedShiftU8 = 8; - static const int64_t fixedOne = (1L << fixedShiftU8); + static const int64_t fixedOneU8 = (1L << fixedShiftU8); + static const int fixedShiftU16 = 16; + static const int64_t fixedOneU16 = (1L << fixedShiftU16); - int64_t v[][9] = { - { fixedOne }, // size 1, sigma 0 - { fixedOne >> 2, fixedOne >> 1, fixedOne >> 2 }, // size 3, sigma 0 - { fixedOne >> 4, fixedOne >> 2, 6 * (fixedOne >> 4), fixedOne >> 2, fixedOne >> 4 }, // size 5, sigma 0 - { fixedOne >> 5, 7 * (fixedOne >> 6), 7 * (fixedOne >> 5), 9 * (fixedOne >> 5), 7 * (fixedOne >> 5), 7 * (fixedOne >> 6), fixedOne >> 5 }, // size 7, sigma 0 + int64_t vU8[][9] = { + { fixedOneU8 }, // size 1, sigma 0 + { fixedOneU8 >> 2, fixedOneU8 >> 1, fixedOneU8 >> 2 }, // size 3, sigma 0 + { fixedOneU8 >> 4, fixedOneU8 >> 2, 6 * (fixedOneU8 >> 4), fixedOneU8 >> 2, fixedOneU8 >> 4 }, // size 5, sigma 0 + { fixedOneU8 >> 5, 7 * (fixedOneU8 >> 6), 7 * (fixedOneU8 >> 5), 9 * (fixedOneU8 >> 5), 7 * (fixedOneU8 >> 5), 7 * (fixedOneU8 >> 6), fixedOneU8 >> 5 }, // size 7, sigma 0 { 4, 13, 30, 51, 60, 51, 30, 13, 4 }, // size 9, sigma 0 #if 1 #define CV_TEST_INACCURATE_GAUSSIAN_BLUR @@ -24,6 +26,14 @@ namespace opencv_test { namespace { #endif }; + int64_t vU16[][9] = { + { fixedOneU16 }, // size 1, sigma 0 + { fixedOneU16 >> 2, fixedOneU16 >> 1, fixedOneU16 >> 2 }, // size 3, sigma 0 + { fixedOneU16 >> 4, fixedOneU16 >> 2, 6 * (fixedOneU16 >> 4), fixedOneU16 >> 2, fixedOneU16 >> 4 }, // size 5, sigma 0 + { fixedOneU16 >> 5, 7 * (fixedOneU16 >> 6), 7 * (fixedOneU16 >> 5), 9 * (fixedOneU16 >> 5), 7 * (fixedOneU16 >> 5), 7 * (fixedOneU16 >> 6), fixedOneU16 >> 5 }, // size 7, sigma 0 + { 4<<8, 13<<8, 30<<8, 51<<8, 60<<8, 51<<8, 30<<8, 13<<8, 4<<8 } // size 9, sigma 0 + }; + template T eval(Mat src, vector kernelx, vector kernely) { @@ -39,8 +49,6 @@ namespace opencv_test { namespace { return saturate_cast((val + fixedRound) >> (fixedShift * 2)); } -TEST(GaussianBlur_Bitexact, Linear8U) -{ struct testmode { int type; @@ -50,34 +58,6 @@ TEST(GaussianBlur_Bitexact, Linear8U) double sigma_y; vector kernel_x; vector kernel_y; - } modes[] = { - { CV_8UC1, Size( 1, 1), Size(3, 3), 0, 0, vector(v[1], v[1]+3), vector(v[1], v[1]+3) }, - { CV_8UC1, Size( 2, 2), Size(3, 3), 0, 0, vector(v[1], v[1]+3), vector(v[1], v[1]+3) }, - { CV_8UC1, Size( 3, 1), Size(3, 3), 0, 0, vector(v[1], v[1]+3), vector(v[1], v[1]+3) }, - { CV_8UC1, Size( 1, 3), Size(3, 3), 0, 0, vector(v[1], v[1]+3), vector(v[1], v[1]+3) }, - { CV_8UC1, Size( 3, 3), Size(3, 3), 0, 0, vector(v[1], v[1]+3), vector(v[1], v[1]+3) }, - { CV_8UC1, Size( 3, 3), Size(5, 5), 0, 0, vector(v[2], v[2]+5), vector(v[2], v[2]+5) }, - { CV_8UC1, Size( 3, 3), Size(7, 7), 0, 0, vector(v[3], v[3]+7), vector(v[3], v[3]+7) }, - { CV_8UC1, Size( 5, 5), Size(3, 3), 0, 0, vector(v[1], v[1]+3), vector(v[1], v[1]+3) }, - { CV_8UC1, Size( 5, 5), Size(5, 5), 0, 0, vector(v[2], v[2]+5), vector(v[2], v[2]+5) }, - { CV_8UC1, Size( 3, 5), Size(5, 5), 0, 0, vector(v[2], v[2]+5), vector(v[2], v[2]+5) }, - { CV_8UC1, Size( 5, 5), Size(5, 5), 0, 0, vector(v[2], v[2]+5), vector(v[2], v[2]+5) }, - { CV_8UC1, Size( 5, 5), Size(7, 7), 0, 0, vector(v[3], v[3]+7), vector(v[3], v[3]+7) }, - { CV_8UC1, Size( 7, 7), Size(7, 7), 0, 0, vector(v[3], v[3]+7), vector(v[3], v[3]+7) }, - { CV_8UC1, Size( 256, 128), Size(3, 3), 0, 0, vector(v[1], v[1]+3), vector(v[1], v[1]+3) }, - { CV_8UC2, Size( 256, 128), Size(3, 3), 0, 0, vector(v[1], v[1]+3), vector(v[1], v[1]+3) }, - { CV_8UC3, Size( 256, 128), Size(3, 3), 0, 0, vector(v[1], v[1]+3), vector(v[1], v[1]+3) }, - { CV_8UC4, Size( 256, 128), Size(3, 3), 0, 0, vector(v[1], v[1]+3), vector(v[1], v[1]+3) }, - { CV_8UC1, Size( 256, 128), Size(5, 5), 0, 0, vector(v[2], v[2]+5), vector(v[2], v[2]+5) }, - { CV_8UC1, Size( 256, 128), Size(7, 7), 0, 0, vector(v[3], v[3]+7), vector(v[3], v[3]+7) }, - { CV_8UC1, Size( 256, 128), Size(9, 9), 0, 0, vector(v[4], v[4]+9), vector(v[4], v[4]+9) }, -#ifdef CV_TEST_INACCURATE_GAUSSIAN_BLUR - { CV_8UC1, Size( 256, 128), Size(3, 3), 1.75, 0.875, vector(v[5], v[5]+3), vector(v[6], v[6]+3) }, - { CV_8UC2, Size( 256, 128), Size(3, 3), 1.75, 0.875, vector(v[5], v[5]+3), vector(v[6], v[6]+3) }, - { CV_8UC3, Size( 256, 128), Size(3, 3), 1.75, 0.875, vector(v[5], v[5]+3), vector(v[6], v[6]+3) }, - { CV_8UC4, Size( 256, 128), Size(3, 3), 1.75, 0.875, vector(v[5], v[5]+3), vector(v[6], v[6]+3) }, - { CV_8UC1, Size( 256, 128), Size(5, 5), 0.375, 0.75, vector(v[7], v[7]+5), vector(v[8], v[8]+5) } -#endif }; int bordermodes[] = { @@ -93,11 +73,12 @@ TEST(GaussianBlur_Bitexact, Linear8U) // BORDER_REFLECT_101 }; - for (int modeind = 0, _modecnt = sizeof(modes) / sizeof(modes[0]); modeind < _modecnt; ++modeind) + template + void checkMode(const testmode& mode) { - int type = modes[modeind].type, depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); - int dcols = modes[modeind].sz.width, drows = modes[modeind].sz.height; - Size kernel = modes[modeind].kernel; + int type = mode.type, depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); + int dcols = mode.sz.width, drows = mode.sz.height; + Size kernel = mode.kernel; int rows = drows + 20, cols = dcols + 20; Mat src(rows, cols, type), refdst(drows, dcols, type), dst; @@ -142,25 +123,93 @@ TEST(GaussianBlur_Bitexact, Linear8U) for (int i = 0; i < dcols; i++) { if (depth == CV_8U) - dst_chan.at(j, i) = eval(src_chan(Rect(i,j,kernel.width,kernel.height)), modes[modeind].kernel_x, modes[modeind].kernel_y); + dst_chan.at(j, i) = eval(src_chan(Rect(i,j,kernel.width,kernel.height)), mode.kernel_x, mode.kernel_y); else if (depth == CV_16U) - dst_chan.at(j, i) = eval(src_chan(Rect(i, j, kernel.width, kernel.height)), modes[modeind].kernel_x, modes[modeind].kernel_y); + dst_chan.at(j, i) = eval(src_chan(Rect(i, j, kernel.width, kernel.height)), mode.kernel_x, mode.kernel_y); else if (depth == CV_16S) - dst_chan.at(j, i) = eval(src_chan(Rect(i, j, kernel.width, kernel.height)), modes[modeind].kernel_x, modes[modeind].kernel_y); + dst_chan.at(j, i) = eval(src_chan(Rect(i, j, kernel.width, kernel.height)), mode.kernel_x, mode.kernel_y); else if (depth == CV_32S) - dst_chan.at(j, i) = eval(src_chan(Rect(i, j, kernel.width, kernel.height)), modes[modeind].kernel_x, modes[modeind].kernel_y); + dst_chan.at(j, i) = eval(src_chan(Rect(i, j, kernel.width, kernel.height)), mode.kernel_x, mode.kernel_y); else CV_Assert(0); } mixChannels(dst_chan, refdst, toFrom, 1); } - cv::GaussianBlur(src_roi, dst, kernel, modes[modeind].sigma_x, modes[modeind].sigma_y, bordermodes[borderind]); + cv::GaussianBlur(src_roi, dst, kernel, mode.sigma_x, mode.sigma_y, bordermodes[borderind]); EXPECT_GE(0, cvtest::norm(refdst, dst, cv::NORM_L1)) - << "GaussianBlur " << cn << "-chan mat " << drows << "x" << dcols << " by kernel " << kernel << " sigma(" << modes[modeind].sigma_x << ";" << modes[modeind].sigma_y << ") failed with max diff " << cvtest::norm(refdst, dst, cv::NORM_INF); + << "GaussianBlur " << cn << "-chan mat " << drows << "x" << dcols << " by kernel " << kernel << " sigma(" << mode.sigma_x << ";" << mode.sigma_y << ") failed with max diff " << cvtest::norm(refdst, dst, cv::NORM_INF); } } + +TEST(GaussianBlur_Bitexact, Linear8U) +{ + testmode modes[] = { + { CV_8UC1, Size( 1, 1), Size(3, 3), 0, 0, vector(vU8[1], vU8[1]+3), vector(vU8[1], vU8[1]+3) }, + { CV_8UC1, Size( 2, 2), Size(3, 3), 0, 0, vector(vU8[1], vU8[1]+3), vector(vU8[1], vU8[1]+3) }, + { CV_8UC1, Size( 3, 1), Size(3, 3), 0, 0, vector(vU8[1], vU8[1]+3), vector(vU8[1], vU8[1]+3) }, + { CV_8UC1, Size( 1, 3), Size(3, 3), 0, 0, vector(vU8[1], vU8[1]+3), vector(vU8[1], vU8[1]+3) }, + { CV_8UC1, Size( 3, 3), Size(3, 3), 0, 0, vector(vU8[1], vU8[1]+3), vector(vU8[1], vU8[1]+3) }, + { CV_8UC1, Size( 3, 3), Size(5, 5), 0, 0, vector(vU8[2], vU8[2]+5), vector(vU8[2], vU8[2]+5) }, + { CV_8UC1, Size( 3, 3), Size(7, 7), 0, 0, vector(vU8[3], vU8[3]+7), vector(vU8[3], vU8[3]+7) }, + { CV_8UC1, Size( 5, 5), Size(3, 3), 0, 0, vector(vU8[1], vU8[1]+3), vector(vU8[1], vU8[1]+3) }, + { CV_8UC1, Size( 5, 5), Size(5, 5), 0, 0, vector(vU8[2], vU8[2]+5), vector(vU8[2], vU8[2]+5) }, + { CV_8UC1, Size( 3, 5), Size(5, 5), 0, 0, vector(vU8[2], vU8[2]+5), vector(vU8[2], vU8[2]+5) }, + { CV_8UC1, Size( 5, 5), Size(5, 5), 0, 0, vector(vU8[2], vU8[2]+5), vector(vU8[2], vU8[2]+5) }, + { CV_8UC1, Size( 5, 5), Size(7, 7), 0, 0, vector(vU8[3], vU8[3]+7), vector(vU8[3], vU8[3]+7) }, + { CV_8UC1, Size( 7, 7), Size(7, 7), 0, 0, vector(vU8[3], vU8[3]+7), vector(vU8[3], vU8[3]+7) }, + { CV_8UC1, Size( 256, 128), Size(3, 3), 0, 0, vector(vU8[1], vU8[1]+3), vector(vU8[1], vU8[1]+3) }, + { CV_8UC2, Size( 256, 128), Size(3, 3), 0, 0, vector(vU8[1], vU8[1]+3), vector(vU8[1], vU8[1]+3) }, + { CV_8UC3, Size( 256, 128), Size(3, 3), 0, 0, vector(vU8[1], vU8[1]+3), vector(vU8[1], vU8[1]+3) }, + { CV_8UC4, Size( 256, 128), Size(3, 3), 0, 0, vector(vU8[1], vU8[1]+3), vector(vU8[1], vU8[1]+3) }, + { CV_8UC1, Size( 256, 128), Size(5, 5), 0, 0, vector(vU8[2], vU8[2]+5), vector(vU8[2], vU8[2]+5) }, + { CV_8UC1, Size( 256, 128), Size(7, 7), 0, 0, vector(vU8[3], vU8[3]+7), vector(vU8[3], vU8[3]+7) }, + { CV_8UC1, Size( 256, 128), Size(9, 9), 0, 0, vector(vU8[4], vU8[4]+9), vector(vU8[4], vU8[4]+9) }, +#ifdef CV_TEST_INACCURATE_GAUSSIAN_BLUR + { CV_8UC1, Size( 256, 128), Size(3, 3), 1.75, 0.875, vector(vU8[5], vU8[5]+3), vector(vU8[6], vU8[6]+3) }, + { CV_8UC2, Size( 256, 128), Size(3, 3), 1.75, 0.875, vector(vU8[5], vU8[5]+3), vector(vU8[6], vU8[6]+3) }, + { CV_8UC3, Size( 256, 128), Size(3, 3), 1.75, 0.875, vector(vU8[5], vU8[5]+3), vector(vU8[6], vU8[6]+3) }, + { CV_8UC4, Size( 256, 128), Size(3, 3), 1.75, 0.875, vector(vU8[5], vU8[5]+3), vector(vU8[6], vU8[6]+3) }, + { CV_8UC1, Size( 256, 128), Size(5, 5), 0.375, 0.75, vector(vU8[7], vU8[7]+5), vector(vU8[8], vU8[8]+5) } +#endif + }; + + for (int modeind = 0, _modecnt = sizeof(modes) / sizeof(modes[0]); modeind < _modecnt; ++modeind) + { + checkMode(modes[modeind]); + } +} + +TEST(GaussianBlur_Bitexact, Linear16U) +{ + testmode modes[] = { + { CV_16UC1, Size( 1, 1), Size(3, 3), 0, 0, vector(vU16[1], vU16[1]+3), vector(vU16[1], vU16[1]+3) }, + { CV_16UC1, Size( 2, 2), Size(3, 3), 0, 0, vector(vU16[1], vU16[1]+3), vector(vU16[1], vU16[1]+3) }, + { CV_16UC1, Size( 3, 1), Size(3, 3), 0, 0, vector(vU16[1], vU16[1]+3), vector(vU16[1], vU16[1]+3) }, + { CV_16UC1, Size( 1, 3), Size(3, 3), 0, 0, vector(vU16[1], vU16[1]+3), vector(vU16[1], vU16[1]+3) }, + { CV_16UC1, Size( 3, 3), Size(3, 3), 0, 0, vector(vU16[1], vU16[1]+3), vector(vU16[1], vU16[1]+3) }, + { CV_16UC1, Size( 3, 3), Size(5, 5), 0, 0, vector(vU16[2], vU16[2]+5), vector(vU16[2], vU16[2]+5) }, + { CV_16UC1, Size( 3, 3), Size(7, 7), 0, 0, vector(vU16[3], vU16[3]+7), vector(vU16[3], vU16[3]+7) }, + { CV_16UC1, Size( 5, 5), Size(3, 3), 0, 0, vector(vU16[1], vU16[1]+3), vector(vU16[1], vU16[1]+3) }, + { CV_16UC1, Size( 5, 5), Size(5, 5), 0, 0, vector(vU16[2], vU16[2]+5), vector(vU16[2], vU16[2]+5) }, + { CV_16UC1, Size( 3, 5), Size(5, 5), 0, 0, vector(vU16[2], vU16[2]+5), vector(vU16[2], vU16[2]+5) }, + { CV_16UC1, Size( 5, 5), Size(5, 5), 0, 0, vector(vU16[2], vU16[2]+5), vector(vU16[2], vU16[2]+5) }, + { CV_16UC1, Size( 5, 5), Size(7, 7), 0, 0, vector(vU16[3], vU16[3]+7), vector(vU16[3], vU16[3]+7) }, + { CV_16UC1, Size( 7, 7), Size(7, 7), 0, 0, vector(vU16[3], vU16[3]+7), vector(vU16[3], vU16[3]+7) }, + { CV_16UC1, Size( 256, 128), Size(3, 3), 0, 0, vector(vU16[1], vU16[1]+3), vector(vU16[1], vU16[1]+3) }, + { CV_16UC2, Size( 256, 128), Size(3, 3), 0, 0, vector(vU16[1], vU16[1]+3), vector(vU16[1], vU16[1]+3) }, + { CV_16UC3, Size( 256, 128), Size(3, 3), 0, 0, vector(vU16[1], vU16[1]+3), vector(vU16[1], vU16[1]+3) }, + { CV_16UC4, Size( 256, 128), Size(3, 3), 0, 0, vector(vU16[1], vU16[1]+3), vector(vU16[1], vU16[1]+3) }, + { CV_16UC1, Size( 256, 128), Size(5, 5), 0, 0, vector(vU16[2], vU16[2]+5), vector(vU16[2], vU16[2]+5) }, + { CV_16UC1, Size( 256, 128), Size(7, 7), 0, 0, vector(vU16[3], vU16[3]+7), vector(vU16[3], vU16[3]+7) }, + { CV_16UC1, Size( 256, 128), Size(9, 9), 0, 0, vector(vU16[4], vU16[4]+9), vector(vU16[4], vU16[4]+9) }, + }; + + for (int modeind = 0, _modecnt = sizeof(modes) / sizeof(modes[0]); modeind < _modecnt; ++modeind) + { + checkMode<16>(modes[modeind]); + } } TEST(GaussianBlur_Bitexact, regression_15015) diff --git a/samples/cpp/tutorial_code/ImgProc/anisotropic_image_segmentation/anisotropic_image_segmentation.cpp b/samples/cpp/tutorial_code/ImgProc/anisotropic_image_segmentation/anisotropic_image_segmentation.cpp index af5a12aa8c..a402e1de4a 100644 --- a/samples/cpp/tutorial_code/ImgProc/anisotropic_image_segmentation/anisotropic_image_segmentation.cpp +++ b/samples/cpp/tutorial_code/ImgProc/anisotropic_image_segmentation/anisotropic_image_segmentation.cpp @@ -81,8 +81,8 @@ void calcGST(const Mat& inputImg, Mat& imgCoherencyOut, Mat& imgOrientationOut, // GST components calculation (stop) // eigenvalue calculation (start) - // lambda1 = J11 + J22 + sqrt((J11-J22)^2 + 4*J12^2) - // lambda2 = J11 + J22 - sqrt((J11-J22)^2 + 4*J12^2) + // lambda1 = 0.5*(J11 + J22 + sqrt((J11-J22)^2 + 4*J12^2)) + // lambda2 = 0.5*(J11 + J22 - sqrt((J11-J22)^2 + 4*J12^2)) Mat tmp1, tmp2, tmp3, tmp4; tmp1 = J11 + J22; tmp2 = J11 - J22; @@ -91,8 +91,10 @@ void calcGST(const Mat& inputImg, Mat& imgCoherencyOut, Mat& imgOrientationOut, sqrt(tmp2 + 4.0 * tmp3, tmp4); Mat lambda1, lambda2; - lambda1 = tmp1 + tmp4; // biggest eigenvalue - lambda2 = tmp1 - tmp4; // smallest eigenvalue + lambda1 = tmp1 + tmp4; + lambda1 = 0.5*lambda1; // biggest eigenvalue + lambda2 = tmp1 - tmp4; + lambda2 = 0.5*lambda2; // smallest eigenvalue // eigenvalue calculation (stop) // Coherency calculation (start) diff --git a/samples/dnn/models.yml b/samples/dnn/models.yml index d177a09aab..7a0bfaf350 100644 --- a/samples/dnn/models.yml +++ b/samples/dnn/models.yml @@ -15,8 +15,9 @@ opencv_fd: rgb: false sample: "object_detection" +# YOLO4 object detection family from Darknet (https://github.com/AlexeyAB/darknet) # YOLO object detection family from Darknet (https://pjreddie.com/darknet/yolo/) -# Might be used for all YOLOv2, TinyYolov2 and YOLOv3 +# Might be used for all YOLOv2, TinyYolov2, YOLOv3, YOLOv4 and TinyYolov4 yolo: model: "yolov3.weights" config: "yolov3.cfg" diff --git a/samples/python/tutorial_code/imgProc/anisotropic_image_segmentation/anisotropic_image_segmentation.py b/samples/python/tutorial_code/imgProc/anisotropic_image_segmentation/anisotropic_image_segmentation.py index 5e2385a1ce..a8a9ccb3cc 100644 --- a/samples/python/tutorial_code/imgProc/anisotropic_image_segmentation/anisotropic_image_segmentation.py +++ b/samples/python/tutorial_code/imgProc/anisotropic_image_segmentation/anisotropic_image_segmentation.py @@ -31,16 +31,16 @@ def calcGST(inputIMG, w): # GST components calculations (stop) # eigenvalue calculation (start) - # lambda1 = J11 + J22 + sqrt((J11-J22)^2 + 4*J12^2) - # lambda2 = J11 + J22 - sqrt((J11-J22)^2 + 4*J12^2) + # lambda1 = 0.5*(J11 + J22 + sqrt((J11-J22)^2 + 4*J12^2)) + # lambda2 = 0.5*(J11 + J22 - sqrt((J11-J22)^2 + 4*J12^2)) tmp1 = J11 + J22 tmp2 = J11 - J22 tmp2 = cv.multiply(tmp2, tmp2) tmp3 = cv.multiply(J12, J12) tmp4 = np.sqrt(tmp2 + 4.0 * tmp3) - lambda1 = tmp1 + tmp4 # biggest eigenvalue - lambda2 = tmp1 - tmp4 # smallest eigenvalue + lambda1 = 0.5*(tmp1 + tmp4) # biggest eigenvalue + lambda2 = 0.5*(tmp1 - tmp4) # smallest eigenvalue # eigenvalue calculation (stop) # Coherency calculation (start)