1 /*M/////////////////////////////////////////////////////////////////////////////////////// 2 // 3 // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. 4 // 5 // By downloading, copying, installing or using the software you agree to this license. 6 // If you do not agree to this license, do not download, install, 7 // copy or use the software. 8 // 9 // 10 // License Agreement 11 // For Open Source Computer Vision Library 12 // 13 // Copyright (C) 2000-2008, Intel Corporation, all rights reserved. 14 // Copyright (C) 2009, Willow Garage Inc., all rights reserved. 15 // Third party copyrights are property of their respective owners. 16 // 17 // Redistribution and use in source and binary forms, with or without modification, 18 // are permitted provided that the following conditions are met: 19 // 20 // * Redistribution's of source code must retain the above copyright notice, 21 // this list of conditions and the following disclaimer. 22 // 23 // * Redistribution's in binary form must reproduce the above copyright notice, 24 // this list of conditions and the following disclaimer in the documentation 25 // and/or other materials provided with the distribution. 26 // 27 // * The name of the copyright holders may not be used to endorse or promote products 28 // derived from this software without specific prior written permission. 29 // 30 // This software is provided by the copyright holders and contributors "as is" and 31 // any express or implied warranties, including, but not limited to, the implied 32 // warranties of merchantability and fitness for a particular purpose are disclaimed. 33 // In no event shall the Intel Corporation or contributors be liable for any direct, 34 // indirect, incidental, special, exemplary, or consequential damages 35 // (including, but not limited to, procurement of substitute goods or services; 36 // loss of use, data, or profits; or business interruption) however caused 37 // and on any theory of liability, whether in contract, strict liability, 38 // or tort (including negligence or otherwise) arising in any way out of 39 // the use of this software, even if advised of the possibility of such damage. 40 // 41 //M*/ 42 43 #include "opencv2/opencv_modules.hpp" 44 45 #ifndef HAVE_OPENCV_CUDEV 46 47 #error "opencv_cudev is required" 48 49 #else 50 51 #include "opencv2/cudev/ptr2d/glob.hpp" 52 53 using namespace cv::cudev; 54 55 void RGB_to_YV12(const GpuMat& src, GpuMat& dst); 56 57 namespace 58 { 59 __device__ __forceinline__ void rgb_to_y(const uchar b, const uchar g, const uchar r, uchar& y) 60 { 61 y = static_cast<uchar>(((int)(30 * r) + (int)(59 * g) + (int)(11 * b)) / 100); 62 } 63 64 __device__ __forceinline__ void rgb_to_yuv(const uchar b, const uchar g, const uchar r, uchar& y, uchar& u, uchar& v) 65 { 66 rgb_to_y(b, g, r, y); 67 u = static_cast<uchar>(((int)(-17 * r) - (int)(33 * g) + (int)(50 * b) + 12800) / 100); 68 v = static_cast<uchar>(((int)(50 * r) - (int)(42 * g) - (int)(8 * b) + 12800) / 100); 69 } 70 71 __global__ void Gray_to_YV12(const GlobPtrSz<uchar> src, GlobPtr<uchar> dst) 72 { 73 const int x = (blockIdx.x * blockDim.x + threadIdx.x) * 2; 74 const int y = (blockIdx.y * blockDim.y + threadIdx.y) * 2; 75 76 if (x + 1 >= src.cols || y + 1 >= src.rows) 77 return; 78 79 // get pointers to the data 80 const size_t planeSize = src.rows * dst.step; 81 GlobPtr<uchar> y_plane = globPtr(dst.data, dst.step); 82 GlobPtr<uchar> u_plane = globPtr(y_plane.data + planeSize, dst.step / 2); 83 GlobPtr<uchar> v_plane = globPtr(u_plane.data + (planeSize / 4), dst.step / 2); 84 85 uchar pix; 86 uchar y_val, u_val, v_val; 87 88 pix = src(y, x); 89 rgb_to_y(pix, pix, pix, y_val); 90 y_plane(y, x) = y_val; 91 92 pix = src(y, x + 1); 93 rgb_to_y(pix, pix, pix, y_val); 94 y_plane(y, x + 1) = y_val; 95 96 pix = src(y + 1, x); 97 rgb_to_y(pix, pix, pix, y_val); 98 y_plane(y + 1, x) = y_val; 99 100 pix = src(y + 1, x + 1); 101 rgb_to_yuv(pix, pix, pix, y_val, u_val, v_val); 102 y_plane(y + 1, x + 1) = y_val; 103 u_plane(y / 2, x / 2) = u_val; 104 v_plane(y / 2, x / 2) = v_val; 105 } 106 107 template <typename T> 108 __global__ void RGB_to_YV12(const GlobPtrSz<T> src, GlobPtr<uchar> dst) 109 { 110 const int x = (blockIdx.x * blockDim.x + threadIdx.x) * 2; 111 const int y = (blockIdx.y * blockDim.y + threadIdx.y) * 2; 112 113 if (x + 1 >= src.cols || y + 1 >= src.rows) 114 return; 115 116 // get pointers to the data 117 const size_t planeSize = src.rows * dst.step; 118 GlobPtr<uchar> y_plane = globPtr(dst.data, dst.step); 119 GlobPtr<uchar> u_plane = globPtr(y_plane.data + planeSize, dst.step / 2); 120 GlobPtr<uchar> v_plane = globPtr(u_plane.data + (planeSize / 4), dst.step / 2); 121 122 T pix; 123 uchar y_val, u_val, v_val; 124 125 pix = src(y, x); 126 rgb_to_y(pix.z, pix.y, pix.x, y_val); 127 y_plane(y, x) = y_val; 128 129 pix = src(y, x + 1); 130 rgb_to_y(pix.z, pix.y, pix.x, y_val); 131 y_plane(y, x + 1) = y_val; 132 133 pix = src(y + 1, x); 134 rgb_to_y(pix.z, pix.y, pix.x, y_val); 135 y_plane(y + 1, x) = y_val; 136 137 pix = src(y + 1, x + 1); 138 rgb_to_yuv(pix.z, pix.y, pix.x, y_val, u_val, v_val); 139 y_plane(y + 1, x + 1) = y_val; 140 u_plane(y / 2, x / 2) = u_val; 141 v_plane(y / 2, x / 2) = v_val; 142 } 143 } 144 145 void RGB_to_YV12(const GpuMat& src, GpuMat& dst) 146 { 147 const dim3 block(32, 8); 148 const dim3 grid(divUp(src.cols, block.x * 2), divUp(src.rows, block.y * 2)); 149 150 switch (src.channels()) 151 { 152 case 1: 153 Gray_to_YV12<<<grid, block>>>(globPtr<uchar>(src), globPtr<uchar>(dst)); 154 break; 155 case 3: 156 RGB_to_YV12<<<grid, block>>>(globPtr<uchar3>(src), globPtr<uchar>(dst)); 157 break; 158 case 4: 159 RGB_to_YV12<<<grid, block>>>(globPtr<uchar4>(src), globPtr<uchar>(dst)); 160 break; 161 } 162 163 CV_CUDEV_SAFE_CALL( cudaGetLastError() ); 164 CV_CUDEV_SAFE_CALL( cudaDeviceSynchronize() ); 165 } 166 167 #endif 168
