Merge pull request #4238 from vladimir-dudnik:d3d11-nv12-interop

modules/core/include/opencv2/core/directx.hpp

@@ -109,19 +109,21 @@ CV_EXPORTS Context& initializeContextFromDirect3DDevice9(IDirect3DDevice9* pDire
 //! @addtogroup core_directx
 //! @{
 
-//! @brief Converts InputArray to ID3D11Texture2D
+//! @brief Converts InputArray to ID3D11Texture2D. If destination texture format is DXGI_FORMAT_NV12 then
+//!        input UMat expected to be in BGR format and data will be downsampled and color-converted to NV12.
 //
-//! @note Note: function does memory copy from src to
+//! @note Note: Destination texture must be allocated by application. Function does memory copy from src to
 //!             pD3D11Texture2D
 //
 //! @param src - source InputArray
 //! @param pD3D11Texture2D - destination D3D11 texture
 CV_EXPORTS void convertToD3D11Texture2D(InputArray src, ID3D11Texture2D* pD3D11Texture2D);
 
-//! @brief Converts ID3D11Texture2D to OutputArray
+//! @brief Converts ID3D11Texture2D to OutputArray. If input texture format is DXGI_FORMAT_NV12 then
+//!        data will be upsampled and color-converted to BGR format.
 //
-//! @note Note: function does memory copy from pD3D11Texture2D
+//! @note Note: Destination matrix will be re-allocated if it has not enough memory to match texture size.
-//!             to dst
+//!             function does memory copy from pD3D11Texture2D to dst
 //
 //! @param pD3D11Texture2D - source D3D11 texture
 //! @param dst             - destination OutputArray

modules/core/src/directx.cpp

@@ -44,6 +44,7 @@
 #include "opencv2/core.hpp"
 #include "opencv2/core/ocl.hpp"
 #include "opencv2/core/directx.hpp"
+#include "opencl_kernels_core.hpp"
 
 #ifdef HAVE_DIRECTX
 #include <vector>
@@ -167,6 +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:
+    case DXGI_FORMAT_NV12: return CV_8UC3;
     default: break;
     }
     return errorType;
@@ -701,6 +703,59 @@ static void __OpenCLinitializeD3D11()
 }
 #endif // defined(HAVE_DIRECTX) && defined(HAVE_OPENCL)
 
+} // namespace directx
+
+
+namespace ocl {
+
+#if defined(HAVE_DIRECTX) && defined(HAVE_OPENCL)
+static
+bool ocl_convert_nv12_to_bgr(
+    cl_mem clImageY,
+    cl_mem clImageUV,
+    cl_mem clBuffer,
+    int step,
+    int cols,
+    int rows)
+{
+    ocl::Kernel k;
+    k.create("YUV2BGR_NV12_8u", cv::ocl::core::cvtclr_dx_oclsrc, "");
+    if (k.empty())
+        return false;
+
+    k.args(clImageY, clImageUV, clBuffer, step, cols, rows);
+
+    size_t globalsize[] = { cols, rows };
+    return k.run(2, globalsize, 0, false);
+}
+
+
+static
+bool ocl_convert_bgr_to_nv12(
+    cl_mem clBuffer,
+    int step,
+    int cols,
+    int rows,
+    cl_mem clImageY,
+    cl_mem clImageUV)
+{
+    ocl::Kernel k;
+    k.create("BGR2YUV_NV12_8u", cv::ocl::core::cvtclr_dx_oclsrc, "");
+    if (k.empty())
+        return false;
+
+    k.args(clBuffer, step, cols, rows, clImageY, clImageUV);
+
+    size_t globalsize[] = { cols, rows };
+    return k.run(2, globalsize, 0, false);
+}
+#endif // HAVE_DIRECTX && HAVE_OPENCL
+
+} // namespace ocl
+
+
+namespace directx {
+
 void convertToD3D11Texture2D(InputArray src, ID3D11Texture2D* pD3D11Texture2D)
 {
     (void)src; (void)pD3D11Texture2D;
@@ -719,33 +774,63 @@ void convertToD3D11Texture2D(InputArray src, ID3D11Texture2D* pD3D11Texture2D)
     Size srcSize = src.size();
     CV_Assert(srcSize.width == (int)desc.Width && srcSize.height == (int)desc.Height);
 
-    using namespace cv::ocl;
-    Context& ctx = Context::getDefault();
-    cl_context context = (cl_context)ctx.ptr();
-
     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 = clCreateFromD3D11Texture2DKHR(context, CL_MEM_WRITE_ONLY, pD3D11Texture2D, 0, &status);
+    cl_mem clImage = 0;
+    cl_mem clImageUV = 0;
+
+    clImage = clCreateFromD3D11Texture2DKHR(context, CL_MEM_WRITE_ONLY, pD3D11Texture2D, 0, &status);
     if (status != CL_SUCCESS)
         CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clCreateFromD3D11Texture2DKHR failed");
 
-    cl_mem clBuffer = (cl_mem)u.handle(ACCESS_READ);
+    if(DXGI_FORMAT_NV12 == desc.Format)
+    {
+        clImageUV = clCreateFromD3D11Texture2DKHR(context, CL_MEM_WRITE_ONLY, pD3D11Texture2D, 1, &status);
+        if (status != CL_SUCCESS)
+            CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clCreateFromD3D11Texture2DKHR failed");
+    }
 
     cl_command_queue q = (cl_command_queue)Queue::getDefault().ptr();
+
     status = clEnqueueAcquireD3D11ObjectsKHR(q, 1, &clImage, 0, NULL, NULL);
     if (status != CL_SUCCESS)
         CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueAcquireD3D11ObjectsKHR failed");
+
+    if(DXGI_FORMAT_NV12 == desc.Format)
+    {
+        status = clEnqueueAcquireD3D11ObjectsKHR(q, 1, &clImageUV, 0, NULL, NULL);
+        if (status != CL_SUCCESS)
+            CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueAcquireD3D11ObjectsKHR 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 = clEnqueueReleaseD3D11ObjectsKHR(q, 1, &clImageUV, 0, NULL, NULL);
+        if (status != CL_SUCCESS)
+            CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueReleaseD3D11ObjectsKHR failed");
+    }
+    else
+    {
         size_t offset = 0; // TODO
-    size_t dst_origin[3] = {0, 0, 0};
+        size_t origin[3] = { 0, 0, 0 };
         size_t region[3] = { u.cols, u.rows, 1 };
-    status = clEnqueueCopyBufferToImage(q, clBuffer, clImage, offset, dst_origin, region, 0, NULL, NULL);
+
+        status = clEnqueueCopyBufferToImage(q, clBuffer, clImage, offset, origin, region, 0, NULL, NULL);
         if (status != CL_SUCCESS)
             CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueCopyBufferToImage failed");
+    }
+
     status = clEnqueueReleaseD3D11ObjectsKHR(q, 1, &clImage, 0, NULL, NULL);
     if (status != CL_SUCCESS)
         CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueReleaseD3D11ObjectsKHR failed");
@@ -757,11 +842,20 @@ void convertToD3D11Texture2D(InputArray src, ID3D11Texture2D* pD3D11Texture2D)
     status = clReleaseMemObject(clImage); // TODO RAII
     if (status != CL_SUCCESS)
         CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clReleaseMem failed");
+
+    if(DXGI_FORMAT_NV12 == desc.Format)
+    {
+        status = clReleaseMemObject(clImageUV);
+        if (status != CL_SUCCESS)
+            CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clReleaseMem failed");
+    }
 #else
     // TODO memcpy
     NO_OPENCL_SUPPORT_ERROR;
 #endif
 }
+
+
 void convertFromD3D11Texture2D(ID3D11Texture2D* pD3D11Texture2D, OutputArray dst)
 {
     (void)pD3D11Texture2D; (void)dst;
@@ -776,10 +870,6 @@ void convertFromD3D11Texture2D(ID3D11Texture2D* pD3D11Texture2D, OutputArray dst
     int textureType = getTypeFromDXGI_FORMAT(desc.Format);
     CV_Assert(textureType >= 0);
 
-    using namespace cv::ocl;
-    Context& ctx = Context::getDefault();
-    cl_context context = (cl_context)ctx.ptr();
-
     // 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();
@@ -788,23 +878,57 @@ void convertFromD3D11Texture2D(ID3D11Texture2D* pD3D11Texture2D, OutputArray dst
     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 = clCreateFromD3D11Texture2DKHR(context, CL_MEM_READ_ONLY, pD3D11Texture2D, 0, &status);
+    cl_mem clImage = 0;
+    cl_mem clImageUV = 0;
+
+    clImage = clCreateFromD3D11Texture2DKHR(context, CL_MEM_READ_ONLY, pD3D11Texture2D, 0, &status);
     if (status != CL_SUCCESS)
         CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clCreateFromD3D11Texture2DKHR failed");
 
-    cl_mem clBuffer = (cl_mem)u.handle(ACCESS_READ);
+    if(DXGI_FORMAT_NV12 == desc.Format)
+    {
+        clImageUV = clCreateFromD3D11Texture2DKHR(context, CL_MEM_READ_ONLY, pD3D11Texture2D, 1, &status);
+        if (status != CL_SUCCESS)
+            CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clCreateFromD3D11Texture2DKHR failed");
+    }
 
     cl_command_queue q = (cl_command_queue)Queue::getDefault().ptr();
+
     status = clEnqueueAcquireD3D11ObjectsKHR(q, 1, &clImage, 0, NULL, NULL);
     if (status != CL_SUCCESS)
         CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueAcquireD3D11ObjectsKHR failed");
+
+    if(DXGI_FORMAT_NV12 == desc.Format)
+    {
+        status = clEnqueueAcquireD3D11ObjectsKHR(q, 1, &clImageUV, 0, NULL, NULL);
+        if (status != CL_SUCCESS)
+            CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueAcquireD3D11ObjectsKHR 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 = clEnqueueReleaseD3D11ObjectsKHR(q, 1, &clImageUV, 0, NULL, NULL);
+        if (status != CL_SUCCESS)
+            CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueReleaseD3D11ObjectsKHR failed");
+    }
+    else
+    {
         size_t offset = 0; // TODO
-    size_t src_origin[3] = {0, 0, 0};
+        size_t origin[3] = { 0, 0, 0 };
         size_t region[3] = { u.cols, u.rows, 1 };
-    status = clEnqueueCopyImageToBuffer(q, clImage, clBuffer, src_origin, region, offset, 0, NULL, NULL);
+
+        status = clEnqueueCopyImageToBuffer(q, clImage, clBuffer, origin, region, offset, 0, NULL, NULL);
         if (status != CL_SUCCESS)
             CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueCopyImageToBuffer failed");
+    }
+
     status = clEnqueueReleaseD3D11ObjectsKHR(q, 1, &clImage, 0, NULL, NULL);
     if (status != CL_SUCCESS)
         CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clEnqueueReleaseD3D11ObjectsKHR failed");
@@ -816,6 +940,13 @@ void convertFromD3D11Texture2D(ID3D11Texture2D* pD3D11Texture2D, OutputArray dst
     status = clReleaseMemObject(clImage); // TODO RAII
     if (status != CL_SUCCESS)
         CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clReleaseMem failed");
+
+    if(DXGI_FORMAT_NV12 == desc.Format)
+    {
+        status = clReleaseMemObject(clImageUV);
+        if (status != CL_SUCCESS)
+            CV_Error(cv::Error::OpenCLApiCallError, "OpenCL: clReleaseMem failed");
+    }
 #else
     // TODO memcpy
     NO_OPENCL_SUPPORT_ERROR;

modules/core/src/opencl/cvtclr_dx.cl

@@ -0,0 +1,207 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
+// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
+// Copyright (C) 2013, OpenCV Foundation, all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// @Authors
+//    Jia Haipeng, jiahaipeng95@gmail.com
+//
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors as is and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the copyright holders or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#ifdef DOUBLE_SUPPORT
+#ifdef cl_amd_fp64
+#pragma OPENCL EXTENSION cl_amd_fp64:enable
+#elif defined cl_khr_fp64
+#pragma OPENCL EXTENSION cl_khr_fp64:enable
+#endif
+#endif
+
+#ifdef INTEL_DEVICE
+#pragma OPENCL FP_CONTRACT ON
+#pragma OPENCL FP_FAST_FMAF ON
+#pragma OPENCL FP_FAST_FMA ON
+#endif
+
+
+static
+__constant
+float c_YUV2RGBCoeffs_420[5] =
+{
+     1.163999557f,
+     2.017999649f,
+    -0.390999794f,
+    -0.812999725f,
+     1.5959997177f
+};
+
+static __constant float CV_8U_MAX         = 255.0f;
+static __constant float CV_8U_HALF        = 128.0f;
+static __constant float BT601_BLACK_RANGE = 16.0f;
+static __constant float CV_8U_SCALE       = 1.0f / 255.0f;
+static __constant float d1                = BT601_BLACK_RANGE / CV_8U_MAX;
+static __constant float d2                = CV_8U_HALF / CV_8U_MAX;
+
+#define NCHANNELS 3
+
+__kernel
+void YUV2BGR_NV12_8u(
+    read_only image2d_t imgY,
+    read_only image2d_t imgUV,
+    __global unsigned char* pBGR,
+   int bgrStep,
+   int cols,
+   int rows)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (x < cols)
+    {
+        if (y < rows)
+        {
+            __global uchar* pDstRow1 = pBGR + mad24(y, bgrStep, mad24(x, NCHANNELS, 0));
+            __global uchar* pDstRow2 = pDstRow1 + bgrStep;
+
+            float4 Y1 = read_imagef(imgY, (int2)(x+0, y+0));
+            float4 Y2 = read_imagef(imgY, (int2)(x+1, y+0));
+            float4 Y3 = read_imagef(imgY, (int2)(x+0, y+1));
+            float4 Y4 = read_imagef(imgY, (int2)(x+1, y+1));
+
+            float4 UV = read_imagef(imgUV, (int2)(x/2, y/2)) - d2;
+
+            __constant float* coeffs = c_YUV2RGBCoeffs_420;
+
+            Y1 = max(0.f, Y1 - d1) * coeffs[0];
+            Y2 = max(0.f, Y2 - d1) * coeffs[0];
+            Y3 = max(0.f, Y3 - d1) * coeffs[0];
+            Y4 = max(0.f, Y4 - d1) * coeffs[0];
+
+            float ruv = fma(coeffs[4], UV.y, 0.0f);
+            float guv = fma(coeffs[3], UV.y, fma(coeffs[2], UV.x, 0.0f));
+            float buv = fma(coeffs[1], UV.x, 0.0f);
+
+            float R1 = (Y1.x + ruv) * CV_8U_MAX;
+            float G1 = (Y1.x + guv) * CV_8U_MAX;
+            float B1 = (Y1.x + buv) * CV_8U_MAX;
+
+            float R2 = (Y2.x + ruv) * CV_8U_MAX;
+            float G2 = (Y2.x + guv) * CV_8U_MAX;
+            float B2 = (Y2.x + buv) * CV_8U_MAX;
+
+            float R3 = (Y3.x + ruv) * CV_8U_MAX;
+            float G3 = (Y3.x + guv) * CV_8U_MAX;
+            float B3 = (Y3.x + buv) * CV_8U_MAX;
+
+            float R4 = (Y4.x + ruv) * CV_8U_MAX;
+            float G4 = (Y4.x + guv) * CV_8U_MAX;
+            float B4 = (Y4.x + buv) * CV_8U_MAX;
+
+            pDstRow1[0*NCHANNELS + 0] = convert_uchar_sat(B1);
+            pDstRow1[0*NCHANNELS + 1] = convert_uchar_sat(G1);
+            pDstRow1[0*NCHANNELS + 2] = convert_uchar_sat(R1);
+
+            pDstRow1[1*NCHANNELS + 0] = convert_uchar_sat(B2);
+            pDstRow1[1*NCHANNELS + 1] = convert_uchar_sat(G2);
+            pDstRow1[1*NCHANNELS + 2] = convert_uchar_sat(R2);
+
+            pDstRow2[0*NCHANNELS + 0] = convert_uchar_sat(B3);
+            pDstRow2[0*NCHANNELS + 1] = convert_uchar_sat(G3);
+            pDstRow2[0*NCHANNELS + 2] = convert_uchar_sat(R3);
+
+            pDstRow2[1*NCHANNELS + 0] = convert_uchar_sat(B4);
+            pDstRow2[1*NCHANNELS + 1] = convert_uchar_sat(G4);
+            pDstRow2[1*NCHANNELS + 2] = convert_uchar_sat(R4);
+        }
+    }
+}
+
+
+static
+__constant float c_RGB2YUVCoeffs_420[8] =
+{
+     0.256999969f,  0.50399971f,   0.09799957f,   -0.1479988098f,
+    -0.2909994125f, 0.438999176f, -0.3679990768f, -0.0709991455f
+};
+
+
+__kernel
+void BGR2YUV_NV12_8u(
+    __global unsigned char* pBGR,
+    int bgrStep,
+    int cols,
+    int rows,
+    write_only image2d_t imgY,
+    write_only image2d_t imgUV)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (x < cols)
+    {
+        if (y < rows)
+        {
+            __global const uchar* pSrcRow1 = pBGR + mad24(y, bgrStep, mad24(x, NCHANNELS, 0));
+            __global const uchar* pSrcRow2 = pSrcRow1 + bgrStep;
+
+            float4 src_pix1 = convert_float4(vload4(0, pSrcRow1 + 0*NCHANNELS)) * CV_8U_SCALE;
+            float4 src_pix2 = convert_float4(vload4(0, pSrcRow1 + 1*NCHANNELS)) * CV_8U_SCALE;
+            float4 src_pix3 = convert_float4(vload4(0, pSrcRow2 + 0*NCHANNELS)) * CV_8U_SCALE;
+            float4 src_pix4 = convert_float4(vload4(0, pSrcRow2 + 1*NCHANNELS)) * CV_8U_SCALE;
+
+            __constant float* coeffs = c_RGB2YUVCoeffs_420;
+
+            float Y1 = fma(coeffs[0], src_pix1.z, fma(coeffs[1], src_pix1.y, fma(coeffs[2], src_pix1.x, d1)));
+            float Y2 = fma(coeffs[0], src_pix2.z, fma(coeffs[1], src_pix2.y, fma(coeffs[2], src_pix2.x, d1)));
+            float Y3 = fma(coeffs[0], src_pix3.z, fma(coeffs[1], src_pix3.y, fma(coeffs[2], src_pix3.x, d1)));
+            float Y4 = fma(coeffs[0], src_pix4.z, fma(coeffs[1], src_pix4.y, fma(coeffs[2], src_pix4.x, d1)));
+
+            float4 UV;
+            UV.x = fma(coeffs[3], src_pix1.z, fma(coeffs[4], src_pix1.y, fma(coeffs[5], src_pix1.x, d2)));
+            UV.y = fma(coeffs[5], src_pix1.z, fma(coeffs[6], src_pix1.y, fma(coeffs[7], src_pix1.x, d2)));
+
+            write_imagef(imgY, (int2)(x+0, y+0), Y1);
+            write_imagef(imgY, (int2)(x+1, y+0), Y2);
+            write_imagef(imgY, (int2)(x+0, y+1), Y3);
+            write_imagef(imgY, (int2)(x+1, y+1), Y4);
+
+            write_imagef(imgUV, (int2)((x/2), (y/2)), UV);
+        }
+    }
+}

samples/directx/d3d10_interop.cpp

@@ -135,7 +135,7 @@ public:
         if (!m_cap.read(m_frame_bgr))
             return -1;
 
-        cv::cvtColor(m_frame_bgr, m_frame_rgba, CV_RGB2BGRA);
+        cv::cvtColor(m_frame_bgr, m_frame_rgba, CV_BGR2RGBA);
 
         UINT subResource = ::D3D10CalcSubresource(0, 0, 1);
 
@@ -166,6 +166,9 @@ public:
             if (m_shutdown)
                 return 0;
 
+            // capture user input once
+            MODE mode = (m_mode == MODE_GPU_NV12) ? MODE_GPU_RGBA : m_mode;
+
             HRESULT r;
             ID3D10Texture2D* pSurface;
 
@@ -177,7 +180,7 @@ public:
 
             m_timer.start();
 
-            switch (m_mode)
+            switch (mode)
             {
                 case MODE_CPU:
                 {
@@ -214,7 +217,7 @@ public:
                     break;
                 }
 
-                case MODE_GPU:
+                case MODE_GPU_RGBA:
                 {
                     // process video frame on GPU
                     cv::UMat u;
@@ -227,7 +230,7 @@ public:
                         cv::blur(u, u, cv::Size(15, 15), cv::Point(-7, -7));
                     }
 
-                    cv::String strMode = cv::format("mode: %s", m_modeStr[MODE_GPU].c_str());
+                    cv::String strMode = cv::format("mode: %s", m_modeStr[MODE_GPU_RGBA].c_str());
                     cv::String strProcessing = m_demo_processing ? "blur frame" : "copy frame";
                     cv::String strTime = cv::format("time: %4.1f msec", m_timer.time(Timer::UNITS::MSEC));
                     cv::String strDevName = cv::format("OpenCL device: %s", m_oclDevName.c_str());

samples/directx/d3d11_interop.cpp

@@ -74,6 +74,14 @@ public:
             throw std::runtime_error("D3D11CreateDeviceAndSwapChain() failed!");
         }
 
+        m_nv12_available = true;
+        UINT fmt = 0;
+        r = m_pD3D11Dev->CheckFormatSupport(DXGI_FORMAT_NV12, &fmt);
+        if (FAILED(r))
+        {
+            m_nv12_available = false;
+        }
+
         r = m_pD3D11SwapChain->GetBuffer(0, __uuidof(ID3D11Texture2D), (LPVOID*)&m_pBackBuffer);
         if (FAILED(r))
         {
@@ -98,24 +106,70 @@ public:
 
         m_pD3D11Ctx->RSSetViewports(1, &viewport);
 
-        D3D11_TEXTURE2D_DESC desc;
+        m_pSurfaceRGBA = 0;
+        m_pSurfaceNV12 = 0;
 
-        desc.Width              = m_width;
+        D3D11_TEXTURE2D_DESC desc_rgba;
-        desc.Height             = m_height;
-        desc.MipLevels          = 1;
-        desc.ArraySize          = 1;
-        desc.Format             = DXGI_FORMAT_R8G8B8A8_UNORM;
-        desc.SampleDesc.Count   = 1;
-        desc.SampleDesc.Quality = 0;
-        desc.BindFlags          = D3D11_BIND_SHADER_RESOURCE;
-        desc.Usage              = D3D11_USAGE_DYNAMIC;
-        desc.CPUAccessFlags     = D3D11_CPU_ACCESS_WRITE;
-        desc.MiscFlags          = 0;
 
-        r = m_pD3D11Dev->CreateTexture2D(&desc, NULL, &m_pSurface);
+        desc_rgba.Width              = m_width;
+        desc_rgba.Height             = m_height;
+        desc_rgba.MipLevels          = 1;
+        desc_rgba.ArraySize          = 1;
+        desc_rgba.Format             = DXGI_FORMAT_R8G8B8A8_UNORM;
+        desc_rgba.SampleDesc.Count   = 1;
+        desc_rgba.SampleDesc.Quality = 0;
+        desc_rgba.BindFlags          = D3D11_BIND_SHADER_RESOURCE;
+        desc_rgba.Usage              = D3D11_USAGE_DYNAMIC;
+        desc_rgba.CPUAccessFlags     = D3D11_CPU_ACCESS_WRITE;
+        desc_rgba.MiscFlags          = 0;
+
+        r = m_pD3D11Dev->CreateTexture2D(&desc_rgba, 0, &m_pSurfaceRGBA);
         if (FAILED(r))
         {
-            throw std::runtime_error("Can't create texture with input image");
+            throw std::runtime_error("Can't create DX texture");
+        }
+
+        if(m_nv12_available)
+        {
+            D3D11_TEXTURE2D_DESC desc_nv12;
+
+            desc_nv12.Width              = m_width;
+            desc_nv12.Height             = m_height;
+            desc_nv12.MipLevels          = 1;
+            desc_nv12.ArraySize          = 1;
+            desc_nv12.Format             = DXGI_FORMAT_NV12;
+            desc_nv12.SampleDesc.Count   = 1;
+            desc_nv12.SampleDesc.Quality = 0;
+            desc_nv12.BindFlags          = D3D11_BIND_SHADER_RESOURCE;
+            desc_nv12.Usage              = D3D11_USAGE_DEFAULT;
+            desc_nv12.CPUAccessFlags     = 0;
+            desc_nv12.MiscFlags          = D3D11_RESOURCE_MISC_SHARED;
+
+            r = m_pD3D11Dev->CreateTexture2D(&desc_nv12, 0, &m_pSurfaceNV12);
+            if (FAILED(r))
+            {
+                throw std::runtime_error("Can't create DX NV12 texture");
+            }
+
+            D3D11_TEXTURE2D_DESC desc_nv12_cpu_copy;
+
+            desc_nv12_cpu_copy.Width              = m_width;
+            desc_nv12_cpu_copy.Height             = m_height;
+            desc_nv12_cpu_copy.MipLevels          = 1;
+            desc_nv12_cpu_copy.ArraySize          = 1;
+            desc_nv12_cpu_copy.Format             = DXGI_FORMAT_NV12;
+            desc_nv12_cpu_copy.SampleDesc.Count   = 1;
+            desc_nv12_cpu_copy.SampleDesc.Quality = 0;
+            desc_nv12_cpu_copy.BindFlags          = 0;
+            desc_nv12_cpu_copy.Usage              = D3D11_USAGE_STAGING;
+            desc_nv12_cpu_copy.CPUAccessFlags     = /*D3D11_CPU_ACCESS_WRITE | */D3D11_CPU_ACCESS_READ;
+            desc_nv12_cpu_copy.MiscFlags          = 0;
+
+            r = m_pD3D11Dev->CreateTexture2D(&desc_nv12_cpu_copy, 0, &m_pSurfaceNV12_cpu_copy);
+            if (FAILED(r))
+            {
+                throw std::runtime_error("Can't create DX NV12 texture");
+            }
         }
 
         // initialize OpenCL context of OpenCV lib from DirectX
@@ -133,31 +187,42 @@ public:
 
 
     // get media data on DX surface for further processing
-    int get_surface(ID3D11Texture2D** ppSurface)
+    int get_surface(ID3D11Texture2D** ppSurface, bool use_nv12)
     {
         HRESULT r;
 
         if (!m_cap.read(m_frame_bgr))
-            throw std::runtime_error("Can't get frame");
+            return -1;
 
-        cv::cvtColor(m_frame_bgr, m_frame_rgba, CV_RGB2BGRA);
+        if (use_nv12)
+        {
+            cv::cvtColor(m_frame_bgr, m_frame_i420, CV_BGR2YUV_I420);
 
+            convert_I420_to_NV12(m_frame_i420, m_frame_nv12, m_width, m_height);
+
+            m_pD3D11Ctx->UpdateSubresource(m_pSurfaceNV12, 0, 0, m_frame_nv12.data, (UINT)m_frame_nv12.step[0], (UINT)m_frame_nv12.total());
+        }
+        else
+        {
+            cv::cvtColor(m_frame_bgr, m_frame_rgba, CV_BGR2RGBA);
+
+            // process video frame on CPU
             UINT subResource = ::D3D11CalcSubresource(0, 0, 1);
 
             D3D11_MAPPED_SUBRESOURCE mappedTex;
-        r = m_pD3D11Ctx->Map(m_pSurface, subResource, D3D11_MAP_WRITE_DISCARD, 0, &mappedTex);
+            r = m_pD3D11Ctx->Map(m_pSurfaceRGBA, subResource, D3D11_MAP_WRITE_DISCARD, 0, &mappedTex);
             if (FAILED(r))
             {
                 throw std::runtime_error("surface mapping failed!");
             }
 
-        cv::Mat m(m_height, m_width, CV_8UC4, mappedTex.pData, (int)mappedTex.RowPitch);
+            cv::Mat m(m_height, m_width, CV_8UC4, mappedTex.pData, mappedTex.RowPitch);
-        // copy video frame data to surface
             m_frame_rgba.copyTo(m);
 
-        m_pD3D11Ctx->Unmap(m_pSurface, subResource);
+            m_pD3D11Ctx->Unmap(m_pSurfaceRGBA, subResource);
+        }
 
-        *ppSurface = m_pSurface;
+        *ppSurface = use_nv12 ? m_pSurfaceNV12 : m_pSurfaceRGBA;
 
         return 0;
     } // get_surface()
@@ -171,10 +236,13 @@ public:
             if (m_shutdown)
                 return 0;
 
+            // capture user input once
+            MODE mode = (m_mode == MODE_GPU_NV12 && !m_nv12_available) ? MODE_GPU_RGBA : m_mode;
+
             HRESULT r;
             ID3D11Texture2D* pSurface = 0;
 
-            r = get_surface(&pSurface);
+            r = get_surface(&pSurface, mode == MODE_GPU_NV12);
             if (FAILED(r))
             {
                 throw std::runtime_error("get_surface() failed!");
@@ -182,7 +250,7 @@ public:
 
             m_timer.start();
 
-            switch (m_mode)
+            switch (mode)
             {
             case MODE_CPU:
             {
@@ -219,7 +287,8 @@ public:
                 break;
             }
 
-                case MODE_GPU:
+            case MODE_GPU_RGBA:
+            case MODE_GPU_NV12:
             {
                 // process video frame on GPU
                 cv::UMat u;
@@ -232,7 +301,7 @@ public:
                     cv::blur(u, u, cv::Size(15, 15), cv::Point(-7, -7));
                 }
 
-                    cv::String strMode = cv::format("mode: %s", m_modeStr[MODE_GPU].c_str());
+                cv::String strMode = cv::format("mode: %s", m_modeStr[mode].c_str());
                 cv::String strProcessing = m_demo_processing ? "blur frame" : "copy frame";
                 cv::String strTime = cv::format("time: %4.1f msec", m_timer.time(Timer::UNITS::MSEC));
                 cv::String strDevName = cv::format("OpenCL device: %s", m_oclDevName.c_str());
@@ -244,6 +313,47 @@ public:
 
                 cv::directx::convertToD3D11Texture2D(u, pSurface);
 
+                if (mode == MODE_GPU_NV12)
+                {
+                    // just for rendering, we need to convert NV12 to RGBA.
+                    m_pD3D11Ctx->CopyResource(m_pSurfaceNV12_cpu_copy, m_pSurfaceNV12);
+
+                    // process video frame on CPU
+                    {
+                        UINT subResource = ::D3D11CalcSubresource(0, 0, 1);
+
+                        D3D11_MAPPED_SUBRESOURCE mappedTex;
+                        r = m_pD3D11Ctx->Map(m_pSurfaceNV12_cpu_copy, subResource, D3D11_MAP_READ, 0, &mappedTex);
+                        if (FAILED(r))
+                        {
+                            throw std::runtime_error("surface mapping failed!");
+                        }
+
+                        cv::Mat frame_nv12(m_height + (m_height / 2), m_width, CV_8UC1, mappedTex.pData, mappedTex.RowPitch);
+                        cv::cvtColor(frame_nv12, m_frame_rgba, CV_YUV2RGBA_NV12);
+
+                        m_pD3D11Ctx->Unmap(m_pSurfaceNV12_cpu_copy, subResource);
+                    }
+
+                    {
+                        UINT subResource = ::D3D11CalcSubresource(0, 0, 1);
+
+                        D3D11_MAPPED_SUBRESOURCE mappedTex;
+                        r = m_pD3D11Ctx->Map(m_pSurfaceRGBA, subResource, D3D11_MAP_WRITE_DISCARD, 0, &mappedTex);
+                        if (FAILED(r))
+                        {
+                            throw std::runtime_error("surface mapping failed!");
+                        }
+
+                        cv::Mat m(m_height, m_width, CV_8UC4, mappedTex.pData, mappedTex.RowPitch);
+                        m_frame_rgba.copyTo(m);
+
+                        m_pD3D11Ctx->Unmap(m_pSurfaceRGBA, subResource);
+                    }
+
+                    pSurface = m_pSurfaceRGBA;
+                }
+
                 break;
             }
 
@@ -267,12 +377,14 @@ public:
         catch (cv::Exception& e)
         {
             std::cerr << "Exception: " << e.what() << std::endl;
+            cleanup();
             return 10;
         }
 
         catch (const std::exception& e)
         {
             std::cerr << "Exception: " << e.what() << std::endl;
+            cleanup();
             return 11;
         }
 
@@ -282,7 +394,9 @@ public:
 
     int cleanup(void)
     {
-        SAFE_RELEASE(m_pSurface);
+        SAFE_RELEASE(m_pSurfaceRGBA);
+        SAFE_RELEASE(m_pSurfaceNV12);
+        SAFE_RELEASE(m_pSurfaceNV12_cpu_copy);
         SAFE_RELEASE(m_pBackBuffer);
         SAFE_RELEASE(m_pD3D11SwapChain);
         SAFE_RELEASE(m_pRenderTarget);
@@ -292,16 +406,74 @@ public:
         return 0;
     } // cleanup()
 
+protected:
+    void convert_I420_to_NV12(cv::Mat& i420, cv::Mat& nv12, int width, int height)
+    {
+        nv12.create(i420.rows, i420.cols, CV_8UC1);
+
+        unsigned char* pSrcY = i420.data;
+        unsigned char* pDstY = nv12.data;
+        size_t srcStep = i420.step[0];
+        size_t dstStep = nv12.step[0];
+
+        {
+            unsigned char* src;
+            unsigned char* dst;
+
+            // copy Y plane
+            for (int i = 0; i < height; i++)
+            {
+                src = pSrcY + i*srcStep;
+                dst = pDstY + i*dstStep;
+
+                for (int j = 0; j < width; j++)
+                {
+                    dst[j] = src[j];
+                }
+            }
+        }
+
+        {
+            // copy U/V planes to UV plane
+            unsigned char* pSrcU;
+            unsigned char* pSrcV;
+            unsigned char* pDstUV;
+
+            size_t uv_offset = height * dstStep;
+
+            for (int i = 0; i < height / 2; i++)
+            {
+                pSrcU = pSrcY + height*width + i*(width / 2);
+                pSrcV = pSrcY + height*width + (height / 2) * (width / 2) + i*(width / 2);
+
+                pDstUV = pDstY + uv_offset + i*dstStep;
+
+                for (int j = 0; j < width / 2; j++)
+                {
+                    pDstUV[j*2 + 0] = pSrcU[j];
+                    pDstUV[j*2 + 1] = pSrcV[j];
+                }
+            }
+        }
+
+        return;
+    }
+
 private:
     ID3D11Device*           m_pD3D11Dev;
     IDXGISwapChain*         m_pD3D11SwapChain;
     ID3D11DeviceContext*    m_pD3D11Ctx;
     ID3D11Texture2D*        m_pBackBuffer;
-    ID3D11Texture2D*        m_pSurface;
+    ID3D11Texture2D*        m_pSurfaceRGBA;
+    ID3D11Texture2D*        m_pSurfaceNV12;
+    ID3D11Texture2D*        m_pSurfaceNV12_cpu_copy;
     ID3D11RenderTargetView* m_pRenderTarget;
     cv::ocl::Context        m_oclCtx;
     cv::String              m_oclPlatformName;
     cv::String              m_oclDevName;
+    bool                    m_nv12_available;
+    cv::Mat                 m_frame_i420;
+    cv::Mat                 m_frame_nv12;
 };
 
 

samples/directx/d3d9_interop.cpp

@@ -108,7 +108,7 @@ public:
         if (!m_cap.read(m_frame_bgr))
             return -1;
 
-        cv::cvtColor(m_frame_bgr, m_frame_rgba, CV_RGB2RGBA);
+        cv::cvtColor(m_frame_bgr, m_frame_rgba, CV_BGR2BGRA);
 
         D3DLOCKED_RECT memDesc = { 0, NULL };
         RECT rc = { 0, 0, m_width, m_height };
@@ -143,6 +143,9 @@ public:
             if (m_shutdown)
                 return 0;
 
+            // capture user input once
+            MODE mode = (m_mode == MODE_GPU_NV12) ? MODE_GPU_RGBA : m_mode;
+
             HRESULT r;
             LPDIRECT3DSURFACE9 pSurface;
 
@@ -154,7 +157,7 @@ public:
 
             m_timer.start();
 
-            switch (m_mode)
+            switch (mode)
             {
                 case MODE_CPU:
                 {
@@ -185,7 +188,7 @@ public:
                     break;
                 }
 
-                case MODE_GPU:
+                case MODE_GPU_RGBA:
                 {
                     // process video frame on GPU
                     cv::UMat u;
@@ -207,7 +210,7 @@ public:
 
             m_timer.stop();
 
-            print_info(pSurface, m_mode, m_timer.time(Timer::UNITS::MSEC), m_oclDevName);
+            print_info(pSurface, mode, m_timer.time(Timer::UNITS::MSEC), m_oclDevName);
 
             // traditional DX render pipeline:
             //   BitBlt surface to backBuffer and flip backBuffer to frontBuffer

samples/directx/d3d9ex_interop.cpp

@@ -108,7 +108,7 @@ public:
         if (!m_cap.read(m_frame_bgr))
             return -1;
 
-        cv::cvtColor(m_frame_bgr, m_frame_rgba, CV_RGB2RGBA);
+        cv::cvtColor(m_frame_bgr, m_frame_rgba, CV_BGR2BGRA);
 
         D3DLOCKED_RECT memDesc = { 0, NULL };
         RECT rc = { 0, 0, m_width, m_height };
@@ -143,6 +143,9 @@ public:
             if (m_shutdown)
                 return 0;
 
+            // capture user input once
+            MODE mode = m_mode == MODE_GPU_NV12 ? MODE_GPU_RGBA : m_mode;
+
             HRESULT r;
             LPDIRECT3DSURFACE9 pSurface;
 
@@ -154,7 +157,7 @@ public:
 
             m_timer.start();
 
-            switch (m_mode)
+            switch (mode)
             {
                 case MODE_CPU:
                 {
@@ -185,7 +188,7 @@ public:
                     break;
                 }
 
-                case MODE_GPU:
+                case MODE_GPU_RGBA:
                 {
                     // process video frame on GPU
                     cv::UMat u;

samples/directx/d3dsample.hpp

@@ -67,7 +67,8 @@ public:
     enum MODE
     {
         MODE_CPU,
-        MODE_GPU
+        MODE_GPU_RGBA,
+        MODE_GPU_NV12
     };
 
     D3DSample(int width, int height, std::string& window_name, cv::VideoCapture& cap) :
@@ -76,7 +77,8 @@ public:
         m_shutdown          = false;
         m_mode              = MODE_CPU;
         m_modeStr[0]        = cv::String("Processing on CPU");
-        m_modeStr[1]        = cv::String("Processing on GPU");
+        m_modeStr[1]        = cv::String("Processing on GPU RGBA");
+        m_modeStr[2]        = cv::String("Processing on GPU NV12");
         m_demo_processing   = false;
         m_cap               = cap;
     }
@@ -104,7 +106,12 @@ protected:
             }
             if (wParam == '2')
             {
-                m_mode = MODE_GPU;
+                m_mode = MODE_GPU_RGBA;
+                return 0;
+            }
+            if (wParam == '3')
+            {
+                m_mode = MODE_GPU_NV12;
                 return 0;
             }
             else if (wParam == VK_SPACE)
@@ -136,7 +143,7 @@ protected:
     bool               m_shutdown;
     bool               m_demo_processing;
     MODE               m_mode;
-    cv::String         m_modeStr[2];
+    cv::String         m_modeStr[3];
     cv::VideoCapture   m_cap;
     cv::Mat            m_frame_bgr;
     cv::Mat            m_frame_rgba;
@@ -151,7 +158,8 @@ static void help()
         "Hot keys: \n"
         "  SPACE - turn processing on/off\n"
         "    1   - process DX surface through OpenCV on CPU\n"
-        "    2   - process DX surface through OpenCV on GPU (via OpenCL)\n"
+        "    2   - process DX RGBA surface through OpenCV on GPU (via OpenCL)\n"
+        "    3   - process DX NV12 surface through OpenCV on GPU (via OpenCL)\n"
         "  ESC   - exit\n\n");
 }