modules/gpu/src/cuda/rgb_to_yv12.cu

   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 bpied warranties, including, but not limited to, the bpied
  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 #if !defined CUDA_DISABLER
  44
  45 #include "opencv2/gpu/device/common.hpp"
  46 #include "opencv2/gpu/device/vec_traits.hpp"
  47
  48 namespace cv { namespace gpu { namespace device
  49 {
  50     namespace video_encoding
  51     {
  52         __device__ __forceinline__ void rgbtoy(const uchar b, const uchar g, const uchar r, uchar& y)
  53         {
  54             y = static_cast<uchar>(((int)(30 * r) + (int)(59 * g) + (int)(11 * b)) / 100);
  55         }
  56
  57         __device__ __forceinline__ void rgbtoyuv(const uchar b, const uchar g, const uchar r, uchar& y, uchar& u, uchar& v)
  58         {
  59             rgbtoy(b, g, r, y);
  60             u = static_cast<uchar>(((int)(-17 * r) - (int)(33 * g) + (int)(50 * b) + 12800) / 100);
  61             v = static_cast<uchar>(((int)(50 * r) - (int)(42 * g) - (int)(8 * b) + 12800) / 100);
  62         }
  63
  64         __global__ void Gray_to_YV12(const PtrStepSzb src, PtrStepb dst)
  65         {
  66             const int x = (blockIdx.x * blockDim.x + threadIdx.x) * 2;
  67             const int y = (blockIdx.y * blockDim.y + threadIdx.y) * 2;
  68
  69             if (x + 1 >= src.cols || y + 1 >= src.rows)
  70                 return;
  71
  72             // get pointers to the data
  73             const size_t planeSize = src.rows * dst.step;
  74             PtrStepb y_plane(dst.data, dst.step);
  75             PtrStepb u_plane(y_plane.data + planeSize, dst.step / 2);
  76             PtrStepb v_plane(u_plane.data + (planeSize / 4), dst.step / 2);
  77
  78             uchar pix;
  79             uchar y_val, u_val, v_val;
  80
  81             pix = src(y, x);
  82             rgbtoy(pix, pix, pix, y_val);
  83             y_plane(y, x) = y_val;
  84
  85             pix = src(y, x + 1);
  86             rgbtoy(pix, pix, pix, y_val);
  87             y_plane(y, x + 1) = y_val;
  88
  89             pix = src(y + 1, x);
  90             rgbtoy(pix, pix, pix, y_val);
  91             y_plane(y + 1, x) = y_val;
  92
  93             pix = src(y + 1, x + 1);
  94             rgbtoyuv(pix, pix, pix, y_val, u_val, v_val);
  95             y_plane(y + 1, x + 1) = y_val;
  96             u_plane(y / 2, x / 2) = u_val;
  97             v_plane(y / 2, x / 2) = v_val;
  98         }
  99
 100         template <typename T>
 101         __global__ void BGR_to_YV12(const PtrStepSz<T> src, PtrStepb dst)
 102         {
 103             const int x = (blockIdx.x * blockDim.x + threadIdx.x) * 2;
 104             const int y = (blockIdx.y * blockDim.y + threadIdx.y) * 2;
 105
 106             if (x + 1 >= src.cols || y + 1 >= src.rows)
 107                 return;
 108
 109             // get pointers to the data
 110             const size_t planeSize = src.rows * dst.step;
 111             PtrStepb y_plane(dst.data, dst.step);
 112             PtrStepb u_plane(y_plane.data + planeSize, dst.step / 2);
 113             PtrStepb v_plane(u_plane.data + (planeSize / 4), dst.step / 2);
 114
 115             T pix;
 116             uchar y_val, u_val, v_val;
 117
 118             pix = src(y, x);
 119             rgbtoy(pix.z, pix.y, pix.x, y_val);
 120             y_plane(y, x) = y_val;
 121
 122             pix = src(y, x + 1);
 123             rgbtoy(pix.z, pix.y, pix.x, y_val);
 124             y_plane(y, x + 1) = y_val;
 125
 126             pix = src(y + 1, x);
 127             rgbtoy(pix.z, pix.y, pix.x, y_val);
 128             y_plane(y + 1, x) = y_val;
 129
 130             pix = src(y + 1, x + 1);
 131             rgbtoyuv(pix.z, pix.y, pix.x, y_val, u_val, v_val);
 132             y_plane(y + 1, x + 1) = y_val;
 133             u_plane(y / 2, x / 2) = u_val;
 134             v_plane(y / 2, x / 2) = v_val;
 135         }
 136
 137         void Gray_to_YV12_caller(const PtrStepSzb src, PtrStepb dst)
 138         {
 139             dim3 block(32, 8);
 140             dim3 grid(divUp(src.cols, block.x * 2), divUp(src.rows, block.y * 2));
 141
 142             Gray_to_YV12<<<grid, block>>>(src, dst);
 143             cudaSafeCall( cudaGetLastError() );
 144
 145             cudaSafeCall( cudaDeviceSynchronize() );
 146         }
 147         template <int cn>
 148         void BGR_to_YV12_caller(const PtrStepSzb src, PtrStepb dst)
 149         {
 150             typedef typename TypeVec<uchar, cn>::vec_type src_t;
 151
 152             dim3 block(32, 8);
 153             dim3 grid(divUp(src.cols, block.x * 2), divUp(src.rows, block.y * 2));
 154
 155             BGR_to_YV12<<<grid, block>>>(static_cast< PtrStepSz<src_t> >(src), dst);
 156             cudaSafeCall( cudaGetLastError() );
 157
 158             cudaSafeCall( cudaDeviceSynchronize() );
 159         }
 160
 161         void YV12_gpu(const PtrStepSzb src, int cn, PtrStepSzb dst)
 162         {
 163             typedef void (*func_t)(const PtrStepSzb src, PtrStepb dst);
 164
 165             static const func_t funcs[] =
 166             {
 167                 0, Gray_to_YV12_caller, 0, BGR_to_YV12_caller<3>, BGR_to_YV12_caller<4>
 168             };
 169
 170             funcs[cn](src, dst);
 171         }
 172     }
 173 }}}
 174
 175 #endif /* CUDA_DISABLER */