Home | History | Annotate | Download | only in src
      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 "precomp.hpp"
     44 
     45 using namespace cv;
     46 using namespace cv::cuda;
     47 
     48 #if !defined (HAVE_CUDA) || defined (CUDA_DISABLER)
     49 
     50 void cv::cuda::cvtColor(InputArray, OutputArray, int, int, Stream&) { throw_no_cuda(); }
     51 
     52 void cv::cuda::demosaicing(InputArray, OutputArray, int, int, Stream&) { throw_no_cuda(); }
     53 
     54 void cv::cuda::swapChannels(InputOutputArray, const int[], Stream&) { throw_no_cuda(); }
     55 
     56 void cv::cuda::gammaCorrection(InputArray, OutputArray, bool, Stream&) { throw_no_cuda(); }
     57 
     58 void cv::cuda::alphaComp(InputArray, InputArray, OutputArray, int, Stream&) { throw_no_cuda(); }
     59 
     60 
     61 #else /* !defined (HAVE_CUDA) */
     62 
     63 #include "cvt_color_internal.h"
     64 
     65 namespace cv { namespace cuda {
     66     namespace device
     67     {
     68         template <int cn>
     69         void Bayer2BGR_8u_gpu(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
     70         template <int cn>
     71         void Bayer2BGR_16u_gpu(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
     72 
     73         template <int cn>
     74         void MHCdemosaic(PtrStepSzb src, int2 sourceOffset, PtrStepSzb dst, int2 firstRed, cudaStream_t stream);
     75     }
     76 }}
     77 
     78 using namespace ::cv::cuda::device;
     79 
     80 namespace
     81 {
     82     typedef void (*gpu_func_t)(const GpuMat& _src, GpuMat& _dst, Stream& stream);
     83 
     84     void BGR_to_RGB(InputArray _src, OutputArray _dst, int, Stream& stream)
     85     {
     86         using namespace cv::cuda::device;
     87         static const gpu_func_t funcs[] = {BGR_to_RGB_8u, 0, BGR_to_RGB_16u, 0, 0, BGR_to_RGB_32f};
     88 
     89         GpuMat src = _src.getGpuMat();
     90 
     91         CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F );
     92         CV_Assert( src.channels() == 3 );
     93 
     94         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), 3));
     95         GpuMat dst = _dst.getGpuMat();
     96 
     97         funcs[src.depth()](src, dst, stream);
     98     }
     99 
    100     void BGR_to_BGRA(InputArray _src, OutputArray _dst, int, Stream& stream)
    101     {
    102         using namespace cv::cuda::device;
    103         static const gpu_func_t funcs[] = {BGR_to_BGRA_8u, 0, BGR_to_BGRA_16u, 0, 0, BGR_to_BGRA_32f};
    104 
    105         GpuMat src = _src.getGpuMat();
    106 
    107         CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F );
    108         CV_Assert( src.channels() == 3 );
    109 
    110         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), 4));
    111         GpuMat dst = _dst.getGpuMat();
    112 
    113         funcs[src.depth()](src, dst, stream);
    114     }
    115 
    116     void BGR_to_RGBA(InputArray _src, OutputArray _dst, int, Stream& stream)
    117     {
    118         using namespace cv::cuda::device;
    119         static const gpu_func_t funcs[] = {BGR_to_RGBA_8u, 0, BGR_to_RGBA_16u, 0, 0, BGR_to_RGBA_32f};
    120 
    121         GpuMat src = _src.getGpuMat();
    122 
    123         CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F );
    124         CV_Assert( src.channels() == 3 );
    125 
    126         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), 4));
    127         GpuMat dst = _dst.getGpuMat();
    128 
    129         funcs[src.depth()](src, dst, stream);
    130     }
    131 
    132     void BGRA_to_BGR(InputArray _src, OutputArray _dst, int, Stream& stream)
    133     {
    134         using namespace cv::cuda::device;
    135         static const gpu_func_t funcs[] = {BGRA_to_BGR_8u, 0, BGRA_to_BGR_16u, 0, 0, BGRA_to_BGR_32f};
    136 
    137         GpuMat src = _src.getGpuMat();
    138 
    139         CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F );
    140         CV_Assert( src.channels() == 4 );
    141 
    142         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), 3));
    143         GpuMat dst = _dst.getGpuMat();
    144 
    145         funcs[src.depth()](src, dst, stream);
    146     }
    147 
    148     void BGRA_to_RGB(InputArray _src, OutputArray _dst, int, Stream& stream)
    149     {
    150         using namespace cv::cuda::device;
    151         static const gpu_func_t funcs[] = {BGRA_to_RGB_8u, 0, BGRA_to_RGB_16u, 0, 0, BGRA_to_RGB_32f};
    152 
    153         GpuMat src = _src.getGpuMat();
    154 
    155         CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F );
    156         CV_Assert( src.channels() == 4 );
    157 
    158         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), 3));
    159         GpuMat dst = _dst.getGpuMat();
    160 
    161         funcs[src.depth()](src, dst, stream);
    162     }
    163 
    164     void BGRA_to_RGBA(InputArray _src, OutputArray _dst, int, Stream& stream)
    165     {
    166         using namespace cv::cuda::device;
    167         static const gpu_func_t funcs[] = {BGRA_to_RGBA_8u, 0, BGRA_to_RGBA_16u, 0, 0, BGRA_to_RGBA_32f};
    168 
    169         GpuMat src = _src.getGpuMat();
    170 
    171         CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F );
    172         CV_Assert( src.channels() == 4 );
    173 
    174         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), 4));
    175         GpuMat dst = _dst.getGpuMat();
    176 
    177         funcs[src.depth()](src, dst, stream);
    178     }
    179 
    180     void BGR_to_BGR555(InputArray _src, OutputArray _dst, int, Stream& stream)
    181     {
    182         GpuMat src = _src.getGpuMat();
    183 
    184         CV_Assert( src.depth() == CV_8U );
    185         CV_Assert( src.channels() == 3 );
    186 
    187         _dst.create(src.size(), CV_8UC2);
    188         GpuMat dst = _dst.getGpuMat();
    189 
    190         cv::cuda::device::BGR_to_BGR555(src, dst, stream);
    191     }
    192 
    193     void BGR_to_BGR565(InputArray _src, OutputArray _dst, int, Stream& stream)
    194     {
    195         GpuMat src = _src.getGpuMat();
    196 
    197         CV_Assert( src.depth() == CV_8U );
    198         CV_Assert( src.channels() == 3 );
    199 
    200         _dst.create(src.size(), CV_8UC2);
    201         GpuMat dst = _dst.getGpuMat();
    202 
    203         cv::cuda::device::BGR_to_BGR565(src, dst, stream);
    204     }
    205 
    206     void RGB_to_BGR555(InputArray _src, OutputArray _dst, int, Stream& stream)
    207     {
    208         GpuMat src = _src.getGpuMat();
    209 
    210         CV_Assert( src.depth() == CV_8U );
    211         CV_Assert( src.channels() == 3 );
    212 
    213         _dst.create(src.size(), CV_8UC2);
    214         GpuMat dst = _dst.getGpuMat();
    215 
    216         cv::cuda::device::RGB_to_BGR555(src, dst, stream);
    217     }
    218 
    219     void RGB_to_BGR565(InputArray _src, OutputArray _dst, int, Stream& stream)
    220     {
    221         GpuMat src = _src.getGpuMat();
    222 
    223         CV_Assert( src.depth() == CV_8U );
    224         CV_Assert( src.channels() == 3 );
    225 
    226         _dst.create(src.size(), CV_8UC2);
    227         GpuMat dst = _dst.getGpuMat();
    228 
    229         cv::cuda::device::RGB_to_BGR565(src, dst, stream);
    230     }
    231 
    232     void BGRA_to_BGR555(InputArray _src, OutputArray _dst, int, Stream& stream)
    233     {
    234         GpuMat src = _src.getGpuMat();
    235 
    236         CV_Assert( src.depth() == CV_8U );
    237         CV_Assert( src.channels() == 4 );
    238 
    239         _dst.create(src.size(), CV_8UC2);
    240         GpuMat dst = _dst.getGpuMat();
    241 
    242         cv::cuda::device::BGRA_to_BGR555(src, dst, stream);
    243     }
    244 
    245     void BGRA_to_BGR565(InputArray _src, OutputArray _dst, int, Stream& stream)
    246     {
    247         GpuMat src = _src.getGpuMat();
    248 
    249         CV_Assert( src.depth() == CV_8U );
    250         CV_Assert( src.channels() == 4 );
    251 
    252         _dst.create(src.size(), CV_8UC2);
    253         GpuMat dst = _dst.getGpuMat();
    254 
    255         cv::cuda::device::BGRA_to_BGR565(src, dst, stream);
    256     }
    257 
    258     void RGBA_to_BGR555(InputArray _src, OutputArray _dst, int, Stream& stream)
    259     {
    260         GpuMat src = _src.getGpuMat();
    261 
    262         CV_Assert( src.depth() == CV_8U );
    263         CV_Assert( src.channels() == 4 );
    264 
    265         _dst.create(src.size(), CV_8UC2);
    266         GpuMat dst = _dst.getGpuMat();
    267 
    268         cv::cuda::device::RGBA_to_BGR555(src, dst, stream);
    269     }
    270 
    271     void RGBA_to_BGR565(InputArray _src, OutputArray _dst, int, Stream& stream)
    272     {
    273         GpuMat src = _src.getGpuMat();
    274 
    275         CV_Assert( src.depth() == CV_8U );
    276         CV_Assert( src.channels() == 4 );
    277 
    278         _dst.create(src.size(), CV_8UC2);
    279         GpuMat dst = _dst.getGpuMat();
    280 
    281         cv::cuda::device::RGBA_to_BGR565(src, dst, stream);
    282     }
    283 
    284     void BGR555_to_RGB(InputArray _src, OutputArray _dst, int, Stream& stream)
    285     {
    286         GpuMat src = _src.getGpuMat();
    287 
    288         CV_Assert( src.depth() == CV_8U );
    289         CV_Assert( src.channels() == 2 );
    290 
    291         _dst.create(src.size(), CV_8UC3);
    292         GpuMat dst = _dst.getGpuMat();
    293 
    294         cv::cuda::device::BGR555_to_RGB(src, dst, stream);
    295     }
    296 
    297     void BGR565_to_RGB(InputArray _src, OutputArray _dst, int, Stream& stream)
    298     {
    299         GpuMat src = _src.getGpuMat();
    300 
    301         CV_Assert( src.depth() == CV_8U );
    302         CV_Assert( src.channels() == 2 );
    303 
    304         _dst.create(src.size(), CV_8UC3);
    305         GpuMat dst = _dst.getGpuMat();
    306 
    307         cv::cuda::device::BGR565_to_RGB(src, dst, stream);
    308     }
    309 
    310     void BGR555_to_BGR(InputArray _src, OutputArray _dst, int, Stream& stream)
    311     {
    312         GpuMat src = _src.getGpuMat();
    313 
    314         CV_Assert( src.depth() == CV_8U );
    315         CV_Assert( src.channels() == 2 );
    316 
    317         _dst.create(src.size(), CV_8UC3);
    318         GpuMat dst = _dst.getGpuMat();
    319 
    320         cv::cuda::device::BGR555_to_BGR(src, dst, stream);
    321     }
    322 
    323     void BGR565_to_BGR(InputArray _src, OutputArray _dst, int, Stream& stream)
    324     {
    325         GpuMat src = _src.getGpuMat();
    326 
    327         CV_Assert( src.depth() == CV_8U );
    328         CV_Assert( src.channels() == 2 );
    329 
    330         _dst.create(src.size(), CV_8UC3);
    331         GpuMat dst = _dst.getGpuMat();
    332 
    333         cv::cuda::device::BGR565_to_BGR(src, dst, stream);
    334     }
    335 
    336     void BGR555_to_RGBA(InputArray _src, OutputArray _dst, int, Stream& stream)
    337     {
    338         GpuMat src = _src.getGpuMat();
    339 
    340         CV_Assert( src.depth() == CV_8U );
    341         CV_Assert( src.channels() == 2 );
    342 
    343         _dst.create(src.size(), CV_8UC4);
    344         GpuMat dst = _dst.getGpuMat();
    345 
    346         cv::cuda::device::BGR555_to_RGBA(src, dst, stream);
    347     }
    348 
    349     void BGR565_to_RGBA(InputArray _src, OutputArray _dst, int, Stream& stream)
    350     {
    351         GpuMat src = _src.getGpuMat();
    352 
    353         CV_Assert( src.depth() == CV_8U );
    354         CV_Assert( src.channels() == 2 );
    355 
    356         _dst.create(src.size(), CV_8UC4);
    357         GpuMat dst = _dst.getGpuMat();
    358 
    359         cv::cuda::device::BGR565_to_RGBA(src, dst, stream);
    360     }
    361 
    362     void BGR555_to_BGRA(InputArray _src, OutputArray _dst, int, Stream& stream)
    363     {
    364         GpuMat src = _src.getGpuMat();
    365 
    366         CV_Assert( src.depth() == CV_8U );
    367         CV_Assert( src.channels() == 2 );
    368 
    369         _dst.create(src.size(), CV_8UC4);
    370         GpuMat dst = _dst.getGpuMat();
    371 
    372         cv::cuda::device::BGR555_to_BGRA(src, dst, stream);
    373     }
    374 
    375     void BGR565_to_BGRA(InputArray _src, OutputArray _dst, int, Stream& stream)
    376     {
    377         GpuMat src = _src.getGpuMat();
    378 
    379         CV_Assert( src.depth() == CV_8U );
    380         CV_Assert( src.channels() == 2 );
    381 
    382         _dst.create(src.size(), CV_8UC4);
    383         GpuMat dst = _dst.getGpuMat();
    384 
    385         cv::cuda::device::BGR565_to_BGRA(src, dst, stream);
    386     }
    387 
    388     void GRAY_to_BGR(InputArray _src, OutputArray _dst, int, Stream& stream)
    389     {
    390         using namespace cv::cuda::device;
    391         static const gpu_func_t funcs[] = {GRAY_to_BGR_8u, 0, GRAY_to_BGR_16u, 0, 0, GRAY_to_BGR_32f};
    392 
    393         GpuMat src = _src.getGpuMat();
    394 
    395         CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F );
    396         CV_Assert( src.channels() == 1 );
    397 
    398         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), 3));
    399         GpuMat dst = _dst.getGpuMat();
    400 
    401         funcs[src.depth()](src, dst, stream);
    402     }
    403 
    404     void GRAY_to_BGRA(InputArray _src, OutputArray _dst, int, Stream& stream)
    405     {
    406         using namespace cv::cuda::device;
    407         static const gpu_func_t funcs[] = {GRAY_to_BGRA_8u, 0, GRAY_to_BGRA_16u, 0, 0, GRAY_to_BGRA_32f};
    408 
    409         GpuMat src = _src.getGpuMat();
    410 
    411         CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F );
    412         CV_Assert( src.channels() == 1 );
    413 
    414         _dst.create(src.size(), CV_MAKETYPE(src.depth(), 4));
    415         GpuMat dst = _dst.getGpuMat();
    416 
    417         funcs[src.depth()](src, dst, stream);
    418     }
    419 
    420     void GRAY_to_BGR555(InputArray _src, OutputArray _dst, int, Stream& stream)
    421     {
    422         GpuMat src = _src.getGpuMat();
    423 
    424         CV_Assert( src.depth() == CV_8U );
    425         CV_Assert( src.channels() == 1 );
    426 
    427         _dst.create(src.size(), CV_8UC2);
    428         GpuMat dst = _dst.getGpuMat();
    429 
    430         cv::cuda::device::GRAY_to_BGR555(src, dst, stream);
    431     }
    432 
    433     void GRAY_to_BGR565(InputArray _src, OutputArray _dst, int, Stream& stream)
    434     {
    435         GpuMat src = _src.getGpuMat();
    436 
    437         CV_Assert( src.depth() == CV_8U );
    438         CV_Assert( src.channels() == 1 );
    439 
    440         _dst.create(src.size(), CV_8UC2);
    441         GpuMat dst = _dst.getGpuMat();
    442 
    443         cv::cuda::device::GRAY_to_BGR565(src, dst, stream);
    444     }
    445 
    446     void BGR555_to_GRAY(InputArray _src, OutputArray _dst, int, Stream& stream)
    447     {
    448         GpuMat src = _src.getGpuMat();
    449 
    450         CV_Assert( src.depth() == CV_8U );
    451         CV_Assert( src.channels() == 2 );
    452 
    453         _dst.create(src.size(), CV_8UC1);
    454         GpuMat dst = _dst.getGpuMat();
    455 
    456         cv::cuda::device::BGR555_to_GRAY(src, dst, stream);
    457     }
    458 
    459     void BGR565_to_GRAY(InputArray _src, OutputArray _dst, int, Stream& stream)
    460     {
    461         GpuMat src = _src.getGpuMat();
    462 
    463         CV_Assert( src.depth() == CV_8U );
    464         CV_Assert( src.channels() == 2 );
    465 
    466         _dst.create(src.size(), CV_8UC1);
    467         GpuMat dst = _dst.getGpuMat();
    468 
    469         cv::cuda::device::BGR565_to_GRAY(src, dst, stream);
    470     }
    471 
    472     void RGB_to_GRAY(InputArray _src, OutputArray _dst, int, Stream& stream)
    473     {
    474         using namespace cv::cuda::device;
    475         static const gpu_func_t funcs[] = {RGB_to_GRAY_8u, 0, RGB_to_GRAY_16u, 0, 0, RGB_to_GRAY_32f};
    476 
    477         GpuMat src = _src.getGpuMat();
    478 
    479         CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F );
    480         CV_Assert( src.channels() == 3 );
    481 
    482         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), 1));
    483         GpuMat dst = _dst.getGpuMat();
    484 
    485         funcs[src.depth()](src, dst, stream);
    486     }
    487 
    488     void BGR_to_GRAY(InputArray _src, OutputArray _dst, int, Stream& stream)
    489     {
    490         using namespace cv::cuda::device;
    491         static const gpu_func_t funcs[] = {BGR_to_GRAY_8u, 0, BGR_to_GRAY_16u, 0, 0, BGR_to_GRAY_32f};
    492 
    493         GpuMat src = _src.getGpuMat();
    494 
    495         CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F );
    496         CV_Assert( src.channels() == 3 );
    497 
    498         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), 1));
    499         GpuMat dst = _dst.getGpuMat();
    500 
    501         funcs[src.depth()](src, dst, stream);
    502     }
    503 
    504     void RGBA_to_GRAY(InputArray _src, OutputArray _dst, int, Stream& stream)
    505     {
    506         using namespace cv::cuda::device;
    507         static const gpu_func_t funcs[] = {RGBA_to_GRAY_8u, 0, RGBA_to_GRAY_16u, 0, 0, RGBA_to_GRAY_32f};
    508 
    509         GpuMat src = _src.getGpuMat();
    510 
    511         CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F );
    512         CV_Assert( src.channels() == 4 );
    513 
    514         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), 1));
    515         GpuMat dst = _dst.getGpuMat();
    516 
    517         funcs[src.depth()](src, dst, stream);
    518     }
    519 
    520     void BGRA_to_GRAY(InputArray _src, OutputArray _dst, int, Stream& stream)
    521     {
    522         using namespace cv::cuda::device;
    523         static const gpu_func_t funcs[] = {BGRA_to_GRAY_8u, 0, BGRA_to_GRAY_16u, 0, 0, BGRA_to_GRAY_32f};
    524 
    525         GpuMat src = _src.getGpuMat();
    526 
    527         CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F );
    528         CV_Assert( src.channels() == 4 );
    529 
    530         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), 1));
    531         GpuMat dst = _dst.getGpuMat();
    532 
    533         funcs[src.depth()](src, dst, stream);
    534     }
    535 
    536     void RGB_to_YUV(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
    537     {
    538         using namespace cv::cuda::device;
    539         static const gpu_func_t funcs[2][2][6] =
    540         {
    541             {
    542                 {RGB_to_YUV_8u, 0, RGB_to_YUV_16u, 0, 0, RGB_to_YUV_32f},
    543                 {RGBA_to_YUV_8u, 0, RGBA_to_YUV_16u, 0, 0, RGBA_to_YUV_32f}
    544             },
    545             {
    546                 {RGB_to_YUV4_8u, 0, RGB_to_YUV4_16u, 0, 0, RGB_to_YUV4_32f},
    547                 {RGBA_to_YUV4_8u, 0, RGBA_to_YUV4_16u, 0, 0, RGBA_to_YUV4_32f}
    548             }
    549         };
    550 
    551         if (dcn <= 0) dcn = 3;
    552 
    553         GpuMat src = _src.getGpuMat();
    554 
    555         CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F );
    556         CV_Assert( src.channels() == 3 || src.channels() == 4 );
    557         CV_Assert( dcn == 3 || dcn == 4 );
    558 
    559         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
    560         GpuMat dst = _dst.getGpuMat();
    561 
    562         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
    563     }
    564 
    565     void BGR_to_YUV(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
    566     {
    567         using namespace cv::cuda::device;
    568         static const gpu_func_t funcs[2][2][6] =
    569         {
    570             {
    571                 {BGR_to_YUV_8u, 0, BGR_to_YUV_16u, 0, 0, BGR_to_YUV_32f},
    572                 {BGRA_to_YUV_8u, 0, BGRA_to_YUV_16u, 0, 0, BGRA_to_YUV_32f}
    573             },
    574             {
    575                 {BGR_to_YUV4_8u, 0, BGR_to_YUV4_16u, 0, 0, BGR_to_YUV4_32f},
    576                 {BGRA_to_YUV4_8u, 0, BGRA_to_YUV4_16u, 0, 0, BGRA_to_YUV4_32f}
    577             }
    578         };
    579 
    580         if (dcn <= 0) dcn = 3;
    581 
    582         GpuMat src = _src.getGpuMat();
    583 
    584         CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F );
    585         CV_Assert( src.channels() == 3 || src.channels() == 4 );
    586         CV_Assert( dcn == 3 || dcn == 4 );
    587 
    588         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
    589         GpuMat dst = _dst.getGpuMat();
    590 
    591         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
    592     }
    593 
    594     void YUV_to_RGB(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
    595     {
    596         using namespace cv::cuda::device;
    597         static const gpu_func_t funcs[2][2][6] =
    598         {
    599             {
    600                 {YUV_to_RGB_8u, 0, YUV_to_RGB_16u, 0, 0, YUV_to_RGB_32f},
    601                 {YUV4_to_RGB_8u, 0, YUV4_to_RGB_16u, 0, 0, YUV4_to_RGB_32f}
    602             },
    603             {
    604                 {YUV_to_RGBA_8u, 0, YUV_to_RGBA_16u, 0, 0, YUV_to_RGBA_32f},
    605                 {YUV4_to_RGBA_8u, 0, YUV4_to_RGBA_16u, 0, 0, YUV4_to_RGBA_32f}
    606             }
    607         };
    608 
    609         if (dcn <= 0) dcn = 3;
    610 
    611         GpuMat src = _src.getGpuMat();
    612 
    613         CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F );
    614         CV_Assert( src.channels() == 3 || src.channels() == 4 );
    615         CV_Assert( dcn == 3 || dcn == 4 );
    616 
    617         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
    618         GpuMat dst = _dst.getGpuMat();
    619 
    620         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
    621     }
    622 
    623     void YUV_to_BGR(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
    624     {
    625         using namespace cv::cuda::device;
    626         static const gpu_func_t funcs[2][2][6] =
    627         {
    628             {
    629                 {YUV_to_BGR_8u, 0, YUV_to_BGR_16u, 0, 0, YUV_to_BGR_32f},
    630                 {YUV4_to_BGR_8u, 0, YUV4_to_BGR_16u, 0, 0, YUV4_to_BGR_32f}
    631             },
    632             {
    633                 {YUV_to_BGRA_8u, 0, YUV_to_BGRA_16u, 0, 0, YUV_to_BGRA_32f},
    634                 {YUV4_to_BGRA_8u, 0, YUV4_to_BGRA_16u, 0, 0, YUV4_to_BGRA_32f}
    635             }
    636         };
    637 
    638         if (dcn <= 0) dcn = 3;
    639 
    640         GpuMat src = _src.getGpuMat();
    641 
    642         CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F );
    643         CV_Assert( src.channels() == 3 || src.channels() == 4 );
    644         CV_Assert( dcn == 3 || dcn == 4 );
    645 
    646         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
    647         GpuMat dst = _dst.getGpuMat();
    648 
    649         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
    650     }
    651 
    652     void RGB_to_YCrCb(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
    653     {
    654         using namespace cv::cuda::device;
    655         static const gpu_func_t funcs[2][2][6] =
    656         {
    657             {
    658                 {RGB_to_YCrCb_8u, 0, RGB_to_YCrCb_16u, 0, 0, RGB_to_YCrCb_32f},
    659                 {RGBA_to_YCrCb_8u, 0, RGBA_to_YCrCb_16u, 0, 0, RGBA_to_YCrCb_32f}
    660             },
    661             {
    662                 {RGB_to_YCrCb4_8u, 0, RGB_to_YCrCb4_16u, 0, 0, RGB_to_YCrCb4_32f},
    663                 {RGBA_to_YCrCb4_8u, 0, RGBA_to_YCrCb4_16u, 0, 0, RGBA_to_YCrCb4_32f}
    664             }
    665         };
    666 
    667         if (dcn <= 0) dcn = 3;
    668 
    669         GpuMat src = _src.getGpuMat();
    670 
    671         CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F );
    672         CV_Assert( src.channels() == 3 || src.channels() == 4 );
    673         CV_Assert( dcn == 3 || dcn == 4 );
    674 
    675         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
    676         GpuMat dst = _dst.getGpuMat();
    677 
    678         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
    679     }
    680 
    681     void BGR_to_YCrCb(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
    682     {
    683         using namespace cv::cuda::device;
    684         static const gpu_func_t funcs[2][2][6] =
    685         {
    686             {
    687                 {BGR_to_YCrCb_8u, 0, BGR_to_YCrCb_16u, 0, 0, BGR_to_YCrCb_32f},
    688                 {BGRA_to_YCrCb_8u, 0, BGRA_to_YCrCb_16u, 0, 0, BGRA_to_YCrCb_32f}
    689             },
    690             {
    691                 {BGR_to_YCrCb4_8u, 0, BGR_to_YCrCb4_16u, 0, 0, BGR_to_YCrCb4_32f},
    692                 {BGRA_to_YCrCb4_8u, 0, BGRA_to_YCrCb4_16u, 0, 0, BGRA_to_YCrCb4_32f}
    693             }
    694         };
    695 
    696         if (dcn <= 0) dcn = 3;
    697 
    698         GpuMat src = _src.getGpuMat();
    699 
    700         CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F );
    701         CV_Assert( src.channels() == 3 || src.channels() == 4 );
    702         CV_Assert( dcn == 3 || dcn == 4 );
    703 
    704         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
    705         GpuMat dst = _dst.getGpuMat();
    706 
    707         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
    708     }
    709 
    710     void YCrCb_to_RGB(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
    711     {
    712         using namespace cv::cuda::device;
    713         static const gpu_func_t funcs[2][2][6] =
    714         {
    715             {
    716                 {YCrCb_to_RGB_8u, 0, YCrCb_to_RGB_16u, 0, 0, YCrCb_to_RGB_32f},
    717                 {YCrCb4_to_RGB_8u, 0, YCrCb4_to_RGB_16u, 0, 0, YCrCb4_to_RGB_32f}
    718             },
    719             {
    720                 {YCrCb_to_RGBA_8u, 0, YCrCb_to_RGBA_16u, 0, 0, YCrCb_to_RGBA_32f},
    721                 {YCrCb4_to_RGBA_8u, 0, YCrCb4_to_RGBA_16u, 0, 0, YCrCb4_to_RGBA_32f}
    722             }
    723         };
    724 
    725         if (dcn <= 0) dcn = 3;
    726 
    727         GpuMat src = _src.getGpuMat();
    728 
    729         CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F );
    730         CV_Assert( src.channels() == 3 || src.channels() == 4 );
    731         CV_Assert( dcn == 3 || dcn == 4 );
    732 
    733         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
    734         GpuMat dst = _dst.getGpuMat();
    735 
    736         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
    737     }
    738 
    739     void YCrCb_to_BGR(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
    740     {
    741         using namespace cv::cuda::device;
    742         static const gpu_func_t funcs[2][2][6] =
    743         {
    744             {
    745                 {YCrCb_to_BGR_8u, 0, YCrCb_to_BGR_16u, 0, 0, YCrCb_to_BGR_32f},
    746                 {YCrCb4_to_BGR_8u, 0, YCrCb4_to_BGR_16u, 0, 0, YCrCb4_to_BGR_32f}
    747             },
    748             {
    749                 {YCrCb_to_BGRA_8u, 0, YCrCb_to_BGRA_16u, 0, 0, YCrCb_to_BGRA_32f},
    750                 {YCrCb4_to_BGRA_8u, 0, YCrCb4_to_BGRA_16u, 0, 0, YCrCb4_to_BGRA_32f}
    751             }
    752         };
    753 
    754         if (dcn <= 0) dcn = 3;
    755 
    756         GpuMat src = _src.getGpuMat();
    757 
    758         CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F );
    759         CV_Assert( src.channels() == 3 || src.channels() == 4 );
    760         CV_Assert( dcn == 3 || dcn == 4 );
    761 
    762         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
    763         GpuMat dst = _dst.getGpuMat();
    764 
    765         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
    766     }
    767 
    768     void RGB_to_XYZ(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
    769     {
    770         using namespace cv::cuda::device;
    771         static const gpu_func_t funcs[2][2][6] =
    772         {
    773             {
    774                 {RGB_to_XYZ_8u, 0, RGB_to_XYZ_16u, 0, 0, RGB_to_XYZ_32f},
    775                 {RGBA_to_XYZ_8u, 0, RGBA_to_XYZ_16u, 0, 0, RGBA_to_XYZ_32f}
    776             },
    777             {
    778                 {RGB_to_XYZ4_8u, 0, RGB_to_XYZ4_16u, 0, 0, RGB_to_XYZ4_32f},
    779                 {RGBA_to_XYZ4_8u, 0, RGBA_to_XYZ4_16u, 0, 0, RGBA_to_XYZ4_32f}
    780             }
    781         };
    782 
    783         if (dcn <= 0) dcn = 3;
    784 
    785         GpuMat src = _src.getGpuMat();
    786 
    787         CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F );
    788         CV_Assert( src.channels() == 3 || src.channels() == 4 );
    789         CV_Assert( dcn == 3 || dcn == 4 );
    790 
    791         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
    792         GpuMat dst = _dst.getGpuMat();
    793 
    794         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
    795     }
    796 
    797     void BGR_to_XYZ(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
    798     {
    799         using namespace cv::cuda::device;
    800         static const gpu_func_t funcs[2][2][6] =
    801         {
    802             {
    803                 {BGR_to_XYZ_8u, 0, BGR_to_XYZ_16u, 0, 0, BGR_to_XYZ_32f},
    804                 {BGRA_to_XYZ_8u, 0, BGRA_to_XYZ_16u, 0, 0, BGRA_to_XYZ_32f}
    805             },
    806             {
    807                 {BGR_to_XYZ4_8u, 0, BGR_to_XYZ4_16u, 0, 0, BGR_to_XYZ4_32f},
    808                 {BGRA_to_XYZ4_8u, 0, BGRA_to_XYZ4_16u, 0, 0, BGRA_to_XYZ4_32f}
    809             }
    810         };
    811 
    812         if (dcn <= 0) dcn = 3;
    813 
    814         GpuMat src = _src.getGpuMat();
    815 
    816         CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F );
    817         CV_Assert( src.channels() == 3 || src.channels() == 4 );
    818         CV_Assert( dcn == 3 || dcn == 4 );
    819 
    820         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
    821         GpuMat dst = _dst.getGpuMat();
    822 
    823         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
    824     }
    825 
    826     void XYZ_to_RGB(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
    827     {
    828         using namespace cv::cuda::device;
    829         static const gpu_func_t funcs[2][2][6] =
    830         {
    831             {
    832                 {XYZ_to_RGB_8u, 0, XYZ_to_RGB_16u, 0, 0, XYZ_to_RGB_32f},
    833                 {XYZ4_to_RGB_8u, 0, XYZ4_to_RGB_16u, 0, 0, XYZ4_to_RGB_32f}
    834             },
    835             {
    836                 {XYZ_to_RGBA_8u, 0, XYZ_to_RGBA_16u, 0, 0, XYZ_to_RGBA_32f},
    837                 {XYZ4_to_RGBA_8u, 0, XYZ4_to_RGBA_16u, 0, 0, XYZ4_to_RGBA_32f}
    838             }
    839         };
    840 
    841         if (dcn <= 0) dcn = 3;
    842 
    843         GpuMat src = _src.getGpuMat();
    844 
    845         CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F );
    846         CV_Assert( src.channels() == 3 || src.channels() == 4 );
    847         CV_Assert( dcn == 3 || dcn == 4 );
    848 
    849         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
    850         GpuMat dst = _dst.getGpuMat();
    851 
    852         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
    853     }
    854 
    855     void XYZ_to_BGR(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
    856     {
    857         using namespace cv::cuda::device;
    858         static const gpu_func_t funcs[2][2][6] =
    859         {
    860             {
    861                 {XYZ_to_BGR_8u, 0, XYZ_to_BGR_16u, 0, 0, XYZ_to_BGR_32f},
    862                 {XYZ4_to_BGR_8u, 0, XYZ4_to_BGR_16u, 0, 0, XYZ4_to_BGR_32f}
    863             },
    864             {
    865                 {XYZ_to_BGRA_8u, 0, XYZ_to_BGRA_16u, 0, 0, XYZ_to_BGRA_32f},
    866                 {XYZ4_to_BGRA_8u, 0, XYZ4_to_BGRA_16u, 0, 0, XYZ4_to_BGRA_32f}
    867             }
    868         };
    869 
    870         if (dcn <= 0) dcn = 3;
    871 
    872         GpuMat src = _src.getGpuMat();
    873 
    874         CV_Assert( src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F );
    875         CV_Assert( src.channels() == 3 || src.channels() == 4 );
    876         CV_Assert( dcn == 3 || dcn == 4 );
    877 
    878         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
    879         GpuMat dst = _dst.getGpuMat();
    880 
    881         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
    882     }
    883 
    884     void RGB_to_HSV(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
    885     {
    886         using namespace cv::cuda::device;
    887         static const gpu_func_t funcs[2][2][6] =
    888         {
    889             {
    890                 {RGB_to_HSV_8u, 0, 0, 0, 0, RGB_to_HSV_32f},
    891                 {RGBA_to_HSV_8u, 0, 0, 0, 0, RGBA_to_HSV_32f},
    892             },
    893             {
    894                 {RGB_to_HSV4_8u, 0, 0, 0, 0, RGB_to_HSV4_32f},
    895                 {RGBA_to_HSV4_8u, 0, 0, 0, 0, RGBA_to_HSV4_32f},
    896             }
    897         };
    898 
    899         if (dcn <= 0) dcn = 3;
    900 
    901         GpuMat src = _src.getGpuMat();
    902 
    903         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
    904         CV_Assert( src.channels() == 3 || src.channels() == 4 );
    905         CV_Assert( dcn == 3 || dcn == 4 );
    906 
    907         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
    908         GpuMat dst = _dst.getGpuMat();
    909 
    910         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
    911     }
    912 
    913     void BGR_to_HSV(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
    914     {
    915         using namespace cv::cuda::device;
    916         static const gpu_func_t funcs[2][2][6] =
    917         {
    918             {
    919                 {BGR_to_HSV_8u, 0, 0, 0, 0, BGR_to_HSV_32f},
    920                 {BGRA_to_HSV_8u, 0, 0, 0, 0, BGRA_to_HSV_32f}
    921             },
    922             {
    923                 {BGR_to_HSV4_8u, 0, 0, 0, 0, BGR_to_HSV4_32f},
    924                 {BGRA_to_HSV4_8u, 0, 0, 0, 0, BGRA_to_HSV4_32f}
    925             }
    926         };
    927 
    928         if (dcn <= 0) dcn = 3;
    929 
    930         GpuMat src = _src.getGpuMat();
    931 
    932         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
    933         CV_Assert( src.channels() == 3 || src.channels() == 4 );
    934         CV_Assert( dcn == 3 || dcn == 4 );
    935 
    936         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
    937         GpuMat dst = _dst.getGpuMat();
    938 
    939         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
    940     }
    941 
    942     void HSV_to_RGB(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
    943     {
    944         using namespace cv::cuda::device;
    945         static const gpu_func_t funcs[2][2][6] =
    946         {
    947             {
    948                 {HSV_to_RGB_8u, 0, 0, 0, 0, HSV_to_RGB_32f},
    949                 {HSV4_to_RGB_8u, 0, 0, 0, 0, HSV4_to_RGB_32f}
    950             },
    951             {
    952                 {HSV_to_RGBA_8u, 0, 0, 0, 0, HSV_to_RGBA_32f},
    953                 {HSV4_to_RGBA_8u, 0, 0, 0, 0, HSV4_to_RGBA_32f}
    954             }
    955         };
    956 
    957         if (dcn <= 0) dcn = 3;
    958 
    959         GpuMat src = _src.getGpuMat();
    960 
    961         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
    962         CV_Assert( src.channels() == 3 || src.channels() == 4 );
    963         CV_Assert( dcn == 3 || dcn == 4 );
    964 
    965         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
    966         GpuMat dst = _dst.getGpuMat();
    967 
    968         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
    969     }
    970 
    971     void HSV_to_BGR(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
    972     {
    973         using namespace cv::cuda::device;
    974         static const gpu_func_t funcs[2][2][6] =
    975         {
    976             {
    977                 {HSV_to_BGR_8u, 0, 0, 0, 0, HSV_to_BGR_32f},
    978                 {HSV4_to_BGR_8u, 0, 0, 0, 0, HSV4_to_BGR_32f}
    979             },
    980             {
    981                 {HSV_to_BGRA_8u, 0, 0, 0, 0, HSV_to_BGRA_32f},
    982                 {HSV4_to_BGRA_8u, 0, 0, 0, 0, HSV4_to_BGRA_32f}
    983             }
    984         };
    985 
    986         if (dcn <= 0) dcn = 3;
    987 
    988         GpuMat src = _src.getGpuMat();
    989 
    990         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
    991         CV_Assert( src.channels() == 3 || src.channels() == 4 );
    992         CV_Assert( dcn == 3 || dcn == 4 );
    993 
    994         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
    995         GpuMat dst = _dst.getGpuMat();
    996 
    997         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
    998     }
    999 
   1000     void RGB_to_HLS(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1001     {
   1002         using namespace cv::cuda::device;
   1003         static const gpu_func_t funcs[2][2][6] =
   1004         {
   1005             {
   1006                 {RGB_to_HLS_8u, 0, 0, 0, 0, RGB_to_HLS_32f},
   1007                 {RGBA_to_HLS_8u, 0, 0, 0, 0, RGBA_to_HLS_32f},
   1008             },
   1009             {
   1010                 {RGB_to_HLS4_8u, 0, 0, 0, 0, RGB_to_HLS4_32f},
   1011                 {RGBA_to_HLS4_8u, 0, 0, 0, 0, RGBA_to_HLS4_32f},
   1012             }
   1013         };
   1014 
   1015         if (dcn <= 0) dcn = 3;
   1016 
   1017         GpuMat src = _src.getGpuMat();
   1018 
   1019         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1020         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1021         CV_Assert( dcn == 3 || dcn == 4 );
   1022 
   1023         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1024         GpuMat dst = _dst.getGpuMat();
   1025 
   1026         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
   1027     }
   1028 
   1029     void BGR_to_HLS(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1030     {
   1031         using namespace cv::cuda::device;
   1032         static const gpu_func_t funcs[2][2][6] =
   1033         {
   1034             {
   1035                 {BGR_to_HLS_8u, 0, 0, 0, 0, BGR_to_HLS_32f},
   1036                 {BGRA_to_HLS_8u, 0, 0, 0, 0, BGRA_to_HLS_32f}
   1037             },
   1038             {
   1039                 {BGR_to_HLS4_8u, 0, 0, 0, 0, BGR_to_HLS4_32f},
   1040                 {BGRA_to_HLS4_8u, 0, 0, 0, 0, BGRA_to_HLS4_32f}
   1041             }
   1042         };
   1043 
   1044         if (dcn <= 0) dcn = 3;
   1045 
   1046         GpuMat src = _src.getGpuMat();
   1047 
   1048         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1049         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1050         CV_Assert( dcn == 3 || dcn == 4 );
   1051 
   1052         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1053         GpuMat dst = _dst.getGpuMat();
   1054 
   1055         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
   1056     }
   1057 
   1058     void HLS_to_RGB(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1059     {
   1060         using namespace cv::cuda::device;
   1061         static const gpu_func_t funcs[2][2][6] =
   1062         {
   1063             {
   1064                 {HLS_to_RGB_8u, 0, 0, 0, 0, HLS_to_RGB_32f},
   1065                 {HLS4_to_RGB_8u, 0, 0, 0, 0, HLS4_to_RGB_32f}
   1066             },
   1067             {
   1068                 {HLS_to_RGBA_8u, 0, 0, 0, 0, HLS_to_RGBA_32f},
   1069                 {HLS4_to_RGBA_8u, 0, 0, 0, 0, HLS4_to_RGBA_32f}
   1070             }
   1071         };
   1072 
   1073         if (dcn <= 0) dcn = 3;
   1074 
   1075         GpuMat src = _src.getGpuMat();
   1076 
   1077         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1078         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1079         CV_Assert( dcn == 3 || dcn == 4 );
   1080 
   1081         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1082         GpuMat dst = _dst.getGpuMat();
   1083 
   1084         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
   1085     }
   1086 
   1087     void HLS_to_BGR(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1088     {
   1089         using namespace cv::cuda::device;
   1090         static const gpu_func_t funcs[2][2][6] =
   1091         {
   1092             {
   1093                 {HLS_to_BGR_8u, 0, 0, 0, 0, HLS_to_BGR_32f},
   1094                 {HLS4_to_BGR_8u, 0, 0, 0, 0, HLS4_to_BGR_32f}
   1095             },
   1096             {
   1097                 {HLS_to_BGRA_8u, 0, 0, 0, 0, HLS_to_BGRA_32f},
   1098                 {HLS4_to_BGRA_8u, 0, 0, 0, 0, HLS4_to_BGRA_32f}
   1099             }
   1100         };
   1101 
   1102         if (dcn <= 0) dcn = 3;
   1103 
   1104         GpuMat src = _src.getGpuMat();
   1105 
   1106         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1107         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1108         CV_Assert( dcn == 3 || dcn == 4 );
   1109 
   1110         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1111         GpuMat dst = _dst.getGpuMat();
   1112 
   1113         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
   1114     }
   1115 
   1116     void RGB_to_HSV_FULL(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1117     {
   1118         using namespace cv::cuda::device;
   1119         static const gpu_func_t funcs[2][2][6] =
   1120         {
   1121             {
   1122                 {RGB_to_HSV_FULL_8u, 0, 0, 0, 0, RGB_to_HSV_FULL_32f},
   1123                 {RGBA_to_HSV_FULL_8u, 0, 0, 0, 0, RGBA_to_HSV_FULL_32f},
   1124             },
   1125             {
   1126                 {RGB_to_HSV4_FULL_8u, 0, 0, 0, 0, RGB_to_HSV4_FULL_32f},
   1127                 {RGBA_to_HSV4_FULL_8u, 0, 0, 0, 0, RGBA_to_HSV4_FULL_32f},
   1128             }
   1129         };
   1130 
   1131         if (dcn <= 0) dcn = 3;
   1132 
   1133         GpuMat src = _src.getGpuMat();
   1134 
   1135         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1136         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1137         CV_Assert( dcn == 3 || dcn == 4 );
   1138 
   1139         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1140         GpuMat dst = _dst.getGpuMat();
   1141 
   1142         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
   1143     }
   1144 
   1145     void BGR_to_HSV_FULL(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1146     {
   1147         using namespace cv::cuda::device;
   1148         static const gpu_func_t funcs[2][2][6] =
   1149         {
   1150             {
   1151                 {BGR_to_HSV_FULL_8u, 0, 0, 0, 0, BGR_to_HSV_FULL_32f},
   1152                 {BGRA_to_HSV_FULL_8u, 0, 0, 0, 0, BGRA_to_HSV_FULL_32f}
   1153             },
   1154             {
   1155                 {BGR_to_HSV4_FULL_8u, 0, 0, 0, 0, BGR_to_HSV4_FULL_32f},
   1156                 {BGRA_to_HSV4_FULL_8u, 0, 0, 0, 0, BGRA_to_HSV4_FULL_32f}
   1157             }
   1158         };
   1159 
   1160         if (dcn <= 0) dcn = 3;
   1161 
   1162         GpuMat src = _src.getGpuMat();
   1163 
   1164         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1165         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1166         CV_Assert( dcn == 3 || dcn == 4 );
   1167 
   1168         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1169         GpuMat dst = _dst.getGpuMat();
   1170 
   1171         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
   1172     }
   1173 
   1174     void HSV_to_RGB_FULL(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1175     {
   1176         using namespace cv::cuda::device;
   1177         static const gpu_func_t funcs[2][2][6] =
   1178         {
   1179             {
   1180                 {HSV_to_RGB_FULL_8u, 0, 0, 0, 0, HSV_to_RGB_FULL_32f},
   1181                 {HSV4_to_RGB_FULL_8u, 0, 0, 0, 0, HSV4_to_RGB_FULL_32f}
   1182             },
   1183             {
   1184                 {HSV_to_RGBA_FULL_8u, 0, 0, 0, 0, HSV_to_RGBA_FULL_32f},
   1185                 {HSV4_to_RGBA_FULL_8u, 0, 0, 0, 0, HSV4_to_RGBA_FULL_32f}
   1186             }
   1187         };
   1188 
   1189         if (dcn <= 0) dcn = 3;
   1190 
   1191         GpuMat src = _src.getGpuMat();
   1192 
   1193         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1194         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1195         CV_Assert( dcn == 3 || dcn == 4 );
   1196 
   1197         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1198         GpuMat dst = _dst.getGpuMat();
   1199 
   1200         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
   1201     }
   1202 
   1203     void HSV_to_BGR_FULL(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1204     {
   1205         using namespace cv::cuda::device;
   1206         static const gpu_func_t funcs[2][2][6] =
   1207         {
   1208             {
   1209                 {HSV_to_BGR_FULL_8u, 0, 0, 0, 0, HSV_to_BGR_FULL_32f},
   1210                 {HSV4_to_BGR_FULL_8u, 0, 0, 0, 0, HSV4_to_BGR_FULL_32f}
   1211             },
   1212             {
   1213                 {HSV_to_BGRA_FULL_8u, 0, 0, 0, 0, HSV_to_BGRA_FULL_32f},
   1214                 {HSV4_to_BGRA_FULL_8u, 0, 0, 0, 0, HSV4_to_BGRA_FULL_32f}
   1215             }
   1216         };
   1217 
   1218         if (dcn <= 0) dcn = 3;
   1219 
   1220         GpuMat src = _src.getGpuMat();
   1221 
   1222         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1223         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1224         CV_Assert( dcn == 3 || dcn == 4 );
   1225 
   1226         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1227         GpuMat dst = _dst.getGpuMat();
   1228 
   1229         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
   1230     }
   1231 
   1232     void RGB_to_HLS_FULL(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1233     {
   1234         using namespace cv::cuda::device;
   1235         static const gpu_func_t funcs[2][2][6] =
   1236         {
   1237             {
   1238                 {RGB_to_HLS_FULL_8u, 0, 0, 0, 0, RGB_to_HLS_FULL_32f},
   1239                 {RGBA_to_HLS_FULL_8u, 0, 0, 0, 0, RGBA_to_HLS_FULL_32f},
   1240             },
   1241             {
   1242                 {RGB_to_HLS4_FULL_8u, 0, 0, 0, 0, RGB_to_HLS4_FULL_32f},
   1243                 {RGBA_to_HLS4_FULL_8u, 0, 0, 0, 0, RGBA_to_HLS4_FULL_32f},
   1244             }
   1245         };
   1246 
   1247         if (dcn <= 0) dcn = 3;
   1248 
   1249         GpuMat src = _src.getGpuMat();
   1250 
   1251         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1252         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1253         CV_Assert( dcn == 3 || dcn == 4 );
   1254 
   1255         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1256         GpuMat dst = _dst.getGpuMat();
   1257 
   1258         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
   1259     }
   1260 
   1261     void BGR_to_HLS_FULL(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1262     {
   1263         using namespace cv::cuda::device;
   1264         static const gpu_func_t funcs[2][2][6] =
   1265         {
   1266             {
   1267                 {BGR_to_HLS_FULL_8u, 0, 0, 0, 0, BGR_to_HLS_FULL_32f},
   1268                 {BGRA_to_HLS_FULL_8u, 0, 0, 0, 0, BGRA_to_HLS_FULL_32f}
   1269             },
   1270             {
   1271                 {BGR_to_HLS4_FULL_8u, 0, 0, 0, 0, BGR_to_HLS4_FULL_32f},
   1272                 {BGRA_to_HLS4_FULL_8u, 0, 0, 0, 0, BGRA_to_HLS4_FULL_32f}
   1273             }
   1274         };
   1275 
   1276         if (dcn <= 0) dcn = 3;
   1277 
   1278         GpuMat src = _src.getGpuMat();
   1279 
   1280         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1281         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1282         CV_Assert( dcn == 3 || dcn == 4 );
   1283 
   1284         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1285         GpuMat dst = _dst.getGpuMat();
   1286 
   1287         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
   1288     }
   1289 
   1290     void HLS_to_RGB_FULL(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1291     {
   1292         using namespace cv::cuda::device;
   1293         static const gpu_func_t funcs[2][2][6] =
   1294         {
   1295             {
   1296                 {HLS_to_RGB_FULL_8u, 0, 0, 0, 0, HLS_to_RGB_FULL_32f},
   1297                 {HLS4_to_RGB_FULL_8u, 0, 0, 0, 0, HLS4_to_RGB_FULL_32f}
   1298             },
   1299             {
   1300                 {HLS_to_RGBA_FULL_8u, 0, 0, 0, 0, HLS_to_RGBA_FULL_32f},
   1301                 {HLS4_to_RGBA_FULL_8u, 0, 0, 0, 0, HLS4_to_RGBA_FULL_32f}
   1302             }
   1303         };
   1304 
   1305         if (dcn <= 0) dcn = 3;
   1306 
   1307         GpuMat src = _src.getGpuMat();
   1308 
   1309         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1310         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1311         CV_Assert( dcn == 3 || dcn == 4 );
   1312 
   1313         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1314         GpuMat dst = _dst.getGpuMat();
   1315 
   1316         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
   1317     }
   1318 
   1319     void HLS_to_BGR_FULL(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1320     {
   1321         using namespace cv::cuda::device;
   1322         static const gpu_func_t funcs[2][2][6] =
   1323         {
   1324             {
   1325                 {HLS_to_BGR_FULL_8u, 0, 0, 0, 0, HLS_to_BGR_FULL_32f},
   1326                 {HLS4_to_BGR_FULL_8u, 0, 0, 0, 0, HLS4_to_BGR_FULL_32f}
   1327             },
   1328             {
   1329                 {HLS_to_BGRA_FULL_8u, 0, 0, 0, 0, HLS_to_BGRA_FULL_32f},
   1330                 {HLS4_to_BGRA_FULL_8u, 0, 0, 0, 0, HLS4_to_BGRA_FULL_32f}
   1331             }
   1332         };
   1333 
   1334         if (dcn <= 0) dcn = 3;
   1335 
   1336         GpuMat src = _src.getGpuMat();
   1337 
   1338         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1339         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1340         CV_Assert( dcn == 3 || dcn == 4 );
   1341 
   1342         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1343         GpuMat dst = _dst.getGpuMat();
   1344 
   1345         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, stream);
   1346     }
   1347 
   1348     void BGR_to_Lab(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1349     {
   1350         using namespace cv::cuda::device;
   1351         static const gpu_func_t funcs[2][2][2] =
   1352         {
   1353             {
   1354                 {BGR_to_Lab_8u, BGR_to_Lab_32f},
   1355                 {BGRA_to_Lab_8u, BGRA_to_Lab_32f}
   1356             },
   1357             {
   1358                 {BGR_to_Lab4_8u, BGR_to_Lab4_32f},
   1359                 {BGRA_to_Lab4_8u, BGRA_to_Lab4_32f}
   1360             }
   1361         };
   1362 
   1363         if (dcn <= 0) dcn = 3;
   1364 
   1365         GpuMat src = _src.getGpuMat();
   1366 
   1367         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1368         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1369         CV_Assert( dcn == 3 || dcn == 4 );
   1370 
   1371         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1372         GpuMat dst = _dst.getGpuMat();
   1373 
   1374         funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream);
   1375     }
   1376 
   1377     void RGB_to_Lab(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1378     {
   1379         using namespace cv::cuda::device;
   1380         static const gpu_func_t funcs[2][2][2] =
   1381         {
   1382             {
   1383                 {RGB_to_Lab_8u, RGB_to_Lab_32f},
   1384                 {RGBA_to_Lab_8u, RGBA_to_Lab_32f}
   1385             },
   1386             {
   1387                 {RGB_to_Lab4_8u, RGB_to_Lab4_32f},
   1388                 {RGBA_to_Lab4_8u, RGBA_to_Lab4_32f}
   1389             }
   1390         };
   1391 
   1392         if (dcn <= 0) dcn = 3;
   1393 
   1394         GpuMat src = _src.getGpuMat();
   1395 
   1396         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1397         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1398         CV_Assert( dcn == 3 || dcn == 4 );
   1399 
   1400         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1401         GpuMat dst = _dst.getGpuMat();
   1402 
   1403         funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream);
   1404     }
   1405 
   1406     void LBGR_to_Lab(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1407     {
   1408         using namespace cv::cuda::device;
   1409         static const gpu_func_t funcs[2][2][2] =
   1410         {
   1411             {
   1412                 {LBGR_to_Lab_8u, LBGR_to_Lab_32f},
   1413                 {LBGRA_to_Lab_8u, LBGRA_to_Lab_32f}
   1414             },
   1415             {
   1416                 {LBGR_to_Lab4_8u, LBGR_to_Lab4_32f},
   1417                 {LBGRA_to_Lab4_8u, LBGRA_to_Lab4_32f}
   1418             }
   1419         };
   1420 
   1421         if (dcn <= 0) dcn = 3;
   1422 
   1423         GpuMat src = _src.getGpuMat();
   1424 
   1425         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1426         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1427         CV_Assert( dcn == 3 || dcn == 4 );
   1428 
   1429         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1430         GpuMat dst = _dst.getGpuMat();
   1431 
   1432         funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream);
   1433     }
   1434 
   1435     void LRGB_to_Lab(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1436     {
   1437         using namespace cv::cuda::device;
   1438         static const gpu_func_t funcs[2][2][2] =
   1439         {
   1440             {
   1441                 {LRGB_to_Lab_8u, LRGB_to_Lab_32f},
   1442                 {LRGBA_to_Lab_8u, LRGBA_to_Lab_32f}
   1443             },
   1444             {
   1445                 {LRGB_to_Lab4_8u, LRGB_to_Lab4_32f},
   1446                 {LRGBA_to_Lab4_8u, LRGBA_to_Lab4_32f}
   1447             }
   1448         };
   1449 
   1450         if (dcn <= 0) dcn = 3;
   1451 
   1452         GpuMat src = _src.getGpuMat();
   1453 
   1454         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1455         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1456         CV_Assert( dcn == 3 || dcn == 4 );
   1457 
   1458         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1459         GpuMat dst = _dst.getGpuMat();
   1460 
   1461         funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream);
   1462     }
   1463 
   1464     void Lab_to_BGR(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1465     {
   1466         using namespace cv::cuda::device;
   1467         static const gpu_func_t funcs[2][2][2] =
   1468         {
   1469             {
   1470                 {Lab_to_BGR_8u, Lab_to_BGR_32f},
   1471                 {Lab4_to_BGR_8u, Lab4_to_BGR_32f}
   1472             },
   1473             {
   1474                 {Lab_to_BGRA_8u, Lab_to_BGRA_32f},
   1475                 {Lab4_to_BGRA_8u, Lab4_to_BGRA_32f}
   1476             }
   1477         };
   1478 
   1479         if (dcn <= 0) dcn = 3;
   1480 
   1481         GpuMat src = _src.getGpuMat();
   1482 
   1483         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1484         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1485         CV_Assert( dcn == 3 || dcn == 4 );
   1486 
   1487         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1488         GpuMat dst = _dst.getGpuMat();
   1489 
   1490         funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream);
   1491     }
   1492 
   1493     void Lab_to_RGB(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1494     {
   1495         using namespace cv::cuda::device;
   1496         static const gpu_func_t funcs[2][2][2] =
   1497         {
   1498             {
   1499                 {Lab_to_RGB_8u, Lab_to_RGB_32f},
   1500                 {Lab4_to_RGB_8u, Lab4_to_RGB_32f}
   1501             },
   1502             {
   1503                 {Lab_to_RGBA_8u, Lab_to_RGBA_32f},
   1504                 {Lab4_to_RGBA_8u, Lab4_to_RGBA_32f}
   1505             }
   1506         };
   1507 
   1508         if (dcn <= 0) dcn = 3;
   1509 
   1510         GpuMat src = _src.getGpuMat();
   1511 
   1512         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1513         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1514         CV_Assert( dcn == 3 || dcn == 4 );
   1515 
   1516         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1517         GpuMat dst = _dst.getGpuMat();
   1518 
   1519         funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream);
   1520     }
   1521 
   1522     void Lab_to_LBGR(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1523     {
   1524         using namespace cv::cuda::device;
   1525         static const gpu_func_t funcs[2][2][2] =
   1526         {
   1527             {
   1528                 {Lab_to_LBGR_8u, Lab_to_LBGR_32f},
   1529                 {Lab4_to_LBGR_8u, Lab4_to_LBGR_32f}
   1530             },
   1531             {
   1532                 {Lab_to_LBGRA_8u, Lab_to_LBGRA_32f},
   1533                 {Lab4_to_LBGRA_8u, Lab4_to_LBGRA_32f}
   1534             }
   1535         };
   1536 
   1537         if (dcn <= 0) dcn = 3;
   1538 
   1539         GpuMat src = _src.getGpuMat();
   1540 
   1541         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1542         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1543         CV_Assert( dcn == 3 || dcn == 4 );
   1544 
   1545         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1546         GpuMat dst = _dst.getGpuMat();
   1547 
   1548         funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream);
   1549     }
   1550 
   1551     void Lab_to_LRGB(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1552     {
   1553         using namespace cv::cuda::device;
   1554         static const gpu_func_t funcs[2][2][2] =
   1555         {
   1556             {
   1557                 {Lab_to_LRGB_8u, Lab_to_LRGB_32f},
   1558                 {Lab4_to_LRGB_8u, Lab4_to_LRGB_32f}
   1559             },
   1560             {
   1561                 {Lab_to_LRGBA_8u, Lab_to_LRGBA_32f},
   1562                 {Lab4_to_LRGBA_8u, Lab4_to_LRGBA_32f}
   1563             }
   1564         };
   1565 
   1566         if (dcn <= 0) dcn = 3;
   1567 
   1568         GpuMat src = _src.getGpuMat();
   1569 
   1570         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1571         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1572         CV_Assert( dcn == 3 || dcn == 4 );
   1573 
   1574         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1575         GpuMat dst = _dst.getGpuMat();
   1576 
   1577         funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream);
   1578     }
   1579 
   1580     void BGR_to_Luv(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1581     {
   1582         using namespace cv::cuda::device;
   1583         static const gpu_func_t funcs[2][2][2] =
   1584         {
   1585             {
   1586                 {BGR_to_Luv_8u, BGR_to_Luv_32f},
   1587                 {BGRA_to_Luv_8u, BGRA_to_Luv_32f}
   1588             },
   1589             {
   1590                 {BGR_to_Luv4_8u, BGR_to_Luv4_32f},
   1591                 {BGRA_to_Luv4_8u, BGRA_to_Luv4_32f}
   1592             }
   1593         };
   1594 
   1595         if (dcn <= 0) dcn = 3;
   1596 
   1597         GpuMat src = _src.getGpuMat();
   1598 
   1599         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1600         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1601         CV_Assert( dcn == 3 || dcn == 4 );
   1602 
   1603         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1604         GpuMat dst = _dst.getGpuMat();
   1605 
   1606         funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream);
   1607     }
   1608 
   1609     void RGB_to_Luv(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1610     {
   1611         using namespace cv::cuda::device;
   1612         static const gpu_func_t funcs[2][2][2] =
   1613         {
   1614             {
   1615                 {RGB_to_Luv_8u, RGB_to_Luv_32f},
   1616                 {RGBA_to_Luv_8u, RGBA_to_Luv_32f}
   1617             },
   1618             {
   1619                 {RGB_to_Luv4_8u, RGB_to_Luv4_32f},
   1620                 {RGBA_to_Luv4_8u, RGBA_to_Luv4_32f}
   1621             }
   1622         };
   1623 
   1624         if (dcn <= 0) dcn = 3;
   1625 
   1626         GpuMat src = _src.getGpuMat();
   1627 
   1628         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1629         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1630         CV_Assert( dcn == 3 || dcn == 4 );
   1631 
   1632         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1633         GpuMat dst = _dst.getGpuMat();
   1634 
   1635         funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream);
   1636     }
   1637 
   1638     void LBGR_to_Luv(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1639     {
   1640         using namespace cv::cuda::device;
   1641         static const gpu_func_t funcs[2][2][2] =
   1642         {
   1643             {
   1644                 {LBGR_to_Luv_8u, LBGR_to_Luv_32f},
   1645                 {LBGRA_to_Luv_8u, LBGRA_to_Luv_32f}
   1646             },
   1647             {
   1648                 {LBGR_to_Luv4_8u, LBGR_to_Luv4_32f},
   1649                 {LBGRA_to_Luv4_8u, LBGRA_to_Luv4_32f}
   1650             }
   1651         };
   1652 
   1653         if (dcn <= 0) dcn = 3;
   1654 
   1655         GpuMat src = _src.getGpuMat();
   1656 
   1657         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1658         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1659         CV_Assert( dcn == 3 || dcn == 4 );
   1660 
   1661         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1662         GpuMat dst = _dst.getGpuMat();
   1663 
   1664         funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream);
   1665     }
   1666 
   1667     void LRGB_to_Luv(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1668     {
   1669         using namespace cv::cuda::device;
   1670         static const gpu_func_t funcs[2][2][2] =
   1671         {
   1672             {
   1673                 {LRGB_to_Luv_8u, LRGB_to_Luv_32f},
   1674                 {LRGBA_to_Luv_8u, LRGBA_to_Luv_32f}
   1675             },
   1676             {
   1677                 {LRGB_to_Luv4_8u, LRGB_to_Luv4_32f},
   1678                 {LRGBA_to_Luv4_8u, LRGBA_to_Luv4_32f}
   1679             }
   1680         };
   1681 
   1682         if (dcn <= 0) dcn = 3;
   1683 
   1684         GpuMat src = _src.getGpuMat();
   1685 
   1686         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1687         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1688         CV_Assert( dcn == 3 || dcn == 4 );
   1689 
   1690         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1691         GpuMat dst = _dst.getGpuMat();
   1692 
   1693         funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream);
   1694     }
   1695 
   1696     void Luv_to_BGR(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1697     {
   1698         using namespace cv::cuda::device;
   1699         static const gpu_func_t funcs[2][2][2] =
   1700         {
   1701             {
   1702                 {Luv_to_BGR_8u, Luv_to_BGR_32f},
   1703                 {Luv4_to_BGR_8u, Luv4_to_BGR_32f}
   1704             },
   1705             {
   1706                 {Luv_to_BGRA_8u, Luv_to_BGRA_32f},
   1707                 {Luv4_to_BGRA_8u, Luv4_to_BGRA_32f}
   1708             }
   1709         };
   1710 
   1711         if (dcn <= 0) dcn = 3;
   1712 
   1713         GpuMat src = _src.getGpuMat();
   1714 
   1715         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1716         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1717         CV_Assert( dcn == 3 || dcn == 4 );
   1718 
   1719         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1720         GpuMat dst = _dst.getGpuMat();
   1721 
   1722         funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream);
   1723     }
   1724 
   1725     void Luv_to_RGB(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1726     {
   1727         using namespace cv::cuda::device;
   1728         static const gpu_func_t funcs[2][2][2] =
   1729         {
   1730             {
   1731                 {Luv_to_RGB_8u, Luv_to_RGB_32f},
   1732                 {Luv4_to_RGB_8u, Luv4_to_RGB_32f}
   1733             },
   1734             {
   1735                 {Luv_to_RGBA_8u, Luv_to_RGBA_32f},
   1736                 {Luv4_to_RGBA_8u, Luv4_to_RGBA_32f}
   1737             }
   1738         };
   1739 
   1740         if (dcn <= 0) dcn = 3;
   1741 
   1742         GpuMat src = _src.getGpuMat();
   1743 
   1744         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1745         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1746         CV_Assert( dcn == 3 || dcn == 4 );
   1747 
   1748         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1749         GpuMat dst = _dst.getGpuMat();
   1750 
   1751         funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream);
   1752     }
   1753 
   1754     void Luv_to_LBGR(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1755     {
   1756         using namespace cv::cuda::device;
   1757         static const gpu_func_t funcs[2][2][2] =
   1758         {
   1759             {
   1760                 {Luv_to_LBGR_8u, Luv_to_LBGR_32f},
   1761                 {Luv4_to_LBGR_8u, Luv4_to_LBGR_32f}
   1762             },
   1763             {
   1764                 {Luv_to_LBGRA_8u, Luv_to_LBGRA_32f},
   1765                 {Luv4_to_LBGRA_8u, Luv4_to_LBGRA_32f}
   1766             }
   1767         };
   1768 
   1769         if (dcn <= 0) dcn = 3;
   1770 
   1771         GpuMat src = _src.getGpuMat();
   1772 
   1773         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1774         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1775         CV_Assert( dcn == 3 || dcn == 4 );
   1776 
   1777         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1778         GpuMat dst = _dst.getGpuMat();
   1779 
   1780         funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream);
   1781     }
   1782 
   1783     void Luv_to_LRGB(InputArray _src, OutputArray _dst, int dcn, Stream& stream)
   1784     {
   1785         using namespace cv::cuda::device;
   1786         static const gpu_func_t funcs[2][2][2] =
   1787         {
   1788             {
   1789                 {Luv_to_LRGB_8u, Luv_to_LRGB_32f},
   1790                 {Luv4_to_LRGB_8u, Luv4_to_LRGB_32f}
   1791             },
   1792             {
   1793                 {Luv_to_LRGBA_8u, Luv_to_LRGBA_32f},
   1794                 {Luv4_to_LRGBA_8u, Luv4_to_LRGBA_32f}
   1795             }
   1796         };
   1797 
   1798         if (dcn <= 0) dcn = 3;
   1799 
   1800         GpuMat src = _src.getGpuMat();
   1801 
   1802         CV_Assert( src.depth() == CV_8U || src.depth() == CV_32F );
   1803         CV_Assert( src.channels() == 3 || src.channels() == 4 );
   1804         CV_Assert( dcn == 3 || dcn == 4 );
   1805 
   1806         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1807         GpuMat dst = _dst.getGpuMat();
   1808 
   1809         funcs[dcn == 4][src.channels() == 4][src.depth() == CV_32F](src, dst, stream);
   1810     }
   1811 
   1812     void RGBA_to_mBGRA(InputArray _src, OutputArray _dst, int, Stream& _stream)
   1813     {
   1814     #if (CUDA_VERSION < 5000)
   1815         (void) _src;
   1816         (void) _dst;
   1817         (void) _stream;
   1818         CV_Error( Error::StsBadFlag, "Unknown/unsupported color conversion code" );
   1819     #else
   1820         GpuMat src = _src.getGpuMat();
   1821 
   1822         CV_Assert( src.type() == CV_8UC4 || src.type() == CV_16UC4 );
   1823 
   1824         _dst.create(src.size(), src.type());
   1825         GpuMat dst = _dst.getGpuMat();
   1826 
   1827         cudaStream_t stream = StreamAccessor::getStream(_stream);
   1828         NppStreamHandler h(stream);
   1829 
   1830         NppiSize oSizeROI;
   1831         oSizeROI.width = src.cols;
   1832         oSizeROI.height = src.rows;
   1833 
   1834         if (src.depth() == CV_8U)
   1835             nppSafeCall( nppiAlphaPremul_8u_AC4R(src.ptr<Npp8u>(), static_cast<int>(src.step), dst.ptr<Npp8u>(), static_cast<int>(dst.step), oSizeROI) );
   1836         else
   1837             nppSafeCall( nppiAlphaPremul_16u_AC4R(src.ptr<Npp16u>(), static_cast<int>(src.step), dst.ptr<Npp16u>(), static_cast<int>(dst.step), oSizeROI) );
   1838 
   1839         if (stream == 0)
   1840             cudaSafeCall( cudaDeviceSynchronize() );
   1841     #endif
   1842     }
   1843 
   1844     void bayer_to_BGR(InputArray _src, OutputArray _dst, int dcn, bool blue_last, bool start_with_green, Stream& stream)
   1845     {
   1846         typedef void (*func_t)(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
   1847         static const func_t funcs[3][4] =
   1848         {
   1849             {0,0,Bayer2BGR_8u_gpu<3>, Bayer2BGR_8u_gpu<4>},
   1850             {0,0,0,0},
   1851             {0,0,Bayer2BGR_16u_gpu<3>, Bayer2BGR_16u_gpu<4>}
   1852         };
   1853 
   1854         if (dcn <= 0) dcn = 3;
   1855 
   1856         GpuMat src = _src.getGpuMat();
   1857 
   1858         CV_Assert( src.type() == CV_8UC1 || src.type() == CV_16UC1 );
   1859         CV_Assert( src.rows > 2 && src.cols > 2 );
   1860         CV_Assert( dcn == 3 || dcn == 4 );
   1861 
   1862         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), dcn));
   1863         GpuMat dst = _dst.getGpuMat();
   1864 
   1865         funcs[src.depth()][dcn - 1](src, dst, blue_last, start_with_green, StreamAccessor::getStream(stream));
   1866     }
   1867     void bayerBG_to_BGR(InputArray src, OutputArray dst, int dcn, Stream& stream)
   1868     {
   1869         bayer_to_BGR(src, dst, dcn, false, false, stream);
   1870     }
   1871     void bayeRGB_to_BGR(InputArray src, OutputArray dst, int dcn, Stream& stream)
   1872     {
   1873         bayer_to_BGR(src, dst, dcn, false, true, stream);
   1874     }
   1875     void bayerRG_to_BGR(InputArray src, OutputArray dst, int dcn, Stream& stream)
   1876     {
   1877         bayer_to_BGR(src, dst, dcn, true, false, stream);
   1878     }
   1879     void bayerGR_to_BGR(InputArray src, OutputArray dst, int dcn, Stream& stream)
   1880     {
   1881         bayer_to_BGR(src, dst, dcn, true, true, stream);
   1882     }
   1883 
   1884     void bayer_to_gray(InputArray _src, OutputArray _dst, bool blue_last, bool start_with_green, Stream& stream)
   1885     {
   1886         typedef void (*func_t)(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
   1887         static const func_t funcs[3] =
   1888         {
   1889             Bayer2BGR_8u_gpu<1>,
   1890             0,
   1891             Bayer2BGR_16u_gpu<1>,
   1892         };
   1893 
   1894         GpuMat src = _src.getGpuMat();
   1895 
   1896         CV_Assert( src.type() == CV_8UC1 || src.type() == CV_16UC1 );
   1897         CV_Assert( src.rows > 2 && src.cols > 2 );
   1898 
   1899         _dst.create(src.size(), CV_MAKE_TYPE(src.depth(), 1));
   1900         GpuMat dst = _dst.getGpuMat();
   1901 
   1902         funcs[src.depth()](src, dst, blue_last, start_with_green, StreamAccessor::getStream(stream));
   1903     }
   1904     void bayerBG_to_gray(InputArray src, OutputArray dst, int /*dcn*/, Stream& stream)
   1905     {
   1906         bayer_to_gray(src, dst, false, false, stream);
   1907     }
   1908     void bayeRGB_to_GRAY(InputArray src, OutputArray dst, int /*dcn*/, Stream& stream)
   1909     {
   1910         bayer_to_gray(src, dst, false, true, stream);
   1911     }
   1912     void bayerRG_to_gray(InputArray src, OutputArray dst, int /*dcn*/, Stream& stream)
   1913     {
   1914         bayer_to_gray(src, dst, true, false, stream);
   1915     }
   1916     void bayerGR_to_gray(InputArray src, OutputArray dst, int /*dcn*/, Stream& stream)
   1917     {
   1918         bayer_to_gray(src, dst, true, true, stream);
   1919     }
   1920 }
   1921 
   1922 ////////////////////////////////////////////////////////////////////////
   1923 // cvtColor
   1924 
   1925 void cv::cuda::cvtColor(InputArray src, OutputArray dst, int code, int dcn, Stream& stream)
   1926 {
   1927     typedef void (*func_t)(InputArray src, OutputArray dst, int dcn, Stream& stream);
   1928     static const func_t funcs[] =
   1929     {
   1930         BGR_to_BGRA,            // CV_BGR2BGRA    =0
   1931         BGRA_to_BGR,            // CV_BGRA2BGR    =1
   1932         BGR_to_RGBA,            // CV_BGR2RGBA    =2
   1933         BGRA_to_RGB,            // CV_RGBA2BGR    =3
   1934         BGR_to_RGB,             // CV_BGR2RGB     =4
   1935         BGRA_to_RGBA,           // CV_BGRA2RGBA   =5
   1936 
   1937         BGR_to_GRAY,            // CV_BGR2GRAY    =6
   1938         RGB_to_GRAY,            // CV_RGB2GRAY    =7
   1939         GRAY_to_BGR,            // CV_GRAY2BGR    =8
   1940         GRAY_to_BGRA,           // CV_GRAY2BGRA   =9
   1941         BGRA_to_GRAY,           // CV_BGRA2GRAY   =10
   1942         RGBA_to_GRAY,           // CV_RGBA2GRAY   =11
   1943 
   1944         BGR_to_BGR565,          // CV_BGR2BGR565  =12
   1945         RGB_to_BGR565,          // CV_RGB2BGR565  =13
   1946         BGR565_to_BGR,          // CV_BGR5652BGR  =14
   1947         BGR565_to_RGB,          // CV_BGR5652RGB  =15
   1948         BGRA_to_BGR565,         // CV_BGRA2BGR565 =16
   1949         RGBA_to_BGR565,         // CV_RGBA2BGR565 =17
   1950         BGR565_to_BGRA,         // CV_BGR5652BGRA =18
   1951         BGR565_to_RGBA,         // CV_BGR5652RGBA =19
   1952 
   1953         GRAY_to_BGR565,         // CV_GRAY2BGR565 =20
   1954         BGR565_to_GRAY,         // CV_BGR5652GRAY =21
   1955 
   1956         BGR_to_BGR555,          // CV_BGR2BGR555  =22
   1957         RGB_to_BGR555,          // CV_RGB2BGR555  =23
   1958         BGR555_to_BGR,          // CV_BGR5552BGR  =24
   1959         BGR555_to_RGB,          // CV_BGR5552RGB  =25
   1960         BGRA_to_BGR555,         // CV_BGRA2BGR555 =26
   1961         RGBA_to_BGR555,         // CV_RGBA2BGR555 =27
   1962         BGR555_to_BGRA,         // CV_BGR5552BGRA =28
   1963         BGR555_to_RGBA,         // CV_BGR5552RGBA =29
   1964 
   1965         GRAY_to_BGR555,         // CV_GRAY2BGR555 =30
   1966         BGR555_to_GRAY,         // CV_BGR5552GRAY =31
   1967 
   1968         BGR_to_XYZ,             // CV_BGR2XYZ     =32
   1969         RGB_to_XYZ,             // CV_RGB2XYZ     =33
   1970         XYZ_to_BGR,             // CV_XYZ2BGR     =34
   1971         XYZ_to_RGB,             // CV_XYZ2RGB     =35
   1972 
   1973         BGR_to_YCrCb,           // CV_BGR2YCrCb   =36
   1974         RGB_to_YCrCb,           // CV_RGB2YCrCb   =37
   1975         YCrCb_to_BGR,           // CV_YCrCb2BGR   =38
   1976         YCrCb_to_RGB,           // CV_YCrCb2RGB   =39
   1977 
   1978         BGR_to_HSV,             // CV_BGR2HSV     =40
   1979         RGB_to_HSV,             // CV_RGB2HSV     =41
   1980 
   1981         0,                      //                =42
   1982         0,                      //                =43
   1983 
   1984         BGR_to_Lab,             // CV_BGR2Lab     =44
   1985         RGB_to_Lab,             // CV_RGB2Lab     =45
   1986 
   1987         bayerBG_to_BGR,         // CV_BayerBG2BGR =46
   1988         bayeRGB_to_BGR,         // CV_BayeRGB2BGR =47
   1989         bayerRG_to_BGR,         // CV_BayerRG2BGR =48
   1990         bayerGR_to_BGR,         // CV_BayerGR2BGR =49
   1991 
   1992         BGR_to_Luv,             // CV_BGR2Luv     =50
   1993         RGB_to_Luv,             // CV_RGB2Luv     =51
   1994 
   1995         BGR_to_HLS,             // CV_BGR2HLS     =52
   1996         RGB_to_HLS,             // CV_RGB2HLS     =53
   1997 
   1998         HSV_to_BGR,             // CV_HSV2BGR     =54
   1999         HSV_to_RGB,             // CV_HSV2RGB     =55
   2000 
   2001         Lab_to_BGR,             // CV_Lab2BGR     =56
   2002         Lab_to_RGB,             // CV_Lab2RGB     =57
   2003         Luv_to_BGR,             // CV_Luv2BGR     =58
   2004         Luv_to_RGB,             // CV_Luv2RGB     =59
   2005 
   2006         HLS_to_BGR,             // CV_HLS2BGR     =60
   2007         HLS_to_RGB,             // CV_HLS2RGB     =61
   2008 
   2009         0,                      // CV_BayerBG2BGR_VNG =62
   2010         0,                      // CV_BayeRGB2BGR_VNG =63
   2011         0,                      // CV_BayerRG2BGR_VNG =64
   2012         0,                      // CV_BayerGR2BGR_VNG =65
   2013 
   2014         BGR_to_HSV_FULL,        // CV_BGR2HSV_FULL = 66
   2015         RGB_to_HSV_FULL,        // CV_RGB2HSV_FULL = 67
   2016         BGR_to_HLS_FULL,        // CV_BGR2HLS_FULL = 68
   2017         RGB_to_HLS_FULL,        // CV_RGB2HLS_FULL = 69
   2018 
   2019         HSV_to_BGR_FULL,        // CV_HSV2BGR_FULL = 70
   2020         HSV_to_RGB_FULL,        // CV_HSV2RGB_FULL = 71
   2021         HLS_to_BGR_FULL,        // CV_HLS2BGR_FULL = 72
   2022         HLS_to_RGB_FULL,        // CV_HLS2RGB_FULL = 73
   2023 
   2024         LBGR_to_Lab,            // CV_LBGR2Lab     = 74
   2025         LRGB_to_Lab,            // CV_LRGB2Lab     = 75
   2026         LBGR_to_Luv,            // CV_LBGR2Luv     = 76
   2027         LRGB_to_Luv,            // CV_LRGB2Luv     = 77
   2028 
   2029         Lab_to_LBGR,            // CV_Lab2LBGR     = 78
   2030         Lab_to_LRGB,            // CV_Lab2LRGB     = 79
   2031         Luv_to_LBGR,            // CV_Luv2LBGR     = 80
   2032         Luv_to_LRGB,            // CV_Luv2LRGB     = 81
   2033 
   2034         BGR_to_YUV,             // CV_BGR2YUV      = 82
   2035         RGB_to_YUV,             // CV_RGB2YUV      = 83
   2036         YUV_to_BGR,             // CV_YUV2BGR      = 84
   2037         YUV_to_RGB,             // CV_YUV2RGB      = 85
   2038 
   2039         bayerBG_to_gray,        // CV_BayerBG2GRAY = 86
   2040         bayeRGB_to_GRAY,        // CV_BayeRGB2GRAY = 87
   2041         bayerRG_to_gray,        // CV_BayerRG2GRAY = 88
   2042         bayerGR_to_gray,        // CV_BayerGR2GRAY = 89
   2043 
   2044         //YUV 4:2:0 formats family
   2045         0,                      // CV_YUV2RGB_NV12 = 90,
   2046         0,                      // CV_YUV2BGR_NV12 = 91,
   2047         0,                      // CV_YUV2RGB_NV21 = 92,
   2048         0,                      // CV_YUV2BGR_NV21 = 93,
   2049 
   2050         0,                      // CV_YUV2RGBA_NV12 = 94,
   2051         0,                      // CV_YUV2BGRA_NV12 = 95,
   2052         0,                      // CV_YUV2RGBA_NV21 = 96,
   2053         0,                      // CV_YUV2BGRA_NV21 = 97,
   2054 
   2055         0,                      // CV_YUV2RGB_YV12 = 98,
   2056         0,                      // CV_YUV2BGR_YV12 = 99,
   2057         0,                      // CV_YUV2RGB_IYUV = 100,
   2058         0,                      // CV_YUV2BGR_IYUV = 101,
   2059 
   2060         0,                      // CV_YUV2RGBA_YV12 = 102,
   2061         0,                      // CV_YUV2BGRA_YV12 = 103,
   2062         0,                      // CV_YUV2RGBA_IYUV = 104,
   2063         0,                      // CV_YUV2BGRA_IYUV = 105,
   2064 
   2065         0,                      // CV_YUV2GRAY_420 = 106,
   2066 
   2067         //YUV 4:2:2 formats family
   2068         0,                      // CV_YUV2RGB_UYVY = 107,
   2069         0,                      // CV_YUV2BGR_UYVY = 108,
   2070         0,                      // //CV_YUV2RGB_VYUY = 109,
   2071         0,                      // //CV_YUV2BGR_VYUY = 110,
   2072 
   2073         0,                      // CV_YUV2RGBA_UYVY = 111,
   2074         0,                      // CV_YUV2BGRA_UYVY = 112,
   2075         0,                      // //CV_YUV2RGBA_VYUY = 113,
   2076         0,                      // //CV_YUV2BGRA_VYUY = 114,
   2077 
   2078         0,                      // CV_YUV2RGB_YUY2 = 115,
   2079         0,                      // CV_YUV2BGR_YUY2 = 116,
   2080         0,                      // CV_YUV2RGB_YVYU = 117,
   2081         0,                      // CV_YUV2BGR_YVYU = 118,
   2082 
   2083         0,                      // CV_YUV2RGBA_YUY2 = 119,
   2084         0,                      // CV_YUV2BGRA_YUY2 = 120,
   2085         0,                      // CV_YUV2RGBA_YVYU = 121,
   2086         0,                      // CV_YUV2BGRA_YVYU = 122,
   2087 
   2088         0,                      // CV_YUV2GRAY_UYVY = 123,
   2089         0,                      // CV_YUV2GRAY_YUY2 = 124,
   2090 
   2091         // alpha premultiplication
   2092         RGBA_to_mBGRA,          // CV_RGBA2mRGBA = 125,
   2093         0,                      // CV_mRGBA2RGBA = 126,
   2094 
   2095         0,                      // CV_COLORCVT_MAX  = 127
   2096     };
   2097 
   2098     CV_Assert( code < 128 );
   2099 
   2100     func_t func = funcs[code];
   2101 
   2102     if (func == 0)
   2103         CV_Error(Error::StsBadFlag, "Unknown/unsupported color conversion code");
   2104 
   2105     func(src, dst, dcn, stream);
   2106 }
   2107 
   2108 ////////////////////////////////////////////////////////////////////////
   2109 // demosaicing
   2110 
   2111 void cv::cuda::demosaicing(InputArray _src, OutputArray _dst, int code, int dcn, Stream& stream)
   2112 {
   2113     CV_Assert( !_src.empty() );
   2114 
   2115     switch (code)
   2116     {
   2117     case cv::COLOR_BayerBG2GRAY: case cv::COLOR_BayerGB2GRAY: case cv::COLOR_BayerRG2GRAY: case cv::COLOR_BayerGR2GRAY:
   2118         bayer_to_gray(_src, _dst, code == cv::COLOR_BayerBG2GRAY || code == cv::COLOR_BayerGB2GRAY, code == cv::COLOR_BayerGB2GRAY || code == cv::COLOR_BayerGR2GRAY, stream);
   2119         break;
   2120 
   2121     case cv::COLOR_BayerBG2BGR: case cv::COLOR_BayerGB2BGR: case cv::COLOR_BayerRG2BGR: case cv::COLOR_BayerGR2BGR:
   2122         bayer_to_BGR(_src, _dst, dcn, code == cv::COLOR_BayerBG2BGR || code == cv::COLOR_BayerGB2BGR, code == cv::COLOR_BayerGB2BGR || code == cv::COLOR_BayerGR2BGR, stream);
   2123         break;
   2124 
   2125     case COLOR_BayerBG2BGR_MHT: case COLOR_BayerGB2BGR_MHT: case COLOR_BayerRG2BGR_MHT: case COLOR_BayerGR2BGR_MHT:
   2126     {
   2127         if (dcn <= 0) dcn = 3;
   2128 
   2129         GpuMat src = _src.getGpuMat();
   2130         const int depth = _src.depth();
   2131 
   2132         CV_Assert( depth == CV_8U );
   2133         CV_Assert( src.channels() == 1 );
   2134         CV_Assert( dcn == 3 || dcn == 4 );
   2135 
   2136         _dst.create(_src.size(), CV_MAKE_TYPE(depth, dcn));
   2137         GpuMat dst = _dst.getGpuMat();
   2138 
   2139         dst.setTo(Scalar::all(0), stream);
   2140 
   2141         Size wholeSize;
   2142         Point ofs;
   2143         src.locateROI(wholeSize, ofs);
   2144         PtrStepSzb srcWhole(wholeSize.height, wholeSize.width, src.datastart, src.step);
   2145 
   2146         const int2 firstRed = make_int2(code == COLOR_BayerRG2BGR_MHT || code == COLOR_BayerGB2BGR_MHT ? 0 : 1,
   2147                                         code == COLOR_BayerRG2BGR_MHT || code == COLOR_BayerGR2BGR_MHT ? 0 : 1);
   2148 
   2149         if (dcn == 3)
   2150             cv::cuda::device::MHCdemosaic<3>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream));
   2151         else
   2152             cv::cuda::device::MHCdemosaic<4>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream));
   2153 
   2154         break;
   2155     }
   2156 
   2157     case COLOR_BayerBG2GRAY_MHT: case COLOR_BayerGB2GRAY_MHT: case COLOR_BayerRG2GRAY_MHT: case COLOR_BayerGR2GRAY_MHT:
   2158     {
   2159         GpuMat src = _src.getGpuMat();
   2160         const int depth = _src.depth();
   2161 
   2162         CV_Assert( depth == CV_8U );
   2163 
   2164         _dst.create(_src.size(), CV_MAKE_TYPE(depth, 1));
   2165         GpuMat dst = _dst.getGpuMat();
   2166 
   2167         dst.setTo(Scalar::all(0), stream);
   2168 
   2169         Size wholeSize;
   2170         Point ofs;
   2171         src.locateROI(wholeSize, ofs);
   2172         PtrStepSzb srcWhole(wholeSize.height, wholeSize.width, src.datastart, src.step);
   2173 
   2174         const int2 firstRed = make_int2(code == COLOR_BayerRG2BGR_MHT || code == COLOR_BayerGB2BGR_MHT ? 0 : 1,
   2175                                         code == COLOR_BayerRG2BGR_MHT || code == COLOR_BayerGR2BGR_MHT ? 0 : 1);
   2176 
   2177         cv::cuda::device::MHCdemosaic<1>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream));
   2178 
   2179         break;
   2180     }
   2181 
   2182     default:
   2183         CV_Error(Error::StsBadFlag, "Unknown / unsupported color conversion code");
   2184     }
   2185 }
   2186 
   2187 ////////////////////////////////////////////////////////////////////////
   2188 // swapChannels
   2189 
   2190 void cv::cuda::swapChannels(InputOutputArray _image, const int dstOrder[4], Stream& _stream)
   2191 {
   2192     GpuMat image = _image.getGpuMat();
   2193 
   2194     CV_Assert( image.type() == CV_8UC4 );
   2195 
   2196     cudaStream_t stream = StreamAccessor::getStream(_stream);
   2197     NppStreamHandler h(stream);
   2198 
   2199     NppiSize sz;
   2200     sz.width  = image.cols;
   2201     sz.height = image.rows;
   2202 
   2203     nppSafeCall( nppiSwapChannels_8u_C4IR(image.ptr<Npp8u>(), static_cast<int>(image.step), sz, dstOrder) );
   2204 
   2205     if (stream == 0)
   2206         cudaSafeCall( cudaDeviceSynchronize() );
   2207 }
   2208 
   2209 ////////////////////////////////////////////////////////////////////////
   2210 // gammaCorrection
   2211 
   2212 void cv::cuda::gammaCorrection(InputArray _src, OutputArray _dst, bool forward, Stream& stream)
   2213 {
   2214 #if (CUDA_VERSION < 5000)
   2215     (void) _src;
   2216     (void) _dst;
   2217     (void) forward;
   2218     (void) stream;
   2219     CV_Error(Error::StsNotImplemented, "This function works only with CUDA 5.0 or higher");
   2220 #else
   2221     typedef NppStatus (*func_t)(const Npp8u* pSrc, int nSrcStep, Npp8u* pDst, int nDstStep, NppiSize oSizeROI);
   2222     typedef NppStatus (*func_inplace_t)(Npp8u* pSrcDst, int nSrcDstStep, NppiSize oSizeROI);
   2223 
   2224     static const func_t funcs[2][5] =
   2225     {
   2226         {0, 0, 0, nppiGammaInv_8u_C3R, nppiGammaInv_8u_AC4R},
   2227         {0, 0, 0, nppiGammaFwd_8u_C3R, nppiGammaFwd_8u_AC4R}
   2228     };
   2229     static const func_inplace_t funcs_inplace[2][5] =
   2230     {
   2231         {0, 0, 0, nppiGammaInv_8u_C3IR, nppiGammaInv_8u_AC4IR},
   2232         {0, 0, 0, nppiGammaFwd_8u_C3IR, nppiGammaFwd_8u_AC4IR}
   2233     };
   2234 
   2235     GpuMat src = _src.getGpuMat();
   2236 
   2237     CV_Assert( src.type() == CV_8UC3 || src.type() == CV_8UC4 );
   2238 
   2239     _dst.create(src.size(), src.type());
   2240     GpuMat dst = _dst.getGpuMat();
   2241 
   2242     NppStreamHandler h(StreamAccessor::getStream(stream));
   2243 
   2244     NppiSize oSizeROI;
   2245     oSizeROI.width = src.cols;
   2246     oSizeROI.height = src.rows;
   2247 
   2248     if (dst.data == src.data)
   2249         funcs_inplace[forward][src.channels()](dst.ptr<Npp8u>(), static_cast<int>(src.step), oSizeROI);
   2250     else
   2251         funcs[forward][src.channels()](src.ptr<Npp8u>(), static_cast<int>(src.step), dst.ptr<Npp8u>(), static_cast<int>(dst.step), oSizeROI);
   2252 
   2253 #endif
   2254 }
   2255 
   2256 ////////////////////////////////////////////////////////////////////////
   2257 // alphaComp
   2258 
   2259 namespace
   2260 {
   2261     template <int DEPTH> struct NppAlphaCompFunc
   2262     {
   2263         typedef typename NPPTypeTraits<DEPTH>::npp_type npp_t;
   2264 
   2265         typedef NppStatus (*func_t)(const npp_t* pSrc1, int nSrc1Step, const npp_t* pSrc2, int nSrc2Step, npp_t* pDst, int nDstStep, NppiSize oSizeROI, NppiAlphaOp eAlphaOp);
   2266     };
   2267 
   2268     template <int DEPTH, typename NppAlphaCompFunc<DEPTH>::func_t func> struct NppAlphaComp
   2269     {
   2270         typedef typename NPPTypeTraits<DEPTH>::npp_type npp_t;
   2271 
   2272         static void call(const GpuMat& img1, const GpuMat& img2, GpuMat& dst, NppiAlphaOp eAlphaOp, cudaStream_t stream)
   2273         {
   2274             NppStreamHandler h(stream);
   2275 
   2276             NppiSize oSizeROI;
   2277             oSizeROI.width = img1.cols;
   2278             oSizeROI.height = img2.rows;
   2279 
   2280             nppSafeCall( func(img1.ptr<npp_t>(), static_cast<int>(img1.step), img2.ptr<npp_t>(), static_cast<int>(img2.step),
   2281                               dst.ptr<npp_t>(), static_cast<int>(dst.step), oSizeROI, eAlphaOp) );
   2282 
   2283             if (stream == 0)
   2284                 cudaSafeCall( cudaDeviceSynchronize() );
   2285         }
   2286     };
   2287 }
   2288 
   2289 void cv::cuda::alphaComp(InputArray _img1, InputArray _img2, OutputArray _dst, int alpha_op, Stream& stream)
   2290 {
   2291     static const NppiAlphaOp npp_alpha_ops[] = {
   2292         NPPI_OP_ALPHA_OVER,
   2293         NPPI_OP_ALPHA_IN,
   2294         NPPI_OP_ALPHA_OUT,
   2295         NPPI_OP_ALPHA_ATOP,
   2296         NPPI_OP_ALPHA_XOR,
   2297         NPPI_OP_ALPHA_PLUS,
   2298         NPPI_OP_ALPHA_OVER_PREMUL,
   2299         NPPI_OP_ALPHA_IN_PREMUL,
   2300         NPPI_OP_ALPHA_OUT_PREMUL,
   2301         NPPI_OP_ALPHA_ATOP_PREMUL,
   2302         NPPI_OP_ALPHA_XOR_PREMUL,
   2303         NPPI_OP_ALPHA_PLUS_PREMUL,
   2304         NPPI_OP_ALPHA_PREMUL
   2305     };
   2306 
   2307     typedef void (*func_t)(const GpuMat& img1, const GpuMat& img2, GpuMat& dst, NppiAlphaOp eAlphaOp, cudaStream_t stream);
   2308     static const func_t funcs[] =
   2309     {
   2310         NppAlphaComp<CV_8U, nppiAlphaComp_8u_AC4R>::call,
   2311         0,
   2312         NppAlphaComp<CV_16U, nppiAlphaComp_16u_AC4R>::call,
   2313         0,
   2314         NppAlphaComp<CV_32S, nppiAlphaComp_32s_AC4R>::call,
   2315         NppAlphaComp<CV_32F, nppiAlphaComp_32f_AC4R>::call
   2316     };
   2317 
   2318     GpuMat img1 = _img1.getGpuMat();
   2319     GpuMat img2 = _img2.getGpuMat();
   2320 
   2321     CV_Assert( img1.type() == CV_8UC4 || img1.type() == CV_16UC4 || img1.type() == CV_32SC4 || img1.type() == CV_32FC4 );
   2322     CV_Assert( img1.size() == img2.size() && img1.type() == img2.type() );
   2323 
   2324     _dst.create(img1.size(), img1.type());
   2325     GpuMat dst = _dst.getGpuMat();
   2326 
   2327     const func_t func = funcs[img1.depth()];
   2328 
   2329     func(img1, img2, dst, npp_alpha_ops[alpha_op], StreamAccessor::getStream(stream));
   2330 }
   2331 
   2332 #endif /* !defined (HAVE_CUDA) */
   2333