From 0fe17eeb68a5b213e95559c790fb07098e625392 Mon Sep 17 00:00:00 2001 From: Brad Kelly Date: Fri, 22 Feb 2019 16:57:50 -0800 Subject: [PATCH 1/8] Implementing AVX512 Support for 1 channel mats for CV_64F format --- modules/imgproc/src/sumpixels.avx512_skx.cpp | 115 +++++++++++++++---- modules/imgproc/src/sumpixels.cpp | 2 +- 2 files changed, 93 insertions(+), 24 deletions(-) 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; From 097fc1a27132fd1453306138494fbbdd5e7427e2 Mon Sep 17 00:00:00 2001 From: iPanda Date: Sun, 24 Mar 2019 23:34:09 +0800 Subject: [PATCH 2/8] Merge pull request #13972 from Mainvooid:add_cuda_support_for_D3D11_interop * Add CUDA support for D3D11 interop. #13888 color_detail.hpp: fixed build error : dynamic initialization is not supported for a __constant__ variable. directx.cpp: Add CUDA support(cl_nv_d3d11_sharing) for D3D11 interop. #13888 Update directx.cpp Format adjustment. Update directx.cpp fix error. Update directx.cpp Format adjustment Update directx.cpp fix trailing whitespace. fix format errors convert indentation to spaces . Trim trailing whitespace. Add information about source of cl_d3d11_ext.h Avoid unrelated changes. Increase compile-time conditional judgment. Increase the judgment of whether the OCL device has the required extensions at compile time. Add compilation option `HAVE_CLNVEXT`.Check CL support in runtime. Check result of `clGetExtensionFunctionAddressForPlatform` for KHR is invalid.It always can get the address(from OpenCL.dll),So I check NV support(from nvopencl64.dll) before KHR when `HAVE_CLNVEXT` is enabled. Delete cl_d3d11_ext.h Modified parameter list fix "cannot open include file: 'CL/cl_d3d11_ext.h'" remove not referenced var fix C2143: syntax error Improve compile-time judgment. dlrectx.cpp Modify the detection order. initializeContextFromD3D11Device: ``` // try with NV(Need to check it first) // try with KHR ``` fix warnig C4100 Revert "fix warnig C4100" This reverts commit 76e5becb67780071d0cbde61cc4f5f807ad7c5ac. fix warning C4100 fix warning C4505 Format alignment Format adjustment and automatically detect header files. Automatically detect header files when users are not configured or configuration errors occur. avoid unrelated changes. Update .cmake Update .cmake * fix build errors * fix warning:defined but not used * Revert "fix warning:defined but not used" This reverts commit 7ab3537cd070f89b15bc2926e4ac9ec74c84a122. * fix warning:defined but not used * fix build error for mac * fix build error for win * optimizing branch judgment * Revert "optimizing branch judgment" This reverts commit 88b72b870ec13fd26f64a5ac374484c5cfe80854. * fix warning C4702: unreachable code * remove unused code * Fix problems that may lead to undefined behavior * Add status check * fix error C2664,C2665 : cannot convert argument * Format adjustment VSCODE will automatically format the indentation to 4 spaces in some situation. * fix error C2440 * fix error C2440 * add cl_d3d11_ext.h * Format adjustment * remove unnecessary checks --- 3rdparty/include/opencl/1.2/CL/cl_d3d11_ext.h | 122 ++++ CMakeLists.txt | 4 + cmake/OpenCVDetectOpenCL.cmake | 9 +- cmake/templates/cvconfig.h.in | 3 + modules/core/src/directx.cpp | 555 +++++++++++++++--- modules/core/src/directx.inc.hpp | 3 + 6 files changed, 608 insertions(+), 88 deletions(-) create mode 100644 3rdparty/include/opencl/1.2/CL/cl_d3d11_ext.h 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 c4a74b2230..59f6fd02d6 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -387,6 +387,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_INTELPERC "Include Intel Perceptual Computing support" OFF VISIBLE_IF WIN32 AND NOT WINRT VERIFY HAVE_INTELPERC) @@ -1614,6 +1617,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 0fcbeaa2fc..f7656e9d71 100644 --- a/cmake/templates/cvconfig.h.in +++ b/cmake/templates/cvconfig.h.in @@ -136,6 +136,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/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 From 714b38360f2fbb01c6c4a77829067472a43ca47f Mon Sep 17 00:00:00 2001 From: Dmitry Kurtaev Date: Mon, 25 Mar 2019 12:49:02 +0300 Subject: [PATCH 3/8] Extend LRN layer support with IE backend --- modules/dnn/src/layers/lrn_layer.cpp | 2 +- modules/dnn/test/test_halide_layers.cpp | 4 ++++ 2 files changed, 5 insertions(+), 1 deletion(-) diff --git a/modules/dnn/src/layers/lrn_layer.cpp b/modules/dnn/src/layers/lrn_layer.cpp index c101d30414..21daa9f65f 100644 --- a/modules/dnn/src/layers/lrn_layer.cpp +++ b/modules/dnn/src/layers/lrn_layer.cpp @@ -91,7 +91,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; } 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); From 41da3ef1d2ee8022368e0ef963104405a48ab8c3 Mon Sep 17 00:00:00 2001 From: Maksim Shabunin Date: Mon, 25 Mar 2019 16:44:08 +0300 Subject: [PATCH 4/8] Fixed cvdef.h for MSVC C users --- modules/core/include/opencv2/core/cvdef.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/modules/core/include/opencv2/core/cvdef.h b/modules/core/include/opencv2/core/cvdef.h index deb08fa003..be7da7a4f6 100644 --- a/modules/core/include/opencv2/core/cvdef.h +++ b/modules/core/include/opencv2/core/cvdef.h @@ -565,7 +565,7 @@ Cv64suf; \****************************************************************************************/ #ifndef CV_CXX_STD_ARRAY -# if __cplusplus >= 201103L || (defined(_MSC_VER) && _MSC_VER >= 1900/*MSVS 2015*/) +# if __cplusplus >= 201103L || (defined(__cplusplus) && defined(_MSC_VER) && _MSC_VER >= 1900/*MSVS 2015*/) # define CV_CXX_STD_ARRAY 1 # include # endif From c443a3d6579b72c7f61db1c566cdea01a71554d0 Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Mon, 25 Mar 2019 19:24:49 +0300 Subject: [PATCH 5/8] cmake: set CMP0068 => NEW --- CMakeLists.txt | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/CMakeLists.txt b/CMakeLists.txt index 59f6fd02d6..0ed2cc0de7 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) From 6686559c7086b5267945753a3c04aa560110da7b Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Tue, 26 Mar 2019 13:11:53 +0300 Subject: [PATCH 6/8] ocl: define CL_SILENCE_DEPRECATION on MacOSX --- .../opencv2/core/opencl/runtime/autogenerated/opencl_core.hpp | 1 + .../src/opencl/runtime/generator/template/opencl_core.hpp.in | 1 + modules/core/src/opencl/runtime/opencl_core.cpp | 1 + samples/opencl/opencl-opencv-interop.cpp | 1 + 4 files changed, 4 insertions(+) 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/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/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 From 3bcff7f941fdf4d2522086becef5e0bfef44cb9a Mon Sep 17 00:00:00 2001 From: Kohei Yoshida Date: Thu, 28 Feb 2019 14:16:27 -0500 Subject: [PATCH 7/8] Make use of explicit output shape for the de-convolution layer. This is relevant to cases where the net is loaded from an ONNX model containing ConvTranspose layers, which may include the output_shape attribute. c.f. https://github.com/onnx/onnx/blob/master/docs/Operators.md#ConvTranspose --- modules/dnn/src/onnx/onnx_importer.cpp | 31 ++++++++++++++++++++++++++ 1 file changed, 31 insertions(+) diff --git a/modules/dnn/src/onnx/onnx_importer.cpp b/modules/dnn/src/onnx/onnx_importer.cpp index d658677d4c..f20f56547b 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") { From 212420516f2f4390d350f59efcfa4a760fca83ab Mon Sep 17 00:00:00 2001 From: Kohei Yoshida Date: Thu, 7 Mar 2019 11:42:57 -0500 Subject: [PATCH 8/8] Add a test case for handling output_shape property for deconvolution layer. --- modules/dnn/test/test_onnx_importer.cpp | 1 + 1 file changed, 1 insertion(+) 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)