diff --git a/3rdparty/include/opencl/1.2/CL/cl_d3d11_ext.h b/3rdparty/include/opencl/1.2/CL/cl_d3d11_ext.h new file mode 100644 index 0000000000..99b2c5bda5 --- /dev/null +++ b/3rdparty/include/opencl/1.2/CL/cl_d3d11_ext.h @@ -0,0 +1,122 @@ +/********************************************************************************** + * Copyright (c) 2008-2009 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. + **********************************************************************************/ + +#ifndef __OPENCL_CL_D3D11_EXT_H +#define __OPENCL_CL_D3D11_EXT_H + +#include +#include +#include + +#ifdef __cplusplus +extern "C" { +#endif + +/****************************************************************************** + * cl_nv_d3d11_sharing */ + +typedef cl_uint cl_d3d11_device_source_nv; +typedef cl_uint cl_d3d11_device_set_nv; + +/******************************************************************************/ + +// Error Codes +#define CL_INVALID_D3D11_DEVICE_NV -1006 +#define CL_INVALID_D3D11_RESOURCE_NV -1007 +#define CL_D3D11_RESOURCE_ALREADY_ACQUIRED_NV -1008 +#define CL_D3D11_RESOURCE_NOT_ACQUIRED_NV -1009 + +// cl_d3d11_device_source_nv +#define CL_D3D11_DEVICE_NV 0x4019 +#define CL_D3D11_DXGI_ADAPTER_NV 0x401A + +// cl_d3d11_device_set_nv +#define CL_PREFERRED_DEVICES_FOR_D3D11_NV 0x401B +#define CL_ALL_DEVICES_FOR_D3D11_NV 0x401C + +// cl_context_info +#define CL_CONTEXT_D3D11_DEVICE_NV 0x401D + +// cl_mem_info +#define CL_MEM_D3D11_RESOURCE_NV 0x401E + +// cl_image_info +#define CL_IMAGE_D3D11_SUBRESOURCE_NV 0x401F + +// cl_command_type +#define CL_COMMAND_ACQUIRE_D3D11_OBJECTS_NV 0x4020 +#define CL_COMMAND_RELEASE_D3D11_OBJECTS_NV 0x4021 + +/******************************************************************************/ + +typedef CL_API_ENTRY cl_int (CL_API_CALL *clGetDeviceIDsFromD3D11NV_fn)( + cl_platform_id platform, + cl_d3d11_device_source_nv d3d_device_source, + void * d3d_object, + cl_d3d11_device_set_nv d3d_device_set, + cl_uint num_entries, + cl_device_id * devices, + cl_uint * num_devices) CL_API_SUFFIX__VERSION_1_0; + +typedef CL_API_ENTRY cl_mem (CL_API_CALL *clCreateFromD3D11BufferNV_fn)( + cl_context context, + cl_mem_flags flags, + ID3D11Buffer * resource, + cl_int * errcode_ret) CL_API_SUFFIX__VERSION_1_0; + +typedef CL_API_ENTRY cl_mem (CL_API_CALL *clCreateFromD3D11Texture2DNV_fn)( + cl_context context, + cl_mem_flags flags, + ID3D11Texture2D * resource, + UINT subresource, + cl_int * errcode_ret) CL_API_SUFFIX__VERSION_1_0; + +typedef CL_API_ENTRY cl_mem (CL_API_CALL *clCreateFromD3D11Texture3DNV_fn)( + cl_context context, + cl_mem_flags flags, + ID3D11Texture3D * resource, + UINT subresource, + cl_int * errcode_ret) CL_API_SUFFIX__VERSION_1_0; + +typedef CL_API_ENTRY cl_int (CL_API_CALL *clEnqueueAcquireD3D11ObjectsNV_fn)( + cl_command_queue command_queue, + cl_uint num_objects, + const cl_mem * mem_objects, + cl_uint num_events_in_wait_list, + const cl_event * event_wait_list, + cl_event * event) CL_API_SUFFIX__VERSION_1_0; + +typedef CL_API_ENTRY cl_int (CL_API_CALL *clEnqueueReleaseD3D11ObjectsNV_fn)( + cl_command_queue command_queue, + cl_uint num_objects, + cl_mem * mem_objects, + cl_uint num_events_in_wait_list, + const cl_event * event_wait_list, + cl_event * event) CL_API_SUFFIX__VERSION_1_0; + +#ifdef __cplusplus +} +#endif + +#endif // __OPENCL_CL_D3D11_H + diff --git a/CMakeLists.txt b/CMakeLists.txt index 441ad70b43..ccf9b5de3a 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -108,6 +108,10 @@ if(POLICY CMP0067) cmake_policy(SET CMP0067 NEW) endif() +if(POLICY CMP0068) + cmake_policy(SET CMP0068 NEW) # CMake 3.9+: `RPATH` settings on macOS do not affect `install_name`. +endif() + include(cmake/OpenCVUtils.cmake) ocv_cmake_reset_hooks() ocv_check_environment_variables(OPENCV_CMAKE_HOOKS_DIR) @@ -368,6 +372,9 @@ OCV_OPTION(WITH_OPENCLAMDBLAS "Include AMD OpenCL BLAS library support" ON OCV_OPTION(WITH_DIRECTX "Include DirectX support" ON VISIBLE_IF WIN32 AND NOT WINRT VERIFY HAVE_DIRECTX) +OCV_OPTION(WITH_OPENCL_D3D11_NV "Include NVIDIA OpenCL D3D11 support" WITH_DIRECTX + VISIBLE_IF WIN32 AND NOT WINRT + VERIFY HAVE_OPENCL_D3D11_NV) OCV_OPTION(WITH_LIBREALSENSE "Include Intel librealsense support" OFF VISIBLE_IF NOT WITH_INTELPERC VERIFY HAVE_LIBREALSENSE) @@ -1570,6 +1577,7 @@ if(WITH_OPENCL OR HAVE_OPENCL) IF HAVE_OPENCL_SVM THEN "SVM" IF HAVE_CLAMDFFT THEN "AMDFFT" IF HAVE_CLAMDBLAS THEN "AMDBLAS" + IF HAVE_OPENCL_D3D11_NV THEN "NVD3D11" ELSE "no extra features") status("") status(" OpenCL:" HAVE_OPENCL THEN "YES (${opencl_features})" ELSE "NO") diff --git a/cmake/OpenCVDetectOpenCL.cmake b/cmake/OpenCVDetectOpenCL.cmake index 629d44a62d..6cd7552fa0 100644 --- a/cmake/OpenCVDetectOpenCL.cmake +++ b/cmake/OpenCVDetectOpenCL.cmake @@ -2,14 +2,19 @@ set(OPENCL_FOUND ON CACHE BOOL "OpenCL library is found") if(APPLE) set(OPENCL_LIBRARY "-framework OpenCL" CACHE STRING "OpenCL library") set(OPENCL_INCLUDE_DIR "" CACHE PATH "OpenCL include directory") -else(APPLE) +else() set(OPENCL_LIBRARY "" CACHE STRING "OpenCL library") set(OPENCL_INCLUDE_DIR "${OpenCV_SOURCE_DIR}/3rdparty/include/opencl/1.2" CACHE PATH "OpenCL include directory") ocv_install_3rdparty_licenses(opencl-headers "${OpenCV_SOURCE_DIR}/3rdparty/include/opencl/LICENSE.txt") -endif(APPLE) +endif() mark_as_advanced(OPENCL_INCLUDE_DIR OPENCL_LIBRARY) if(OPENCL_FOUND) + + if(WITH_OPENCL_D3D11_NV AND EXISTS "${OPENCL_INCLUDE_DIR}/CL/cl_d3d11_ext.h") + set(HAVE_OPENCL_D3D11_NV ON) + endif() + if(OPENCL_LIBRARY) set(HAVE_OPENCL_STATIC ON) set(OPENCL_LIBRARIES "${OPENCL_LIBRARY}") diff --git a/cmake/templates/cvconfig.h.in b/cmake/templates/cvconfig.h.in index 271bf275ef..1f7981cd41 100644 --- a/cmake/templates/cvconfig.h.in +++ b/cmake/templates/cvconfig.h.in @@ -100,6 +100,9 @@ #cmakedefine HAVE_OPENCL_STATIC #cmakedefine HAVE_OPENCL_SVM +/* NVIDIA OpenCL D3D Extensions support */ +#cmakedefine HAVE_OPENCL_D3D11_NV + /* OpenEXR codec */ #cmakedefine HAVE_OPENEXR diff --git a/modules/core/include/opencv2/core/opencl/runtime/autogenerated/opencl_core.hpp b/modules/core/include/opencv2/core/opencl/runtime/autogenerated/opencl_core.hpp index fdaf469ff1..28618a1f3a 100644 --- a/modules/core/include/opencv2/core/opencl/runtime/autogenerated/opencl_core.hpp +++ b/modules/core/include/opencv2/core/opencl/runtime/autogenerated/opencl_core.hpp @@ -96,6 +96,7 @@ #define clWaitForEvents clWaitForEvents_ #if defined __APPLE__ +#define CL_SILENCE_DEPRECATION #include #else #include diff --git a/modules/core/src/directx.cpp b/modules/core/src/directx.cpp index 99c27ec705..f8b641d5cf 100644 --- a/modules/core/src/directx.cpp +++ b/modules/core/src/directx.cpp @@ -48,13 +48,13 @@ #ifdef HAVE_DIRECTX #include -# include "directx.inc.hpp" +#include "directx.inc.hpp" #else // HAVE_DIRECTX #define NO_DIRECTX_SUPPORT_ERROR CV_Error(cv::Error::StsBadFunc, "OpenCV was build without DirectX support") #endif #ifndef HAVE_OPENCL -# define NO_OPENCL_SUPPORT_ERROR CV_Error(cv::Error::StsBadFunc, "OpenCV was build without OpenCL support") +#define NO_OPENCL_SUPPORT_ERROR CV_Error(cv::Error::StsBadFunc, "OpenCV was build without OpenCL support") #endif // HAVE_OPENCL namespace cv { namespace directx { @@ -168,7 +168,7 @@ int getTypeFromDXGI_FORMAT(const int iDXGI_FORMAT) //case DXGI_FORMAT_BC7_TYPELESS: //case DXGI_FORMAT_BC7_UNORM: //case DXGI_FORMAT_BC7_UNORM_SRGB: -#ifdef HAVE_DIRECTX_NV12 +#ifdef HAVE_DIRECTX_NV12 //D3DX11 should support DXGI_FORMAT_NV12. case DXGI_FORMAT_NV12: return CV_8UC3; #endif default: break; @@ -256,75 +256,70 @@ Context& initializeContextFromD3D11Device(ID3D11Device* pD3D11Device) CV_Error(cv::Error::OpenCLInitError, "OpenCL: No available platforms"); std::vector platforms(numPlatforms); - status = clGetPlatformIDs(numPlatforms, &platforms[0], NULL); - if (status != CL_SUCCESS) - CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't get number of platforms"); - - // TODO Filter platforms by name from OPENCV_OPENCL_DEVICE + size_t exts_len; + cv::AutoBuffer extensions; + bool is_support_cl_khr_d3d11_sharing = false; +#ifdef HAVE_OPENCL_D3D11_NV + bool is_support_cl_nv_d3d11_sharing = false; +#endif + for (int i = 0; i < (int)numPlatforms; i++) + { + status = clGetPlatformIDs(numPlatforms, &platforms[i], NULL); + if (status != CL_SUCCESS) + CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't get number of platforms"); + status = clGetPlatformInfo(platforms[i], CL_PLATFORM_EXTENSIONS, 0, NULL, &exts_len); + if (status != CL_SUCCESS) + CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't get length of CL_PLATFORM_EXTENSIONS"); + extensions.resize(exts_len); + status = clGetPlatformInfo(platforms[i], CL_PLATFORM_EXTENSIONS, exts_len, static_cast(extensions.data()), NULL); + if (status != CL_SUCCESS) + CV_Error(cv::Error::OpenCLInitError, "OpenCL: No available CL_PLATFORM_EXTENSIONS"); + if (strstr(extensions.data(), "cl_khr_d3d11_sharing")) + is_support_cl_khr_d3d11_sharing = true; +#ifdef HAVE_OPENCL_D3D11_NV + if (strstr(extensions.data(), "cl_nv_d3d11_sharing")) + is_support_cl_nv_d3d11_sharing = true; +#endif + } +#ifdef HAVE_OPENCL_D3D11_NV + if (!is_support_cl_nv_d3d11_sharing && !is_support_cl_khr_d3d11_sharing) + CV_Error(cv::Error::OpenCLInitError, "OpenCL: No supported extensions"); +#else + if (!is_support_cl_khr_d3d11_sharing) + CV_Error(cv::Error::OpenCLInitError, "OpenCL: No supported extensions"); +#endif int found = -1; cl_device_id device = NULL; cl_uint numDevices = 0; cl_context context = NULL; - // try with CL_PREFERRED_DEVICES_FOR_D3D11_KHR - for (int i = 0; i < (int)numPlatforms; i++) +#ifdef HAVE_OPENCL_D3D11_NV + if (is_support_cl_nv_d3d11_sharing) { - clGetDeviceIDsFromD3D11KHR_fn clGetDeviceIDsFromD3D11KHR = (clGetDeviceIDsFromD3D11KHR_fn) - clGetExtensionFunctionAddressForPlatform(platforms[i], "clGetDeviceIDsFromD3D11KHR"); - if (!clGetDeviceIDsFromD3D11KHR) - continue; - - device = NULL; - numDevices = 0; - status = clGetDeviceIDsFromD3D11KHR(platforms[i], CL_D3D11_DEVICE_KHR, pD3D11Device, - CL_PREFERRED_DEVICES_FOR_D3D11_KHR, 1, &device, &numDevices); - if (status != CL_SUCCESS) - continue; - if (numDevices > 0) - { - cl_context_properties properties[] = { - CL_CONTEXT_PLATFORM, (cl_context_properties)platforms[i], - CL_CONTEXT_D3D11_DEVICE_KHR, (cl_context_properties)(pD3D11Device), - CL_CONTEXT_INTEROP_USER_SYNC, CL_FALSE, - NULL, NULL - }; - context = clCreateContext(properties, 1, &device, NULL, NULL, &status); - if (status != CL_SUCCESS) - { - clReleaseDevice(device); - } - else - { - found = i; - break; - } - } - } - if (found < 0) - { - // try with CL_ALL_DEVICES_FOR_D3D11_KHR + // try with CL_PREFERRED_DEVICES_FOR_D3D11_NV for (int i = 0; i < (int)numPlatforms; i++) { - clGetDeviceIDsFromD3D11KHR_fn clGetDeviceIDsFromD3D11KHR = (clGetDeviceIDsFromD3D11KHR_fn) - clGetExtensionFunctionAddressForPlatform(platforms[i], "clGetDeviceIDsFromD3D11KHR"); - if (!clGetDeviceIDsFromD3D11KHR) + clGetDeviceIDsFromD3D11NV_fn clGetDeviceIDsFromD3D11NV = (clGetDeviceIDsFromD3D11NV_fn) + clGetExtensionFunctionAddressForPlatform(platforms[i], "clGetDeviceIDsFromD3D11NV"); + if (!clGetDeviceIDsFromD3D11NV) continue; device = NULL; numDevices = 0; - status = clGetDeviceIDsFromD3D11KHR(platforms[i], CL_D3D11_DEVICE_KHR, pD3D11Device, - CL_ALL_DEVICES_FOR_D3D11_KHR, 1, &device, &numDevices); + status = clGetDeviceIDsFromD3D11NV(platforms[i], CL_D3D11_DEVICE_NV, pD3D11Device, + CL_PREFERRED_DEVICES_FOR_D3D11_NV, 1, &device, &numDevices); if (status != CL_SUCCESS) continue; if (numDevices > 0) { cl_context_properties properties[] = { CL_CONTEXT_PLATFORM, (cl_context_properties)platforms[i], - CL_CONTEXT_D3D11_DEVICE_KHR, (cl_context_properties)(pD3D11Device), - CL_CONTEXT_INTEROP_USER_SYNC, CL_FALSE, - NULL, NULL + CL_CONTEXT_D3D11_DEVICE_NV, (cl_context_properties)(pD3D11Device), + //CL_CONTEXT_INTEROP_USER_SYNC, CL_FALSE, + 0 }; + context = clCreateContext(properties, 1, &device, NULL, NULL, &status); if (status != CL_SUCCESS) { @@ -338,9 +333,127 @@ Context& initializeContextFromD3D11Device(ID3D11Device* pD3D11Device) } } if (found < 0) - CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't create context for DirectX interop"); - } + { + // try with CL_ALL_DEVICES_FOR_D3D11_NV + for (int i = 0; i < (int)numPlatforms; i++) + { + clGetDeviceIDsFromD3D11NV_fn clGetDeviceIDsFromD3D11NV = (clGetDeviceIDsFromD3D11NV_fn) + clGetExtensionFunctionAddressForPlatform(platforms[i], "clGetDeviceIDsFromD3D11NV"); + if (!clGetDeviceIDsFromD3D11NV) + continue; + device = NULL; + numDevices = 0; + status = clGetDeviceIDsFromD3D11NV(platforms[i], CL_D3D11_DEVICE_NV, pD3D11Device, + CL_ALL_DEVICES_FOR_D3D11_NV, 1, &device, &numDevices); + if (status != CL_SUCCESS) + continue; + if (numDevices > 0) + { + cl_context_properties properties[] = { + CL_CONTEXT_PLATFORM, (cl_context_properties)platforms[i], + CL_CONTEXT_D3D11_DEVICE_NV, (cl_context_properties)(pD3D11Device), + //CL_CONTEXT_INTEROP_USER_SYNC, CL_FALSE, + 0 + }; + context = clCreateContext(properties, 1, &device, NULL, NULL, &status); + if (status != CL_SUCCESS) + { + clReleaseDevice(device); + } + else + { + found = i; + break; + } + } + } + } + } +#endif + if (is_support_cl_khr_d3d11_sharing) + { + if (found < 0) + { + // try with CL_PREFERRED_DEVICES_FOR_D3D11_KHR + for (int i = 0; i < (int)numPlatforms; i++) + { + clGetDeviceIDsFromD3D11KHR_fn clGetDeviceIDsFromD3D11KHR = (clGetDeviceIDsFromD3D11KHR_fn) + clGetExtensionFunctionAddressForPlatform(platforms[i], "clGetDeviceIDsFromD3D11KHR"); + if (!clGetDeviceIDsFromD3D11KHR) + continue; + + device = NULL; + numDevices = 0; + + status = clGetDeviceIDsFromD3D11KHR(platforms[i], CL_D3D11_DEVICE_KHR, pD3D11Device, + CL_PREFERRED_DEVICES_FOR_D3D11_KHR, 1, &device, &numDevices); + + if (status != CL_SUCCESS) + continue; + if (numDevices > 0) + { + cl_context_properties properties[] = { + CL_CONTEXT_PLATFORM, (cl_context_properties)platforms[i], + CL_CONTEXT_D3D11_DEVICE_KHR, (cl_context_properties)(pD3D11Device), + CL_CONTEXT_INTEROP_USER_SYNC, CL_FALSE, + NULL, NULL + }; + context = clCreateContext(properties, 1, &device, NULL, NULL, &status); + if (status != CL_SUCCESS) + { + clReleaseDevice(device); + } + else + { + found = i; + break; + } + } + } + } + if (found < 0) + { + // try with CL_ALL_DEVICES_FOR_D3D11_KHR + for (int i = 0; i < (int)numPlatforms; i++) + { + clGetDeviceIDsFromD3D11KHR_fn clGetDeviceIDsFromD3D11KHR = (clGetDeviceIDsFromD3D11KHR_fn) + clGetExtensionFunctionAddressForPlatform(platforms[i], "clGetDeviceIDsFromD3D11KHR"); + if (!clGetDeviceIDsFromD3D11KHR) + continue; + + device = NULL; + numDevices = 0; + status = clGetDeviceIDsFromD3D11KHR(platforms[i], CL_D3D11_DEVICE_KHR, pD3D11Device, + CL_ALL_DEVICES_FOR_D3D11_KHR, 1, &device, &numDevices); + if (status != CL_SUCCESS) + continue; + if (numDevices > 0) + { + cl_context_properties properties[] = { + CL_CONTEXT_PLATFORM, (cl_context_properties)platforms[i], + CL_CONTEXT_D3D11_DEVICE_KHR, (cl_context_properties)(pD3D11Device), + CL_CONTEXT_INTEROP_USER_SYNC, CL_FALSE, + NULL, NULL + }; + context = clCreateContext(properties, 1, &device, NULL, NULL, &status); + if (status != CL_SUCCESS) + { + clReleaseDevice(device); + } + else + { + found = i; + break; + } + } + } + } + } + if (found < 0) + { + CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't create context for DirectX interop"); + } Context& ctx = Context::getDefault(false); initializeContextFromHandle(ctx, platforms[found], context, device); @@ -679,29 +792,85 @@ Context& initializeContextFromDirect3DDevice9(IDirect3DDevice9* pDirect3DDevice9 } // namespace cv::ocl #if defined(HAVE_DIRECTX) && defined(HAVE_OPENCL) + +#ifdef HAVE_OPENCL_D3D11_NV +clCreateFromD3D11Texture2DNV_fn clCreateFromD3D11Texture2DNV = NULL; +clEnqueueAcquireD3D11ObjectsNV_fn clEnqueueAcquireD3D11ObjectsNV = NULL; +clEnqueueReleaseD3D11ObjectsNV_fn clEnqueueReleaseD3D11ObjectsNV = NULL; +#endif clCreateFromD3D11Texture2DKHR_fn clCreateFromD3D11Texture2DKHR = NULL; clEnqueueAcquireD3D11ObjectsKHR_fn clEnqueueAcquireD3D11ObjectsKHR = NULL; clEnqueueReleaseD3D11ObjectsKHR_fn clEnqueueReleaseD3D11ObjectsKHR = NULL; -static void __OpenCLinitializeD3D11() +static bool __OpenCLinitializeD3D11() { using namespace cv::ocl; static cl_platform_id initializedPlatform = NULL; cl_platform_id platform = (cl_platform_id)Platform::getDefault().ptr(); - if (initializedPlatform != platform) + + bool useCLNVEXT = false; + size_t exts_len; + cl_int status = clGetPlatformInfo(platform, CL_PLATFORM_EXTENSIONS, 0, NULL, &exts_len); + if (status != CL_SUCCESS) + CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't get length of CL_PLATFORM_EXTENSIONS"); + cv::AutoBuffer extensions(exts_len); + status = clGetPlatformInfo(platform, CL_PLATFORM_EXTENSIONS, exts_len, static_cast(extensions.data()), NULL); + if (status != CL_SUCCESS) + CV_Error(cv::Error::OpenCLInitError, "OpenCL: No available CL_PLATFORM_EXTENSIONS"); + bool is_support_cl_khr_d3d11_sharing = false; + if (strstr(extensions.data(), "cl_khr_d3d11_sharing")) + is_support_cl_khr_d3d11_sharing = true; +#ifdef HAVE_OPENCL_D3D11_NV + bool is_support_cl_nv_d3d11_sharing = false; + if (strstr(extensions.data(), "cl_nv_d3d11_sharing")) + is_support_cl_nv_d3d11_sharing = true; + if (!is_support_cl_nv_d3d11_sharing && !is_support_cl_khr_d3d11_sharing) + CV_Error(cv::Error::OpenCLInitError, "OpenCL: No supported extensions"); +#else + if (!is_support_cl_khr_d3d11_sharing) + CV_Error(cv::Error::OpenCLInitError, "OpenCL: No supported extensions"); +#endif + +#ifdef HAVE_OPENCL_D3D11_NV + if (is_support_cl_nv_d3d11_sharing) { - clCreateFromD3D11Texture2DKHR = (clCreateFromD3D11Texture2DKHR_fn) - clGetExtensionFunctionAddressForPlatform(platform, "clCreateFromD3D11Texture2DKHR"); - clEnqueueAcquireD3D11ObjectsKHR = (clEnqueueAcquireD3D11ObjectsKHR_fn) - clGetExtensionFunctionAddressForPlatform(platform, "clEnqueueAcquireD3D11ObjectsKHR"); - clEnqueueReleaseD3D11ObjectsKHR = (clEnqueueReleaseD3D11ObjectsKHR_fn) - clGetExtensionFunctionAddressForPlatform(platform, "clEnqueueReleaseD3D11ObjectsKHR"); - initializedPlatform = platform; + if (initializedPlatform != platform) + { + clCreateFromD3D11Texture2DNV = (clCreateFromD3D11Texture2DNV_fn) + clGetExtensionFunctionAddressForPlatform(platform, "clCreateFromD3D11Texture2DNV"); + clEnqueueAcquireD3D11ObjectsNV = (clEnqueueAcquireD3D11ObjectsNV_fn) + clGetExtensionFunctionAddressForPlatform(platform, "clEnqueueAcquireD3D11ObjectsNV"); + clEnqueueReleaseD3D11ObjectsNV = (clEnqueueReleaseD3D11ObjectsNV_fn) + clGetExtensionFunctionAddressForPlatform(platform, "clEnqueueReleaseD3D11ObjectsNV"); + initializedPlatform = platform; + } + if (clCreateFromD3D11Texture2DNV && clEnqueueAcquireD3D11ObjectsNV && clEnqueueReleaseD3D11ObjectsNV) + { + useCLNVEXT = true; + } } - if (!clCreateFromD3D11Texture2DKHR || !clEnqueueAcquireD3D11ObjectsKHR || !clEnqueueReleaseD3D11ObjectsKHR) + else +#endif { - CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't find functions for D3D11"); + if (is_support_cl_khr_d3d11_sharing) + { + if (initializedPlatform != platform) + { + clCreateFromD3D11Texture2DKHR = (clCreateFromD3D11Texture2DKHR_fn) + clGetExtensionFunctionAddressForPlatform(platform, "clCreateFromD3D11Texture2DKHR"); + clEnqueueAcquireD3D11ObjectsKHR = (clEnqueueAcquireD3D11ObjectsKHR_fn) + clGetExtensionFunctionAddressForPlatform(platform, "clEnqueueAcquireD3D11ObjectsKHR"); + clEnqueueReleaseD3D11ObjectsKHR = (clEnqueueReleaseD3D11ObjectsKHR_fn) + clGetExtensionFunctionAddressForPlatform(platform, "clEnqueueReleaseD3D11ObjectsKHR"); + initializedPlatform = platform; + } + if (!clCreateFromD3D11Texture2DKHR || !clEnqueueAcquireD3D11ObjectsKHR || !clEnqueueReleaseD3D11ObjectsKHR) + { + CV_Error(cv::Error::OpenCLInitError, "OpenCL: Can't find functions for D3D11"); + } + } } + return useCLNVEXT; } #endif // defined(HAVE_DIRECTX) && defined(HAVE_OPENCL) @@ -762,14 +931,9 @@ bool ocl_convert_bgr_to_nv12( namespace directx { -void convertToD3D11Texture2D(InputArray src, ID3D11Texture2D* pD3D11Texture2D) +#if defined(HAVE_DIRECTX) && defined(HAVE_OPENCL) +static void __convertToD3D11Texture2DKHR(InputArray src, ID3D11Texture2D* pD3D11Texture2D) { - CV_UNUSED(src); CV_UNUSED(pD3D11Texture2D); -#if !defined(HAVE_DIRECTX) - NO_DIRECTX_SUPPORT_ERROR; -#elif defined(HAVE_OPENCL) - __OpenCLinitializeD3D11(); - D3D11_TEXTURE2D_DESC desc = { 0 }; pD3D11Texture2D->GetDesc(&desc); @@ -797,7 +961,6 @@ void convertToD3D11Texture2D(InputArray src, ID3D11Texture2D* pD3D11Texture2D) #ifdef HAVE_DIRECTX_NV12 cl_mem clImageUV = 0; #endif - clImage = clCreateFromD3D11Texture2DKHR(context, CL_MEM_WRITE_ONLY, pD3D11Texture2D, 0, &status); if (status != CL_SUCCESS) CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clCreateFromD3D11Texture2DKHR failed"); @@ -863,22 +1026,108 @@ void convertToD3D11Texture2D(InputArray src, ID3D11Texture2D* pD3D11Texture2D) CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clReleaseMem failed"); } #endif +} +#endif -#else - // TODO memcpy - NO_OPENCL_SUPPORT_ERROR; +#if defined(HAVE_OPENCL_D3D11_NV) +static void __convertToD3D11Texture2DNV(InputArray src, ID3D11Texture2D* pD3D11Texture2D) +{ + D3D11_TEXTURE2D_DESC desc = { 0 }; + pD3D11Texture2D->GetDesc(&desc); + + int srcType = src.type(); + int textureType = getTypeFromDXGI_FORMAT(desc.Format); + CV_Assert(textureType == srcType); + + Size srcSize = src.size(); + CV_Assert(srcSize.width == (int)desc.Width && srcSize.height == (int)desc.Height); + + UMat u = src.getUMat(); + + // TODO Add support for roi + CV_Assert(u.offset == 0); + CV_Assert(u.isContinuous()); + + cl_mem clBuffer = (cl_mem)u.handle(ACCESS_READ); + + using namespace cv::ocl; + Context& ctx = Context::getDefault(); + cl_context context = (cl_context)ctx.ptr(); + + cl_int status = 0; + cl_mem clImage = 0; +#ifdef HAVE_DIRECTX_NV12 + cl_mem clImageUV = 0; +#endif + clImage = clCreateFromD3D11Texture2DNV(context, CL_MEM_WRITE_ONLY, pD3D11Texture2D, 0, &status); + if (status != CL_SUCCESS) + CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clCreateFromD3D11Texture2DNV failed"); + +#ifdef HAVE_DIRECTX_NV12 + if (DXGI_FORMAT_NV12 == desc.Format) + { + clImageUV = clCreateFromD3D11Texture2DNV(context, CL_MEM_WRITE_ONLY, pD3D11Texture2D, 1, &status); + if (status != CL_SUCCESS) + CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clCreateFromD3D11Texture2DNV failed"); + } +#endif + cl_command_queue q = (cl_command_queue)Queue::getDefault().ptr(); + status = clEnqueueAcquireD3D11ObjectsNV(q, 1, &clImage, 0, NULL, NULL); + if (status != CL_SUCCESS) + CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueAcquireD3D11ObjectsNV failed"); + +#ifdef HAVE_DIRECTX_NV12 + if(DXGI_FORMAT_NV12 == desc.Format) + { + status = clEnqueueAcquireD3D11ObjectsNV(q, 1, &clImageUV, 0, NULL, NULL); + if (status != CL_SUCCESS) + CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueAcquireD3D11ObjectsNV failed"); + + if(!ocl::ocl_convert_bgr_to_nv12(clBuffer, (int)u.step[0], u.cols, u.rows, clImage, clImageUV)) + CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: ocl_convert_bgr_to_nv12 failed"); + + status = clEnqueueReleaseD3D11ObjectsNV(q, 1, &clImageUV, 0, NULL, NULL); + if (status != CL_SUCCESS) + CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueReleaseD3D11ObjectsNV failed"); + } + else +#endif + { + size_t offset = 0; // TODO + size_t origin[3] = { 0, 0, 0 }; + size_t region[3] = { (size_t)u.cols, (size_t)u.rows, 1 }; + + status = clEnqueueCopyBufferToImage(q, clBuffer, clImage, offset, origin, region, 0, NULL, NULL); + if (status != CL_SUCCESS) + CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueCopyBufferToImage failed"); + } + + status = clEnqueueReleaseD3D11ObjectsNV(q, 1, &clImage, 0, NULL, NULL); + if (status != CL_SUCCESS) + CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueReleaseD3D11ObjectsNV failed"); + + status = clFinish(q); // TODO Use events + if (status != CL_SUCCESS) + CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clFinish failed"); + + status = clReleaseMemObject(clImage); // TODO RAII + if (status != CL_SUCCESS) + CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clReleaseMem failed"); + +#ifdef HAVE_DIRECTX_NV12 + if(DXGI_FORMAT_NV12 == desc.Format) + { + status = clReleaseMemObject(clImageUV); + if (status != CL_SUCCESS) + CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clReleaseMem failed"); + } #endif } +#endif - -void convertFromD3D11Texture2D(ID3D11Texture2D* pD3D11Texture2D, OutputArray dst) +#if defined(HAVE_DIRECTX) && defined(HAVE_OPENCL) +static void __convertFromD3D11Texture2DKHR(ID3D11Texture2D* pD3D11Texture2D, OutputArray dst) { - CV_UNUSED(pD3D11Texture2D); CV_UNUSED(dst); -#if !defined(HAVE_DIRECTX) - NO_DIRECTX_SUPPORT_ERROR; -#elif defined(HAVE_OPENCL) - __OpenCLinitializeD3D11(); - D3D11_TEXTURE2D_DESC desc = { 0 }; pD3D11Texture2D->GetDesc(&desc); @@ -968,10 +1217,144 @@ void convertFromD3D11Texture2D(ID3D11Texture2D* pD3D11Texture2D, OutputArray dst CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clReleaseMem failed"); } #endif +} +#endif -#else - // TODO memcpy +#if defined(HAVE_OPENCL_D3D11_NV) +static void __convertFromD3D11Texture2DNV(ID3D11Texture2D* pD3D11Texture2D, OutputArray dst) +{ + D3D11_TEXTURE2D_DESC desc = { 0 }; + pD3D11Texture2D->GetDesc(&desc); + + int textureType = getTypeFromDXGI_FORMAT(desc.Format); + CV_Assert(textureType >= 0); + + // TODO Need to specify ACCESS_WRITE here somehow to prevent useless data copying! + dst.create(Size(desc.Width, desc.Height), textureType); + UMat u = dst.getUMat(); + + // TODO Add support for roi + CV_Assert(u.offset == 0); + CV_Assert(u.isContinuous()); + + cl_mem clBuffer = (cl_mem)u.handle(ACCESS_READ); + + using namespace cv::ocl; + Context& ctx = Context::getDefault(); + cl_context context = (cl_context)ctx.ptr(); + + cl_int status = 0; + cl_mem clImage = 0; + + clImage = clCreateFromD3D11Texture2DNV(context, CL_MEM_READ_ONLY, pD3D11Texture2D, 0, &status); + if (status != CL_SUCCESS) + CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clCreateFromD3D11Texture2DNV failed"); + +#ifdef HAVE_DIRECTX_NV12 + cl_mem clImageUV = 0; + if(DXGI_FORMAT_NV12 == desc.Format) + { + clImageUV = clCreateFromD3D11Texture2DNV(context, CL_MEM_READ_ONLY, pD3D11Texture2D, 1, &status); + if (status != CL_SUCCESS) + CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clCreateFromD3D11Texture2DNV failed"); + } +#endif + + cl_command_queue q = (cl_command_queue)Queue::getDefault().ptr(); + status = clEnqueueAcquireD3D11ObjectsNV(q, 1, &clImage, 0, NULL, NULL); + if (status != CL_SUCCESS) + CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueAcquireD3D11ObjectsNV failed"); + +#ifdef HAVE_DIRECTX_NV12 + if (DXGI_FORMAT::DXGI_FORMAT_NV12 == desc.Format) + { + status = clEnqueueAcquireD3D11ObjectsNV(q, 1, &clImageUV, 0, NULL, NULL); + if (status != CL_SUCCESS) + CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueAcquireD3D11ObjectsNV failed"); + + if (!ocl::ocl_convert_nv12_to_bgr(clImage, clImageUV, clBuffer, (int)u.step[0], u.cols, u.rows)) + CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: ocl_convert_nv12_to_bgr failed"); + + status = clEnqueueReleaseD3D11ObjectsNV(q, 1, &clImageUV, 0, NULL, NULL); + if (status != CL_SUCCESS) + CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueReleaseD3D11ObjectsNV failed"); + } + else +#endif + { + size_t offset = 0; // TODO + size_t origin[3] = { 0, 0, 0 }; + size_t region[3] = { (size_t)u.cols, (size_t)u.rows, 1 }; + + status = clEnqueueCopyImageToBuffer(q, clImage, clBuffer, origin, region, offset, 0, NULL, NULL); + if (status != CL_SUCCESS) + CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueCopyImageToBuffer failed"); + } + + status = clEnqueueReleaseD3D11ObjectsNV(q, 1, &clImage, 0, NULL, NULL); + if (status != CL_SUCCESS) + CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueReleaseD3D11ObjectsNV failed"); + + status = clFinish(q); // TODO Use events + if (status != CL_SUCCESS) + CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clFinish failed"); + + status = clReleaseMemObject(clImage); // TODO RAII + if (status != CL_SUCCESS) + CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clReleaseMem failed"); + +#ifdef HAVE_DIRECTX_NV12 + if(DXGI_FORMAT_NV12 == desc.Format) + { + status = clReleaseMemObject(clImageUV); + if (status != CL_SUCCESS) + CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clReleaseMem failed"); + } +#endif +} +#endif + +void convertToD3D11Texture2D(InputArray src, ID3D11Texture2D* pD3D11Texture2D) +{ + CV_UNUSED(src); CV_UNUSED(pD3D11Texture2D); +#if !defined(HAVE_DIRECTX) + NO_DIRECTX_SUPPORT_ERROR; +#elif !defined(HAVE_OPENCL) NO_OPENCL_SUPPORT_ERROR; +#else + + bool useCLNVEXT = __OpenCLinitializeD3D11(); + if(!useCLNVEXT){ + __convertToD3D11Texture2DKHR(src,pD3D11Texture2D); + } +#ifdef HAVE_OPENCL_D3D11_NV + else + { + __convertToD3D11Texture2DNV(src,pD3D11Texture2D); + } +#endif +#endif +} + +void convertFromD3D11Texture2D(ID3D11Texture2D* pD3D11Texture2D, OutputArray dst) +{ + CV_UNUSED(pD3D11Texture2D); CV_UNUSED(dst); +#if !defined(HAVE_DIRECTX) + NO_DIRECTX_SUPPORT_ERROR; +#elif !defined(HAVE_OPENCL) + NO_OPENCL_SUPPORT_ERROR; +#else + + bool useCLNVEXT = __OpenCLinitializeD3D11(); + if(!useCLNVEXT){ + __convertFromD3D11Texture2DKHR(pD3D11Texture2D,dst); + } +#ifdef HAVE_OPENCL_D3D11_NV + else + { + __convertFromD3D11Texture2DNV(pD3D11Texture2D,dst); + } +#endif #endif } diff --git a/modules/core/src/directx.inc.hpp b/modules/core/src/directx.inc.hpp index a8e7a17144..cf8571616c 100644 --- a/modules/core/src/directx.inc.hpp +++ b/modules/core/src/directx.inc.hpp @@ -48,6 +48,9 @@ #include "opencv2/core/opencl/runtime/opencl_core.hpp" #include +#ifdef HAVE_OPENCL_D3D11_NV +#include +#endif #include #include #endif // HAVE_OPENCL diff --git a/modules/core/src/opencl/runtime/generator/template/opencl_core.hpp.in b/modules/core/src/opencl/runtime/generator/template/opencl_core.hpp.in index e04fb0b883..a2b5296ff2 100644 --- a/modules/core/src/opencl/runtime/generator/template/opencl_core.hpp.in +++ b/modules/core/src/opencl/runtime/generator/template/opencl_core.hpp.in @@ -5,6 +5,7 @@ @CL_REMAP_ORIGIN@ #if defined __APPLE__ +#define CL_SILENCE_DEPRECATION #include #else #include diff --git a/modules/core/src/opencl/runtime/opencl_core.cpp b/modules/core/src/opencl/runtime/opencl_core.cpp index 1d761390cc..9885e4d647 100644 --- a/modules/core/src/opencl/runtime/opencl_core.cpp +++ b/modules/core/src/opencl/runtime/opencl_core.cpp @@ -47,6 +47,7 @@ #if defined(HAVE_OPENCL_STATIC) #if defined __APPLE__ +#define CL_SILENCE_DEPRECATION #include #else #include diff --git a/modules/dnn/src/layers/lrn_layer.cpp b/modules/dnn/src/layers/lrn_layer.cpp index 3b7ce25ddf..cc94329c8b 100644 --- a/modules/dnn/src/layers/lrn_layer.cpp +++ b/modules/dnn/src/layers/lrn_layer.cpp @@ -92,7 +92,7 @@ public: virtual bool supportBackend(int backendId) CV_OVERRIDE { if (backendId == DNN_BACKEND_INFERENCE_ENGINE) - return (bias == 1) && (preferableTarget != DNN_TARGET_MYRIAD || type == SPATIAL_NRM); + return bias == 1; return backendId == DNN_BACKEND_OPENCV || backendId == DNN_BACKEND_HALIDE || (backendId == DNN_BACKEND_VKCOM && haveVulkan() && (size % 2 == 1) && (type == CHANNEL_NRM)); diff --git a/modules/dnn/src/onnx/onnx_importer.cpp b/modules/dnn/src/onnx/onnx_importer.cpp index c29c192241..26228af5b1 100644 --- a/modules/dnn/src/onnx/onnx_importer.cpp +++ b/modules/dnn/src/onnx/onnx_importer.cpp @@ -591,6 +591,37 @@ void ONNXImporter::populateNet(Net dstNet) } layerParams.set("num_output", layerParams.blobs[0].size[1] * layerParams.get("group", 1)); layerParams.set("bias_term", node_proto.input_size() == 3); + + if (layerParams.has("output_shape")) + { + const DictValue& outShape = layerParams.get("output_shape"); + + if (outShape.size() != 4) + CV_Error(Error::StsNotImplemented, "Output shape must have 4 elements."); + + const int strideY = layerParams.get("stride_h", 1); + const int strideX = layerParams.get("stride_w", 1); + const int outH = outShape.getIntValue(2); + const int outW = outShape.getIntValue(3); + + if (layerParams.get("pad_mode") == "SAME") + { + layerParams.set("adj_w", (outW - 1) % strideX); + layerParams.set("adj_h", (outH - 1) % strideY); + } + else if (layerParams.get("pad_mode") == "VALID") + { + if (!layerParams.has("kernel_h") || !layerParams.has("kernel_w")) + CV_Error(Error::StsNotImplemented, + "Required attributes 'kernel_h' and 'kernel_w' are not present."); + + int kernelH = layerParams.get("kernel_h"); + int kernelW = layerParams.get("kernel_w"); + + layerParams.set("adj_w", (outW - kernelW) % strideX); + layerParams.set("adj_h", (outH - kernelH) % strideY); + } + } } else if (layer_type == "Transpose") { diff --git a/modules/dnn/test/test_halide_layers.cpp b/modules/dnn/test/test_halide_layers.cpp index bb605c05bc..0e70dc4c49 100644 --- a/modules/dnn/test/test_halide_layers.cpp +++ b/modules/dnn/test/test_halide_layers.cpp @@ -228,6 +228,10 @@ TEST_P(LRN, Accuracy) Backend backendId = get<0>(get<5>(GetParam())); Target targetId = get<1>(get<5>(GetParam())); + if ((inSize.width == 5 || inSize.height == 5) && targetId == DNN_TARGET_MYRIAD && + nrmType == "ACROSS_CHANNELS") + throw SkipTestException("This test case is disabled"); + LayerParams lp; lp.set("norm_region", nrmType); lp.set("local_size", localSize); diff --git a/modules/dnn/test/test_onnx_importer.cpp b/modules/dnn/test/test_onnx_importer.cpp index acc8337304..6d556f3e03 100644 --- a/modules/dnn/test/test_onnx_importer.cpp +++ b/modules/dnn/test/test_onnx_importer.cpp @@ -73,6 +73,7 @@ TEST_P(Test_ONNX_layers, Deconvolution) testONNXModels("deconvolution"); testONNXModels("two_deconvolution"); testONNXModels("deconvolution_group"); + testONNXModels("deconvolution_output_shape"); } TEST_P(Test_ONNX_layers, Dropout) diff --git a/modules/imgproc/src/sumpixels.avx512_skx.cpp b/modules/imgproc/src/sumpixels.avx512_skx.cpp index 511123b4f8..2c009a352c 100644 --- a/modules/imgproc/src/sumpixels.avx512_skx.cpp +++ b/modules/imgproc/src/sumpixels.avx512_skx.cpp @@ -13,7 +13,6 @@ namespace { // Anonymous namespace to avoid exposing the implementation classes // NOTE: Look at the bottom of the file for the entry-point function for external callers // -// TODO: Add support for 1 channel input (WIP: currently hitting hardware glassjaw) template class IntegralCalculator; template @@ -191,51 +190,55 @@ public: } - // The calculate_integral function referenced here must be implemented in the templated derivatives - // because the algorithm depends heavily on the number of channels in the image - // This is the incomplete definition (just the prototype) here. - // static CV_ALWAYS_INLINE - __m512d calculate_integral(__m512i src_longs, const __m512d above_values, __m512i &accumulator); - - - static CV_ALWAYS_INLINE - __m512i read_64_bytes(const __m512i *srcs, __mmask64 data_mask) { + __m512i read_64_bytes(const __m512i *srcs, const __mmask64 data_mask) { return _mm512_maskz_loadu_epi8(data_mask, srcs); } static CV_ALWAYS_INLINE - __m128i extract_lower_16bytes(__m512i src_64byte_chunk) { + __m128i extract_lower_16bytes(const __m512i src_64byte_chunk) { return _mm512_extracti64x2_epi64(src_64byte_chunk, 0x0); } static CV_ALWAYS_INLINE - __m512i convert_lower_8bytes_to_longs(__m128i src_16bytes) { + __m512i convert_lower_8bytes_to_longs(const __m128i src_16bytes) { return _mm512_cvtepu8_epi64(src_16bytes); } static CV_ALWAYS_INLINE - __m512i square_m512(__m512i src_longs) { + __m512i square_m512(const __m512i src_longs) { return _mm512_mullo_epi64(src_longs, src_longs); } static CV_ALWAYS_INLINE - __m128i shift_right_8_bytes(__m128i src_16bytes) { + __m128i shift_right_8_bytes(const __m128i src_16bytes) { return _mm_maskz_compress_epi64(2, src_16bytes); } - static CV_ALWAYS_INLINE - __m512i shift_right_16_bytes(__m512i src_64byte_chunk) { + __m512i shift_right_16_bytes(const __m512i src_64byte_chunk) { return _mm512_maskz_compress_epi64(0xFC, src_64byte_chunk); } + static CV_ALWAYS_INLINE + __m512i m512_hadd(const __m512i a){ + return _mm512_add_epi64(_mm512_maskz_compress_epi64(0xAA, a), _mm512_maskz_compress_epi64(0x55, a)); + } + + + // The calculate_integral function referenced here must be implemented in the templated derivatives + // because the algorithm depends heavily on the number of channels in the image + // This is the incomplete definition (just the prototype) here. + // + static CV_ALWAYS_INLINE + __m512d calculate_integral(const __m512i src_longs, const __m512d above_values, __m512i &accumulator); + }; @@ -246,7 +249,7 @@ public: // // The function prototype that needs to be implemented is: // -// __m512d calculate_integral(__m512i src_longs, const __m512d above_values, __m512i &accumulator){ ... } +// __m512d calculate_integral(const __m512i src_longs, const __m512d above_values, __m512i &accumulator){ ... } // // Description of parameters: // INPUTS: @@ -265,12 +268,72 @@ public: // Below here is the channel specific implementation // +//======================================== +// 1 Channel Integral Implementation +//======================================== +template<> +CV_ALWAYS_INLINE +__m512d IntegralCalculator < 1 > ::calculate_integral(const __m512i src_longs, const __m512d above_values, __m512i &accumulator) +{ + // One channel support is implemented differently than 2, 3, or 4 channel + // One channel support has more horizontal operations that cannot be made vertical without losing performance + // The logical operations needed look like: + // Vertical LANES : |7|6|5|4|3|2|1|0| + // src_longs : |H|G|F|E|D|C|B|A| + // shift_by_1 : + |G|F|E|D|C|B|A| | + // shift_by_2 : + |F|E|D|C|B|A| | | + // shift_by_3 : + |E|D|C|B|A| | | | + // shift_by_4 : + |D|C|B|A| | | | | + // shift_by_5 : + |C|B|A| | | | | | + // shift_by_6 : + |B|A| | | | | | | + // shift_by_7 : + |A| | | | | | | | + // carry_over_idxs : + |7|7|7|7|7|7|7|7| (index position of result from previous iteration) + // = integral + // + // If we do this vertically we end up losing performance because of the number of operations. We will instead + // do a horizontal add tree to create the vertical sections we need as a tree + // Vertical Lanes: | 7 | 6 | 5 | 4 | 3 | 2 | 1 | 0 | + // src_longs: | H | G | F | E | D | C | B | A | + // horiz_sum_1: | | | | | G+H | E+F | C+D | A+B | + // horiz_sum_2: | | | | | | | EFGH | ABCD | + // + const __m512i horiz_sum_1 = m512_hadd(src_longs); // indexes for the permutes below (3,2,1,0) = (GH, EF, CD, AB) + const __m512i horiz_sum_2 = m512_hadd(horiz_sum_1); // indexes for the permutes below (9, 8) = (EFGH, ABCD) + + // Then we can use the partial sums by looking at the vertical stacks above and realize that, for example + // ABCD appears vertically in lanes 7, 6, 5, 4, and 3 so we will permute the values so that all partial products + // appear in the right lanes. and sum them up along with the carry over value from the accumulator. So we setup + // the lanes like: + // Vertical Lanes: | 7 | 6 | 5 | 4 | 3 | 2 | 1 | 0 | + // s1 : | 0 | G | 0 | E | 0 | C | 0 | A | + // s2 : | ABCD | ABCD | ABCD | ABCD | ABCD | AB | AB | 0 | + // s3 : | EFGH | EF | EF | 0 | 0 | 0 | 0 | 0 | + // +------+------+------+------+------+------+------+------+ + // sum : | A..H | A..G | A..F | A..E | A..D | A..C | A..B | A | Integral :-) + // + const __m512i s1 = _mm512_maskz_mov_epi64(0x55, src_longs); // 0 G 0 E 0 D 0 C 0 A + const __m512i s2 = _mm512_permutex2var_epi64(horiz_sum_1, _mm512_set_epi64(8,8,8,8,8,0,0,4), horiz_sum_2); + const __m512i s3 = _mm512_permutex2var_epi64(horiz_sum_1, _mm512_set_epi64(9,2,2,4,4,4,4,4), horiz_sum_2); + + // Now we use the rolling sum from the previous iteration from accumulator and replicate it into carry_over + // And sum everything up into the accumulator + // + const __m512i carry_over = _mm512_permutex2var_epi64(accumulator, _mm512_set_epi64(7,7,7,7,7,7,7,7), accumulator); + accumulator = _mm512_add_epi64(_mm512_add_epi64(s2, s3), _mm512_add_epi64(carry_over, s1)); + + // Convert to double precision and store + // + __m512d integral_pd = _mm512_add_pd(_mm512_cvtepu64_pd(accumulator), above_values); + return integral_pd; +} + + //======================================== // 2 Channel Integral Implementation //======================================== template<> CV_ALWAYS_INLINE -__m512d IntegralCalculator < 2 > ::calculate_integral(__m512i src_longs, const __m512d above_values, __m512i &accumulator) +__m512d IntegralCalculator < 2 > ::calculate_integral(const __m512i src_longs, const __m512d above_values, __m512i &accumulator) { __m512i carryover_idxs = _mm512_set_epi64(7, 6, 7, 6, 7, 6, 7, 6); @@ -300,12 +363,13 @@ __m512d IntegralCalculator < 2 > ::calculate_integral(__m512i src_longs, const _ return integral_pd; } + //======================================== // 3 Channel Integral Implementation //======================================== template<> CV_ALWAYS_INLINE -__m512d IntegralCalculator < 3 > ::calculate_integral(__m512i src_longs, const __m512d above_values, __m512i &accumulator) +__m512d IntegralCalculator < 3 > ::calculate_integral(const __m512i src_longs, const __m512d above_values, __m512i &accumulator) { __m512i carryover_idxs = _mm512_set_epi64(6, 5, 7, 6, 5, 7, 6, 5); @@ -338,7 +402,7 @@ __m512d IntegralCalculator < 3 > ::calculate_integral(__m512i src_longs, const _ //======================================== template<> CV_ALWAYS_INLINE -__m512d IntegralCalculator < 4 > ::calculate_integral(__m512i src_longs, const __m512d above_values, __m512i &accumulator) +__m512d IntegralCalculator < 4 > ::calculate_integral(const __m512i src_longs, const __m512d above_values, __m512i &accumulator) { __m512i carryover_idxs = _mm512_set_epi64(7, 6, 5, 4, 7, 6, 5, 4); @@ -376,18 +440,23 @@ void calculate_integral_avx512(const uchar *src, size_t _srcstep, int width, int height, int cn) { switch(cn){ + case 1: { + IntegralCalculator< 1 > calculator; + calculator.calculate_integral_avx512(src, _srcstep, sum, _sumstep, sqsum, _sqsumstep, width, height); + break; + } case 2: { - IntegralCalculator<2> calculator; + IntegralCalculator< 2 > calculator; calculator.calculate_integral_avx512(src, _srcstep, sum, _sumstep, sqsum, _sqsumstep, width, height); break; } case 3: { - IntegralCalculator<3> calculator; + IntegralCalculator< 3 > calculator; calculator.calculate_integral_avx512(src, _srcstep, sum, _sumstep, sqsum, _sqsumstep, width, height); break; } case 4: { - IntegralCalculator<4> calculator; + IntegralCalculator< 4 > calculator; calculator.calculate_integral_avx512(src, _srcstep, sum, _sumstep, sqsum, _sqsumstep, width, height); } } diff --git a/modules/imgproc/src/sumpixels.cpp b/modules/imgproc/src/sumpixels.cpp index ca8b99afde..8531fc61f4 100755 --- a/modules/imgproc/src/sumpixels.cpp +++ b/modules/imgproc/src/sumpixels.cpp @@ -77,7 +77,7 @@ struct Integral_SIMD { #if CV_TRY_AVX512_SKX CV_UNUSED(_tiltedstep); // TODO: Add support for 1 channel input (WIP) - if (CV_CPU_HAS_SUPPORT_AVX512_SKX && !tilted && ((cn >= 2) && (cn <= 4))){ + if (CV_CPU_HAS_SUPPORT_AVX512_SKX && !tilted && (cn <= 4)){ opt_AVX512_SKX::calculate_integral_avx512(src, _srcstep, sum, _sumstep, sqsum, _sqsumstep, width, height, cn); return true; diff --git a/samples/opencl/opencl-opencv-interop.cpp b/samples/opencl/opencl-opencv-interop.cpp index 816550feaf..d3b15668a9 100644 --- a/samples/opencl/opencl-opencv-interop.cpp +++ b/samples/opencl/opencl-opencv-interop.cpp @@ -19,6 +19,7 @@ #define CL_USE_DEPRECATED_OPENCL_2_0_APIS // eliminate build warning #ifdef __APPLE__ +#define CL_SILENCE_DEPRECATION #include #else #include